ECE408 Lecture 18 Atomic Operations and Histogramming

这次回顾ECE408 Lecture 18，这次介绍了原子操作和直方图。

课程主页：

https://wiki.illinois.edu/wiki/display/ECE408

搬运视频：

https://www.youtube.com/playlist?list=PL6RdenZrxrw-UKfRL5smPfFFpeqwN3Dsz

目标

理解原子操作
- 并行计算中的读取-修改-写入；
- 并行程序中“临界区”的原始形式；
- 在CUDA中使用原子操作；
- 为什么原子操作会降低内存系统的吞吐量；
- 如何避免某些并行算法中的原子操作;
直方图作为原子操作的示例应用
- 基本的直方图算法；
- 私有化；

一种常见模式

多个客户预订机票；
每个客户：
- 调出飞行座位图；
- 决定座位；
- 更新座位图，将座位标记为已占用；
糟糕的结果：多名乘客最终预订了同一个座位；

Read-Modify-Write操作

如果Mem[x]最初为0，那么在线程1和2完成后Mem[x]的值是多少？

每个线程在它们的Old变量中得到什么？

答案可能因数据竞争而有所不同。为了避免数据竞争，我们需要使用原子操作。

分析

场景1：

Thread 1 Old = 0;
Thread 2 Old = 0;
Mem[x] = 2;

场景2：

Thread 1 Old = 0;
Thread 2 Old = 0;
Mem[x] = 2;

场景3：

Thread 1 Old = 0;
Thread 2 Old = 0;
Mem[x] = 1;

场景4：

Thread 1 Old = 0;
Thread 2 Old = 0;
Mem[x] = 1;

原子操作——确保正确的结果

一般的原子操作

通常由处理器指令集中的单个指令在内存位置地址上执行；
- 读取旧值，计算新值，并将新值写入该位置；
硬件确保在原子操作完成之前没有其他线程可以访问该位置；
- 访问该位置的任何其他线程通常会停止或保留在队列中，直到轮到它；
- 所有线程串行执行原子操作；

CUDA的原子操作

被翻译成单条指令的函数调用（也叫intrinsics ）
- 原子add、sub、inc、dec、min、max、exch（交换）、CAS（比较和交换）；
- 阅读CUDA C编程指南了解更多详情；
原子添加int atomicAdd(int* address, int val);读取全局或共享内存中地址指向的32位字old，计算(old + val)，并将结果存储回同一地址的内存。该函数返回old。

CUDA中更多的原子加法

Unsigned 32-bit: unsigned int atomicAdd(unsigned int* address, unsigned int val);
Unsigned 64-bit: unsigned long long int atomicAdd(unsigned long long int* address, unsigned long long int val);
单精度浮点数加法: float atomicAdd(float* address, float val);

与原子建立同步

我们如何为block构建__syncthreads()；
我们如何为整个grid创建__syncthreads()？
- 为什么这不是一个好主意？
我们如何创建临界区？即，每个block的每个线程执行一段特定的代码？
我们如何为每个grid创建一个临界区？
- 为什么这与grid的__syncthreads()没有相同的问题？

原子Compare and Swap

int atomicCAS(int *address, int compare, int val)
{
    int old = *address;
    if (old == compare)
    	*address = val;
    return old;
}

直方图

一种从大型数据集中提取显着特征和模式的方法；
- 用于图像中对象识别的特征提取；
- 信用卡交易中的欺诈检测；
- 在天体物理学中关联恒星物体运动；
基本直方图——对于数据集中的每个元素，使用值来确定要递增的“bin”；

直方图示例

在句子“Programming Massively Parallel Processors”中构建每个字母的频率直方图
A(4)、C(1)、E(1)、G(1)；
如何并行执行此操作？

更好的方法

合并读取
- 以strided模式为每个线程分配输入；
- 相邻的线程处理相邻的输入字母；

所有线程移动到输入的下一部分：

Kernel

__global__
void histo_kernel(unsigned char *buffer,int size, unsigned int *histo)
{
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    
    // stride is total number of threads
    int stride = blockDim.x * gridDim.x;
    
    while (i < size) {
        atomicAdd( &(histo[buffer[i]]), 1);
        i += stride;
    }
}