1

Cuda中有错误吗？我在 GTX580 上运行了以下代码，最后 r1 为零。我希望它是一种由于携带传播？我已经使用 Cuda Toolkit 4.2.9 和 5.5 测试了代码，并使用“nvcc -arch=sm_20 bug.cu -o bug && ./bug”来编译和运行它。

#include <stdio.h>
#include <cuda.h>

__global__ void bug()
{
  unsigned int r1 = 0;
  unsigned int r2 = 0;

  asm( "\n\t"
       "sub.cc.u32 %0, 0, 1;\n\t"
       "addc.cc.u32 %1, 0, 0;\n\t"
     : "=r"(r1), "=r"(r2) );

  printf("r1 >> %04X\n", r1);
  printf("r2 >> %04X\n", r2);

}

int main(void)
{
  float *a_d;
  cudaMalloc((void **) &a_d, 1);

  bug <<< 1,1 >>> ();

  cudaFree(a_d);
}

Output
r1 >> FFFFFFFF
r2 >> 0000
4

2 回答 2

3

我相信您正在对PTX ISA 文档CC.CF中引用的标志做出一些可能无效的假设。

请注意，我从未看到该位的特定状态（例如0或1）的定义。此外，我没有找到“带入/带出”和“借入/借出”的定义之间的任何映射

换句话说，我认为您假设此标志中的“借用”状态与“携带”状态相同。换句话说，您假设如下：

CF:  
0    =  (NO CARRY) or (NO BORROW)
1    =  (CARRY) or (BORROW)

但是从来没有给出这样的真值表或映射。此外，手册指出：

条件码寄存器 ... 主要用于计算扩展精度整数加法、减法和乘法的直线代码序列。

我认为您的代码不满足意图，我也不认为上述真值表假设CC.CF是有效的。

事实上，我认为正在发生的是这样的真值表：

CF:  
0    =  (CARRY) or (NO BORROW)
1    =  (NO CARRY) or (BORROW)

（这里的 0 和 1 是任意的；手册中也没有定义。）

我尝试过的所有代码示例（大约 6 个案例，包括您的案例）都符合我上面给出的定义。

话虽如此，我认为依赖它是不明智的，因为它大多是无证的。计算机体系结构的一个安全规则是，未记录的行为将来可能会发生变化。

于 2013-08-09T16:42:31.703 回答
-1

我想我找到了解释。PTX 手册中有一条关于 sub.cc 指令的注释：“无符号整数和有符号整数的行为相同。”

于 2013-08-09T23:51:57.027 回答