库达- PTX携带传播

Question

我想在CUDA PTX中添加两个32位无符号整数，我也想要处理进位传播。我正在使用下面的代码来完成这个任务，但是结果并不像预期的那样。

根据文档，add.cc.u32 d, a, b执行整数加法，并将执行值写入条件代码寄存器，即CC.CF。

另一方面，addc.cc.u32 d, a, b使用进位执行整数加法，并将执行值写入条件代码寄存器。这个指令的语义应该是

d = a + b + CC.CF。我也尝试了addc.u32 d, a, b，没有什么不同。

#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime_api.h>
#include "device_launch_parameters.h"
#include <cuda.h>

typedef unsigned int u32;
#define TRY_CUDA_CALL(x) \
do \
  { \
    cudaError_t err; \
    err = x; \
    if(err != cudaSuccess) \
  { \
    printf("Error %08X: %s at %s in line %d\n", err, cudaGetErrorString(err), __FILE__, __LINE__); \
    exit(err); \
  } \
} while(0)


__device__ u32
__uaddo(u32 a, u32 b) {
    u32 res;
    asm("add.cc.u32 %0, %1, %2; /* inline */ \n\t" 
        : "=r" (res) : "r" (a) , "r" (b));
    return res;
}

__device__ u32
__uaddc(u32 a, u32 b) {
    u32 res;
    asm("addc.cc.u32 %0, %1, %2; /* inline */ \n\t" 
        : "=r" (res) : "r" (a) , "r" (b));
    return res;
}

__global__ void testing(u32* s)
{
    u32 a, b;

    a = 0xffffffff;
    b = 0x2;
    
    s[0] = __uaddo(a,b);
    s[0] = __uaddc(0,0);

}

int main()
{
    u32 *s_dev;
    u32 *s;
    s = (u32*)malloc(sizeof(u32));
    TRY_CUDA_CALL(cudaMalloc((void**)&s_dev, sizeof(u32)));
    testing<<<1,1>>>(s_dev);
    TRY_CUDA_CALL( cudaMemcpy(s, s_dev, sizeof(u32), cudaMemcpyDeviceToHost) );
    
    printf("s = %d;\n",s[0]);
    
    
    return 1;
}

据我所知，如果结果不适合变量，您将得到一个进位，如果符号位损坏了，则会出现溢出，但我正在处理无符号值。

上面的代码试图将0xFFFFFFFF添加到0x2中，当然结果并不适合32位，那么为什么在__uaddc(0,0)调用之后我没有得到1呢？

编辑

Nvidia Geforce GT 520 GT

Windows 7终极版，64位

2012

CUDA 7.0

Answer 1

影响asm()语句的唯一数据依赖项是变量绑定显式表示的数据依赖项。请注意，可以绑定寄存器操作数，但不能绑定条件代码。因为在这段代码中，__uaddo(a, b)的结果立即被覆盖，所以编译器确定它对可观察的结果没有贡献，因此是“死代码”，可以消除。通过检查生成的用于cuobjdump --dump-sass版本构建的机器代码(SASS)，可以轻松地检查这一点。

如果我们的代码稍有不同，不允许编译器彻底消除__uaddo()的代码，那么仍然存在一个问题，即编译器可以在为__uaddo()和__uaddc()生成的代码之间安排它喜欢的任何指令，并且这些指令可能会破坏由于__uaddo()而产生的任何进位标志设置。

因此，如果计划在多字算术中使用进位标志，则必须在同一个asm()语句中同时执行进位生成指令和进位消耗指令。在这个答案中可以找到一个有用的示例，它演示了如何添加128位操作数。或者，如果必须使用两个单独的asm()语句，一个可以将前面一个的进位标志设置导出到一个C变量中，然后从那里导入到后续的asm()语句中。我想不出很多情况，这是可行的，因为使用进位标志的性能优势可能会丢失。

Answer 2

因此，正如@njuffa已经说过的那样，来自其他源代码的其他指令可以在两个调用之间修改CC.CF寄存器，并且无法保证获得寄存器的预期值。

作为一种可能的解决方案，可以使用__add32函数：

__device__ uint2 __add32 (u32 a, u32 b)
{
    uint2 res;
    asm ("add.cc.u32      %0, %2, %3;\n\t"
         "addc.u32        %1, 0, 0;\n\t"
         : "=r"(res.x), "=r"(res.y)
         : "r"(a), "r"(b));
    return res;
}

res.y将有可能的进位，res.x将有加法的结果。