如何在CUDA中打包(有效)？

Question

我有一个字节数组,其中每个字节都是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);
}

这是正确有效的吗？

Answer 1

该__ballot()经投票功能来此非常方便.假设您可以重新定义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]);