device function pointers

by

function pointers are allowed on Fermi.
This is how you could do it:

typedef double (*func)(double x);

__device__ double func1(double x)
{
return x+1.0f;
}

__device__ double func2(double x)
{
return x+2.0f;
}

__device__ double func3(double x)
{
return x+3.0f;
}

__device__ func pfunc1 = func1;
__device__ func pfunc2 = func2;
__device__ func pfunc3 = func3;

__global__ void test_kernel(func* f, int n)
{
  double x;

  for(int i=0;i<n;++i){
   x=f[i](2.0);
   printf("%g\n",x);
  }
}

int main(void)
{
  int N = 5;
  func* h_f;
  func* d_f;
  h_f = (func*)malloc(N*sizeof(func));
  cudaMalloc((void**)&d_f,N*sizeof(func));

  cudaMemcpyFromSymbol( &h_f[0], pfunc1, sizeof(func));
  cudaMemcpyFromSymbol( &h_f[1], pfunc1, sizeof(func));
  cudaMemcpyFromSymbol( &h_f[2], pfunc2, sizeof(func));
  cudaMemcpyFromSymbol( &h_f[3], pfunc3, sizeof(func));
  cudaMemcpyFromSymbol( &h_f[4], pfunc3, sizeof(func));

  cudaMemcpy(d_f,h_f,N*sizeof(func),cudaMemcpyHostToDevice);

  test_kernel<<<1,1>>>(d_f,N);

  cudaFree(d_f);
  free(h_f);

  return 0;
}

More Related Contents:

  1. What is the canonical way to check for errors using the CUDA runtime API?
  2. Unspecified launch failure on Memcpy
  3. Different CUDA versions shown by nvcc and NVIDIA-smi
  4. Thrust inside user written kernels
  5. GPU Emulator for CUDA programming without the hardware [closed]
  6. Using GPU from a docker container?
  7. How to measure the inner kernel time in NVIDIA CUDA?
  8. How do CUDA blocks/warps/threads map onto CUDA cores?
  9. Understanding CUDA grid dimensions, block dimensions and threads organization (simple explanation) [closed]
  10. What can I do against ‘CUDA driver version is insufficient for CUDA runtime version’?
  11. How to get the CUDA version?
  12. Can anyone provide sample code demonstrating the use of 16 bit floating point in cuda?
  13. Copy an object to device?
  14. What kind of variables consume registers in CUDA?
  15. In CUDA, what is memory coalescing, and how is it achieved?
  16. CUDA: How many concurrent threads in total?
  17. CUDA apps time out & fail after several seconds – how to work around this?
  18. multi-GPU basic usage
  19. How to use 2D Arrays in CUDA?
  20. How to create a CUDA context?
  21. CUBLAS: Incorrect inversion for matrix with zero pivot
  22. Using maximum shared memory in Cuda
  23. Why does my CUDA kernel crash (unspecified launch failure) with a different dataset size?
  24. Cuda Mutex, why deadlock?
  25. Inconsistency of IDs between ‘nvidia-smi -L’ and cuDeviceGetName()
  26. CUDA allocation alignment is 256 bytes – seriously?
  27. printf() in my CUDA kernel doesn’t result produce any output
  28. Multiply Rectangular Matrices in CUDA
  29. Which Compute Capability is supported by which CUDA versions?
  30. Using an array of device function pointers

Leave a Comment