结构中的 CUDA 设备函数指针，没有静态指针或符号副本

Question

如果可能的话，我的预期程序流程如下所示：

typedef struct structure_t
{
  [...]
  /* device function pointer. */
  __device__ float (*function_pointer)(float, float, float[]);
  [...]
} structure;

[...]

/* function to be assigned. */
__device__ float
my_function (float a, float b, float c[])
{
  /* do some stuff on the device. */
  [...]
}

void
some_structure_initialization_function (structure *st)
{
  /* assign. */
  st->function_pointer = my_function;
  [...]
}

这是不可能的，并且在编译过程中会出现一个熟悉的错误，即 __device__ 在结构中的位置。

 error: attribute "device" does not apply here

stackoverflow 上有一些类似类型问题的示例，但它们都涉及在结构外使用静态指针。例如device function pointers as struct members 和device function pointers。我之前在其他代码中采用了类似的方法并取得了成功，在这些代码中我很容易使用静态设备指针并在任何结构之外定义它们。目前虽然这是一个问题。它被编写为各种 API，用户可以定义一个或两个或几十个需要包含设备函数指针的结构。因此，在结构之外定义静态设备指针是一个主要问题。

我相当肯定答案存在于我上面链接的帖子中，通过使用符号副本，但我无法成功使用它们。

Answer 1

您尝试做的事情是可能的，但是您在声明和定义将保存和使用函数指针的结构的方式上犯了一些错误。

这是不可能的，并且在编译期间以熟悉的错误结束 关于 __device__ 在结构中的位置。

 error: attribute "device" does not apply here

这只是因为您试图将内存空间分配给结构或类数据成员，这在 CUDA 中是非法的。所有类或结构数据成员的内存空间都是在定义或实例化类时隐式设置的。所以只是略有不同（更具体）：

typedef float (* fp)(float, float, float4);

struct functor
{
    float c0, c1;
    fp f;

    __device__ __host__
    functor(float _c0, float _c1, fp _f) : c0(_c0), c1(_c1), f(_f) {};

    __device__ __host__
    float operator()(float4 x) { return f(c0, c1, x); };
};

__global__
void kernel(float c0, float c1, fp f, const float4 * x, float * y, int N)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;

    struct functor op(c0, c1, f);
    for(int i = tid; i < N; i += blockDim.x * gridDim.x) {
        y[i] = op(x[i]);
    }
}

完全有效。当functor 的实例在设备代码中实例化时，functor 中的函数指针 fp 隐含为 __device__ 函数。如果它在宿主代码中实例化，则函数指针将隐含地成为宿主函数。在内核中，作为参数传递的设备函数指针用于实例化functor 实例。完全合法。

我相信我说没有直接的方法可以在主机代码中获取__device__ 函数的地址是正确的，因此您仍然需要一些静态声明和符号操作。这可能在 CUDA 5 中有所不同，但我还没有测试过。如果我们用几个__device__ 函数和一些支持主机代码来充实上面的设备代码：

__device__ __host__ 
float f1 (float a, float b, float4 c)
{
    return a + (b * c.x) +  (b * c.y) + (b * c.z) + (b * c.w);
}

__device__ __host__
float f2 (float a, float b, float4 c)
{
    return a + b + c.x + c.y + c.z + c.w;
}

__constant__ fp function_table[] = {f1, f2};

int main(void)
{
    const float c1 = 1.0f, c2 = 2.0f;
    const int n = 20;
    float4 vin[n];
    float vout1[n], vout2[n];
    for(int i=0, j=0; i<n; i++) {
        vin[i].x = j++; vin[i].y = j++;
        vin[i].z = j++; vin[i].w = j++;
    }

    float4 * _vin;
    float * _vout1, * _vout2;
    size_t sz4 = sizeof(float4) * size_t(n);
    size_t sz1 = sizeof(float) * size_t(n);
    cudaMalloc((void **)&_vin, sz4);
    cudaMalloc((void **)&_vout1, sz1);
    cudaMalloc((void **)&_vout2, sz1);
    cudaMemcpy(_vin, &vin[0], sz4, cudaMemcpyHostToDevice);

    fp funcs[2];
    cudaMemcpyFromSymbol(&funcs, "function_table", 2 * sizeof(fp));

    kernel<<<1,32>>>(c1, c2, funcs[0], _vin, _vout1, n);
    cudaMemcpy(&vout1[0], _vout1, sz1, cudaMemcpyDeviceToHost); 

    kernel<<<1,32>>>(c1, c2, funcs[1], _vin, _vout2, n);
    cudaMemcpy(&vout2[0], _vout2, sz1, cudaMemcpyDeviceToHost); 

    struct functor func1(c1, c2, f1), func2(c1, c2, f2); 
    for(int i=0; i<n; i++) {
        printf("%2d %6.f %6.f (%6.f,%6.f,%6.f,%6.f ) %6.f %6.f %6.f %6.f\n", 
                i, c1, c2, vin[i].x, vin[i].y, vin[i].z, vin[i].w,
                vout1[i], func1(vin[i]), vout2[i], func2(vin[i]));
    }

    return 0;
}

您将获得一个完全可编译且可运行的示例。这里两个__device__ 函数和一个静态函数表为宿主代码提供了在运行时检索__device__ 函数指针的机制。每个__device__ 函数调用内核一次并显示结果，以及从主机 代码（并因此在主机上运行）实例化和调用的完全相同的仿函数和函数以进行比较：

$ nvcc -arch=sm_30 -Xptxas="-v" -o function_pointer function_pointer.cu 

ptxas info    : Compiling entry function '_Z6kernelffPFfff6float4EPKS_Pfi' for 'sm_30'
ptxas info    : Function properties for _Z6kernelffPFfff6float4EPKS_Pfi
    16 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Function properties for _Z2f1ff6float4
    24 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Function properties for _Z2f2ff6float4
    24 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 16 registers, 356 bytes cmem[0], 16 bytes cmem[3]

$ ./function_pointer 
 0      1      2 (     0,     1,     2,     3 )     13     13      9      9
 1      1      2 (     4,     5,     6,     7 )     45     45     25     25
 2      1      2 (     8,     9,    10,    11 )     77     77     41     41
 3      1      2 (    12,    13,    14,    15 )    109    109     57     57
 4      1      2 (    16,    17,    18,    19 )    141    141     73     73
 5      1      2 (    20,    21,    22,    23 )    173    173     89     89
 6      1      2 (    24,    25,    26,    27 )    205    205    105    105
 7      1      2 (    28,    29,    30,    31 )    237    237    121    121
 8      1      2 (    32,    33,    34,    35 )    269    269    137    137
 9      1      2 (    36,    37,    38,    39 )    301    301    153    153
10      1      2 (    40,    41,    42,    43 )    333    333    169    169
11      1      2 (    44,    45,    46,    47 )    365    365    185    185
12      1      2 (    48,    49,    50,    51 )    397    397    201    201
13      1      2 (    52,    53,    54,    55 )    429    429    217    217
14      1      2 (    56,    57,    58,    59 )    461    461    233    233
15      1      2 (    60,    61,    62,    63 )    493    493    249    249
16      1      2 (    64,    65,    66,    67 )    525    525    265    265
17      1      2 (    68,    69,    70,    71 )    557    557    281    281
18      1      2 (    72,    73,    74,    75 )    589    589    297    297
19      1      2 (    76,    77,    78,    79 )    621    621    313    313

如果我正确理解了您的问题，上述示例应该为您提供了在设备代码中实现您的想法所需的几乎所有设计模式。