[clang] 876b5ea - [OpenMP][Clang] Allow passing target features in ISA trait for metadirective clause
Saiyedul Islam via cfe-commits
cfe-commits at lists.llvm.org
Tue Jan 11 21:26:14 PST 2022
Author: Saiyedul Islam
Date: 2022-01-12T05:24:49Z
New Revision: 876b5ea96bf5890074aec61bc2c6a37b2cdc0617
URL: https://github.com/llvm/llvm-project/commit/876b5ea96bf5890074aec61bc2c6a37b2cdc0617
DIFF: https://github.com/llvm/llvm-project/commit/876b5ea96bf5890074aec61bc2c6a37b2cdc0617.diff
LOG: [OpenMP][Clang] Allow passing target features in ISA trait for metadirective 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
Added:
clang/test/OpenMP/metadirective_device_isa_codegen.cpp
clang/test/OpenMP/metadirective_device_isa_codegen_amdgcn.cpp
Modified:
clang/include/clang/Basic/DiagnosticParseKinds.td
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/lib/Parse/ParseOpenMP.cpp
clang/test/OpenMP/metadirective_messages.cpp
Removed:
################################################################################
diff --git a/clang/include/clang/Basic/DiagnosticParseKinds.td b/clang/include/clang/Basic/DiagnosticParseKinds.td
index 193dff8b9c8f6..770ddb3ab16f6 100644
--- a/clang/include/clang/Basic/DiagnosticParseKinds.td
+++ b/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">,
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 1c8cd79910add..0a659688d82e0 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/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">;
diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp
index 0089da741b597..de3d58baf84c9 100644
--- a/clang/lib/Parse/ParseOpenMP.cpp
+++ b/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>());
diff --git a/clang/test/OpenMP/metadirective_device_isa_codegen.cpp b/clang/test/OpenMP/metadirective_device_isa_codegen.cpp
new file mode 100644
index 0000000000000..a1954c7b8bf1f
--- /dev/null
+++ b/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
diff --git a/clang/test/OpenMP/metadirective_device_isa_codegen_amdgcn.cpp b/clang/test/OpenMP/metadirective_device_isa_codegen_amdgcn.cpp
new file mode 100644
index 0000000000000..c6a5fb0f9dad3
--- /dev/null
+++ b/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
diff --git a/clang/test/OpenMP/metadirective_messages.cpp b/clang/test/OpenMP/metadirective_messages.cpp
index 5c89c495731f1..b342a094a7870 100644
--- a/clang/test/OpenMP/metadirective_messages.cpp
+++ b/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}}
+ ;
}
More information about the cfe-commits
mailing list