如何在 CUDA 中(有效地)打包位?

How to pack bits (efficiently) in CUDA?

我有一个字节数组,其中每个字节为 0 或 1。现在我想将这些值打包成位,以便 8 个原始字节占用 1 个目标字节,原始字节 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);
}

这样正确有效吗?

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

严格来说,if 条件甚至不是必需的,因为 warp 的所有线程都会将相同的数据写入相同的地址。所以高度优化的版本就是

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

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

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