[clang] [OpenMP] Fix defaultmap(firstprivate:pointer) handling for implicit c… (PR #169622)
Sairudra More via cfe-commits
cfe-commits at lists.llvm.org
Wed Nov 26 02:19:27 PST 2025
https://github.com/Saieiei updated https://github.com/llvm/llvm-project/pull/169622
>From 1e6111d2a09339fe296bd302bf7d5f97523a5765 Mon Sep 17 00:00:00 2001
From: Sairudra More <moresair at pe31.hpc.amslabs.hpecorp.net>
Date: Wed, 26 Nov 2025 04:18:53 -0600
Subject: [PATCH] [OpenMP] Add OMP_MAP_LITERAL flag for firstprivate pointers
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 78 ++++++--
.../OpenMP/target_defaultmap_codegen_01.cpp | 4 +-
.../target_firstprivate_pointer_codegen.cpp | 169 ++++++++++++++++++
clang/test/OpenMP/target_map_codegen_26.cpp | 2 +-
4 files changed, 237 insertions(+), 16 deletions(-)
create mode 100644 clang/test/OpenMP/target_firstprivate_pointer_codegen.cpp
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index a8255ac74cfcf..aa81b63e10de6 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"
@@ -1039,7 +1040,6 @@ CGOpenMPRuntime::CGOpenMPRuntime(CodeGenModule &CGM)
hasRequiresUnifiedSharedMemory(), /*HasRequiresDynamicAllocators*/ false);
Config.setDefaultTargetAS(
CGM.getContext().getTargetInfo().getTargetAddressSpace(LangAS::Default));
- Config.setRuntimeCC(CGM.getRuntimeCC());
OMPBuilder.setConfig(Config);
OMPBuilder.initialize();
@@ -7211,6 +7211,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 +8992,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 +9573,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,
@@ -9593,6 +9630,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 +9640,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 +9660,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_defaultmap_codegen_01.cpp b/clang/test/OpenMP/target_defaultmap_codegen_01.cpp
index 0936aa08e21e7..42b6fa6c5fc16 100644
--- a/clang/test/OpenMP/target_defaultmap_codegen_01.cpp
+++ b/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 (){
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..326bc812d7d33
--- /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 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
diff --git a/clang/test/OpenMP/target_map_codegen_26.cpp b/clang/test/OpenMP/target_map_codegen_26.cpp
index 2bc1092685ac3..c5fa463fa736f 100644
--- a/clang/test/OpenMP/target_map_codegen_26.cpp
+++ b/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
More information about the cfe-commits
mailing list