6 files changed, 116 insertions, 9 deletions
diff --git a/clang/test/OpenMP/declare_target_ast_print.cpp b/clang/test/OpenMP/declare_target_ast_print.cpp
index 613926d5f61..510162f808d 100644
--- a/clang/test/OpenMP/declare_target_ast_print.cpp
+++ b/clang/test/OpenMP/declare_target_ast_print.cpp
@@ -2,6 +2,10 @@
 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -I %S/Inputs -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -std=c++11 -include-pch %t -fsyntax-only -I %S/Inputs -verify %s -ast-print | FileCheck %s
 
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -I %S/Inputs -ast-print %s | FileCheck %s --check-prefix=CHECK --check-prefix=OMP50
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -I %S/Inputs -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -I %S/Inputs -verify %s -ast-print | FileCheck %s --check-prefix=CHECK --check-prefix=OMP50
+
 // RUN: %clang_cc1 -verify -fopenmp-simd -I %S/Inputs -ast-print %s | FileCheck %s
 // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -I %S/Inputs -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -std=c++11 -include-pch %t -fsyntax-only -I %S/Inputs -verify %s -ast-print | FileCheck %s
@@ -10,7 +14,29 @@
 #ifndef HEADER
 #define HEADER
 
+#if _OPENMP == 201811
+void bar();
+#pragma omp declare target to(bar) device_type(any)
+// OMP50: #pragma omp declare target{{$}}
+// OMP50: void bar();
+// OMP50: #pragma omp end declare target{{$}}
+void baz();
+#pragma omp declare target to(baz) device_type(nohost)
+// OMP50: #pragma omp declare target device_type(nohost){{$}}
+// OMP50: void baz();
+// OMP50: #pragma omp end declare target{{$}}
+void bazz();
+#pragma omp declare target to(bazz) device_type(host)
+// OMP50: #pragma omp declare target device_type(host){{$}}
+// OMP50: void bazz();
+// OMP50: #pragma omp end declare target{{$}}
+#endif // _OPENMP
+
 int out_decl_target = 0;
+#if _OPENMP == 201811
+#pragma omp declare target (out_decl_target)
+#endif // _OPENMP
+
 // CHECK: #pragma omp declare target{{$}}
 // CHECK: int out_decl_target = 0;
 // CHECK: #pragma omp end declare target{{$}}
diff --git a/clang/test/OpenMP/declare_target_codegen.cpp b/clang/test/OpenMP/declare_target_codegen.cpp
index c21dbffc5b6..e24e58e10aa 100644
--- a/clang/test/OpenMP/declare_target_codegen.cpp
+++ b/clang/test/OpenMP/declare_target_codegen.cpp
@@ -3,6 +3,14 @@
 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t
 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -o - | FileCheck %s
 
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - -fopenmp-version=50 -DOMP5 | FileCheck %s --check-prefix HOST5
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc -fopenmp-version=50 -DOMP5
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -fopenmp-version=50 -DOMP5 | FileCheck %s --check-prefix DEV5
+
+// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - -fopenmp-version=50 -DOMP5 | FileCheck %s --check-prefix SIMD-ONLY
+// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc -fopenmp-version=50 -DOMP5
+// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -fopenmp-version=50 -DOMP5 | FileCheck %s --check-prefix SIMD-ONLY
+
 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o -| FileCheck %s --check-prefix SIMD-ONLY
 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t
@@ -32,13 +40,13 @@
 // CHECK-DAG: [[STAT:@.+stat]] = internal global %struct.S zeroinitializer,
 // CHECK-DAG: [[STAT_REF:@.+]] = internal constant %struct.S* [[STAT]]
 // CHECK-DAG: @out_decl_target ={{ dso_local | }}global i32 0,
-// CHECK-DAG: @llvm.used = appending global [2 x i8*] [i8* bitcast (void ()* @__omp_offloading__{{.+}}_globals_l[[@LINE+80]]_ctor to i8*), i8* bitcast (void ()* @__omp_offloading__{{.+}}_stat_l[[@LINE+81]]_ctor to i8*)],
+// CHECK-DAG: @llvm.used = appending global [2 x i8*] [i8* bitcast (void ()* @__omp_offloading__{{.+}}_globals_l[[@LINE+84]]_ctor to i8*), i8* bitcast (void ()* @__omp_offloading__{{.+}}_stat_l[[@LINE+85]]_ctor to i8*)],
 // CHECK-DAG: @llvm.compiler.used = appending global [1 x i8*] [i8* bitcast (%struct.S** [[STAT_REF]] to i8*)],
 
 // CHECK-DAG: define {{.*}}i32 @{{.*}}{{foo|bar|baz2|baz3|FA|f_method}}{{.*}}()
 // CHECK-DAG: define {{.*}}void @{{.*}}TemplateClass{{.*}}(%class.TemplateClass* %{{.*}})
 // CHECK-DAG: define {{.*}}i32 @{{.*}}TemplateClass{{.*}}f_method{{.*}}(%class.TemplateClass* %{{.*}})
-// CHECK-DAG: define {{.*}}void @__omp_offloading__{{.*}}_globals_l[[@LINE+74]]_ctor()
+// CHECK-DAG: define {{.*}}void @__omp_offloading__{{.*}}_globals_l[[@LINE+78]]_ctor()
 
 #ifndef HEADER
 #define HEADER
@@ -68,6 +76,10 @@ int hhh = 0;
 #pragma omp declare target link(eee, fff, ggg, hhh)
 
 int out_decl_target = 0;
+#ifdef OMP5
+#pragma omp declare target(out_decl_target)
+#endif
+
 #pragma omp declare target
 void lambda () {
 #ifdef __cpp_lambdas
@@ -224,4 +236,18 @@ int main() {
 // CHECK-DAG: !{i32 1, !"ccc", i32 0, i32 {{[0-9]+}}}
 // CHECK-DAG: !{{{.+}}virtual_foo
 
+#ifdef OMP5
+void host_fun() {}
+#pragma omp declare target to(host_fun) device_type(host)
+void device_fun() {}
+#pragma omp declare target to(device_fun) device_type(nohost)
+// HOST5-NOT: define {{.*}}void {{.*}}device_fun{{.*}}
+// HOST5: define {{.*}}void {{.*}}host_fun{{.*}}
+// HOST5-NOT: define {{.*}}void {{.*}}device_fun{{.*}}
+
+// DEV5-NOT: define {{.*}}void {{.*}}host_fun{{.*}}
+// DEV5: define {{.*}}void {{.*}}device_fun{{.*}}
+// DEV5-NOT: define {{.*}}void {{.*}}host_fun{{.*}}
+#endif // OMP5
+
 #endif // HEADER
diff --git a/clang/test/OpenMP/declare_target_messages.cpp b/clang/test/OpenMP/declare_target_messages.cpp
index fb5bd34c4a8..f416eb6ac12 100644
--- a/clang/test/OpenMP/declare_target_messages.cpp
+++ b/clang/test/OpenMP/declare_target_messages.cpp
@@ -1,7 +1,10 @@
-// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -fnoopenmp-use-tls -ferror-limit 100 -o - %s
-// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp5 -fopenmp -fopenmp-version=50 -fnoopenmp-use-tls -ferror-limit 100 -o - %s
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp45 -fopenmp -fnoopenmp-use-tls -ferror-limit 100 -o - %s
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp5,host5 -fopenmp -fopenmp-version=50 -fnoopenmp-use-tls -ferror-limit 100 -o - %s
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp5,dev5 -fopenmp -fopenmp-is-device -fopenmp-targets=x86_64-apple-macos10.7.0 -aux-triple x86_64-apple-macos10.7.0 -fopenmp-version=50 -fnoopenmp-use-tls -ferror-limit 100 -o - %s
 
-// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp-simd -fnoopenmp-use-tls -ferror-limit 100 -o - %s
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp5,host5 -fopenmp-simd -fopenmp-version=50 -fnoopenmp-use-tls -ferror-limit 100 -o - %s
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp5,host5 -fopenmp-simd -fopenmp-is-device -fopenmp-version=50 -fnoopenmp-use-tls -ferror-limit 100 -o - %s
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp45 -fopenmp-simd -fnoopenmp-use-tls -ferror-limit 100 -o - %s
 
 #pragma omp end declare target // expected-error {{unexpected OpenMP directive '#pragma omp end declare target'}}
 
@@ -14,17 +17,23 @@ __thread int t; // expected-note {{defined as threadprivate or thread local}}
 void f();
 #pragma omp end declare target shared(a) // expected-warning {{extra tokens at the end of '#pragma omp end declare target' are ignored}}
 
-#pragma omp declare target map(a) // expected-error {{unexpected 'map' clause, only 'to' or 'link' clauses expected}}
+#pragma omp declare target map(a) // omp45-error {{unexpected 'map' clause, only 'to' or 'link' clauses expected}} omp5-error {{unexpected 'map' clause, only 'to', 'link' or 'device_type' clauses expected}}
 
 #pragma omp declare target to(foo1) // expected-error {{use of undeclared identifier 'foo1'}}
 
 #pragma omp declare target link(foo2) // expected-error {{use of undeclared identifier 'foo2'}}
 
+#pragma omp declare target to(f) device_type(any) device_type(any) device_type(host) // omp45-error {{unexpected 'device_type' clause, only 'to' or 'link' clauses expected}} omp5-warning 2 {{more than one 'device_type' clause is specified}} omp5-error {{'device_type(host)' does not match previously specified 'device_type(any)' for the same declaration}}
+
 void c();
 
 void func() {} // expected-note {{'func' defined here}}
 
-#pragma omp declare target link(func) allocate(a) // expected-error {{function name is not allowed in 'link' clause}} expected-error {{unexpected 'allocate' clause, only 'to' or 'link' clauses expected}}
+#pragma omp declare target link(func) allocate(a) // expected-error {{function name is not allowed in 'link' clause}} omp45-error {{unexpected 'allocate' clause, only 'to' or 'link' clauses expected}} omp5-error {{unexpected 'allocate' clause, only 'to', 'link' or 'device_type' clauses expected}}
+
+void bar();
+void baz() {bar();}
+#pragma omp declare target(bar) // omp5-warning {{declaration marked as declare target after first use, it may lead to incorrect results}}
 
 extern int b;
 
@@ -152,4 +161,30 @@ namespace {
 #pragma omp declare target to(x) to(x) // expected-error {{'x' appears multiple times in clauses on the same declare target directive}}
 #pragma omp declare target link(x) // expected-error {{'x' must not appear in both clauses 'to' and 'link'}}
 
+void bazz() {}
+#pragma omp declare target to(bazz) device_type(nohost) // omp45-error {{unexpected 'device_type' clause, only 'to' or 'link' clauses expected}}
+void bazzz() {bazz();}
+#pragma omp declare target to(bazzz) device_type(nohost) // omp45-error {{unexpected 'device_type' clause, only 'to' or 'link' clauses expected}}
+void any() {bazz();}
+void host1() {bazz();}
+#pragma omp declare target to(host1) device_type(host) // omp45-error {{unexpected 'device_type' clause, only 'to' or 'link' clauses expected}} dev5-note 2 {{marked as 'device_type(host)' here}}
+void host2() {bazz();}
+#pragma omp declare target to(host2)
+void device() {host1();}
+#pragma omp declare target to(device) device_type(nohost) // omp45-error {{unexpected 'device_type' clause, only 'to' or 'link' clauses expected}} host5-note 2 {{marked as 'device_type(nohost)' here}}
+void host3() {host1();}
+#pragma omp declare target to(host3)
+
+#pragma omp declare target
+void any1() {any();}
+void any2() {host1();} // dev5-error {{function with 'device_type(host)' is not available on device}}
+void any3() {device();} // host5-error {{function with 'device_type(nohost)' is not available on host}}
+void any4() {any2();}
+#pragma omp end declare target
+
+void any5() {any();}
+void any6() {host1();} // dev5-error {{function with 'device_type(host)' is not available on device}}
+void any7() {device();} // host5-error {{function with 'device_type(nohost)' is not available on host}}
+void any8() {any2();}
+
 #pragma omp declare target // expected-error {{expected '#pragma omp end declare target'}} expected-note {{to match this '#pragma omp declare target'}}
diff --git a/clang/test/OpenMP/nvptx_allocate_messages.cpp b/clang/test/OpenMP/nvptx_allocate_messages.cpp
index e6fb83f6346..5b4a99623ac 100644
--- a/clang/test/OpenMP/nvptx_allocate_messages.cpp
+++ b/clang/test/OpenMP/nvptx_allocate_messages.cpp
@@ -58,7 +58,7 @@ template <class T> T foo() {
   #pragma omp allocate(v) allocator(omp_cgroup_mem_alloc)
   v = ST<T>::m;
 #if defined(DEVICE) && !defined(REQUIRES)
-// expected-error@+2 2 {{expected an allocator expression inside of the target region; provide an allocator expression or use 'requires' directive with the 'dynamic_allocators' clause}}
+// expected-error@+2 {{expected an allocator expression inside of the target region; provide an allocator expression or use 'requires' directive with the 'dynamic_allocators' clause}}
 #endif // DEVICE && !REQUIRES
 #pragma omp parallel private(v) allocate(v)
   v = 0;
diff --git a/clang/test/OpenMP/nvptx_asm_delayed_diags.c b/clang/test/OpenMP/nvptx_asm_delayed_diags.c
index 3d347046e59..2f82abd7477 100644
--- a/clang/test/OpenMP/nvptx_asm_delayed_diags.c
+++ b/clang/test/OpenMP/nvptx_asm_delayed_diags.c
@@ -2,6 +2,10 @@
 // RUN: %clang_cc1 -verify -fopenmp -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized
 // RUN: %clang_cc1 -verify -DDIAGS -DIMMEDIATE -fopenmp -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized
 // RUN: %clang_cc1 -verify -DDIAGS -DDELAYED -fopenmp -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized
+// RUN: %clang_cc1 -fopenmp -x c -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -fopenmp-version=50 -emit-llvm-bc %s -o %t-x86-host.bc
+// RUN: %clang_cc1 -verify=expected,omp5 -fopenmp -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -fopenmp-version=50 %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,omp5 -DDIAGS -DOMP5 -DIMMEDIATE -fopenmp -fopenmp-version=50 -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,omp5 -DDIAGS -DOMP5 -DDELAYED -fopenmp -fopenmp-version=50 -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized
 // REQUIRES: x86-registered-target
 // REQUIRES: nvptx-registered-target
 
@@ -9,6 +13,22 @@
 // expected-no-diagnostics
 #endif // DIAGS
 
+#ifdef OMP5
+void bar(int r) {
+#ifdef IMMEDIATE
+// omp5-error@+4 {{invalid input constraint 'mx' in asm}}
+#endif // IMMEDIATE
+  __asm__("PR3908 %[lf] %[xx] %[li] %[r]"
+          : [ r ] "+r"(r)
+          : [ lf ] "mx"(0), [ li ] "mr"(0), [ xx ] "x"((double)(0)));
+}
+#ifdef IMMEDIATE
+#pragma omp declare target to(bar) device_type(nohost)
+#else
+#pragma omp declare target to(bar) device_type(host)
+#endif // IMMEDIATE
+#endif // OMP5
+
 void foo(int r) {
 #ifdef IMMEDIATE
 // expected-error@+4 {{invalid input constraint 'mx' in asm}}
diff --git a/clang/test/OpenMP/target_vla_messages.cpp b/clang/test/OpenMP/target_vla_messages.cpp
index cf6a024c6ba..38a0a660775 100644
--- a/clang/test/OpenMP/target_vla_messages.cpp
+++ b/clang/test/OpenMP/target_vla_messages.cpp
@@ -47,7 +47,7 @@ void target_template(int arg) {
 #pragma omp target
   {
 #ifdef NO_VLA
-    // expected-error@+2 2 {{variable length arrays are not supported for the current target}}
+    // expected-error@+2 {{variable length arrays are not supported for the current target}}
 #endif
     T vla[arg];
   }