diff --git a/clang/test/OpenMP/requires_ast_print.cpp b/clang/test/OpenMP/requires_ast_print.cpp --- a/clang/test/OpenMP/requires_ast_print.cpp +++ b/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 diff --git a/clang/test/OpenMP/requires_messages.cpp b/clang/test/OpenMP/requires_messages.cpp --- a/clang/test/OpenMP/requires_messages.cpp +++ b/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 @@ #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 @@ #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}} diff --git a/clang/test/OpenMP/requires_target_messages.cpp b/clang/test/OpenMP/requires_target_messages.cpp --- a/clang/test/OpenMP/requires_target_messages.cpp +++ b/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 @@ #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}} diff --git a/clang/test/OpenMP/target_ast_print.cpp b/clang/test/OpenMP/target_ast_print.cpp --- a/clang/test/OpenMP/target_ast_print.cpp +++ b/clang/test/OpenMP/target_ast_print.cpp @@ -342,7 +342,18 @@ // 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 @@ 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 @@ // 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 @@ // 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 @@ // 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) diff --git a/clang/test/OpenMP/target_device_codegen.cpp b/clang/test/OpenMP/target_device_codegen.cpp --- a/clang/test/OpenMP/target_device_codegen.cpp +++ b/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 @@ // 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 diff --git a/llvm/include/llvm/Frontend/OpenMP/OMP.td b/llvm/include/llvm/Frontend/OpenMP/OMP.td --- a/llvm/include/llvm/Frontend/OpenMP/OMP.td +++ b/llvm/include/llvm/Frontend/OpenMP/OMP.td @@ -590,7 +590,15 @@ let allowedClauses = [ VersionedClause, VersionedClause, - VersionedClause, + // 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, VersionedClause, VersionedClause ];