算法设计:直方图统计

直方图频数统计,也可以看成一个字典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

详细参照图示:

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程序设计(二)的更多相关文章

  1. CUDA程序设计(一)

    为什么需要GPU 几年前我启动并主导了一个项目,当时还在谷歌,这个项目叫谷歌大脑.该项目利用谷歌的计算基础设施来构建神经网络. 规模大概比之前的神经网络扩大了一百倍,我们的方法是用约一千台电脑.这确实 ...

  2. CUDA程序设计(三)

    算法设计:基数排序 CUDA程序里应当尽量避免递归,因而在迭代排序算法里,基数排序通常作为首选. 1.1 串行算法实现 十进制位的基数排序需要考虑数位对齐问题,比较麻烦.通常实现的是二进制位的基数排序 ...

  3. JavaScript高级程序设计(二):在HTML中使用JavaScript

    一.使用<script>元素 1.<script>元素定义了6个属性: async:可选.表示应该立即下载脚本,但不应该妨碍页面中的其他操作,比如下载其他资源或等待加载其他脚本 ...

  4. python基础16 ----面向对象程序设计二

    一.继承与派生 1.继承的定义:继承是一种创建新类的方式,即在类中提取共同的部分创建出一个类,这样的类称为父类,也可称为基类和超类,新建的类称为派生类或子类. 2.单继承:就相当于子类继承了一个父类. ...

  5. javascript 高级程序设计 二

    这里我们直接进入主题: 在JS刚刚开始的时候,必须面临一个问题,那就是如何使的JS的加载和执行不会影响web核心语言HTML的展示效果,和HTML和谐共存. 在这个背景下<script>标 ...

  6. 实验7 shell程序设计二(1)

    编写一个shell过程完成如下功能(必须在脚本中使用函数)1.程序接收3个参数:$1/$2和$3,合并两个文件$1/$2为$3,并显示,三个文件均为文本文件.2.如果文件$3不存在,那么先报告缺少$3 ...

  7. cuda测试二维block的使用

    #include "cuda_runtime.h" #include <stdio.h> #include <stdlib.h> #include < ...

  8. Linux命令(十一)——Shell程序设计二(循环控制语句)

    1.if语句 (1)两路分支的if语句 (2)多路条件判断分支的if语句 2.测试语句 (1)文件测试 (2)字符串测试 (3)数值测试 (4)用逻辑操作符进行组合的测试语句 3.case语句 4.f ...

  9. 周会材料:高并发程序设计&lt;二&gt;

    第三章 JDK并发包https://www.cnblogs.com/sean-zeng/p/11957569.html JDK内部提供了大量实用的API和框架.本章主要介绍这些JDK内部功能,主要分为 ...

随机推荐

  1. VC++ error C2248: &ldquo;CObject::CObject&rdquo;: 无法访问 private 成员(在&ldquo;CObject&rdquo;类中声明)

    在使用诸如:CArray或是 CList等类时,经常会出现此错误 此错误的原因是由于自定义的类的数组项时 有一个操作如  Add()  在这个操作中,实际上需要一个 = 操作,但是这个 =操作在 自定 ...

  2. 一个null引发的错误

    写程序时,发生一个unrecognized selector的错误,很显然,这是一个经典错误,运行时找不到可调用的方法. 系统提示的reason是 -[NSNull length] 错误 经过查找,原 ...

  3. 【BZOJ】1180: [CROATIAN2009]OTOCI &amp; 2843: 极地旅行社(lct)

    http://www.lydsy.com/JudgeOnline/problem.php?id=1180 今天状态怎么这么不好..................................... ...

  4. python学习之字符串

    最近在学习python,随手做些记录,方便以后回顾 #字符串是不可再改变的序列aa='abcd'#aa[2:]='ff' #报错,不可直接赋值#字符串格式化:使用格式化操作符即百分号%来实现print ...

  5. 将JSON对象转化为数组对象

    package web.helper; import java.util.ArrayList; import net.sf.json.JSONArray; import web.model.Abstr ...

  6. [LeetCode] Split Concatenated Strings 分割串联字符串

    Given a list of strings, you could concatenate these strings together into a loop, where for each st ...

  7. kubernetes进阶之一:简单例子

    kubernetes 从一个简单例子开始 参考 <kubernetes 权威指南>一节的 从一个简单例子开始,操作实录. 一.Java Web 应用结构 二.启动MySql服务 1.首先为 ...

  8. Android开发 MVP模式的规范记录(个人总结)

    前言 首先,这篇文章不在讲解什么是mvp模式,如果需要请自行搜索mvp模式文章了解.这个文章里我只记录mvp模式的创建和mvp各自层的界限.另外这个博客属于个人使用mvp模式后一些经验总结与记录.并不 ...

  9. Liferay7 BPM门户开发之25: Liferay7应用程序配置(APPLICATION CONFIGURATION)

    首先有几个概念需要明确.1.第一个概念是这里的应用程序配置不是写XML之类的配置文件,是类似字典的类型化配置这意味着应用程序配置不只是一个字符串键值对的列表.值还可以有类型,如整数列表,字符串列表,一 ...

  10. ceph存储osd启动异常处理和正常启停操作

    机器角色:cloudstack虚拟机的宿主机:ceph存储机器. 事件:ceph存储的物理机器由于内存异常,需要停机更换,仅仅是把该物理机上面的虚拟机迁移走,同时启动了停机维护,然后就直接关机.结果造 ...