CUDA编程 基础与实践 学习笔记（二）

环境安装过程略

2.1-2.2
NVCC是CUDA的编译器驱动，将C++代码给C++编译器，例如g++,自己编译剩下的部分。
GPU只是一个device,需要gpu来调度。CUDA源文件拓展名为.cu 一个典型的CUDA代码如下：

CUDA核函数两个特点：

1. 被限定词__global__来修饰
2. 必须返回void
   这两个位置可以混，但一般还是__global__开头。
   
   在上述程序中，
3. <<<grid, block>>>是标准写法，grid规定了block的数量，而block规定了线程的数量。一个GPU有很多核心，例如Tesla V100有5120个核心，可以支持很多线程。一维的时候可以直接这样定义，多维时需要使用dim3这种结构体来定义，(x, y, z)。这三个维度并不是独立的，相当于x是个位，y是十位，z是百位的关系。
4. 可以使用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$, 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为次版本号。

可以一次指定编译多个架构的可执行程序，这里不是重点，先略过。