diff --git a/clang/include/clang/Basic/DiagnosticParseKinds.td b/clang/include/clang/Basic/DiagnosticParseKinds.td --- a/clang/include/clang/Basic/DiagnosticParseKinds.td +++ b/clang/include/clang/Basic/DiagnosticParseKinds.td @@ -1376,7 +1376,7 @@ "%select{set|selector|property}0; " "%select{set|selector|property}0 skipped">, InGroup; -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 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/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; 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 --- a/clang/lib/Parse/ParseOpenMP.cpp +++ b/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 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()); diff --git a/clang/test/OpenMP/metadirective_device_isa_codegen.cpp b/clang/test/OpenMP/metadirective_device_isa_codegen.cpp new file mode 100644 --- /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 --- /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 --- a/clang/test/OpenMP/metadirective_messages.cpp +++ b/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}} + ; }