[OpenMP] Fix defaultmap(firstprivate:pointer) handling for implicit captures

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.

Author: Sairudra More

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,13 +9632,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).
@@ -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();
     }

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]

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 (){

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

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

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

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 (){

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

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 x i64] zeroinitializer
-// CK27: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i64] [i64 544]
+// CK27: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i64] [i64 800]
 
 // CK27-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK27: [[SIZE01:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer
@@ -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