▶ P46。SPMD (Single-Program Multiple-Data) 单程序多数据，CUDA使用的并行编程风格。并行处理单元在数据的多个部分执行相同程序，但这些处理单元不用同时执行限购通的指令；SIMD (Single-Instruction Multiple-Data) 单指令多数据，SM 对 warps 使用的调度方式，所有并行处理单元在任何时候都执行相同的指令。
▶ P87。内存变量类型：
| 变量类型 | 前缀名 | 声明位置 | 存储器 | 作用域 | 生命周期 | 说明 |
| 全局变量 | (__global__) | 核函数外 | 全局存储器 | 网格 | 程序 | 被整个Grid和CPU使用，访问速度很慢 |
| 数组以外的自动变量 |
(__local__) | 核函数内 | 寄存器 | 线程 | 核函数调用 | 若从全局内存中产生副本，则速度很慢 |
| 数组自动变量 | (__local__) | 核函数内 | 局部存储器 | 线程 | 核函数调用 | |
| 共享内存变量 | __shared__ | 核函数内、外均可 | 线程块 | 核函数调用 | 线程块共享，被同一线程快中的所有线程同时读写 | |
| 常量内存变量 | __constant__ | 核函数外 | 常数存储器 | 网格 | 程序 | 从CPU中赋值后GPU可读不可写 |
常量内存加速原理：半线程束广播（临近内存集合整体读取）；不存在缓存一致性问题，可以放入 L1 缓存；GPU缓存易命中。
运行速度比较，__constant__ > __local__ > __share__ >>__global__ >> host memory
▶ 有关线程和线程块的层次与坐标计算
1 Device ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━┓ 2 含多个 Grid，一次核函数调用看作一个 Grid 3 4 Grid 1 ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━至多 65535 个 Block，至多三维 5 指定方法： 6 一维可用 unsigned int a; 7 一般可用 dim3 blocksize(x) 或 dim3 blocksize(x, y) 或 dim3 blocksize(x, y, z); 8 或直接在核函数调用时使用 dim3(x) 或 dim3(x, y) 或 dim3(x, y, z); 9 10 下标范围 11 一维：blockIdx.x，取值范围 [0, gridDim.x - 1] 12 二维：blockIdx.x 和 blockIdx.y，取值范围 [0, gridDim.x - 1], [0 ~gridDim.y - 1] 13 三维：blockIdx.x 和 blockIdx.y 和 blockIdx.z，取值范围 [0, gridDim.x - 1], [0, gridDim.y - 1], [0, gridDim.z - 1] 14 15 Block 1 ━━━━━━━━━━━━━━━━至多1024个Thread，至多三维 16 指定方法： 17 一维可用 unsigned int a; 18 一般可用 dim3 threadsize(x) 或 dim3 threadsize(x, y) 或 dim3 threadsize(x, y, z); 19 或直接在核函数调用时使用 dim3(x) 或 dim3(x, y) 或 dim3(x, y, z); 20 21 下标范围 22 一维：threadIdx.x，取值范围 [0, blockDim.x - 1] 23 二维；threadIdx.x 和 threadIdx.y，取值范围 [0, blockDim.x - 1], [0, blockDim.y - 1] 24 三维；threadIdx.x 和 threadIdx.y 和 threadIdx.z，取值范围 [0, blockDim.x - 1] 和 [0, blockDim.y - 1], [0, blockDim.z - 1] 25 26 Thread 1 ━━━━━━━━按照block、thread编号进行偏移： 27 一维： 28 id = blockIdx.x * blockDim.x + threadIdx.x; 29 跳转： 30 id += gridDim.x * blockDim.x; 31 在同一个 thread 中跨步吞吐更多的索引，跨步等于一次核函数调用时所有线程块的线程总数 32 33 二维： 34 idx = blockIdx.x * blockDim.x + threadIdx.x; 35 idy = blockIdx.y * blockDim.y + threadIdx.y; 36 id = idy * gridDim.x * blockDim.x + idx; 37 行 (x) 优先存储，列 (y) 在高位上（y相邻的元素存储位置不相邻） 38 跳转： 39 id += gridDim.x * gridDim.y * blockDim.x * blockDim.y; 40 41 三维： 42 idx = blockIdx.x * blockDim.x + threadIdx.x; 43 idy = blockIdx.y * blockDim.y + threadIdx.y; 44 idz = blockIdx.z * blockDim.z + threadIdx.z; 45 id = (idz * gridDim.y * blockDim.y + idy) * gridDim.x * blockDim.x + idx; 46 层 (z) 在最高位上 47 跳转： 48 id += gridDim.x * gridDim.y * gridDim.z * blockDim.x * blockDim.y * blockDim.z 49 50 Thread 2 ━━━━━━━━ 51 52 ... 53 54 Block 2 ━━━━━━━━━━━━━━━━ 55 56 ... 57 58 Grid 2 ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ 59 60 ... 61 62 Device ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━┛
● 数组遍历测试
1 #include <stdio.h> 2 #include <time.h> 3 #include "cuda_runtime.h" 4 #include "device_launch_parameters.h" 5 6 __global__ void print() 7 { 8 int idx = blockIdx.x * blockDim.x + threadIdx.x; 9 int idy = blockIdx.y * blockDim.y + threadIdx.y; 10 int idz = blockIdx.z * blockDim.z + threadIdx.z; 11 int id = idz * gridDim.x * blockDim.x * gridDim.y * blockDim.y + idy * gridDim.x * blockDim.x + idx; 12 printf("gridDim %2d-%2d-%2d, block %2d-%2d-%2d, " 13 "blockDim %2d-%2d-%2d, thread %2d-%2d-%2d, " 14 "x=%2d,y=%2d,z=%2d,total %2d ", 15 gridDim.x, gridDim.y, gridDim.z, blockIdx.x, blockIdx.y, blockIdx.z, 16 blockDim.x, blockDim.y, blockDim.z, threadIdx.x, threadIdx.y, threadIdx.z, 17 idx, idy, idz, id); 18 return; 19 } 20 21 int main() 22 { 23 print << <dim3(2, 3, 4), dim3(2, 3, 4) >> >(); 24 cudaDeviceSynchronize(); 25 26 getchar(); 27 return 0; 28 }
▶ P62。Fortran和Matlab采用的是是列优先的存储方式，在与C程序交互的过程中要注意转置。
▶ P112。访问全局内存时，若同一个Warp中所有线程都执行同一条指令来访问全局内存中的连续地址，访问效率最高。
▶ P116。当数据放入共享内存中，不管是以行方式还是列方式访问都不会造成性能差异。
▶ P120。线程块并行和线程并行的关系。线程调度为细粒度，效率高；线程块并行为粗粒度，每次调度都要重新分配资源。在采用分治法解决问题时，先将大规模的问题分解为几个小规模问题，分别用线程块实现，线程块内再将任务细化为线程并行。线程块之间粗粒度，块内细粒度，可以充分利用硬件资源，降低线程并行的计算复杂度。应尽量安排每一个线程完成较多工作，并采用更少的线程数。
▶ P128。浮点数相关。
● 负数的补码是其绝对值数按位取反再加 1。
● 余码的好处是无符号比较器可以用来比较有符号数，无符号比较器比较小，速度快。
● 单精度浮点数，1位符号位，8位阶码，23为位尾数；双精度浮点数，1为符号位，11为阶码，52位尾数。
● NaN 有两种。SignalingNaN 作为运算操作的输入会引发异常，中断执行，通常用于标记未初始化的数据；quietNaN 可以作为运算操作的输入，结果仍为 quietNaN，不会中断执行。
● 由于计算舍入造成的误差，称为 0.5D ULP (Unit in the Last Place) 误差。
● 串行 / 并行算法在精确性上并没有绝对优势（可以分别举出精确性串行高于并行，或并行高于串行的例子）。应根据需要决定是否接受算法和计算结果。选择算法时可以先对数据进行排序、分组，这对串行和并行程序都能提高精确性。