用户
 找回密码
 立即注册
91金属狂人 该用户已被删除
发表于 2013-10-29 16:56:40
67938
#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);
}
使用道具 举报 回复
发表于 2013-10-29 18:15:31
前几天已经详细回答了你的问题了,请不要重复发帖。

使用道具 举报 回复 支持 反对
发表于 2013-10-29 20:34:29
横扫千军 发表于 2013-10-29 18:15
前几天已经详细回答了你的问题了,请不要重复发帖。

先谢谢版主了,那个程序是分64类的,已经明白了,这一次不是上次的问题,而是分256类时,一个warp内可能有相同的类,写向相同的地址,而形成bank冲突,下面这3句代码应该是为了解决冲突的。
但看得不怎么明白。帮请教一下版主。
count = s_WarpHist[data] & TAG_MASK;   
        count = threadTag | (count + 1);   
         s_WarpHist[data] = count
使用道具 举报 回复 支持 反对
发表于 2013-10-30 18:51:37
91金属狂人 发表于 2013-10-29 20:34
先谢谢版主了,那个程序是分64类的,已经明白了,这一次不是上次的问题,而是分256类时,一个warp内可能 ...

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

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

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

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

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

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

感谢您的莅临。
使用道具 举报 回复 支持 反对
发表于 2013-10-30 18:54:44
本帖最后由 ice 于 2013-10-30 18:59 编辑
横扫千军 发表于 2013-10-30 18:51
楼主您好,深表歉意,昨日错误的以为您是在重复发帖了。

我来说下这个,这个可以实现在一个warp内,不使 ...

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

该补充可能有助您理解此代码,
感谢您的来访。
--------------------------------------------------------------------------
代为修正笔误一处。
使用道具 举报 回复 支持 反对
发表于 2013-11-1 08:28:11
横扫千军 发表于 2013-10-30 18:51
楼主您好,深表歉意,昨日错误的以为您是在重复发帖了。

我来说下这个,这个可以实现在一个warp内,不使 ...

恩,谢谢版主,我明白你说的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;
    }
使用道具 举报 回复 支持 反对
发表于 2013-11-1 15:04:30
91金属狂人 发表于 2013-11-1 08:28
恩,谢谢版主,我明白你说的4点的意思了,但下面第3小点有一些不明白。   
(1)s_WarpHist[data] = cou ...

楼主您好,

感谢您的认真阅读。

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

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

感谢您的来访。
使用道具 举报 回复 支持 反对
发表于 2013-11-1 20:50:22
玫瑰幻想 发表于 2013-11-1 15:04
楼主您好,

感谢您的认真阅读。

哦,彻底明白了,原来这里这一句对每一个线程做了一个标记
const uint tag = threadIdx.x << (UINT_BITS - LOG2_WARP_SIZE);   因此要用
#define TAG_MASK ( (1U << (UINT_BITS - LOG2_WARP_SIZE)) - 1U )来抹除上个成功者的痕迹,因此
sum += .....;那行也有个AND。
谢谢热心的版主们了。
使用道具 举报 回复 支持 反对
发表于 2013-11-1 20:51:59
91金属狂人 发表于 2013-11-1 20:50
哦,彻底明白了,原来这里这一句对每一个线程做了一个标记
const uint tag = threadIdx.x  ...

感谢您的来访,祝您周末愉快。
使用道具 举报 回复 支持 反对
发新帖
您需要登录后才可以回帖 登录 | 立即注册