[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 01:22:32 PST 2025
https://github.com/Saieiei created https://github.com/llvm/llvm-project/pull/169622
…aptures
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.
>From e1f4a91429578ad6aae67625529b614ea83bee41 Mon Sep 17 00:00:00 2001
From: Sairudra More <moresair at pe31.hpc.amslabs.hpecorp.net>
Date: Wed, 26 Nov 2025 03:20:07 -0600
Subject: [PATCH] [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.
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 78 ++++++--
clang/test/OpenMP/target_codegen.cpp | 2 +-
.../OpenMP/target_defaultmap_codegen_01.cpp | 4 +-
clang/test/OpenMP/target_depend_codegen.cpp | 2 +-
.../target_firstprivate_pointer_codegen.cpp | 169 ++++++++++++++++++
clang/test/OpenMP/target_map_codegen_01.cpp | 6 +-
clang/test/OpenMP/target_map_codegen_09.cpp | 4 +-
clang/test/OpenMP/target_map_codegen_10.cpp | 3 +-
clang/test/OpenMP/target_map_codegen_26.cpp | 4 +-
.../OpenMP/target_parallel_depend_codegen.cpp | 2 +-
.../target_parallel_for_depend_codegen.cpp | 2 +-
...arget_parallel_for_simd_depend_codegen.cpp | 2 +-
.../OpenMP/target_simd_depend_codegen.cpp | 2 +-
.../OpenMP/target_teams_depend_codegen.cpp | 2 +-
...target_teams_distribute_depend_codegen.cpp | 2 +-
...distribute_parallel_for_depend_codegen.cpp | 2 +-
...ibute_parallel_for_simd_depend_codegen.cpp | 2 +-
...t_teams_distribute_simd_depend_codegen.cpp | 4 +-
18 files changed, 258 insertions(+), 34 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..be86d65a74897 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"
@@ -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,6 +9632,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 +9642,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 +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();
}
diff --git a/clang/test/OpenMP/target_codegen.cpp b/clang/test/OpenMP/target_codegen.cpp
index ff126fbe4d02c..d2c0004204016 100644
--- a/clang/test/OpenMP/target_codegen.cpp
+++ b/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]
diff --git a/clang/test/OpenMP/target_defaultmap_codegen_01.cpp b/clang/test/OpenMP/target_defaultmap_codegen_01.cpp
index 0936aa08e21e7..652794e6b9d9e 100644
--- a/clang/test/OpenMP/target_defaultmap_codegen_01.cpp
+++ b/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 (){
diff --git a/clang/test/OpenMP/target_depend_codegen.cpp b/clang/test/OpenMP/target_depend_codegen.cpp
index 73ffa120452c1..b9c30f3e73dd7 100644
--- a/clang/test/OpenMP/target_depend_codegen.cpp
+++ b/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]]
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..eab131b814841
--- /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 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
diff --git a/clang/test/OpenMP/target_map_codegen_01.cpp b/clang/test/OpenMP/target_map_codegen_01.cpp
index 9f3553d2377cb..2285e9b7964fa 100644
--- a/clang/test/OpenMP/target_map_codegen_01.cpp
+++ b/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){
diff --git a/clang/test/OpenMP/target_map_codegen_09.cpp b/clang/test/OpenMP/target_map_codegen_09.cpp
index eebd811b50c3d..2b825562d1802 100644
--- a/clang/test/OpenMP/target_map_codegen_09.cpp
+++ 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 (){
diff --git a/clang/test/OpenMP/target_map_codegen_10.cpp b/clang/test/OpenMP/target_map_codegen_10.cpp
index 5f8f4dc7771a7..67c531169390b 100644
--- a/clang/test/OpenMP/target_map_codegen_10.cpp
+++ b/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){
diff --git a/clang/test/OpenMP/target_map_codegen_26.cpp b/clang/test/OpenMP/target_map_codegen_26.cpp
index 2bc1092685ac3..00a42c017d7ee 100644
--- a/clang/test/OpenMP/target_map_codegen_26.cpp
+++ 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
diff --git a/clang/test/OpenMP/target_parallel_depend_codegen.cpp b/clang/test/OpenMP/target_parallel_depend_codegen.cpp
index 52264ee79e123..106c34518cce6 100644
--- a/clang/test/OpenMP/target_parallel_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_parallel_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 [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: @{{.*}} = weak constant i8 0
// TCHECK: @{{.+}} = weak constant [[ENTTY]]
diff --git a/clang/test/OpenMP/target_parallel_for_depend_codegen.cpp b/clang/test/OpenMP/target_parallel_for_depend_codegen.cpp
index aec4feda15cf0..ae01116b83101 100644
--- a/clang/test/OpenMP/target_parallel_for_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_parallel_for_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 [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: @{{.*}} = weak constant i8 0
// TCHECK: @{{.+}} = weak constant [[ENTTY]]
diff --git a/clang/test/OpenMP/target_parallel_for_simd_depend_codegen.cpp b/clang/test/OpenMP/target_parallel_for_simd_depend_codegen.cpp
index 62f772ea79156..2e8731e766e11 100644
--- a/clang/test/OpenMP/target_parallel_for_simd_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_parallel_for_simd_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 [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: @{{.*}} = weak constant i8 0
// TCHECK: @{{.+}} = weak constant [[ENTTY]]
diff --git a/clang/test/OpenMP/target_simd_depend_codegen.cpp b/clang/test/OpenMP/target_simd_depend_codegen.cpp
index d813173dffbef..4a0196b996d20 100644
--- a/clang/test/OpenMP/target_simd_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_simd_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 [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: @{{.*}} = weak constant i8 0
// TCHECK: @{{.+}} = weak constant [[ENTTY]]
diff --git a/clang/test/OpenMP/target_teams_depend_codegen.cpp b/clang/test/OpenMP/target_teams_depend_codegen.cpp
index b2280b80995fb..65c3d88103b76 100644
--- a/clang/test/OpenMP/target_teams_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_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 [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: @{{.*}} = weak constant i8 0
// TCHECK: @{{.+}} = weak constant [[ENTTY]]
diff --git a/clang/test/OpenMP/target_teams_distribute_depend_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_depend_codegen.cpp
index f9306d3091ee8..de5e1f3ff7991 100644
--- a/clang/test/OpenMP/target_teams_distribute_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_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 [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: @{{.*}} = weak constant i8 0
// TCHECK: @{{.+}} = weak constant [[ENTTY]]
diff --git a/clang/test/OpenMP/target_teams_distribute_parallel_for_depend_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_parallel_for_depend_codegen.cpp
index 4d93515c738ed..db5fa66cf707b 100644
--- a/clang/test/OpenMP/target_teams_distribute_parallel_for_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_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 [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: @{{.*}} = weak constant i8 0
// TCHECK: @{{.+}} = weak constant [[ENTTY]]
diff --git a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_depend_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_depend_codegen.cpp
index ed760d7e2e000..c462bd21f6aa7 100644
--- a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_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 [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: @{{.*}} = weak constant i8 0
// TCHECK: @{{.+}} = weak constant [[ENTTY]]
diff --git a/clang/test/OpenMP/target_teams_distribute_simd_depend_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_simd_depend_codegen.cpp
index 9335c65bb2296..a7f5ac38abba8 100644
--- a/clang/test/OpenMP/target_teams_distribute_simd_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_simd_depend_codegen.cpp
@@ -53,11 +53,11 @@
// TCHECK: [[ENTTY:%.+]] = type { i64, i16, i16, i32, ptr, ptr, i64, i64, ptr }
// OMP45-DAG: [[SIZET:@.+]] = private unnamed_addr constant [2 x i64] [i64 0, i64 4]
-// OMP45-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 544, i64 800]
+// OMP45-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 800, i64 800]
// OMP45-DAG: @{{.*}} = weak constant i8 0
// OMP50-DAG: [[SIZET:@.+]] = private unnamed_addr constant [3 x i64] [i64 0, i64 4, i64 1]
-// OMP50-DAG: [[MAPT:@.+]] = private unnamed_addr constant [3 x i64] [i64 544, i64 800, i64 800]
+// OMP50-DAG: [[MAPT:@.+]] = private unnamed_addr constant [3 x i64] [i64 800, i64 800, i64 800]
// OMP50-DAG: @{{.*}} = weak constant i8 0
More information about the cfe-commits
mailing list