6 files changed, 87 insertions, 3 deletions
diff --git a/clang/test/CodeGenCUDA/openmp-target.cu b/clang/test/CodeGenCUDA/openmp-target.cu
new file mode 100644
index 00000000000..869bde60729
--- /dev/null
+++ b/clang/test/CodeGenCUDA/openmp-target.cu
@@ -0,0 +1,20 @@
+// REQUIRES: x86-registered-target
+// REQUIRES: nvptx-registered-target
+
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm \
+// RUN:            -fopenmp -fopenmp-version=50 -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm \
+// RUN:            -fopenmp -fopenmp-version=50 -o - -x c++ %s | FileCheck %s
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device \
+// RUN:            -emit-llvm -o - %s | FileCheck -check-prefixes=DEV %s
+
+// CHECK: declare{{.*}}@_Z7nohost1v()
+// DEV-NOT: _Z7nohost1v
+void nohost1() {}
+#pragma omp declare target to(nohost1) device_type(nohost)
+
+// CHECK: declare{{.*}}@_Z7nohost2v()
+// DEV-NOT: _Z7nohost2v
+void nohost2() {nohost1();}
+#pragma omp declare target to(nohost2) device_type(nohost)
+
diff --git a/clang/test/OpenMP/declare_target_messages.cpp b/clang/test/OpenMP/declare_target_messages.cpp
index f416eb6ac12..cc6558debde 100644
--- a/clang/test/OpenMP/declare_target_messages.cpp
+++ b/clang/test/OpenMP/declare_target_messages.cpp
@@ -162,10 +162,10 @@ namespace {
 #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}}
+#pragma omp declare target to(bazz) device_type(nohost) // omp45-error {{unexpected 'device_type' clause, only 'to' or 'link' clauses expected}} host5-note {{marked as 'device_type(nohost)' here}}
 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 any() {bazz();} // host5-error {{function with 'device_type(nohost)' is not available on host}}
 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();}
diff --git a/clang/test/SemaCUDA/call-device-fn-from-host.cu b/clang/test/SemaCUDA/call-device-fn-from-host.cu
index ba1ce860204..5d506d65ea5 100644
--- a/clang/test/SemaCUDA/call-device-fn-from-host.cu
+++ b/clang/test/SemaCUDA/call-device-fn-from-host.cu
@@ -1,5 +1,7 @@
 // RUN: %clang_cc1 %s --std=c++11 -triple x86_64-unknown-linux -emit-llvm -o - \
 // RUN:   -verify -verify-ignore-unexpected=note
+// RUN: %clang_cc1 %s --std=c++11 -triple x86_64-unknown-linux -emit-llvm -o - \
+// RUN:   -verify -verify-ignore-unexpected=note -fopenmp
 
 // Note: This test won't work with -fsyntax-only, because some of these errors
 // are emitted during codegen.
diff --git a/clang/test/SemaCUDA/host-device-constexpr.cu b/clang/test/SemaCUDA/host-device-constexpr.cu
index 6625d722c19..6d81034fd34 100644
--- a/clang/test/SemaCUDA/host-device-constexpr.cu
+++ b/clang/test/SemaCUDA/host-device-constexpr.cu
@@ -1,5 +1,10 @@
 // RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs %s
-// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs %s -fcuda-is-device
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs %s \
+// RUN:            -fcuda-is-device
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs \
+// RUN:            -fopenmp %s
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs \
+// RUN:            -fopenmp %s -fcuda-is-device
 
 #include "Inputs/cuda.h"
 
diff --git a/clang/test/SemaCUDA/openmp-static-func.cu b/clang/test/SemaCUDA/openmp-static-func.cu
new file mode 100644
index 00000000000..f1f0c488749
--- /dev/null
+++ b/clang/test/SemaCUDA/openmp-static-func.cu
@@ -0,0 +1,14 @@
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only \
+// RUN:            -verify -fopenmp %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only \
+// RUN:            -verify -fopenmp -x hip %s
+// expected-no-diagnostics
+
+// Tests there is no assertion in Sema::markKnownEmitted when fopenmp is used
+// with CUDA/HIP host compilation.
+
+static void f() {}
+
+static void g() { f(); }
+
+static void h() { g(); }
diff --git a/clang/test/SemaCUDA/openmp-target.cu b/clang/test/SemaCUDA/openmp-target.cu
new file mode 100644
index 00000000000..2775dc1e2c5
--- /dev/null
+++ b/clang/test/SemaCUDA/openmp-target.cu
@@ -0,0 +1,43 @@
+// RUN: %clang_cc1 -triple x86_64 -verify=expected,dev \
+// RUN:            -verify-ignore-unexpected=note \
+// RUN:            -fopenmp -fopenmp-version=50 -o - %s
+// RUN: %clang_cc1 -triple x86_64 -verify -verify-ignore-unexpected=note\
+// RUN:            -fopenmp -fopenmp-version=50 -o - -x c++ %s
+// RUN: %clang_cc1 -triple x86_64 -verify=dev -verify-ignore-unexpected=note\
+// RUN:            -fcuda-is-device -o - %s
+
+#if __CUDA__
+#include "Inputs/cuda.h"
+__device__ void cu_devf();
+#endif
+
+void bazz() {}
+#pragma omp declare target to(bazz) device_type(nohost)
+void bazzz() {bazz();}
+#pragma omp declare target to(bazzz) device_type(nohost)
+void any() {bazz();} // expected-error {{function with 'device_type(nohost)' is not available on host}}
+void host1() {bazz();}
+#pragma omp declare target to(host1) device_type(host)
+void host2() {bazz();}
+#pragma omp declare target to(host2)
+void device() {host1();}
+#pragma omp declare target to(device) device_type(nohost)
+void host3() {host1();}
+#pragma omp declare target to(host3)
+
+#pragma omp declare target
+void any1() {any();}
+void any2() {host1();}
+void any3() {device();} // expected-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();}
+void any7() {device();} // expected-error {{function with 'device_type(nohost)' is not available on host}}
+void any8() {any2();}
+
+#if __CUDA__
+void cu_hostf() { cu_devf(); } // dev-error {{no matching function for call to 'cu_devf'}}
+__device__ void cu_devf2() { cu_hostf(); } // dev-error{{no matching function for call to 'cu_hostf'}}
+#endif