OpenACC+CUDA in C code called from Fortran issue

[caplanr]

Hi,

I am having an issue with OpenACC and CUDA interacting.

I have a Fortran code where an array x is put onto the GPU with
!$acc enter data create(x).

Later in the code, this array is sent into a C code that calls CUDA libraries
using an !$acc host_data use_device(x) region and C_LOC as:

!$acc host_data use_device(x)
call routine_in_c_code(C_LOC(x(1)))
!$acc end host_data

Now, in the C code, the routine has the signature:

void routine_in_c_code(double* x){

If I use the x array in CUDA library calls, it works fine:

cusparseDnVecSetValues(DenseVecX,x);

However now I need to modify the array in the C code itself.

I wanted to avoid writing a CUDA kernel, so instead I make an OpenACC loop in the C code:

#pragma acc parallel loop default(present)
for (int i=0;i<N;i++){
<CODE THAT USES x[i]>;
}

The code compiles fine (using nvc and -acc=gpu) but the run crashes with:

FATAL ERROR: data in PRESENT clause was not found on device 1: name=x host:0x7f5255400000

Any ideas?

My initial thought is that even though x is a device pointer,
the OpenACC runtime in the C code does not see it in the present table?

– Ron

[caplanr]

Hi again,

Problem solved!

I just didn’t know my OpenACC well enough…

I needed:
#pragma acc parallel loop deviceptr(x)

I suppose this is why deviceptr is a thing.

I still am wondering though - since I did in fact send the x array to the GPU with !$acc enter data create, shouldn’t it have been in the present table and not needed deviceptr?

Or are the present tables for the Fortran and C code separate?

– Ron

[MatColgrove]

@caplanr

I still am wondering though - since I did in fact send the x array to the GPU with !$acc enter data create , shouldn’t it have been in the present table and not needed deviceptr ?

The present table associates the host pointer to the device pointer but here you’re passing the raw device pointer in directly and why “deviceptr” is needed. Otherwise it’s assumed to be a host pointer, hence the look-up fails.

Though, you probably don’t need to use “host_data” any longer since you’re no longer calling CUDA. In which case you can then remove ‘deviceptr’ as well.

Or are the present tables for the Fortran and C code separate?

No, it’s the same OpenACC runtime library for both.

-Mat