- c - 在位数组中找到第一个零
- linux - Unix 显示有关匹配两种模式之一的文件的信息
- 正则表达式替换多个文件
- linux - 隐藏来自 xtrace 的命令
我一直在 GPU 上开发加密算法,目前坚持使用一种算法来执行大整数加法。大整数通常表示为一堆 32 位字。
例如,我们可以用一个线程来添加两个32位的字。为简单起见,假设要添加的数字具有相同的长度和每个 block 的线程数 == 字数。然后:
__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 while propagate[i] == 1 意味着进位将从位置传播(i - 1) 到 (i + 1)。我们的目标是计算用于更新结果和 s[0..n-1] 的函数 carryout[0..n-1]。 carryout 可以递归计算如下:
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-1
carry = carryout[n-1];
现在可以非常简单地证明进位函数确实是二元关联的,因此可以应用并行前缀和计算。为了在 CUDA 上实现这一点,我们可以将标志“生成”和“传播”合并到一个变量中,因为它们是互斥的,即:
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 < i。下面的示例代码显示了 2048 字整数的大加法。这里我使用了 512 个线程的 CUDA block :
// 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, a, b) \
asm volatile("addc.cc.u32 %0, %1, %2;" : "=r"(c) : "r"(a) , "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 128-bit chunks with carry flag
UADDC(a.y, a.y, b.y)
UADDC(a.z, a.z, b.z)
UADDC(a.w, a.w, b.w)
UADDC(cf, 0, 0) // save carry-out
// memory consumption: 49 * N_THIDS / 64
// use "alternating" data layout for each pair of warps
volatile short *scan = (volatile short *)(r + 16 + thid_in_warp +
49 * (thid / 64)) + ((thid / 32) & 1);
scan[-32] = -1; // put identity element
if(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 -1
scan[0] = max((int)cf, ps) & cf; // update carry flags within warps
cf = scan[-2];
if(thid_in_warp == 0)
cf = ps;
if((int)cf < 0)
cf = 0;
UADDO(a.x, a.x, cf) // propagate carry flag if needed
UADDC(a.y, a.y, 0)
UADDC(a.z, a.z, 0)
UADDC(a.w, a.w, 0)
((uint4 *)g_R)[thid] = a;
}
请注意,宏 UADDO/UADDC 可能不再是必需的,因为 CUDA 4.0 具有相应的内在函数(但我不完全确定)。
另请注意,虽然并行缩减非常快,但如果您需要连续添加几个大整数,最好使用一些冗余表示(上面的评论中建议),即首先累加结果在 64 位字中进行加法运算,然后在“一次扫描”中在最后执行一次进位传播。
关于c - CUDA 大整数加法,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/12957116/
我正在尝试将父 div 标记的最小宽度设置为内部所有子项的等效宽度。有办法吗? 例如, #sidebar{ width: 325px; } #content{ width: 500
我正在其中一个脚本中做一些附加操作,下面是一些简化的代码: foreach($entry in $arr){ ... switch($entry.AccessRights) { "GenericRea
float 在我的 Java/JOGL (OpenGL for Java) 程序中没有按预期计算。在绘制方法中,当调用每一帧(每秒 60 帧)时,我尝试修改对象的位置。所有值都是浮点值。 float
我正在尝试使用 C 中的结构为一个项目进行复杂的 vector 加法和点积。我已经编写了代码,但是,虽然它的编译没有问题,但一旦我运行我的程序,它就会停止工作。我还有该程序的其他部分,但这只是相关部分
这个问题已经有答案了: Use of java.math.MathContext (5 个回答) 已关闭 8 年前。 首先,我的搜索能力可能没有我希望的那么好,所以也许这种问题已经存在了。如果是的话请
PFB 说明问题的示例代码片段: var x=0.323; var cumulativeVal = 0; for(i=0;i<30;i++){
这个查询的每一步在 PostgreSQL 中的执行顺序是什么? SELECT SUM(field1)+SUM(field2)+SUM(field3)-SUM(field4); 据我所知,加法/减法是按
我正在尝试熟悉 Java 多线程应用程序。我试图想出一个可以很好地并行化的简单应用程序。我认为 vector 加法是一个很好的应用。但是,在我的 Linux 服务器(有 4 个内核)上运行时,我没有得
我在进行简单的加法并将值保存在变量中时遇到问题。 基本上我有以下代码: var accsen; var lowsev = parseInt(accsen); var hisev
所以我最近几个小时一直在解决一个问题,似乎无法阻止我的程序崩溃。问题是创建一个程序,该程序采用任意大小的矩阵,并且能够使用运算符重载将一个矩阵加到另一个矩阵上。当我尝试添加我类(class)的两个对象
我正在尝试添加以下内容,但它一直连接并返回一个字符串。 var nums = [1.99, 5.11, 2.99]; var total = 0; nums.forEach(f
我在网上搜索了数据仓库中加法、半加法和非加法度量之间的区别。我找到了一些结果,但我很难理解这些差异,因为它们不是一个例子。您能否通过示例向我更多地解释加法、半加法和非加法措施之间的区别。 最佳答案 T
%{control.current + #displayRows} 最终是我需要执行的语句。我将其放在 s:if 标记中,并使用 test 来查看该值是否在特定范围内。 最终,我得到的是字符串连接而不
请帮助我解释为什么下面的代码会得到奇怪的输出......为什么 getName() 得到 null。 输出: 列表检查:null:1 public class ListTest { public st
我需要通过字典生成校验和。键和值。 是否有任何简单的方法以迭代方式完成此任务。 foreach(dic.Keys 中的变量项) 校验和 += 校验和(dic[item]) + 校验和(item); 在
我想计算平均销售产品数量。表: pieces | date | status ------------------------------------------- 1
我正在尝试从 mysql 获取 INT 值并进行添加,最后更新数据库。不过这个好像没有更新?我该如何解决这个问题? $resultSecond = mysql_query("SELECT * FROM
我遇到了一个奇怪的问题。 有一张图片,我只需要重新计算非零像素。我想通过 numpy 来完成,因为我处理了数千张图像并且我需要它的速度。 这是一个维度较低的简化示例。 假设我有以下矩阵: [[0,
我不确定下一步该做什么。它们只是文本字段中的美元金额。我正在尝试将它们加在一起。 NSString *checkAmount = [checkAmountInput.text substringFro
我正在测试我的一些代码,在 javascript 中我添加了 .1+.2 ,它给了我 .30000000000000004 而不是 .3 。我不明白这一点。但是当我添加 .1+.3 时,它给了我 .4
我是一名优秀的程序员,十分优秀!