r358503 - [OPENMP][NVPTX]Run combined constructs with if clause in SPMD mode.

[Alexey Bataev via cfe-commits]

Author: abataev
Date: Tue Apr 16 08:39:12 2019
New Revision: 358503

URL: http://llvm.org/viewvc/llvm-project?rev=358503&view=rev
Log:
[OPENMP][NVPTX]Run combined constructs with if clause in SPMD mode.

Combined constructs with parallel and if clauses without modifiers may
be executed in SPMD mode since if the condition is true for the target
region, it is also true for parallel region and the threads must be run
in parallel.

Modified:
    cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
    cfe/trunk/test/OpenMP/nvptx_SPMD_codegen.cpp

Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp?rev=358503&r1=358502&r2=358503&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp Tue Apr 16 08:39:12 2019
@@ -717,10 +717,12 @@ getDataSharingMode(CodeGenModule &CGM) {
 /// Check if the parallel directive has an 'if' clause with non-constant or
 /// false condition.
 static bool hasParallelIfClause(ASTContext &Ctx,
-                                const OMPExecutableDirective &D) {
+                                const OMPExecutableDirective &D,
+                                bool StandaloneParallel) {
   for (const auto *C : D.getClausesOfKind<OMPIfClause>()) {
     OpenMPDirectiveKind NameModifier = C->getNameModifier();
-    if (NameModifier != OMPD_parallel && NameModifier != OMPD_unknown)
+    if (NameModifier != OMPD_parallel &&
+        (!StandaloneParallel || NameModifier != OMPD_unknown))
       continue;
     const Expr *Cond = C->getCondition();
     bool Result;
@@ -744,7 +746,7 @@ static bool hasNestedSPMDDirective(ASTCo
     switch (D.getDirectiveKind()) {
     case OMPD_target:
       if (isOpenMPParallelDirective(DKind) &&
-          !hasParallelIfClause(Ctx, *NestedDir))
+          !hasParallelIfClause(Ctx, *NestedDir, /*StandaloneParallel=*/true))
         return true;
       if (DKind == OMPD_teams) {
         Body = NestedDir->getInnermostCapturedStmt()->IgnoreContainers(
@@ -756,14 +758,14 @@ static bool hasNestedSPMDDirective(ASTCo
                 dyn_cast_or_null<OMPExecutableDirective>(ChildStmt)) {
           DKind = NND->getDirectiveKind();
           if (isOpenMPParallelDirective(DKind) &&
-              !hasParallelIfClause(Ctx, *NND))
+              !hasParallelIfClause(Ctx, *NND, /*StandaloneParallel=*/true))
             return true;
         }
       }
       return false;
     case OMPD_target_teams:
       return isOpenMPParallelDirective(DKind) &&
-             !hasParallelIfClause(Ctx, *NestedDir);
+             !hasParallelIfClause(Ctx, *NestedDir, /*StandaloneParallel=*/true);
     case OMPD_target_simd:
     case OMPD_target_parallel:
     case OMPD_target_parallel_for:
@@ -837,7 +839,7 @@ static bool supportsSPMDExecutionMode(AS
   case OMPD_target_parallel_for_simd:
   case OMPD_target_teams_distribute_parallel_for:
   case OMPD_target_teams_distribute_parallel_for_simd:
-    return !hasParallelIfClause(Ctx, D);
+    return !hasParallelIfClause(Ctx, D, /*StandaloneParallel=*/false);
   case OMPD_target_simd:
   case OMPD_target_teams_distribute:
   case OMPD_target_teams_distribute_simd:

Modified: cfe/trunk/test/OpenMP/nvptx_SPMD_codegen.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/nvptx_SPMD_codegen.cpp?rev=358503&r1=358502&r2=358503&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/nvptx_SPMD_codegen.cpp (original)
+++ cfe/trunk/test/OpenMP/nvptx_SPMD_codegen.cpp Tue Apr 16 08:39:12 2019
@@ -8,6 +8,8 @@
 #ifndef HEADER
 #define HEADER
 
+int a;
+
 // CHECK-NOT: @__omp_offloading_{{.+}}_exec_mode = weak constant i8 1
 // CHECK-DAG: [[DISTR_LIGHT:@.+]] = private unnamed_addr constant %struct.ident_t { i32 0, i32 2050, i32 3, i32 0, i8* getelementptr inbounds
 // CHECK-DAG: [[FOR_LIGHT:@.+]] = private unnamed_addr constant %struct.ident_t { i32 0, i32 514, i32 3, i32 0, i8* getelementptr inbounds
@@ -43,7 +45,7 @@ void foo() {
 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
 // CHECK-DAG: [[DISTR_FULL]]
 // CHECK-DAG: [[FULL]]
-#pragma omp target teams distribute parallel for simd
+#pragma omp target teams distribute parallel for simd if(a)
   for (int i = 0; i < 10; ++i)
     ;
 #pragma omp target teams distribute parallel for simd schedule(static)
@@ -301,7 +303,7 @@ int a;
 // CHECK-DAG: [[FULL]]
 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
 // CHECK-DAG: [[FULL]]
-#pragma omp target parallel for
+#pragma omp target parallel for if(a)
   for (int i = 0; i < 10; ++i)
     ;
 #pragma omp target parallel for schedule(static)
@@ -346,7 +348,7 @@ int a;
 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
 // CHECK-DAG: [[FULL]]
 // CHECK-DAG: [[BAR_FULL]]
-#pragma omp target parallel
+#pragma omp target parallel if(a)
 #pragma omp for simd
   for (int i = 0; i < 10; ++i)
     ;