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}} + ; } _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
