PGI User Forum
 SearchSearch   MemberlistMemberlist     RegisterRegister   ProfileProfile    Log inLog in 

CUDA-x86.

Linking CUDA fortran compiled code with ifort

 
Post new topic   Reply to topic    PGI User Forum Forum Index -> Programming and Compiling
View previous topic :: View next topic  
Author Message
wiersma



Joined: 16 May 2013
Posts: 29

PostPosted: Thu May 16, 2013 2:15 pm    Post subject: Linking CUDA fortran compiled code with ifort Reply with quote

Hi all,

Please bear with me if this is obvious - I wouldn't describe myself as an expert in compilers or with Fortran by any stretch of the imagination.

As a part of evaluating whether CUDA would suit my needs, I tried just dropping one of the matrix multiplication routines (say, here http://geco.mines.edu/software/pg10/gpu/pgicudaforug.pdf) into my code suite and see what happens. I'm having difficulty compiling though.

Normally my code is arranged in modules by file, so would have a file named mod1.f90:

Code:

module mod1
!lblah blah blah
end module mod1


And then test.f90 containing:
Code:

program test
use mod1
end program test


Which I compile using
Code:

ifort -c mod1.f90
ifort -c test.f90
ifort -o test mod1.o test.o


Simple right? Now if I have a cuda fortran module named, say cmod1.cuf:

Code:

module cmod1
use cudafor
!lblah blah blah
end module cmod1


I try compiling:
Code:

pgf90 -c cmod1.cuf
ifort -c test.f90
ifort -o test cmod1.o test.o


But get
Quote:
test.f90(4): error #7013: This module file was not generated by any release of this compiler. [cmod1]
use cmod1
-------^

When trying to compile test.f90. So ifort doesn't like pgf90 compiled modules? Do I have to compile everything with pgf90, because our suite uses some ifort specific things.

(Sorry if this was incredibly verbose, I just wanted to be clear :)).
Back to top
View user's profile
mkcolg



Joined: 30 Jun 2004
Posts: 6134
Location: The Portland Group Inc.

PostPosted: Thu May 16, 2013 5:19 pm    Post subject: Reply with quote

Hi wiersma,

Quote:
When trying to compile test.f90. So ifort doesn't like pgf90 compiled modules? Do I have to compile everything with pgf90, because our suite uses some ifort specific things.
Several F90 features such as modules and allocatable arrays are not compatible between Fortran compilers.

What you'll need to do is put the CUDA Fortran routines into a library and then have the main Fortran program call the routines via a F77 or C wrapper.

Hope this helps,
Mat
Back to top
View user's profile
wiersma



Joined: 16 May 2013
Posts: 29

PostPosted: Tue May 21, 2013 1:45 pm    Post subject: Reply with quote

Hi all,

Thanks for the reply, but now I'm having difficulty setting up a library.

So now I have a pretty basic example:

cmod1.cuf:
Code:

module cmod1
use cudafor

contains

subroutine mmul( A, B, C )
    real, dimension(:,:) :: A, B, C
end subroutine mmul
end module cmod1


test.f90:
Code:

program test

    implicit none
    real  :: A(10,10), B(10,10), C(10,10)
    A = 0.
    B = 0.
    C = 0.
    call mmul(A,B,C)
   
end program test


The commands:
Code:

pgf90 -c cmod1.cuf
ar rcvs libcmod1.a cmod1.o
ifort -c test.f90
ifort test.o -L. -lcmod1 -o test


give:
Code:

test.o: In function `MAIN__':
test.f90:(.text+0x750): undefined reference to `mmul_'


Sorry if this is basic library stuff - I've tried multiple configurations, but can't seem to figure this out.
Back to top
View user's profile
mkcolg



Joined: 30 Jun 2004
Posts: 6134
Location: The Portland Group Inc.

PostPosted: Wed May 22, 2013 11:49 am    Post subject: Reply with quote

Hi wiersma,

So I had never actually tried calling CUDA Fortran code from Intel compiled code. Turned out to be a bit more tricky then I thought. We do some magic during the initialization of the program that's necessary to get the device working properly. Hence, you have to link with the PGI driver to use CUDA Fortran. Here's the steps I did:

Code:
% cat prog.f90
! just call the main test program
program testp
   call test()
end program testp

% cat test.f90
! change the Intel main program to a subroutine
subroutine test
    implicit none
    real  :: A(10,10), B(10,10), C(10,10)
    A = 10.1
    B = 0.11
    C = 0.
    call cmod1_mmul(A,B,C,10,10)
    print *, A(1,2), B(2,4), C(3,6)
   
end subroutine test

% cat cmod1.cuf
! Here's a basic CUDA Fortran kernel

module cmod1
use cudafor

contains

attributes(global) subroutine kernel( A, B, C, N, M )
 real, device :: A(N,M), B(N,M), C(N,M)
 integer, value :: N, M
 integer :: i, j, tx, ty
 ! Get the thread indices
 tx = threadidx%x
 ty = threadidx%y
 i = (blockidx%x-1) * blockdim%x + tx
 j = (blockidx%y-1) * blockdim%y + ty
 if (i .le. N .and. j .le. M) then
   C(i,j) = A(i,j)+ B(i,j)
 endif
end subroutine kernel

! need to pass in arrays a automatic or assumed-shaped
subroutine mmul( A, B, C, N, M)
    implicit none
    real, dimension(n,m) :: A, B, C
    real, device, allocatable, dimension(:,:) :: Ad, Bd, Cd
    integer i,N,M,ierr
    type(dim3) :: dimGrid, dimBlock

    allocate(Ad(N,M), Bd(N,M), Cd(N,M))
    Ad=A
    Bd=B
    dimGrid = dim3( N/16, M/16, 1 )
    dimBlock = dim3( 16, 16, 1 )
    call kernel<<<dimGrid,dimBlock>>>( Ad, Bd, Cd, N, M )
    C=Cd
    deallocate(Ad, Bd, Cd)

end subroutine mmul
end module cmod1

% pgf90 -c cmod1.cuf
% ifort -c test.f90
% pgf90 -c prog.f90
% pgf90 -Mcuda prog.o test.o cmod1.o -L/opt/intel/composerxe-2013.3.163/compiler/lib/intel64/ -lifport -lifcore -limf
% a.out
 N=           10  M=           10  ierr=            0
   10.10000      0.1100000       10.21000   


You can also wrap-up "cmod1.o" into a library.

- Mat
Back to top
View user's profile
wiersma



Joined: 16 May 2013
Posts: 29

PostPosted: Mon May 27, 2013 2:09 pm    Post subject: Reply with quote

Works great! Thanks a bunch and I'm glad to know I wasn't missing something completely obvious :).
Back to top
View user's profile
Display posts from previous:   
Post new topic   Reply to topic    PGI User Forum Forum Index -> Programming and Compiling All times are GMT - 7 Hours
Page 1 of 1

 
Jump to:  
You cannot post new topics in this forum
You cannot reply to topics in this forum
You cannot edit your posts in this forum
You cannot delete your posts in this forum
You cannot vote in polls in this forum


Powered by phpBB © phpBB Group