为什么没有溢出的加法设置 CC.CF 为 1?
Why is addition without overflow set CC.CF to 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 documentation 中的任何地方都找不到关于 PTX 调用 CC.CF
标志的实际生成方式的信息。查看生成的机器代码 (SASS) 我看到减法是通过加法实现的,并且使用扩展标志 CC.X
.
根据一些快速实验,这个 .X
标志总是 似乎 是加法器的正常进位输出。由于 a-b
= a+~b+1
,如果 a >= b
将设置减法 .X
。它代表加法器的进位输出,它是 x86 风格的减法借位的补码,在 a < b
.
时设置
换句话说,GPU 的扩展算术指令似乎使用了 ARM 和 PowerPC 架构用于其扩展算术指令的相同约定。 carry flag 上的维基百科文章介绍了在减法过程中处理标志的两种设计方案。
在问题的代码中,add.cc.u32
清除了 CC.CF
,这向后续的 subc.u32
发出了借位发生的信号,导致它计算 a+~b
。
您可能希望向 NVIDIA 提交增强请求,以澄清有关 CC.CF
生成和处理细节的 PTX 文档。
我有下一个代码
#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 documentation 中的任何地方都找不到关于 PTX 调用 CC.CF
标志的实际生成方式的信息。查看生成的机器代码 (SASS) 我看到减法是通过加法实现的,并且使用扩展标志 CC.X
.
根据一些快速实验,这个 .X
标志总是 似乎 是加法器的正常进位输出。由于 a-b
= a+~b+1
,如果 a >= b
将设置减法 .X
。它代表加法器的进位输出,它是 x86 风格的减法借位的补码,在 a < b
.
换句话说,GPU 的扩展算术指令似乎使用了 ARM 和 PowerPC 架构用于其扩展算术指令的相同约定。 carry flag 上的维基百科文章介绍了在减法过程中处理标志的两种设计方案。
在问题的代码中,add.cc.u32
清除了 CC.CF
,这向后续的 subc.u32
发出了借位发生的信号,导致它计算 a+~b
。
您可能希望向 NVIDIA 提交增强请求,以澄清有关 CC.CF
生成和处理细节的 PTX 文档。