如何在数据自动化系统中有效地打包位元？

Question

我有一个字节数组，每个字节要么是0，要么是1。现在我想把这些值打包成比特，这样8个原始字节占据了一个目标字节，原始字节0变成了位0，字节1变成了位1，等等。到目前为止，内核中有以下内容：

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

这是正确和有效的吗？

Answer 1

__ballot()的扭曲投票功能在这方面非常有用。假设您可以将pOutput重新定义为uint32_t类型，并且您的块大小是翘曲尺寸(32)的倍数：

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

严格地说，if条件甚至是不必要的，因为翘曲的所有线程都会将相同的数据写入相同的地址。所以一个高度优化的版本

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

Answer 2

对于每个线程有两个位，使用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;

您必须测试这是否仍然比您的传统解决方案更快。