# How to copy a host function pointer to device in CUDA

**URL:** <https://forums.developer.nvidia.com/t/how-to-copy-a-host-function-pointer-to-device-in-cuda/217771>\
**Category:** CUDA Programming and Performance\
**Tags:** cuda\
**Created:** [June 15, 2022, 2:04pm UTC](https://forums.developer.nvidia.com/t/how-to-copy-a-host-function-pointer-to-device-in-cuda/217771 "2022-06-15T14:04:06Z")\
**Posts on this page:** 7\
**Page:** 1

<div class="post-metadata">

**Author:** ![Narrow](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Narrow](https://forums.developer.nvidia.com/u/Narrow)\
**Post date:** [June 15, 2022, 2:04pm UTC](https://forums.developer.nvidia.com/t/how-to-copy-a-host-function-pointer-to-device-in-cuda/217771/1 "2022-06-15T14:04:06Z")

</div>

Hello everyone.  
I recently started porting a cpp matrix library to cuda code in order to accelerate the matrix operations.

in the old cpu-based code, i had a function in the matrix class, .apply(double (\*func)(double)), that applied a function to each element of the matrix;

```cpp
Matrix Matrix::apply(double (*func)(double)) const
{
    Matrix applied(this->rows, this->cols);

    for (int i = 0; i < this->rows; i++)
        for (int j = 0; j < this->cols; j++)
            applied.data[i][j] = func(this->data[i][j]);
    
    return applied;
}

```

Now, how would i be able to sort of “copy” the function pointer from host to device, in order to give it to a kernel that does the exact same thing only parallelized between multiple threads like this:

```auto
void applyWrapper(double** A, double (*hostFuncPtr)(double), size_t rows, size_t cols
{
     auto deviceFuncPtr = /* copy the function pointed by hostFuncPtr into a device function pointer */
     applyKernel<<<1, 1>>>(A, deviceFuncPtr, rows, cols);
}

void applyKernel(double** A, double (*devFuncPtr)(double), size_t rows, size_t cols)
{
    size_t idx = threadIdx.x;
    size_t stride = blockDim.x;

    for (size_t i = idx; i < A_cols; i += stride)
        for (size_t j = 0; j < A_rows; j++)
            A[j][i] = func(A[j][i]);
}

```

---

<div class="post-metadata">

**Author:** ![TomNVIDIA](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/tomnvidia/32/14181_2.png) [@TomNVIDIA](https://forums.developer.nvidia.com/u/TomNVIDIA)\
**Post date:** [June 15, 2022, 3:10pm UTC](https://forums.developer.nvidia.com/t/how-to-copy-a-host-function-pointer-to-device-in-cuda/217771/2 "2022-06-15T15:10:43Z")

</div>

Hi,

This is the Networking category, this issue should be posted in one of the CUDA forums. I will move it over for you.

---

<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:** [June 15, 2022, 7:20pm UTC](https://forums.developer.nvidia.com/t/how-to-copy-a-host-function-pointer-to-device-in-cuda/217771/3 "2022-06-15T19:20:40Z")

</div>

cuda sample codes include a sample demonstrating function pointer usage

any function pointer you want to use in device code had better be obtained from a properly decorated device function. You cannot pass a pointer that points to a host function and expect to do anything sensible with it in device code.

The cuda sample code will make this evident.

---

<div class="post-metadata">

**Author:** ![jccporter](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@jccporter](https://forums.developer.nvidia.com/u/jccporter)\
**Post date:** [July 7, 2024, 4:33pm UTC](https://forums.developer.nvidia.com/t/how-to-copy-a-host-function-pointer-to-device-in-cuda/217771/4 "2024-07-07T16:33:06Z")

</div>

Would you kindly be able to provide a link to this sample code?

---

<div class="post-metadata">

**Author:** ![rs277](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@rs277](https://forums.developer.nvidia.com/u/rs277)\
**Post date:** [July 7, 2024, 6:45pm UTC](https://forums.developer.nvidia.com/t/how-to-copy-a-host-function-pointer-to-device-in-cuda/217771/5 "2024-07-07T18:45:19Z")

</div>

> **[cuda-samples/Samples/2\_Concepts\_and\_Techniques/FunctionPointers at master ·...](https://github.com/NVIDIA/cuda-samples/tree/master/Samples/2_Concepts_and_Techniques/FunctionPointers)**
>
> Samples for CUDA Developers which demonstrates features in CUDA Toolkit - NVIDIA/cuda-samples

---

<div class="post-metadata">

**Author:** ![Curefab](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Curefab](https://forums.developer.nvidia.com/u/Curefab)\
**Post date:** [July 8, 2024, 4:38pm UTC](https://forums.developer.nvidia.com/t/how-to-copy-a-host-function-pointer-to-device-in-cuda/217771/6 "2024-07-08T16:38:42Z")

</div>

If the function pointer at some point is known at compile-time, you could create a template parameter. For example replace the function pointer with a class type with a suitable member function.

---

<div class="post-metadata">

**Author:** ![warpdivergence77](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@warpdivergence77](https://forums.developer.nvidia.com/u/warpdivergence77)\
**Post date:** [July 10, 2024, 12:37pm UTC](https://forums.developer.nvidia.com/t/how-to-copy-a-host-function-pointer-to-device-in-cuda/217771/7 "2024-07-10T12:37:12Z")

</div>

You can’t pass a pointer to host function to the device, the address spaces are separate. Trying to call that in device code would be nonsensical.