[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
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);
}
原文地址：https://blog.csdn.net/mingshili/article/details/143685349
免责声明：本站文章内容转载自网络资源，如本站内容侵犯了原著者的合法权益，可联系本站删除。更多内容请关注自学内容网（zxcms.com）！