CUDA 笔记

贡献者： addis

1. 笔记 1

• Nvidia 的一个官方入门 ppt cuda-c-basics.pdf
• 一本很好的入门教材是 CUDA by examples.
• CUDA 的编译器叫做 nvcc, 用法与 g++ 大致相同。
• 在函数定义前面加 __global__ 就可以定义为在 GPU 中执行的函数（device code）. 在 CPU 中执行的叫做 host code.
• host code 中调用 device code 函数用 fun<<<m,n>>>().
• 指针也分为 device pointer 和 host pointer. 前者不能在 host code 中 dereference, 后者不能在 device code 中 dereference, 但两者都可在 device 和 host 间传递。
• memory management: cudaMalloc(); 用于在 GPU 中分配空间
• nvprof ./<exename> 可用于测试程序的用时情况（profiling）

2. 笔记 2

• 所有有 CUDA 语法的源文件都用 .cu 后缀，否则会报错（包括 include 的头文件中的 cuda 语法）!
• cudaGetDeviceCount(int *deviceCount) 可以返回 GPU 的数量。
• cudaSetDevice(int device) 用于切换 GPU, 其中 device 是 GPU 的编号，从 0 开始。默认的 GPU 编号是 0.
• cudaGetDeviceProperties(struct cudaDeviceProp * prop, int device) 可以获取 GPU 信息。prop 中常用的 member 如下： prop.name 是 GPU 型号（名称）, prop.major 是 capability 的编号，每个编号是一种构架，3: Kepler, 5: Maxwell, 6: Pascal, 7: Volta, prop.minor 是 capability 的小数点后的编号，prop.multiProcessorCount 是 multiprocessor 的数量，总核数等于该数量乘以每个 mp 中的核数（与构架有关）, prop.deviceOverlap 表示 GPU 是否有同时执行 kernel 和传数据的能力，只有这个能力存在，用多个 stream 才有可能加速。 prop.totalGlobalMem 和 prop.totalConstMem 分别是 GPU 内存的大小和 const 空间的大小。prop.canMapHostMemory 决定 GPU 是否能使用 zero-copy memory. 剩下的见 cudaDeviceProp 的定义。
• kernel function 如果在 runtime 时出错了，应该是会终止而不会与任何报错的，所以在 nvprof 里面应该就看不到 kernel 了。但如果能看到，应该就是执行成功了。
• 许多 cuda*() 函数都会返回错误类型 cudaError_t, cudaGetErrorString(cudaError_t err) 可以返回错误的描述。
• cudaDeviceSynchronize() 函数在使 host 程序等待所有 device 完成任务。
• __syncthreads() 函数在 kernel 中的所有 thread 同步。
• pascal 6.0: 64 cores/Multiprocessor, 6.1: 128 cores/Multiprocessor. volta 7.0: 64 cores/Multiprocessor
• pascal 构架的 shared memory 是 48kb/block (6144 个 double), volta 构架是 96kb/block. 如果超出，编译的时候会报错。
• 在 global scope 声明的 const (不是 __constant__) 变量既可以在 host code 中也可以在 device code 中使用。
• 在 global scope 声明的变量前面加 __device__ , 则该 global 变量只能用于 device code, 不能在 host 中使用，但可以在 host 中用 cudaMemcpyToSymbol() 对其赋值。如果改用 __cnstant__, 则只能在 device code 中读取该变量，但同样可以在 host 中用 cudaMemcpyToSymbol(). __constant__ 空间一般只有 64kb, 但有可能将速度变快（见 cuda by examples）.
• __device__, __constant__, 支持特定的 class object, 这个 class 必须要有一个空的 constructor 或者没有 constructor (默认).
• 如果 class 的成员函数需要在 kernel 中执行，前面要加 __device__, 如果要在 kernel 和 host 中执行，前面要加 __host__ __device__.
• kernel 里面是可以用 new 和 delete 的，也可用 malloc() 和 free() ! 使用的是 device 空间的储存，在一块专门的区域叫做 heap. 如果不够的话，new 会返回 nullptr.（貌似不一定！）
• heap size 默认是 8M, 但也可以手动设置大小。

  cudaDeviceGetLimit(size_t* size, cudaLimitMallocHeapSize)  // 获取 heap size
  cudaDeviceSetLimit(cudaLimitMallocHeapSize, size_t size) // 设置 heap size
  

• 如果每个线程都用 new，貌似消耗的 heap 空间会随 block 数量不断增长，而不是与 cuda core 的数量成正比。而且 heap 不够的时候不会报错，而是整个 kernel 都不执行。
• 在 Tesla V100 上，用 cudaMalloc() 代替 new, 速度要快许多， 然而 Quadro P5000 反而会慢，但是慢的不多。
• 要在 cuda 中使用 complex, 用 #include "cuda_complex.h" , 见我的 CUDATest/template.
• 执行一个 parallel for 类似的问题时，最好是安排每个线程执行 for(i=ind; i<N; i+=Nblock*Nthread)，而千万不要企图用大量的 block 使得每个线程只用执行一个 i.
• 如果每个线程都是完全独立的话，貌似在 Tesla V100 上用 <<<2*Ncore/32,32>>> 所需的时间最短，总线程数是 2*Ncore = 10240. (这个结论是用 cn3Dz 得出的).
• 用 Even 给 GPU 计时：

  cudaEven_t start, stop;
  cudaEventCreate(&start); cudaEventCreate(&stop);
  cudaEventRecord(start, 0);
  // do some work on the GPU
  cudaEventRecord(stop, 0);
  cudaEventSynchronize(stop); // let CPU wait for the event to finish
  cudaEventElapsedtime(&elapsedTime, start, stop); // time unit is "ms"
  cudaEventDestroy(start); cudaEventDestroy(stop);
  

  3. 不同的 Memory

• malloc() 分配 pageable host memory, 而 cudaHostAlloc(void** pointer, size_t size, cudaHostAllocDefault) 分配 page-locked host memory (或 pinned memory). 后者不会被 swap 到硬盘中。要释放后者分配的内存，用 cudaFreeHost(void * pointer) 即可。page-locked 内存在 CPU 和 GPU 间的传输速度大概是 pageable 内存的两倍。
• 如果用 cudaHostAlloc(void** pointer, size_t size, cudaHostAllocMapped), 分配出来的内存就是 zero-copy memory, 因为这种内存可以直接在 kernel code 中读取。
• 要使用 zero-copy memory, 开始时需要在 host 程序中用 cudaSetDeviceFlags(cudaDeviceMapHost).
• 当 zero-copy memory 内存只用于 GPU 时，选项 cudaHostAllocMapped | cudaHostAllocWriteCombined 可以提升性能。
• cudaHostGetDevicePointer(void **pDevice, void *pHost, 0) 函数用于将 zero-copy memory 从 Host 中的指针得到 Device 的指针。
• 调用 kernel 后，zero-copy memory 会变成 undefined. 使用 cudaThreadSynchronize() 即可恢复。

4. Unified Memory

• cudaMallocManaged(void **pointer, size_t size);
• cudaFree(void *pointer);
• 每次调用完 kernel 以后需要 cudaDeviceSynchronize(), 否则 CPU 无法读取 Unified Memory.
• 至于多线程多 GPU 同时运行的情况，还有待研究。

5. cuda-gdb

• 基本用法与 gdb 一样 cuda-gdb name.x
• 编译的时候，给 nvcc 加上 -g 和 -G. 其中 -g 是 CPU 调试，-G 是 GPU 调试。
• 在 kernel 中，切换不同的线程用 cuda block <number> tread <number> 或者 cuda block (1,2,3) tread (4,5,6) 指定三维线程。
• 也可以单独用 cuda block <number> 和 cuda thread <number> .
• 要显示当前线程编号，用 cuda block thread
• 显示当前 device 信息用 info cuda devices
• 显示正在运行的 kernel 用 info cuda kernels

6. 多 GPU

• 首先创建 cuda 工程（也可以用 CUDATest/template）.
• 为了在 CUDA 中使用 openMP, 在 Project Property > CUDA C/C++ > Command Line > Additional Options 中加入 -Xcompiler "/openmp" 即可。然后插入工程中唯一一个 ".cu" 源代码。这个代码是 CUDA by Examples 中 11 章最后一个例程的简化版。这个代码很有希望可以改成多 GPU 的 cuCn3D 工程。

  #ifdef _MSC_VER
  #include "cuda_runtime.h"
  #include "device_launch_parameters.h"
  #endif
  #include <iostream>
  using namespace std;
  __global__
  void add1(int *x, int thread) { x[thread] += 1; }
  
  int main()
  {
  	int i, N = 3;
  	int *x, *dev_x;
  	cudaSetDevice(0);
  	cudaSetDeviceFlags(cudaDeviceMapHost);
  	cudaHostAlloc((void**)&x, N*sizeof(int),
         cudaHostAllocWriteCombined | cudaHostAllocPortable | 
         cudaHostAllocMapped);
  	for (i = 0; i < N; ++i)
  		x[i] = i;
  	cudaHostGetDevicePointer(&dev_x, x, 0);
  
  	#pragma omp parallel for
  	for (i = 0; i < N; ++i) {
  		if (i != 0) {
  			cudaSetDevice(i);
  			cudaSetDeviceFlags(cudaDeviceMapHost);
  		}
  		add1<<<1,1>>>(dev_x, i);
  		cudaThreadSynchronize();
  	}
  	cout << "the result is: " << x[0] << " " << x[1]
          << " " << x[2] << endl;
  }