环境安装过程略
2.1-2.2
NVCC是CUDA的编译器驱动,将C++代码给C++编译器,例如g++,自己编译剩下的部分。
GPU只是一个device,需要gpu来调度。CUDA源文件拓展名为.cu 一个典型的CUDA代码如下:
CUDA核函数两个特点:
- 被限定词__global__来修饰
- 必须返回void
这两个位置可以混,但一般还是__global__开头。
在上述程序中, - <<<grid, block>>>是标准写法,grid规定了block的数量,而block规定了线程的数量。一个GPU有很多核心,例如Tesla V100有5120个核心,可以支持很多线程。一维的时候可以直接这样定义,多维时需要使用dim3这种结构体来定义,(x, y, z)。这三个维度并不是独立的,相当于x是个位,y是十位,z是百位的关系。
- 可以使用printf,但不支持cout. 亲测cout无效。
疑问:没太搞懂这个cudaDeviceSynchronize()的作用
2.3 线程组织
一个疑问:在我目前的认知下,一个小任务会分配到一个线程,而总的线程数为grid*block。那么线程又如何分配到实际的计算核心上呢?
简单的维度和索引的规定。GridDim和blockDim分别规定了grid中有多少个block,一个block有多少个thread. 而blockIdx和threadIdx则是该线程的索引,取值范围在0到各自的dim-1。线程块之间的计算是独立的,完成顺序无关。
多维网络构造语法:
dim3 grid(gx, gy);
dim3 block(bx, by);
gpu<<<grid, block>>>();
多维网络本质上还是一维的,
第一种是一维顺序索引,第二种复合索引,第二种更符合矩阵的运算。
线程束(warp):一个内建变量,所有GPU架构的warpsize都是32, 相当于在block内部增加了一个维度。
线程定义的数量限制:grid(
2
31
−
1
2^{31}-1
231−1, 65535,65535), block(1024,1024,32).一个block最多只能有1024个线程。切记。
nvcc编译cuda代码时,会自动将代码分为主机代码和device代码,主机代码完全支持C++,host代码部分支持C++。nvcc先将设备代码编译为PTX(parallel thread execution)伪汇编代码,这一步需要通过-arch=compute_xx来指定一个虚拟架构,以确保cuda功能。再将其编译为二进制的cubin目标代码,需要使用-code=sm_xx来指定一个真实架构。虚拟架构版本号<=真实架构版本号。 例如:
其中,compute_35中的3为主版本号,5为次版本号。
可以一次指定编译多个架构的可执行程序,这里不是重点,先略过。