MindSpore算子开发——CUDA代码优化

CUDA编程入门：谭升的博客

对于cuda编程，核心是如何高效率地利用多线程，每个线程完成一个小任务，最终实现完成一个算子任务。

一个核函数只能有一个grid，一个grid可以有很多个块，每个块可以有很多的线程。

每个线程通过线程号唯一标识，线程号地获得通过：

    int64_t ourOutputPoint = threadIdx.x + blockIdx.x * blockDim.x;
    int64_t plane = blockIdx.y;
    int64_t batch = blockIdx.z;

因为一个grid里面的block是三维的，对于NCHW类型的Tensor，通常pytorch会把N维放在blockIdx.z，C维放在blockIdx.y，从而高效地利用线程。

声明grid和block的时候，有可能因为shape太大导致声明的线程块不足，导致报错，于是要设置一个最大block和threads数：(代码来自pytorch)

    dim3 grid(
      (H*W + 127) / 128, // ceil(outputPlaneSize / 128)
      C,
      N);
    dim3 block(H*W > 128 ? 128 : H*W);

有可能因为元素太多（亿万级），导致线程数不足以完成如此多任务，于是可以让每个线程完成两个及以上的任务：（代码来自mindspore）

template <>
__global__ void Fractionalmaxpool3dwithfixedksize) {
  for (size_t pos = blockIdx.x * blockDim.x + threadIdx.x; pos < outer_size; pos += blockDim.x * gridDim.x) {
    单个线程的任务
  }
  return;
}

在计算某些算子时，由于GPU的地址空间不会初始化为0，所以可能需要手动初始化：（代码来自mindspore）

template <typename T>
__global__ void InitOutput(T *output, const int64_t outer_size) {
  T zero = 0;
  for (size_t id = blockIdx.x * blockDim.x + threadIdx.x; id < outer_size; id += blockDim.x * gridDim.x) {
    output[id] = zero;
  }
  return;
}

在多个线程争夺同一个地址时，善用原子加。