Обработка массива CUDA как единого объекта

Я бьюсь головой о стену, пытаясь решить эту проблему уже около месяца, и ни мои навыки C, ни мое гугл-фу не были достаточно сильны, чтобы найти решение.

Один из моих любимых сторонних проектов был и остается попыткой найти палиндром для числа 196 с помощью метода «обратить и добавить»:

196 + 691 = 887

887 + 788 = 1675

И так далее до тех пор, пока результат не будет одинаковым.

В последнее время я предпочитал использовать cuda, но я снова и снова сталкивался с одной и той же проблемой. нести.

Поддерживая, я представляю число в памяти как массив беззнаковых char, каждая цифра - это один char - так по сути распакован bcd.

__device__ __align__(4) unsigned char DigitArray[1024 * 1024];

Генерация частичной суммы легко выполняется параллельно. Я сохраняю текущее количество цифр в памяти устройства, а затем на каждой итерации:

__device__ int DigitCount;

__global__ void PartialSums()
{
    int idx = GlobalThreadIndex();
    int rev = DigitCount - (1 +  idx);
    unsigned char sum = DigitArray[idx];
    __threadfence();
    if(rev >= 0)
    { 
        sum += DigitArray[rev];
    }
    DigitArray[idx] = sum;
}

Все очень хорошо и хорошо.

Теперь о переносе.
В идеальном мире я бы хотел, чтобы произошло следующее:

__device__ unsigned int SumScratch[1024*256];

__global__ void Carry()
{
    int idx = GlobalThreadIndex();
    SumScratch[idx] = 0xF6F6F6F6;
    __threadfence();
    unsigned int * ptr = (unsigned int *)(DigitArray + (idx * size of(unsigned char));
    SumScratch[idx] += *ptr;
    __threadfence();
    unsigned int cMask = __vcmples(SumScratch[idx], 0x0A0A0A0A);
    unsigned int nCMask = ~cMask;
    *ptr = __vadd4((SumScratch[idx] & cMask), __vsub4((SumScratch[idx] & nCMask), (OxF6F6F6F6 & nCMask)) & nCMask);

}

В этом совершенном мире линия

SumScratch[idx] += *ptr;

Будет переполняться в следующий байт, если старший байт в *ptr будет больше 9.

Однако этого не происходит, поэтому указанную строку можно заменить на:

unsigned int val = *ptr;
unsigned int ret = 0;
unsigned int carryOut = 0;

asm("{"
    "add.cc.u32 %0, %2, %3;"
    "addc.cc.u32 %1, 0, 0;"
    "}"
    : "=r"(ret), "=r"(carryOut)
    : "r"(val), "r"(OxF6F6F6F6)
);

SumScratch[idx] = 0;
__threadfence();
atomicAdd(&(SumScratch[idx]), ret);
atomicadd(&(SumScratch[idx+1]), carryOut);

Выполнил все симд инструкции по маскировке.

Цель этого существа, если бы у вас было: (от наибольшего знака к наименьшему)

0x00090401 0x09090909 0x10081204

Тогда самый значащий байт наименее достаточного int при добавлении к F6 приведет к переносу в средние суммы int (добавление pist к F6... все байты равны FF), в результате чего все его биты перевернутся на 0, и он перенесет out в самый значимый int.

Итак, на самом деле я хочу рассматривать весь массив так, как если бы это была всего лишь одна двоичная последовательность, и позволять битам просто переворачиваться.

Любые мысли или идеи будут очень признательны.