如何在CUDA中高效地打包比特？

我有一个字节数组，其中每个字节是0或1.现在我想打包这些值的位，使8个原始字节占据1个目标字节，原始字节0进入位0，字节1进入位1，等到目前为止，我在内核中的以下内容：如何在CUDA中高效地打包比特？

const uint16_t tid = threadIdx.x; 
__shared__ uint8_t packing[cBlockSize]; 

// ... Computation of the original bytes in packing[tid] 
__syncthreads(); 

if ((tid & 4) == 0) 
{ 
    packing[tid] |= packing[tid | 4] << 4; 
} 
if ((tid & 6) == 0) 
{ 
    packing[tid] |= packing[tid | 2] << 2; 
} 
if ((tid & 7) == 0) 
{ 
    pOutput[(tid + blockDim.x*blockIdx.x)>>3] = packing[tid] | (packing[tid | 1] << 1); 
}

这是正确和有效的？

来源

2016-09-14 Serge Rogatch

这不能工作。这是一场记忆赛。在CUDA – talonmies

@talonmies中没有并行比特大小的事务，我认为没有竞争，因为处理同一字节的线程属于同一个warp。 –

在同一个warp中不能保证安全性，没有两个线程可以同时修改同一个字节而不会导致竞争 – talonmies

__ballot() warp-voting函数为此非常方便。假定可以重新定义pOutput是uint32_t类型的，并且您的块大小是warp大小的倍数（32）：

unsigned int target = __ballot(packing[tid]); 
if (tid % warpSize == 0) { 
    pOutput[(tid + blockDim.x*blockIdx.x)/warpSize] = target; 
}

严格地说，如果条件甚至没有必要的，因为所有warp的线程会将相同的数据写入相同的地址。因此，一个高度优化的版本也只是

pOutput[(tid + blockDim.x*blockIdx.x)/warpSize] = __ballot(packing[tid]);

来源

2016-09-14 10:57:20 tera

伟大的解决方案，谢谢。然后我不需要共享内存和'__syncthreads（）'。 –

好的问题的答案。 – harrism

据我所知，这不能扩展到打包2位值？尽管我们可以使用2个'__ballot'调用来获取2个独立的32位变量中的低位和高位，然后对位进行交织（例如，使用http://stackoverflow.com/questions/39490345/interleave-bits-efficiently）在CUDA上比我在问题中给出的算法更昂贵。为了打包2位值，该算法需要在if（（tid＆6）== 0）'内写入输出。 –

对于每个线程两位，使用uint2 *pOutput

int lane = tid % warpSize; 
uint2 target; 
target.x = __ballot(__shfl(packing[tid], lane/2)    & (lane & 1) + 1)); 
target.y = __ballot(__shfl(packing[tid], lane/2 + warpSize/2) & (lane & 1) + 1)); 
pOutput[(tid + blockDim.x*blockIdx.x)/warpSize] = target;

你必须标杆，这是否仍是比你的传统解决方案快。

来源

2016-09-15 15:12:13 tera

如何在CUDA中高效地打包比特？

回答

相关问题