Static array of pointers to __device__ functions

Accelerated Computing CUDA CUDA Programming and Performance
1

I’ve recently written code that references a static array of pointers to device functions, as shown in the extract below

template<int DIM>
__device__ void helperKernel( int x )
{
    int buffer[DIM];  // this is real reason for using this 'pattern'
    // ...
}

typedef void (*HelperKernel)( int );
__device__ static HelperKernel helper[] = {
                                                0 ,
                                                &helperKernel<1> ,
                                                &helperKernel<2> ,                                            
                                                &helperKernel<3>
                                            };

__global__ void demo( int x , int dim )
{
    if( dim > 0 && dim < 4 )
    {
        (helper[dim])( x );
    }
}

The code runs (or appears to run) successfully but I’m not certain if its valid CUDA code. Does CUDA permit static arrays of pointers to device functions? If so, out of interest, when would this static array be initialized?

2

a device function pointer, captured into a device variable should work. I don’t think the use of an array has any bearing on that statement.

not sure what you are indicating with the use of the static keyword here (its not obviously necessary to me)

The CUDA programming guide claims adherence to a particular C++ version, and lists restrictions/deviations. You can review those. I don’t think this falls into any of the documented restrictions, and in my experience this approach works

3

thanks txbob.