__device__函数指针

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;

}