[CUDA] cuda kernel开发记录

1. kernel基本书写

# 基本步骤
分配host内存，并进行数据初始化；
分配device内存，并从host将数据拷贝到device上；
调用CUDA的 kernel 函数在device上完成指定的运算；
将device上的运算结果拷贝到host上；
释放device和host上分配的内存。

2. grid-block设置

// 1维时，可以直接使用int来表示
const int block_size = 128;
const int grid_size = (size + block_size - 1 ) / block_size;

// 多维时，可以使用dim3数据类型
dim3 grid_dim1, block_dim1(32, 32);
grid_dim1.x = (kNTotal + 32 - 1) / 32;
grid_dim1.y = (kMTotal + 32 - 1) / 32;

3. device 使用

• device一些struct用法

template<typename T>
struct NonZeroOp
{
    __host__ __device__ __forceinline__ bool operator()(const T& a) const {
      return (a!=T(0));
    }
};

• device属性设置

extern __attribute__((device)) __attribute__((cudart_builtin)) cudaError_t cudaMalloc(void **devPtr, size_t size);
等同于
extern __host__ __device__ cudaError_t cudaMalloc(void **devPtr, size_t size);

4. launch_bounds

• launch_bounds的使用

5. kernel问题排查

• kernel中invalid argument错误，一般是kernel配置的参数问题，这个需要确定grid size, block size是否为0； 如果为零则会出invalid的错误。
• 如果是一些stream爆出错误，则考虑是否stream上的kernel有问题，需要通过每个kernel调用后加入cudaGetLastError或者cudaPeekAtLastError() 来确定是哪个函数。
• 所以以后写kernel函数，最好在调用函数后面加上cudaPeekAtLastError() 保证kernel出错能及时报出问题。

CUDA_CHECK(cudaPeekAtLastError()); // 不会清理错误flag状态。
CUDA_CHECK(cudaGetLastError()); // 会清理错误flag状态。

6. CUDA_KERNEL_LOOP的使用

6.1 基本写法

• 一般写kernel函数时，最好多使用CUDA_KERNEL_LOOP

• 注意__global__ void kernel中 的N不能是引用

// template <typename T> \
// __global__ void ##name_kernel(T* buf, const int N) { \
//     int tid = threadIdx.x + blockIdx.x * blockDim.x; \
//     buf[tid] = op(buf[tid]); \
// }

• 注意基本写法index 通过循环来，这样保证一个block的thread读取的连续数据

// 利用这种宏来保证kernel数量小于处理数据的数量时，也能处理全数据。
#define CUDA_KERNEL_LOOP(i, n)                                 \
  for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < (n); \
       i += blockDim.x * gridDim.x)

6.2 使用注意事项

• 但写loop cuda kernel的时候，需要注意最好kernel的个数按照数据赋值的地址的数量进行设置，这样防止地址访问冲突。
• 但是如果kernel loop的过程中，存在两个或两个以上的kernel会访问一个地址，尤其累加或累乘操作，需要注意用cuda提供的原子操作，防止多个kernel对同一个地址同时写，从而导致结果不正确的问题。

template <typename scalar_t>
__global__ void devoxelize_forward_kernel2(int c, int N,
                                          const int *__restrict__ indices,
                                          const scalar_t *__restrict__ weight,
                                          const scalar_t *__restrict__ feat,
                                          scalar_t *__restrict__ out)
{
  // index is for indices or weights
  CUDA_KERNEL_LOOP(index, N*c) {
	  int i = index / N;
	  int j = index % N;
	
	  if (i < 8) {
	    const int indices_ = *(indices + index);
	    const scalar_t weight_ = *(weight + index);
	    const scalar_t *feat_ = feat + indices_ * c;
	
	    scalar_t cur_feat;
	    for(int k = 0; k < c; k++) {
	      cur_feat = 0;
	      if (indices_ >= 0) cur_feat = feat_[k];
	      // before: out[j * c + k] += weight_ * cur_feat;
	      // fix the bug, conflict.
	      atomicAdd(out + j * c + k, weight_ * cur_feat);  
	    }
	  }
  }
}

7. kernel中打印GPU数据

• 当debug cuda kernel的时候 打印kernel中一些关键值的变化很重要，对排查问题很有帮助，但是cuda kernel只能用printf打印，注意打印float的时候，要小数点多一些，因为有效非零值会小数点后几位才有值。
• 另外gpu上的数据只能用kernel进行封装printf的方式打印； 另一种方法就是将gpu数据copy到cpu后，再打印。

template <typename Type>
__global__ void PrintKernel(const Type* data, int start, int end) {
  for (int i = start; i < end; ++i) {
    if (std::is_floating_point<Type>::value) {
      printf("| %.7f ", static_cast<float>(data[i]));
    } else {
      printf("| %.1f ", static_cast<float>(data[i]));
    }
  }
  printf("\n");
}
template <typename Type>
void Print(const Type* data, int start, int end) {
  PrintGpuDataKernel<Type><<<1, 1, 0>>>(data, start, end);
}