在 cuda 设备函数中使用省略号

Question

我正在尝试将一些 C 代码移植到 cuda 内核。我移植的代码普遍使用省略号。当我尝试在如下所示的设备函数中使用省略号时，我收到一条错误消息，指出设备函数中不允许使用省略号。

__device__ int add(int a, ...){}

但是，cuda 支持在主机和设备函数中使用 printf，并在其自己的代码中使用省略号，如下面的 common_functions.h 中。

extern "C"
{
extern _CRTIMP __host__ __device__ __device_builtin__ __cudart_builtin__ int     __cdecl printf(const char*, ...);
extern _CRTIMP __host__ __device__ __device_builtin__ __cudart_builtin__ int     __cdecl fprintf(FILE*, const char*, ...);
extern _CRTIMP __host__ __device__ __cudart_builtin__ void*   __cdecl malloc(size_t) __THROW;
extern _CRTIMP __host__ __device__ __cudart_builtin__ void    __cdecl free(void*) __THROW;

}

有没有办法在设备函数中使用省略号？

我不想硬编码最大数量的参数，然后更改所有调用。
我也不想编写自定义可变参数函数方法。

我还尝试创建一个 PTX 文件，我可以用它来替换省略号的用法，因为 ISA PTX 文档似乎具有处理可变参数的工具（请注意，文档说它不支持它们，然后提供了一个段落支持函数和例子。也许，有一个错字？）。在下面定义的过程中，我一直得到一个简单的 PTX 文件，但在最后一条评论中遇到了可执行问题。我打算阅读 nvcc 编译器文档来尝试理解。

How can I call a ptx function from CUDA C?

我在 Ubuntu 12.04 上使用 GTX660，我认为它是 3.0 级和 cuda 5.0 工具包。

关于下面提到的“魔法”的更新：

在我看来，编译器中一定发生了一些特别的事情来限制省略号的使用并做一些特别的事情。当我如下调用 printf 时：

printf("The result = %i from adding %i numbers.", result, 2);

我很惊讶在 ptx 中发现了这个：

.extern .func  (.param .b32 func_retval0) vprintf
(
.param .b64 vprintf_param_0,
.param .b64 vprintf_param_1
)

以后

    add.u64     %rd2, %SP, 0;
st.u32  [%SP+0], %r5;
mov.u32     %r6, 2;
st.u32  [%SP+4], %r6;
// Callseq Start 1
{
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64    [param0+0], %rd1;
.param .b64 param1;
st.param.b64    [param1+0], %rd2;
.param .b32 retval0;
call.uni (retval0), 
vprintf, 
(
param0, 
param1
);

在我看来，编译器接受 printf 的省略号，但随后交换对 vprintf 的调用并手动动态创建 va_list。 va_list 是设备函数中的有效类型。

Answer 1

正如@JaredHoberock 所说（我想他不会介意我创建一个答案）：

__device__ 函数不能有省略号参数；这就是您收到编译器错误消息的原因。

内置的printf函数是一种特殊情况，并不表示一般支持省略号。

可以提到一些替代方案，但据我所知，没有一个支持真正的通用变量参数。例如，正如 Jared 所说，您可以简单地定义一些参数，其中一些/大部分都指定了默认值，因此不需要显式传递它们。

您也可以像cuPrintf sample code 中所做的那样使用模板玩游戏来尝试模拟可变参数，但这也不能任意扩展，我不认为。