第三章 CUDA简介

Posted 2020-10-10 爨爨爨好

? cudaMemcpy()不能用于过GPU系统中不同GPU之间的数据复制。

 

? CUDA在用的是SPMD (Single-Program Multiple-Data)

 

? 函数声明类型：

格式	调用位置	执行位置
__device__ foo()	设备	设备
__global__ foo()	主机	设备
__host__ foo()	主机	主机

 

 

 

 

 

可以同时声明为 __device__ __host__ foo() ，生成两份拷贝，可以在主机或者设备上执行。

 

? 内存变量类型：

变量类型	前缀名	声明位置	作用域	生命周期	说明
全局变量	(__global__)	核函数外	文件作用域	程序	被整个Grid和CPU使用
局部变量	(__local__)	核函数内	函数作用域	核函数调用	若从__global__中逐个产生副本，则速度取决于__global__的速度很慢
共享内存变量	__share__	核函数内、外均可	函数作用域	核函数调用	线程块共享，被同一线程快中的所有线程同时读写
常量内存变量	__constant__	核函数外	文件作用域	程序	从CPU中赋值后GPU可读不可写

 

 

 

 

 

常量内存加速原理：半线程束广播（临近内存集合整体读取）；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(a) 或 dim3 blocksize(a, b) 或 dim3 blocksize(a, b, c);
 8         或直接在核函数调用时使用 dim3(a) 或 dim3(a, b) 或 dim3(a, b, c);
 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(a) 或 dim3 threadsize(a, b) 或 dim3 threadsize(a, b, c);
19             或直接在核函数调用时使用 dim3(a) 或 dim3(a, b) 或 dim3(a, b, c);
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                     不改变blocksize和threadsize的条件下，在同一个thread中跨步吞吐更多的索引，
32                     需要的跨步恰好为一次核函数调用的线程总数
33 
34                 二维：
35                     idx = blockIdx.x * blockDim.x + threadIdx.x;
36                     idy = blockIdx.y * blockDim.y + threadIdx.y;
37                     id = idy * gridDim.x * blockDim.x + idx;
38                     行（x坐标）优先存储，列（y坐标）在高位上（y相邻的元素存储位置不相邻）
39                 跳转：
40                     id += gridDim.x * gridDim.y * blockDim.x * blockDim.y;
41                 
42                 三维：
43                     idx = blockIdx.x * blockDim.x + threadIdx.x;
44                     idy = blockIdx.y * blockDim.y + threadIdx.y;
45                     idz = blockIdx.z * blockDim.z + threadIdx.z; 
46                     id = (idz * gridDim.y * blockDim.y + idy) * gridDim.x * blockDim.x + idx;
47                     层（z坐标）在最高位上
48                 跳转：
49                     id += gridDim.x * gridDim.y * gridDim.z * blockDim.x * blockDim.y * blockDim.z
50 
51             Thread 2 ━━━━━━━━
52 
53             ...
54 
55         Block 2 ━━━━━━━━━━━━━━━━
56 
57         ...
58 
59     Grid 2 ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━
60 
61     ...
62 
63 Device ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━┛

 

? 线程块并行和线程并行的关系
        线程并行为细粒度，调度效率高；线程块并行为粗粒度，每次调度都要重新分配资源。在采用分治法解决问题时，先将大规模的问题分解为几个小规模问题，分别用线程块实现，线程块内在将任务细化为线程并行。线程块之间粗粒度，块内细粒度，可以充分利用硬件资源，降低线程并行的计算复杂度。

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