saiislam updated this revision to Diff 398897.
saiislam added a comment.
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.
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/include/clang/Basic/DiagnosticSemaKinds.td
clang/lib/Parse/ParseOpenMP.cpp
clang/test/OpenMP/metadirective_device_isa_codegen.cpp
clang/test/OpenMP/metadirective_device_isa_codegen_amdgcn.cpp
clang/test/OpenMP/metadirective_messages.cpp
Index: clang/test/OpenMP/metadirective_messages.cpp
===================================================================
--- clang/test/OpenMP/metadirective_messages.cpp
+++ clang/test/OpenMP/metadirective_messages.cpp
@@ -17,4 +17,6 @@
;
#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}}
+ ;
}
Index: clang/test/OpenMP/metadirective_device_isa_codegen_amdgcn.cpp
===================================================================
--- /dev/null
+++ 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
Index: clang/test/OpenMP/metadirective_device_isa_codegen.cpp
===================================================================
--- /dev/null
+++ 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
Index: clang/lib/Parse/ParseOpenMP.cpp
===================================================================
--- clang/lib/Parse/ParseOpenMP.cpp
+++ clang/lib/Parse/ParseOpenMP.cpp
@@ -2214,7 +2214,7 @@
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 @@
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>());
Index: clang/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -10810,11 +10810,6 @@
"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">;
Index: clang/include/clang/Basic/DiagnosticParseKinds.td
===================================================================
--- clang/include/clang/Basic/DiagnosticParseKinds.td
+++ clang/include/clang/Basic/DiagnosticParseKinds.td
@@ -1376,7 +1376,7 @@
"%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">,
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
