CUDA 学习 2

• 开发小技巧
• • 调试
  • 性能
  • 错误处理
• 函数说明
• • 核函数
  • 内置变量

开发小技巧

调试

当进行调试的时候可以把核函数设置成单线程：

kernelName<<<1,1>>>(argument list)

性能

调整网格和线程块大小可以得到不同的性能。

错误处理

建议在CUDA开发时加上错误信息提示处理，提高排错效率，在release版本可以去除这部分，例如添加下宏：

#define CHECK(call)\
{\const cudaError_t error=call;\if(error!=cudaSuccess)\{\printf("ERROR: %s:%d,",__FILE__,__LINE__);\printf("code:%d,reason:%s\n",error,cudaGetErrorString(error));\exit(1);\}\
}//使用示例
CHECK(cudaMalloc((float**)&a, size));

函数说明

核函数

kernelName<<>>(argument list)

• Dg：int型或者dim3类型(x,y,z)，用于定义一个Grid中Block是如何组织的，如果是int型，则表示一维组织结构
• Db：int型或者dim3类型(x,y,z)，用于定义一个Block中Thread是如何组织的，如果是int型，则表示一维组织结构
• Ns：size_t类型，可缺省，默认为0； 用于设置每个block除了静态分配的共享内存外，最多能动态分配的共享内存大小，单位为byte。 0表示不需要动态分配。
• S：cudaStream_t类型，可缺省，默认为0。 表示该核函数位于哪个流。

//一个Grid中有3x2x1=6个Block，在(x,y,z)三个方向上的排布方式分别是3、2、1
//一个Block中有4x3x1=12个Thread，在(x,y,z)三个方向上的排布方式分别是4、3、1
dim3 grid(3,2,1), block(4,3,1);
kernelName<<>>(...);//一个Grid中有5个Block，在(x,y,z)三个方向上的排布方式分别是5、1、1
//一个Block中有8个Thread，在(x,y,z)三个方向上的排布方式分别是8、1、1
kernelName<<<5,8>>>(...);

注：grid(3,2)表示第一维度有3个索引值，第二维度有2个索引值，即2行3列

内置变量

• threadIdx.[x, y, z]表示Block内Thread的编号
• blockIdx.[x, y, z]表示Gird内Block的编号
• blockDim.[x, y, z]表示Block的维度，也就是Block中每个方向上的Thread的数目
• gridDim.[x, y, z]表示Gird的维度，也就是Grid中每个方向上Block的数目

//各维度索引计算公式//一维Grid 一维Block
int blockId = blockIdx.x; 
int threadId = blockIdx.x *blockDim.x + threadIdx.x;//一维Grid 二维Block
int blockId = blockIdx.x; 
int threadId = blockIdx.x * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx.x;//一维Grid 三维Block
int blockId = blockIdx.x; 
int threadId = blockIdx.x * blockDim.x * blockDim.y * blockDim.z + threadIdx.z * blockDim.y * blockDim.x + threadIdx.y * blockDim.x + threadIdx.x;//二维Grid 一维Block
int blockId = blockIdx.y * gridDim.x + blockIdx.x;  
int threadId = blockId * blockDim.x + threadIdx.x;//二维Grid 二维Block
int blockId = blockIdx.x + blockIdx.y * gridDim.x;  
int threadId = blockId * (blockDim.x * blockDim.y) + (threadIdx.y * blockDim.x) + threadIdx.x;  //二维Grid 三维Block
int blockId = blockIdx.x + blockIdx.y * gridDim.x;  
int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)  + (threadIdx.z * (blockDim.x * blockDim.y))  + (threadIdx.y * blockDim.x) + threadIdx.x;  //三维Grid 一维Block
int blockId = blockIdx.x + blockIdx.y * gridDim.x + gridDim.x * gridDim.y * blockIdx.z;  
int threadId = blockId * blockDim.x + threadIdx.x;  //三维Grid 二维Block
int blockId = blockIdx.x + blockIdx.y * gridDim.x + gridDim.x * gridDim.y * blockIdx.z;  
int threadId = blockId * (blockDim.x * blockDim.y) + (threadIdx.y * blockDim.x) + threadIdx.x;  //三维Grid 三维Block
int blockId = blockIdx.x + blockIdx.y * gridDim.x + gridDim.x * gridDim.y * blockIdx.z;  
int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)  + (threadIdx.z * (blockDim.x * blockDim.y))  + (threadIdx.y * blockDim.x) + threadIdx.x;//二维
int blockId = blockIdx.x + blockId.y * gridDim.x;
int threadId = blockId * (blockDim.x * blockDim.y) + (threadIdx.y *blockDim.x) + threadIdx.x;//三维
int blockId = blockIdx.x + blockIdx.y * gridDim.x + gridDim.x * gridDim.y * blockIdx.z;
int threadIc = blockId * (blockDim.x * blockDim.y * blockDim.z) + (threadIdx.z * (blockDim.x * blockDim.y)) + (threadIdx.y * blockDim.x) + threadIdx.x;