Undefined symbol for cuda kernel function with pgfortran

I try to compile a cuda/fortran example code :

$ cat mmul_orig.cuf
attributes(GLOBAL) subroutine mmul_kernel( A, B, C, n, m, l)
(…)
end subroutine mmul_kernel

!-------------------------------------------------------------------------------
!-------------------------------------------------------------------------------
! The host routine to drive the matrix multiplication
subroutine mmul( N,M,L, A, B, C )

use cudafor
(…)
call mmul_kernel<<<dimGrid,dimBlock>>>( Adev, Bdev, Cdev, N, M, L )
(…)
end subroutine mmul
$ pgfortran -Mcuda=cc13 -c mmul_orig.cuf

I get an object file with an undefined symbol :
$ nm mmul_orig.o
(…)
(…)
0000000000000280 T mmul_
U mmul_kernel_
0000000000000000 T mmul_kernel___entry

If I change the source file and defined the kernel function in a fortran module ‘kernel’, there is no undefined symbol anymore :

$ cat mmul.cuf
module kernel
use cudafor
contains
attributes(GLOBAL) subroutine mmul_kernel( A, B, C, n, m, l)
(…)
end subroutine mmul_kernel
end module kernel

!-------------------------------------------------------------------------------
!-------------------------------------------------------------------------------
! The host routine to drive the matrix multiplication
subroutine mmul( N,M,L, A, B, C )

use cudafor
use kernel
(…)
call mmul_kernel<<<dimGrid,dimBlock>>>( Adev, Bdev, Cdev, N, M, L )
(…)
end subroutine mmul
$ pgfortran -Mcuda=cc13 -c mmul.cuf
$ nm mmul.o
(…)
0000000000000000 T kernel_
0000000000000010 T kernel_mmul_kernel___entry
0000000000000290 T mmul_
(…)


What did I do wrong in mmul_orig.cuf to get this undefined symbol “mmul_kernel_” ?
Regards,

Olivier

Hi Olivier,

CUDA Fortran kernel routines must have explicit interfaces. If you added one for mmul_kernel in mmul, the code will link correctly.

However, I recommend that you put your CUDA Fortran Kernels into a module. Modules have the benefit of creating an interface for you (and is why yours linked correctly in the second example).

Hope this helps,
Mat