1

我正在尝试在 CUDA 中实现我自己的 64 位随机播放功能。但是，如果我这样做：

static __inline__ __device__ double __shfl_xor(double var, int laneMask, int width=warpSize)
{
    int hi, lo;
    asm volatile( "mov.b64 { %0, %1 }, %2;" : "=r"(lo), "=r"(hi) : "d"(var) );
    hi = __shfl_xor( hi, laneMask, width );
    lo = __shfl_xor( lo, laneMask, width );
    return __hiloint2double( hi, lo );
}

对 __shfl_xor 的所有后续调用都将从这个 64 位版本实例化，无论参数的类型是什么。例如，如果我正在做

int a;
a = __shfl_xor( a, 16 );

它仍然会使用双重版本。解决方法可能是使用不同的函数名称。但是由于我是从模板函数中调用这个 shuffle 函数，所以使用不同的名称意味着我必须为 64 位浮点制作不同的版本，这不是很整洁。

那么如何在重载 __shfl_xor(double,...) 函数的同时仍然确保可以适当地调用 __shfl_xor(int,...) 呢？

4

1 回答 1

2

所有整数类型和浮点数都可以向上转换为双精度。当在内置函数和您的专用双函数之间进行选择时，这里的编译器可能会为所有类型选择您的。

您是否尝试过创建一个具有不同名称的函数并使用它来创建您的专用双变量和其他类型的虚拟变量？

例如：

static __inline__ __device__ double foo_shfl_xor(double var, int laneMask, int width=warpSize)
{
    // Your double shuffle implementation
}

static __inline__ __device__ int foo_shfl_xor(int var, int laneMask, int width=warpSize)
{
    // For every non-double data type you use
    // Just call the original shuffle function
    return __shfl_xor(var, laneMask, width);
}

// Your code that uses shuffle
double d;
int a;
foo_shfl_xor(d, ...); // Calls your custom shuffle
foo_shfl_xor(a, ...); // Calls default shuffle
于 2013-04-12T06:31:19.290 回答