c++ - 没有静态指针或符号拷贝的结构中的 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

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

This is not possible, and ends in a familiar error during compilation regarding the placement of __device__ in the structure.

 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__ 函数都会调用一次内核并显示结果，以及从 host 代码实例化和调用的完全相同的仿函数和函数(因此在主机上运行)比较:

$ 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

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