CUDA - PTX进行传播

Question

我想在CUDA PTX中添加两个32位无符号整数，并且我也希望处理进位传播。我使用下面的代码来做到这一点，但结果并不如预期。
根据documentation，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 520mx
Windows 7旗舰版，64位
的Visual Studio 2012
CUDA 7.0

Answer 1

影响asm()语句的唯一数据依赖性是那些由变量绑定明确表示的数据依赖项。请注意，您可以绑定寄存器操作数，但不能绑定条件代码。由于在这段代码中，__uaddo(a, b)的结果立即被覆盖，编译器确定它不会对可观察结果做出贡献，因此是“死代码”并且可以被消除。通过使用cuobjdump --dump-sass检查生成的机器代码（SASS）来发布版本，可以轻松检查。

如果我们有略有不同的代码不允许编译器，以消除代码__uaddo()顾左右而言他，仍然会有编译器可以安排它喜欢的__uaddo()和__uaddc()生成的代码之间的任何指令的问题，这样的由于__uaddo()，指令可能会破坏进位标志的任何设置。因此，如果打算将进位标志用于多字算术，则进位产生和进位消耗指令必须发生在相同的asm()语句中。在this answer中可以找到一个工作示例，其中显示了如何添加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相加的结果。