介绍

以向量加法为例介绍了CUDA程序得基本结构，使用xmake进行项目组织及编译。
相关代码：https://github.com/JilinLi4/LearningCUDA

一、 相关API

1. 设备CUDARuntime 初始化

在CUDA运行时API中，没有显示地初始化设备的函数。在第一次调用一个和设备管理及版本查询无关的运行时API函数时，设备将自动地初始化。

2. 设备内存的分配及释放

申请显存

cudaError_t cudaMalloc(void **address, size_t size);

• address 待分配GPU内存的指针。
• size 分配内存的字节数
• 返回错误码，如果成功，返回 cudaSuccess

释放显存

cudaError_t cudaFree(void* address);

3. 主机与设备之间数据的传递

cudaError_t cudaMemcpy(void                *dst,const void          *src,size_t              count,enum cudaMemcpyKind kind
)

• dst 是目标地址
• src 是源地址
• count 复制数据的字节数
• kind 一个枚举类型的变量，取如下几个值：
  • cudaMemcpyHostToHost 从主机复制到主机
  • cudaMemcpyHostToDevice 从主机复制到设备
  • cudaMemcpyDevicetoHost 从设备复制到主机
  • cudaMemcpyDefault, 根据指针dst和src所指地址自行判断数据传输方向。

二、 核函数中数据与线程的对应

1. 使用多个线程的核函数

核函数中允许指派很多线程。因为一个GPU中有几千个计算核心，而总的线程数必须至少等于计算核心数才能充分利用GPU中的全部计算资源。
所以根据需要，在调用核函数时可以指定多个线程。比如：

func<<<2, 4>>>();

func指定了2个线程块，每个线程块中包含4个线程。总的线程数是 2×4=82 \times 4 = 82×4=8。也就是说，该程序中的核函数中代码执行方式是：“单指令-多线程”，即每一个线程都执行同一串指令。

2 使用线程索引

线程组织结构是由调用函数配置的。如：

func<<>>();

一般来说：

• gridSize≤231−1\text{gridSize} \le 2^{31}-1gridSize≤231−1
• blockSize≤1024\text{blockSize} \le 1024blockSize≤1024
  每个线程在核函数中都有一个唯一的身份表示。由于我们用两个参数指定线程数目。

3 核函数中的数据与线程对应

将有关数据从主机传至设备之后，就可以调用核函数在设备中进行计算了。使用以下线程布局:

• blockSize=128\text{blockSize} = 128blockSize=128
• gridSize=108/128\text{gridSize} = 10^8 / 128gridSize=108/128
  在这种情况下，核函数的数据与线程对应关系为：

const int n = blockDim.x * blockIdx.x + threadIdx.x;

4 是否需要使用 if 判断 n是否越界？

核函数没有使用参数N，当N是blockDim.x的整数倍时，不会引起问题，因为核函数的线程数刚好等于数组元素个数。然而，当N不是 blockDim.x的整数倍时，就有可能发生错误。例如将 N改为： 108+110^8 + 1108+1 而 block_size依然是128,如果griadSize的计算方式不变，那么必将有一个元素无法处理。所以我们应该将原来的计算方式：
gridSize=N/blockSize\text{gridSize} = N / \text{blockSize}gridSize=N/blockSize 改为
gridSize=(N+blockSize−1)/blockSize\text{gridSize} = (N + blockSize - 1) / \text{blockSize}gridSize=(N+blockSize−1)/blockSize

这样一来也会引入越界的风险，所以需要将线程索引的计算方式改为：

const int n = blockDim.x * blockIdx.x + threadIdx.x;
if(n >= N) return;
z[n] = x[n] + y[n];

三、 自定义核函数

1. 核函数的要求

• 返回值必须是void,可以使用return 关键字，但是不可返回任何值
• 必须使用限定符 global 也可以加上一些其他C++的限定符，如：static 限定符的次序任意
• 函数名支持重载
• 不可以向核函数传递非指针变量
• 除非使用统一内存编程机制，否则传给核函数的指针必须只想设备内存
• 核函数不可成为一个类的成员。通常做法是用一个包装函数调用核函数，将包装函数定义为一个类。
• 核函数之间在算力 > 3.5 开始，可以调用其他核函数，也可以调用自己。
• 核函数必须指定线程布局配置。

2. 函数执行空间标识符

• __ global__ 称为核函数，一般由主机调用，也可以被核函数包括自己调用。
• __ device__ 称为设备函数，只能被核函数或其他设备函数调用，在设备中执行。
• __host __ 主机段普通的C++函数，在主机中被调用，在主机执行。可省略，存在的意义在于，可以使用**host**和 **device**修饰同一个函数，表示即可在设备端调用也可以在主机端调用，减少代码冗余。
• 不能使用__global__ 和 __device__修饰统一个函数
• 不能使用__host__和 __device__修饰统一个函数
• 编译器决定把设备函数当做内联函数或非内联函数，但可以使用 __noinline__建议一个设备函数为非内联函数，也可以用 __forceinline__建议一个设备为内联函数。

3. 实例代码

3.1 code1 返回值

double __device__ add1_device(const double x, const double y) {return x + y;
}void __global__ add1(const double *x, const double *y, double *z, const int N) {const int n = blockDim.x * blockIdx.x + threadIdx.x;if(n < N) {z[n] = add1_device(x[n], y[n]);}
}

3.2 code2 传指针

void __device__ add2_device(const double x, const double y, double *z) {*z = *x + *y;
}void __global__ add1(const double *x, const double *y, double *z, const int N) {const int n = blockDim.x * blockIdx.x + threadIdx.x;if(n < N) {add1_device(x[n], y[n], &z[n]);}
}

3.3 code3 传引用

void __device__ add3_device(const double x, const double y, double &z) {z = x + y;
}void __global__ add1(const double *x, const double *y, double *z, const int N) {const int n = blockDim.x * blockIdx.x + threadIdx.x;if(n < N) {add1_device(x[n], y[n], z[n]);}
}

CUDA 错误检查

bool checkRuntime(cudaError_t e, const char* call, int line, const char *file){if (e != cudaSuccess) {INFOE("CUDA Runtime error %s # %s, code = %s [ %d ] in file %s:%d", call, cudaGetErrorString(e), cudaGetErrorName(e), e, file, line);return false;}return true;}

错误信息实例：

CUDA Runtime error cudaMemcpy(d_x, h_x, M, cudaMemcpyDeviceToDevice) # invalid argument, code = cudaErrorInvalidValue [ 1 ] in file D:\work_dir\experiment\LearningCUDA\src\add.cu:29