Hi all,
when troubleshooting building/running SPEC HPC 2021 with GCC with OpenMP
offloading,
specifically 534.hpgmgfv_t, an issue encountered in the benchmark was:
when the benchmark was initializing and creating its data environment on the
GPU,
it was trying to map array sections where the base-pointer is actually NULL:
...
for (block=0;block<3;++block) {
#pragma omp target enter data
map(to:level->restriction[shape].blocks[block][:length])
// level->restriction[shape].blocks[block] == NULL for some values of index
'block'
...
The benchmark appears to be assuming that such NULL base-pointers would simply
be
silently ignored, and the program would just keep running.
(BTW, the above case needs this patch to compile:
https://gcc.gnu.org/pipermail/gcc-patches/2022-February/590658.html
which is still awaiting review :) )
What we currently do in libgomp, however, is that we issue an error and call
gomp_fatal():
libgomp/target.c:gomp_attach_pointer():
...
if ((void *) target == NULL)
{
- gomp_mutex_unlock (&devicep->lock);
- gomp_fatal ("attempt to attach null pointer");
+ n->aux->attach_count[idx] = 0; // proposed change
attached in patch
+ return;
...
Some quick testing shows that clang/LLVM behaves mostly the same as GCC.
OTOH, nVidia HPC SDK (PGI) does appear to silently go on without bailing out.
(I have not verified if 534.hpgmgfv_t fully works with PGI, just observed how
their
runtime handles NULL base-pointers)
I don't see any explicit description of this case in the OpenMP specifications,
just simply
"The corresponding pointer variable becomes an attached pointer", lack of
description on how
this is to be handled.
So WDYGT? Should libgomp behavior be adjusted here, or should SPEC benchmark
source be adjusted?
(The attached patch to adjust libgomp attach behavior has been regtested
without regressions, FWIW)
Thanks,
Chung-Lin
2022-03-09 Chung-Lin Tang <clt...@codesourcery.com>
libgomp/ChangeLog:
* target.c (gomp_attach_pointer): When pointer is NULL,
return instead of calling gomp_fatal.
diff --git a/libgomp/target.c b/libgomp/target.c
index 9017458885e..0e8bbd83c20 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -796,8 +796,8 @@ gomp_attach_pointer (struct gomp_device_descr *devicep,
if ((void *) target == NULL)
{
- gomp_mutex_unlock (&devicep->lock);
- gomp_fatal ("attempt to attach null pointer");
+ n->aux->attach_count[idx] = 0;
+ return;
}
s.host_start = target + bias;
