Cuda Fortran

Hey,

I tried to program my first CUDA Source with the PGI Cuda Fortran Compiler . The compute capability of the my card ist 1.0 (GeForce 8800 GTS). In the fortran settings i have set the compute capabilty to 1.0.
In the Emulation Mode my program works, but if i want to built the source without the emulation mode i get a windows message that my program does not work. I hope anyone can help me. My Code:

program prog

use m_ConDet

implicit none

real,allocatable,dimension(:,:) :: RPP
real,allocatable,dimension(:) :: test
integer,allocatable,dimension(:,:):: ContactList
integer :: NumRPP = 4
integer :: i,ii,NumPart

open(1,file=‘Partikel.txt’,status=‘UNKNOWN’)

read(1,*)

i=0
Do
read(1,*,END=100) ii
i=i+1
Enddo
100 rewind(1)

NumPart = i

allocate(RPP(NumRPP,NumPart))
allocate(ContactList(NumPart,NumPart))

ContactList = 1

read(1,)
Do i=1,NumPart
read(1,
)RPP(1:4,i)
enddo

close(1)

Call h_ConDet(NumPart,NumRPP,RPP,ContactList)

open(1,file=‘ContactList.txt’,status=‘UNKNOWN’)

Do i=1,NumPart
write(1,’(10i20.8)’)ContactList(i,:)
enddo
close(1)
end program prog

module m_ConDet

use cudafor

contains

!-----------kernel subroutine---------------
attributes(global) subroutine k_ConDet(NumPart,NumRPP,RPPdev,ContactListdev)

integer :: i,ii
integer,value :: NumRPP,NumPart
integer :: Px=1,Py=2,Pz=3,Rad=4
real :: RPPdev(NumRPP,NumPart)
integer :: ContactListdev(NumPart,NumPart)

i=(blockidx%x-1)blockdim%x+threadidx%x


if(i.lt.NumPart)then
do ii=1,NumPart
if(.not.ii.eq.i)then
dx = RPPdev(Px,i)-RPPdev(Px,ii)
dy = RPPdev(Py,i)-RPPdev(Py,ii)
dz = RPPdev(Pz,i)-RPPdev(Pz,ii)
d = sqrt(dx
dx+dydy+dzdz)
If(d.le.(RPPdev(Rad,i)+RPPdev(Rad,ii)))then
ContactListdev(i,ii) = 1
Endif
endif
enddo
endif

end subroutine k_ConDet

!-----------host subroutine---------------
subroutine h_ConDet(NumPart,NumRPP,RPP,ContactList)

implicit none

integer :: NumIPP,NumRPP,NumPart
type(dim3) :: dimGrid,dimBlock
real,dimension(:,:) :: RPP
integer,dimension(:,:) :: ContactList
real,device,allocatable,dimension(:,:) :: RPPdev
integer,device,allocatable,dimension(:,:) :: ContactListdev

allocate(RPPdev(NumRPP,NumPart),ContactListdev(NumPart,NumPart))

RPPdev = RPP

ContactListdev = 0

dimGrid = dim3(1,1,1)
dimBlock = dim3(NumPart/2,NumPart/2,1)

call k_ConDet<<<dimGrid,dimBlock>>>(NumPart,NumRPP,RPPdev,ContactListdev)

ContactList(:,:)=ContactListdev(:,:)

deallocate(RPPdev,ContactListdev)

end subroutine h_ConDet

end module m_ConDet

Hi JensWiese95812,

The error I’m getting is a compiler issue where were not handling the constant propagation correctly for your data initialized variables. I have added a problem report (TPR#16730) and sent it off to our engineers for further investigation.

The work around is to not data initialize these variables and instead initialize them in the subroutine body. For example:

attributes(global) subroutine k_ConDet(NumPart,NumRPP,RPPdev,ContactListdev)

integer :: i,ii
integer :: Px,Py,Pz,Rad
integer,value :: NumRPP,NumPart
real :: RPPdev(NumRPP,NumPart)
integer :: ContactListdev(NumPart,NumPart)

Px=1;Py=2;Pz=3;Rad=4
... continues

Hope this helps,
Mat

Thanks, that was quite easy. Now I can built my source, even if the code will not do what i want ;-).

Hi,

now i have a new probelm. If i want to compile my source (is the same source like before expecting the problem before) i get a Assembler message:

" Warning: rest of line ignored; first ignored character is `.’ "

Has anybody an idea where my mistake is!?

Second question: Is that the best way to allocate static arrays on the device? I think it is not quite easy to understand what kind of array allocation is the best one and where i allocate my arrays.

Thanks for your support

Hi JensWiese95812,

" Warning: rest of line ignored; first ignored character is `.’ "

I’m not able to recreate this with the above code. Have you made other changes? Can you please compile the code with “-Mkeepasm” and then review the resulting assembly file for the offending line? This may give us some clues as to the exact error.

Second question: Is that the best way to allocate static arrays on the device?

Sorry, I’m not clear on what you’re asking. Do you mean fixed size arrays (versus the dynamically allocated arrays you currently use) or do you mean global arrays that keep their values from one iteration to another?

For the former, you can use fixed size device arrays in the same manner as allocatable.

For the later, you can declare your device arrays in the module’s definition section to give them global scope. Currently, only fixed size arrays can be used, however, we should have support for allocatable arrays very soon.

  • Mat

Hello,

TPR 16730 has been corrected in an early 11.* release.


regards,
dave