[llvm] [openmp] [OpenMP][Offload] Add `LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS` to treat `attach(auto)` as `attach(always)`. (PR #172382)
Abhinav Gaba via llvm-commits
llvm-commits at lists.llvm.org
Tue Dec 16 12:53:40 PST 2025
https://github.com/abhinavgaba updated https://github.com/llvm/llvm-project/pull/172382
>From ea98e0369dbdc8af16f90d1928fde5ed33e7639f Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Mon, 15 Dec 2025 15:35:36 -0800
Subject: [PATCH 1/3] [Offload] Add `LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS`
to treat attach(auto) as attach(always).
This is needed as a way to support older code that was expecting
unconditional attachment to happen for cases like:
```c
int *p;
int x;
#pragma omp targret enter data map(p) // (A)
#pragma omp target enter data map(x) // (B)
p = &x;
// By default, this does NOT attach p and x
#pragma omp target enter data map(p[0:0]) // (C)
```
When the environment variable is set, such maps, where both the pointer
and the pointee already have corresponding copies on the device, but are
not attached to one another, will be attached as-if OpenMP 6.1 TR14's
`attach(always)` map-type-modifier was specified on `(C)`.
---
offload/include/OpenMP/Mapping.h | 11 +++++++++++
offload/libomptarget/omptarget.cpp | 10 +++++++++-
offload/test/lit.cfg | 4 ++++
openmp/docs/design/Runtimes.rst | 18 ++++++++++++++++++
4 files changed, 42 insertions(+), 1 deletion(-)
diff --git a/offload/include/OpenMP/Mapping.h b/offload/include/OpenMP/Mapping.h
index 45bd9c6e7da8b..d40e2b188076a 100644
--- a/offload/include/OpenMP/Mapping.h
+++ b/offload/include/OpenMP/Mapping.h
@@ -33,6 +33,10 @@ class MappingConfig {
MappingConfig() {
BoolEnvar ForceAtomic = BoolEnvar("LIBOMPTARGET_MAP_FORCE_ATOMIC", true);
UseEventsForAtomicTransfers = ForceAtomic;
+
+ BoolEnvar TreatAttachAutoAsAlwaysEnvar(
+ "LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS", false);
+ TreatAttachAutoAsAlways = TreatAttachAutoAsAlwaysEnvar;
}
public:
@@ -44,6 +48,13 @@ class MappingConfig {
/// Flag to indicate if we use events to ensure the atomicity of
/// map clauses or not. Can be modified with an environment variable.
bool UseEventsForAtomicTransfers = true;
+
+ /// Flag to indicate if attach(auto) should be treated as attach(always).
+ /// This forces pointer attachments to occur between a pointer an a pointee,
+ /// for something like `map(p[:])` even when both were already present on the
+ /// device before encountering the construct. Can be modified with
+ /// an environment variable.
+ bool TreatAttachAutoAsAlways = false;
};
/// Information about shadow pointers.
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index 021caff159919..87a2a41898c46 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -798,7 +798,15 @@ int processAttachEntries(DeviceTy &Device, AttachInfoTy &AttachInfo,
<< ", PtrSize=" << PtrSize << ", MapType=0x"
<< llvm::utohexstr(MapType);
- const bool IsAttachAlways = MapType & OMP_TGT_MAPTYPE_ALWAYS;
+ bool IsAttachAlways = MapType & OMP_TGT_MAPTYPE_ALWAYS;
+
+ // Treat ATTACH(auto) as ATTACH(always) if environment variable is set
+ if (!IsAttachAlways && MappingConfig::get().TreatAttachAutoAsAlways) {
+ IsAttachAlways = true;
+ ODBG(ODT_Mapping) << "ATTACH(auto) will be treated as ATTACH(always) "
+ << " because LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS "
+ << " is true";
+ }
// Lambda to check if a pointer was newly allocated
auto WasNewlyAllocated = [&](void *Ptr, const char *PtrName) {
diff --git a/offload/test/lit.cfg b/offload/test/lit.cfg
index c0290bfdab3ff..6f54cbe8064d2 100644
--- a/offload/test/lit.cfg
+++ b/offload/test/lit.cfg
@@ -21,6 +21,10 @@ if 'ROCR_VISIBLE_DEVICES' in os.environ:
if 'LIBOMPTARGET_DEBUG' in os.environ:
config.environment['LIBOMPTARGET_DEBUG'] = os.environ['LIBOMPTARGET_DEBUG']
+# Allow running tests with attach auto treated as always
+if 'LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS' in os.environ:
+ config.environment['LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS'] = os.environ['LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS']
+
# Allow running the tests with nextgen plugins when available
if 'LIBOMPTARGET_NEXTGEN_PLUGINS' in os.environ:
config.environment['LIBOMPTARGET_NEXTGEN_PLUGINS'] = os.environ['LIBOMPTARGET_NEXTGEN_PLUGINS']
diff --git a/openmp/docs/design/Runtimes.rst b/openmp/docs/design/Runtimes.rst
index 1b6f30ae73a33..d46ec5ba5293c 100644
--- a/openmp/docs/design/Runtimes.rst
+++ b/openmp/docs/design/Runtimes.rst
@@ -735,6 +735,7 @@ variables is defined below.
* ``LIBOMPTARGET_STACK_SIZE=<Num>``
* ``LIBOMPTARGET_SHARED_MEMORY_SIZE=<Num>``
* ``LIBOMPTARGET_MAP_FORCE_ATOMIC=[TRUE/FALSE] (default TRUE)``
+ * ``LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS=[TRUE/FALSE] (default FALSE)``
* ``LIBOMPTARGET_JIT_OPT_LEVEL={0,1,2,3} (default 3)``
* ``LIBOMPTARGET_JIT_SKIP_OPT=[TRUE/FALSE] (default FALSE)``
* ``LIBOMPTARGET_JIT_REPLACEMENT_OBJECT=<in:Filename> (object file)``
@@ -1088,6 +1089,23 @@ value of the ``LIBOMPTARGET_MAP_FORCE_ATOMIC`` environment variable.
The default behavior of LLVM 14 is to force atomic maps clauses, prior versions
of LLVM did not.
+LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS
+"""""""""""""""""""""""""""""""""""""""""
+
+By default, OpenMP attach operations only perform pointer attachment
+when mapping an expression with a base-pointer/base-referring-pointer,
+when either the pointer or the pointee was newly allocated on a
+map-entering directive (aka ``attach(auto)`` as per OpenMP 6.1 TR14).
+
+When ``LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS`` is set to ``true``,
+ATTACH map entries without the ALWAYS flag are implicitly treated as if
+the ALWAYS flag was set. This forces pointer attachments to occur even when
+the pointee/pointer were not newly allocated (similar to OpenMP 6.1
+TR14's ``attach(always)`` map-type-modifier), thereby treating
+``attach(auto))`` as ``attach(always)``. This can be used for
+experimentation, or as a workaround for programs compiled without
+``-fopenmp-version=61``.
+
.. _libomptarget_jit_opt_level:
LIBOMPTARGET_JIT_OPT_LEVEL
>From 7e7c2dd38923fb62ca13c88f2f353b426ff96cec Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Tue, 16 Dec 2025 11:55:29 -0800
Subject: [PATCH 2/3] Do the debug print before starting individual attach
entries.
---
offload/libomptarget/omptarget.cpp | 16 +++++++---------
1 file changed, 7 insertions(+), 9 deletions(-)
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index 87a2a41898c46..509b41f5d4695 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -778,6 +778,11 @@ int processAttachEntries(DeviceTy &Device, AttachInfoTy &AttachInfo,
ODBG(ODT_Mapping) << "Processing " << AttachInfo.AttachEntries.size();
+ bool TreatAttachAutoAsAlways = MappingConfig::get().TreatAttachAutoAsAlways;
+ if (TreatAttachAutoAsAlways)
+ ODBG(ODT_Mapping) << "Treating ATTACH(auto) as ATTACH(always) because "
+ << "LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS is true";
+
int Ret = OFFLOAD_SUCCESS;
bool IsFirstPointerAttachment = true;
for (size_t EntryIdx = 0; EntryIdx < AttachInfo.AttachEntries.size();
@@ -798,15 +803,8 @@ int processAttachEntries(DeviceTy &Device, AttachInfoTy &AttachInfo,
<< ", PtrSize=" << PtrSize << ", MapType=0x"
<< llvm::utohexstr(MapType);
- bool IsAttachAlways = MapType & OMP_TGT_MAPTYPE_ALWAYS;
-
- // Treat ATTACH(auto) as ATTACH(always) if environment variable is set
- if (!IsAttachAlways && MappingConfig::get().TreatAttachAutoAsAlways) {
- IsAttachAlways = true;
- ODBG(ODT_Mapping) << "ATTACH(auto) will be treated as ATTACH(always) "
- << " because LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS "
- << " is true";
- }
+ bool IsAttachAlways =
+ (MapType & OMP_TGT_MAPTYPE_ALWAYS) || TreatAttachAutoAsAlways;
// Lambda to check if a pointer was newly allocated
auto WasNewlyAllocated = [&](void *Ptr, const char *PtrName) {
>From 8ec55c4bd1ba3a51366eb0d703a4a5cc09f313bc Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Tue, 16 Dec 2025 12:53:04 -0800
Subject: [PATCH 3/3] Add a compile+run test.
---
.../mapping/map_ptr_then_ptee_then_attach.c | 54 +++++++++++++++++++
1 file changed, 54 insertions(+)
create mode 100644 offload/test/mapping/map_ptr_then_ptee_then_attach.c
diff --git a/offload/test/mapping/map_ptr_then_ptee_then_attach.c b/offload/test/mapping/map_ptr_then_ptee_then_attach.c
new file mode 100644
index 0000000000000..512edded6daef
--- /dev/null
+++ b/offload/test/mapping/map_ptr_then_ptee_then_attach.c
@@ -0,0 +1,54 @@
+// RUN: %libomptarget-compile-generic
+//
+// RUN: env LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS=1 \
+// RUN: env LIBOMPTARGET_DEBUG=1 \
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefix=DEBUG
+//
+// RUN: env LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS=1 \
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefix=CHECK
+
+// Ensure that under LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS, the pointer
+// attachment for map(p[0:0]) happens as if the user had specified the
+// attach(always) map-type-modifier.
+
+#include <omp.h>
+#include <stdio.h>
+
+int x[10];
+int *p;
+
+void f1() {
+#pragma omp target enter data map(to : p)
+#pragma omp target enter data map(to : x)
+
+ p = &x[0];
+ int **p_mappedptr = (int **)omp_get_mapped_ptr(&p, omp_get_default_device());
+ int *x0_mappedptr =
+ (int *)omp_get_mapped_ptr(&x[0], omp_get_default_device());
+ int *p0_deviceaddr = NULL;
+
+ printf("p_mappedptr %s null\n", p_mappedptr == (int **)NULL ? "==" : "!=");
+ printf("x0_mappedptr %s null\n", x0_mappedptr == (int *)NULL ? "==" : "!=");
+ // CHECK: p_mappedptr != null
+ // CHECK: x0_mappedptr != null
+
+#pragma omp target enter data map(to : p[0 : 0]) // Implies: attach(auto)
+ // clang-format off
+ // DEBUG: omptarget --> Treating ATTACH(auto) as ATTACH(always) because LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS is true
+ // clang-format on
+
+#pragma omp target map(present, alloc : p) map(from : p0_deviceaddr)
+ {
+ p0_deviceaddr = &p[0];
+ }
+
+ printf("p0_deviceaddr %s x0_mappedptr\n",
+ p0_deviceaddr == x0_mappedptr ? "==" : "!=");
+ // CHECK: p0_deviceaddr == x0_mappedptr
+
+#pragma omp target exit data map(delete : x, p)
+}
+
+int main() { f1(); }
More information about the llvm-commits
mailing list