CUDA程序设计(二)

算法设计：直方图统计

直方图频数统计，也可以看成一个字典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卡，带宽速度为100GB\s。

但是上述代码的处理速度仅有1GB\s，缩水了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.5TB\s。

不同的是，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最无奈的地方：