cuda5

直方图

parallel histogram

A simple parallel histogram algorithm

将输入分成几部分
让每个线程获取输入的一部分
每个线程遍历其部分。
对于每个字母，增加适当的 bin 计数器

输入分区影响内存访问效率

• 分段分区导致内存访问效率低下
• 相邻线程不访问相邻内存位置
• 访问未合并
• DRAM 带宽利用率低

解决

更改为交错分区
所有线程处理元素的连续部分
他们都移动到下一部分并重复
内存访问被合并

---

数据竞争

发生在读-修改-写过程中，导致无法预计的错误。

多个线程同时操作一些一样的变量，造成了竞争，出错

使用原子操作可以避免。

原子操作

由单个硬件指令对存储器位置执行读-修改-写操作，硬件确保当前原子操作完成之前没有其他线程可以 完成读-修改-写操作，会维护一个队列。

   int atomicAdd(int* address, int val);
//原子加  addr+val,写回addr

more

Unsigned 32-bit integer atomic add
unsigned int atomicAdd(unsigned int* address,
   unsigned int val);

Unsigned 64-bit integer atomic add
unsigned long long int atomicAdd(unsigned long long
   int* address, unsigned long long int val);

Single-precision floating-point atomic add (capability > 2.0)
float atomicAdd(float* address, float val);

基本直方图内核

__global__ void histo_kernel(unsigned char *buffer,long size, unsigned int *histo) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;

// stride is total number of threads
    int stride = blockDim.x * gridDim.x;

 // All threads handle blockDim.x * gridDim.x
   // consecutive elements
   whileint alphabet_position = buffer[i] – “a”;
      if (alphabet_position >= 0 && alpha_position < 26)
		atomicAdd(&(histo[alphabet_position/4]), 1);
       i += stride;
   }
}

原子操作 DRAM

对 DRAM 位置的原子操作从读取开始，其延迟为数百个周期
原子操作以写入同一位置结束，延迟数百个周期
在这整个过程中，没有其他人可以访问该位置

延迟决定吞吐量

同一 DRAM 位置上原子操作的吞吐量是应用程序可以执行原子操作的速率。
特定位置上的原子操作速率受读取-修改-写入序列的总延迟限制，对于全局存储器 (DRAM) 位置通常超过 1000 个周期。
这意味着，如果许多线程尝试在同一位置（争用）上执行原子操作，则内存吞吐量将减少到 < 一个内存通道峰值带宽的 1/1000！

Fermi L2 缓存 原子操作

中等延迟，约为 DRAM 延迟的 1/10
在所有块之间共享
全局内存原子的“免费改进”

共享内存原子操作

非常短的延迟
每个线程块私有
需要程序员的算法工作（稍后详述）手动编程

---

私有化副本有缺点

创建和初始化需要开销，将私有化写入最终副本需要开销

但访问和串行化代价更小，总体性能提高 10 倍以上

如果直方图太大无法私有化可以部分私有化

私有化的成本和收益

成本

创建和初始化私有副本的开销
将私人副本的内容累积到最终副本的开销

受益

访问私有副本和最终副本时的争用和序列化要少得多
整体性能通常可以提高 10 倍以上

直方图的共享内存原子操作（共享内存需要私有化

每个线程子集都在同一个块中
吞吐量远高于 DRAM (100x) 或 L2 (10x) 原子
减少争用——只有同一块中的线程才能访问共享内存变量
这是共享内存的一个非常重要的用例！

一些其他的东西

私有化是用于并行化应用程序的强大且常用的技术

私有直方图大小需要很小
适合共享内存

如果直方图太大而无法私有化怎么办？
有时可以部分私有化输出直方图并使用范围测试转到全局内存或共享内存

一些数据集在局部区域有大量相同的数据值
一个简单而有效的优化是每个线程在更新直方图的相同元素时将连续更新聚合为单个更新