【GPU高性能编程CUDA实战】第5章线程协作

5.2 并行线程块的分解

add<<<N, 1>>>(dev_a, dev_b, dev_c);

第一个参数：启动的线程块数量。
第二个参数：每个线程块中创建的线程数量。

5.2.1 矢量求和：重新回顾

1. 使用线程实现GPU上的矢量求和

代码：add_loop_blocks.cu

// 启动N个线程块，每个线程块对应一个线程
// add<<<N, 1>>>(dev_a, dev_b, dev_c);

// 改为启动N个线程，所有线程都在一个线程块内。
add<<<1, N>>>(dev_a, dev_b, dev_c);

// 线程块索引，改为线程索引
int tid = blockIdx.x;
// -->>
int tid = threadIdx.x;

####　2. 在GPU上对更长的矢量求和
（1）线程块数量限制：65535
（2）线程最量：maxThreadsPerBlock

• blockDim 线程块中每一维的线程数量。三维的线程数组
• gridDim 线程格中每一维的线程格数量。二维的线程块

3. 在GPU上对任意长度的矢量求和

代码：add_loop_long_blocks.cu

5.3 共享内存和同步

__share__       使用这个变量驻留在共享内存中。

对于GPU上启动的每一个线程块，CUDA C编译器将创建该 变量的一个副本，线程块中的每个线程都共享这块内存，，但线程却无法看到也不能修改其他线程块的变量副本。

5.3.1 点积运算

代码：dot.cu

__syncthreads(); // 确保线程块中的每个线程都执行完__syncthreads()前面的语句后，才会执行下一条语句。

5.3.3 基于共享内存的位图

代码：shared_bitmap.cu