[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

2533

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

2533

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

2533

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

2533

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

2533

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

2533

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/lib/Parse/ParseOpenMP.cpp

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) {};
+    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/amdgcn_target_codegen.cpp

 // REQUIRES: amdgpu-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 -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
 
 #define N 1000
 
 int test_amdgcn_target_tid_threads() {
 // CHECK-LABEL: define weak amdgpu_kernel void @{{.*}}test_amdgcn_target_tid_threads

 // CHECK: call i32 @__kmpc_target_init(%struct.ident_t* addrspacecast (%struct.ident_t addrspace(1)* @1 to %struct.ident_t*), i8 2, i1 false, i1 false)
 #pragma omp target simd
   for (int i = 0; i < N; i++) {
     arr[i] = 1;
   }
   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

clang/test/OpenMP/metadirective_implementation_codegen.c

 // 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
 #define HEADER
 
 void bar();
 
 void foo() {

 #pragma omp metadirective when(implementation = {extension(match_any)} \
                                : parallel) default(parallel for)
   for (int i = 0; i < 100; i++)
     ;
 #pragma omp metadirective when(implementation = {extension(match_none)} \
                                : 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.(
 // CHECK: @bar
 // CHECK: ret void
 
 // CHECK: define internal void @.omp_outlined..1(
 // CHECK: @bar

 // CHECK: call void @__kmpc_for_static_fini
 // CHECK: ret void
 
 // CHECK: define internal void @.omp_outlined..5(
 // NO-CHECK: call void @__kmpc_for_static_init
 // NO-CHECK: call void @__kmpc_for_static_fini
 // CHECK: ret void
 
+// SUPPORTEDISA: define internal void @.omp_outlined..6(
+// SUPPORTEDISA: @bar
+// SUPPORTEDISA: ret void
 #endif

clang/test/OpenMP/metadirective_implementation_codegen.cpp

 // 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
 #define HEADER
 
 void bar();
 
 void foo() {

 #pragma omp metadirective when(implementation = {extension(match_any)} \
                                : parallel) default(parallel for)
   for (int i = 0; i < 100; i++)
     ;
 #pragma omp metadirective when(implementation = {extension(match_none)} \
                                : 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()
 // CHECK: @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OUTLINED_2:@.+]] to void
 // CHECK: @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OUTLINED_3:@.+]] to void
 // CHECK: @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OUTLINED_4:@.+]] to void
 // 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]](
 // CHECK: @_Z3barv
 // CHECK: ret void
 
 // CHECK: define internal void [[OUTLINED_3]](
 // CHECK: @_Z3barv

 // CHECK: call void @__kmpc_for_static_fini
 // CHECK: ret void
 
 // CHECK: define internal void [[OUTLINED_7]](
 // NO-CHECK: call void @__kmpc_for_static_init
 // NO-CHECK: call void @__kmpc_for_static_fini
 // CHECK: ret void
 
+// SUPPORTEDISA: define internal void @.omp_outlined..6(
+// SUPPORTEDISA: @_Z3barv
+// SUPPORTEDISA: ret void
 #endif