为什么没有溢出的加法将 CC.CF 设置为 1？

如何解决为什么没有溢出的加法将 CC.CF 设置为 1？

[ { product: A },{ product: B,info: "reusable,environmental friendly...",quality: "good" },{ product: C,quality: "very good" } ]

此代码打印

#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;\n\t"
    "subc.u32 %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 操作使用了 subc.u32 的倒数。但从文档来看，它不应该是那样的。

解决方法

我无法在 PTX documentation 中的任何地方找到有关 PTX 调用 CC.CF 标志的实际生成方式的信息。查看生成的机器代码 (SASS)，我发现减法是通过加法实现的，并且使用了扩展标志 CC.X。

根据一些快速实验，这个 .X 标志总是似乎是加法器的正常进位。由于 a-b = a+~b+1，减法时 .X 将被设置为 a >= b。它表示来自加法器的进位，它是 x86 风格的减法借位的补码，在 a < b 时设置。

换句话说，GPU 的扩展算术指令似乎使用了 ARM 和 PowerPC 架构为其扩展算术指令使用的相同约定。维基百科关于 carry flag 的文章涵盖了在减法过程中处理标志的两种设计方案。

在问题的代码中，add.cc.u32 清除 CC.CF，这向后续 subc.u32 发出信号，表明发生了借用，从而导致它计算 a+~b。>

您可能希望向 NVIDIA 提交增强请求，以澄清有关 CC.CF 生成和处理细节的 PTX 文档。