# Calling a CUDA kernel in an external library

**URL:** <https://discourse.julialang.org/t/calling-a-cuda-kernel-in-an-external-library/46505>\
**Category:** GPU\
**Created:** [September 12, 2020, 10:47am UTC](https://discourse.julialang.org/t/calling-a-cuda-kernel-in-an-external-library/46505 "2020-09-12T10:47:50Z")\
**Posts on this page:** 7\
**Page:** 1

<div class="post-metadata">

**Author:** ![Niclas\_Wiberg](https://sea2.discourse-cdn.com/julialang/user_avatar/discourse.julialang.org/niclas_wiberg/32/19063_2.png) [@Niclas\_Wiberg](https://discourse.julialang.org/u/Niclas_Wiberg)\
**Post date:** [September 12, 2020, 10:47am UTC](https://discourse.julialang.org/t/calling-a-cuda-kernel-in-an-external-library/46505/1 "2020-09-12T10:47:50Z")

</div>

Is it possible to call a precompiled CUDA kernel from Julia?

I am writing CUDA kernels in C++ and compiling with nvcc, for use in a C++ application. I use Julia to test the code. Currently I write host-side wrapping functions that I call from Julia using ccall. I would like to avoid the wrapping functions and call the kernels directly from Julia.

Is this possible?

---

<div class="post-metadata">

**Author:** ![maleadt](https://sea2.discourse-cdn.com/julialang/user_avatar/discourse.julialang.org/maleadt/32/10097_2.png) [@maleadt](https://discourse.julialang.org/u/maleadt)\
**Post date:** [September 13, 2020, 10:03am UTC](https://discourse.julialang.org/t/calling-a-cuda-kernel-in-an-external-library/46505/2 "2020-09-13T10:03:04Z")

</div>

If you can compile to a PTX file (`nvcc -pxt`), you can load the file with `CuModule`, look up a kernel, and `cudacall` it. So pretty low level. It should also be possible to extract the PTX code from the fatbin, but we don’t have any functionality for that.

---

<div class="post-metadata">

**Author:** ![Niclas\_Wiberg](https://sea2.discourse-cdn.com/julialang/user_avatar/discourse.julialang.org/niclas_wiberg/32/19063_2.png) [@Niclas\_Wiberg](https://discourse.julialang.org/u/Niclas_Wiberg)\
**Post date:** [September 14, 2020, 2:54pm UTC](https://discourse.julialang.org/t/calling-a-cuda-kernel-in-an-external-library/46505/3 "2020-09-14T14:54:51Z")

</div>

Thanks, I might try that. Do you know if the low-level code is generated using Nvidias ptxas in that case, or with some other technique?

I am also considering using cudaLaunchKernel(…) from the CUDA runtime, which might allow me to use an already compiled kernel. Do you see any pitfalls with this?

---

<div class="post-metadata">

**Author:** ![maleadt](https://sea2.discourse-cdn.com/julialang/user_avatar/discourse.julialang.org/maleadt/32/10097_2.png) [@maleadt](https://discourse.julialang.org/u/maleadt)\
**Post date:** [September 14, 2020, 3:03pm UTC](https://discourse.julialang.org/t/calling-a-cuda-kernel-in-an-external-library/46505/4 "2020-09-14T15:03:06Z")

</div>

> [@Niclas\_Wiberg](#):
>
> I am also considering using cudaLaunchKernel(…) from the CUDA runtime, which might allow me to use an already compiled kernel. Do you see any pitfalls with this?

That won’t work, it expects a certain layout of the binary in order to look up the compiled function. The Julia binary doesn’t have that, and we don’t emit our code like that. Unless you want to call `cudaLaunchKernel` from a pre-compiled binary, but then you’re in the same boat (having to `ccall` an external library).

---

<div class="post-metadata">

**Author:** ![Niclas\_Wiberg](https://sea2.discourse-cdn.com/julialang/user_avatar/discourse.julialang.org/niclas_wiberg/32/19063_2.png) [@Niclas\_Wiberg](https://discourse.julialang.org/u/Niclas_Wiberg)\
**Post date:** [September 14, 2020, 3:23pm UTC](https://discourse.julialang.org/t/calling-a-cuda-kernel-in-an-external-library/46505/5 "2020-09-14T15:23:09Z")

</div>

Hmm, I’m confused.

A small example:

```julia
__global__
void example_kernel(int *data)
{
    data[3] = data[0] + data[1] * data[2];
}

extern "C"
void example_wrapper(int *data)
{
    int *data_cuda;
    size_t size = 4 * sizeof(int);
    cudaMalloc(&data_cuda, size);
    cudaMemcpy(data_cuda, data, size, cudaMemcpyHostToDevice);
    example_kernel<<<1,1>>>(data_cuda);
    cudaDeviceSynchronize();
    cudaMemcpy(data, data_cuda, size, cudaMemcpyDeviceToHost);
    cudaFree(data_cuda);
}

```

I compile this using `nvcc -Xcompiler -fPIC --shared example_kernel.cu -o example_kernel.so`  
And then in Julia:

```julia
julia> data=Cint[1,2,3,0]
4-element Array{Int32,1}:
 1
 2
 3
 0

julia> ccall((:example_wrapper,"./example_kernel.so"),Cvoid,(Ptr{Cint},),pointer(data))

julia> data
4-element Array{Int32,1}:
 1
 2
 3
 7

```

The shared object `example_kernel.so` should now contain the kernel code in the format that the CUDA runtime expects. So I would guess that `cudaLaunchKernel(...)` should be able to launch it, if I’m able to invoke it from within Julia.

Am I wrong?

---

<div class="post-metadata">

**Author:** ![maleadt](https://sea2.discourse-cdn.com/julialang/user_avatar/discourse.julialang.org/maleadt/32/10097_2.png) [@maleadt](https://discourse.julialang.org/u/maleadt)\
**Post date:** [September 14, 2020, 3:34pm UTC](https://discourse.julialang.org/t/calling-a-cuda-kernel-in-an-external-library/46505/6 "2020-09-14T15:34:37Z")

</div>

> [@Niclas\_Wiberg](#):
>
> So I would guess that `cudaLaunchKernel(...)` should be able to launch it, if I’m able to invoke it from within Julia.

With what arguments? `cudaLaunchKernel` takes a function pointer, which is resolved within the executing application, and AFAIK depends on the executable having specific symbols and state set-up.

---

<div class="post-metadata">

**Author:** ![Niclas\_Wiberg](https://sea2.discourse-cdn.com/julialang/user_avatar/discourse.julialang.org/niclas_wiberg/32/19063_2.png) [@Niclas\_Wiberg](https://discourse.julialang.org/u/Niclas_Wiberg)\
**Post date:** [September 14, 2020, 5:48pm UTC](https://discourse.julialang.org/t/calling-a-cuda-kernel-in-an-external-library/46505/7 "2020-09-14T17:48:44Z")

</div>

> [@maleadt](#):
>
> With what arguments? `cudaLaunchKernel` takes a function pointer, which is resolved within the executing application, and AFAIK depends on the executable having specific symbols and state set-up.

Fair point, I don’t know how to get that function pointer. Maybe I can create a single C function that does it for me. Will investigate and come back. Thanks for the feedback.