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相加的結果。