diff --git a/clang/include/clang/Basic/DiagnosticParseKinds.td b/clang/include/clang/Basic/DiagnosticParseKinds.td
--- a/clang/include/clang/Basic/DiagnosticParseKinds.td
+++ b/clang/include/clang/Basic/DiagnosticParseKinds.td
@@ -1379,6 +1379,9 @@
               "spelling or consider restricting the context selector with the "
               "'arch' selector further">,
       InGroup<SourceUsesOpenMP>;
+def remark_unknown_declare_variant_isa_trait
+    : Remark<"isa trait '%0' is not a valid feature of the target '%1'">,
+      InGroup<RemarkBackendPlugin>;
 def note_omp_declare_variant_ctx_options
     : Note<"context %select{set|selector|property}0 options are: %1">;
 def warn_omp_declare_variant_expected
diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp
--- a/clang/lib/Parse/ParseOpenMP.cpp
+++ b/clang/lib/Parse/ParseOpenMP.cpp
@@ -2529,7 +2529,12 @@
     TPA.Revert();
     // End of the first iteration. Parser is reset to the start of metadirective
 
-    TargetOMPContext OMPCtx(ASTContext, /* DiagUnknownTrait */ nullptr,
+    std::function<void(StringRef)> DiagUnknownTrait =
+        [this, Loc](StringRef ISATrait) {
+          Diag(Loc, diag::remark_unknown_declare_variant_isa_trait)
+              << ISATrait << this->getTargetInfo().getTriple().getArchName();
+        };
+    TargetOMPContext OMPCtx(ASTContext, std::move(DiagUnknownTrait),
                             /* CurrentFunctionDecl */ nullptr,
                             ArrayRef<llvm::omp::TraitProperty>());
 
diff --git a/clang/test/OpenMP/metadirective_device_isa_codegen.cpp b/clang/test/OpenMP/metadirective_device_isa_codegen.cpp
new file mode 100644
--- /dev/null
+++ b/clang/test/OpenMP/metadirective_device_isa_codegen.cpp
@@ -0,0 +1,81 @@
+// REQUIRES: amdgpu-registered-target
+// REQUIRES: x86-registered-target
+
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple amdgcn-amd-amdhsa -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -target-cpu gfx906 -o - | FileCheck %s -check-prefix=AMDGPUISA
+
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope -target-cpu x86-64| FileCheck %s -check-prefixes=X86_64ISA
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+int amdgcn_device_isa_selected() {
+  int threadCount = 0;
+
+#pragma omp target map(tofrom \
+                       : threadCount)
+  {
+#pragma omp metadirective                     \
+    when(device = {isa("flat-address-space")} \
+         : parallel) default(single)
+    threadCount++;
+  }
+
+  return threadCount;
+}
+
+// AMDGPUISA: define weak amdgpu_kernel void @__omp_offloading_{{.*}}amdgcn_device_isa_selected
+// AMDGPUISA: user_code.entry:
+// AMDGPUISA: call void @__kmpc_parallel_51
+// AMDGPUISA-NOT: call i32 @__kmpc_single
+// AMDGPUISA: ret void
+
+int amdgcn_device_isa_not_selected() {
+  int threadCount = 0;
+
+#pragma omp target map(tofrom \
+                       : threadCount)
+  {
+#pragma omp metadirective                                      \
+    when(device = {isa("sse")}                                 \
+         : parallel)                                           \
+        when(device = {isa("another-unsupported-gpu-feature")} \
+             : parallel) default(single)
+    threadCount++;
+  }
+
+  return threadCount;
+}
+// AMDGPUISA: define weak amdgpu_kernel void @__omp_offloading_{{.*}}amdgcn_device_isa_not_selected
+// AMDGPUISA: user_code.entry:
+// AMDGPUISA: call i32 @__kmpc_single
+// AMDGPUISA-NOT: call void @__kmpc_parallel_51
+// AMDGPUISA: ret void
+
+void bar();
+
+void x86_64_device_isa_selected() {
+#pragma omp metadirective when(device = {isa("sse2")} \
+                               : parallel) default(single)
+  bar();
+}
+// X86_64ISA-LABEL: void @_Z26x86_64_device_isa_selectedv()
+// X86_64ISA: ...) @__kmpc_fork_call{{.*}}@.omp_outlined..1
+// X86_64ISA: ret void
+
+// X86_64ISA: define internal void @.omp_outlined..1(
+// X86_64ISA: @_Z3barv
+// X86_64ISA: ret void
+
+void x86_64_device_isa_not_selected() {
+#pragma omp metadirective when(device = {isa("some-unsupported-feature")} \
+                               : parallel) default(single)
+  bar();
+}
+// X86_64ISA-LABEL: void @_Z30x86_64_device_isa_not_selectedv()
+// X86_64ISA: call i32 @__kmpc_single
+// X86_64ISA:  @_Z3barv
+// X86_64ISA: call void @__kmpc_end_single
+// X86_64ISA: ret void
+#endif
diff --git a/clang/test/OpenMP/metadirective_device_isa_messages.c b/clang/test/OpenMP/metadirective_device_isa_messages.c
new file mode 100644
--- /dev/null
+++ b/clang/test/OpenMP/metadirective_device_isa_messages.c
@@ -0,0 +1,14 @@
+// REQUIRES: x86-registered-target
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple x86_64-unknown-linux -emit-llvm-only -target-cpu znver1 %s -Rremark-backend-plugin
+
+#ifndef HEADER
+#define HEADER
+
+void bar();
+
+void foo() {
+#pragma omp metadirective when(device = {isa("some-unsupported-feature")} : parallel) default(single) // expected-remark {{isa trait 'some-unsupported-feature' is not a valid feature of the target 'x86_64'}}
+  bar();
+}
+
+#endif