not enough memory

I am new to PVF, and have some problems with the memory management. The compiler doesn’t give and error message but when I run the exe, the allocatable array that requires more than 2gb creates problem and it gives me the not enough memory message.
Is there any way to increase the memory?
I run the same code with Intel Fortran and it runs perfectly fine.

Thanks

Hi bguler,

Try adding the “-Mlarge_array” flags under your project’s property page’s “Fortran->Command Line” options. This flag enables allocate able arrays larger than 2GB.

Hope this helps,
Mat

I did that, but still I get the message: Allocate: xxx bytes requested; not enough memory. But the strange thing is that if I write a separate code with only creating such a big array only, with the flag you mentioned I don’t get any problem. But within my code, I keep receiving this message although the code works fine with Intel fortran compiler???

An update on the last message…

The flag doesn’t work for creating arrays bigger than 2g at all, even if I write a very simple code with allocating an integer array of size close to 2g: I still get the same message.

program main

INTEGER, parameter :: N=471500000


INTEGER, dimension(:), allocatable :: A


allocate(A(N))


end program

Hi Bguler,

Can you check to make sure that you’re targeting 64-bit? I.e make sure your configuration manager is using x64 not win32. I just checked and your code runs fine for me in PVF 2010.

  • Mat

Yes, that was the problem. Thank you very much…

However, now I have another problem. The code was compiling fine when I target 32-bit. But with 64-bit it creates problems. I gives me the “Internal compiler error. pgnvd job exited with nonzero status code 0” message.
Any thoughts?

Hi bguler,

‘pgnvd’ is the driver that handles the CUDA portions of the code. This message means that it failed for some reason. Are there more error messages above? Can you post a reproducing example?

Thanks,
Mat

The code is long, but just in case I posted it. It works fine with 32-bit platform, but with 64-bit it gives the same error message. And on the output file, it gives this message in several places for vsolvelast subroutine: error: too few arguments in function call

I appreciate your help.

Thanks

module parameters


REAL, parameter :: delta=-0.0 !default cost to individual: utility cost
REAL, parameter :: kappa=0.22 !default cost to lender and seller: proportional to house price
REAL, parameter :: chi=0.0 !prepayment penalty, proportional to outstanding debt
REAL, parameter :: sigma=2.0 !utility parameter for risk aversion
REAL, parameter :: gamma=0.198 !relative utility benefit of homeownership
REAL, parameter :: teta1=0.1 !hh loss in selling the house
REAL, parameter :: teta2=0.0 !lender’s price loss in selling the house
REAL, parameter :: teta3=0.0 !transaction cost of purchasing a house
REAL, parameter :: rep1=0.35, rep2=0.2 !replacement ratio in retirement
REAL, parameter :: omega=0.1 !last period loss of house price
REAL, parameter :: ra=0.032 !annual risk-free interest rate
REAL, DIMENSION(2) :: beta1 !household discount factor
parameter (beta1=(/0.8,0.96/))
real, dimension(2) :: type
parameter (type=(/0.3,0.7/))
REAL, parameter :: rmmaxa=0.06 !max mortgage interest rate
REAL, parameter :: mk=0.0044 !markup for mortgage interest rate (annual)
REAL, DIMENSION(3) :: coef !coefficients for deterministic part of earning process; quadratic in age
parameter (coef=(/-0.559, 0.056, -0.001/))
REAL, parameter :: stda_eta=0.34 !std deviation for persistent component of earnings
!REAL, parameter :: stda_eps=0.15 !std deviation for transitory component of earnings
REAL, parameter :: muw=-2.794 !mean for wealth distribution (lognormal)
REAL, parameter :: sigmaw=1.784 !std deviation for wealth distribution (lognormal)
REAL, parameter :: rho=0.8 !persistency of the earning shock
REAL, parameter :: pwra=3.0, pwrr=3.0, pwrd=1.5, pwrp=1.5 !for discretization purposes
REAL, parameter :: tolbrent = 1.0e-07
REAL, parameter :: qqa=0.05
INTEGER, parameter :: age=36 !life-cycle horizon until retirement
INTEGER, parameter :: ret=18
INTEGER, parameter :: Nsim=50000
INTEGER, parameter :: Na=128,Nnu=7,Nr=16,Nd=16, Np=40, Npr=3 !number of grid points
!Na:asset, Neps:transitory shock to earning,Nnu:persistent shock to earning
!Nd:debt,Nr:mort.int.rate
REAL, DIMENSION(3) :: eps
parameter (eps=(/0.0,0.264,0.8218/))
REAL, DIMENSION(Npr) :: Pp1 !array for expense shock prob.
parameter (Pp1=(/0.9244,0.071,0.0046/))

!REAL, DIMENSION(1) :: eps
!parameter (eps=(/1.0/))
REAL, parameter :: infinity=1.0e+15, ftol=1.0e-1
REAL, parameter :: mytol=1.0e-3,epsC=1.0e-3,epsA=0.0
REAL, parameter :: prba=0.17 !annual probability of returning to the housing market for defaulter
INTEGER, parameter :: itermax=1000
REAL, parameter :: conv=0.0 !convexity parameter for updating the mortgage rate function. Close to
! 1 gives more weight on the old rate

end module

!============================================================================================================!

module global_parameters

USE parameters

implicit none

REAL, DIMENSION(Na,age+1) :: agrid !grid points for asset
REAL, DIMENSION(Nnu) :: nugrid !grid points for permanent comp
REAL, DIMENSION(Nr) :: rmgrid !grid points for mortgage interest rate
REAL, DIMENSION(Nd) :: dgrid !grid points for debt
REAL, DIMENSION(Np) :: pgrid !grid points for downpayment
REAL :: beta(2),ph,pr,r,rmmax,rmk !household discount rate,house and rental prices
REAL :: amin,amax(age+1),prb,meany,qq
real, dimension(Npr) :: Pp
REAL, DIMENSION(Nnu,Nnu) :: Pnu !transition matrix for persistent shock
REAL, DIMENSION(age) :: ydet !deterministic part of earnings
INTEGER :: cons

end module


module global

USE parameters


real, dimension(:,:,:,:,:), allocatable :: ar !policy function for renter
real, dimension(:,:,:,:,:), allocatable :: ad !policy function for defaulter
integer, dimension(:,:,:,:,:), allocatable :: str !indicator for renter, 1 for renting, 2 for purchasing


end module

!==================================================================================================!

module gpu_mod
use cudafor

real, device :: beta_dev(2),ydet_dev,rep1_dev,rep2_dev,meany_dev,ph_dev,&
pr_dev,epsC_dev,gamma_dev,r_dev,epsA_dev,tolbrent_dev,teta1_dev,kappa_dev,&
sigma_dev,eps_dev(3),qq_dev,rmk_dev,&
rmmax_dev,prb_dev
integer, device :: Na_dev,Nd_dev,Nr_dev,Nnu_dev,Npr_dev,ret_dev,rem,Np_dev

real, device, dimension(:), allocatable :: dg_dev,rg_dev,ng_dev,pg_dev,Pp_dev
real, device, dimension(:,:), allocatable :: Pnu_dev, ag_dev

REAL, device, DIMENSION(:,:,:,:), ALLOCATABLE :: vr,vd
real, device, dimension(:,:,:,:), allocatable :: vrp,vdp
real, device, dimension(:,:,:,:), allocatable :: ard !policy function for renter
real, device, dimension(:,:,:,:), allocatable :: add !policy function for defaulter
integer, device, dimension(:,:,:,:), allocatable :: strd !indicator for renter, 1 for renting, 2 for purchasing

contains

!==============================================================================================!

!kernel for the last period

attributes(global) subroutine vlast_kernel()

implicit none

real :: income,cash,vr1,ar1,vr2,ar2,alf,alfa,alfa1,alfa2,det
real :: debt,cash1,vd1,ad1,vh1,vh2,ah1,ah2
integer :: i,j,k,s,b,n

i=threadidx%x
j=blockidx%x
k=ceiling(real(j)/Nnu_dev)
s=ceiling(real(j-(k-1)*Nnu_dev)/Npr_dev)
b=j-(k-1)*Nnu_dev-(s-1)*Npr_dev

alf=1/(1+r_dev)
alfa=(beta_dev(b)(1+r_dev))(1/sigma_dev)/(1+r_dev)
alfa1=(1-alfa
(ret_dev+1))/(1-alfa)
alfa2=(1-alf**(ret_dev+1))/(1-alf)
det=(beta_dev(b)
(1+r_dev))**((1-sigma_dev)/sigma_dev)*beta_dev(b)

income=exp(ydet_dev+ng_dev(k))
cash=ag_dev(i,1)(1+r_dev)+income(1-eps_dev(s))
call vsolvelast(cash,0,vr1,ar1,income,alf,alfa1,alfa2,det,b)
call vsolvelast(cash-ph_dev+pr_dev,1,vr2,ar2,income,alf,alfa1,alfa2,det,b)

if (vr1.ge.vr2) then
vr(b,s,k,i)=vr1
ard(b,s,k,i)=ar1
strd(b,s,k,i)=1
else
vr(b,s,k,i)=vr2
ard(b,s,k,i)=ar2
strd(b,s,k,i)=2
end if

call vsolvelast(cash,0,vd1,ad1,income,alf,alfa1,alfa2,det,b)

end subroutine vlast_kernel

!==============================================================================================


!This part solves for a homeowner who decided to return to the rental market by prepaying the mortgage
attributes(device) subroutine vsolvelast(cash,x,vh1,ap,income,alf,alfa1,alfa2,det,b)

implicit none

integer, value, intent(in) :: b,x
REAL, device, INTENT(IN) :: cash,income,alf,alfa1,alfa2,det
REAL, device, INTENT(OUT) :: vh1,ap
real :: ax,bx,cx,uterm,cons,vht,vmax
INTEGER :: i,amax

!CHECK TO SEE IF THE INTERVAL IS LARGE ENOUGH FOR THE FOLLOWING MAXIMIZATION ARGUMENTS TO MAKE SENSE…
!OTHERWISE SET VALUE TO -INFINITY
if (cash.lt.epsC_dev) then
vh1=-1.0e+15
ap=epsA_dev
else
vmax=-1.0e+15
amax=1
do i=1,Na_dev
if (cash-ag_dev(i,2).lt.epsC_dev) then
vht=-1.0e+15
else
cons=(incomerep1_dev+meany_devrep2_dev-pr_dev*(1-x))alfa2/alfa1+&
(ag_dev(i,2)
(1+r_dev)+ph_devxalfret_dev)/alfa1
if (cons.lt.epsC_dev) then
uterm=-1.0e+15
else
uterm=util(cons,xgamma_dev)(1-det
(ret_dev+1))/(1-det)
end if
vht=util(cash-ag_dev(i,2),x*gamma_dev)+beta_dev(b)*uterm
end if
if (vht.gt.vmax) then
vmax=vht
amax=i
end if
end do

ap=ag_dev(amax,2)
vh1=vmax

end if
end subroutine

!============================================================================================================!

!utility function
attributes(device) function util(c,x)

implicit none
real, device :: util,c,x
!real :: sigma
!sigma=2.0
!util=((c*(1+x))**(1-sigma))/(1-sigma)
util=-1.0/(c*(1+x))
end function


end module !end of gpu module


!===========================================================================================================!

module procedures

implicit none

contains

!same as matlab, creates an array of size n, equally distributed between l and k

function linspace(l,k,n)

implicit none
REAL, INTENT(IN) :: l, k
REAL :: d
INTEGER, INTENT(IN) :: n
INTEGER :: i
REAL :: linspace(n)

d = (k-l)/(n-1)
linspace(1) = Dble(l)
linspace(n) =DBLE(k)
do i = 2,n-1
linspace(i) = linspace(i-1) + d
end do

end function

end module


!============================================================================================================!

subroutine initial
USE parameters
USE global_parameters
USE procedures
!USE bspline, ONLY: dbsnak

implicit none
REAL :: rholong(36/age),std_eta,std_eps
REAL :: ydet1(36),mtr(Nnu),inita(Na),x,y,temp1,temp2,ydet2(age)
INTEGER :: i,j,k

cons=36/age
if (MOD(36,age)/=0) then
PRINT*,‘age should be dividend of 45’
pause
end if

beta=beta1cons
r=(1+ra)cons-1
rholong=(/(rho
(i-1),i=1,cons)/)
std_eta=SQRT(SUM(rholong))*stda_eta
rmmax=(1+rmmaxa+mk)
(cons)-1
rmk=(1+ra+mk)**(cons)-1
prb=1-(1-prba)**cons
qq=1-(1-qqa)**cons

do i=2,Npr
Pp(i)=1-(1-Pp1(i))**cons
end do
Pp(1)=1.0-sum(Pp(2:Npr))

!call adda_cooper(Nnu,0.0,std_eta,rhocons,nugrid,Pnu)
!call tauchen(Nnu,0.0,std_eta,rho
cons,nugrid,Pnu)
nugrid=0.0
Pnu=1/Nnu

do i=1,36
ydet1(i)=coef(1)+coef(2)*i+coef(3)*i**2
end do

ydet=0.0

!if (cons.ne.1) then
do i=1,age
do j=cons*(i-1)+1,cons*i
ydet(i)=ydet(i)+ydet1(j)
end do
end do
!else
!ydet=ydet1
!end if

meany=sum(exp(ydet))/age


!ph=meany4.1/real(cons)(1+ra)/(1-3.1ra)
ph=meany
4.1/REAL(cons) !house price in terms of the peak of the average earning profile
!pr=0.031phREAL(cons) !annual rental price as a percentage of house price
pr=ra/(1+ra)phreal(cons)

amax(1)=EXP(ydet(1))
do i=2,age+1
amax(i)=amax(i-1)*(1+r)+EXP(ydet(i-1)+maxval(nugrid))*0.3
end do

amin=epsA

do i=1,age+1
agrid(:,i)=(linspace(0.0,(amax(i)-amin)**(1/pwra),Na))pwra+amin
end do
rmgrid=(linspace(0.0,(rmmax-rmk)
(1/pwrr),Nr))pwrr+rmk
dgrid=(linspace(0.0,(1.0-0.05)
(1/pwrd),Nd))pwrd+0.05_8
pgrid=(linspace(0.0,(1.0-0.05)
(1/pwrp),Np))**pwrp+0.05_8

end subroutine initial

!============================================================================================================!

subroutine mainsolve
USE parameters
USE global_parameters
USE global
USE procedures
use cudafor
use gpu_mod

implicit none



integer :: m,istat,q
type(dim3) :: dimGrid, dimBlock


allocate(vrp(2,Npr,Nnu,Na),vdp(2,Npr,Nnu,Na))

ALLOCATE(vr(2,Npr,Nnu,Na),vd(2,Npr,Nnu,na))

allocate(ard(2,Npr,Nnu,Na),add(2,Npr,Nnu,Na),strd(2,Npr,Nnu,Na))

allocate(dg_dev(Nd),rg_dev(Nr),pg_dev(Np),ng_dev(Nnu),Pp_dev(Npr))
allocate(Pnu_dev(Nnu,Nnu),ag_dev(Na,2))


dg_dev=dgrid
ng_dev=nugrid
pg_dev=pgrid
rg_dev=rmgrid
Pp_dev=Pp
Pnu_dev=Pnu
beta_dev=beta
eps_dev=eps
rep1_dev=rep1
rep2_dev=rep2
meany_dev=meany
ph_dev=ph
pr_dev=pr
epsC_dev=epsC
epsA_dev=epsA
gamma_dev=gamma
r_dev=r
tolbrent_dev=tolbrent
teta1_dev=teta1
kappa_dev=kappa
ret_dev=ret
sigma_dev=sigma
Na_dev=Na
Nnu_dev=Nnu
Nd_dev=Nd
Np_dev=Np
Npr_dev=Npr
Nr_dev=Nr
qq_dev=qq
prb_dev=prb
rmk_dev=rmk
rmmax_dev=rmmax


dimGrid = dim3(NnuNpr2,1,1)
dimBlock = dim3(Na,1,1)

!loop for mortgage interest rate evaluation: uses present value condition and solves household problem

!the last period mortgage interest rate is trivial: rm=r
m=age
print*, ‘age=’,m

ydet_dev=ydet(m)
ag_dev=agrid(1:Na,m:m+1)



call vlast_kernel<<<dimGrid,dimBlock>>>()
istat=cudaThreadsynchronize()
vrp=vr
vdp=vd
ar(:,:,:,:,m)=ard
ad(:,:,:,:,m)=add
str(:,:,:,:,m)=strd

!deallocate all the arrays!!!
deallocate(vr,vd)
deallocate(vrp,vdp)
deallocate(ard,add)

end subroutine mainsolve

program main

use gpu_mod
use global
!USE parameters
!USE procedures

implicit none

integer idevice, istat

call initial

idevice=0

istat=cudaSetDevice(idevice)

allocate(ar(2,Npr,Nnu,Na,age),ad(2,Npr,Nnu,Na,age),str(2,Npr,Nnu,Na,age))

call mainsolve
!call distribution

deallocate(ar,ad,str)

end program

Hi bguler,

Which version of the compiler are you using? The code compiles and runs fine for me with 10.6 and 10.8. In 10.4 and earlier we did not support device to device copies so your code would get a syntax error. In 10.5, I does see a similar ICE due to a missing identifier “__cuda_powif” but this occurs in both 32-bit and 64-bit.

If you are using 10.6 or 10.8, can you please post your command line options and the full error that you’re seeing?

Thanks,
Mat

% pgf90 test1.f90 -Mcuda -V10.8; a.out
 age=           36

a% pgf90 test1.f90 -Mcuda -V10.5
/tmp/pgcudafor_Tig6rCzBi1Q.gpu(41): error: identifier "__cuda_powif" is undefined

/tmp/pgcudafor_Tig6rCzBi1Q.gpu(90): error: identifier "__cuda_powif" is undefined

2 errors detected in the compilation of "/tmp/pgnvd5UigRyvdK_lK.nv0".
PGF90-F-0000-Internal compiler error. pgnvd job exited with nonzero status code       0 (test1.f90: 211)
PGF90/x86-64 Linux 10.5-0: compilation aborted

% pgf90 test1.f90 -Mcuda -V10.4
PGF90-S-0155-more than one device-resident object in assignment  (test1.f90: 404)
PGF90-S-0155-more than one device-resident object in assignment  (test1.f90: 405)
  0 inform,   0 warnings,   2 severes, 0 fatal for mainsolve

I have version 10.8. That shouldn’t be the problem. Here is the command line options:

-g -Bstatic -Mbackslash -Mcuda=cuda3.0 -I"c:\program files\pgi\win64\10.8\include" -I"C:\Program Files\PGI\Microsoft Open Tools 10\include" -I"C:\Program Files\PGI\Microsoft Open Tools 10\PlatformSDK\include" -tp=core2-64 -ta=nvidia,wait,cuda3.0 -Minform=warn

As for the error list, the only error is written as:
Error 2 Internal compiler error. pgnvd job exited with nonzero status code 0


And here is the output file:

------ Rebuild All started: Project: mortgage_cuda, Configuration: Debug x64 ------
Deleting intermediate and output files for project ‘mortgage_cuda’, configuration ‘Debug’
Compiling Project …
test2.f90
C:\Users\bguler\AppData\Local\Temp\pgcudafor2aK8YSMXoX6ux.gpu(41): error: too few arguments in function call

C:\Users\bguler\AppData\Local\Temp\pgcudafor2aK8YSMXoX6ux.gpu(43): error: too few arguments in function call

C:\Users\bguler\AppData\Local\Temp\pgcudafor2aK8YSMXoX6ux.gpu(54): error: too few arguments in function call

3 errors detected in the compilation of “C:\Users\bguler\AppData\Local\Temp\pgnvd2a0qZC0SmIfrT.nv0”.
C:\Users\bguler\Documents\Visual Studio 2008\Projects\mortgage_cuda\mortgage_cuda\test2.f90(213) : error F0000 : Internal compiler error. pgnvd job exited with nonzero status code 0
PGF90/x86-64 Windows 10.8-0: compilation aborted
test2.f90
C:\Users\bguler\AppData\Local\Temp\pgcudafor2aNalb1MXoQ6Kz.gpu(41): error: too few arguments in function call

C:\Users\bguler\AppData\Local\Temp\pgcudafor2aNalb1MXoQ6Kz.gpu(43): error: too few arguments in function call

C:\Users\bguler\AppData\Local\Temp\pgcudafor2aNalb1MXoQ6Kz.gpu(54): error: too few arguments in function call

3 errors detected in the compilation of “C:\Users\bguler\AppData\Local\Temp\pgnvd2a3qabL0HczLOq.nv0”.
C:\Users\bguler\Documents\Visual Studio 2008\Projects\mortgage_cuda\mortgage_cuda\test2.f90(213) : error F0000 : Internal compiler error. pgnvd job exited with nonzero status code 0
PGF90/x86-64 Windows 10.8-0: compilation aborted
mortgage_cuda build failed.
Build log was saved at “file://C:\Users\bguler\Documents\Visual Studio 2008\Projects\mortgage_cuda\mortgage_cuda\x64\Debug\BuildLog.htm”

========== Rebuild All: 0 succeeded, 1 failed, 0 skipped ==========

Hi bguler,

Thanks that helped. I was using Linux and the error does seem to only appear on Win64. It does appear to be a compiler issue so I have sent a report to our engineers (TPR#17204) for further investigation.

Hopefully, we can have it fixed soon.

Best Regards,
Mat

Hi Matt,

I have a related question. I kept modifying the code to see where the problem is, and I started to have some other weird problems. Below is a slightly modified version of the same code. I can compile it with win32 but with win64 configuration it fails to compile.
More important than that, even it is working with win32, it gives wrong results! The simple max operation fails to compute the result in the v_last_kernel subroutine: vr(b,s,k,i)=max(vr1,vr2)
It gives 0 although it shouldn’t be. I printed vr1 and vr2 separately and they look OK, but once I try to take the max of these two numbers, it fails to compute the result, and returns 0 instead. I don’t understand what this is all about? Is it some stupid bug in my code that I cannot see, or is it something related to the error message I got with win64?

Thanks





module parameters

REAL, parameter :: kappa=0.22 !default cost to lender and seller: proportional to house price
REAL, parameter :: sigma=2.0 !utility parameter for risk aversion
REAL, parameter :: gamma=0.198 !relative utility benefit of homeownership
REAL, parameter :: teta1=0.1 !hh loss in selling the house
REAL, parameter :: rep1=0.35, rep2=0.2 !replacement ratio in retirement
REAL, parameter :: omega=0.1_8 !last period loss of house price
REAL, parameter :: ra=0.032 !annual risk-free interest rate
REAL, DIMENSION(2) :: beta1 !household discount factor
parameter (beta1=(/0.8,0.96/))
REAL, DIMENSION(3) :: coef !coefficients for deterministic part of earning process; quadratic in age
parameter (coef=(/-0.558984685666261, 0.056358450306784, -0.001111973415131/)) !persistency of the earning shock
REAL, parameter :: pwra=3.0
REAL, parameter :: tolbrent = 1.0e-07
INTEGER, parameter :: age=36 !life-cycle horizon until retirement
INTEGER, parameter :: ret=18
INTEGER, parameter :: Na=256,Nnu=1, Npr=3 !number of grid points

REAL, DIMENSION(3) :: eps
parameter (eps=(/0.0,0.264,0.8218/))
REAL, DIMENSION(Npr) :: Pp1 !array for expense shock prob.
parameter (Pp1=(/0.9244,0.071,0.0046/))

REAL, parameter :: infinity=1.0e+15, ftol=1.0e-1
REAL, parameter :: mytol=1.0e-3,epsC=1.0e-3,epsA=0.0

end module

!============================================================================================================!

module global_parameters

USE parameters

implicit none

REAL, DIMENSION(Na,age+1) :: agrid !grid points for asset
REAL, DIMENSION(Nnu) :: nugrid !grid points for permanent comp
REAL :: beta(2),ph,pr,r !household discount rate,house and rental prices
REAL :: amin,amax(age+1),meany
REAL, DIMENSION(Nnu,Nnu) :: Pnu !transition matrix for persistent shock
real, dimension(Npr) :: Pp
REAL, DIMENSION(age) :: ydet !deterministic part of earnings
INTEGER :: cons

end module


module global

USE parameters

real, dimension(:,:,:,:), allocatable :: vrf
end module

!==================================================================================================!

module gpu_mod
use cudafor

real, device :: beta_dev(2),ydet_dev,rep1_dev,rep2_dev,meany_dev,ph_dev,pr_dev,&
epsC_dev,gamma_dev,r_dev,epsA_dev,tolbrent_dev,teta1_dev,kappa_dev,sigma_dev,&
eps_dev(3)
integer, device :: Na_dev,Nnu_dev,ret_dev,Npr_dev

real, device, dimension(:), allocatable :: ng_dev,Pp_dev
real, device, dimension(:,:), allocatable :: Pnu_dev, ag_dev

REAL, device, DIMENSION(:,:,:,:), ALLOCATABLE :: vr

contains

!==============================================================================================!

!kernel for the last period

attributes(global) subroutine vlast_kernel()

implicit none

real :: alf,alfa1,alfa2
real :: income,cash,vr1,vr2,alfa,det,vr3
real :: debt,cash1,vd1,vh1,vh2
integer :: ar1,ad1,ah1,ar2,ah2
integer :: i,k,b,s,j

i=threadidx%x
j=blockidx%x
k=ceiling(real(j)/(Npr_dev*2))
s=ceiling(real(j-(k-1)Npr_dev2)/2)
b=j-(k-1)Npr_dev2-(s-1)*2

alf=1.0/(1.0+r_dev)
alfa=(beta_dev(b)(1+r_dev))(1/sigma_dev)/(1+r_dev)
alfa1=(1-alfa
(ret_dev+1))/(1-alfa)
alfa2=(1-alf**(ret_dev+1))/(1-alf)
det=(beta_dev(b)
(1+r_dev))**((1-sigma_dev)/sigma_dev)beta_dev(b)
income=exp(ydet_dev+ng_dev(k))
cash=ag_dev(i,1)
(1+r_dev)+income*(1-eps_dev(s))
call vsolvelast(cash,0,income,alf,alfa1,alfa2,det,b,vr1,ar1)
cash1=cash-ph_dev+pr_dev
call vsolvelast(cash1,1,income,alf,alfa1,alfa2,det,b,vr2,ar2)

vr(b,s,k,i)=max(vr1,vr2)
!vr(b,s,k,i)=vr1
!vr(b,s,k,i)=vr2

end subroutine vlast_kernel


!==============================================================================================

!This part solves for a homeowner who decided to return to the rental market by prepaying the mortgage
attributes(device) subroutine vsolvelast(cash,x,income,alf,alfa1,alfa2,det,b,vh1,ap)

implicit none

integer, intent(in) :: b,x
REAL, INTENT(IN) :: cash,income,alf,alfa1,alfa2,det
REAL, INTENT(OUT) :: vh1
integer, intent(out) :: ap
real :: uterm,cons,vht,alf2
INTEGER :: i

!CHECK TO SEE IF THE INTERVAL IS LARGE ENOUGH FOR THE FOLLOWING MAXIMIZATION ARGUMENTS TO MAKE SENSE…
!OTHERWISE SET VALUE TO -INFINITY
vh1=-1.0e+15
ap=1

alf2=alfret_dev
if (cash.ge.epsC_dev) then
do i=1,Na_dev
if (cash-ag_dev(i,2).ge.epsC_dev) then
cons=(incomerep1_dev+meany_devrep2_dev-pr_dev*(1.0-x))alfa2/alfa1+&
(ag_dev(i,2)
(1+r_dev)+alf2xph_dev)/alfa1
if (cons.ge.epsC_dev) then
uterm=util(cons,xgamma_dev)(1-det
(ret_dev+1))/(1-det)
vht=util(cash-ag_dev(i,2),x*gamma_dev)+beta_dev(b)*uterm
if (vht.gt.vh1) then
vh1=vht
ap=i
end if
end if
end if
end do
end if
end subroutine


!utility function
attributes(device) function util(c,x)

implicit none
real, intent(in) :: c,x
real :: util
!real :: sigma
!sigma=2.0
!util=((c*(1+x))**(1-sigma_dev))/(1-sigma_dev)
util=-1.0/(c*(1.0+x))
end function

end module !end of gpu module


!===========================================================================================================!

module procedures

implicit none

contains

function linspace(l,k,n)

implicit none
REAL, INTENT(IN) :: l, k
REAL :: d
INTEGER, INTENT(IN) :: n
INTEGER :: i
REAL :: linspace(n)

d = (k-l)/(n-1)
linspace(1) = Dble(l)
linspace(n) =DBLE(k)
do i = 2,n-1
linspace(i) = linspace(i-1) + d
end do

end function

end module


!============================================================================================================!

subroutine initial
USE parameters
USE global_parameters
USE procedures

implicit none
REAL :: ydet1(36),inita(Na),x,y
INTEGER :: i,j,k

cons=36/age
if (MOD(36,age)/=0) then
PRINT*,‘age should be dividend of 45’
pause
end if

beta=beta1**cons
r=(1+ra)**cons-1

do i=2,Npr
Pp(i)=1-(1-Pp1(i))**cons
end do
Pp(1)=1.0-sum(Pp(2:Npr))

nugrid=0.0

Pnu=1.0/Nnu

do i=1,36
ydet1(i)=coef(1)+coef(2)*i+coef(3)*i**2
end do

ydet=0.0

do i=1,age
do j=cons*(i-1)+1,consi
ydet(i)=ydet(i)+ydet1(j)
end do
end do


meany=sum(exp(ydet))/age


!ph=meany
4.1/real(cons)(1+ra)/(1-3.1ra)
ph=meany4.1/REAL(cons) !house price in terms of the peak of the average earning profile
!pr=0.031
ph*REAL(cons) !annual rental price as a percentage of house price
pr=ra/(1+ra)phreal(cons)

amax(1)=EXP(ydet(1))
do i=2,age+1
amax(i)=amax(i-1)*(1+r)+EXP(ydet(i-1)+maxval(nugrid))*0.3
end do

amin=epsA

do i=1,age+1
agrid(:,i)=(linspace(0.0,(amax(i)-amin)**(1/pwra),Na))**pwra+amin
end do

end subroutine initial

!============================================================================================================!

subroutine mainsolve
USE parameters
USE global_parameters
USE global
USE procedures
use cudafor
use gpu_mod

implicit none



integer :: m,istat,q
type(dim3) :: dimGrid, dimBlock


ALLOCATE(vr(2,Npr,Nnu,Na))

allocate(ng_dev(Nnu))
allocate(Pnu_dev(Nnu,Nnu),ag_dev(Na,2))


ng_dev=nugrid
Pnu_dev=Pnu
Pp_dev=Pp
beta_dev=beta
eps_dev=eps
rep1_dev=rep1
rep2_dev=rep2
meany_dev=meany
ph_dev=ph
pr_dev=pr
epsC_dev=epsC
epsA_dev=epsA
gamma_dev=gamma
r_dev=r
tolbrent_dev=tolbrent
teta1_dev=teta1
kappa_dev=kappa
ret_dev=ret
sigma_dev=sigma
Na_dev=Na
Nnu_dev=Nnu
Npr_dev=Npr

dimGrid = dim3(Npr2Nnu,1,1)
dimBlock = dim3(Na,1,1)

!loop for mortgage interest rate evaluation: uses present value condition and solves household problem

!the last period mortgage interest rate is trivial: rm=r
m=age
print*, ‘age=’,m

ydet_dev=ydet(m)
ag_dev=agrid(1:Na,m:m+1)

call vlast_kernel<<<dimGrid,dimBlock>>>()
istat=cudaThreadsynchronize()

vrf=vr

open(2,file=‘vr.txt’,status=‘replace’)
write(2,*) vrf
close(2)


!deallocate all the arrays!!!
deallocate(vr)

end subroutine mainsolve

!============================================================================================================!
program main

use gpu_mod
use global
use cudafor

implicit none

integer :: idevice, istat

call initial

idevice=0

istat=cudaSetDevice(idevice)

allocate(vrf(2,Npr,Nnu,Na))

call mainsolve


deallocate(vrf)

end program

TPR 17204 ICE “too few arguments in function call” and CUDAFE seg fault

has been corrected in the 11.0 release, which is available now.

Thanks for your participation.

regards,
dave