参加【Nvidia CUDA线上训练营2023春】笔记——尝试写自己的第一个CUDA程序

一、知识点梳理

基本的CUDA编程语法及如何使用Makefile构建编译体系

二、Makefile语法

利用nvcc编译可执行文件

!/usr/local/cuda/bin/nvcc -arch=compute_72 -code=sm_72 hello_cuda.cu -o hello_cuda -run   
//自行替换文件名，另外-arch和-code的参数需要根据硬件资源替换

利用Makefile编译可执行文件

//代码示例
TEST_SOURCE = hello_cuda.cu

TARGETBIN := ./hello_cuda


CC = /usr/local/cuda/bin/nvcc

$(TARGETBIN):$(TEST_SOURCE)
    $(CC)  $(TEST_SOURCE) -o $(TARGETBIN)

.PHONY:clean
clean:
    -rm -rf $(TARGETBIN)
    -rm -rf *.o
    
//需要注意的是.PHONY参数的作用是当有文件名为clean时，Makefile会先响应clean命令

三、CUDA语法

• gridDim表示一个grid中包含多少个block

• blockDim表示一个block中包含多少个线程

• threadIdx.x 是执行当前kernel函数的线程在block中的x方向的序号

• blockIdx.x 是执行当前kernel函数的线程所在block，在grid中的x方向的序号

线程编号为index = blockIdx.x * blockDim.x + threadIdx.x

CUDA代码实验示例

#include <math.h>
#include <stdio.h>

void __global__ add(const double *x, const double *y, double *z, int count)
//__global__为CUDA函数执行环境标识符（Function execution space specifiers）。
//__global__修饰的函数是核函数，这些函数 在GPU上执行 ，但是需要 在CPU端调用
//__global__修饰的函数必须采用void返回值，并且需要在调用时制定 运行的参数 (也就是<<<>>>里的block数和线程数)，这个格式在下面就可以看到
//同时__global___函数是异步的，这也代表着函数没被执行完就返回了控制权，所以测量核函数的时间需要同步操作才能获得准确的结果。
//此外还有__host__和__device__.__host__定义了主机端的函数，就是在主机端执行，在主机端调用的函数，也就是我们正常的c/c++的函数，如果不加任何的修饰符，默认就是__host__ 函数，这些函数只为主机端编译;__device__函数是在GPU端调用且在GPU端执行的的函数__global___和__device__不能同时使用。
//另外需要注意的是，__global__和__host__不能同时使用，__device__和__host__可以同时使用。
{
    const int n = blockDim.x * blockIdx.x + threadIdx.x;    //通过线程索引为每个线程分配各自的数据，这也是GPU只能用于简单逻辑，大数据量的场景，因为对线程批量化处理才高效

    具体的计算逻辑
    if( n < count)
    {
    z[n] = x[n] + y[n];          
    }

}


//在GPU里计算后再回到CPU进行验证，如果全对则打印Pass，如果有错则打印Errors
void check(const double *z, const int N)
{
    bool error = false;
    for (int n = 0; n < N; ++n)
    {
        if (fabs(z[n] - 3) > (1.0e-10))
        {
            error = true;
        }
    }
    printf("%s\n", error ? "Errors" : "Pass");
}



//主函数
int main(void)
{
    const int N = 1000;    //对总线程数进行定义
    const int M = sizeof(double) * N;   //定义数组大小
    //为在Host中数组申请内存
    double *h_x = (double*) malloc(M);
    double *h_y = (double*) malloc(M);
    double *h_z = (double*) malloc(M);

    //初始化计算数据
    for (int n = 0; n < N; ++n)
    {
        h_x[n] = 1;
        h_y[n] = 2;
    }
    
    double *d_x, *d_y, *d_z;
    //为在Device中用到的变量申请显存
    cudaMalloc((void **)&d_x, M);
    cudaMalloc((void **)&d_y, M);
    cudaMalloc((void **)&d_z, M);
    //将数据从Host内存复制到Device显存，格式为cudaMemcpy(1, 2, SIZE, cudaMemcpyHostToDevice)，1为Device变量，2为Host变量，SIZE为数据大小，cudaMemcpyHostToDevice为关键字，按字面意思理解就行
    cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);
    cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice);
    
    //规定thread数量和block数量
    const int block_size = 128;
    const int grid_size = (N + block_size - 1) / block_size;

    //调用Device函数，格式为函数名<<<block数量, 每个block的thread数量>>>(要传的参数);
    add<<<grid_size, block_size>>>(d_x, d_y, d_z, N);
    
    //将数据从Device显存复制到Host内存，格式为cudaMemcpy(1, 2, SIZE, cudaMemcpyDeviceToHost)，1为Host变量，2为Device变量，SIZE为数据大小，cudaMemcpyDeviceToHost为关键字，按字面意思理解就行
    cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost);
    check(h_z, N);
    
    //释放Host内存
    free(h_x);
    free(h_y);
    free(h_z);
    //释放Device内存
    cudaFree(d_x);
    cudaFree(d_y);
    cudaFree(d_z);
    return 0;
}

大家应该看到我上面的注释说到Host和Device，一般而言，Host指的是CPU，Device指的是GPU。 在CUDA程序构架中，主程序还是由 CPU 来执行，而当遇到数据并行处理的部分，CUDA 就会将程序编译成 GPU 能执行的程序，并传送到GPU。

在完成一次编译后，可以利用nvprof进行查看程序性能

!sudo /usr/local/cuda/bin/nvprof  ./hello_cuda
//自行替换文件

//此为打印信息，Profiling result是GPU（kernel函数）上运行的时间，API calls是在cpu上测量的程序调用API的时间
==17405== NVPROF is profiling process 17405, command: ./hello_cuda
==17405== Warning: Unified Memory Profiling is not supported on the underlying platform. System requirements for unified memory can be found at: http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#um-requirements
==17405== Profiling application: ./hello_cuda
==17405== Profiling result:
No kernels were profiled.
            Type  Time(%)      Time     Calls       Avg       Min       Max  Name
      API calls:   99.96%  575.20ms         1  575.20ms  575.20ms  575.20ms  cudaLaunchKernel
                    0.03%  170.26us        97  1.7550us     938ns  37.553us  cuDeviceGetAttribute
                    0.00%  16.876us         1  16.876us  16.876us  16.876us  cudaDeviceSynchronize
                    0.00%  12.812us         1  12.812us  12.812us  12.812us  cuDeviceTotalMem
                    0.00%  8.6980us         3  2.8990us  1.8220us  4.0630us  cuDeviceGetCount
                    0.00%  4.2200us         2  2.1100us  1.4590us  2.7610us  cuDeviceGet
                    0.00%  2.3960us         1  2.3960us  2.3960us  2.3960us  cuDeviceGetName
                    0.00%  1.6140us         1  1.6140us  1.6140us  1.6140us  cuDeviceGetUuid

四、总结

CUDA编程入门相对来说不是很难，都遵循一套流程，很容易就可以进行更改，新手入门最重要的就是转变代码编写思维，由原来的串行思路转为并行思路，充分发挥GPU加速计算的优点。