您尝试做的事情是可能的,但是您在声明和定义将保存和使用函数指针的结构的方式上犯了一些错误。
这是不可能的,并且在编译过程中会出现一个熟悉的错误,即 __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]);
}
}
是完全有效的。当在设备代码中实例化的实例时,函数指针fp
infunctor
隐含地是一个__device__
函数。functor
如果它在宿主代码中实例化,则函数指针将隐含地成为宿主函数。在内核中,作为参数传递的设备函数指针用于实例化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
如果我正确理解了您的问题,那么上面的示例应该为您提供了在设备代码中实现您的想法所需的几乎所有设计模式。