《大规模并行处理器编程实战》前十章

▶ 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

▶ 有关线程和线程块的层次与坐标计算

Device ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━┓
含多个 Grid，一次核函数调用看作一个 Grid

    Grid 1 ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━至多 65535 个 Block，至多三维
        指定方法：
        一维可用 unsigned int a;
        一般可用 dim3 blocksize(x) 或 dim3 blocksize(x, y) 或 dim3 blocksize(x, y, z);
        或直接在核函数调用时使用 dim3(x) 或 dim3(x, y) 或 dim3(x, y, z);

        下标范围
        一维：blockIdx.x，取值范围 [0, gridDim.x - 1]
        二维：blockIdx.x 和 blockIdx.y，取值范围 [0, gridDim.x - 1], [0 ~gridDim.y - 1]
        三维：blockIdx.x 和 blockIdx.y 和 blockIdx.z，取值范围 [0, gridDim.x - 1], [0, gridDim.y - 1], [0, gridDim.z - 1]

        Block 1 ━━━━━━━━━━━━━━━━至多1024个Thread，至多三维
            指定方法：
            一维可用 unsigned int a;
            一般可用 dim3 threadsize(x) 或 dim3 threadsize(x, y) 或 dim3 threadsize(x, y, z);
            或直接在核函数调用时使用 dim3(x) 或 dim3(x, y) 或 dim3(x, y, z);
            
            下标范围
            一维：threadIdx.x，取值范围 [0, blockDim.x - 1]
            二维；threadIdx.x 和 threadIdx.y，取值范围 [0, blockDim.x - 1], [0, blockDim.y - 1]
            三维；threadIdx.x 和 threadIdx.y 和 threadIdx.z，取值范围 [0, blockDim.x - 1] 和 [0, blockDim.y - 1], [0, blockDim.z - 1]

            Thread 1 ━━━━━━━━按照block、thread编号进行偏移：
                一维：
                    id = blockIdx.x * blockDim.x + threadIdx.x;
                跳转：
                    id += gridDim.x * blockDim.x;
                    在同一个 thread 中跨步吞吐更多的索引，跨步等于一次核函数调用时所有线程块的线程总数

                二维：
                    idx = blockIdx.x * blockDim.x + threadIdx.x;
                    idy = blockIdx.y * blockDim.y + threadIdx.y;
                    id = idy * gridDim.x * blockDim.x + idx;
                    行 (x) 优先存储，列 (y) 在高位上（y相邻的元素存储位置不相邻）
                跳转：
                    id += gridDim.x * gridDim.y * blockDim.x * blockDim.y;
                
                三维：
                    idx = blockIdx.x * blockDim.x + threadIdx.x;
                    idy = blockIdx.y * blockDim.y + threadIdx.y;
                    idz = blockIdx.z * blockDim.z + threadIdx.z; 
                    id = (idz * gridDim.y * blockDim.y + idy) * gridDim.x * blockDim.x + idx;
                    层 (z) 在最高位上
                跳转：
                    id += gridDim.x * gridDim.y * gridDim.z * blockDim.x * blockDim.y * blockDim.z

            Thread 2 ━━━━━━━━

            ...

        Block 2 ━━━━━━━━━━━━━━━━

        ...

    Grid 2 ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━

    ...

Device ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━┛

 ● 数组遍历测试

#include <stdio.h>
#include <time.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"

__global__ void print()
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int idy = blockIdx.y * blockDim.y + threadIdx.y;
    int idz = blockIdx.z * blockDim.z + threadIdx.z;
    int id  = idz * gridDim.x * blockDim.x * gridDim.y * blockDim.y + idy * gridDim.x * blockDim.x + idx;
    printf("gridDim %2d-%2d-%2d, block %2d-%2d-%2d, "
            "blockDim %2d-%2d-%2d, thread %2d-%2d-%2d, "
            "x=%2d,y=%2d,z=%2d,total %2d
",
        gridDim.x, gridDim.y, gridDim.z, blockIdx.x, blockIdx.y, blockIdx.z,
        blockDim.x, blockDim.y, blockDim.z, threadIdx.x, threadIdx.y, threadIdx.z,
        idx, idy, idz, id);
    return;
}

int main()
{
    print << <dim3(2, 3, 4), dim3(2, 3, 4) >> >();
    cudaDeviceSynchronize();

    getchar();
    return 0;
}

▶ P62。Fortran和Matlab采用的是是列优先的存储方式，在与C程序交互的过程中要注意转置。

▶ P112。访问全局内存时，若同一个Warp中所有线程都执行同一条指令来访问全局内存中的连续地址，访问效率最高。

▶ P116。当数据放入共享内存中，不管是以行方式还是列方式访问都不会造成性能差异。

▶ P120。线程块并行和线程并行的关系。线程调度为细粒度，效率高；线程块并行为粗粒度，每次调度都要重新分配资源。在采用分治法解决问题时，先将大规模的问题分解为几个小规模问题，分别用线程块实现，线程块内再将任务细化为线程并行。线程块之间粗粒度，块内细粒度，可以充分利用硬件资源，降低线程并行的计算复杂度。应尽量安排每一个线程完成较多工作，并采用更少的线程数。

▶ P128。浮点数相关。

● 负数的补码是其绝对值数按位取反再加 1。

● 余码的好处是无符号比较器可以用来比较有符号数，无符号比较器比较小，速度快。

● 单精度浮点数，1位符号位，8位阶码，23为位尾数；双精度浮点数，1为符号位，11为阶码，52位尾数。

● NaN 有两种。SignalingNaN 作为运算操作的输入会引发异常，中断执行，通常用于标记未初始化的数据；quietNaN 可以作为运算操作的输入，结果仍为 quietNaN，不会中断执行。

● 由于计算舍入造成的误差，称为 0.5_D ULP (Unit in the Last Place) 误差。

● 串行 / 并行算法在精确性上并没有绝对优势（可以分别举出精确性串行高于并行，或并行高于串行的例子）。应根据需要决定是否接受算法和计算结果。选择算法时可以先对数据进行排序、分组，这对串行和并行程序都能提高精确性。