Shared variables and OpenACC

0
votes

In OpenMP, one can use shared variables in a loop by 

#pragma omp parallel for shared(foo) private(bar)

In OpenACC we have a private clause, but no shared clause. There are data clauses such as copy, copyin, copyout on the other hand.

Sometimes, we use accelerators, which have their own private memory but can as well access the common memory.

In such a case we may want the accelerator to avoid copying the data to its own private memory and operate on the instance in the common memory.

How can we tell OpenACC not to copy the data?

1
openacc

1 Answers

1
votes

Note that arrays are shared by default.

The "create" data clause will create the data on the device but not perform a copy.

If you want to use data already created on the device, such as via a call to cudaMalloc or acc_malloc, you can use the "deviceptr" data clause to tell the compiler to use the pointer's address in device code rather than look-up the device pointer in the present table using the host address.

If you want to associate a device variable with a host variable, you can use the API call "acc_map_data".

It sounds like you have larger pool of memory that you then want to re-use. In this case, you can create the memory pool using an OpenACC data clause, cudaMalloc, or acc_malloc. If using a data clause, then call "acc_deviceptr" to get the device pointer address. Next you can use "acc_map_data" to associate a host pointer with the device pointer. Note that the mapped data can be a subset of the larger device pool and you can map to an offset, i.e. "devptr+offset".

For an example of using "acc_map_data" see: https://github.com/rmfarber/ParallelProgrammingWithOpenACC/blob/master/Chapter05/acc_map.c