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,9 @@
     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) {};
+    TargetOMPContext OMPCtx(ASTContext, std::move(DiagUnknownTrait),
                             /* CurrentFunctionDecl */ nullptr,
                             ArrayRef<llvm::omp::TraitProperty>());
 
diff --git a/clang/test/OpenMP/amdgcn_target_codegen.cpp b/clang/test/OpenMP/amdgcn_target_codegen.cpp
--- a/clang/test/OpenMP/amdgcn_target_codegen.cpp
+++ b/clang/test/OpenMP/amdgcn_target_codegen.cpp
@@ -2,6 +2,7 @@
 
 // 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 -o - | FileCheck %s
+// 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=AMDGPUTARGET
 // expected-no-diagnostics
 #ifndef HEADER
 #define HEADER
@@ -35,4 +36,46 @@
   return arr[0];
 }
 
+// AMDGPUTARGET: define weak amdgpu_kernel void @__omp_offloading_{{.*}}test_amdgcn_device_isa_selected
+// AMDGPUTARGET: user_code.entry:
+// AMDGPUTARGET: call void @__kmpc_parallel_51
+// AMDGPUTARGET-NOT: call i32 @__kmpc_single
+// AMDGPUTARGET: ret void
+int test_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;
+}
+
+// AMDGPUTARGET: define weak amdgpu_kernel void @__omp_offloading_{{.*}}test_amdgcn_device_isa_not_selected
+// AMDGPUTARGET: user_code.entry:
+// AMDGPUTARGET: call i32 @__kmpc_single
+// AMDGPUTARGET-NOT: call void @__kmpc_parallel_51
+// AMDGPUTARGET: ret void
+int test_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;
+}
+
 #endif
diff --git a/clang/test/OpenMP/metadirective_implementation_codegen.c b/clang/test/OpenMP/metadirective_implementation_codegen.c
--- a/clang/test/OpenMP/metadirective_implementation_codegen.c
+++ b/clang/test/OpenMP/metadirective_implementation_codegen.c
@@ -1,6 +1,7 @@
 // RUN: %clang_cc1 -verify -fopenmp -x c -triple x86_64-unknown-linux -emit-llvm %s -o - | FileCheck %s
 // RUN: %clang_cc1 -verify -fopenmp -x c -triple aarch64-unknown-linux -emit-llvm %s -o - | FileCheck %s
 // RUN: %clang_cc1 -verify -fopenmp -x c -triple ppc64le-unknown-linux -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple x86_64-unknown-linux -emit-llvm -target-cpu x86-64 %s -o - | FileCheck %s -check-prefixes=SUPPORTEDISA,UNSUPPORTEDISA
 // expected-no-diagnostics
 
 #ifndef HEADER
@@ -35,10 +36,20 @@
                                : parallel) default(parallel for)
   for (int i = 0; i < 100; i++)
     ;
+#pragma omp metadirective when(device = {isa("sse2")} \
+                               : parallel) default(single)
+  bar();
+#pragma omp metadirective when(device = {isa("some-unsupported-feature")} \
+                               : parallel) default(single)
+  bar();
 }
 
 // CHECK: void @foo()
 // CHECK-COUNT-6: ...) @__kmpc_fork_call(
+// SUPPORTEDISA: ...) @__kmpc_fork_call(
+// UNSUPPORTEDISA: call i32 @__kmpc_single
+// UNSUPPORTEDISA: @bar
+// UNSUPPORTEDISA: call void @__kmpc_end_single
 // CHECK: ret void
 
 // CHECK: define internal void @.omp_outlined.(
@@ -68,4 +79,7 @@
 // NO-CHECK: call void @__kmpc_for_static_fini
 // CHECK: ret void
 
+// SUPPORTEDISA: define internal void @.omp_outlined..6(
+// SUPPORTEDISA: @bar
+// SUPPORTEDISA: ret void
 #endif
diff --git a/clang/test/OpenMP/metadirective_implementation_codegen.cpp b/clang/test/OpenMP/metadirective_implementation_codegen.cpp
--- a/clang/test/OpenMP/metadirective_implementation_codegen.cpp
+++ b/clang/test/OpenMP/metadirective_implementation_codegen.cpp
@@ -1,6 +1,7 @@
 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope | FileCheck %s
 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple aarch64-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope | FileCheck %s
 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple ppc64le-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope | FileCheck %s
+// 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=SUPPORTEDISA,UNSUPPORTEDISA
 // expected-no-diagnostics
 
 #ifndef HEADER
@@ -35,6 +36,12 @@
                                : parallel) default(parallel for)
   for (int i = 0; i < 100; i++)
     ;
+#pragma omp metadirective when(device = {isa("sse2")} \
+                               : parallel) default(single)
+  bar();
+#pragma omp metadirective when(device = {isa("some-unsupported-feature")} \
+                               : parallel) default(single)
+  bar();
 }
 
 // CHECK-LABEL: void @_Z3foov()
@@ -44,6 +51,10 @@
 // CHECK: @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OUTLINED_5:@.+]] to void
 // CHECK: @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OUTLINED_6:@.+]] to void
 // CHECK: @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OUTLINED_7:@.+]] to void
+// SUPPORTEDISA: ...) @__kmpc_fork_call(
+// UNSUPPORTEDISA: call i32 @__kmpc_single
+// UNSUPPORTEDISA:  @_Z3barv
+// UNSUPPORTEDISA: call void @__kmpc_end_single
 // CHECK: ret void
 
 // CHECK: define internal void [[OUTLINED_2]](
@@ -73,4 +84,7 @@
 // NO-CHECK: call void @__kmpc_for_static_fini
 // CHECK: ret void
 
+// SUPPORTEDISA: define internal void @.omp_outlined..6(
+// SUPPORTEDISA: @_Z3barv
+// SUPPORTEDISA: ret void
 #endif