https://forums.developer.nvidia.com/t/how-can-i-use-device-function-pointer-in-cuda/14405/8
device double step_d[8]={0.5,0.25,0.125,0.06250,0.03125,0.015625,0.0078125,0.0039065};
但是编译出现这样的错误:expression must have integral or enum type
__device__定义的数组/变量是device端的全局数组/变量,无需通过参数传递,直接在kernel中可以访问。您可以在kernel中通过下标来选择您需要的数据。
以及,此类数组/变量称为Symbol,有一些特殊处理。
如果您想使用kernel参数的方法访问,那么您需要使用API函数cudaGetSymbolAddress()来获得Symbol的实际地址,然后作为参数传递给kernel使用。
#include <stdio.h>
#include <stdlib.h>
#define N 5
__device__ float add_func (float x, float y)
{
return x + y;
}
__device__ float mul_func (float x, float y)
{
return x * y;
}
__device__ float div_func (float x, float y)
{
return x / y;
}
typedef float (*op_func) (float, float);
__device__ op_func func[3] = { add_func, mul_func, div_func };
__device__ char* op_name[3] = { "add", "mul", "div" };
__device__ void op_array (const float *a, const float *b, float *res, int op, int n)
{
for (int i = 0; i < N; i++) {
res[i] = func[op](a[i], b[i]);
}
}
__global__ void kernel (void)
{
float x[N];
float y[N];
float res[N];
for (int i = 0; i < N; i++) {
x[i] = (float)(10 + i);
}
for (int i = 0; i < N; i++) {
y[i] = (float)(100 + i);
}
for (int op = 0; op < 3; op++) {
printf ("\nop=%s\n", op_name[op]);
op_array (x, y, res, op, N);
for (int i = 0; i < N; i++) {
printf ("res = % 16.9e\n", res[i]);
}
}
}
int main (void)
{
kernel<<<1,1>>>();
cudaThreadSynchronize();
return EXIT_SUCCESS;
}