CUDA based Spline Interpolation - PyCUDA

[kalosu]

Hello there friends from PyCUDA,

I am totally new to this community and also to CUDA. My experience so far with CUDA is really limited to just a few simple kernels that I wrote using numba cuda.

What I am trying to do right now is to use CUDA to perform spline interpolations of a 2D function (function of x and y).

While searching online, I have found this project:

https://github.com/DannyRuijters/CubicInterpolationCUDA/tree/master

from where apparently it is possible to use CUDA to perform bi cubic spline interpolation.

In theory the project is "ready to be used". However, as said above, my experience with CUDA is really limited so I am not really sure on how to approach using this project and integrating it into my own application.

What I would like to do is to use PyCUDA to load the cuda kernel via SourceModule I guess.

However, if you refer to the documentation from the interpolation project, you will see that they use something called textures that I am not really familiar with and I am also not so sure about how to use these from PyCUDA.

After asking chatGPT, it suggested to use the following CUDA code and then use SourceModule to interface this code:

texture<float, 2, cudaReadModeElementType> tex;

__device__ float interpolate_bicubic_2d(float x, float y) {
    return tex2D(tex, x, y);
}

extern "C" __global__ void interpolate_points_2d(float* output, float* coords, int num_points) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < num_points) {
        float x = coords[2 * idx];
        float y = coords[2 * idx + 1];
        output[idx] = interpolate_bicubic_2d(x, y);
    }
}

extern "C" void create_texture_from_2d_array(cudaArray* texArray, cudaChannelFormatDesc channelDesc, cudaExtent extent, cudaMemcpy3DParms copyParams) {
    cudaMallocArray(&texArray, &channelDesc, extent.width, extent.height);
    cudaMemcpyToArray(texArray, 0, 0, copyParams.srcPtr.ptr, copyParams.srcPtr.pitch * extent.height, cudaMemcpyDeviceToDevice);
    tex.addressMode[0] = cudaAddressModeClamp;
    tex.addressMode[1] = cudaAddressModeClamp;
    tex.filterMode = cudaFilterModeLinear;
    tex.normalized = false;
    cudaBindTextureToArray(tex, texArray, channelDesc);
}

How can I understand what is going on here? Do you think this approach is the correct one for interfacing to the interpolation kernel function from the projection I mentioned (in this case, it would be the function text2D)?

Do you maybe have some suggestion on what would be the recommended "workflow" to get this thing working?

Thanks a lot for your comments!!