Cuda 基础

Cuda

Compute Unified Device Architecture，英伟达并行计算框架；
是暴露英伟达开发的 GPU 接口（有强计算形的 Task 可以考虑使用/新开发 Cuda 加速）；

一、理解 cuda 加速基本思想

线性代数矩阵乘法，A * B = C，C 中每个结果的计算互不干扰，可以并行；用 cpu 也可以并行：std::thread ，每一个 thread 计算一个 $ro{w}_i\ast co{l}_j$，cpu 代价高，涉及线程调度，资源切换，GPU 就这样诞生；

二、查看 Cuda 配置

Linux 查看显卡配置： lspci | grep -i vga，一般出厂服务器都有内嵌显卡，没有独立显卡，自己买；
vga：video graphic array

查看配置：
工具：deviceQuery：显存容量、显存带宽、核心频率、核等都决定性能；

GPU 定义了一些概念，用来管理 GPU 内的核、线程等；主要概念有 3 个：Grid > Block > Thread
概念：

• Grid：每个 Grid 中包含一堆 Block，gridDim 定义每个 grid 中 block 的个数；
  • gridDim 被设定为三维：gridDim.x * gridDim.y * gridDim.z，一般情况下二维计算 gridDim.z = 1
• Block：每个 Block 中包含一堆 Thread，blockDim 定义每个 block 中包含的 Thread 的个数；
  • blockDim 被设定为三维： blockDim.x * blockDim.y * blockDim.z，一般情况下二维计算 blockDim.z = 1
  • blockIdx.x、blockIdx.y、blockIdx.z 可获取当前线程块在网格中的索引；
• Thread：底层 thread，最小单位
  • threadIdx.x、threadIdx.y、threadIdx.z 可获取当前线程在其所在 Block 中的索引；

简单 case：

1. 线程数计算

假设有一个二维的网格，其中 gridDim.x = 16 和 gridDim.y = 8，
而每个线程块的大小是 blockDim.x = 16 和 blockDim.y = 16。
那么总共16 * 8 = 128个线程块，并且每个线程块包含 16 * 16 = 256个线程。
因此，整个网格将包含128 * 256 = 32768个线程。

2. 线程索引

int threadId = threadIdx.x + blockDim.x * threadIdx.y; // 在线程块内的线程ID
int blockId = blockIdx.x + gridDim.x * blockIdx.y;     // 在网格内的线程块ID

// 如果需要一个全局唯一的线程ID，可以结合线程块ID和线程ID
int globalThreadId = threadId + blockId * (blockDim.x * blockDim.y);

3. 线程和计算

假设 计算 A * B = C；A_heigh, A_width, B_height, B_width(A_width == B_height);
C 的维度就等于 A_height * B_width， 
那就至少需要 A_height * B_width 个线程来算，
那就至少需要 (A_height * B_width - 1) / thread_count + 1 个 block，
每个 block 有 thread_count 的 thread

三、Cuda 编程

在 CUDA 编程中, __host__、__device__和__global__是三个重要的限定符，它们用于指定函数在哪里执行（CPU或GPU）以及如何执行。

1. 宏修饰函数

1. host

• 定义：用于定义在主机（CPU）上执行的函数。
• 应用：在 CPU 上运行的普通 C/C++函数，通常用于处理GPU的初始化、内存分配、数据传输等任务。

2. device

• 定义：用于定义在 GPU 上执行的函数，但只能从 GPU 上的其他函数（例如__global__函数或其他__device__函数）调用。
• 应用：实现并行操作中的小函数，可以在多个线程中调用以实现更细粒度的并行化。

3. global

• 定义：代码里有__global__修饰的函数是核函数，必须运行在 GPU 上。
• 应用：实现主要的并行操作，通过启动大量线程来执行。每个线程都会执行一次__global__函数
  其运行时参数就是放在三重尖括号<<< >>>之中的值，运行参数写法：

1. kernelFunc<<<b, t>>> (arg1, arg2); 
	b：部署需要的 block 数；
	t：每个 block 的线程数；
2. 三位向量写法：暂略

双重修饰符：host ，__device__可以来定义一个既可以在主机上执行又可以在设备上执行的函数。

函数类型：__global__函数的返回类型必须是void，因为它们不能返回任何值到主机代码。而__host__和__device__函数则可以有返回值。

内存管理：由于 GPU 和 CPU 有各自的内存空间，所以在CUDA编程中，需要考虑如何有效地管理这两者的内存，包括内存分配、数据传输和内存释放等。

2. 宏修饰变量

__device__、__shared__与__constant__也是 cuda 的宏，用于修饰变量。三种变量都不会被声明在 CPU 中，而是在 GPU 中。

1. device
   这个修饰符用于声明一个变量，该变量位于 GPU 的内存（显存）中，可用于全局通信。这样的变量可以在多个线程中访问，但每次访问的都是同一个内存位置。使用__device__修饰的变量在 GPU 上分配空间，并且它们的数据生命周期与整个 CUDA 程序相同。__host__函数无法直接访问__device__变量，__device__函数和__global__函数可以直接访问它们，只需要注意不要线程写冲突。

2. shared
   这个修饰符用于在核函数（kernel）内部声明一个共享内存变量。共享内存是GPU中每个线程块（block）私有的、高速的存储区域。在一个核函数中，每个线程块都会获得这个变量的一个副本，但线程块内的所有线程都可以访问这个副本。快内共享；变量声明时不能初始化，但可以对它进行赋值；

3. constant
   这个修饰符用于声明一个常量，该常量位于设备（GPU）的全局内存中，并且在 kernel 运行期间不能被修改。然而，在主机代码中，可以使用cudaMemcpyToSymbol() 函数来修改这个常量的值。但通常不推荐这样做。它在__device__函数和__global__函数中的访问权限是只读的，这样它就可以被放在高速缓存中，极大地提升访问效率。声明方法和C/C++不同：声明时赋初值是无效的，必须在__host__函数中通过 cudaMemcpyToSymbol() 函数传递给它；当然，__host__函数内部也可以用 cudaMemcpyFromSymbol() 函数获取到它的值。

而如果变量前面没有修饰，那就是寄存器变量（就像C/C++里的寄存器变量），如果是在__device__函数或者__global__函数内，那么每个线程分别持有一个该变量，不会共享，对其读取和修改也只会发生在该线程内。

不过要注意，__device__和__constant__只能声明在全局变量区域，__shared__变量只能声明在核函数内部，类的成员变量和其他函数内的局部变量是无法被上述关键词修饰的。

3. 显存操作接口

host 函数/cpu函数中对显存申请释放初始化拷贝等：

cudaError_t cudaMalloc(void **devPtr, size_t size);
cudaError_t cudaFree(void *devPtr);
cudaError_t cudaMemset(void *devPtr, int value, size_t count);
cudaError_t cudaMemcpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind);
cudaError_t cudaMemcpyToSymbol(const void *symbol, const void *src, size_t count, size_t offset = 0, enum cudaMemcpyKind kind = cudaMemcpyHostToDevice);
cudaError_t cudaMemcpyFromSymbol(void *dst, const void *symbol, size_t count, size_t offset = 0, enum cudaMemcpyKind kind = cudaMemcpyDeviceToHost);

cudaMemcpy 重点区别一下：

enum __device_builtin__ cudaMemcpyKind
{
   cudaMemcpyHostToHost          =   0,      /**< Host   -> Host */
   cudaMemcpyHostToDevice        =   1,      /**< Host   -> Device */
   cudaMemcpyDeviceToHost        =   2,      /**< Device -> Host */
   cudaMemcpyDeviceToDevice      =   3,      /**< Device -> Device */
   cudaMemcpyDefault             =   4       /**< Direction of the transfer is inferred from the pointer values. Requires unified virtual addressing */
};

• cudaMemcpyHostToHost 就是从主机端拷贝到主机端，即此时 cudaMemcp y等价于 memcpy，不属于I/O，耗时最短；
• cudaMemcpyHostToDevice 则是从主机端传送到设备端，即源数据在内存中，目标指针指向了一段显存范围，属于I/O，消耗时间较长；
• cudaMemcpyDeviceToHost 则是从设备端传送到主机端，即源数据在显存中，目标指针指向了一段内存范围，同样属于I/O，消耗时间较长；
• cudaMemcpyDeviceToDevice 则是从设备端拷贝到设备端，CPU只给显卡发送一个信号，不涉及数据交互，因此不属于I/O，不会消耗太多时间。大多数情况下可以异步执行。

而 cudaMemcpyToSymbol() 和cudaMemcpyFromSymbol() 两个函数，前文也提到了，是用来初始化__device__显存全局变量和__constant__显存常量的。虽然函数有五个变量，但后两个变量我们一般只用其初始值，所以写法通常为：

__constant__ int arr[N];

__host__ void init() {
	int a[N] = {9, 8, 7, 6, 5, 4, 3, 2, 1, 0};
	int b[N];
	assert(cudaSuccess == cudaMemcpyToSymbol(arr, a, sizeof(a)));
	assert(cudaSuccess == cudaMemcpyFromSymbol(b, arr, sizeof(b)));
}

设备(GPU）端（global 和 device 函数）就直接用 malloc、free、memset 和 memcpy 即可。