r373374 - [OPENMP]Fix PR43330: OpenMP target: Mapping of partial arrays fails.
Alexey Bataev via cfe-commits
cfe-commits at lists.llvm.org
Tue Oct 1 11:18:04 PDT 2019
Author: abataev
Date: Tue Oct 1 11:18:03 2019
New Revision: 373374
URL: http://llvm.org/viewvc/llvm-project?rev=373374&view=rev
Log:
[OPENMP]Fix PR43330: OpenMP target: Mapping of partial arrays fails.
Fixed calculation the size of the array sections.
Modified:
cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp
cfe/trunk/test/OpenMP/target_map_codegen.cpp
Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp?rev=373374&r1=373373&r2=373374&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp Tue Oct 1 11:18:03 2019
@@ -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);
}
Modified: cfe/trunk/test/OpenMP/target_map_codegen.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_map_codegen.cpp?rev=373374&r1=373373&r2=373374&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/target_map_codegen.cpp (original)
+++ cfe/trunk/test/OpenMP/target_map_codegen.cpp Tue Oct 1 11:18:03 2019
@@ -1323,172 +1323,176 @@ void implicit_maps_template_type_capture
// 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]
More information about the cfe-commits
mailing list