From 842e99f500dec2f96beb4a9b69cf5ce156578eef Mon Sep 17 00:00:00 2001 From: Alexey Bataev Date: Wed, 19 Dec 2018 18:16:37 +0000 Subject: [PATCH] [OPENMP]Mark the loop as started when initialized. Need to mark the loop as started when the initialization statement is found. It is required to prevent possible incorrect loop iteraton variable detection during template instantiation and fix the compiler crash during the codegen. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@349657 91177308-0d34-0410-b5e6-96231b3b80d8 --- lib/Sema/SemaOpenMP.cpp | 1 + ...target_teams_distribute_parallel_for_codegen.cpp | 13 ++++++++++++- 2 files changed, 13 insertions(+), 1 deletion(-) diff --git a/lib/Sema/SemaOpenMP.cpp b/lib/Sema/SemaOpenMP.cpp index c207ba9bd3..4fe8d3dd59 100644 --- a/lib/Sema/SemaOpenMP.cpp +++ b/lib/Sema/SemaOpenMP.cpp @@ -4649,6 +4649,7 @@ void Sema::ActOnOpenMPLoopInitialization(SourceLocation ForLoc, Stmt *Init) { unsigned AssociatedLoops = DSAStack->getAssociatedLoops(); if (AssociatedLoops > 0 && isOpenMPLoopDirective(DSAStack->getCurrentDirective())) { + DSAStack->loopStart(); OpenMPIterationSpaceChecker ISC(*this, ForLoc); if (!ISC.checkAndSetInit(Init, /*EmitDiags=*/false)) { if (ValueDecl *D = ISC.getLoopDecl()) { diff --git a/test/OpenMP/nvptx_target_teams_distribute_parallel_for_codegen.cpp b/test/OpenMP/nvptx_target_teams_distribute_parallel_for_codegen.cpp index 5f003dd6cc..2ecb9ca0e2 100644 --- a/test/OpenMP/nvptx_target_teams_distribute_parallel_for_codegen.cpp +++ b/test/OpenMP/nvptx_target_teams_distribute_parallel_for_codegen.cpp @@ -53,6 +53,13 @@ tx ftemplate(int n) { } } +#pragma omp target teams distribute parallel for collapse(2) + for(int i = 0; i < n; i++) { + for(int j = 0; j < n; j++) { + c[i][j] = i + j; + } + } + #pragma omp target teams distribute parallel for map(a, v[:N]) for(int i = 0; i < n; i++) a[i] = v[i]; @@ -212,7 +219,11 @@ int bar(int n){ // CHECK: call void @__kmpc_for_static_fini( // CHECK: ret void -// CHECK: define weak void @__omp_offloading_{{.*}}_l56(i[[SZ:64|32]] %{{[^,]+}}, [1000 x i32]* dereferenceable{{.*}}, i32* %{{[^)]+}}) +// CHECK: define weak void @__omp_offloading_{{.*}}_l56(i[[SZ:64|32]] %{{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{.*}}) +// CHECK: call void [[OUTLINED:@__omp_outlined.*]](i32* %{{.+}}, i32* %{{.+}}, i[[SZ]] %{{.*}}, i[[SZ]] %{{.*}}, i[[SZ]] %{{.*}}, [10 x [10 x i32]]* %{{.*}}) +// CHECK: define internal void [[OUTLINED]](i32* noalias %{{.*}}, i32* noalias %{{.*}} i[[SZ]] %{{.+}}, i[[SZ]] %{{.+}}, i[[SZ]] %{{.+}}, [10 x [10 x i32]]* dereferenceable{{.*}}) + +// CHECK: define weak void @__omp_offloading_{{.*}}_l63(i[[SZ:64|32]] %{{[^,]+}}, [1000 x i32]* dereferenceable{{.*}}, i32* %{{[^)]+}}) // CHECK: call void [[OUTLINED:@__omp_outlined.*]](i32* %{{.+}}, i32* %{{.+}}, i[[SZ]] %{{.*}}, i[[SZ]] %{{.*}}, i[[SZ]] %{{.*}}, [1000 x i32]* %{{.*}}, i32* %{{.*}}) // CHECK: define internal void [[OUTLINED]](i32* noalias %{{.*}}, i32* noalias %{{.*}} i[[SZ]] %{{.+}}, i[[SZ]] %{{.+}}, i[[SZ]] %{{.+}}, [1000 x i32]* dereferenceable{{.*}}, i32* %{{.*}}) -- 2.40.0