[SYCL] Add sycl-single-task implict property on single_task

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<SYCLAddIRAttributesFunctionAttr>();
-  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<AttributeCommonInfo *>{
            D->getAttr<SYCLReqdWorkGroupSizeAttr>(),

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 <typename... Pairs> struct Wrapper {
+  template <typename KernelName, typename KernelType>
+  [[__sycl_detail__::add_ir_attributes_function(Pairs::name..., Pairs::value...)]] __attribute__((sycl_kernel)) void kernel_single_task(const KernelType &kernelFunc) {
+    kernelFunc();
+  }
+};
+
+template <typename... Pairs> struct WrapperWithImplicit {
+  template <typename KernelName, typename KernelType>
+  [[__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 <typename... Pairs> struct WrapperWithFilter {
+  template <typename KernelName, typename KernelType>
+  [[__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 <const char *... Strs> struct Wrapper {
+template <typename... Pairs> struct WrapperWithImplicitAndFilter {
   template <typename KernelName, typename KernelType>
-  [[__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<AttrName1, AttrVal1> NonemptyWrapper;
+  Wrapper<NameValuePair> NonemptyWrapper;
+  WrapperWithImplicit<> EmptyWrapperWithImplicit;
+  WrapperWithImplicit<NameValuePair> NonemptyWrapperWithImplicit;
+  WrapperWithFilter<> EmptyWrapperWithFilter;
+  WrapperWithFilter<NameValuePair> NonemptyWrapperWithFilter;
+  WrapperWithImplicitAndFilter<> EmptyWrapperWithImplicitAndFilter;
+  WrapperWithImplicitAndFilter<NameValuePair> NonemptyWrapperWithImplicitAndFilter;
 
   EmptyWrapper.kernel_single_task<class EK1>([]() [[sycl::reqd_work_group_size(1)]] {});
   EmptyWrapper.kernel_single_task<class EK2>([]() [[sycl::reqd_work_group_size(1,2)]] {});
@@ -30,6 +59,39 @@ int main() {
   EmptyWrapper.kernel_single_task<class EK9>([]() [[sycl::device_has(sycl::aspect::cpu)]] {});
   EmptyWrapper.kernel_single_task<class EK10>([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {});
 
+  EmptyWrapperWithImplicit.kernel_single_task<class EKWI1>([]() [[sycl::reqd_work_group_size(1)]] {});
+  EmptyWrapperWithImplicit.kernel_single_task<class EKWI2>([]() [[sycl::reqd_work_group_size(1,2)]] {});
+  EmptyWrapperWithImplicit.kernel_single_task<class EKWI3>([]() [[sycl::reqd_work_group_size(1,2,3)]] {});
+  EmptyWrapperWithImplicit.kernel_single_task<class EKWI4>([]() [[sycl::work_group_size_hint(1)]] {});
+  EmptyWrapperWithImplicit.kernel_single_task<class EKWI5>([]() [[sycl::work_group_size_hint(1,2)]] {});
+  EmptyWrapperWithImplicit.kernel_single_task<class EKWI6>([]() [[sycl::work_group_size_hint(1,2,3)]] {});
+  EmptyWrapperWithImplicit.kernel_single_task<class EKWI7>([]() [[sycl::reqd_sub_group_size(1)]] {});
+  EmptyWrapperWithImplicit.kernel_single_task<class EKWI8>([]() [[sycl::device_has()]] {});
+  EmptyWrapperWithImplicit.kernel_single_task<class EKWI9>([]() [[sycl::device_has(sycl::aspect::cpu)]] {});
+  EmptyWrapperWithImplicit.kernel_single_task<class EKWI10>([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {});
+
+  EmptyWrapperWithFilter.kernel_single_task<class EKWF1>([]() [[sycl::reqd_work_group_size(1)]] {});
+  EmptyWrapperWithFilter.kernel_single_task<class EKWF2>([]() [[sycl::reqd_work_group_size(1,2)]] {});
+  EmptyWrapperWithFilter.kernel_single_task<class EKWF3>([]() [[sycl::reqd_work_group_size(1,2,3)]] {});
+  EmptyWrapperWithFilter.kernel_single_task<class EKWF4>([]() [[sycl::work_group_size_hint(1)]] {});
+  EmptyWrapperWithFilter.kernel_single_task<class EKWF5>([]() [[sycl::work_group_size_hint(1,2)]] {});
+  EmptyWrapperWithFilter.kernel_single_task<class EKWF6>([]() [[sycl::work_group_size_hint(1,2,3)]] {});
+  EmptyWrapperWithFilter.kernel_single_task<class EKWF7>([]() [[sycl::reqd_sub_group_size(1)]] {});
+  EmptyWrapperWithFilter.kernel_single_task<class EKWF8>([]() [[sycl::device_has()]] {});
+  EmptyWrapperWithFilter.kernel_single_task<class EKWF9>([]() [[sycl::device_has(sycl::aspect::cpu)]] {});
+  EmptyWrapperWithFilter.kernel_single_task<class EKWF10>([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {});
+
+  EmptyWrapperWithImplicitAndFilter.kernel_single_task<class EKWIF1>([]() [[sycl::reqd_work_group_size(1)]] {});
+  EmptyWrapperWithImplicitAndFilter.kernel_single_task<class EKWIF2>([]() [[sycl::reqd_work_group_size(1,2)]] {});
+  EmptyWrapperWithImplicitAndFilter.kernel_single_task<class EKWIF3>([]() [[sycl::reqd_work_group_size(1,2,3)]] {});
+  EmptyWrapperWithImplicitAndFilter.kernel_single_task<class EKWIF4>([]() [[sycl::work_group_size_hint(1)]] {});
+  EmptyWrapperWithImplicitAndFilter.kernel_single_task<class EKWIF5>([]() [[sycl::work_group_size_hint(1,2)]] {});
+  EmptyWrapperWithImplicitAndFilter.kernel_single_task<class EKWIF6>([]() [[sycl::work_group_size_hint(1,2,3)]] {});
+  EmptyWrapperWithImplicitAndFilter.kernel_single_task<class EKWIF7>([]() [[sycl::reqd_sub_group_size(1)]] {});
+  EmptyWrapperWithImplicitAndFilter.kernel_single_task<class EKWIF8>([]() [[sycl::device_has()]] {});
+  EmptyWrapperWithImplicitAndFilter.kernel_single_task<class EKWIF9>([]() [[sycl::device_has(sycl::aspect::cpu)]] {});
+  EmptyWrapperWithImplicitAndFilter.kernel_single_task<class EKWIF10>([]() [[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<class NEK1>([]() [[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<class NEK9>([]() [[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<class NEK10>([]() [[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<class NEKWI1>([]() [[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<class NEKWI2>([]() [[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<class NEKWI3>([]() [[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<class NEKWI4>([]() [[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<class NEKWI5>([]() [[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<class NEKWI6>([]() [[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<class NEKWI7>([]() [[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<class NEKWI8>([]() [[sycl::device_has()]] {});
+  // expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}}
+  NonemptyWrapperWithImplicit.kernel_single_task<class NEKWI9>([]() [[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<class NEKWI10>([]() [[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<class NEKWF1>([]() [[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<class NEKWF2>([]() [[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<class NEKWF3>([]() [[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<class NEKWF4>([]() [[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<class NEKWF5>([]() [[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<class NEKWF6>([]() [[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<class NEKWF7>([]() [[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<class NEKWF8>([]() [[sycl::device_has()]] {});
+  // expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}}
+  NonemptyWrapperWithFilter.kernel_single_task<class NEKWF9>([]() [[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<class NEKWF10>([]() [[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<class NEKWIF1>([]() [[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<class NEKWIF2>([]() [[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<class NEKWIF3>([]() [[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<class NEKWIF4>([]() [[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<class NEKWIF5>([]() [[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<class NEKWIF6>([]() [[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<class NEKWIF7>([]() [[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<class NEKWIF8>([]() [[sycl::device_has()]] {});
+  // expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}}
+  NonemptyWrapperWithImplicitAndFilter.kernel_single_task<class NEKWIF9>([]() [[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<class NEKWIF10>([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {});
 }

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<Metadata *, 1> MD{MDVal};
+    return std::pair<std::string, MDNode *>("max_global_work_dim",
+                                            MDNode::get(Ctx, MD));
+  }
+
   auto getIpInterface = [](const char *Name, LLVMContext &Ctx,
                            const Attribute &Attr) {
     // generate either:

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}

sycl/include/sycl/handler.hpp

@@ -1157,7 +1157,9 @@ class __SYCL_EXPORT handler {
   template <typename KernelName, typename KernelType, typename... Props>
 #ifdef __SYCL_DEVICE_ONLY__
   [[__sycl_detail__::add_ir_attributes_function(
+      "sycl-single-task",
       ext::oneapi::experimental::detail::PropertyMetaInfo<Props>::name...,
+      nullptr,
       ext::oneapi::experimental::detail::PropertyMetaInfo<Props>::value...)]]
 #endif
   __SYCL_KERNEL_ATTR__ void
@@ -1174,7 +1176,9 @@ class __SYCL_EXPORT handler {
   template <typename KernelName, typename KernelType, typename... Props>
 #ifdef __SYCL_DEVICE_ONLY__
   [[__sycl_detail__::add_ir_attributes_function(
+      "sycl-single-task",
       ext::oneapi::experimental::detail::PropertyMetaInfo<Props>::name...,
+      nullptr,
       ext::oneapi::experimental::detail::PropertyMetaInfo<Props>::value...)]]
 #endif
   __SYCL_KERNEL_ATTR__ void