From b9b8ca233458f67c14a1fe259dd507430cc85ab1 Mon Sep 17 00:00:00 2001 From: Alexey Bataev Date: Tue, 1 Oct 2019 18:18:03 +0000 Subject: [PATCH] [OPENMP]Fix PR43330: OpenMP target: Mapping of partial arrays fails. Fixed calculation the size of the array sections. llvm-svn: 373374 --- clang/lib/CodeGen/CGOpenMPRuntime.cpp | 35 +++++++-- clang/test/OpenMP/target_map_codegen.cpp | 126 ++++++++++++++++++------------- 2 files changed, 102 insertions(+), 59 deletions(-) diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index 75971d8..64bfa40 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -7249,9 +7249,11 @@ private: OAE->getBase()->IgnoreParenImpCasts()) .getCanonicalType(); - // If there is no length associated with the expression, that means we - // are using the whole length of the base. - if (!OAE->getLength() && OAE->getColonLoc().isValid()) + // If there is no length associated with the expression and lower bound is + // not specified too, that means we are using the whole length of the + // base. + if (!OAE->getLength() && OAE->getColonLoc().isValid() && + !OAE->getLowerBound()) return CGF.getTypeSize(BaseTy); llvm::Value *ElemSize; @@ -7265,13 +7267,30 @@ private: // If we don't have a length at this point, that is because we have an // array section with a single element. - if (!OAE->getLength()) + if (!OAE->getLength() && OAE->getColonLoc().isInvalid()) return ElemSize; - llvm::Value *LengthVal = CGF.EmitScalarExpr(OAE->getLength()); - LengthVal = - CGF.Builder.CreateIntCast(LengthVal, CGF.SizeTy, /*isSigned=*/false); - return CGF.Builder.CreateNUWMul(LengthVal, ElemSize); + if (const Expr *LenExpr = OAE->getLength()) { + llvm::Value *LengthVal = CGF.EmitScalarExpr(OAE->getLength()); + LengthVal = CGF.EmitScalarConversion( + LengthVal, OAE->getLength()->getType(), + CGF.getContext().getSizeType(), OAE->getLength()->getExprLoc()); + return CGF.Builder.CreateNUWMul(LengthVal, ElemSize); + } + assert(!OAE->getLength() && OAE->getColonLoc().isValid() && + OAE->getLowerBound() && "expected array_section[lb:]."); + // Size = sizetype - lb * elemtype; + llvm::Value *LengthVal = CGF.getTypeSize(BaseTy); + llvm::Value *LBVal = CGF.EmitScalarExpr(OAE->getLowerBound()); + LBVal = CGF.EmitScalarConversion(LBVal, OAE->getLowerBound()->getType(), + CGF.getContext().getSizeType(), + OAE->getLowerBound()->getExprLoc()); + LBVal = CGF.Builder.CreateNUWMul(LBVal, ElemSize); + llvm::Value *Cmp = CGF.Builder.CreateICmpUGT(LengthVal, LBVal); + llvm::Value *TrueVal = CGF.Builder.CreateNUWSub(LengthVal, LBVal); + LengthVal = CGF.Builder.CreateSelect( + Cmp, TrueVal, llvm::ConstantInt::get(CGF.SizeTy, 0)); + return LengthVal; } return CGF.getTypeSize(ExprTy); } diff --git a/clang/test/OpenMP/target_map_codegen.cpp b/clang/test/OpenMP/target_map_codegen.cpp index ca66802..0a8198a 100644 --- a/clang/test/OpenMP/target_map_codegen.cpp +++ b/clang/test/OpenMP/target_map_codegen.cpp @@ -1323,172 +1323,176 @@ void implicit_maps_template_type_capture (int a){ // SIMD-ONLY18-NOT: {{__kmpc|__tgt}} #ifdef CK19 -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1510.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1514.region_id = weak constant i8 0 // CK19: [[SIZE00:@.+]] = private {{.*}}constant [1 x i64] [i64 4] // CK19: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i64] [i64 32] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1531.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1535.region_id = weak constant i8 0 // CK19: [[SIZE00n:@.+]] = private {{.*}}constant [1 x i64] [i64 4] // CK19: [[MTYPE00n:@.+]] = private {{.*}}constant [1 x i64] [i64 32] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1553.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1557.region_id = weak constant i8 0 // CK19: [[SIZE01:@.+]] = private {{.*}}constant [1 x i64] [i64 400] // CK19: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 33] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1572.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1576.region_id = weak constant i8 0 // CK19: [[SIZE02:@.+]] = private {{.*}}constant [1 x i64] [i64 240] // CK19: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i64] [i64 34] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1591.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1595.region_id = weak constant i8 0 // CK19: [[SIZE03:@.+]] = private {{.*}}constant [1 x i64] [i64 240] // CK19: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i64] [i64 35] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1610.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1614.region_id = weak constant i8 0 // CK19: [[SIZE04:@.+]] = private {{.*}}constant [1 x i64] [i64 400] // CK19: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i64] [i64 32] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1629.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1633.region_id = weak constant i8 0 // CK19: [[SIZE05:@.+]] = private {{.*}}constant [1 x i64] [i64 4] // CK19: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i64] [i64 33] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1652.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1656.region_id = weak constant i8 0 // CK19: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i64] [i64 35] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1675.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1679.region_id = weak constant i8 0 // CK19: [[MTYPE07:@.+]] = private {{.*}}constant [1 x i64] [i64 32] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1694.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1698.region_id = weak constant i8 0 // CK19: [[SIZE08:@.+]] = private {{.*}}constant [1 x i64] [i64 4] // CK19: [[MTYPE08:@.+]] = private {{.*}}constant [1 x i64] [i64 35] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1715.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1719.region_id = weak constant i8 0 // CK19: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] [i64 {{8|4}}] // CK19: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 34] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1736.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1740.region_id = weak constant i8 0 // CK19: [[SIZE10:@.+]] = private {{.*}}constant [1 x i64] [i64 240] // CK19: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i64] [i64 35] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1757.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1761.region_id = weak constant i8 0 // CK19: [[SIZE11:@.+]] = private {{.*}}constant [1 x i64] [i64 240] // CK19: [[MTYPE11:@.+]] = private {{.*}}constant [1 x i64] [i64 32] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1778.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1782.region_id = weak constant i8 0 // CK19: [[SIZE12:@.+]] = private {{.*}}constant [1 x i64] [i64 4] // CK19: [[MTYPE12:@.+]] = private {{.*}}constant [1 x i64] [i64 33] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1803.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1807.region_id = weak constant i8 0 // CK19: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i64] [i64 32] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1828.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1832.region_id = weak constant i8 0 // CK19: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i64] [i64 33] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1849.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1853.region_id = weak constant i8 0 // CK19: [[SIZE15:@.+]] = private {{.*}}constant [1 x i64] [i64 4] // CK19: [[MTYPE15:@.+]] = private {{.*}}constant [1 x i64] [i64 34] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1883.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1887.region_id = weak constant i8 0 // CK19: [[MTYPE16:@.+]] = private {{.*}}constant [2 x i64] [i64 800, i64 33] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1909.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1913.region_id = weak constant i8 0 // CK19: [[SIZE17:@.+]] = private {{.*}}constant [2 x i64] [i64 {{8|4}}, i64 240] // CK19: [[MTYPE17:@.+]] = private {{.*}}constant [2 x i64] [i64 800, i64 34] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1935.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1939.region_id = weak constant i8 0 // CK19: [[SIZE18:@.+]] = private {{.*}}constant [2 x i64] [i64 {{8|4}}, i64 240] // CK19: [[MTYPE18:@.+]] = private {{.*}}constant [2 x i64] [i64 800, i64 35] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1967.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1971.region_id = weak constant i8 0 // CK19: [[MTYPE19:@.+]] = private {{.*}}constant [2 x i64] [i64 800, i64 32] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1993.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l1997.region_id = weak constant i8 0 // CK19: [[SIZE20:@.+]] = private {{.*}}constant [2 x i64] [i64 {{8|4}}, i64 4] // CK19: [[MTYPE20:@.+]] = private {{.*}}constant [2 x i64] [i64 800, i64 33] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2025.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2029.region_id = weak constant i8 0 // CK19: [[MTYPE21:@.+]] = private {{.*}}constant [2 x i64] [i64 800, i64 35] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2051.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2055.region_id = weak constant i8 0 // CK19: [[SIZE22:@.+]] = private {{.*}}constant [2 x i64] [i64 {{8|4}}, i64 4] // CK19: [[MTYPE22:@.+]] = private {{.*}}constant [2 x i64] [i64 800, i64 35] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2070.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2074.region_id = weak constant i8 0 // CK19: [[SIZE23:@.+]] = private {{.*}}constant [1 x i64] [i64 4] // CK19: [[MTYPE23:@.+]] = private {{.*}}constant [1 x i64] [i64 39] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2092.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2096.region_id = weak constant i8 0 // CK19: [[SIZE24:@.+]] = private {{.*}}constant [1 x i64] [i64 480] // CK19: [[MTYPE24:@.+]] = private {{.*}}constant [1 x i64] [i64 35] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2113.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2117.region_id = weak constant i8 0 // CK19: [[SIZE25:@.+]] = private {{.*}}constant [1 x i64] [i64 16] // CK19: [[MTYPE25:@.+]] = private {{.*}}constant [1 x i64] [i64 35] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2134.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2138.region_id = weak constant i8 0 // CK19: [[SIZE26:@.+]] = private {{.*}}constant [1 x i64] [i64 24] // CK19: [[MTYPE26:@.+]] = private {{.*}}constant [1 x i64] [i64 35] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2155.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2159.region_id = weak constant i8 0 // CK19: [[SIZE27:@.+]] = private {{.*}}constant [1 x i64] [i64 4] // CK19: [[MTYPE27:@.+]] = private {{.*}}constant [1 x i64] [i64 35] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2200.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2204.region_id = weak constant i8 0 // CK19: [[SIZE28:@.+]] = private {{.*}}constant [3 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 16] // CK19: [[MTYPE28:@.+]] = private {{.*}}constant [3 x i64] [i64 35, i64 16, i64 19] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2245.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2249.region_id = weak constant i8 0 // CK19: [[SIZE29:@.+]] = private {{.*}}constant [3 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 4] // CK19: [[MTYPE29:@.+]] = private {{.*}}constant [3 x i64] [i64 35, i64 16, i64 19] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2301.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2305.region_id = weak constant i8 0 // CK19: [[MTYPE30:@.+]] = private {{.*}}constant [4 x i64] [i64 800, i64 800, i64 800, i64 35] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2345.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2349.region_id = weak constant i8 0 // CK19: [[SIZE31:@.+]] = private {{.*}}constant [4 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 {{8|4}}, i64 40] // CK19: [[MTYPE31:@.+]] = private {{.*}}constant [4 x i64] [i64 800, i64 800, i64 800, i64 35] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2368.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2372.region_id = weak constant i8 0 // CK19: [[SIZE32:@.+]] = private {{.*}}constant [1 x i64] [i64 13728] // CK19: [[MTYPE32:@.+]] = private {{.*}}constant [1 x i64] [i64 35] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2387.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2391.region_id = weak constant i8 0 // CK19: [[SIZE33:@.+]] = private {{.*}}constant [1 x i64] [i64 13728] // CK19: [[MTYPE33:@.+]] = private {{.*}}constant [1 x i64] [i64 35] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2406.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2410.region_id = weak constant i8 0 // CK19: [[SIZE34:@.+]] = private {{.*}}constant [1 x i64] [i64 13728] // CK19: [[MTYPE34:@.+]] = private {{.*}}constant [1 x i64] [i64 35] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2431.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2435.region_id = weak constant i8 0 // CK19: [[MTYPE35:@.+]] = private {{.*}}constant [1 x i64] [i64 35] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2452.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2456.region_id = weak constant i8 0 // CK19: [[SIZE36:@.+]] = private {{.*}}constant [1 x i64] [i64 208] // CK19: [[MTYPE36:@.+]] = private {{.*}}constant [1 x i64] [i64 35] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2492.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2496.region_id = weak constant i8 0 // CK19: [[MTYPE37:@.+]] = private {{.*}}constant [3 x i64] [i64 800, i64 800, i64 35] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2534.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2538.region_id = weak constant i8 0 // CK19: [[MTYPE38:@.+]] = private {{.*}}constant [3 x i64] [i64 800, i64 800, i64 35] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2576.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2580.region_id = weak constant i8 0 // CK19: [[MTYPE39:@.+]] = private {{.*}}constant [3 x i64] [i64 800, i64 800, i64 35] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2618.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2622.region_id = weak constant i8 0 // CK19: [[MTYPE40:@.+]] = private {{.*}}constant [3 x i64] [i64 800, i64 800, i64 35] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2653.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2657.region_id = weak constant i8 0 // CK19: [[SIZE41:@.+]] = private {{.*}}constant [3 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 208] // CK19: [[MTYPE41:@.+]] = private {{.*}}constant [3 x i64] [i64 800, i64 800, i64 35] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2698.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2702.region_id = weak constant i8 0 // CK19: [[SIZE42:@.+]] = private {{.*}}constant [3 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 104] // CK19: [[MTYPE42:@.+]] = private {{.*}}constant [3 x i64] [i64 35, i64 16, i64 19] -// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2723.region_id = weak constant i8 0 +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2727.region_id = weak constant i8 0 // CK19: [[MTYPE43:@.+]] = private {{.*}}constant [1 x i64] [i64 35] +// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l2746.region_id = weak constant i8 0 +// CK19: [[SIZE44:@.+]] = private {{.*}}constant [1 x i64] [i64 320] +// CK19: [[MTYPE44:@.+]] = private {{.*}}constant [1 x i64] [i64 34] + // CK19-LABEL: explicit_maps_single{{.*}}( void explicit_maps_single (int ii){ // Map of a scalar. @@ -2725,6 +2729,25 @@ void explicit_maps_single (int ii){ marras[1][2][3]++; } + // Region 44 + // CK19-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE44]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE44]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [100 x i32]** + // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** + // CK19-DAG: store [100 x i32]* [[VAR0:%.+]], [100 x i32]** [[CBP0]] + // CK19-DAG: store i32* [[SEC0:%[^,]+]], i32** [[CP0]] + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} 20 + + // CK19: call void [[CALL44:@.+]]([100 x i32]* {{[^,]+}}) + #pragma omp target map(from:arra[20:]) + { + arra[50]++; + } + } // CK19: define {{.+}}[[CALL00]] @@ -2771,6 +2794,7 @@ void explicit_maps_single (int ii){ // CK19: define {{.+}}[[CALL41]] // CK19: define {{.+}}[[CALL42]] // CK19: define {{.+}}[[CALL43]] +// CK19: define {{.+}}[[CALL44]] #endif ///==========================================================================/// @@ -2790,19 +2814,19 @@ void explicit_maps_single (int ii){ // SIMD-ONLY19-NOT: {{__kmpc|__tgt}} #ifdef CK20 -// CK20-LABEL: @.__omp_offloading_{{.*}}explicit_maps_references_and_function_args{{.*}}_l2832.region_id = weak constant i8 0 +// CK20-LABEL: @.__omp_offloading_{{.*}}explicit_maps_references_and_function_args{{.*}}_l2856.region_id = weak constant i8 0 // CK20: [[SIZE00:@.+]] = private {{.*}}constant [1 x i64] [i64 4] // CK20: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i64] [i64 33] -// CK20-LABEL: @.__omp_offloading_{{.*}}explicit_maps_references_and_function_args{{.*}}_l2853.region_id = weak constant i8 0 +// CK20-LABEL: @.__omp_offloading_{{.*}}explicit_maps_references_and_function_args{{.*}}_l2877.region_id = weak constant i8 0 // CK20: [[SIZE01:@.+]] = private {{.*}}constant [1 x i64] [i64 20] // CK20: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 33] -// CK20-LABEL: @.__omp_offloading_{{.*}}explicit_maps_references_and_function_args{{.*}}_l2871.region_id = weak constant i8 0 +// CK20-LABEL: @.__omp_offloading_{{.*}}explicit_maps_references_and_function_args{{.*}}_l2895.region_id = weak constant i8 0 // CK20: [[SIZE02:@.+]] = private {{.*}}constant [1 x i64] [i64 4] // CK20: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i64] [i64 34] -// CK20-LABEL: @.__omp_offloading_{{.*}}explicit_maps_references_and_function_args{{.*}}_l2892.region_id = weak constant i8 0 +// CK20-LABEL: @.__omp_offloading_{{.*}}explicit_maps_references_and_function_args{{.*}}_l2916.region_id = weak constant i8 0 // CK20: [[SIZE03:@.+]] = private {{.*}}constant [1 x i64] [i64 12] // CK20: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i64] [i64 34] @@ -5275,11 +5299,11 @@ void map_with_deep_copy() { // SIMD-ONLY18-NOT: {{__kmpc|__tgt}} #ifdef CK31 -// CK31-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l5305.region_id = weak constant i8 0 +// CK31-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l5329.region_id = weak constant i8 0 // CK31: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4] // CK31: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i64] [i64 1059] -// CK31-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l5324.region_id = weak constant i8 0 +// CK31-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l5348.region_id = weak constant i8 0 // CK31: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4] // CK31: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 1063] -- 2.7.4