Hello,
I’m using PGF90 11.7 with CUDA FORTRAN on Linux (OpenSuse 11.3, x86_64, 2.6.34.8, NVIDIA Quadro 4000, CUDA 4.0, drivers 275.09.07).
I have worked on small programs for a few weeks now, and I haven’t run into much trouble (except for some already raised on this forum). But now that I’m trying to convert an existing program to CUDA, I’m having strange errors: the generated C code contains multiple declarations, and thus cannot be compiled by pgnvd.
SUBROUTINE prodscal( rnorm, a, b )
USE numerics
IMPLICIT NONE
REAL(rp), DEVICE, INTENT(OUT) :: rnorm
REAL(rp), DEVICE, DIMENSION( sx-1:ex+1, sy-1:ey+1, sz-1:ez+1), INTENT(IN) :: a
REAL(rp), DEVICE, DIMENSION( sx-1:ex+1, sy-1:ey+1, sz-1:ez+1), INTENT(IN) :: b
INTEGER :: i, j, k
!$cuf kernel do <<<*,*>>>
do k = 1, (sx+ex+3)*(sy+ey+3)*(sz+ez+3)
rnorm = rnorm + a(k) * b(k)
end do
RETURN
END SUBROUTINE prodscal
If I do
pgf90 -Mcuda -ta=nvidia:4.0 -Minfo -c prodscal.f90
I get:
<errors about multiple declarations in a /tmp/<random>.gpu file>
PGF90-F-0000-Internal compiler error. pgnvd job exited with nonzero status code 0 (prodscal.f90: 32)
PGF90/x86-64 Linux 11.7-0: compilation aborted
The generated .gpu file contains (excerpt)
prodscal_35_gpu(
int tc1,
signed char* p1,
signed char* p7,
int u1,
int u2,
int u3,
int x2,
int x3,
int x4,
int x5,
int x6,
signed char* p7)
{
In another version, I even get this (it’s a generated header):
struct DT1_50{int m0;int m4;int m8;};
struct DT1_56{long long m0;};
struct DT1_62{long long m0;};
struct DT1_68{long long m0;};
struct DT1_74{struct DT1_68 m0;};
struct DT1_50{int m0;int m4;int m8;};
struct DT1_56{long long m0;};
struct DT1_62{long long m0;};
struct DT1_68{long long m0;};
struct DT1_74{struct DT1_68 m0;};
extern "C" __global__ void prodscal(signed char* _prnorm,signed char* _pa,signed char* _pb)
which obviously results in
/tmp/pgcudaforA2feohoyjhk8.h(6): error: invalid redeclaration of type name "DT1_50"
(1): here
My code may not be straightforward or even correct, but even if it’s the case, such problems should be detected before code generation, that’s why I’m reporting this.
Antoine
PS: The original code comes from CRIHAN [fr].