diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 820d6d62d027b..bc2ee224d6386 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -8007,10 +8007,29 @@ static bool checkSYCLAddIRAttributesMergeability(const AddIRAttrT &NewAttr, void Sema::CheckSYCLAddIRAttributesFunctionAttrConflicts(Decl *D) { const auto *AddIRFuncAttr = D->getAttr(); - if (!AddIRFuncAttr || AddIRFuncAttr->args_size() == 0 || + + // If there is no such attribute there is nothing to check. If there are + // dependent arguments we cannot know the actual number of arguments so we + // defer the check. + if (!AddIRFuncAttr || hasDependentExpr(AddIRFuncAttr->args_begin(), AddIRFuncAttr->args_size())) return; + // If there are no name-value pairs in the attribute it will not have an + // effect and we can skip the check. The filter is ignored. + size_t NumArgsWithoutFilter = + AddIRFuncAttr->args_size() - (AddIRFuncAttr->hasFilterList() ? 1 : 0); + if (NumArgsWithoutFilter == 0) + return; + + // "sycl-single-task" is present on all single_task invocations, implicitly + // added by the SYCL headers. It can only conflict with max_global_work_dim, + // but the value will be the same so there is no need for a warning. + if (NumArgsWithoutFilter == 2 && + AddIRFuncAttr->getAttributeNameValuePairs(Context)[0].first == + "sycl-single-task") + return; + // If there are potentially conflicting attributes, we issue a warning. for (const auto *Attr : std::vector{ D->getAttr(), diff --git a/clang/test/SemaSYCL/attr-add-ir-attributes-function-conflict.cpp b/clang/test/SemaSYCL/attr-add-ir-attributes-function-conflict.cpp index 90ad47a4ae46f..4a2dd626b4b37 100644 --- a/clang/test/SemaSYCL/attr-add-ir-attributes-function-conflict.cpp +++ b/clang/test/SemaSYCL/attr-add-ir-attributes-function-conflict.cpp @@ -5,19 +5,48 @@ #include "sycl.hpp" -constexpr const char AttrName1[] = "Attr1"; -constexpr const char AttrVal1[] = "Val1"; +struct NameValuePair { + static constexpr const char *name = "Attr1"; + static constexpr const int value = 1; +}; + +template struct Wrapper { + template + [[__sycl_detail__::add_ir_attributes_function(Pairs::name..., Pairs::value...)]] __attribute__((sycl_kernel)) void kernel_single_task(const KernelType &kernelFunc) { + kernelFunc(); + } +}; + +template struct WrapperWithImplicit { + template + [[__sycl_detail__::add_ir_attributes_function("sycl-single-task", Pairs::name..., 0, Pairs::value...)]] __attribute__((sycl_kernel)) void kernel_single_task(const KernelType &kernelFunc) { + kernelFunc(); + } +}; + +template struct WrapperWithFilter { + template + [[__sycl_detail__::add_ir_attributes_function({"some-filter-string"}, Pairs::name..., Pairs::value...)]] __attribute__((sycl_kernel)) void kernel_single_task(const KernelType &kernelFunc) { + kernelFunc(); + } +}; -template struct Wrapper { +template struct WrapperWithImplicitAndFilter { template - [[__sycl_detail__::add_ir_attributes_function(Strs...)]] __attribute__((sycl_kernel)) void kernel_single_task(const KernelType &kernelFunc) { + [[__sycl_detail__::add_ir_attributes_function({"some-filter-string"}, "sycl-single-task", Pairs::name..., 0, Pairs::value...)]] __attribute__((sycl_kernel)) void kernel_single_task(const KernelType &kernelFunc) { kernelFunc(); } }; int main() { Wrapper<> EmptyWrapper; - Wrapper NonemptyWrapper; + Wrapper NonemptyWrapper; + WrapperWithImplicit<> EmptyWrapperWithImplicit; + WrapperWithImplicit NonemptyWrapperWithImplicit; + WrapperWithFilter<> EmptyWrapperWithFilter; + WrapperWithFilter NonemptyWrapperWithFilter; + WrapperWithImplicitAndFilter<> EmptyWrapperWithImplicitAndFilter; + WrapperWithImplicitAndFilter NonemptyWrapperWithImplicitAndFilter; EmptyWrapper.kernel_single_task([]() [[sycl::reqd_work_group_size(1)]] {}); EmptyWrapper.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2)]] {}); @@ -30,6 +59,39 @@ int main() { EmptyWrapper.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu)]] {}); EmptyWrapper.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {}); + EmptyWrapperWithImplicit.kernel_single_task([]() [[sycl::reqd_work_group_size(1)]] {}); + EmptyWrapperWithImplicit.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2)]] {}); + EmptyWrapperWithImplicit.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2,3)]] {}); + EmptyWrapperWithImplicit.kernel_single_task([]() [[sycl::work_group_size_hint(1)]] {}); + EmptyWrapperWithImplicit.kernel_single_task([]() [[sycl::work_group_size_hint(1,2)]] {}); + EmptyWrapperWithImplicit.kernel_single_task([]() [[sycl::work_group_size_hint(1,2,3)]] {}); + EmptyWrapperWithImplicit.kernel_single_task([]() [[sycl::reqd_sub_group_size(1)]] {}); + EmptyWrapperWithImplicit.kernel_single_task([]() [[sycl::device_has()]] {}); + EmptyWrapperWithImplicit.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu)]] {}); + EmptyWrapperWithImplicit.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {}); + + EmptyWrapperWithFilter.kernel_single_task([]() [[sycl::reqd_work_group_size(1)]] {}); + EmptyWrapperWithFilter.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2)]] {}); + EmptyWrapperWithFilter.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2,3)]] {}); + EmptyWrapperWithFilter.kernel_single_task([]() [[sycl::work_group_size_hint(1)]] {}); + EmptyWrapperWithFilter.kernel_single_task([]() [[sycl::work_group_size_hint(1,2)]] {}); + EmptyWrapperWithFilter.kernel_single_task([]() [[sycl::work_group_size_hint(1,2,3)]] {}); + EmptyWrapperWithFilter.kernel_single_task([]() [[sycl::reqd_sub_group_size(1)]] {}); + EmptyWrapperWithFilter.kernel_single_task([]() [[sycl::device_has()]] {}); + EmptyWrapperWithFilter.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu)]] {}); + EmptyWrapperWithFilter.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {}); + + EmptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::reqd_work_group_size(1)]] {}); + EmptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2)]] {}); + EmptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2,3)]] {}); + EmptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::work_group_size_hint(1)]] {}); + EmptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::work_group_size_hint(1,2)]] {}); + EmptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::work_group_size_hint(1,2,3)]] {}); + EmptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::reqd_sub_group_size(1)]] {}); + EmptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::device_has()]] {}); + EmptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu)]] {}); + EmptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {}); + // expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}} NonemptyWrapper.kernel_single_task([]() [[sycl::reqd_work_group_size(1)]] {}); // expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}} @@ -50,4 +112,67 @@ int main() { NonemptyWrapper.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu)]] {}); // expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}} NonemptyWrapper.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {}); + + // expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithImplicit.kernel_single_task([]() [[sycl::reqd_work_group_size(1)]] {}); + // expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithImplicit.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2)]] {}); + // expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithImplicit.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2,3)]] {}); + // expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithImplicit.kernel_single_task([]() [[sycl::work_group_size_hint(1)]] {}); + // expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithImplicit.kernel_single_task([]() [[sycl::work_group_size_hint(1,2)]] {}); + // expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithImplicit.kernel_single_task([]() [[sycl::work_group_size_hint(1,2,3)]] {}); + // expected-warning@+1 {{kernel has both attribute 'reqd_sub_group_size' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithImplicit.kernel_single_task([]() [[sycl::reqd_sub_group_size(1)]] {}); + // expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithImplicit.kernel_single_task([]() [[sycl::device_has()]] {}); + // expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithImplicit.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu)]] {}); + // expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithImplicit.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {}); + + // expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithFilter.kernel_single_task([]() [[sycl::reqd_work_group_size(1)]] {}); + // expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithFilter.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2)]] {}); + // expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithFilter.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2,3)]] {}); + // expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithFilter.kernel_single_task([]() [[sycl::work_group_size_hint(1)]] {}); + // expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithFilter.kernel_single_task([]() [[sycl::work_group_size_hint(1,2)]] {}); + // expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithFilter.kernel_single_task([]() [[sycl::work_group_size_hint(1,2,3)]] {}); + // expected-warning@+1 {{kernel has both attribute 'reqd_sub_group_size' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithFilter.kernel_single_task([]() [[sycl::reqd_sub_group_size(1)]] {}); + // expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithFilter.kernel_single_task([]() [[sycl::device_has()]] {}); + // expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithFilter.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu)]] {}); + // expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithFilter.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {}); + + // expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::reqd_work_group_size(1)]] {}); + // expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2)]] {}); + // expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2,3)]] {}); + // expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::work_group_size_hint(1)]] {}); + // expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::work_group_size_hint(1,2)]] {}); + // expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::work_group_size_hint(1,2,3)]] {}); + // expected-warning@+1 {{kernel has both attribute 'reqd_sub_group_size' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::reqd_sub_group_size(1)]] {}); + // expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::device_has()]] {}); + // expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu)]] {}); + // expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}} + NonemptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {}); } diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp index cdffd1d0019ef..1d7fd512ba868 100644 --- a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp +++ b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp @@ -19,6 +19,7 @@ #include "llvm/IR/IntrinsicInst.h" #include "llvm/IR/Module.h" #include "llvm/IR/Operator.h" +#include "llvm/TargetParser/Triple.h" using namespace llvm; @@ -209,6 +210,19 @@ attributeToExecModeMetadata(Module &M, const Attribute &Attr) { MDNode::get(Ctx, MD)); } + // The sycl-single-task attribute currently only has an effect when targeting + // SPIR FPGAs, in which case it will generate a "max_global_work_dim" MD node + // with a 0 value, similar to applying [[intel::max_global_work_dim(0)]] to + // a SYCL single_target kernel. + if (AttrKindStr == "sycl-single-task" && + Triple(M.getTargetTriple()).getSubArch() == Triple::SPIRSubArch_fpga) { + IntegerType *Ty = Type::getInt32Ty(Ctx); + Metadata *MDVal = ConstantAsMetadata::get(Constant::getNullValue(Ty)); + SmallVector MD{MDVal}; + return std::pair("max_global_work_dim", + MDNode::get(Ctx, MD)); + } + auto getIpInterface = [](const char *Name, LLVMContext &Ctx, const Attribute &Attr) { // generate either: diff --git a/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/fpga-single-task-property.ll b/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/fpga-single-task-property.ll new file mode 100644 index 0000000000000..85175e2b8efe0 --- /dev/null +++ b/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/fpga-single-task-property.ll @@ -0,0 +1,15 @@ +; RUN: opt -passes=compile-time-properties --mtriple=spir64_fpga-unknown-unknown %s -S | FileCheck %s --check-prefix CHECK-FPGA-IR +; RUN: opt -passes=compile-time-properties %s -S | FileCheck %s --check-prefix CHECK-DEFAULT-IR + +; CHECK-DEFAULT-IR-NOT: !max_global_work_dim + +; CHECK-FPGA-IR-DAG: @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE9TheKernel0"() #0 {{.*}}!max_global_work_dim ![[MaxGlobWorkDim:[0-9]+]] +; Function Attrs: convergent norecurse +define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE9TheKernel0"() #0 { +entry: + ret void +} + +attributes #0 = { convergent norecurse "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="fpga_single_task_property.cpp" "uniform-work-group-size"="true" "sycl-single-task" } + +; CHECK-FPGA-IR-DAG: ![[MaxGlobWorkDim]] = !{i32 0} diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index a12fd31fed127..0cf58142e43a5 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -1157,7 +1157,9 @@ class __SYCL_EXPORT handler { template #ifdef __SYCL_DEVICE_ONLY__ [[__sycl_detail__::add_ir_attributes_function( + "sycl-single-task", ext::oneapi::experimental::detail::PropertyMetaInfo::name..., + nullptr, ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] #endif __SYCL_KERNEL_ATTR__ void @@ -1174,7 +1176,9 @@ class __SYCL_EXPORT handler { template #ifdef __SYCL_DEVICE_ONLY__ [[__sycl_detail__::add_ir_attributes_function( + "sycl-single-task", ext::oneapi::experimental::detail::PropertyMetaInfo::name..., + nullptr, ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] #endif __SYCL_KERNEL_ATTR__ void diff --git a/sycl/test/check_device_code/fpga_single_task_max_global_work_dim.cpp b/sycl/test/check_device_code/fpga_single_task_max_global_work_dim.cpp new file mode 100644 index 0000000000000..057430ae8f921 --- /dev/null +++ b/sycl/test/check_device_code/fpga_single_task_max_global_work_dim.cpp @@ -0,0 +1,34 @@ +// RUN: %clangxx -fsycl-device-only -fsycl-targets=spir64_fpga -S -emit-llvm %s -o - | FileCheck %s -check-prefix=CHECK-FPGA +// RUN: %clangxx -fsycl-device-only -S -emit-llvm %s -o - | FileCheck %s -check-prefix=CHECK-DEFAULT + +// Tests that single_task implicitly adds the max_global_work_dim when AOT +// compiling for FPGA. +// Additionally it checks that existing attributes do not cause conflicts. + +#include + +int main() { + sycl::queue Q; + // CHECK-FPGA: spir_kernel void @_ZTSZ4mainE7Kernel1() {{.*}} !max_global_work_dim ![[MAX_GLOBAL_WORK_DIM:[0-9]+]] + // CHECK-DEFAULT-NOT: spir_kernel void @_ZTSZ4mainE7Kernel1() {{.*}} !max_global_work_dim + Q.single_task([]() {}); + // CHECK-FPGA: spir_kernel void @_ZTSZ4mainE7Kernel2() {{.*}} !max_global_work_dim ![[MAX_GLOBAL_WORK_DIM]] + // CHECK-DEFAULT: spir_kernel void @_ZTSZ4mainE7Kernel2() {{.*}} !max_global_work_dim ![[MAX_GLOBAL_WORK_DIM:[0-9]+]] + Q.single_task([]() [[intel::max_global_work_dim(0)]] {}); + // CHECK-FPGA: spir_kernel void @_ZTSZ4mainE7Kernel3() {{.*}} !max_global_work_dim ![[MAX_GLOBAL_WORK_DIM]] + // CHECK-DEFAULT: spir_kernel void @_ZTSZ4mainE7Kernel3() {{.*}} !max_global_work_dim ![[MAX_GLOBAL_WORK_DIM]] + Q.single_task( + []() [[sycl::work_group_size_hint(1), intel::max_global_work_dim(0)]] {}); + // CHECK-FPGA: spir_kernel void @_ZTSZ4mainE7Kernel4() {{.*}} !max_global_work_dim ![[MAX_GLOBAL_WORK_DIM]] + // CHECK-DEFAULT: spir_kernel void @_ZTSZ4mainE7Kernel4() {{.*}} !max_global_work_dim ![[MAX_GLOBAL_WORK_DIM]] + Q.single_task( + []() [[intel::max_global_work_dim(0), sycl::reqd_work_group_size(1)]] {}); + // CHECK-FPGA: spir_kernel void @_ZTSZ4mainE7Kernel5() {{.*}} !max_global_work_dim ![[MAX_GLOBAL_WORK_DIM]] + // CHECK-DEFAULT: spir_kernel void @_ZTSZ4mainE7Kernel5() {{.*}} !max_global_work_dim ![[MAX_GLOBAL_WORK_DIM]] + Q.single_task( + []() [[sycl::work_group_size_hint(1), intel::max_global_work_dim(0), + sycl::reqd_work_group_size(1)]] {}); + return 0; +} + +// CHECK-FPGA: ![[MAX_GLOBAL_WORK_DIM:[0-9]+]] = !{i32 0}