View previous topic :: View next topic 
Author 
Message 
steve.xu
Joined: 20 Feb 2012 Posts: 25

Posted: Wed Mar 14, 2012 4:41 am Post subject: cuda fortran donot support an array of negative index?? 


hi,everyone
I got the following code and the value of "rhs_h(1)" cannot be changed by the kernel.
module mod_cuda_test
use cudafor
implicit none
integer,parameter::kind_int=4
integer,parameter::kind_real=4
contains
attributes(global) subroutine calc_I3(rhs_d)
implicit none
real(kind_real),intent(inout)::rhs_d(:)
rhs_d(1)=3.0
end subroutine calc_I3
end module mod_cuda_test
program test_kernel
use cudafor
use mod_cuda_test
implicit none
real(kind_real),allocatable::rhs_h(:)
real(kind_real),allocatable,device::rhs_d(:)
allocate(rhs_d(1:10))
allocate(rhs_h(1:10))
rhs_h=2.0
rhs_d=rhs_h
call calc_I3<<<1>>>(rhs_d)
rhs_h=rhs_d
write(*,*),"rhs_d(1) is",rhs_h(1)
end program
when i run the code on Geforce 9400 (both driver and runtime version are 3.2) with pgi 11.8, rhs_h(1) is still 2.0.
But if i change
allocate(rhs_d(1:10))
allocate(rhs_h(1:10))
to
allocate(rhs_d(1:10))
allocate(rhs_h(1:10))
rhs_h(1) is changed to 3.0
Anybody know if cuda fortran donot support an array of negative index??? 

Back to top 


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

Posted: Wed Mar 14, 2012 11:11 am Post subject: 


Hi Steve,
Quote:  Anybody know if cuda fortran donot support an array of negative index???  Negative indices are supported, however this looks like a compiler error when passing in a array with a negative lower bound as a argument. I sent a report to our engineers (TPR#18545) for further investigation.
The work around is to declare "rhs_d" as a module variable.
Code: 
% cat neg2.cuf
module mod_cuda_test
use cudafor
implicit none
integer,parameter::kind_int=4
integer,parameter::kind_real=4
real(kind_real),allocatable,device::rhs_d(:)
contains
attributes(global) subroutine calc_I3()
implicit none
rhs_d(1)=3.0
end subroutine calc_I3
end module mod_cuda_test
program test_kernel
use cudafor
use mod_cuda_test
implicit none
real(kind_real),allocatable::rhs_h(:)
allocate(rhs_d(1:10))
allocate(rhs_h(1:10))
rhs_h=2.0
rhs_d=rhs_h
call calc_I3<<<1>>>()
rhs_h=rhs_d
write(*,*),"rhs_d(1) is",rhs_h(1)
end program
% pgf90 neg2.cuf V12.3 ; a.out
rhs_d(1) is 3.000000

Thanks,
Mat 

Back to top 


steve.xu
Joined: 20 Feb 2012 Posts: 25

Posted: Thu Mar 15, 2012 2:38 am Post subject: 


thanks Mat.
But i still cannot get the right result when rhs_d is changed to a module device variable. Fortunatly if the index range of rhs_d is declared explicitly in the definition of kernel function, i can see it takes effect. Below is my code:
module mod_cuda_test
use cudafor
implicit none
integer,parameter::kind_int=4
integer,parameter::kind_real=4
contains
attributes(global) subroutine calc_I3(rhs_d,n1,n2)
implicit none
integer(kind_int),value::n1,n2
real(kind_real),intent(inout)::rhs_d(n1:n2)
rhs_d(1)=3.0
end subroutine calc_I3
end module mod_cuda_test
program test_kernel
use cudafor
use mod_cuda_test
implicit none
real(kind_real),allocatable::rhs_h(:)
real(kind_real),allocatable,device::rhs_d(:)
integer(kind_int)::n1,n2
n1=1
n2=10
allocate(rhs_d(n1:n2))
allocate(rhs_h(n1:n2))
rhs_h=2.0
rhs_d=rhs_h
call calc_I3<<<1>>>(rhs_d,n1,n2)
rhs_h=rhs_d
write(*,*),"after call cacl_I"
write(*,*),"rhs_d(1) is",rhs_h(1)
end program 

Back to top 


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

Posted: Thu Apr 12, 2012 2:53 pm Post subject: 


Hi Steve,
FYI, TPR#18545 has been closed. The engineer investigating the issue determined that the original program was in error. He notes:
Quote:  Assumedshape dummy arrays do not assume the lower bound of the actual argument. The assumedshape array gets the extent (the shape) from the actual argument, but the lower bound is defined by the declaration. The correct program would declare the argument as "real(kind_real),intent(inout)::rhs_d(1:)"
I quote from "Fortran 95/2003 explained" page 100, section 6.3 "Assumedshape arrays" "When the shape is declared by the dimension clause, each dimension has the form: [lowerbound]: where lowerbound is an integer expression that may depend on module data or the other arguments (...). If lowerbound is omitted, the default value is 1. Note that it is the shape that is passed, and not the upper and lower bounds." 
So your solution is of passing both the lower and upper bounds is correct.
 Mat 

Back to top 


steve.xu
Joined: 20 Feb 2012 Posts: 25

Posted: Mon Apr 16, 2012 12:29 am Post subject: 


Thanks Mat!
mkcolg wrote:  Hi Steve,
FYI, TPR#18545 has been closed. The engineer investigating the issue determined that the original program was in error. He notes:
Quote:  Assumedshape dummy arrays do not assume the lower bound of the actual argument. The assumedshape array gets the extent (the shape) from the actual argument, but the lower bound is defined by the declaration. The correct program would declare the argument as "real(kind_real),intent(inout)::rhs_d(1:)"
I quote from "Fortran 95/2003 explained" page 100, section 6.3 "Assumedshape arrays" "When the shape is declared by the dimension clause, each dimension has the form: [lowerbound]: where lowerbound is an integer expression that may depend on module data or the other arguments (...). If lowerbound is omitted, the default value is 1. Note that it is the shape that is passed, and not the upper and lower bounds." 
So your solution is of passing both the lower and upper bounds is correct.
 Mat 


Back to top 




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
