llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: Sairudra More (Saieiei) <details> <summary>Changes</summary> …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. --- Patch is 28.14 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/169622.diff 18 Files Affected: - (modified) clang/lib/CodeGen/CGOpenMPRuntime.cpp (+66-12) - (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) - (added) 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 (+2-1) - (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 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 ... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/169622 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
