Cuda Program Basic

CUDA编程水平高低的不同，会导致几十上百倍的性能差距。但是这篇将聚焦于CUDA的编程语法，编译与运行。

编程语法

函数前缀

与函数调用设备有关

函数前缀名称	作用	目的
global	指定函数是CPU上调用，GPU上执行	GPU设置规模参数<<<1,12>>>
device	指定函数是GPU上调用，GPU上执行	GPU内执行函数
host	指定函数是CPU上调用，CPU上执行(最正常的函数，平常就省略不写)	CPU内执行函数

• 如果一个函数不加修饰，默认他是 _device_ 函数，正如上面的 main 一样。
• functions that are decorated with both __host__ and __device__ labels will be compiled to run on both, the host and the device.

变量修饰符

变量修饰符	作用
device	数据存放在显存中，所有的线程都可以访问，而且CPU也可以通过运行时库访问
shared	数据存放在共享存储器在，只有在所在的块内的线程可以访问，其它块内的线程不能访问
constant	数据存放在常量存储器中，可以被所有的线程访问，也可以被CPU通过运行时库访问
Texture	纹理内存（Texture Memory）也是一种只读内存。
/	没有限定符，那表示它存放在寄存器或者本地存储器中，在寄存器中的数据只归线程所有，其它线程不可见。

SMEM 静态与动态声明

// array with a fixed size
__shared__ float s_in[34];
// allocate the array dynamically,
extern __shared__ float s_in[];

动态的s_in大小，在kernel的第三个参数指定smemSize字节数

int smemSize = (TPB + 2)*sizeof(float);
ddKernel <<< (n+TPB-1)/TPB, TPB, smemSize>>> (args)

配置运算符

执行配置运算符 <<< >>>，用来传递内核函数的执行参数。执行配置有四个参数，

第一个参数声明网格的大小，

第二个参数声明块的大小，

第三个参数声明动态分配的共享存储器大小，默认为 0，

最后一个参数声明执行的流，默认为 0.

add<<<grid,block>>>(a,b);

stream

CUDA内置变量

变量	意义
gridDim	gridDim 是一个包含三个元素 x,y,z 的结构体，分别表示网格在x,y,z 三个方向上的尺寸(一般只有2维度)
blockDim	blockDim 也是一个包含三个元素 x,y,z 的结构体，分别表示块在x,y,z 三个方向上的尺寸
blockIdx	blockIdx 也是一个包含三个元素 x,y,z 的结构体，分别表示当前线程块在网格中 x,y,z 三个方向上的索引
threadIdx	是一个包含三个元素 x,y,z 的结构体，分别表示当前线程在其所在块中 x,y,z 三个方向上的索引
warpSize	在计算能力为 1.0 的设备中，这个值是24，在 1.0 以上的设备中，这个值是 32

三维的举例

__global__ void kernel() {  
   printf("Block (%d,%d,%d) of (%d,%d,%d), Thread (%d,%d,%d) of (%d,%d,%d)\n",  
          blockIdx.x, blockIdx.y, blockIdx.z,  
          gridDim.x, gridDim.y, gridDim.z,  
          threadIdx.x, threadIdx.y, threadIdx.z,  
          blockDim.x, blockDim.y, blockDim.z);  
}  
  
int main() {  
   kernel<<<dim3(2, 1, 1), dim3(2, 2, 2)>>>();  
   cudaDeviceSynchronize();  
   return 0;  
}

Block (0,0,0) of (2,1,1), Thread (0,0,0) of (2,2,2)  
Block (0,0,0) of (2,1,1), Thread (1,0,0) of (2,2,2)  
Block (0,0,0) of (2,1,1), Thread (0,1,0) of (2,2,2)  
Block (0,0,0) of (2,1,1), Thread (1,1,0) of (2,2,2)  
Block (0,0,0) of (2,1,1), Thread (0,0,1) of (2,2,2)  
Block (0,0,0) of (2,1,1), Thread (1,0,1) of (2,2,2)  
Block (0,0,0) of (2,1,1), Thread (0,1,1) of (2,2,2)  
Block (0,0,0) of (2,1,1), Thread (1,1,1) of (2,2,2)  
Block (1,0,0) of (2,1,1), Thread (0,0,0) of (2,2,2)  
Block (1,0,0) of (2,1,1), Thread (1,0,0) of (2,2,2)  
Block (1,0,0) of (2,1,1), Thread (0,1,0) of (2,2,2)  
Block (1,0,0) of (2,1,1), Thread (1,1,0) of (2,2,2)  
Block (1,0,0) of (2,1,1), Thread (0,0,1) of (2,2,2)  
Block (1,0,0) of (2,1,1), Thread (1,0,1) of (2,2,2)  
Block (1,0,0) of (2,1,1), Thread (0,1,1) of (2,2,2)  
Block (1,0,0) of (2,1,1), Thread (1,1,1) of (2,2,2)

二维的例子,最后一个维度都是 0, 我们使用结果的时候不使用 z 维度即可

__global__ void kernel() {  
   printf("Block (%d,%d,%d) of (%d,%d,%d), Thread (%d,%d,%d) of (%d,%d,%d)\n",  
          blockIdx.x, blockIdx.y, blockIdx.z,  
          gridDim.x, gridDim.y, gridDim.z,  
          threadIdx.x, threadIdx.y, threadIdx.z,  
          blockDim.x, blockDim.y, blockDim.z);  
}  
  
int main() {  
   kernel<<<dim3(2, 3, 1), dim3(2, 1, 1)>>>();  
   cudaDeviceSynchronize();  
   return 0;  
}

Block (1,2,0) of (2,3,1), Thread (0,0,0) of (2,1,1)  
Block (1,2,0) of (2,3,1), Thread (1,0,0) of (2,1,1)  
Block (0,2,0) of (2,3,1), Thread (0,0,0) of (2,1,1)  
Block (0,2,0) of (2,3,1), Thread (1,0,0) of (2,1,1)  
Block (0,1,0) of (2,3,1), Thread (0,0,0) of (2,1,1)  
Block (0,1,0) of (2,3,1), Thread (1,0,0) of (2,1,1)  
Block (1,0,0) of (2,3,1), Thread (0,0,0) of (2,1,1)  
Block (1,0,0) of (2,3,1), Thread (1,0,0) of (2,1,1)  
Block (0,0,0) of (2,3,1), Thread (0,0,0) of (2,1,1)  
Block (0,0,0) of (2,3,1), Thread (1,0,0) of (2,1,1)  
Block (1,1,0) of (2,3,1), Thread (0,0,0) of (2,1,1)  
Block (1,1,0) of (2,3,1), Thread (1,0,0) of (2,1,1)

常用函数

调用 GPU 的函数声明和定义不要分离，写在同一个文件里。分开(如：CUDA_SEPARABLE_COMPILATION)可能影响内联导致性能损失。

访存

__host____device__cudaError_t  cudaMalloc ( void** devPtr, size_t size )
cudaMallocPitch() //分配二维数组空间并自动对齐
//在显存中为待运算的数据以及需要存放结果的变量开辟显存空间。
__host____device__cudaError_t cudaFree ( void* devPtr )
__host__cudaError_t cudaMemcpy ( void* dst, const void* src, size_t count, cudaMemcpyKind kind )

• where kind specifies the direction of the copy, and must be one of cudaMemcpyHostToHost, cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, cudaMemcpyDeviceToDevice, or cudaMemcpyDefault.
• Passing cudaMemcpyDefault is recommended, in which case the type of transfer is inferred from the pointer values. However, cudaMemcpyDefault is only allowed on systems that support unified virtual addressing.
• Calling cudaMemcpy() with dst and src pointers that do not match the direction of the copy results in an undefined behavior.

cudaMemcpy可以自动实现同步工作，可以省去cudaDeviceSynchronize。

可以通过 cudaMallocManaged(&a, sizeof(int) * 12)申请在 Host 和 Device 上都直接使用的Unified Memory。性能多数情况会损失。

同步

__host____device__cudaError_t  cudaDeviceSynchronize ( void )
//Wait for compute device to finish.

__syncthreads() //block内线程快速同步

Stream

CUDA enables developers to define independent streams of commands, where it is assumed that commands in different streams do not depend on each other.

字符打印输出

很明显CPU和GPU打印是异步的，需要同步。

而且cuda暂时不支持cout等流输出语句。

Debug打印

cudaError_t是不能理解的输出。 cuda samples 里面提供了 helper_cuda.h 头文件解决问题。 Debug 的时候也可以直接把 gridDim 改成 1, 更方便

# CMakeLists.txt
target_include_directories(hello PUBLIC /usr/local/cuda/samples/common/inc)

checkCudaErrors(cudaDeviceSynchronize());

时间统计打印

cudaEvent_t begin, end;
cudaEventCreate(&begin);
cudaEventCreate(&end);

cudaEventRecord(begin);

// do sth

cudaEventRecord(end);
cudaEventSynchronize (end);

float elapsedTime;
cudaEventElapsedTime (&elapsed, begin, end);
elapsedTime /= 1000;

cudaEventDestroy (end);
cudaEventDestroy (begin);

return elapsedTime;

函数指针和lambda算子

template <class Func>  
__global__ void kernel(int *arr, int n, Func func) {  
   for (int i = blockDim.x * blockIdx.x + threadIdx.x;  
        i < n; i += blockDim.x * gridDim.x) {  
       func(arr, i);  
   }  
}  
  
struct funcop1 {  
   __device__ void operator()(int *arr, int i) {  
       arr[i] = i;  
   }  
};  
  
struct funcop2 {  
   __device__ void operator()(int *arr, int i) {  
       printf("%d %f\n", arr[i], sinf(arr[i]));  
   }  
};

//使用
kernel<<<grid_dim, block_dim>>>(arr, n, funcop1{}); 
kernel<<<grid_dim, block_dim>>>(arr, n, funcop2{});

// lambda算子
template <class Func>  
__global__ void kernel(int n, Func func) {  
   for (int i = blockDim.x * blockIdx.x + threadIdx.x;  
        i < n; i += blockDim.x * gridDim.x) {  
       func(i);  
   }  
}  

kernel<<<grid_dim, block_dim>>>(n, [=] __device__ (int i) {  
       arr[i] = i;  
   });
// 或者
kernel<<<grid_dim, block_dim>>>(n, [=] __device__ (int i) {  
       printf("%d, %f\n", i, sinf(arr[i]));  
   });

// lambda算子例子2
template <class Func>  
__global__ void kernel(int n, Func func) {  
   for (int i = blockDim.x * blockIdx.x + threadIdx.x;  
        i < n; i += blockDim.x * gridDim.x) {  
       func(i);  
   }  
}  

kernel<<<grid_dim, block_dim>>>(n, [x = x_dev.data(), y = y_dev.data()] __device__ (int index){  
       x[index] = x[index] + y[index];  
   });

cuda 容器的实现——thrust

STL 容器 cuda 并没有很好的适配和实现，CUDA对应的叫做thrust 库被称为： Template library for CUDA
ref1 and ref2

thrust::host_vector<float> x_host(n);
thrust::generate(x_host.begin(), x_host.end(), []{return std::rand() / 3.0;});

thrust::device_vector<float> x_dev(n); 
x_dev = x_host;

全局变量传递

GPU计算的全局变量 sum最后传递到CPU的 result里

__device__ float sum = 0;
...

int main() {
    float result = 0;
...

cudaMemcpyFromSymbol(&result, sum, sizeof(float), 0, cudaMemcpyDeviceToHost);

常见原子操作

atomicAdd (dst, src)
atomicSub(dst, src)
atomicOr(dst, src)
atomicAnd(dst, src)
atomicXor(dst, src)
atomicMax(dst, src)
atomicMin(dst, src)

他们都有返回值，返回违背更改前的数值。

也可以通过 atomicCAS自定义原子操作。但是前面的原子操作有特殊设计的，会基于blockDim和gridDim,并行各块串行执行然后规约。

单卡多GPU的实现

int gpu_numbers = cudaGetDeviceCount();
int *pointers[gpu_numbers];

for (int index = 0; index < gpu_numbers; ++index) {
   cudaSetDevice(index);
   cudaMalloc(&pointers[index], size);
}//在各自卡上声明空间

for (int indexi = 0; indexi < gpu_numbers; ++indexi) {
   cudaSetDevice(indexi); //设置当前卡
   for (int indexj = 0; indexj < gpu_numbers; ++indexj) {
      if (indexi == indexj)
            continue;
      cudaDeviceEnablePeerAccess(indexj, 0); //打通indexj与当前卡的访问
   }
}

for (int index = 1; index < gpu_numbers; ++index) {
   cudaMemcpyAsync(pointers[0], pointers[index], size, cudaMemcpyDeviceToDevice); //非阻塞memoryCopy，在这里实现device0到其他的广播
}

指定某卡运行程序

通过环境变量实现

export CUDA_VISIBLE_DEVICES=1
export CUDA_VISIBLE_DEVICES=0,1 # 多卡
CUDA_VISIBLE_DEVICES=1 ./cuda_executable

GPU 编译器

相对于CPU编译器简单一些

可能要手动循环展开, 消除分支，GPU分支预测几乎没有

#pragma unroll 一句即可展开

nvcc优化选项

target_compile_options(${exe}  PUBLIC $<$<COMPILE_LANGUAGE:CUDA>:
   -Xptxas 
   -O3 
   -v 
   --use_fast_math
 >)

fast math

–-use_fast_math对于频繁的数学函数：三角函数、快速傅立叶变换、幂次、根号有5~15%的效率提升。

ECC

ECC(error correcting code, 错误检查和纠正)能够提高数据的正确性，随之而来的是可用内存的减少和性能上的损失。对于Tesla系列伺服器该功能默认开启。

通过命令 nvidia-smi -i n可查看第n个个显卡的简要信息（详细信息可通过 nvidia-smi -q -i 0获取），其中有一项是volatile Uncorr。

通过 nvidia-smi -i n -e 0/1 可关闭(0)/开启(1)第n号GPU的ECC模式。

通过实践，关闭ECC程序的性能能得到13%~15%的提升。

测试运行

现有cuda 是兼容 C++17 语法的，可以减少移植工作量

export CUDA_ROOT=/usr/local/cuda/bin
export PATH=$CUDA_ROOT:$PATH
which nvcc
nvcc -V
nvcc src.cu -o a.out
./a.out

发现版本太老了不支持更新的gcc，自己安装最新cuda

CUDA实例

CUDA项目

https://github.com/Kirrito-k423/StencilAcc

一维的例子

先将数据分成多个block,每个block里面进行第一遍归约。

第二个for的作用

for 循环中的算法就是将数组的后一半加到前一半上去,然后再在前一半中的后一半加到前一半的前一半中…

这中被称为“对数归约”,循环完成后一个block 中的和是sPartials[0]的值.

接着，将这个值导出到out中.

杂项

GPU线程的创建与调度

shared memory In Stencil Computing

问题

2. thread 和硬件的关系？
3. shared memory位置和cache的关系（根据GA100，L1 data cache=shared memory）
   1. 联合访问搬数据，没有cache line的概念吗？
4. shared memory VS streaming Multiprocessor
   1. https://blog.csdn.net/qq_41598072/article/details/82877655
   2. https://blog.csdn.net/junparadox/article/details/50540602

参考文献

实例：手写 CUDA 算子，让 Pytorch 提速 20 倍

https://docs.nvidia.com/cuda/cuda-c-programming-guide/#function-parameters

例子代码:

https://github.com/chivier/cutests

https://chivier.github.io/2022/02/20/2022/2202-CudaProgramming/

https://chivier.github.io/2022/04/11/2022/2204-GPU%E7%A8%8B%E5%BA%8F%E4%BC%98%E5%8C%96%E6%96%B9%E6%B3%95/

https://comzyh.com/blog/archives/967/

https://itlanyan.com/cuda-enable-disable-ecc/

[^1]: 并行计算课程-CUDA 密码pa22