# Array of function pointers assignment

**URL:** <https://forums.developer.nvidia.com/t/array-of-function-pointers-assignment/208952>\
**Category:** CUDA Programming and Performance\
**Created:** [March 21, 2022, 1:52pm UTC](https://forums.developer.nvidia.com/t/array-of-function-pointers-assignment/208952 "2022-03-21T13:52:49Z")\
**Posts on this page:** 8\
**Page:** 1

<div class="post-metadata">

**Author:** ![chanokin](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@chanokin](https://forums.developer.nvidia.com/u/chanokin)\
**Post date:** [March 21, 2022, 1:52pm UTC](https://forums.developer.nvidia.com/t/array-of-function-pointers-assignment/208952/1 "2022-03-21T13:52:50Z")

</div>

Hi,

I’m trying to build a framework for simulations and would like to keep it sufficiently generic for ‘end users’ to add functions. To that end, I think having a global function pointers array would be ideal. So, I’m trying to do the following (pseudo-coded):

```auto
typedef func_ptr_t X(*)(Y, Z)

__device__ func_ptr_t *functions[2];

__device__ X f(Y y, Z z);

main(){
    functions[0] = f;
}

```

This clearly doesn’t work because of the ` __device__ ` variables not being allowed to be used on the host side. So, I added a kernel to ‘assign’ the functions.

```auto
typedef func_ptr_t X(*)(Y, Z)

__device__ func_ptr_t *functions[2];

__device__ X f(Y y, Z z);

__device__ assign_funcs(index, func_ptr_t fptr){
    functions[index] = fptr;
}

main(){
   assign_funcs<<<1, 1>>>(0, f);
}

```

This now compiles but, when I try to call the function, I get an `invalid program counter` error.

I’m not sure what I’m doing wrong or if this is the right way of doing things. I’ve read some of the posts on similar problems but I can’t grasp what my mistake is. Would anyone care to lend a hand?

---

<div class="post-metadata">

**Author:** ![striker159](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@striker159](https://forums.developer.nvidia.com/u/striker159)\
**Post date:** [March 21, 2022, 8:38pm UTC](https://forums.developer.nvidia.com/t/array-of-function-pointers-assignment/208952/2 "2022-03-21T20:38:52Z")

</div>

It’s not possible like that. See [Programming Guide :: CUDA Toolkit Documentation](https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#functions)  
`It is not allowed to take the address of a __device__ function in host code. `

But it’s possible to pass ` __device__ ` lambdas to kernels. You may get something working if you use nvstd::function instead of raw function pointers.

---

<div class="post-metadata">

**Author:** ![chanokin](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@chanokin](https://forums.developer.nvidia.com/u/chanokin)\
**Post date:** [March 22, 2022, 10:31am UTC](https://forums.developer.nvidia.com/t/array-of-function-pointers-assignment/208952/3 "2022-03-22T10:31:56Z")

</div>

So I could do a lambda wrapper to the device function?

---

<div class="post-metadata">

**Author:** ![Robert\_Crovella](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/robert_crovella/32/14043_2.png) [@Robert\_Crovella](https://forums.developer.nvidia.com/u/Robert_Crovella)\
**Post date:** [March 22, 2022, 2:19pm UTC](https://forums.developer.nvidia.com/t/array-of-function-pointers-assignment/208952/4 "2022-03-22T14:19:18Z")

</div>

You can get what you have working by capturing the device pointer in device code, copying that pointer to host code, and then dispatch with that pointer.

See [here](https://stackoverflow.com/questions/31057870/passing-host-function-as-a-function-pointer-in-global-or-device-function/31058123#31058123)

You can also use lambdas. However every lambda has a unique type. So its difficult or impossible to create an array of lambdas in a useful fashion.

---

<div class="post-metadata">

**Author:** ![chanokin](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@chanokin](https://forums.developer.nvidia.com/u/chanokin)\
**Post date:** [March 23, 2022, 9:47am UTC](https://forums.developer.nvidia.com/t/array-of-function-pointers-assignment/208952/5 "2022-03-23T09:47:38Z")

</div>

So it would be something like this?

```auto
typedef func_ptr_t X(*)(Y, Z)

__device__ func_ptr_t *functions[2];

__device__ X f(Y y, Z z);

__device__ assign_funcs(index, func_ptr_t fptr){
    functions[index] = fptr;
}

main(){
  func_ptr_t f_h = null;

  cudaMemcpyFromSymbol(&f_h, &f, sizeof(void *));

   assign_funcs<<<1, 1>>>(0, f_h);
}

```

---

<div class="post-metadata">

**Author:** ![Robert\_Crovella](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/robert_crovella/32/14043_2.png) [@Robert\_Crovella](https://forums.developer.nvidia.com/u/Robert_Crovella)\
**Post date:** [March 23, 2022, 1:38pm UTC](https://forums.developer.nvidia.com/t/array-of-function-pointers-assignment/208952/6 "2022-03-23T13:38:47Z")

</div>

Not quite. You’re still _taking the address of a device function in host code_ here:

```
cudaMemcpyFromSymbol(&f_h, &f, sizeof(void *));
                           ^^

```

and additionally, `cudaMemcpyFromSymbol` is designed to copy data from a device variable, not from a device function address.

so _capture the device function address in device code_, into a device variable. Then copy that variable to host code. The link I gave you already has such an example. [here](https://stackoverflow.com/questions/30002353/function-pointer-to-other-kernel-as-kernel-arg-in-cuda) it is again.

---

<div class="post-metadata">

**Author:** ![chanokin](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@chanokin](https://forums.developer.nvidia.com/u/chanokin)\
**Post date:** [March 23, 2022, 3:57pm UTC](https://forums.developer.nvidia.com/t/array-of-function-pointers-assignment/208952/7 "2022-03-23T15:57:56Z")

</div>

You Sir, are a genius! Thanks for your help. Just for completion’s sake, here’s the full pseudocode.

```auto
// define a function type
typedef func_ptr_t X(*)(Y, Z)

// device side storage for this typeof functions
__device__ func_ptr_t *functions[2];

// an instance of this function type
__device__ X my_func(Y y, Z z);

// Cuda magic where we get the address of the 
// function straight to a pointer
__device__ func_ptr_t capture = my_func;

// actually store the (host-stored) pointer to device address
__global__ assign_funcs(index, func_ptr_t fptr){
    functions[index] = fptr;
}

main(){
  // temporary storage for function address
  func_ptr_t address_stored_on_host;
  
  // copy device function address to temporary storage
  cudaMemcpyFromSymbol(&address_stored_on_host, 
                       capture, sizeof(func_ptr_t));

   // set the address on the function addresses array
   assign_funcs<<<1, 1>>>(0, address_stored_on_host);
}

```

---

<div class="post-metadata">

**Author:** ![system](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/system/32/68080_2.png) [@system](https://forums.developer.nvidia.com/u/system)\
**Post date:** [April 12, 2022, 10:59am UTC](https://forums.developer.nvidia.com/t/array-of-function-pointers-assignment/208952/8 "2022-04-12T10:59:44Z")

</div>

This topic was automatically closed 14 days after the last reply. New replies are no longer allowed.