1

大约一个月来，我一直在努力解决这个问题，但我的 C 技能和我的 google-fu 都没有足够强大，无法提出解决方案。

我最喜欢的一个副项目已经并且继续尝试通过 reverse and add 方法找到数字 196 的回文：

196 + 691 = 887

887 + 788 = 1675

一直如此，直到结果从前到后读取相同。

最近，我选择的方法是使用 cuda，但我一遍又一遍地遇到同样的问题。携带。

备份，我将内存中的数字表示为一个无符号字符数组，每个数字都是一个字符 - 所以本质上是解压的 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);

接下来是所有用于屏蔽的 simd 指令。

如果你有这个存在的目标：（最多 Sig 到最少）

0x00090401 0x09090909 0x10081204

然后，最不充足的 int 的最高有效字节，当添加到 F6 时，将导致中间 int 的总和（pist 添加到 F6...所有字节都是 FF）导致其所有位翻转为 0 并携带进入最重要的 int。

因此，实际上，我想将整个数组视为一个二进制序列，并允许位翻转。

任何想法或想法将不胜感激。

4

0 回答 0