[llvm-branch-commits] [clang] cefbc48 - Revert "[OpenMP] Fix firstprivate pointer handling in target regions (#167879)"
via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Fri Nov 21 18:57:31 PST 2025
Author: dpalermo
Date: 2025-11-21T20:57:27-06:00
New Revision: cefbc487d482c607b3292d4356f2eeac8783c719
URL: https://github.com/llvm/llvm-project/commit/cefbc487d482c607b3292d4356f2eeac8783c719
DIFF: https://github.com/llvm/llvm-project/commit/cefbc487d482c607b3292d4356f2eeac8783c719.diff
LOG: Revert "[OpenMP] Fix firstprivate pointer handling in target regions (#167879)"
This reverts commit 622f72f4bef8b177e1e4f318465260fbdb7711ef.
Added:
Modified:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/test/OpenMP/target_codegen.cpp
clang/test/OpenMP/target_defaultmap_codegen_01.cpp
clang/test/OpenMP/target_depend_codegen.cpp
clang/test/OpenMP/target_map_codegen_01.cpp
clang/test/OpenMP/target_map_codegen_09.cpp
clang/test/OpenMP/target_map_codegen_10.cpp
clang/test/OpenMP/target_map_codegen_26.cpp
clang/test/OpenMP/target_parallel_depend_codegen.cpp
clang/test/OpenMP/target_parallel_for_depend_codegen.cpp
clang/test/OpenMP/target_parallel_for_simd_depend_codegen.cpp
clang/test/OpenMP/target_simd_depend_codegen.cpp
clang/test/OpenMP/target_teams_depend_codegen.cpp
clang/test/OpenMP/target_teams_distribute_depend_codegen.cpp
clang/test/OpenMP/target_teams_distribute_parallel_for_depend_codegen.cpp
clang/test/OpenMP/target_teams_distribute_parallel_for_simd_depend_codegen.cpp
clang/test/OpenMP/target_teams_distribute_simd_depend_codegen.cpp
Removed:
clang/test/OpenMP/target_firstprivate_pointer_codegen.cpp
################################################################################
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 65a7daca9fcf1..a8255ac74cfcf 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -28,7 +28,6 @@
#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"
@@ -7212,9 +7211,6 @@ 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<
@@ -8993,10 +8989,6 @@ 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())
@@ -9574,35 +9566,6 @@ 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
- if (FirstPrivateDecls.count(VD))
- return true;
-
- // 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,
@@ -9630,9 +9593,6 @@ 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.
@@ -9640,13 +9600,6 @@ 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).
@@ -9660,31 +9613,26 @@ 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();
- bool IsFirstprivate = isEffectivelyFirstprivate(VD, ElementType);
+ auto I = FirstPrivateDecls.find(VD);
CombinedInfo.Exprs.push_back(VD->getCanonicalDecl());
CombinedInfo.BasePointers.push_back(CV);
CombinedInfo.DevicePtrDecls.push_back(nullptr);
CombinedInfo.DevicePointers.push_back(DeviceInfoTy::None);
-
- // 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);
+ 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));
} 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 d2c0004204016..ff126fbe4d02c 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 800, i64 800]
+// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 544, 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 652794e6b9d9e..0936aa08e21e7 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 800]
+// CK14: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 544]
// 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 800]
+// CK22-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 544]
// 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 b9c30f3e73dd7..73ffa120452c1 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 800, i64 800, i64 3]
+// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [3 x i64] [i64 544, 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
deleted file mode 100644
index 326bc812d7d33..0000000000000
--- a/clang/test/OpenMP/target_firstprivate_pointer_codegen.cpp
+++ /dev/null
@@ -1,169 +0,0 @@
-// 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_01.cpp b/clang/test/OpenMP/target_map_codegen_01.cpp
index 2285e9b7964fa..9f3553d2377cb 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_LITERAL | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 800
+// Map types: OMP_MAP_PRIVATE_VAL | 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_LITERAL | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 800
-// CK2: [[TYPES2:@.+]] = {{.+}}constant [1 x i64] [i64 800]
+// Map types: OMP_MAP_IS_PTR | OMP_MAP_IMPLICIT = 544
+// CK2: [[TYPES2:@.+]] = {{.+}}constant [1 x i64] [i64 544]
// 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 2b825562d1802..eebd811b50c3d 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_IS_PTR | OMP_MAP_LITERAL | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 800
-// CK10-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 800]
+// Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544
+// CK10-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 544]
// 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 67c531169390b..5f8f4dc7771a7 100644
--- a/clang/test/OpenMP/target_map_codegen_10.cpp
+++ b/clang/test/OpenMP/target_map_codegen_10.cpp
@@ -32,8 +32,7 @@
// 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]
-// 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_5-DAG: [[TYPES:@.+]] = {{.+}}constant [2 x i64] [i64 547, i64 544]
// 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 00a42c017d7ee..2bc1092685ac3 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 800]
+// CK27: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i64] [i64 544]
// 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 288]
+// CK27: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i64] [i64 32]
// 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 106c34518cce6..52264ee79e123 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 800, i64 800]
+// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 544, 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 ae01116b83101..aec4feda15cf0 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 800, i64 800]
+// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 544, 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 2e8731e766e11..62f772ea79156 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 800, i64 800]
+// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 544, 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 4a0196b996d20..d813173dffbef 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 800, i64 800]
+// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 544, 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 65c3d88103b76..b2280b80995fb 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 800, i64 800]
+// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 544, 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 de5e1f3ff7991..f9306d3091ee8 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 800, i64 800]
+// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 544, 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 db5fa66cf707b..4d93515c738ed 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 800, i64 800]
+// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 544, 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 c462bd21f6aa7..ed760d7e2e000 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 800, i64 800]
+// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 544, 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 a7f5ac38abba8..9335c65bb2296 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 800, i64 800]
+// OMP45-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 544, 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 800, i64 800, i64 800]
+// OMP50-DAG: [[MAPT:@.+]] = private unnamed_addr constant [3 x i64] [i64 544, i64 800, i64 800]
// OMP50-DAG: @{{.*}} = weak constant i8 0
More information about the llvm-branch-commits
mailing list