[OpenMP][Offload] Add `LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS` to treat `attach(auto)` as `attach(always)`. #172382

abhinavgaba · 2025-12-15T23:45:13Z

This is needed as a way to support older code that was expecting unconditional attachment to happen for cases like:

  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).

…ach(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)`.

abhinavgaba · 2025-12-15T23:46:13Z

offload/include/OpenMP/Mapping.h

    UseEventsForAtomicTransfers = ForceAtomic;
+
+    BoolEnvar TreatAttachAutoAsAlwaysEnvar(
+        "LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS", false);


We can rename this to LIBOMPTARGET_ATTACH_FORCE_ALWAYS to align with the above, if that's preferable. That is shorter, but maybe not as clear.

adurang · 2025-12-16T13:54:13Z

offload/libomptarget/omptarget.cpp

+    // 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) "


Won't this happen at every single attach? Isn't that too verbose? Maybe it's better to do just print it once when TreatAttachAutoAsAlways is set to true?

There are no other debug prints in Mapping.h. That means unless we want to include debug headers from Mapping.h, we would move the constructor of MappingConfig out to the cpp file. Or, we can print the message once per construct, before looping over individual attach entries.

ronlieb · 2025-12-16T15:45:44Z

This workaround seems to work for the failures i saw.

ronlieb

need this to land to fix the parent PR.
request not essential but could we default to the old method ?

abhinavgaba · 2025-12-16T18:26:53Z

request not essential but could we default to the old method ?

I don't think it's a good idea to make the incorrect implementation the default. This is a workaround. For specHPC, I sent a source level patch to @mjklemm, who forwarded it to @kparzysz, and we'll try to get the benchmark fixed. And until then, this environment variable can be used as a workaround.

ronlieb · 2025-12-16T18:31:43Z

i have the workaround , will bring it up to commiitee tomorrow at our wed meeting. land you pr as is when ready

kparzysz · 2025-12-16T18:38:35Z

Ron has verified that this helps with the problem. Can we merge it?

…tach-as-attach-always

abhinavgaba · 2025-12-16T20:56:41Z

Ron has verified that this helps with the problem. Can we merge it?

I wanted to add a test first. It's now added. I'll mark it as ready and merge if there are no objections in the next couple of hours.

llvmbot · 2025-12-16T20:57:24Z

@llvm/pr-subscribers-offload

Author: Abhinav Gaba (abhinavgaba)

Changes

This is needed as a way to support older code that was expecting unconditional attachment to happen for cases like:

  int *p;
  int x;

  #pragma omp targret enter data map(p) // (A)
  #pragma omp target enter data map(x)  // (B)
  p = &amp;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).

Full diff: https://github.com/llvm/llvm-project/pull/172382.diff

5 Files Affected:

(modified) offload/include/OpenMP/Mapping.h (+11)
(modified) offload/libomptarget/omptarget.cpp (+7-1)
(modified) offload/test/lit.cfg (+4)
(added) offload/test/mapping/map_ptr_then_ptee_then_attach.c (+54)
(modified) openmp/docs/design/Runtimes.rst (+18)

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 973ab39520ff1..35c2c662a3884 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -779,6 +779,11 @@ int processAttachEntries(DeviceTy &Device, AttachInfoTy &AttachInfo,
   ODBG(ODT_Mapping) << "Processing " << AttachInfo.AttachEntries.size()
                     << " deferred ATTACH map entries";
 
+  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();
@@ -799,7 +804,8 @@ 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) || TreatAttachAutoAsAlways;
 
     // 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/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(); }
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

abhinavgaba commented Dec 15, 2025

View reviewed changes

abhinavgaba mentioned this pull request Dec 15, 2025

[OpenMP][Clang] Use ATTACH map-type for list-items with base-pointers. #153683

Merged

adurang reviewed Dec 16, 2025

View reviewed changes

ronlieb self-requested a review December 16, 2025 17:44

ronlieb approved these changes Dec 16, 2025

View reviewed changes

abhinavgaba added 4 commits December 16, 2025 11:55

Do the debug print before starting individual attach entries.

7e7c2dd

Merge remote-tracking branch 'upstream/main' into env-var-to-treat-at…

f05372e

…tach-as-attach-always

Merge remote-tracking branch 'upstream/main' into env-var-to-treat-at…

cafe050

…tach-as-attach-always

Add a compile+run test.

8ec55c4

abhinavgaba marked this pull request as ready for review December 16, 2025 20:56

llvmbot added openmp:libomp OpenMP host runtime offload labels Dec 16, 2025

abhinavgaba merged commit c62cd28 into llvm:main Dec 16, 2025
15 checks passed

abhinavgaba deleted the env-var-to-treat-attach-as-attach-always branch December 17, 2025 19:58

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

[OpenMP][Offload] Add `LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS` to treat `attach(auto)` as `attach(always)`. #172382

[OpenMP][Offload] Add `LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS` to treat `attach(auto)` as `attach(always)`. #172382

Uh oh!

abhinavgaba commented Dec 15, 2025

Uh oh!

abhinavgaba Dec 15, 2025

Uh oh!

adurang Dec 16, 2025

Uh oh!

abhinavgaba Dec 16, 2025

Uh oh!

ronlieb commented Dec 16, 2025

Uh oh!

ronlieb left a comment

Uh oh!

abhinavgaba commented Dec 16, 2025

Uh oh!

ronlieb commented Dec 16, 2025

Uh oh!

kparzysz commented Dec 16, 2025

Uh oh!

abhinavgaba commented Dec 16, 2025

Uh oh!

llvmbot commented Dec 16, 2025

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

[OpenMP][Offload] Add LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS to treat attach(auto) as attach(always). #172382

[OpenMP][Offload] Add LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS to treat attach(auto) as attach(always). #172382

Uh oh!

Conversation

abhinavgaba commented Dec 15, 2025

Uh oh!

abhinavgaba Dec 15, 2025

Choose a reason for hiding this comment

Uh oh!

adurang Dec 16, 2025

Choose a reason for hiding this comment

Uh oh!

abhinavgaba Dec 16, 2025

Choose a reason for hiding this comment

Uh oh!

ronlieb commented Dec 16, 2025

Uh oh!

ronlieb left a comment

Choose a reason for hiding this comment

Uh oh!

abhinavgaba commented Dec 16, 2025

Uh oh!

ronlieb commented Dec 16, 2025

Uh oh!

kparzysz commented Dec 16, 2025

Uh oh!

abhinavgaba commented Dec 16, 2025

Uh oh!

llvmbot commented Dec 16, 2025

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

[OpenMP][Offload] Add `LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS` to treat `attach(auto)` as `attach(always)`. #172382

[OpenMP][Offload] Add `LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS` to treat `attach(auto)` as `attach(always)`. #172382