[OpenMP][Clang] Move partial support of reverse offload to a future version

OpenMP Spec 5.2 requires unimplemented requires clauses to produce
compile time error termination. Moving current partial support of
reverse_offload to a distant future version 9.9 so that existing
code can be tested and maintained until a complete implementation
is available.

Reviewed By: ABataev

Differential Revision: https://reviews.llvm.org/D119256

clang/test/OpenMP/requires_ast_print.cpp

@@ -5,6 +5,14 @@
 // RUN: %clang_cc1 -verify -fopenmp-simd -ast-print %s | FileCheck %s
 // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=99 -DOMP99 -ast-print %s | FileCheck --check-prefixes=CHECK,REV %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=99 -DOMP99 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=99 -DOMP99 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck --check-prefixes=CHECK,REV %s
+
+// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=99 -DOMP99 -ast-print %s | FileCheck --check-prefixes=CHECK,REV %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=99 -DOMP99 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=99 -DOMP99 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck --check-prefixes=CHECK,REV %s
 // expected-no-diagnostics
 
 #ifndef HEADER
@@ -16,8 +24,10 @@
 #pragma omp requires unified_shared_memory
 // CHECK:#pragma omp requires unified_shared_memory
 
+#ifdef OMP99
 #pragma omp requires reverse_offload
-// CHECK:#pragma omp requires reverse_offload
+// REV:#pragma omp requires reverse_offload
+#endif
 
 #pragma omp requires dynamic_allocators
 // CHECK:#pragma omp requires dynamic_allocators

clang/test/OpenMP/requires_messages.cpp

@@ -1,9 +1,10 @@
 // RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100  %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=99 -DOMP99 -verify=expected,rev -ferror-limit 100  %s -Wuninitialized
 
 int a;
-#pragma omp requires unified_address allocate(a) // expected-note {{unified_address clause previously used here}} expected-note {{unified_address clause previously used here}} expected-note {{unified_address clause previously used here}} expected-note {{unified_address clause previously used here}} expected-note {{unified_address clause previously used here}} expected-note{{unified_address clause previously used here}} expected-error {{unexpected OpenMP clause 'allocate' in directive '#pragma omp requires'}}
+#pragma omp requires unified_address allocate(a) // rev-note {{unified_address clause previously used here}} expected-note {{unified_address clause previously used here}} expected-note {{unified_address clause previously used here}} expected-note {{unified_address clause previously used here}} expected-note {{unified_address clause previously used here}} expected-note{{unified_address clause previously used here}} expected-error {{unexpected OpenMP clause 'allocate' in directive '#pragma omp requires'}}
 
-#pragma omp requires unified_shared_memory // expected-note {{unified_shared_memory clause previously used here}} expected-note{{unified_shared_memory clause previously used here}}
+#pragma omp requires unified_shared_memory // rev-note {{unified_shared_memory clause previously used here}} expected-note{{unified_shared_memory clause previously used here}}
 
 #pragma omp requires unified_shared_memory, unified_shared_memory // expected-error {{Only one unified_shared_memory clause can appear on a requires directive in a single translation unit}} expected-error {{directive '#pragma omp requires' cannot contain more than one 'unified_shared_memory' clause}}
 
@@ -11,15 +12,17 @@ int a;
 
 #pragma omp requires unified_address, unified_address // expected-error {{Only one unified_address clause can appear on a requires directive in a single translation unit}} expected-error {{directive '#pragma omp requires' cannot contain more than one 'unified_address' clause}}
 
-#pragma omp requires reverse_offload // expected-note {{reverse_offload clause previously used here}} expected-note {{reverse_offload clause previously used here}}
+#ifdef OMP99
+#pragma omp requires reverse_offload // rev-note {{reverse_offload clause previously used here}} rev-note {{reverse_offload clause previously used here}}
 
-#pragma omp requires reverse_offload, reverse_offload // expected-error {{Only one reverse_offload clause can appear on a requires directive in a single translation unit}} expected-error {{directive '#pragma omp requires' cannot contain more than one 'reverse_offload' clause}}
+#pragma omp requires reverse_offload, reverse_offload // rev-error {{Only one reverse_offload clause can appear on a requires directive in a single translation unit}} rev-error {{directive '#pragma omp requires' cannot contain more than one 'reverse_offload' clause}}
+#endif
 
-#pragma omp requires dynamic_allocators // expected-note {{dynamic_allocators clause previously used here}} expected-note {{dynamic_allocators clause previously used here}}
+#pragma omp requires dynamic_allocators // rev-note {{dynamic_allocators clause previously used here}} expected-note {{dynamic_allocators clause previously used here}}
 
 #pragma omp requires dynamic_allocators, dynamic_allocators // expected-error {{Only one dynamic_allocators clause can appear on a requires directive in a single translation unit}} expected-error {{directive '#pragma omp requires' cannot contain more than one 'dynamic_allocators' clause}}
 
-#pragma omp requires atomic_default_mem_order(seq_cst) // expected-note {{atomic_default_mem_order clause previously used here}} expected-note {{atomic_default_mem_order clause previously used here}} expected-note {{atomic_default_mem_order clause previously used here}} expected-note {{atomic_default_mem_order clause previously used here}} expected-note {{atomic_default_mem_order clause previously used here}}
+#pragma omp requires atomic_default_mem_order(seq_cst) // rev-note {{atomic_default_mem_order clause previously used here}} expected-note {{atomic_default_mem_order clause previously used here}} expected-note {{atomic_default_mem_order clause previously used here}} expected-note {{atomic_default_mem_order clause previously used here}} expected-note {{atomic_default_mem_order clause previously used here}}
 
 #pragma omp requires atomic_default_mem_order(acq_rel) // expected-error {{Only one atomic_default_mem_order clause can appear on a requires directive in a single translation unit}}
 
@@ -47,7 +50,9 @@ int a;
 
 #pragma omp requires invalid_clause unified_address // expected-warning {{extra tokens at the end of '#pragma omp requires' are ignored}} expected-error {{expected at least one clause on '#pragma omp requires' directive}}
 
-#pragma omp requires unified_shared_memory, unified_address, reverse_offload, dynamic_allocators, atomic_default_mem_order(seq_cst) // expected-error {{Only one unified_shared_memory clause can appear on a requires directive in a single translation unit}} expected-error{{Only one unified_address clause can appear on a requires directive in a single translation unit}} expected-error{{Only one reverse_offload clause can appear on a requires directive in a single translation unit}} expected-error{{Only one dynamic_allocators clause can appear on a requires directive in a single translation unit}} expected-error {{Only one atomic_default_mem_order clause can appear on a requires directive in a single translation unit}}
+#ifdef OMP99
+#pragma omp requires unified_shared_memory, unified_address, reverse_offload, dynamic_allocators, atomic_default_mem_order(seq_cst) // rev-error {{Only one unified_shared_memory clause can appear on a requires directive in a single translation unit}} rev-error{{Only one unified_address clause can appear on a requires directive in a single translation unit}} rev-error{{Only one reverse_offload clause can appear on a requires directive in a single translation unit}} rev-error{{Only one dynamic_allocators clause can appear on a requires directive in a single translation unit}} rev-error {{Only one atomic_default_mem_order clause can appear on a requires directive in a single translation unit}}
+#endif
 
 namespace A {
   #pragma omp requires unified_address // expected-error {{Only one unified_address clause can appear on a requires directive in a single translation unit}}

clang/test/OpenMP/requires_target_messages.cpp

@@ -1,8 +1,9 @@
 // RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100  %s
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=99 -DOMP99 -verify=expected,rev -ferror-limit 100  %s
 
 void foo2() {
   int a;
-  #pragma omp target // expected-note 4 {{'target' previously encountered here}}
+  #pragma omp target // expected-note 3 {{'target' previously encountered here}} rev-note {{'target' previously encountered here}}
   {
     a = a + 1;
   }
@@ -11,5 +12,7 @@ void foo2() {
 #pragma omp requires atomic_default_mem_order(seq_cst)
 #pragma omp requires unified_address //expected-error {{'target' region encountered before requires directive with 'unified_address' clause}}
 #pragma omp requires unified_shared_memory //expected-error {{'target' region encountered before requires directive with 'unified_shared_memory' clause}}
+#ifdef OMP99
 #pragma omp requires reverse_offload //expected-error {{'target' region encountered before requires directive with 'reverse_offload' clause}}
+#endif
 #pragma omp requires dynamic_allocators //expected-error {{'target' region encountered before requires directive with 'dynamic_allocators' clause}}

clang/test/OpenMP/target_ast_print.cpp

@@ -342,7 +342,18 @@ int main (int argc, char **argv) {
 // RUN: %clang_cc1 -DOMP5 -verify -fopenmp-simd -fopenmp-version=50 -ast-print %s | FileCheck %s --check-prefix OMP5
 // RUN: %clang_cc1 -DOMP5 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
 // RUN: %clang_cc1 -DOMP5 -fopenmp-simd -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s --check-prefix OMP5
+
+// RUN: %clang_cc1 -DOMP5 -verify -fopenmp -fopenmp-version=99 -DOMP99 -ast-print %s | FileCheck %s --check-prefixes=OMP5,REV
+// RUN: %clang_cc1 -DOMP5 -fopenmp -fopenmp-version=99 -DOMP99 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -DOMP5 -fopenmp -fopenmp-version=99 -DOMP99 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s --check-prefixes=OMP5,REV
+
+// RUN: %clang_cc1 -DOMP5 -verify -fopenmp-simd -fopenmp-version=99 -DOMP99 -ast-print %s | FileCheck %s --check-prefixes=OMP5,REV
+// RUN: %clang_cc1 -DOMP5 -fopenmp-simd -fopenmp-version=99 -DOMP99 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -DOMP5 -fopenmp-simd -fopenmp-version=99 -DOMP99 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s --check-prefixes=OMP5,REV
+
+#ifdef OMP99
 #pragma omp requires reverse_offload
+#endif
 typedef void **omp_allocator_handle_t;
 extern const omp_allocator_handle_t omp_null_allocator;
 extern const omp_allocator_handle_t omp_default_mem_alloc;
@@ -370,8 +381,10 @@ T tmain(T argc, T *argv) {
   foo();
 #pragma omp target if (target:argc > 0) device(device_num: C)
   foo();
+#ifdef OMP99
 #pragma omp target if (C) device(ancestor: argc)
   foo();
+#endif
 #pragma omp target map(i)
   foo();
 #pragma omp target map(a[0:10], i)
@@ -475,8 +488,8 @@ T tmain(T argc, T *argv) {
 // OMP5-NEXT: foo();
 // OMP5-NEXT: #pragma omp target if(target: argc > 0) device(device_num: C)
 // OMP5-NEXT: foo()
-// OMP5-NEXT: #pragma omp target if(C) device(ancestor: argc)
-// OMP5-NEXT: foo()
+// REV: #pragma omp target if(C) device(ancestor: argc)
+// REV-NEXT: foo()
 // OMP5-NEXT: #pragma omp target map(tofrom: i)
 // OMP5-NEXT: foo()
 // OMP5-NEXT: #pragma omp target map(tofrom: a[0:10],i)
@@ -571,8 +584,8 @@ T tmain(T argc, T *argv) {
 // OMP5-NEXT: foo();
 // OMP5-NEXT: #pragma omp target if(target: argc > 0)
 // OMP5-NEXT: foo()
-// OMP5-NEXT: #pragma omp target if(5)
-// OMP5-NEXT: foo()
+// REV: #pragma omp target if(5)
+// REV-NEXT: foo()
 // OMP5-NEXT: #pragma omp target map(tofrom: i)
 // OMP5-NEXT: foo()
 // OMP5-NEXT: #pragma omp target map(tofrom: a[0:10],i)
@@ -667,8 +680,8 @@ T tmain(T argc, T *argv) {
 // OMP5-NEXT: foo();
 // OMP5-NEXT: #pragma omp target if(target: argc > 0) device(device_num: 1)
 // OMP5-NEXT: foo()
-// OMP5-NEXT: #pragma omp target if(1) device(ancestor: argc)
-// OMP5-NEXT: foo()
+// REV: #pragma omp target if(1) device(ancestor: argc)
+// REV-NEXT: foo()
 // OMP5-NEXT: #pragma omp target map(tofrom: i)
 // OMP5-NEXT: foo()
 // OMP5-NEXT: #pragma omp target map(tofrom: a[0:10],i)

clang/test/OpenMP/target_device_codegen.cpp

@@ -6,12 +6,25 @@
 // RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=99 -DOMP99 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefixes=CHECK,REV
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=99 -DOMP99 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=99 -DOMP99 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,REV
+
+// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=99 -DOMP99 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefixes=SIMD-ONLY0
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=99 -DOMP99 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=99 -DOMP99 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefixes=SIMD-ONLY0
+
 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
 
 // expected-no-diagnostics
 #ifndef HEADER
 #define HEADER
+
+#ifdef OMP99
 #pragma omp requires reverse_offload
+#endif
+
 void foo(int n) {
 
   // CHECK:       [[N:%.+]] = load i32, i32* [[N_ADDR:%.+]],
@@ -40,11 +53,14 @@ void foo(int n) {
   // CHECK:       [[END]]
   #pragma omp target device(device_num: n)
   ;
-  // CHECK-NOT:   call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}},
-  // CHECK:       call void @__omp_offloading_{{.+}}_l46()
-  // CHECK-NOT:   call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}},
+
+#ifdef OMP99
+  // REV-NOT:   call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}},
+  // REV:       call void @__omp_offloading_{{.+}}_l61()
+  // REV-NOT:   call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}},
   #pragma omp target device(ancestor: n)
   ;
+#endif
 }
 
 #endif

llvm/include/llvm/Frontend/OpenMP/OMP.td

@@ -590,7 +590,15 @@ def OMP_Requires : Directive<"requires"> {
   let allowedClauses = [
     VersionedClause<OMPC_UnifiedAddress>,
     VersionedClause<OMPC_UnifiedSharedMemory>,
-    VersionedClause<OMPC_ReverseOffload>,
+    // OpenMP 5.2 Spec: If an implementation is not supporting a requirement
+    // (reverse offload in this case) then it should give compile-time error
+    // termination.
+    // Seeting supported version for reverse_offload to a distant future version
+    // 9.9 so that its partial support can be tested in the meantime.
+    //
+    // TODO: Correct this supprted version number whenever complete
+    // implementation of reverse_offload is available.
+    VersionedClause<OMPC_ReverseOffload, 99>,
     VersionedClause<OMPC_DynamicAllocators>,
     VersionedClause<OMPC_AtomicDefaultMemOrder>
   ];