CUDA程序设计(二)

算法设计:直方图统计

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 = ; 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.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[];
__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] = ;
__syncthreads();
atomicAdd(&cache[val], );
__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

详细参照图示:

CUDA程序设计(二)

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] = ;
__syncthreads();
for (u32 i = ,offset=; i < N; i ++,offset+=THREAD)
{
char val = hist_data[tid+offset];
atomicAdd(&cache[val], );
}
__syncthreads();
atomicAdd(&bin_data[threadIdx.x], cache[threadIdx.x]);
}

增大N,会增加在共享内存上的冲突,而减少在全局内存上的冲突,获得加速。

N增大一定情况后,加速衰减直至0,遇到I/O瓶颈。这是CUDA最无奈的地方:

CUDA程序设计(二)

上一篇:Couldn't open CUDA library cublas64_80.dll etc. tensorflow-gpu on windows


下一篇:win7注册表损坏的修复方法