2

我想在 CUDA PTX 中添加两个 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

4

2 回答 2

2

影响asm()语句的唯一数据依赖关系是那些由变量绑定明确表达的依赖关系。请注意，您可以绑定寄存器操作数，但不能绑定条件代码。由于在此代码中的结果__uaddo(a, b)立即被覆盖，编译器确定它对可观察结果没有贡献，因此是“死代码”并且可以被消除。这很容易通过检查生成的机器代码 (SASS) 以使用cuobjdump --dump-sass.

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

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

于 2016-03-26T21:13:29.583 回答
0

因此，正如@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加法的结果。

于 2016-03-26T21:10:52.040 回答