r318834 - [OPENMP] Codegen for `target teams` directive.
Alexey Bataev via cfe-commits
cfe-commits at lists.llvm.org
Wed Nov 22 06:25:55 PST 2017
Author: abataev
Date: Wed Nov 22 06:25:55 2017
New Revision: 318834
URL: http://llvm.org/viewvc/llvm-project?rev=318834&view=rev
Log:
[OPENMP] Codegen for `target teams` directive.
Added codegen of the clauses for `target teams` directive.
Modified:
cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp
cfe/trunk/lib/Sema/SemaOpenMP.cpp
cfe/trunk/test/OpenMP/nvptx_target_teams_codegen.cpp
cfe/trunk/test/OpenMP/target_teams_codegen.cpp
Modified: cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp?rev=318834&r1=318833&r2=318834&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp Wed Nov 22 06:25:55 2017
@@ -3812,11 +3812,20 @@ static void emitTargetTeamsRegion(CodeGe
const OMPTargetTeamsDirective &S) {
auto *CS = S.getCapturedStmt(OMPD_teams);
Action.Enter(CGF);
- auto &&CodeGen = [CS](CodeGenFunction &CGF, PrePostActionTy &) {
- // TODO: Add support for clauses.
+ // Emit teams region as a standalone region.
+ auto &&CodeGen = [&S, CS](CodeGenFunction &CGF, PrePostActionTy &Action) {
+ CodeGenFunction::OMPPrivateScope PrivateScope(CGF);
+ (void)CGF.EmitOMPFirstprivateClause(S, PrivateScope);
+ CGF.EmitOMPPrivateClause(S, PrivateScope);
+ CGF.EmitOMPReductionClauseInit(S, PrivateScope);
+ (void)PrivateScope.Privatize();
+ Action.Enter(CGF);
CGF.EmitStmt(CS->getCapturedStmt());
+ CGF.EmitOMPReductionClauseFinal(S, /*ReductionKind=*/OMPD_teams);
};
emitCommonOMPTeamsDirective(CGF, S, OMPD_teams, CodeGen);
+ emitPostUpdateForReductionClause(
+ CGF, S, [](CodeGenFunction &) -> llvm::Value * { return nullptr; });
}
void CodeGenFunction::EmitOMPTargetTeamsDeviceFunction(
Modified: cfe/trunk/lib/Sema/SemaOpenMP.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOpenMP.cpp?rev=318834&r1=318833&r2=318834&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaOpenMP.cpp (original)
+++ cfe/trunk/lib/Sema/SemaOpenMP.cpp Wed Nov 22 06:25:55 2017
@@ -7182,6 +7182,16 @@ StmtResult Sema::ActOnOpenMPTargetTeamsD
// longjmp() and throw() must not violate the entry/exit criteria.
CS->getCapturedDecl()->setNothrow();
+ for (int ThisCaptureLevel = getOpenMPCaptureLevels(OMPD_target_teams);
+ ThisCaptureLevel > 1; --ThisCaptureLevel) {
+ CS = cast<CapturedStmt>(CS->getCapturedStmt());
+ // 1.2.2 OpenMP Language Terminology
+ // Structured block - An executable statement with a single entry at the
+ // top and a single exit at the bottom.
+ // The point of exit cannot be a branch out of the structured block.
+ // longjmp() and throw() must not violate the entry/exit criteria.
+ CS->getCapturedDecl()->setNothrow();
+ }
getCurFunction()->setHasBranchProtectedScope();
return OMPTargetTeamsDirective::Create(Context, StartLoc, EndLoc, Clauses,
Modified: cfe/trunk/test/OpenMP/nvptx_target_teams_codegen.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/nvptx_target_teams_codegen.cpp?rev=318834&r1=318833&r2=318834&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/nvptx_target_teams_codegen.cpp (original)
+++ cfe/trunk/test/OpenMP/nvptx_target_teams_codegen.cpp Wed Nov 22 06:25:55 2017
@@ -121,7 +121,9 @@ int bar(int n){
// CHECK: [[ACV:%.+]] = load i[[SZ]], i[[SZ]]* [[AC]], align
// CHECK: store i[[SZ]] [[ACV]], i[[SZ]]* [[A_ADDR_T:%.+]], align
// CHECK: [[CONV2:%.+]] = bitcast i[[SZ]]* [[A_ADDR_T]] to i8*
- // CHECK: store i8 49, i8* [[CONV2]], align
+ // CHECK: [[LD_CONV2:%.+]] = load i8, i8* [[CONV2]],
+ // CHECK: store i8 [[LD_CONV2]], i8* [[A_PRIV:%[^,]+]],
+ // CHECK: store i8 49, i8* [[A_PRIV]], align
// CHECK: br label {{%?}}[[TERMINATE:.+]]
//
// CHECK: [[TERMINATE]]
@@ -207,7 +209,9 @@ int bar(int n){
// CHECK: [[ACV:%.+]] = load i[[SZ]], i[[SZ]]* [[AC]], align
// CHECK: store i[[SZ]] [[ACV]], i[[SZ]]* [[AA_ADDR_T:%.+]], align
// CHECK: [[CONV2:%.+]] = bitcast i[[SZ]]* [[AA_ADDR_T]] to i16*
- // CHECK: store i16 1, i16* [[CONV2]], align
+ // CHECK: [[LD_CONV2:%.+]] = load i16, i16* [[CONV2]],
+ // CHECK: store i16 [[LD_CONV2]], i16* [[A_PRIV:%[^,]+]],
+ // CHECK: store i16 1, i16* [[A_PRIV]], align
// CHECK: br label {{%?}}[[TERMINATE:.+]]
//
// CHECK: [[TERMINATE]]
Modified: cfe/trunk/test/OpenMP/target_teams_codegen.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_teams_codegen.cpp?rev=318834&r1=318833&r2=318834&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/target_teams_codegen.cpp (original)
+++ cfe/trunk/test/OpenMP/target_teams_codegen.cpp Wed Nov 22 06:25:55 2017
@@ -38,7 +38,9 @@
// code, only 6 will have mapped arguments, and only 4 have all-constant map
// sizes.
-// CHECK-DAG: [[SIZET2:@.+]] = private unnamed_addr constant [1 x i{{32|64}}] [i[[SZ:32|64]] 2]
+// CHECK-DAG: [[SIZET:@.+]] = private unnamed_addr constant [3 x i[[SZ]]] [i[[SZ]] 2, i[[SZ]] 4, i[[SZ]] 4]
+// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [3 x i64] [i64 288, i64 288, i64 288]
+// CHECK-DAG: [[SIZET2:@.+]] = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 2]
// CHECK-DAG: [[MAPT2:@.+]] = private unnamed_addr constant [1 x i64] [i64 288]
// CHECK-DAG: [[SIZET3:@.+]] = private unnamed_addr constant [2 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2]
// CHECK-DAG: [[MAPT3:@.+]] = private unnamed_addr constant [2 x i64] [i64 288, i64 288]
@@ -95,14 +97,34 @@ int foo(int n) {
double cn[5][n];
TT<long long, char> d;
- // CHECK: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 0, i8** null, i8** null, i[[SZ]]* null, i64* null, i32 0, i32 0)
+ // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 3, i8** [[BP:%[^,]+]], i8** [[P:%[^,]+]], i[[SZ]]* getelementptr inbounds ([3 x i[[SZ]]], [3 x i[[SZ]]]* [[SIZET]], i32 0, i32 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[MAPT]], i32 0, i32 0), i32 {{[^,]+}}, i32 {{[^)]+}})
+ // CHECK-DAG: [[BP]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BPR:%[^,]+]], i32 0, i32 0
+ // CHECK-DAG: [[P]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PR:%[^,]+]], i32 0, i32 0
+ // CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BPR]], i32 0, i32 [[IDX0:[0-9]+]]
+ // CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PR]], i32 0, i32 [[IDX0]]
+ // CHECK-DAG: [[CBPADDR0:%.+]] = bitcast i8** [[BPADDR0]] to i[[SZ]]*
+ // CHECK-DAG: [[CPADDR0:%.+]] = bitcast i8** [[PADDR0]] to i[[SZ]]*
+ // CHECK-DAG: store i[[SZ]] %{{.+}}, i[[SZ]]* [[CBPADDR0]]
+ // CHECK-DAG: store i[[SZ]] %{{.+}}, i[[SZ]]* [[CPADDR0]]
+ // CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BPR]], i32 0, i32 [[IDX1:[0-9]+]]
+ // CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PR]], i32 0, i32 [[IDX1]]
+ // CHECK-DAG: [[CBPADDR1:%.+]] = bitcast i8** [[BPADDR1]] to i[[SZ]]*
+ // CHECK-DAG: [[CPADDR1:%.+]] = bitcast i8** [[PADDR1]] to i[[SZ]]*
+ // CHECK-DAG: store i[[SZ]] %{{.+}}, i[[SZ]]* [[CBPADDR1]]
+ // CHECK-DAG: store i[[SZ]] %{{.+}}, i[[SZ]]* [[CPADDR1]]
+ // CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BPR]], i32 0, i32 [[IDX1:[0-9]+]]
+ // CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PR]], i32 0, i32 [[IDX1]]
+ // CHECK-DAG: [[CBPADDR2:%.+]] = bitcast i8** [[BPADDR2]] to i[[SZ]]*
+ // CHECK-DAG: [[CPADDR2:%.+]] = bitcast i8** [[PADDR2]] to i[[SZ]]*
+ // CHECK-DAG: store i[[SZ]] %{{.+}}, i[[SZ]]* [[CBPADDR2]]
+ // CHECK-DAG: store i[[SZ]] %{{.+}}, i[[SZ]]* [[CPADDR2]]
// CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
// CHECK: [[FAIL]]
- // CHECK: call void [[HVT0:@.+]]()
+ // CHECK: call void [[HVT0:@.+]](i[[SZ]] {{[^,]+}}, i[[SZ]] {{[^,]+}}, i[[SZ]] {{[^)]+}})
// CHECK-NEXT: br label %[[END]]
// CHECK: [[END]]
- #pragma omp target teams
+ #pragma omp target teams num_teams(a) thread_limit(a) firstprivate(aa)
{
}
@@ -301,11 +323,12 @@ int foo(int n) {
// Check that the offloading functions are emitted and that the arguments are
// correct and loaded correctly for the target regions in foo().
-// CHECK: define internal void [[HVT0]]()
-// CHECK: call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC]], i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OMP_OUTLINED:@.+]] to void (i32*, i32*, ...)*))
+// CHECK: define internal void [[HVT0]](i[[SZ]] {{[^,]+}}, i[[SZ]] {{[^,]+}}, i[[SZ]] {{[^)]+}})
+// CHECK: call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i[[SZ]])* [[OMP_OUTLINED:@.+]] to void (i32*, i32*, ...)*), i[[SZ]] {{[^)]+}})
//
//
-// CHECK: define internal {{.*}}void [[OMP_OUTLINED]](i32* noalias %.global_tid., i32* noalias %.bound_tid.)
+// CHECK: define internal {{.*}}void [[OMP_OUTLINED]](i32* noalias %.global_tid., i32* noalias %.bound_tid., i[[SZ]] {{[^)]+}})
+// CHECK: alloca i16,
// CHECK: ret void
// CHECK-NEXT: }
More information about the cfe-commits
mailing list