Fortran cuda kernels for pycuda

[w9cf]

I have been using nvfortran to write Fortran cuda kernels called from Fortran. Our group would like to use some these kernels in python code using pycuda. I have successfully done this by compiling and linking the Fortran code with the keep option, and then loading the cubin file. I think it would be useful if that feature could be included with the pycuda as an option using pycuda.compiler.SourceModule or equivalent.

Is this something that would be of interest?

Here is a minimal example of what I am currently doing:
Start with examples/hello_gpu.py, and change the mod = SourceModule line from

mod = SourceModule("""
__global__ void multiply_them(float *dest, float *a, float *b)
{
  const int i = threadIdx.x;
  dest[i] = a[i] * b[i];
}
""")

to

mod = drv.module_from_file("kernel.cubin")

Write the cuda kernel in fortran with a dummy main program. The file name is multiply.f90

module cudakernelmod

   use, intrinsic :: iso_fortran_env
   implicit none

contains

   attributes (global) subroutine multiply_them(dest,a,b) bind(c)
   real(kind=real32), intent(out) :: dest(*)
   real(kind=real32), intent(in) :: a(*),b(*)
   integer(kind=int32) :: i
   i=threadIdx%x
   dest(i)=a(i)*b(i)
   end subroutine multiply_them

end module cudakernelmod

program nullprog
end program nullprog

Compile the code to an executable keeping all the intermediate files

nvfortran -cuda -O3 -fast -gpu:ccnative,keep multiply.f90

This creates a bunch of files including, for compute capability 7.5, pgcuda.75.cubin which I rename to kernel.cubin

mv pgcuda.??.cubin kernel.cubin

Running

python hello_gpu_fortran.py

prints out the array of zeroes as in the original example.

[inducer]

That sounds like a useful extension. I would be happy to review a PR along those lines.