《CUDA编程》5.获得GPU加速的关键

news/2024/10/6 18:00:19 标签: c++, cuda

从本章起，将关注CDUA程序的性能，即执行速度

1 用CUDA事件计时

在前几章中，使用的是C++的<time.h>库进行程序运行计时，CUDA也提供了一种基于CUDA event的计时方式，用来给一段CUDA代码进行计时，这里只介绍基于cudaEvent_t的计时方式，下面是一代码框架：

#include <cuda_runtime.h>
#include "error_check.cuh"

cudaEvent_t start, stop;

CHECK(cudaEventCreate(&start));
CHECK(cudaEventCreate(&stop));
CHECK(cudaEventRecord(start));
cudaEventQuery(start);//不能使用CHECK，因为可能返回cudaErrorNotReady，但并不是代码报错

/*需要计时的代码块*/

CHECK(cudaEventRecord(stop));
CHECK(cudaEventSynchronize(stop));
float time;
CHECK(cudaEventElapsedTime(&time, start, stop));
printf("time: %f ms\n", time);

CHECK(cudaEventDestroy(start));
CHECK(cudaEventDestroy(stop));
  1. 首先定义两个cudaEvent_t的变量start和stop，并用cudaEventCreate()初始化。
  2. 在需要计时的代码块运行之前，把start传入cudaEventRecord()。
  3. 若是在TCC驱动模式的GPU中，cudaEventQuery(start);可以省略；若是处于WDDM驱动模式中，则必须保留。关于这两种模式，后面会讨论。
  4. 在需要计时的代码块运行之后，把stop传入cudaEventRecord()。
  5. 使用cudaEventSynchronize(stop)让主机等待事件stop被记录完毕
  6. 调用 cudaEventElapsedTime() 函数计算 start 和 stop 这两个事件之间的时间差（单位是 ms）并输出到屏幕。
  7. 调用 cudaEventDestroy() 函数销毁 start 和 stop 这两个 CUDA 事件。

1.1 举例说明

下面是一段利用cudaEvent_t进行计时的代码，在调用核函数add前后进行计时，注意，为了能够进行单精度和双精度的比较，需要定义一个宏变量，以便在编译时选择精度

#include <cuda.h>
#include <cuda_runtime.h>
#include <math.h>
#include <stdio.h>
#include "error_check.cuh"
#ifdef USE_DP
    typedef double real;
    const real EPSILON = 1.0e-15;
#else
    typedef float real; 
    const real EPSILON = 1.0e-6f;
#endif 


const real EPS = 1.0e-15;
const real a = 1.23;
const real b = 2.34;
const real c = 3.57;

// 希望 add 函数在 GPU 上执行
__global__ void add(const real* x, const real* y, real* z);
void check(const real* z, const int N);

int main(void) {
    const int N = 100000000; // 定义数组的长度为 10 的 8 次方
    const int M = sizeof(real) * N; // 每个数组所需的字节数

    // 分配host内存
    real* h_x = (real*)malloc(M);
    real* h_y = (real*)malloc(M);
    real* h_z = (real*)malloc(M);


    for (int n = 0; n < N; ++n) {
        h_x[n] = a;
        h_y[n] = b;
    }

    //分配device内存
    real* d_x, * d_y, * d_z;
    CHECK(cudaMalloc((void**)&d_x, M));
    CHECK(cudaMalloc((void**)&d_y, M));
    CHECK(cudaMalloc((void**)&d_z, M));

    // 将数据从主机复制到设备上
    CHECK(cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice));
    


    const int block_size = 128;
    // 计算网格尺寸，确保所有元素都能被处理
    const int grid_size = (N + block_size - 1) / block_size;

    cudaEvent_t start, stop;
    CHECK(cudaEventCreate(&start));
    CHECK(cudaEventCreate(&stop));
    CHECK(cudaEventRecord(start));


    // 调用内核函数在设备中进行计算
    add << <grid_size, block_size >> > (d_x, d_y, d_z);

    CHECK(cudaEventRecord(stop));
    CHECK(cudaEventSynchronize(stop));

    // 将计算结果从设备复制回主机
    CHECK(cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost));
    check(h_z, N);
    float elapsed_time;
    CHECK(cudaEventElapsedTime(&elapsed_time, start, stop));
    printf("Elapsed time: %f ms\n", elapsed_time);

    // 释放内存
    CHECK(cudaEventDestroy(start));
    CHECK(cudaEventDestroy(stop));
    free(h_x);
    free(h_y);
    free(h_z);
    CHECK(cudaFree(d_x));
    CHECK(cudaFree(d_y));
    CHECK(cudaFree(d_z));
    return 0;
}

__global__ void add(const real* x, const real* y, real* z) {
    const int n = blockIdx.x * blockDim.x + threadIdx.x;
    z[n] = x[n] + y[n];
}

void check(const real* z, const int N) {
    bool has_error = false;
    for (int n = 0; n < N; ++n) {
        if (fabs(z[n] - c) > EPS) {
            has_error = true;
        }
    }
    printf("Has error: %d\n", has_error);
}

①单精度编译过程和输出结果

nvcc -o singel_cuda addFunction.cu -arch=sm_75

运行后输出如下：
在这里插入图片描述
②双精度编译过程和输出结果

nvcc -DUSE_DP -o double_cuda addFunction.cu -arch=sm_75

运行后输出如下：
在这里插入图片描述

观察结果，我们发现单精度的运行时间是10.624ms；双精度的运行时间是21.490ms

1.2 该计算任务并不适合使用GPU进行加速

我们把1.1代码中的数据复制步骤也加入到计时当中，观察耗时情况：
①单精度输出结果
在这里插入图片描述
①双精度输出结果
在这里插入图片描述
观察发现，核函数的运行时间连整体运行时间的10%都没有，若是算上CPU和GPU之间的传输时间，把该程序放入GPU中运算的性能，可能还不如直接在CPU上运行。

这里可以使用CUDA自带的nvprof工具对程序进行性能分析：

nvprof .\singel_cuda.exe

输出结果如下：
在这里插入图片描述
根据分析结果可以得出

  • Host-to-Device (HtoD) 内存复制：
    占总 GPU 时间的 58.98%，耗时 127.74 毫秒。
  • Device-to-Host (DtoH) 内存复制：
    占总 GPU 时间的 36.13%，耗时 78.259 毫秒。
  • CUDA 内核函数 add 的执行：
    占总 GPU 时间的 4.89%，耗时 10.590 毫秒

这意味着在总的 GPU 活动时间中，大约 95.11% 的时间都花在了内存复制上，而只有 4.89% 的时间用于实际的计算。所以这样的任务，其实是不适合使用GPU进行 “加速” 的，那么什么任务才能真正的发挥GPU加速能力呢？

2 影响GPU加速的关键因素

2.1 数据传输比例

从上一个例子我们可以得出，如果一任务仅仅是计算两个数组的和，那么用GPU可能比用CPU还慢，因为花在CPU与GPU之间传输的时间比计算时间还要多太多。

所以一个适合使用GPU加速的任务，一定是数据传输占比时间少的任务，尽量让一些操作在GPU中完成，避免过多数据经过PCIe传输，例如做10000次数组相加，只在开头和结尾进行数据传输（H to D/D to H)

下面是把上面代码的add操作重复1000次：

    // 调用内核函数在设备中进行计算 1000 次
    for (int i = 0; i < 1000; ++i) {
        add << <grid_size, block_size >> > (d_x, d_y, d_z);
    }

运行性能分析，结果输出如下：
在这里插入图片描述
性能分析结果显示，add函数运行耗时占比97.96%，数据传输占比是2.04%，所以这样的任务就适合使用GPU加速

2.2 算术强度（arithmetic intensity）

算术强度： 计算过程中浮点运算次数与读写内存字节数的比例。

在上述例子中，我们只进行了加法运算，接下来我们修改核函数，进行一些更复杂的数学运算，代码如下：

__global__ void add(const real* x, const real* y, real* z) {
    const int n = blockIdx.x * blockDim.x + threadIdx.x;
    real x_val = x[n];
    real y_val = y[n];

    // 复杂的数学运算
    real result = sin(x_val) + cos(y_val) + exp(x_val * y_val) / (1.0 + x_val * y_val);
    result += log10(x_val + 1.0) * sqrt(y_val);

    // 最终结果
    z[n] = result;
}

性能分析结果如下：
在这里插入图片描述
性能分析结果显示，运算时间占比为31.71%，传输数据时间占比为68.29%，比之前只做一次加法运算的操作更适合用GPU加速（之前运算时间占比是4.89%）

2.3 并行规模

并行规模： 指的是在GPU上同时执行的线程数量。

因为GPU上的线程是可以并行执行的，在设计核函数时，尽量让设备中的所有线程都要参与到计算之中，可以最大程度的加速CUDA程序，下面是两幅图：
在这里插入图片描述
N是指放到GPU上进行运算的数据规模。

  1. 左图N在3次方和4次方时，耗时差距不大，是因为这两个时候，GPU的线程还有空闲，即所有数据都有线程在处理。从5次方开始，因为数据规模以及超过线程数量了，所以需要排队等待计算，故而随着N的增加，耗时成比例增加
  2. 有图是和CPU相比，GPU的运算加速比。可以发现在3次方和4次方时，加速比增大，因为GPU有大量的线程可以并行，远远比CPU运算快。在达到5次方时，线程以及利用完毕，也得排队等待计算，所以加速比几乎不变

3 总结

在编写CUDA程序时，一定要做到以下三点：

  1. 减少主机和设备之间的数据传输时间占比、也要减少数据传输次数
  2. 提高核函数的算术强度
  3. 增大核函数的并行规模

附：下面给出CUDA自带的数学函数库网站

http://docs.nvidia.com/cuda/cuda-math-api
包含幂函数、三角函数、指数函数、对数函数等，在编写代码时，要注意单精度和双精度的使用范围，例如有的计算精度不高的计算可以使用单精度，可以大大提升CUDA程序性能


http://www.niftyadmin.cn/n/5691967.html

相关文章

《数据结构》--队列【各种实现，算法推荐】

《数据结构》--队列【各种实现，算法推荐】

一、认识队列 队列是一种常见的数据结构&#xff0c;按照先进先出&#xff08;FIFO&#xff0c;First In First Out&#xff09;的原则排列数据。也就是说&#xff0c;最早进入队列的元素最先被移除。队列主要支持两种基本操作&#xff1a; 入队&#xff08;enqueue&#xff0…
阅读更多...
ElasticSearch备考 -- Multi match

ElasticSearch备考 -- Multi match

一、题目 索引task有3个字段a、b、c&#xff0c;写一个查询去匹配这三个字段为mom&#xff0c;其中b的字段评分比a、c字段大一倍&#xff0c;将他们的分数相加作为最后的总分数 二、思考 通过题目要求对多个字段进行匹配查询&#xff0c;可以考虑multi match、bool query操作。…
阅读更多...
结合大语言模型的机械臂抓取操作简单介绍

结合大语言模型的机械臂抓取操作简单介绍

一、大语言模型与机械臂抓取的基本操作 1. 大语言模型简介 大语言模型是基于深度学习技术构建的自然语言处理模型&#xff0c;能够生成、理解和处理文本信息。这些模型通过训练大量的文本数据&#xff0c;学习语法、上下文和常识&#xff0c;能够执行多种任务&#xff0c;如文…
阅读更多...
涂色问题 乘法原理（2024CCPC 山东省赛 C）

涂色问题 乘法原理（2024CCPC 山东省赛 C）

//*下午打得脑子连着眼睛一起疼 很多很基础的题目都没有做出来&#xff0c;规律题也找得很慢。比如下面这题&#xff0c;一定要多做&#xff0c;下次看到就直接写。 原题链接&#xff1a;https://codeforces.com/group/w6iGs8kreW/contest/555584/problem/C C. Colorful Segm…
阅读更多...
如何在 SQL 中创建一个新的数据库？

如何在 SQL 中创建一个新的数据库？

在SQL中创建一个新的数据库&#xff0c;首先你需要有一个可以执行SQL语句的环境。 这通常意味着你已经有了一个数据库管理系统&#xff08;DBMS&#xff09;&#xff0c;如MySQL、PostgreSQL、Oracle或Microsoft SQL Server等。 不同的DBMS可能有不同的细节&#xff0c;但基本…
阅读更多...
每日读则推(五)——Rose

每日读则推(五)——Rose

n/v. 照顾 n.护理 We sponsored the care these dogs needed and most of them found forever homes! One n.赞助者,赞助商 v.赞助(活动、节目等),为慈善活动捐资 lovely senior, Rose, is still waiting for hers. Shes mostly blind and needs eye medication, …
阅读更多...
讯飞星火编排创建智能体学习（五）：变量和文本拼接

讯飞星火编排创建智能体学习（五）：变量和文本拼接

引言 在讯飞星火编排创建智能体学习&#xff08;四&#xff09;&#xff1a;网页读取-CSDN博客中&#xff0c;我介绍了如何用网页读取功能从网上搜索车次信息。其中&#xff0c;我使用用大模型节点从文本中提取车次并合成了所需要的URL&#xff0c;今天介绍一下如何用变量和文…
阅读更多...
k8s的简介和部署

k8s的简介和部署

一、k8s简介 在部署应用程序的方式上面&#xff0c;主要经历了三个阶段&#xff1a; 传统部署:互联网早期&#xff0c;会直接将应用程序部署在物理机上优点:简单&#xff0c;不需要其它技术的参与缺点:不能为应用程序定义资源使用边界&#xff0c;很难合理地分配计算资源&…
阅读更多...
最新文章

Copyright @ 2022~2023