Hi,
I am testing pgi13.1 which seems to support the host_data directive (there are no more error at compile time).
I tried it with a small test code which uses the host_data directive (see below), but it seems to only work when the -g flag is used:
> nvcc -c simple.cu
> pgf90 -acc -ta=nvidia,cuda5.0 -L$CUDALIB -lcudart -lcuda simple.o -o test.exe test.f90
> ./test.exe
call to cuStreamSynchronize returned error 700: Launch failed array a=0x69d860 b=0x69db80 n=100
srun: error: castor26: task 0: Exited with exit code 1
with -g
> nvcc -c simple.cu
> pgf90 -g -acc -ta=nvidia,cuda5.0 -L$CUDALIB -lcudart -lcuda simple.o -o
> ./test.exe
array a=0x900100200 b=0x900100000 n=100
3.000000 12.00000
3.000000 12.00000
Here is the code:
> cat simple.cu
extern "C" {
__global__ void simple_add(float* a, float* b, int n) {
int i = threadIdx.x+ blockIdx.x * blockDim.x;
if ( i < n ) {
a[i] = a[i] + b[i];
}
}
void my_cuda_(float* a, float* b, int* p_n) {
int n = *p_n;
printf( "array a=%p b=%p n=%d\n", a, b, n );
cudaThreadSynchronize( );
simple_add<<<(n+127)/128, 128>>>(a, b, n);
cudaThreadSynchronize( );
}
};
----------
> cat test.f90
subroutine test(a,b,n)
integer n
real a(n),b(n)
!dir$ inlinenever test
!$acc data present(a,b)
!$acc host_data use_device(a,b)
call my_cuda(a,b,n)
!$acc end host_data
!$acc end data
end subroutine test
program driver
integer n
real a(100),b(100),act_a(100)
n = 100
do i=1,n
b(i) = real(i,4)
end do
a = 1.0
act_a = a + b + 1.0
!$acc data copy(a),copyin(b)
!$acc parallel
do i=1,n
a(i) = a(i) + 1.0
end do
!$acc end parallel
call test(a,b,n)
!$acc end data
print *, act_a(1), act_a(10)
print *, a(1),a(10)
end program driver
Am I doing something wrong ?
Thanks,
Xavier