带CUDA的大整数加法

概述我一直在GPU上开发一种加密算法,目前还在使用算法来执行大的整数加法.大整数以一般的方式表示为一堆32位字. 例如,我们可以使用一个线程来添加两个32位字.为了简单起见,假设 要添加的数字具有相同的长度和每个块的线程数==字数.然后： __global__ void add_kernel(int *C, const int *A, const int *B) { int x = A[th 我一直在GPU上开发一种加密算法,目前还在使用算法来执行大的整数加法.大整数以一般的方式表示为一堆32位字.

例如,我们可以使用一个线程来添加两个32位字.为了简单起见,假设
要添加的数字具有相同的长度和每个块的线程数==字数.然后：

__global__ voID add_kernel(int *C,const int *A,const int *B) {     int x = A[threadIDx.x];     int y = B[threadIDx.x];     int z = x + y;     int carry = (z < x);     /** do carry propagation in parallel somehow ? */     ............     z = z + newcarry; // update the resulting words after carry propagation     C[threadIDx.x] = z; }

我很确定有一种方法可以通过一些棘手的减少程序进行传播,但无法理解.

我看了一下CUDA thrust extensions,但是大整数包似乎还没有实现.
也许有人可以给我一个提示如何在CUDA上这样做？

解决方法 你是对的,进位传播可以通过前缀和计算完成,但是定义这个 *** 作的二进制函数有点棘手,并证明它是关联的(并行前缀和需要).事实上,这个算法(理论上)在 Carry-lookahead adder中被使用.

假设我们有两个大的整数a [0..n-1]和b [0..n-1].
然后我们计算(i = 0..n-1)：

s[i] = a[i] + b[i]l;carryin[i] = (s[i] < a[i]);

我们定义了两个功能：

generate[i] = carryin[i];propagate[i] = (s[i] == 0xffffffff);

具有相当直观的意义：generate [i] == 1表示进位是在…生成的
位置i当传播[i] == 1表示进位将从位置传播
(i-1)至(i 1).我们的目标是计算用于更新所得和s [0..n-1]的函数进位[0..n-1].递归可以递归计算如下：

carryout[i] = generate[i] OR (propagate[i] AND carryout[i-1])carryout[0] = 0

这里carryout [i] == 1如果在位置i产生进位,或者它有时较早生成并传播到位置i.最后,我们更新结果总和：

s[i] = s[i] + carryout[i-1];  for i = 1..n-1carry = carryout[n-1];

现在证明进位函数确实是二进制关联是非常简单的,因此并行前缀和计算适用.为了在CUDA上实现这一点,我们可以将两个标志’generate’和’propagate’合并在一个变量中,因为它们是互斥的,即：

cy[i] = (s[i] == -1u ? -1u : 0) | carryin[i];

换一种说法,

cy[i] = 0xffffffff  if propagate[i]cy[i] = 1           if generate[i]cy[u] = 0           otherwise

然后,可以验证以下公式计算进位功能的前缀和：

cy[i] = max((int)cy[i],(int)cy[k]) & cy[i];

对于所有k <一世.下面的示例代码显示2048个字整数的大量加法.这里我用了512个线程的CUDA块：

// add & output carry flag#define UADDO(c,a,b) \      asm volatile("add.cc.u32 %0,%1,%2;" : "=r"(c) : "r"(a),"r"(b));// add with carry & output carry flag#define UADDC(c,b) \      asm volatile("addc.cc.u32 %0,"r"(b));#define WS 32__global__ voID bignum_add(unsigned *g_R,const unsigned *g_A,const unsigned *g_B) {extern __shared__ unsigned shared[];unsigned *r = shared; const unsigned N_THIDS = 512;unsigned thID = threadIDx.x,thID_in_warp = thID & WS-1;unsigned ofs,cf;uint4 a = ((const uint4 *)g_A)[thID],b = ((const uint4 *)g_B)[thID];UADDO(a.x,a.x,b.x) // adding @R_603_5028@ chunks with carry flagUADDC(a.y,a.y,b.y)UADDC(a.z,a.z,b.z)UADDC(a.w,a.w,b.w)UADDC(cf,0) // save carry-out// memory consumption: 49 * N_THIDS / 64// use "alternating" data layout for each pair of warpsvolatile short *scan = (volatile short *)(r + 16 + thID_in_warp +        49 * (thID / 64)) + ((thID / 32) & 1);scan[-32] = -1; // put IDentity elementif(a.x == -1u && a.x == a.y && a.x == a.z && a.x == a.w)    // this indicates that carry will propagate through the number    cf = -1u;// "Hillis-and-Steele-style" reduction scan[0] = cf;cf = max((int)cf,(int)scan[-2]) & cf;scan[0] = cf;cf = max((int)cf,(int)scan[-4]) & cf;scan[0] = cf;cf = max((int)cf,(int)scan[-8]) & cf;scan[0] = cf;cf = max((int)cf,(int)scan[-16]) & cf;scan[0] = cf;cf = max((int)cf,(int)scan[-32]) & cf;scan[0] = cf;int *postscan = (int *)r + 16 + 49 * (N_THIDS / 64);if(thID_in_warp == WS - 1) // scan leading carry-outs once again    postscan[thID >> 5] = cf;__syncthreads();if(thID < N_THIDS / 32) {    volatile int *t = (volatile int *)postscan + thID;    t[-8] = -1; // load IDentity symbol    cf = t[0];    cf = max((int)cf,(int)t[-1]) & cf;    t[0] = cf;    cf = max((int)cf,(int)t[-2]) & cf;    t[0] = cf;    cf = max((int)cf,(int)t[-4]) & cf;    t[0] = cf;}__syncthreads();cf = scan[0];int ps = postscan[(int)((thID >> 5) - 1)]; // postscan[-1] equals to -1scan[0] = max((int)cf,ps) & cf; // update carry flags within warpscf = scan[-2];if(thID_in_warp == 0)    cf = ps;if((int)cf < 0)    cf = 0;UADDO(a.x,cf) // propagate carry flag if neededUADDC(a.y,0)UADDC(a.z,0)UADDC(a.w,0)((uint4 *)g_R)[thID] = a;}

请注意,由于CUDA 4.0具有相应的内在函数(但我并不完全确定),因此可能不需要UADDO / UADDC宏.

还要指出的是,尽管并行还原相当快,但如果您需要在一行中添加几个大整数,则可能会使用一些冗余表示(上面的注释中提到的),即首先将添加的结果累加在64位字,然后在“一次扫描”中最后执行一个进位传播.

总结

以上是内存溢出为你收集整理的带CUDA的大整数加法全部内容，希望文章能够帮你解决带CUDA的大整数加法所遇到的程序开发问题。

如果觉得内存溢出网站内容还不错，欢迎将内存溢出网站推荐给程序员好友。