#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);
}
|