c - 大整数加法码

Question

我正在尝试使用以下代码在CUDA中实现大整数加法

__global__ void add(unsigned *A, unsigned *B, unsigned *C/*output*/, int radix){

    int id = blockIdx.x * blockDim.x + threadIdx.x; 
    A[id ] = A[id] + B[id]; 

    C[id ] =  A[id]/radix;
    __syncthreads();
    A[id] =  A[id]%radix + ((id>0)?C[id -1]:0);
    __syncthreads();
    C[id] = A[id];
}

但它不能正常工作，而且我现在也不怎么处理额外的进位位。谢谢

Answer 1

TL; DR建立了一个超前进位加法器，其中每个单独的加法器都将取模数基而不是模2

添加需要进位进位

您的模型中的问题是您有涟漪携带。请参见Rippling carry adders。

如果您使用的是FPGA，那将不是问题，因为它们具有专用的逻辑来快速完成操作（carry chains, they're cool）。但是，a，您使用的是GPU！

也就是说，对于给定的id，只有在计算了所有具有较小A[id]+B[id]值的总和后，您才知道输入进位（因此是要对A[id]+B[id]+1还是id求和）。实际上，最初，您只知道第一个进位。

    A[3]+B[3] + ?     A[2]+B[2] + ?     A[1]+B[1] + ?     A[0]+B[0] + 0       
        |                 |                 |                 |
        v                 v                 v                 v
       C[3]              C[2]              C[1]              C[0]

表征进位输出

每个总和还具有进位输出，该输出不在图形上。因此，您必须将这种较大方案中的加法考虑为具有3个输入和2个输出的函数： (C, c_out) = add(A, B, c_in)

为了不等待O（n）完成总和（其中n是您的总和切入的项数），您可以预先计算每个 id处的所有可能结果。这不是那么大的工作量，因为 A和 B不会改变，只有进位。因此，您有2个可能的输出： (c_out0, C) = add(A, B, 0)和 (c_out1, C') = add(A, B, 1)。

现在，有了所有这些结果，我们需要基本实现 carry lookahead unit。

为此，我们需要弄清楚每个和的进位输出 P和 G的功能：


P下列所有定义

传播
“如果进位进位，那么进位额将超出此总数”
c_out1 && !c_out0
A + B == radix-1

G下列所有定义

生成
“无论进位进位，进位都会从这笔款项中扣除”
c_out1 && c_out0
c_out0
A + B >= radix



换句话说，就是 c_out = G or (P and c_in)。因此，现在我们有了一个算法的开始，该算法可以轻松地直接为每个id告知进位输出作为其进位输入的函数：


在每个 id上计算 C[id] = A[id]+B[id]+0
获取 G[id] = C[id] > radix -1
获取 P[id] = C[id] == radix-1


对数树

现在，即使树状结构在GPU上令人讨厌，我们也可以在O（log（n））中完成，但是比等待时间短。实际上，从彼此相邻的两个添加项中，我们可以得到一个组 G和一个组 P：

对于 id和 id+1：


step = 2
if id % step == 0, do steps 6 through 10, otherwise, do nothing
group_P = P[id] and P[id+step/2]
group_G = (P[id+step/2] and G[id]) or G[id+step/2]
c_in[id+step/2] = G[id] or (P[id] and c_in[id])
step = step * 2
if step < n, go to 5


最后（对树的每个级别重复执行步骤5-10，每次使用更少的 id时），所有内容都将按照您计算出的 P和 G以及 c_in[0]这是 0。在Wikipedia页面上，有按4而不是2进行分组的公式，这将为您提供O（log_4（n））而不是O（log_2（n））的答案。

因此，算法的结尾：


在每个 id处获取 c_in[id]
返回 (C[id]+c_in[id]) % radix


利用硬件

在最后一部分中，我们真正要做的是模仿带有逻辑的超前进位加法器的电路。但是，我们的硬件中已经有添加器，它们可以做类似的事情（根据定义）。

让我们将基于基数的 P和 G定义替换为基于 2的定义，就像我们硬件中的逻辑一样，在每个阶段模拟2位 a和 b的总和：if （异或）和 P = a ^ b（逻辑与）。换句话说， G = a & b和 a = P or G。因此，如果我们创建一个 b = G整数和 intP整数，其中每个位分别是我们根据每个 intG的总和（将我们限制为64个总和）计算出的 P和 G，则加法< cc>具有与我们精心设计的逻辑方案完全相同的进位传播。

我猜想形成这些整数的归约运算仍然是对数运算，但这是可以预期的。

有趣的是，总和的每一位都是其进位输入的函数。实际上，总和的每一位最终都是3位 id的函数。


如果在那个位 (intP | intG) + intG，那么 a+b+c_in % 2，因此 P == 1
否则， a + b == 1是 a+b+c_in % 2 == !c_in或 a+b，而 0


因此，我们可以简单地形成 2为 a+b+c_in % 2 == c_in的整数（或位数组） int_cin = ((P|G)+G) ^ P。

因此，我们的算法有了另一种结局，替换了步骤4和更高版本：


在每个 ^处，将 xor和 id移 P： G和 id
进行“或”归约以获得 P = P << id和 G = G << id，它们是 intG 0..63的所有 intP和 OR的 P
计算（一次） G
在每个 id处，获取`c_in = int_cin＆（1 << id）吗？ 1：0；
返回 int_cin = ((P|G)+G) ^ P




PS：另外，如果 id大，请注意数组中的整数溢出。如果不是这样的话，我猜整个事情就没有意义了...

PPS：在替代结尾中，如果您有64个以上的项目，则以它们的 (C[id]+c_in) % radix和 radix为特征，就好像 P是 G，然后在更高级别上重新运行相同的步骤（还原，获取 radix），然后返回较低的级别，将 2^64与 c_in