[OpenMP][Clang] Allow passing target features in ISA trait for metadi…

…rective clause

Passing any feature in the device-isa trait which is not supported by the host
was causing a compilation failure.

Differential Revision: https://reviews.llvm.org/D116549

clang/include/clang/Basic/DiagnosticParseKinds.td

@@ -1376,7 +1376,7 @@ def warn_omp_declare_variant_string_literal_or_identifier
               "%select{set|selector|property}0; "
               "%select{set|selector|property}0 skipped">,
       InGroup<OpenMPClauses>;
-def warn_unknown_begin_declare_variant_isa_trait
+def warn_unknown_declare_variant_isa_trait
     : Warning<"isa trait '%0' is not known to the current target; verify the "
               "spelling or consider restricting the context selector with the "
               "'arch' selector further">,

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -10810,11 +10810,6 @@ def err_omp_non_lvalue_in_map_or_motion_clauses: Error<
   "expected addressable lvalue in '%0' clause">;
 def err_omp_var_expected : Error<
   "expected variable of the '%0' type%select{|, not %2}1">;
-def warn_unknown_declare_variant_isa_trait
-    : Warning<"isa trait '%0' is not known to the current target; verify the "
-              "spelling or consider restricting the context selector with the "
-              "'arch' selector further">,
-      InGroup<SourceUsesOpenMP>;
 def err_omp_non_pointer_type_array_shaping_base : Error<
   "expected expression with a pointer to a complete type as a base of an array "
   "shaping operation">;

clang/lib/Parse/ParseOpenMP.cpp

@@ -2214,7 +2214,7 @@ Parser::DeclGroupPtrTy Parser::ParseOpenMPDeclarativeDirectiveWithExtDecl(
                                                           StringRef ISATrait) {
       // TODO Track the selector locations in a way that is accessible here to
       // improve the diagnostic location.
-      Diag(Loc, diag::warn_unknown_begin_declare_variant_isa_trait) << ISATrait;
+      Diag(Loc, diag::warn_unknown_declare_variant_isa_trait) << ISATrait;
     };
     TargetOMPContext OMPCtx(
         ASTCtx, std::move(DiagUnknownTrait),
@@ -2551,7 +2551,13 @@ Parser::ParseOpenMPDeclarativeOrExecutableDirective(ParsedStmtContext StmtCtx) {
     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) {
+      // TODO Track the selector locations in a way that is accessible here to
+      // improve the diagnostic location.
+      Diag(Loc, diag::warn_unknown_declare_variant_isa_trait) << ISATrait;
+    };
+    TargetOMPContext OMPCtx(ASTContext, std::move(DiagUnknownTrait),
                             /* CurrentFunctionDecl */ nullptr,
                             ArrayRef<llvm::omp::TraitProperty>());
 

clang/test/OpenMP/metadirective_device_isa_codegen.cpp

@@ -0,0 +1,32 @@
+// RUN: %clang_cc1 -verify -w -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope | FileCheck %s
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+void bar();
+
+void x86_64_device_isa_selected() {
+#pragma omp metadirective when(device = {isa("sse2")} \
+                               : parallel) default(single)
+  bar();
+}
+// CHECK-LABEL: void @_Z26x86_64_device_isa_selectedv()
+// CHECK: ...) @__kmpc_fork_call{{.*}}@.omp_outlined.
+// CHECK: ret void
+
+// CHECK: define internal void @.omp_outlined.(
+// CHECK: @_Z3barv
+// CHECK: ret void
+
+void x86_64_device_isa_not_selected() {
+#pragma omp metadirective when(device = {isa("some-unsupported-feature")} \
+                               : parallel) default(single)
+  bar();
+}
+// CHECK-LABEL: void @_Z30x86_64_device_isa_not_selectedv()
+// CHECK: call i32 @__kmpc_single
+// CHECK:  @_Z3barv
+// CHECK: call void @__kmpc_end_single
+// CHECK: ret void
+#endif

clang/test/OpenMP/metadirective_device_isa_codegen_amdgcn.cpp

@@ -0,0 +1,53 @@
+// REQUIRES: amdgpu-registered-target
+
+// RUN: %clang_cc1 -fopenmp -x c++ -w -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++ -w -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
+// 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;
+}
+
+// CHECK: define weak amdgpu_kernel void @__omp_offloading_{{.*}}amdgcn_device_isa_selected
+// CHECK: user_code.entry:
+// CHECK: call void @__kmpc_parallel_51
+// CHECK-NOT: call i32 @__kmpc_single
+// CHECK: 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;
+}
+// CHECK: define weak amdgpu_kernel void @__omp_offloading_{{.*}}amdgcn_device_isa_not_selected
+// CHECK: user_code.entry:
+// CHECK: call i32 @__kmpc_single
+// CHECK-NOT: call void @__kmpc_parallel_51
+// CHECK: ret void
+
+#endif

clang/test/OpenMP/metadirective_messages.cpp

@@ -17,4 +17,6 @@ void foo() {
   ;
 #pragma omp metadirective when(device = {arch(nvptx)} : parallel default() // expected-error {{expected ',' or ')' in 'when' clause}} expected-error {{expected expression}}
   ;
+#pragma omp metadirective when(device = {isa("some-unsupported-feature")} : parallel) default(single) // expected-warning {{isa trait 'some-unsupported-feature' is not known to the current target; verify the spelling or consider restricting the context selector with the 'arch' selector further}}
+  ;
 }