cuda - CUDA - PTX キャリー伝搬

Question

CUDA PTX に 2 つの 32 ビット符号なし整数を追加したいのですが、キャリーの伝播も処理したいと考えています。以下のコードを使用してそれを実行していますが、結果は期待どおりではありません。
ドキュメントによると、add.cc.u32 d, a, bは整数加算を実行し、キャリーアウト値を条件コード レジスタに書き込みますCC.CF。
一方、キャリーインでaddc.cc.u32 d, a, b整数加算を行い、キャリーアウトした値をコンディションコードレジスタに書き込みます。この命令のセマンティクスは. 私も違いなしで試しました。
d = a + b + CC.CFaddc.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 Ultimate、64 ビット
Visual Studio 2012
CUDA 7.0