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 <[email protected]> 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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
