昨天试着在GPU上输出了"HELLO WORD!"以及cuda的存储问题,那么今天先开始看第一个:线程管理:
核函数在主机端启动的时候,他的执行会移动到设备上去,这个时候,设备上会产生大量的线程,而且每一个线程都在执行由核函数指定的语句,所以,怎么样去组织线程是CUDA的一个关键部分,CUDA明确了线程层次抽象的概念,以便于我们组织线程,这里我们展示书上的一个两层的线程层次结构:
由一个内核启动所产生的线程统称为一个网格,同意网格中的所有线程共享相同的全局存储空间,一个网格由多个线程块构成,一个线程块内又包含了一组线程,同一线程块内的线程写作可以通过以下方式来实现:
同步
共享内存
但是,不同块内的线程不能够协作!!!
线程依靠两个坐标变量来找:
blockidx:线程块在线程网格内的索引
threadidx:线程块内的线程索引
这两个变量是核函数中需要预先初始化的,当执行一个核函数的时候,CUDA运行时就会给每个线程分配以上两个变量,基于这种思想,我们就可以很轻松的把数据给分配到不同的线程中去,
CUDA是可以组织三维的网格和块的:
blockIdx.x
blockIdx.y
blockIdx.z
threadIdx.x
threadIdx.y
threadIdx.z
他的结构是一个包含二维块的二维网格,网格和快的维度由下面两个内置变量指定:
blockDim:线程快的维度:每个线程块中的线程数
gridDim:线程格的维度,每个线程格中的线程数来表示
和上面一样也是三维的;
通常,一个线程格会被组装成线程块的二维数组形式,一个线程块会被组织成线程的三维数组形式。
CUDA程序中有两种不同的网格和块变量:
手动定义的dim3数据类型(仅主机端可见)
预定义的uint3数据类型(仅设备端可见)
主机端作为内核调用的一部分,可以用dim3来定义一个网格和块的维度,执行核函数的时候,CUDA运行时会自动生成相应的内置预初始化的网格,块和线程变量,在核函数内均可以被访问且都是uint3类型。网格大小是块的倍数,跟着书上的学着写一下:checkDimension.cu文件:
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include
#include
#include
__global__ void checkDimension() {printf("threadindex:(%d,%d,%d) blockindex:(%d,%d,%d) blockdimindex:(%d,%d,%d)""griddim:(%d,%d,%d)\n", threadIdx.x, threadIdx.y, threadIdx.z, blockIdx.x, blockIdx.y, blockIdx.z, blockDim.x, blockDim.y, blockDim.z, gridDim.x, gridDim.y, gridDim.z);}
int main() {int nElem = 6;dim3 block(3);dim3 grid((nElem + block.x - 1) / block.x);printf("grid.x%d,grid.y%d,grid.z%d\n", grid.x, grid.y, grid.z);printf("block.x%d,block.y%d,block.z%d\n", block.x, block.y, block.z);checkDimension << > > ();cudaDeviceReset();return 0;
}
对于一个给定的数据大小,确认网格和块尺寸的一般步骤是:
确定块的大小
在已知数据大小和块大小的基础上计算网格维度
要确定块的尺寸,通常要考虑的是
内核的性能特性
GPU资源的限制
上一篇:NeRF in the Wild