llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: None (dpalermo) <details> <summary>Changes</summary> Reverts llvm/llvm-project#<!-- -->167879 This PR is causing assertions in the check-offload tests: https://lab.llvm.org/staging/#/builders/105 https://lab.llvm.org/staging/#/builders/105/builds/37057 --- Patch is 27.79 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/169143.diff 18 Files Affected: - (modified) clang/lib/CodeGen/CGOpenMPRuntime.cpp (+12-64) - (modified) clang/test/OpenMP/target_codegen.cpp (+1-1) - (modified) clang/test/OpenMP/target_defaultmap_codegen_01.cpp (+2-2) - (modified) clang/test/OpenMP/target_depend_codegen.cpp (+1-1) - (removed) clang/test/OpenMP/target_firstprivate_pointer_codegen.cpp (-169) - (modified) clang/test/OpenMP/target_map_codegen_01.cpp (+3-3) - (modified) clang/test/OpenMP/target_map_codegen_09.cpp (+2-2) - (modified) clang/test/OpenMP/target_map_codegen_10.cpp (+1-2) - (modified) clang/test/OpenMP/target_map_codegen_26.cpp (+2-2) - (modified) clang/test/OpenMP/target_parallel_depend_codegen.cpp (+1-1) - (modified) clang/test/OpenMP/target_parallel_for_depend_codegen.cpp (+1-1) - (modified) clang/test/OpenMP/target_parallel_for_simd_depend_codegen.cpp (+1-1) - (modified) clang/test/OpenMP/target_simd_depend_codegen.cpp (+1-1) - (modified) clang/test/OpenMP/target_teams_depend_codegen.cpp (+1-1) - (modified) clang/test/OpenMP/target_teams_distribute_depend_codegen.cpp (+1-1) - (modified) clang/test/OpenMP/target_teams_distribute_parallel_for_depend_codegen.cpp (+1-1) - (modified) clang/test/OpenMP/target_teams_distribute_parallel_for_simd_depend_codegen.cpp (+1-1) - (modified) clang/test/OpenMP/target_teams_distribute_simd_depend_codegen.cpp (+2-2) ``````````diff 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] zero... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/169143 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
