[OpenMP] Fix defaultmap(firstprivate:pointer) handling (llvm#169622)

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]>

Author: 2 people

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<
@@ -9007,6 +9011,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())
@@ -9584,6 +9592,36 @@ class MappableExprsHandler {
     }
   }
 
+  /// Check if a variable should be treated as firstprivate due to explicit
+  /// firstprivate clause or defaultmap(firstprivate:...).
+  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
+
+    // Check defaultmap(firstprivate:scalar) for scalar types
+    if (DefaultmapFirstprivateKinds.count(OMPC_DEFAULTMAP_scalar)) {
+      if (Type->isScalarType())
+        return true;
+    }
+
+    // Check defaultmap(firstprivate:pointer) for pointer types
+    if (DefaultmapFirstprivateKinds.count(OMPC_DEFAULTMAP_pointer)) {
+      if (Type->isAnyPointerType())
+        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,
@@ -9611,13 +9649,23 @@ 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.
         CombinedInfo.Types.push_back(
             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).
@@ -9631,26 +9679,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();
     }

clang/test/OpenMP/target_defaultmap_codegen_01.cpp

@@ -734,8 +734,8 @@ void explicit_maps_single (){
 // CK14-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0
 
 // 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]
+// Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_LITERAL | OMP_MAP_IMPLICIT = 800
+// CK14: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 800]
 
 // CK14-LABEL: explicit_maps_single{{.*}}(
 void explicit_maps_single (){

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 800
+///==========================================================================
+
+// CHECK-DAG: @.offload_maptypes{{.*}} = private unnamed_addr constant [1 x i64] [i64 800]
+// 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 800 (TARGET_PARAM | LITERAL | IS_PTR)
+  #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) → map type 800
+///==========================================================================
+
+// CHECK-DAG: @.offload_maptypes{{.*}} = private unnamed_addr constant [1 x i64] [i64 800]
+// 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).
+  // However, the pointer still gets 800 because in OpenMP 5.0+, pointers without explicit
+  // data-sharing attributes are implicitly firstprivate and lowered as IS_PTR|LITERAL|TARGET_PARAM.
+  // This is the default pointer behavior, NOT due to the scalar defaultmap.
+  #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

clang/test/OpenMP/target_map_codegen_26.cpp

@@ -52,7 +52,7 @@
 // CK27-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK27-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK27: [[SIZE05:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer
-// CK27: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i64] [i64 32]
+// CK27: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i64] [i64 288]
 
 // CK27-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK27-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0