SDK例子中的histogram256

加速计算 CUDA
1

#define USE_SMEM_ATOMICS 0
#if(!USE_SMEM_ATOMICS)
#define TAG_MASK ( (1U << (UINT_BITS - LOG2_WARP_SIZE)) - 1U )
inline device void addByte(volatile uint *s_WarpHist, uint data, uint threadTag)
{
uint count;
do
{
count = s_WarpHist[data] & TAG_MASK; 说到底就是不明白这个函数addByte，请版主指教一下。
count = threadTag | (count + 1); 详细代码在C:\ProgramData\NVIDIA Corporation\CUDA Samples\v5.5
s_WarpHist[data] = count; \3_Imaging\histogram的histogram256.cu
}
while (s_WarpHist[data] != count);
}

inline device void addWord(uint *s_WarpHist, uint data, uint tag)
{
addByte(s_WarpHist, (data >> 0) & 0xFFU, tag); 这个函数的目的是用来分类，分成255类,当一个warp中有相同的类
addByte(s_WarpHist, (data >> 8) & 0xFFU, tag); 时，会写在相同的位置，会引起bank冲突，应该是为了解决这个问
addByte(s_WarpHist, (data >> 16) & 0xFFU, tag); 题。
addByte(s_WarpHist, (data >> 24) & 0xFFU, tag);
}

2

前几天已经详细回答了你的问题了，请不要重复发帖。

3

先谢谢版主了，那个程序是分64类的，已经明白了，这一次不是上次的问题，而是分256类时，一个warp内可能有相同的类，写向相同的地址，而形成bank冲突，下面这3句代码应该是为了解决冲突的。
但看得不怎么明白。帮请教一下版主。
count = s_WarpHist[data] & TAG_MASK;
count = threadTag | (count + 1);
s_WarpHist[data] = count

4

楼主您好，深表歉意，昨日错误的以为您是在重复发帖了。

我来说下这个，这个可以实现在一个warp内，不使用内建的原子操作而进行安全的原子性写入的一个过程。

该过程其实很容易理解，
（1）warp内的多个线程可能存在竞争的写入一个地址。
（2）根据Warp的行为规则，你知道只有1个线程会成功。
（3）warp内的每个线程都检测成功的是否是本人。
（4）如果不是，重新尝试写入。

你要注意到，这里的s_WarpHist中的元素实际上分成了2部分，
高位部分是写入者标示，
而低位部分则是普通的计数。

每个线程通过读取高位的成功者的标识，来判断自己是否是竞争胜利的那个。
如果不是，则下一次尝试开始：
（1）抹除上个成功者的痕迹（上个成功者将不再参与竞争了），
（2）将值+1, 并涂抹上自己的痕迹。
（3）竞争性写入
（4）读取回来看看胜利者是否是自己，如果是，退出竞争。如果不是，本次写入没成功，下次竞争开始。

这个实际上和您认为的bank conflict无关。
昨日错误的判断了您的意图，并将贵贴移动到水区，深表歉意。

感谢您的莅临。

5

您需要注意的是，这个利用了warp的2个特性：
（1）warp内的32个线程写入同一个地址，只有1个线程会成功，其他31个线程的值将被硬件丢弃。（是哪个线程是未定义的）
（2）warp内的指令执行是lock-step的，当前执行的指令只有固定的1到2条。warp内的1个线程，要么参与执行这1到2条指令，要么处于不执行的状态，而不会去执行这以外的其他的指令。

该补充可能有助您理解此代码，
感谢您的来访。

代为修正笔误一处。

6

恩，谢谢版主，我明白你说的4点的意思了，但下面第3小点有一些不明白。
（1）s_WarpHist[data] = count; 这一句把计算写入共享内存。
（2）while (s_WarpHist[data] != count); 再检查刚才有没有写成功，因为每个线程写的count不同的

（3）count = s_WarpHist[data] & TAG_MASK; 这里不明白为什么取出计数后要先与TAG_MASK？TAG_MASK这个数字代表什么意思？

（4）count = threadTag | (count + 1); 这里是每个线程在计数时作一个标记。

下面这段代码有一行不明白，估计与上面的TAG_MASK有关，下面这段代码的意思是并行求每列的和
s_Hist的大小是WARP_COUNT=6行乘以HISTOGRAM256_BIN_COUNT=256列，

for (uint bin = threadIdx.x; bin < HISTOGRAM256_BIN_COUNT; bin += HISTOGRAM256_THREADBLOCK_SIZE)
{
uint sum = 0;

for (uint i = 0; i < WARP_COUNT; i++) 这个循环就列求和了，我不明白的地方就是
{ 为什么 s_Hist的内容要与上TAG_MASK才是计数？估计与上面那个有关吧？
sum += s_Hist[bin + i * HISTOGRAM256_BIN_COUNT] & TAG_MASK;
}

d_PartialHistograms[blockIdx.x * HISTOGRAM256_BIN_COUNT + bin] = sum;
}

7

楼主您好，

感谢您的认真阅读。

每个计数值的高位，实际上有5个bit的值, 用来标记是谁写入的。这个值和正常的计数无关，所以要用and操作，进行抹掉。也就是我中文说的，要抹掉上个成功者的信息。这样剩下的是纯粹的计数值，可以+1，并用or操作摸上准备写入的竞争者（之一）的值。

无论是每次竞争者的尝试写入，
还是最终的要获取某warp的计数值们，
高位的写入者的编号信息是无用的，需要摸掉。
所以你也看到的sum += …;那行也有个AND.

感谢您的来访。

8

哦，彻底明白了，原来这里这一句对每一个线程做了一个标记
const uint tag = threadIdx.x << (UINT_BITS - LOG2_WARP_SIZE); 因此要用
define TAG_MASK ( (1U << (UINT_BITS - LOG2_WARP_SIZE)) - 1U )来抹除上个成功者的痕迹，因此
sum += …;那行也有个AND。
谢谢热心的版主们了。

9

感谢您的来访，祝您周末愉快。