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

Question

我有下一个代码

#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 的倒数。但从文档来看，它不应该是这样的。

Answer 1

我在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 上的 Wikipedia 文章涵盖了在减法期间处理标志的两种设计替代方案。

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

您可能希望向 NVIDIA 提交增强请求，以阐明 PTX 文档中有关 CC.CF 生成和处理的详细信息。