算法设计：直方图统计
直方图频数统计，也可以看成一个字典Hash计数。用处不是很多，但是涉及CUDA核心操作：全局内存、共享内存、原子函数。
1.1 基本串行算法
这只是一个C语言练习题。
#define MAXN 1005 #define u32 unsigned int __host__ void count(char *hist_data, u32 *bin_data) { for (u32 i = 0; i < MAXN; i++) bin_data[hist_data[i]]++; }
1.2 基于数据分解的并行算法
1.2.1 多线程访存冲突
__global__ void gpu_count1(char *hist_data, u32 *bin_data) { u32 x = blockDim.x*blockIdx.x + threadIdx.x; u32 y = blockDim.y*blockIdx.y + threadIdx.y; u32 tid = x + y*blockDim.x*gridDim.x; /*这是错的*/ bin_data[hist_data[tid]]++; }
多线程情况下，大量相同的hist_data[tid]对bin_data的同一位置同时Read。
结果就是，只有第一个Read是成功的，后续总线周期全部请求失败。
1.2.2 原子函数
原子函数是CUDA默认提供的一些基本函数，包含：
☻算术运算：atomicAdd、atomicSub
☻比较运算：atomicMax、atomicMin
☻位运算：atomicAnd、atomicOr、atomicXor
原子函数为访存提供了傻瓜式的自动阻塞功能。
在相同位置上的并行冲突访问，会被阻塞分解为串行访问。
如上述错误的统计代码应该改成：
atmoicAdd(&bin_data[hist_data[tid]], 1);
1.2.3 性能分析
上述代码使用的是全局内存，也就是GPU的片外显存。一块标准GTX卡，带宽速度为100GBs。
但是上述代码的处理速度仅有1GBs，缩水了100倍。
主要问题也很明显，atomic为了避开访存冲突，将大规模并行退化至大规模串行。GPU利用率很低。
访存冲突域：整个显存。
假设有7个线程块，每个线程块中的线程在bin_data[0]上访存冲突20次，那么阻塞出的串行队列长度为140。
1.3 基于模型分解的并行算法
1.3.1 共享内存
Shared Memory是CUDA中最特殊的一类存储体，有两大特性：
☻线程块内所有线程共享
☻每个存储体与一级Cache级联映射，Cache速度大概是存储体的10倍
共享内存的块内共享机制，意味着你开了256的数组，且有5个线程块，那么会创建5个大小为256的副本数组。
每个副本只在块内使用。仍然隶属于片外显存，速度仍然受制于显存带宽。
同CPU一样，GPU每个SM阵列都有一个64KB的一级Cache。Cache带宽约1.5TBs。
不同的是，CPU中全体内存与Cache相连，GPU中只有共享内存与Cache相连，全局内存无权进入Cache。
Cache的好处就是访存的 ”时间局部性" 原理：如果一个信息项正在被访问，那么在近期它很可能还会被再次访问。
这正是访存冲突的另一个角度，如果将冲突域的一部分转为共享内存，那么不仅不会减速，反而会得到Cache的加速。
1.3.2 降解冲突域
__shared__ u32 cache[256]; __global__ void gpu_count2(char *hist_data, u32 *bin_data) { u32 x = blockDim.x*blockIdx.x + threadIdx.x; u32 y = blockDim.y*blockIdx.y + threadIdx.y; u32 tid = x + y*blockDim.x*gridDim.x; char val = hist_data[tid]; cache[threadIdx.x] = 0; __syncthreads(); atomicAdd(&cache[val], 1); __syncthreads(); atomicAdd(&bin_data[threadIdx.x], cache[threadIdx.x]); }
代码的重点是 __syncthreads() ，这是个让块内线程同步的函数：
跑的快的线程在断点处被锁住，等待全部线程执行完毕后，再跳转到下一行代码。
线程锁是多线程必备武器，参照一个笑话：
前苏联某官员去视察植树造林的情况，现场他看到一个人在远处挖坑，其后不远另一个人在把刚挖出的坑逐个填上。
上面这个笑话如果发生在程序中就是线程调度的问题，种树这个任务有三个线程：挖坑线程，种树线程和填坑线程。
后面的线程必须等前一个线程完成才能进行，而不是按时间顺序来进行，否则一旦一个线程出错就会出现上面荒谬的结果。
用线程锁来处理两个线程先后执行的情况在程序中，和种树一样，很多任务也必须以确定的先后秩序执行。
--------------------------------------------------------------------------------------------------------
上述代码，为每个线程块开了一块共享内存，假若按照1.2.3那样假设：7个线程块，bin_data[0]上冲突20次。
由于atomicAdd(&cache[val], 1)仅仅作用于自己的块内，所以7个线程块，最长冲突队列长度=20
而下面atomicAdd(&bin_data[threadIdx.x], cache[threadIdx.x])仅仅是7个线程块拼凑，最长冲突队列长度=7
详细参照图示：
1.3.3 平衡线程块个数与线程块内计算压力
1.3.2中代码，线程块中每个线程仅仅负责统计一个值，如果减少线程块数，而增加单线程处理量：
#define THREAD 256 #define N 5 __global__ void gpu_count2(char *hist_data, u32 *bin_data) { u32 x = blockDim.x*blockIdx.x + threadIdx.x; u32 y = blockDim.y*blockIdx.y + threadIdx.y; u32 tid = x + y*blockDim.x*gridDim.x; cache[threadIdx.x] = 0; __syncthreads(); for (u32 i = 0,offset=0; i < N; i ++,offset+=THREAD) { char val = hist_data[tid+offset]; atomicAdd(&cache[val], 1); } __syncthreads(); atomicAdd(&bin_data[threadIdx.x], cache[threadIdx.x]); }
增大N，会增加在共享内存上的冲突，而减少在全局内存上的冲突，获得加速。
N增大一定情况后，加速衰减直至0，遇到I/O瓶颈。这是CUDA最无奈的地方：
