Acc data delete() pragma not working, explicit acc_delete() required

Accelerated Computing HPC Compilers nvc, nvc++ and nvfortran
1

The initial pragma was compiling cleanly, but was triggering a runtime error:

Present table errors:

descP_[:descN_] lives at 0x7ffa9eb66010 size 61447872 partially present in

host:0x7ffa9eb2c010 device:0x7ff858000000 size:61682352 presentcount:0+1 line:469 name:descP_[:descN_] file:/…/tree_build_gpu.hpp

FATAL ERROR: variable in data clause is partially present on the device: name=descP_[:descN_]

Replacing the pragma with three explicit calls to acc_delete() as shown below fixed the problem. Why wouldn’t the pragma work?

@@ -47,2 +47,5 @@
 #include <vector>
+#if defined(_OPENACC)
+#include <openacc.h>
+#endif

@@ -473,3 +476,5 @@ class GpuResidentData {
 #if defined(_OPENACC)
-               #pragma acc exit data delete(descP_[0:descN_], probP_[0:descN_], bdP_[0:bdN_])
+               acc_delete(descP_, descN_ * sizeof(*descP_));
+               acc_delete(probP_, descN_ * sizeof(*probP_));
+               acc_delete(bdP_,   bdN_   * sizeof(*bdP_));
 #endif

Using nvc++ 25.11 on Linux.

2

Hi UnitedMarupials,

A partially present error means that there is a size mismatch between the what’s listed in the internal “present” table (which keeps track of the host to device address mapping) and what’s included in the directive.

The present table lists the size as “61447872” but is using the size “61682352” in the delete. If the value of “descN_” changed, that could trigger this.

Exactly why it works with “acc_delete”, I don’t know. They are functionally equivalent, so I would expect the same error.

If you can provide a reproducing example, I can investigate.

Otherwise, you can try removing the size of the array from the delete so the compiler uses the size listed in the present table. i.e.

#pragma acc exit data delete(descP_, probP_, bdP_)

-Mat