cuda - 为什么没有溢出的加法将 CC.CF 设置为 1?
问题描述
我有下一个代码
#include <stdio.h>
#include <cuda.h>
#include <cuda_runtime.h>
__global__ void cuda_test() {
int result;
asm(
".reg .u32 r1;\n\t"
"add.cc.u32 r1, 0, 0;\n\t"
"subc.u32 %0, 0, 0; \n\t"
:"=r"(result)
);
printf("r= %x\n", result);
}
int main() {
cuda_test<<<1, 1>>>();
cudaDeviceSynchronize();
return 0;
}
此代码打印
r= ffffffff
为什么?据我了解,操作add.cc.u32 r1, 0, 0
必须将进位标志设置为0
. 我的印象是该subc.u32
操作使用CC.CF
. 但从文档来看,它不应该是这样的。
解决方案
我在PTX 文档中的任何地方都找不到有关PTX 所谓的CC.CF
标志是如何实际生成的信息。查看生成的机器代码(SASS),我看到减法是通过加法实现的,并使用扩展标志CC.X
。
根据一些快速实验,这个.X
标志似乎总是加法器的正常执行。由于a-b
= ,如果 将设置a+~b+1
减法。它表示加法器的进位,它是 x86 风格的减法借位的补码,在 时设置。.X
a >= b
a < b
换句话说,GPU 的扩展算术指令似乎使用了 ARM 和 PowerPC 架构用于其扩展算术指令的相同约定。关于进位标志的维基百科文章涵盖了在减法期间处理标志的两种设计方案。
在问题中的代码中,add.cc.u32
clearsCC.CF
向随后subc.u32
发出借位发生的信号,导致其计算a+~b
。
CC.CF
您可能希望向 NVIDIA 提交增强请求,以阐明有关生成和处理细节的 PTX 文档。