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的数量，而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<<>>();
多维网络本质上还是一维的，


第一种是一维顺序索引，第二种复合索引，第二种更符合矩阵的运算。

线程束(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为次版本号。

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