Как упаковать биты (эффективно) в CUDA?

Question

Как упаковать биты (эффективно) в 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);
}

Это правильно и эффективно?

2

bit-packing c++cuda parallel-processing

Решение

Другие решения

Для двух бит на поток, используя 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;

Вам нужно будет оценить, является ли это все еще быстрее, чем ваше обычное решение.

1

Источник

Accepted Answer

__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]);

8