[OpenMP][Clang] Allow passing target features in ISA trait for metadirective clause

[saiislam]

Summary

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

[jdoerfert]

Can you make the test check for the diagnose message? Also, do we have a test to verify an isa trait is properly handled?

clang/lib/Parse/ParseOpenMP.cpp

2555

Why doesn't this diagnose nothing?

[saiislam]

In D116549#3218281, @jdoerfert wrote:

Can you make the test check for the diagnose message? Also, do we have a test to verify an isa trait is properly handled?

I don't see a point in adding a diagnostic message when a feature is not found valid for a certain target. Because if it satisfies one of the host or device compilation, then it will definitely fail in the other.
I have added new tests to check that isa trait is handled properly in the next revision.

clang/lib/Parse/ParseOpenMP.cpp

2555

Because an isa-feature will fail at least once, for either host compilation or device compilation. So, no point in always giving a warning.

[saiislam]

Added target specific tests for ISA traits, for CPU as well as GPU.

[jdoerfert]

clang/lib/Parse/ParseOpenMP.cpp

2555

That is debatable.

First, if I compile for a single architecture there is no device compilation and it should warn.
Second, if I place the metadirective into a declare variant function or add a kind(...) selector to it it will also not warn even if you have multiple architectures.

[saiislam]

Added diagnostic remarks for when ISA trait is not selected.

[saiislam]

Fixed the lit test failing in pre-check build bot.

[jdoerfert]

clang/lib/Parse/ParseOpenMP.cpp

2555

ASTContext &Context = getASTContext();
std::function<void(StringRef)> DiagUnknownTrait = [this,
┊ ┊ ┊ ┊ ┊ ┊ ┊ ┊ ┊ ┊                                CE](StringRef ISATrait) {
┊ // TODO Track the selector locations in a way that is accessible here to
┊ // improve the diagnostic location.
┊ Diag(CE->getBeginLoc(), diag::warn_unknown_declare_variant_isa_trait)
┊ ┊ ┊ << ISATrait;                                                   
};
TargetOMPContext OMPCtx(Context, std::move(DiagUnknownTrait),                                                                                                                                                                              
┊ ┊ ┊ ┊ ┊ ┊ ┊ ┊ ┊ ┊     getCurFunctionDecl(), DSAStack->getConstructTraits());

Already exists (SemaOpenMP). Why do we need a second, different diagnostic?

[saiislam]

clang/lib/Parse/ParseOpenMP.cpp

2555

Isn't giving a remark better than a warning, when we know in many cases this will be hit during a normal (expected) compilation for target offload?
Remark diagnostic will give ample handle for understanding the flow without the need to explicitly deal with this warning during compilation of user programs.

I am fine changing it to a warning if you feel strongly about this.

[jdoerfert]

clang/lib/Parse/ParseOpenMP.cpp

2555

1. Having two diagnostics for the same thing is generally bad and there doesn't seem to be a reason why we would need/want to treat metadirective and declare variant differently. Thus, if we want to move to remarks we should move the second diagnostic as well. Cleaning up existing code is more important than adding new code.
2. I still don't see why this is "normal" or "expected" at all. The warning reads:

isa trait 'foo' is not known to the current target; verify the spelling or consider restricting the context selector with the 'arch' selector further.
Hence, device={isa("flat..."), arch(amdgcn)} should not cause a warning even if you compile it for the host, an nvidia gpu, or anything other than AMDGCN.
FWIW, this is [supposed to be] ensured by the positioning of the isa trait selector in OMPKinds.def:

// Note that we put isa last so that the other conditions are checked first.
// This allows us to issue warnings wrt. isa only if we match otherwise.
__OMP_TRAIT_SELECTOR(device, isa, true)

[saiislam]

1. Used a common diagnostic warning warn_unknown_declare_variant_isa_trait for ParseOpenMP and SemaOpenMP for decalre variant and metadirectives.
2. Split lit codegen tests into two files, one requiring amdgpu-registered target and another for host only.
3. Added warning message lit test at an appropriate place.

[jdoerfert]

LG, thanks for the adjustment. Hope you are happy with the result.

clang/include/clang/Basic/DiagnosticParseKinds.td

def err_omp_declare_variant_duplicate_nested_trait : Error<
   " in selector '%1' and set '%2' with different score">;
 def err_omp_declare_variant_nested_user_condition : Error<
   "nested user conditions in OpenMP context selector not supported (yet)">;
 def warn_omp_declare_variant_string_literal_or_identifier
     : Warning<"expected identifier or string literal describing a context "
               "%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">,
       InGroup<SourceUsesOpenMP>;
 def note_omp_declare_variant_ctx_options
     : Note<"context %select{set|selector|property}0 options are: %1">;
 def warn_omp_declare_variant_expected
     : Warning<"expected '%0' after the %1; '%0' assumed">,

clang/include/clang/Basic/DiagnosticSemaKinds.td

 def err_omp_flush_order_clause_and_list : Error<
   "'flush' directive with memory order clause '%0' cannot have the list">;
 def note_omp_flush_order_clause_here : Note<
   "memory order clause '%0' is specified here">;
 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">;
 def err_omp_reduction_task_not_parallel_or_worksharing : Error<
   "'reduction' clause with 'task' modifier allowed only on non-simd parallel or"
   " worksharing constructs">;
 def err_omp_expected_array_alloctraits : Error<
   "expected constant sized array of 'omp_alloctrait_t' elements, not %0">;

clang/lib/Parse/ParseOpenMP.cpp

case OMPD_begin_declare_variant: {
 
     VariantMatchInfo VMI;
     TI.getAsVariantMatchInfo(ASTCtx, VMI);
 
     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_begin_declare_variant_isa_trait) << ISATrait;
+      Diag(Loc, diag::warn_unknown_declare_variant_isa_trait) << ISATrait;
     };
     TargetOMPContext OMPCtx(
         ASTCtx, std::move(DiagUnknownTrait),
         /* CurrentFunctionDecl */ nullptr,
         /* ConstructTraits */ ArrayRef<llvm::omp::TraitProperty>());
 
     if (isVariantApplicableInContext(VMI, OMPCtx, /* DeviceSetOnly */ true)) {
       Actions.ActOnOpenMPBeginDeclareVariant(Loc, TI);
while (Tok.isNot(tok::annot_pragma_openmp_end)) {
       TI.getAsVariantMatchInfo(ASTContext, VMI);
 
       VMIs.push_back(VMI);
     }
 
     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>());
 
     // A single match is returned for OpenMP 5.0
     int BestIdx = getBestVariantMatchForContext(VMIs, OMPCtx);
 
     int Idx = 0;
     // In OpenMP 5.0 metadirective is either replaced by another directive or

clang/test/OpenMP/metadirective_device_isa_codegen.cpp

+// 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

+// 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

 #pragma omp metadirective when(device{arch(nvptx)}) // expected-error {{missing ':' in when clause}} expected-error {{expected expression}} expected-warning {{expected '=' after the context set name "device"; '=' assumed}}
   ;
 #pragma omp metadirective when(device{arch(nvptx)}: ) default() // expected-warning {{expected '=' after the context set name "device"; '=' assumed}}
   ;
 #pragma omp metadirective when(device = {arch(nvptx)} : ) default(xyz) // expected-error {{expected an OpenMP directive}} expected-error {{use of undeclared identifier 'xyz'}}
   ;
 #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}}
+  ;
 }