CUDA并行笔记

内存共享
1、 多个 thread 可以组成一个 block，block 中的 thread 可并行执行且可存取同一块共享的显存。但要注意的是每个 block 中的 thread 是有数量限制的。
2、多个 block 可以组成一个 grid，但 grid 中的 block 无法共享显存( 只能共享些别的信息 )，因此可合作度并不高。
3、查看计算机共享内存大小的代码

int dev=0;
cudaDeviceProp devProp;
cudaGetDeviceProperties(&devProp,dev);
cout<<devProp.sharedMemPerBlock;  //单位是字节B
cout<<devProp.maxThreadsPerBlock;  //每个block的最大thread数量

内存结构图

CUDA 程序执行模式
1. CUDA 程序利用并行化来替代内存 cache，即一个 thread 需要等待内存则 GPU 会切换到另一个 thread 执行。
2. CUDA 程序对于 “分支预处理” 的实现也是采用和 1 类似的方式。
块内同步函数：__syncthreads ()
线程调用此函数后，该线程所属块中的所有线程均运行到这个调用点后才会继续往下运行。
“device_functions.h” 头文件包含 __syncthreads ()函数

声明共享内存 (数组)

__ shared__ type name[size]; 

__shared__内存空间说明符（可选地与__device__一起使用）声明了一个变量：
‣驻留在线程块的共享内存空间中，
‣具有该块的生命周期，
‣每个块具有一个不同的对象，
‣仅仅可被一个块中的所有线程访问，
‣没有常量地址。
在共享内存中将变量声明为外部数组时，例如

 extern __shared__ float shared [];

数组的大小在启动时确定。

使用共享内存的步骤：

1. 声明共享内存(使用__shared__)
2. 多个线程各自处理，将结果存到共享内存上
3. 同步函数同步、等待（使用__syncthreads ()）
4. 选其中一个内存单元（通常为第一个进行累加）
5. 将共享内存中求得的和赋给传入的变量
   For the parallel workloads, at points in the algorithm where parallelism is broken because some threads need to synchronize in order to share data with each other, there are two cases: Either these threads belong to the same block, in which case they should use __syncthreads() and share data through shared memory within the same kernel invocation, or they belong to different blocks, in which case they must share data through global memory using two separate kernel invocations, one for writing to and one for reading from global memory. The second case is much less optimal since it adds the overhead of extra kernel invocations and global memory traffic. Its occurrence should therefore be minimized by mapping the algorithm to the CUDA programming model in such a way that the computations that require inter-thread communication are performed within a single thread block as much as possible.

参考文章：https://www.cnblogs.com/1024incn/p/4605502.html