Skip to content

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

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 1 commit into
base: main
Choose a base branch
from
+237 −16
Open

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

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
78 changes: 65 additions & 13 deletions clang/lib/CodeGen/CGOpenMPRuntime.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"
Expand Down Expand Up @@ -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();
Expand Down Expand Up @@ -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<
Expand Down Expand Up @@ -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())
Expand Down Expand Up @@ -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,
Expand Down Expand Up @@ -9593,13 +9630,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).
Expand All @@ -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();
}
Expand Down
4 changes: 2 additions & 2 deletions clang/test/OpenMP/target_defaultmap_codegen_01.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 (){
Expand Down
169 changes: 169 additions & 0 deletions clang/test/OpenMP/target_firstprivate_pointer_codegen.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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
2 changes: 1 addition & 1 deletion clang/test/OpenMP/target_map_codegen_26.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down