Using a __device__ function pointer. Problems using a pointer to a __device__ function.

I am trying to make a kernel that invokes a device function via a pointer. It works well as long as the function and its caller reside in the same source (.cu) file, but breaks if they are in different files. Here is the full example code.


#ifndef FuncPointer_h

#define FuncPointer_h

typedef float (*op_func) (float, float);

struct FuncPointer {


    op_func fptr;


#endif // FuncPointer_h

#include <cstdio>

#include "FuncPointer.h"

/// start of

__device__ float add_func (float x, float y)


    return x + y;


__device__ op_func func = add_func;

FuncPointer::FuncPointer() {

    cudaMemcpyFromSymbol(&fptr, func, sizeof(func));


/// end of

__global__ void kernel (FuncPointer* p)


    float x=100, y=10, result=0;

    result = p->fptr(x, y);

    printf ("result = %f\n", result);


int main () 


    FuncPointer fp;

	FuncPointer* dev_fp;

	cudaMalloc(&dev_fp, sizeof(FuncPointer));

	cudaMemcpy(dev_fp, &fp, 

		sizeof(FuncPointer), cudaMemcpyHostToDevice);






This works as expected.

Note the code section between [font=“Courier New”]/// start of[/font] and [font=“Courier New”]/// end of[/font].

If I move this code from into another file and link them together, the execution stops with the “unspecified launch error” message.

What is wrong with calling a device function from another file by pointer?

A similar question was asked in this post, but never answered.

Calling a device function defined in a different compilation unit requires linking of device code, so references to other compilation units can be resolved. Up to and including CUDA 4.2, there is no support for linking of device code. In this case the function pointer is a device variable not accessible from outside the compilation unit it is defined in.

As was announced at GTC, CUDA 5.0 will provide for (static) linking of device code. See for example this presentation by our chief technologist for GPU computing, Mark Harris:

A CUDA 5.0 preview is available to registered developers. I believe (but have not checked) that the new linker is part of the preview. Please note that the stability and maturity of the preview should not be assume to be on par with that of release candidates.

Thanks for the explanation; I will try CUDA 5.0 preview out!

BTW, I see the same problem if I try to invoke a virtual function that is defined in a different compilation unit. For the same reason, obviously.