Skip to content

Conversation

@Saieiei
Copy link
Contributor

@Saieiei Saieiei commented Nov 26, 2025

This fixes a bug where pointers from defaultmap(firstprivate:pointer) were incorrectly treated as firstprivate literals, causing OMP_MAP_LITERAL to be set. This prevented the runtime from performing device address lookup.
Realted PR #167879

@llvmbot llvmbot added clang Clang issues not falling into any other category clang:codegen IR generation bugs: mangling, exceptions, etc. clang:openmp OpenMP related changes to Clang labels Nov 26, 2025
@llvmbot
Copy link
Member

llvmbot commented Nov 26, 2025

@llvm/pr-subscribers-clang-codegen

@llvm/pr-subscribers-clang

Author: Sairudra More (Saieiei)

Changes

…aptures

This fixes a bug where pointers from defaultmap(firstprivate:pointer) were incorrectly treated as firstprivate literals, causing OMP_MAP_LITERAL to be set. This prevented the runtime from performing device address lookup.


Patch is 28.14 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/169622.diff

18 Files Affected:

  • (modified) clang/lib/CodeGen/CGOpenMPRuntime.cpp (+66-12)
  • (modified) clang/test/OpenMP/target_codegen.cpp (+1-1)
  • (modified) clang/test/OpenMP/target_defaultmap_codegen_01.cpp (+2-2)
  • (modified) clang/test/OpenMP/target_depend_codegen.cpp (+1-1)
  • (added) clang/test/OpenMP/target_firstprivate_pointer_codegen.cpp (+169)
  • (modified) clang/test/OpenMP/target_map_codegen_01.cpp (+3-3)
  • (modified) clang/test/OpenMP/target_map_codegen_09.cpp (+2-2)
  • (modified) clang/test/OpenMP/target_map_codegen_10.cpp (+2-1)
  • (modified) clang/test/OpenMP/target_map_codegen_26.cpp (+2-2)
  • (modified) clang/test/OpenMP/target_parallel_depend_codegen.cpp (+1-1)
  • (modified) clang/test/OpenMP/target_parallel_for_depend_codegen.cpp (+1-1)
  • (modified) clang/test/OpenMP/target_parallel_for_simd_depend_codegen.cpp (+1-1)
  • (modified) clang/test/OpenMP/target_simd_depend_codegen.cpp (+1-1)
  • (modified) clang/test/OpenMP/target_teams_depend_codegen.cpp (+1-1)
  • (modified) clang/test/OpenMP/target_teams_distribute_depend_codegen.cpp (+1-1)
  • (modified) clang/test/OpenMP/target_teams_distribute_parallel_for_depend_codegen.cpp (+1-1)
  • (modified) clang/test/OpenMP/target_teams_distribute_parallel_for_simd_depend_codegen.cpp (+1-1)
  • (modified) clang/test/OpenMP/target_teams_distribute_simd_depend_codegen.cpp (+2-2)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index a8255ac74cfcf..be86d65a74897 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -28,6 +28,7 @@
 #include "clang/Basic/SourceManager.h"
 #include "clang/CodeGen/ConstantInitBuilder.h"
 #include "llvm/ADT/ArrayRef.h"
+#include "llvm/ADT/SmallSet.h"
 #include "llvm/ADT/SmallVector.h"
 #include "llvm/ADT/StringExtras.h"
 #include "llvm/Bitcode/BitcodeReader.h"
@@ -7211,6 +7212,9 @@ class MappableExprsHandler {
   /// firstprivate, false otherwise.
   llvm::DenseMap<CanonicalDeclPtr<const VarDecl>, bool> FirstPrivateDecls;
 
+  /// Set of defaultmap clause kinds that use firstprivate behavior.
+  llvm::SmallSet<OpenMPDefaultmapClauseKind, 4> DefaultmapFirstprivateKinds;
+
   /// Map between device pointer declarations and their expression components.
   /// The key value for declarations in 'this' is null.
   llvm::DenseMap<
@@ -8989,6 +8993,10 @@ class MappableExprsHandler {
           FirstPrivateDecls.try_emplace(VD, /*Implicit=*/true);
       }
     }
+    // Extract defaultmap clause information.
+    for (const auto *C : Dir.getClausesOfKind<OMPDefaultmapClause>())
+      if (C->getDefaultmapModifier() == OMPC_DEFAULTMAP_MODIFIER_firstprivate)
+        DefaultmapFirstprivateKinds.insert(C->getDefaultmapKind());
     // Extract device pointer clause information.
     for (const auto *C : Dir.getClausesOfKind<OMPIsDevicePtrClause>())
       for (auto L : C->component_lists())
@@ -9566,6 +9574,37 @@ class MappableExprsHandler {
     }
   }
 
+  /// Check if a variable should be treated as firstprivate literal.
+  /// Returns true ONLY for explicit firstprivate clauses, not for implicit
+  /// captures via defaultmap(firstprivate:pointer). Implicitly captured
+  /// pointers need runtime lookup to get their device addresses.
+  bool isEffectivelyFirstprivate(const VarDecl *VD, QualType Type) const {
+    // Check explicit firstprivate clauses (not implicit from defaultmap)
+    auto I = FirstPrivateDecls.find(VD);
+    if (I != FirstPrivateDecls.end() && !I->getSecond())
+      return true;  // Explicit firstprivate only
+
+    // For non-pointer types, defaultmap(firstprivate:...) should also
+    // be treated as firstprivate literals since they're passed by value
+    if (Type->isAnyPointerType())
+      return false;  // Pointers from defaultmap need runtime lookup
+
+    // Check defaultmap(firstprivate:scalar) for scalar types
+    if (DefaultmapFirstprivateKinds.count(OMPC_DEFAULTMAP_scalar)) {
+      if (Type->isScalarType())
+        return true;
+    }
+
+    // Check defaultmap(firstprivate:aggregate) for aggregate types
+    if (DefaultmapFirstprivateKinds.count(OMPC_DEFAULTMAP_aggregate)) {
+      if (Type->isAggregateType())
+        return true;
+    }
+
+    // Check defaultmap(firstprivate:all) for all types
+    return DefaultmapFirstprivateKinds.count(OMPC_DEFAULTMAP_all);
+  }
+
   /// Generate the default map information for a given capture \a CI,
   /// record field declaration \a RI and captured value \a CV.
   void generateDefaultMapInfo(const CapturedStmt::Capture &CI,
@@ -9593,6 +9632,9 @@ class MappableExprsHandler {
       CombinedInfo.DevicePtrDecls.push_back(nullptr);
       CombinedInfo.DevicePointers.push_back(DeviceInfoTy::None);
       CombinedInfo.Pointers.push_back(CV);
+      bool IsFirstprivate =
+          isEffectivelyFirstprivate(VD, RI.getType().getNonReferenceType());
+
       if (!RI.getType()->isAnyPointerType()) {
         // We have to signal to the runtime captures passed by value that are
         // not pointers.
@@ -9600,6 +9642,13 @@ class MappableExprsHandler {
             OpenMPOffloadMappingFlags::OMP_MAP_LITERAL);
         CombinedInfo.Sizes.push_back(CGF.Builder.CreateIntCast(
             CGF.getTypeSize(RI.getType()), CGF.Int64Ty, /*isSigned=*/true));
+      } else if (IsFirstprivate) {
+        // Firstprivate pointers should be passed by value (as literals)
+        // without performing a present table lookup at runtime.
+        CombinedInfo.Types.push_back(
+            OpenMPOffloadMappingFlags::OMP_MAP_LITERAL);
+        // Use zero size for pointer literals (just passing the pointer value)
+        CombinedInfo.Sizes.push_back(llvm::Constant::getNullValue(CGF.Int64Ty));
       } else {
         // Pointers are implicitly mapped with a zero size and no flags
         // (other than first map that is added for all implicit maps).
@@ -9613,26 +9662,31 @@ class MappableExprsHandler {
       assert(CI.capturesVariable() && "Expected captured reference.");
       const auto *PtrTy = cast<ReferenceType>(RI.getType().getTypePtr());
       QualType ElementType = PtrTy->getPointeeType();
-      CombinedInfo.Sizes.push_back(CGF.Builder.CreateIntCast(
-          CGF.getTypeSize(ElementType), CGF.Int64Ty, /*isSigned=*/true));
-      // The default map type for a scalar/complex type is 'to' because by
-      // default the value doesn't have to be retrieved. For an aggregate
-      // type, the default is 'tofrom'.
-      CombinedInfo.Types.push_back(getMapModifiersForPrivateClauses(CI));
       const VarDecl *VD = CI.getCapturedVar();
-      auto I = FirstPrivateDecls.find(VD);
+      bool IsFirstprivate = isEffectivelyFirstprivate(VD, ElementType);
       CombinedInfo.Exprs.push_back(VD->getCanonicalDecl());
       CombinedInfo.BasePointers.push_back(CV);
       CombinedInfo.DevicePtrDecls.push_back(nullptr);
       CombinedInfo.DevicePointers.push_back(DeviceInfoTy::None);
-      if (I != FirstPrivateDecls.end() && ElementType->isAnyPointerType()) {
-        Address PtrAddr = CGF.EmitLoadOfReference(CGF.MakeAddrLValue(
-            CV, ElementType, CGF.getContext().getDeclAlign(VD),
-            AlignmentSource::Decl));
-        CombinedInfo.Pointers.push_back(PtrAddr.emitRawPointer(CGF));
+
+      // For firstprivate pointers, pass by value instead of dereferencing
+      if (IsFirstprivate && ElementType->isAnyPointerType()) {
+        // Treat as a literal value (pass the pointer value itself)
+        CombinedInfo.Pointers.push_back(CV);
+        // Use zero size for pointer literals
+        CombinedInfo.Sizes.push_back(llvm::Constant::getNullValue(CGF.Int64Ty));
+        CombinedInfo.Types.push_back(
+            OpenMPOffloadMappingFlags::OMP_MAP_LITERAL);
       } else {
+        CombinedInfo.Sizes.push_back(CGF.Builder.CreateIntCast(
+            CGF.getTypeSize(ElementType), CGF.Int64Ty, /*isSigned=*/true));
+        // The default map type for a scalar/complex type is 'to' because by
+        // default the value doesn't have to be retrieved. For an aggregate
+        // type, the default is 'tofrom'.
+        CombinedInfo.Types.push_back(getMapModifiersForPrivateClauses(CI));
         CombinedInfo.Pointers.push_back(CV);
       }
+      auto I = FirstPrivateDecls.find(VD);
       if (I != FirstPrivateDecls.end())
         IsImplicit = I->getSecond();
     }
diff --git a/clang/test/OpenMP/target_codegen.cpp b/clang/test/OpenMP/target_codegen.cpp
index ff126fbe4d02c..d2c0004204016 100644
--- a/clang/test/OpenMP/target_codegen.cpp
+++ b/clang/test/OpenMP/target_codegen.cpp
@@ -78,7 +78,7 @@
 // code and have mapped arguments, and only 6 have all-constant map sizes.
 
 // CHECK-DAG: [[SIZET:@.+]] = private unnamed_addr constant [2 x i64] [i64 0, i64 4]
-// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 544, i64 800]
+// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 800, i64 800]
 // CHECK-DAG: [[SIZET2:@.+]] = private unnamed_addr constant [1 x i{{32|64}}] [i64 2]
 // CHECK-DAG: [[MAPT2:@.+]] = private unnamed_addr constant [1 x i64] [i64 800]
 // CHECK-DAG: [[SIZET3:@.+]] = private unnamed_addr constant [2 x i64] [i64 4, i64 2]
diff --git a/clang/test/OpenMP/target_defaultmap_codegen_01.cpp b/clang/test/OpenMP/target_defaultmap_codegen_01.cpp
index 0936aa08e21e7..652794e6b9d9e 100644
--- a/clang/test/OpenMP/target_defaultmap_codegen_01.cpp
+++ b/clang/test/OpenMP/target_defaultmap_codegen_01.cpp
@@ -735,7 +735,7 @@ void explicit_maps_single (){
 
 // CK14: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer
 // Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544
-// CK14: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 544]
+// CK14: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 800]
 
 // CK14-LABEL: explicit_maps_single{{.*}}(
 void explicit_maps_single (){
@@ -1235,7 +1235,7 @@ void implicit_maps_struct (int a){
 
 // CK22-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] zeroinitializer
 // Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544
-// CK22-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 544]
+// CK22-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 800]
 
 // CK22-LABEL: implicit_maps_pointer{{.*}}(
 void implicit_maps_pointer (){
diff --git a/clang/test/OpenMP/target_depend_codegen.cpp b/clang/test/OpenMP/target_depend_codegen.cpp
index 73ffa120452c1..b9c30f3e73dd7 100644
--- a/clang/test/OpenMP/target_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_depend_codegen.cpp
@@ -44,7 +44,7 @@
 // TCHECK: [[ENTTY:%.+]] = type { i64, i16, i16, i32, ptr, ptr, i64, i64, ptr }
 
 // CHECK-DAG: [[SIZET:@.+]] = private unnamed_addr constant [3 x i64] [i64 0, i64 4, i64 {{16|12}}]
-// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [3 x i64] [i64 544, i64 800, i64 3]
+// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [3 x i64] [i64 800, i64 800, i64 3]
 // CHECK-DAG: @{{.*}} = weak constant i8 0
 
 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
diff --git a/clang/test/OpenMP/target_firstprivate_pointer_codegen.cpp b/clang/test/OpenMP/target_firstprivate_pointer_codegen.cpp
new file mode 100644
index 0000000000000..eab131b814841
--- /dev/null
+++ b/clang/test/OpenMP/target_firstprivate_pointer_codegen.cpp
@@ -0,0 +1,169 @@
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+/// ========================================================================
+/// Test: Firstprivate pointer handling in OpenMP target regions
+/// ========================================================================
+///
+/// This test verifies that pointers with firstprivate semantics get the
+/// OMP_MAP_LITERAL flag, enabling the runtime to pass pointer values directly
+/// without performing present table lookups.
+///
+/// Map type values:
+///   288 = OMP_MAP_TARGET_PARAM (32) + OMP_MAP_LITERAL (256)
+///         Used for explicit firstprivate(ptr)
+///
+///   800 = OMP_MAP_TARGET_PARAM (32) + OMP_MAP_LITERAL (256) + OMP_MAP_IS_PTR (512)
+///         Used for implicit firstprivate pointers (e.g., from defaultmap clauses)
+///         Note: 512 is OMP_MAP_IS_PTR, not IMPLICIT. Implicitness is tracked separately.
+///
+///   544 = OMP_MAP_TARGET_PARAM (32) + OMP_MAP_IS_PTR (512)
+///         Incorrect behavior - missing LITERAL flag, causes runtime present table lookup
+///
+
+///==========================================================================
+/// Test 1: Explicit firstprivate(pointer) → map type 288
+///==========================================================================
+
+// CHECK-DAG: @.offload_maptypes{{[^.]*}} = private unnamed_addr constant [1 x i64] [i64 288]
+// CHECK-DAG: @.offload_sizes{{[^.]*}} = private unnamed_addr constant [1 x i64] zeroinitializer
+
+void test1_explicit_firstprivate() {
+  double *ptr = nullptr;
+  
+  // Explicit firstprivate should generate map type 288
+  // (TARGET_PARAM | LITERAL, no IS_PTR flag for explicit clauses)
+  #pragma omp target firstprivate(ptr)
+  {
+    if (ptr) ptr[0] = 1.0;
+  }
+}
+
+///==========================================================================
+/// Test 2: defaultmap(firstprivate:pointer) → map type 544
+///==========================================================================
+
+// CHECK-DAG: @.offload_maptypes{{.*}} = private unnamed_addr constant [1 x i64] [i64 544]
+// CHECK-DAG: @.offload_sizes{{.*}} = private unnamed_addr constant [1 x i64] zeroinitializer
+
+void test2_defaultmap_firstprivate_pointer() {
+  double *ptr = nullptr;
+  
+  // defaultmap(firstprivate:pointer) creates implicit firstprivate
+  // Should generate map type 544 (TARGET_PARAM | TO) - pointers need runtime lookup, not literal
+  #pragma omp target defaultmap(firstprivate:pointer)
+  {
+    if (ptr) ptr[0] = 2.0;
+  }
+}
+
+///==========================================================================
+/// Test 3: defaultmap(firstprivate:scalar) with double → map type 800
+///==========================================================================
+
+// CHECK-DAG: @.offload_maptypes{{.*}} = private unnamed_addr constant [1 x i64] [i64 800]
+
+void test3_defaultmap_scalar_double() {
+  double d = 3.0;
+  
+  // OpenMP's "scalar" category excludes pointers but includes arithmetic types
+  // Double gets implicit firstprivate → map type 800
+  #pragma omp target defaultmap(firstprivate:scalar)
+  {
+    d += 1.0;
+  }
+}
+
+///==========================================================================
+/// Test 4: Pointer with defaultmap(firstprivate:scalar) → gets 544 (not firstprivate)
+///==========================================================================
+
+// CHECK-DAG: @.offload_maptypes{{.*}} = private unnamed_addr constant [1 x i64] [i64 544]
+// CHECK-DAG: @.offload_sizes{{.*}} = private unnamed_addr constant [1 x i64] zeroinitializer
+
+void test4_pointer_with_scalar_defaultmap() {
+  double *ptr = nullptr;
+  
+  // Note: defaultmap(firstprivate:scalar) does NOT apply to pointers (scalar excludes pointers).
+  // The pointer gets implicit map type 544 (OMP_MAP_TO | OMP_MAP_TARGET_PARAM) which means
+  // it will perform runtime lookup to get the device address, NOT treat it as a literal.
+  // This is the correct behavior - pointers need device address translation.
+  #pragma omp target defaultmap(firstprivate:scalar)
+  {
+    if (ptr) ptr[0] = 4.0;
+  }
+}
+
+///==========================================================================
+/// Test 5: Multiple pointers with explicit firstprivate → all get 288
+///==========================================================================
+
+// CHECK-DAG: @.offload_maptypes{{.*}} = private unnamed_addr constant [3 x i64] [i64 288, i64 288, i64 288]
+// CHECK-DAG: @.offload_sizes{{.*}} = private unnamed_addr constant [3 x i64] zeroinitializer
+
+void test5_multiple_firstprivate() {
+  int *a = nullptr;
+  float *b = nullptr;
+  double *c = nullptr;
+  
+  // All explicit firstprivate pointers get map type 288
+  #pragma omp target firstprivate(a, b, c)
+  {
+    if (a) a[0] = 6;
+    if (b) b[0] = 7.0f;
+    if (c) c[0] = 8.0;
+  }
+}
+
+///==========================================================================
+/// Test 6: Pointer to const with firstprivate → map type 288
+///==========================================================================
+
+// CHECK-DAG: @.offload_maptypes{{.*}} = private unnamed_addr constant [1 x i64] [i64 288]
+// CHECK-DAG: @.offload_sizes{{.*}} = private unnamed_addr constant [1 x i64] zeroinitializer
+
+void test6_const_pointer() {
+  const double *const_ptr = nullptr;
+  
+  // Const pointer with explicit firstprivate → 288
+  #pragma omp target firstprivate(const_ptr)
+  {
+    if (const_ptr) {
+      double val = const_ptr[0];
+      (void)val;
+    }
+  }
+}
+
+///==========================================================================
+/// Test 7: Pointer-to-pointer with firstprivate → map type 288
+///==========================================================================
+
+// CHECK-DAG: @.offload_maptypes{{.*}} = private unnamed_addr constant [1 x i64] [i64 288]
+// CHECK-DAG: @.offload_sizes{{.*}} = private unnamed_addr constant [1 x i64] zeroinitializer
+
+void test7_pointer_to_pointer() {
+  int **pp = nullptr;
+  
+  // Pointer-to-pointer with explicit firstprivate → 288
+  #pragma omp target firstprivate(pp)
+  {
+    if (pp && *pp) (*pp)[0] = 9;
+  }
+}
+
+///==========================================================================
+/// Verification: The key fix is that firstprivate pointers now include
+/// the LITERAL flag (256), which tells the runtime to pass the pointer
+/// value directly instead of performing a present table lookup.
+///
+/// Before fix: Pointers got 544 (TARGET_PARAM | IS_PTR) → runtime lookup
+/// After fix:  Pointers get 288 or 800 (includes LITERAL) → direct pass
+///==========================================================================
+
+#endif // HEADER
diff --git a/clang/test/OpenMP/target_map_codegen_01.cpp b/clang/test/OpenMP/target_map_codegen_01.cpp
index 9f3553d2377cb..2285e9b7964fa 100644
--- a/clang/test/OpenMP/target_map_codegen_01.cpp
+++ b/clang/test/OpenMP/target_map_codegen_01.cpp
@@ -38,12 +38,12 @@
 // CK2-LABEL: @.__omp_offloading_{{.*}}implicit_maps_reference{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 
 // CK2: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 4]
-// Map types: OMP_MAP_PRIVATE_VAL | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 800
+// Map types: OMP_MAP_LITERAL | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 800
 // CK2: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 800]
 // CK2-LABEL: @.__omp_offloading_{{.*}}implicit_maps_reference{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK2: [[SIZES2:@.+]] = {{.+}}constant [1 x i64] zeroinitializer
-// Map types: OMP_MAP_IS_PTR | OMP_MAP_IMPLICIT = 544
-// CK2: [[TYPES2:@.+]] = {{.+}}constant [1 x i64] [i64 544]
+// Map types: OMP_MAP_IS_PTR | OMP_MAP_LITERAL | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 800
+// CK2: [[TYPES2:@.+]] = {{.+}}constant [1 x i64] [i64 800]
 
 // CK2-LABEL: implicit_maps_reference{{.*}}(
 void implicit_maps_reference (int a, int *b){
diff --git a/clang/test/OpenMP/target_map_codegen_09.cpp b/clang/test/OpenMP/target_map_codegen_09.cpp
index eebd811b50c3d..2b825562d1802 100644
--- a/clang/test/OpenMP/target_map_codegen_09.cpp
+++ b/clang/test/OpenMP/target_map_codegen_09.cpp
@@ -36,8 +36,8 @@
 // CK10-LABEL: @.__omp_offloading_{{.*}}implicit_maps_pointer{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 
 // CK10-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] zeroinitializer
-// Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544
-// CK10-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 544]
+// Map types: OMP_MAP_IS_PTR | OMP_MAP_LITERAL | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 800
+// CK10-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 800]
 
 // CK10-LABEL: implicit_maps_pointer{{.*}}(
 void implicit_maps_pointer (){
diff --git a/clang/test/OpenMP/target_map_codegen_10.cpp b/clang/test/OpenMP/target_map_codegen_10.cpp
index 5f8f4dc7771a7..67c531169390b 100644
--- a/clang/test/OpenMP/target_map_codegen_10.cpp
+++ b/clang/test/OpenMP/target_map_codegen_10.cpp
@@ -32,7 +32,8 @@
 // CK11_5-DAG: [[SIZES:@.+]] = {{.+}}constant [2 x i64] [i64 16, i64 0]
 // Map types: OMP_MAP_TO  | OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 547
 // CK11_4-DAG: [[TYPES:@.+]] = {{.+}}constant [2 x i64] [i64 547, i64 547]
-// CK11_5-DAG: [[TYPES:@.+]] = {{.+}}constant [2 x i64] [i64 547, i64 544]
+// Map types: OMP_MAP_TO  | OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 547, OMP_MAP_IS_PTR | OMP_MAP_LITERAL | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 800
+// CK11_5-DAG: [[TYPES:@.+]] = {{.+}}constant [2 x i64] [i64 547, i64 800]
 
 // CK11-LABEL: implicit_maps_double_complex{{.*}}(
 void implicit_maps_double_complex (int a, int *b){
diff --git a/clang/test/OpenMP/target_map_codegen_26.cpp b/clang/test/OpenMP/target_map_codegen_26.cpp
index 2bc1092685ac3..00a42c017d7ee 100644
--- a/clang/test/OpenMP/target_map_codegen_26.cpp
+++ b/clang/test/OpenMP/target_map_codegen_26.cpp
@@ -35,7 +35,7 @@
 
 // CK27-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK27: [[SIZE00:@.+]] = private {{.*}}constant [1 ...
[truncated]

@github-actions
Copy link

github-actions bot commented Nov 26, 2025

✅ With the latest revision this PR passed the C/C++ code formatter.

@Saieiei Saieiei force-pushed the fix-openmp-firstprivate-pointer-v2 branch from e1f4a91 to 1e6111d Compare November 26, 2025 10:19
@Saieiei Saieiei changed the title [OpenMP] Fix defaultmap(firstprivate:pointer) handling for implicit c… [OpenMP] Fix defaultmap(firstprivate:pointer) handling Nov 26, 2025
@Saieiei
Copy link
Contributor Author

Saieiei commented Nov 28, 2025

Hi @dpalermo, I've created a new PR with the fix for the firstprivate pointer handling issue. The previous test failures (check-offload) that caused the revert have all been resolved. I've validated with check-clang, check-openmp, and check-offload. All passing cleanly now. Could you please review and approve when you get a chance? Thanks!
The Previous PR #167879

@Saieiei
Copy link
Contributor Author

Saieiei commented Dec 3, 2025

Hi @alexey-bataev , this PR replaces the previous PR that was reverted. We’ve fixed the test failures, and everything is passing now. Could you please review and approve? Thanks!

hasRequiresUnifiedSharedMemory(), /*HasRequiresDynamicAllocators*/ false);
Config.setDefaultTargetAS(
CGM.getContext().getTargetInfo().getTargetAddressSpace(LangAS::Default));
Config.setRuntimeCC(CGM.getRuntimeCC());
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I was asked by @dpalermo to have a little look at the PR, it all LGTM, except I'm wondering why this line is removed? Seems like it might be an artifact from a downstream change you have, but I could be wrong.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It’s an unintended rebase artifact, the branch was based on an upstream before #168608 was merged. I have rebased on the latest main and updated the patch to keep this line. Thanks for catching this!

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

No problem at all, I've done the same a few times! Thank you for the fix.

@Saieiei Saieiei force-pushed the fix-openmp-firstprivate-pointer-v2 branch from 1e6111d to def235f Compare December 5, 2025 04:48
@chandraghale chandraghale merged commit 9e9e64a into llvm:main Dec 6, 2025
10 checks passed
@llvm-ci
Copy link
Collaborator

llvm-ci commented Dec 6, 2025

LLVM Buildbot has detected a new failure on builder llvm-clang-x86_64-gcc-ubuntu running on sie-linux-worker3 while building clang at step 6 "test-build-unified-tree-check-all".

Full details are available at: https://lab.llvm.org/buildbot/#/builders/174/builds/28510

Here is the relevant piece of the build log for the reference
Step 6 (test-build-unified-tree-check-all) failure: test (failure)
******************** TEST 'SanitizerCommon-asan-x86_64-Linux :: Linux/soft_rss_limit_mb_test.cpp' FAILED ********************
Exit Code: 1

Command Output (stdout):
--
# RUN: at line 2
/home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/build/./bin/clang  --driver-mode=g++ -gline-tables-only -fsanitize=address  -m64 -funwind-tables  -I/home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/compiler-rt/test -ldl -O2 /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/compiler-rt/test/sanitizer_common/TestCases/Linux/soft_rss_limit_mb_test.cpp -o /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/build/runtimes/runtimes-bins/compiler-rt/test/sanitizer_common/asan-x86_64-Linux/Linux/Output/soft_rss_limit_mb_test.cpp.tmp
# executed command: /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/build/./bin/clang --driver-mode=g++ -gline-tables-only -fsanitize=address -m64 -funwind-tables -I/home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/compiler-rt/test -ldl -O2 /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/compiler-rt/test/sanitizer_common/TestCases/Linux/soft_rss_limit_mb_test.cpp -o /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/build/runtimes/runtimes-bins/compiler-rt/test/sanitizer_common/asan-x86_64-Linux/Linux/Output/soft_rss_limit_mb_test.cpp.tmp
# note: command had no output on stdout or stderr
# RUN: at line 5
env ASAN_OPTIONS=soft_rss_limit_mb=220:quarantine_size=1:allocator_may_return_null=1      /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/build/runtimes/runtimes-bins/compiler-rt/test/sanitizer_common/asan-x86_64-Linux/Linux/Output/soft_rss_limit_mb_test.cpp.tmp 2>&1 | FileCheck /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/compiler-rt/test/sanitizer_common/TestCases/Linux/soft_rss_limit_mb_test.cpp -check-prefix=CHECK_MAY_RETURN_1
# executed command: env ASAN_OPTIONS=soft_rss_limit_mb=220:quarantine_size=1:allocator_may_return_null=1 /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/build/runtimes/runtimes-bins/compiler-rt/test/sanitizer_common/asan-x86_64-Linux/Linux/Output/soft_rss_limit_mb_test.cpp.tmp
# note: command had no output on stdout or stderr
# executed command: FileCheck /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/compiler-rt/test/sanitizer_common/TestCases/Linux/soft_rss_limit_mb_test.cpp -check-prefix=CHECK_MAY_RETURN_1
# note: command had no output on stdout or stderr
# RUN: at line 6
env ASAN_OPTIONS=soft_rss_limit_mb=220:quarantine_size=1:allocator_may_return_null=0 not  /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/build/runtimes/runtimes-bins/compiler-rt/test/sanitizer_common/asan-x86_64-Linux/Linux/Output/soft_rss_limit_mb_test.cpp.tmp 2>&1 | FileCheck /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/compiler-rt/test/sanitizer_common/TestCases/Linux/soft_rss_limit_mb_test.cpp -check-prefix=CHECK_MAY_RETURN_0 --implicit-check-not="returned null"
# executed command: env ASAN_OPTIONS=soft_rss_limit_mb=220:quarantine_size=1:allocator_may_return_null=0 not /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/build/runtimes/runtimes-bins/compiler-rt/test/sanitizer_common/asan-x86_64-Linux/Linux/Output/soft_rss_limit_mb_test.cpp.tmp
# note: command had no output on stdout or stderr
# executed command: FileCheck /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/compiler-rt/test/sanitizer_common/TestCases/Linux/soft_rss_limit_mb_test.cpp -check-prefix=CHECK_MAY_RETURN_0 '--implicit-check-not=returned null'
# note: command had no output on stdout or stderr
# RUN: at line 10
env ASAN_OPTIONS=soft_rss_limit_mb=220:quarantine_size=1:allocator_may_return_null=0:can_use_proc_maps_statm=0 not  /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/build/runtimes/runtimes-bins/compiler-rt/test/sanitizer_common/asan-x86_64-Linux/Linux/Output/soft_rss_limit_mb_test.cpp.tmp 2>&1 | FileCheck /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/compiler-rt/test/sanitizer_common/TestCases/Linux/soft_rss_limit_mb_test.cpp -check-prefix=CHECK_MAY_RETURN_0 --implicit-check-not="returned null"
# executed command: env ASAN_OPTIONS=soft_rss_limit_mb=220:quarantine_size=1:allocator_may_return_null=0:can_use_proc_maps_statm=0 not /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/build/runtimes/runtimes-bins/compiler-rt/test/sanitizer_common/asan-x86_64-Linux/Linux/Output/soft_rss_limit_mb_test.cpp.tmp
# note: command had no output on stdout or stderr
# executed command: FileCheck /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/compiler-rt/test/sanitizer_common/TestCases/Linux/soft_rss_limit_mb_test.cpp -check-prefix=CHECK_MAY_RETURN_0 '--implicit-check-not=returned null'
# .---command stderr------------
# | �[1m/home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/compiler-rt/test/sanitizer_common/TestCases/Linux/soft_rss_limit_mb_test.cpp:72:24: �[0m�[0;1;31merror: �[0m�[1mCHECK_MAY_RETURN_0: expected string not found in input
�[0m# | �[1m�[0m// CHECK_MAY_RETURN_0: Some of the malloc calls returned non-null:
# | �[0;1;32m                       ^
�[0m# | �[0;1;32m�[0m�[1m<stdin>:1:24: �[0m�[0;1;30mnote: �[0m�[1mscanning from here
�[0m# | �[1m�[0m[0] allocating 32 times
# | �[0;1;32m                       ^
�[0m# | �[0;1;32m�[0m�[1m<stdin>:9:55: �[0m�[0;1;30mnote: �[0m�[1mpossible intended match here
�[0m# | �[1m�[0m==3605013==HINT: if you don't care about these errors you may set allocator_may_return_null=1
# | �[0;1;32m                                                      ^
�[0m# | �[0;1;32m�[0m
# | Input file: <stdin>
# | Check file: /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/compiler-rt/test/sanitizer_common/TestCases/Linux/soft_rss_limit_mb_test.cpp
# | 
# | -dump-input=help explains the following input dump.
# | 
# | Input was:
# | <<<<<<
# | �[1m�[0m�[0;1;30m            1: �[0m�[1m�[0;1;46m[0] �[0mallocating 32 times�[0;1;46m �[0m
# | �[0;1;32mcheck:71           ^~~~~~~~~~~~~~~~~~~
�[0m# | �[0;1;32m�[0m�[0;1;32mnot:imp1       X~~~
�[0m# | �[0;1;32m�[0m�[0;1;31mcheck:72'0                            X error: no match found
�[0m# | �[0;1;31m�[0m�[0;1;30m            2: �[0m�[1m�[0;1;46m [0] �[0m
# | �[0;1;31mcheck:72'0     ~~~~~
...

honeygoyal pushed a commit to honeygoyal/llvm-project that referenced this pull request Dec 9, 2025
This fixes a bug where pointers from defaultmap(firstprivate:pointer)
were incorrectly treated as firstprivate literals, causing
OMP_MAP_LITERAL to be set. This prevented the runtime from performing
device address lookup.
Realted PR llvm#167879

Co-authored-by: Sairudra More <[email protected]>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

clang:codegen IR generation bugs: mangling, exceptions, etc. clang:openmp OpenMP related changes to Clang clang Clang issues not falling into any other category

Projects

None yet

Development

Successfully merging this pull request may close these issues.

6 participants