[OpenMP][Sema] Emit a Sema error when -fopenmp-target-xteam-scan is missing

Emit an error when the 'scan' directive is nested inside the
'target teams distribute parallel for' directive while the flag
'-fopenmp-target-xteam-scan' has not been turned on to enable the
Xteam Scan Codegen. Added a test for the same.

Change-Id: Ib8a8eab177c7dbb2b5fa3f964465e5b5d74db662

Author: animeshk-amd

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -11657,6 +11657,8 @@ def note_omp_previous_inscan_reduction : Note<
   "'reduction' clause with 'inscan' modifier is used here">;
 def err_omp_multivar_xteam_scan_unsupported : Error<
   "multiple list items are not yet supported with the 'inclusive' or the 'exclusive' clauses that appear with the 'scan' directive">;  
+def err_omp_xteam_scan_prohibited : Error<
+  "'scan' directive is not supported inside target regions. Use flag '-fopenmp-target-xteam-scan' to enable it">;  
 def err_omp_expected_predefined_allocator : Error<
   "expected one of the predefined allocators for the variables with the static "
   "storage: 'omp_default_mem_alloc', 'omp_large_cap_mem_alloc', "

clang/lib/Sema/SemaOpenMP.cpp

@@ -5400,6 +5400,13 @@ static bool checkNestingOfRegions(Sema &SemaRef, const DSAStackTy *Stack,
     Recommend = ShouldBeInTargetRegion;
   } else if (CurrentRegion == OMPD_scan) {
     if (SemaRef.LangOpts.OpenMP >= 50) {
+      // Make sure that the '-fopenmp-target-xteam-scan' flag is passed to
+      // enable the Xteam-Scan Codegen, if the 'scan' directive is found to be
+      // nested inside the 'target teams distribute parallel for' directive
+      if (ParentRegion == OMPD_target_teams_distribute_parallel_for &&
+          !SemaRef.getLangOpts().OpenMPTargetXteamScan)
+        SemaRef.Diag(StartLoc, diag::err_omp_xteam_scan_prohibited)
+            << getOpenMPDirectiveName(CurrentRegion) << Recommend;
       // OpenMP spec 5.0 and 5.1 require scan to be directly enclosed by for,
       // simd, or for simd. This has to take into account combined directives.
       // In 5.2 this seems to be implied by the fact that the specified

clang/test/OpenMP/Inputs/nesting_of_regions.cpp

@@ -9271,7 +9271,7 @@ void foo() {
   }
 #pragma omp target teams distribute parallel for
   for (int i = 0; i < 10; ++i) {
-#pragma omp scan // omp45-error {{region cannot be closely nested inside 'target teams distribute parallel for' region}} omp50-error {{exactly one of 'inclusive' or 'exclusive' clauses is expected}} omp51-error {{exactly one of 'inclusive' or 'exclusive' clauses is expected}} 
+#pragma omp scan // omp45-error {{region cannot be closely nested inside 'target teams distribute parallel for' region}} omp50-error {{exactly one of 'inclusive' or 'exclusive' clauses is expected}} omp51-error {{exactly one of 'inclusive' or 'exclusive' clauses is expected}} omp50-error {{'scan' directive is not supported inside target regions. Use flag '-fopenmp-target-xteam-scan' to enable it}} omp51-error {{'scan' directive is not supported inside target regions. Use flag '-fopenmp-target-xteam-scan' to enable it}}
     bar();
   }
 #pragma omp target teams distribute parallel for
@@ -18547,7 +18547,7 @@ void foo() {
   }
 #pragma omp target teams distribute parallel for
   for (int i = 0; i < 10; ++i) {
-#pragma omp scan // omp45-error {{region cannot be closely nested inside 'target teams distribute parallel for' region}} omp50-error {{exactly one of 'inclusive' or 'exclusive' clauses is expected}} omp51-error {{exactly one of 'inclusive' or 'exclusive' clauses is expected}} 
+#pragma omp scan // omp45-error {{region cannot be closely nested inside 'target teams distribute parallel for' region}} omp50-error {{exactly one of 'inclusive' or 'exclusive' clauses is expected}} omp51-error {{exactly one of 'inclusive' or 'exclusive' clauses is expected}} omp50-error {{'scan' directive is not supported inside target regions. Use flag '-fopenmp-target-xteam-scan' to enable it}} omp51-error {{'scan' directive is not supported inside target regions. Use flag '-fopenmp-target-xteam-scan' to enable it}}
     bar();
   }
 #pragma omp target teams distribute parallel for

clang/test/OpenMP/xteam_scan_messages.cpp

@@ -1,6 +1,10 @@
-// RUN: %clang_cc1 -verify -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-target-xteam-scan %s -Wuninitialized
 
-// RUN: %clang_cc1 -verify -fopenmp-simd %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-target-xteam-scan %s -Wuninitialized
+
+// RUN: %clang_cc1 -verify=missing-flag,expected -fopenmp %s -Wuninitialized
+
+// RUN: %clang_cc1 -verify=missing-flag,expected -fopenmp-simd %s -Wuninitialized
 
 #define NUM_TEAMS 256
 #define NUM_THREADS 256
@@ -16,7 +20,9 @@ int main() {
   for(int i = 0; i < N; i++) {
     sum1 += in[i];  
     sum2 += 2*in[i];  
-    #pragma omp scan inclusive(sum1,sum2) // expected-error {{multiple list items are not yet supported with the 'inclusive' or the 'exclusive' clauses that appear with the 'scan' directive}}
+    // missing-flag-error@+2 {{'scan' directive is not supported inside target regions. Use flag '-fopenmp-target-xteam-scan' to enable it}}
+    // expected-error@+1 {{multiple list items are not yet supported with the 'inclusive' or the 'exclusive' clauses that appear with the 'scan' directive}}
+    #pragma omp scan inclusive(sum1,sum2)
     out1[i] = sum1; 
     out2[i] = sum2; 
   }
@@ -25,10 +31,13 @@ int main() {
   for(int i = 0; i < N; i++) {
     out1[i] = sum1; 
     out2[i] = sum2; 
-    #pragma omp scan exclusive(sum1,sum2) // expected-error {{multiple list items are not yet supported with the 'inclusive' or the 'exclusive' clauses that appear with the 'scan' directive}}
+    // missing-flag-error@+2 {{'scan' directive is not supported inside target regions. Use flag '-fopenmp-target-xteam-scan' to enable it}}
+    // expected-error@+1 {{multiple list items are not yet supported with the 'inclusive' or the 'exclusive' clauses that appear with the 'scan' directive}}
+    #pragma omp scan exclusive(sum1,sum2) 
     sum1 += in[i];  
     sum2 += 2*in[i];  
   }
 
+
   return 0;
 }