Listed below is a basic code where I am trying to use a C++ wrapper to link in some CUDA Fortran code.
/* test_c.cpp */
#include <stdio.h>
#include <iostream>
extern "C" {
extern void fill_(int *);
}
int main(int argc, char **argv)
{
int a[100*100];
fill_(a);
return 0;
}
C test_f.CUF
module kernel
contains
!============================================================!
attributes(global) SUBROUTINE dGPURUN(VAL)
implicit none
INTEGER*4 VAL(:,:)
CALL dGPUFILL(VAL)
RETURN
END
!============================================================!
attributes(device) SUBROUTINE dGPUFILL(VAL)
implicit none
INTEGER*4 VAL(:,:),ROW,COL
ROW = threadidx%x +(blockidx%x-1)*blockdim%x
COL = threadidx%y +(blockidx%y-1)*blockdim%y
IF(ROW.LE.100.AND.COL.LE.100) THEN
VAL(ROW,COL) = ROW + COL
ENDIF
RETURN
END
end module kernel
!*************************************************************!
!*************************************************************!
SUBROUTINE FILL(VAL)
use cudafor
use kernel
implicit none
INTEGER*4 VAL(100,100)
INTEGER*4, device, dimension(100,100) :: dVAL
type(dim3) :: grid, tBlock
tBlock = dim3(32,32,1)
grid = dim3(ceiling(100./tBlock%x), ceiling(100./tBlock%y), 1)
dVAL = VAL
CALL dGPURUN<<<grid,tBlock>>>(dVAL)
VAL = dVAL
RETURN
END
When I try to run this I get the following error:
[shipmanb1@beehive TEST]$ echo $LD_LIBRARY_PATH
/opt/pgi/linux86-64/2014/cuda/6.0/lib64
[shipmanb1@beehive TEST]$ pgfortran -Mcuda -Mfixed -Mextend -c test_f.CUF
[shipmanb1@beehive TEST]$ pgc++ -c test_c.cpp
[shipmanb1@beehive TEST]$ pgc++ test_c.o test_f.o -o test.x -Mcuda -pgf90libs \
> -L/opt/pgi/linux86-64/14.7/lib -lcudafor -L/opt/pgi/linux86-64/2014/cuda/6.0/lib64 \
> -lcudart
[shipmanb1@beehive TEST]$ ./test.x
No CUDA device code available
But if I modify things to be just Fortran then it appears to work correctly:
C test.CUF
module kernel
contains
!============================================================!
attributes(global) SUBROUTINE dGPURUN(VAL)
implicit none
INTEGER*4 VAL(:,:)
CALL dGPUFILL(VAL)
RETURN
END
!============================================================!
attributes(device) SUBROUTINE dGPUFILL(VAL)
implicit none
INTEGER*4 VAL(:,:),ROW,COL
ROW = threadidx%x +(blockidx%x-1)*blockdim%x
COL = threadidx%y +(blockidx%y-1)*blockdim%y
IF(ROW.LE.100.AND.COL.LE.100) THEN
VAL(ROW,COL) = ROW + COL
ENDIF
RETURN
END
end module kernel
!*************************************************************!
!*************************************************************!
program test
use cudafor
use kernel
implicit none
INTEGER*4 VAL(100,100)
INTEGER*4, device, dimension(100,100) :: dVAL
type(dim3) :: grid, tBlock
tBlock = dim3(32,32,1)
grid = dim3(ceiling(100./tBlock%x), ceiling(100./tBlock%y), 1)
dVAL = VAL
CALL dGPURUN<<<grid,tBlock>>>(dVAL)
VAL = dVAL
write(*,*) 'VAL(3,3)=',VAL(3,3)
end program
[shipmanb1@beehive TEST]$ pgfortran -Mcuda -Mfixed -Mextend test.CUF -o test.x
[shipmanb1@beehive TEST]$ ./test.x
VAL(3,3)= 6
Is it possible to have a C++ code link in a CUDA Fortran code?
Thanks!