CUDA 编程经验分享

最新推荐文章于 2023-08-29 19:42:58 发布

不正道

最新推荐文章于 2023-08-29 19:42:58 发布

阅读量556

点赞数

分类专栏：踩坑记录文章标签：经验分享 c++

本文链接：https://blog.csdn.net/weixin_49069890/article/details/126973952

版权

踩坑记录专栏收录该内容

9 篇文章 0 订阅

订阅专栏

CUDA简介
- CUDA是Nvidia公司推出的计算平台，也是GPU的接口，当然也只适用于Nvidia的GPU。
- CUDA可对GPU编程，利用GPU多核心的特点，开多线程并行处理数据，大大提高程序运算速度。
硬件平台
- 使用CUDA编程必须拥有Nvidia显卡，且该显卡支持CUDA。
- 计算平台可分为通用平台和异构平台。通用平台包括电脑主机、服务器、笔记本；异构平台包括各类GPU开发板。
  异构平台
  通用平台
- 通用平台的内存和显存的物理空间是分开的，异构平台内存和显存使用同一块物理空间。在两种平台上分别运行下面的代码，会发现：通用平台打印的地址不一样，而异构平台打印的地址一样。也就是说通用平台会开辟100byte的内存和100byte的显存，这两块空间会自动互相拷贝以保持同步；异构平台只会开辟一段100byte的空间，这段空间即使内存又是显存。
```
__global__ void Check_GPU_Addr(int *A)
{
    printf("GPU memary address: %p\n", A);
}

void Check_CPU_Addr(int *A)
{
    printf("GPU memary address: %p\n", A);
}

int main(int arvc, char **argv)
{
    int *A;
    cudaMallocManaged(&A, 100);
    
    Check_CPU_Addr(A);
    Check_GPU_Addr<<<1, 1>>>(A);
    cudaDeviceSynchronize();

    return 0;
}
```
软件环境
- 需要Windows或Linux系统。不支持虚拟机。
- 需要CUDA开发套件（软件）。
- 使用C/C++ 语言编程，当然用Python的小伙伴也可以自行探索。
- 使用nvcc编译器。
- IDE无所谓，但想要使用框架就麻烦了。例如用QT就需要用Qmake，由于核函数的存在Qmake无法编译.cu文件。非要结合的话，需要将核函数封装成C/C++的函数。这一步的坑放到后面讲。
核函数
- CUDA编程核心就是写核函数。核函数是一个入口，核函数中的代码由GPU执行。
- 核函数由__global__声明，返回值必须为void类型。
核函数执行过程
- 核函数<<<grid_size, block_size, dynamic_size, stream>>>中的四个参数：分别为线程块数、每块中线程数、动态内存大小、流号。前两个参数是dim3类型，默认值是{1， 1， 1}；第三个参数uint32类型，是分配给这个核函数的动态共享内存大小（不是零拷贝内存），等同于共享内存的堆的概念；最后一个参数是流号。
```
// dim3结构类型，此结构体构造时会给x,y,z都赋初值1
struct __device_builtin__ dim3
{
    unsigned int x, y, z;
};
```
  核函数线程分配示意图
- 核函数是可以嵌套核函数的（也叫动态并行）。核函数嵌套需要满足以下条件：
  1. 显卡版本不能太低，（sm_61以前的）。
  2. 编译flag添加 -rdc=true。
  3. 编译flag添加 -arch或-gencode，如-arch=sm_86，-gencode=arch=compute_75,code=compute_75，这关系到代码在设备上的兼容性问题，具体可参考CUDA版本——设备架构——gencode匹配关系。当然，如果你不清楚的话，可以把所有的gencode都加上，就行这样：
```
#编译FLAG
-gencode=arch=compute_61,code=sm_61 \
-gencode=arch=compute_61,code=compute_61 \
-gencode=arch=compute_70,code=sm_70 \ 
-gencode=arch=compute_70,code=compute_70 \
-gencode=arch=compute_75,code=sm_75 \
-gencode=arch=compute_75,code=compute_75 \
-gencode=arch=compute_80,code=sm_80 \
-gencode=arch=compute_80,code=compute_80 \
-gencode=arch=compute_86,code=sm_86 \
-gencode=arch=compute_86,code=compute_86
```
```
// 核函数嵌套
__global__ void kernal_parent(int *A, int *B)
{
    kernal_children<<<gridDim.x, blockDim.x>>>(A, B);
}
```
- 核函数内部不能使用cudaMalloc()或cudaMallocManaged()也不能用cudaMemcpy()，因为这些都不是设备函数。想要在核函数内申请内存可直接使用malloc()或new，这时申请到的是设备内存。内存拷贝使用memcpy()。
```
	// 核函数内内存操作函数
	__global__ void kernal_malloc()
	{
	    char buf[] = "你好啊，我叫赛利亚~";
	    char *mem = new char[100];
	    // *mem = (char *)malloc(100);
	    memcpy(mem, buf, 100);
	    
	    printf("%s\n", mem);
	}
```
线程同步
- 线程同步的意义：因为核函数是非阻塞的，设备中所有线程和主机是同时运行的。主机想要拿到设备的计算结果就必须等待设备完成计算。当然主机也可趁这段时间做点爱做的事~，比如启动其他核函数或拷贝点数据什么的。
- 使用阻塞等待的方式同步线程（查询的方式没用过，不了解），常用线程同步函数有：
  1. __syncthreads()核函数内使用，此函数解释参考__syncthreads()同步方式，意思是等待能到达该点的线程都到达即为同步成功。但我不这么认为，例如下面的这段代码，说明了并不存在这么智能的方式。
```
// 由于条件判断的存在，多线程不会达到同一个同步点，此核函数永远无法跳出。
__global__ void kernal_sync()
{
    if(threadIdx.x == 0){
        printf("%d\n", threadIdx.x);
        __syncthreads();
    }else{
        printf("%d\n", threadIdx.x);
        __syncthreads();
    }
}
```
    而说它是等待核函数内所有线程到达该同步点也不对，例如下面：
```
// 这个函数是可以正常跳出的
__global__ void kernal_sync()
{
    if(threadIdx.x == 0){
        printf("%d\n", threadIdx.x);
        // __syncthreads();
    }else{
        printf("%d\n", threadIdx.x);
        __syncthreads();
    }
}
```
    总结来说：虽然我不确定它是怎么同步的，但我知道，绝对避免此函数在产生线程分化的地方使用。
    2. cudaDeviceSynchronize()等待所有线程到达该同步点。(因为没用过多GPU并行运算，所以不清楚是单个设备的所有线程，还是所有设备的所有线程。)
```
// 所有线程同步
kernal_sync<<<1, 5>>>();
cudaDeviceSynchronize();
```
    1. cudaStreamSynchronize()等待某个流的所有线程到达该同步点。
```
// 流同步
cudaStream_t stream;
cudaStreamCreate(&stream); //分配stream
kernal_sync<<<1, 5>>>();
cudaStreamDestroy(stream); 
```
返回值检查
- 一般核函数、cudaMalloc()、cudaMemcpy()等函数是不会提示段错误的，对于可能出现错误的地方需要检查函数返回值来确认工作状态。
- 在/usr/local/cuda/samples/common/inc/helper_cuda.h中有两个函数可以方便检查返回值，这两个函数在状态正常（返回值为0）时不打印任何东西。
  1. getLastCudaError()此函数等价于perror()，是专门用来检查返回全局errno的函数，会打印用户设置的字符串和错误类型。
```
// 检查全局设备errno，若错误则打印错误类型。
kernal_malloc<<<1, 5>>>();
cudaDeviceSynchronize();
getLastCudaError("核函数执行时出错");
```
  2. checkCudaErrors()此函数用于将错误码翻译成错误类型并打印。
```
// 检查返回值，若错误则打印错误类型。
cudaError_t err;
int *A;
err = cudaMalloc(&A, 10*sizeof(int));
checkCudaErrors(err);

int *B;
checkCudaErrors(cudaMalloc(&B, 10*sizeof(int)));
```
内存管理
- 在CUDA编程中，内存主要分为：主机内存，和设备内存两大类。设备内存有可分为：寄存器、本地内存、共享内存、全局内存、常量内存、纹理内存。内存详细解释参考这里
- 主机和设备逻辑上不能访问对方内存，除非使用零拷贝内存。所以需要互相拷贝，这里就不详细说明了。
- 零拷贝内存（统一内存）的使用大大简化了内存操作，让程序猿不至于频繁的做内存拷贝的这等下流事。零拷贝内存在通用平台和异构平台上是有区别的，详见第2节。零拷贝内存的申请方式有两种：
  1. 使用cudaMallocManaged()函数动态申请一块空间作为零拷贝内存。
```
// 申请零拷贝内存
int *A;
cudaMallocManaged(&A, 10*sizeof(int));
```
  2. 使用__managed__来声明全局的零拷贝内存，这种方式不推荐用于申请大空间。要注意的是，这种方式只能在全局变量区域定义，放在函数中这样写是不合法的。
```
// 声明零拷贝内存
__managed__ int A[10];

// 错误示例
void func(void)
{
	__managed__ int A[10];
}
```
- 关于函数传参
- 关于核函数的参数传递
nvcc混合编译