saiislam updated this revision to Diff 398179.
saiislam marked an inline comment as done.
saiislam added a comment.
Added diagnostic remarks for when ISA trait is not selected.
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D116549/new/
https://reviews.llvm.org/D116549
Files:
clang/include/clang/Basic/DiagnosticParseKinds.td
clang/lib/Parse/ParseOpenMP.cpp
clang/test/OpenMP/metadirective_device_isa_codegen.cpp
clang/test/OpenMP/metadirective_device_isa_messages.c
Index: clang/test/OpenMP/metadirective_device_isa_messages.c
===================================================================
--- /dev/null
+++ clang/test/OpenMP/metadirective_device_isa_messages.c
@@ -0,0 +1,15 @@
+// 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
Index: clang/test/OpenMP/metadirective_device_isa_codegen.cpp
===================================================================
--- /dev/null
+++ 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
Index: clang/lib/Parse/ParseOpenMP.cpp
===================================================================
--- clang/lib/Parse/ParseOpenMP.cpp
+++ 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>());
Index: clang/include/clang/Basic/DiagnosticParseKinds.td
===================================================================
--- clang/include/clang/Basic/DiagnosticParseKinds.td
+++ 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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
