CUDA - PTX 进位传播
CUDA - PTX carry propagation
我想在 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？
编辑
英伟达 Geforce GT 520mx
Windows 7 旗舰版，64 位
Visual Studio 2012
CUDA 7.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相加的结果
影响 asm() 语句的唯一数据依赖项是那些由变量绑定明确表达的数据依赖项。请注意，您可以绑定寄存器操作数，但不能绑定条件代码。由于在此代码中 __uaddo(a, b) 的结果立即被覆盖，编译器确定它对可观察结果没有贡献，因此是 "dead code" 并且可以消除。这很容易通过检查生成的机器代码 (SASS) 来检查发布版本 cuobjdump --dump-sass.
如果我们的代码略有不同，不允许编译器完全消除 __uaddo() 的代码，仍然会存在编译器可以在为 [ 生成的代码之间安排它喜欢的任何指令的问题=13=] 和 __uaddc()，并且由于 __uaddo().
，此类指令可能会破坏进位标志的任何设置
因此，如果计划将进位标志用于多字运算，则进位生成指令和进位消耗指令必须出现在同一个 asm() 语句中。在 this answer 中可以找到一个有效的示例，它显示了如何添加 128 位操作数。或者，如果必须使用两个单独的 asm() 语句 ，则可以将进位标志设置从较早的语句导出到 C 变量中，然后将其导入后续的 asm() 那里的声明。我想不出很多情况下这会实用，因为使用进位标志的性能优势可能会丢失。
我想在 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？
编辑
英伟达 Geforce GT 520mx
Windows 7 旗舰版，64 位
Visual Studio 2012
CUDA 7.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相加的结果
影响 asm() 语句的唯一数据依赖项是那些由变量绑定明确表达的数据依赖项。请注意，您可以绑定寄存器操作数，但不能绑定条件代码。由于在此代码中 __uaddo(a, b) 的结果立即被覆盖，编译器确定它对可观察结果没有贡献，因此是 "dead code" 并且可以消除。这很容易通过检查生成的机器代码 (SASS) 来检查发布版本 cuobjdump --dump-sass.
如果我们的代码略有不同，不允许编译器完全消除 __uaddo() 的代码，仍然会存在编译器可以在为 [ 生成的代码之间安排它喜欢的任何指令的问题=13=] 和 __uaddc()，并且由于 __uaddo().
因此，如果计划将进位标志用于多字运算，则进位生成指令和进位消耗指令必须出现在同一个 asm() 语句中。在 this answer 中可以找到一个有效的示例，它显示了如何添加 128 位操作数。或者，如果必须使用两个单独的 asm() 语句 ，则可以将进位标志设置从较早的语句导出到 C 变量中，然后将其导入后续的 asm() 那里的声明。我想不出很多情况下这会实用，因为使用进位标志的性能优势可能会丢失。