Skip to content

Commit 9170a5d

Browse files
steffenlarsensteffenlarsen
authored
[SYCL] Add sycl-single-task implict property on single_task (#8190)
This commit adds the "sycl-single-task" LLVM IR attribute to all SYCL kernels originating from single_task calls. Leveraging the existing mechanism for creating attributes for the SYCL kernel compile-time properties extension, applying the add_ir_attributes_function to SYCL kernel wrappers, this new attribute acts like a kernel compile-time property. To avoid this implicit attribute from causing property/attribute conflict warnings, a special rule is made in the frontend to skip the conflict check if it is the only value in the add_ir_attributes_function attribute. Additionally, this commit adds special handling of "sycl-single-task" when targeting "spir64_fpga", in which case it will cause a "max_global_work_dim" metadata node to be created with value "0" on functions where it is present. --------- Signed-off-by: Larsen, Steffen <[email protected]>
1 parent f8a16f7 commit 9170a5d

File tree

6 files changed

+217
-6
lines changed

6 files changed

+217
-6
lines changed

‎clang/lib/Sema/SemaDeclAttr.cpp

Lines changed: 20 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -8007,10 +8007,29 @@ static bool checkSYCLAddIRAttributesMergeability(const AddIRAttrT &NewAttr,
80078007

80088008
void Sema::CheckSYCLAddIRAttributesFunctionAttrConflicts(Decl *D) {
80098009
const auto *AddIRFuncAttr = D->getAttr<SYCLAddIRAttributesFunctionAttr>();
8010-
if (!AddIRFuncAttr || AddIRFuncAttr->args_size() == 0 ||
8010+
8011+
// If there is no such attribute there is nothing to check. If there are
8012+
// dependent arguments we cannot know the actual number of arguments so we
8013+
// defer the check.
8014+
if (!AddIRFuncAttr ||
80118015
hasDependentExpr(AddIRFuncAttr->args_begin(), AddIRFuncAttr->args_size()))
80128016
return;
80138017

8018+
// If there are no name-value pairs in the attribute it will not have an
8019+
// effect and we can skip the check. The filter is ignored.
8020+
size_t NumArgsWithoutFilter =
8021+
AddIRFuncAttr->args_size() - (AddIRFuncAttr->hasFilterList() ? 1 : 0);
8022+
if (NumArgsWithoutFilter == 0)
8023+
return;
8024+
8025+
// "sycl-single-task" is present on all single_task invocations, implicitly
8026+
// added by the SYCL headers. It can only conflict with max_global_work_dim,
8027+
// but the value will be the same so there is no need for a warning.
8028+
if (NumArgsWithoutFilter == 2 &&
8029+
AddIRFuncAttr->getAttributeNameValuePairs(Context)[0].first ==
8030+
"sycl-single-task")
8031+
return;
8032+
80148033
// If there are potentially conflicting attributes, we issue a warning.
80158034
for (const auto *Attr : std::vector<AttributeCommonInfo *>{
80168035
D->getAttr<SYCLReqdWorkGroupSizeAttr>(),

‎clang/test/SemaSYCL/attr-add-ir-attributes-function-conflict.cpp

Lines changed: 130 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -5,19 +5,48 @@
55

66
#include "sycl.hpp"
77

8-
constexpr const char AttrName1[] = "Attr1";
9-
constexpr const char AttrVal1[] = "Val1";
8+
struct NameValuePair {
9+
static constexpr const char *name = "Attr1";
10+
static constexpr const int value = 1;
11+
};
12+
13+
template <typename... Pairs> struct Wrapper {
14+
template <typename KernelName, typename KernelType>
15+
[[__sycl_detail__::add_ir_attributes_function(Pairs::name..., Pairs::value...)]] __attribute__((sycl_kernel)) void kernel_single_task(const KernelType &kernelFunc) {
16+
kernelFunc();
17+
}
18+
};
19+
20+
template <typename... Pairs> struct WrapperWithImplicit {
21+
template <typename KernelName, typename KernelType>
22+
[[__sycl_detail__::add_ir_attributes_function("sycl-single-task", Pairs::name..., 0, Pairs::value...)]] __attribute__((sycl_kernel)) void kernel_single_task(const KernelType &kernelFunc) {
23+
kernelFunc();
24+
}
25+
};
26+
27+
template <typename... Pairs> struct WrapperWithFilter {
28+
template <typename KernelName, typename KernelType>
29+
[[__sycl_detail__::add_ir_attributes_function({"some-filter-string"}, Pairs::name..., Pairs::value...)]] __attribute__((sycl_kernel)) void kernel_single_task(const KernelType &kernelFunc) {
30+
kernelFunc();
31+
}
32+
};
1033

11-
template <const char *... Strs> struct Wrapper {
34+
template <typename... Pairs> struct WrapperWithImplicitAndFilter {
1235
template <typename KernelName, typename KernelType>
13-
[[__sycl_detail__::add_ir_attributes_function(Strs...)]] __attribute__((sycl_kernel)) void kernel_single_task(const KernelType &kernelFunc) {
36+
[[__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) {
1437
kernelFunc();
1538
}
1639
};
1740

1841
int main() {
1942
Wrapper<> EmptyWrapper;
20-
Wrapper<AttrName1, AttrVal1> NonemptyWrapper;
43+
Wrapper<NameValuePair> NonemptyWrapper;
44+
WrapperWithImplicit<> EmptyWrapperWithImplicit;
45+
WrapperWithImplicit<NameValuePair> NonemptyWrapperWithImplicit;
46+
WrapperWithFilter<> EmptyWrapperWithFilter;
47+
WrapperWithFilter<NameValuePair> NonemptyWrapperWithFilter;
48+
WrapperWithImplicitAndFilter<> EmptyWrapperWithImplicitAndFilter;
49+
WrapperWithImplicitAndFilter<NameValuePair> NonemptyWrapperWithImplicitAndFilter;
2150

2251
EmptyWrapper.kernel_single_task<class EK1>([]() [[sycl::reqd_work_group_size(1)]] {});
2352
EmptyWrapper.kernel_single_task<class EK2>([]() [[sycl::reqd_work_group_size(1,2)]] {});
@@ -30,6 +59,39 @@ int main() {
3059
EmptyWrapper.kernel_single_task<class EK9>([]() [[sycl::device_has(sycl::aspect::cpu)]] {});
3160
EmptyWrapper.kernel_single_task<class EK10>([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {});
3261

62+
EmptyWrapperWithImplicit.kernel_single_task<class EKWI1>([]() [[sycl::reqd_work_group_size(1)]] {});
63+
EmptyWrapperWithImplicit.kernel_single_task<class EKWI2>([]() [[sycl::reqd_work_group_size(1,2)]] {});
64+
EmptyWrapperWithImplicit.kernel_single_task<class EKWI3>([]() [[sycl::reqd_work_group_size(1,2,3)]] {});
65+
EmptyWrapperWithImplicit.kernel_single_task<class EKWI4>([]() [[sycl::work_group_size_hint(1)]] {});
66+
EmptyWrapperWithImplicit.kernel_single_task<class EKWI5>([]() [[sycl::work_group_size_hint(1,2)]] {});
67+
EmptyWrapperWithImplicit.kernel_single_task<class EKWI6>([]() [[sycl::work_group_size_hint(1,2,3)]] {});
68+
EmptyWrapperWithImplicit.kernel_single_task<class EKWI7>([]() [[sycl::reqd_sub_group_size(1)]] {});
69+
EmptyWrapperWithImplicit.kernel_single_task<class EKWI8>([]() [[sycl::device_has()]] {});
70+
EmptyWrapperWithImplicit.kernel_single_task<class EKWI9>([]() [[sycl::device_has(sycl::aspect::cpu)]] {});
71+
EmptyWrapperWithImplicit.kernel_single_task<class EKWI10>([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {});
72+
73+
EmptyWrapperWithFilter.kernel_single_task<class EKWF1>([]() [[sycl::reqd_work_group_size(1)]] {});
74+
EmptyWrapperWithFilter.kernel_single_task<class EKWF2>([]() [[sycl::reqd_work_group_size(1,2)]] {});
75+
EmptyWrapperWithFilter.kernel_single_task<class EKWF3>([]() [[sycl::reqd_work_group_size(1,2,3)]] {});
76+
EmptyWrapperWithFilter.kernel_single_task<class EKWF4>([]() [[sycl::work_group_size_hint(1)]] {});
77+
EmptyWrapperWithFilter.kernel_single_task<class EKWF5>([]() [[sycl::work_group_size_hint(1,2)]] {});
78+
EmptyWrapperWithFilter.kernel_single_task<class EKWF6>([]() [[sycl::work_group_size_hint(1,2,3)]] {});
79+
EmptyWrapperWithFilter.kernel_single_task<class EKWF7>([]() [[sycl::reqd_sub_group_size(1)]] {});
80+
EmptyWrapperWithFilter.kernel_single_task<class EKWF8>([]() [[sycl::device_has()]] {});
81+
EmptyWrapperWithFilter.kernel_single_task<class EKWF9>([]() [[sycl::device_has(sycl::aspect::cpu)]] {});
82+
EmptyWrapperWithFilter.kernel_single_task<class EKWF10>([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {});
83+
84+
EmptyWrapperWithImplicitAndFilter.kernel_single_task<class EKWIF1>([]() [[sycl::reqd_work_group_size(1)]] {});
85+
EmptyWrapperWithImplicitAndFilter.kernel_single_task<class EKWIF2>([]() [[sycl::reqd_work_group_size(1,2)]] {});
86+
EmptyWrapperWithImplicitAndFilter.kernel_single_task<class EKWIF3>([]() [[sycl::reqd_work_group_size(1,2,3)]] {});
87+
EmptyWrapperWithImplicitAndFilter.kernel_single_task<class EKWIF4>([]() [[sycl::work_group_size_hint(1)]] {});
88+
EmptyWrapperWithImplicitAndFilter.kernel_single_task<class EKWIF5>([]() [[sycl::work_group_size_hint(1,2)]] {});
89+
EmptyWrapperWithImplicitAndFilter.kernel_single_task<class EKWIF6>([]() [[sycl::work_group_size_hint(1,2,3)]] {});
90+
EmptyWrapperWithImplicitAndFilter.kernel_single_task<class EKWIF7>([]() [[sycl::reqd_sub_group_size(1)]] {});
91+
EmptyWrapperWithImplicitAndFilter.kernel_single_task<class EKWIF8>([]() [[sycl::device_has()]] {});
92+
EmptyWrapperWithImplicitAndFilter.kernel_single_task<class EKWIF9>([]() [[sycl::device_has(sycl::aspect::cpu)]] {});
93+
EmptyWrapperWithImplicitAndFilter.kernel_single_task<class EKWIF10>([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {});
94+
3395
// expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}}
3496
NonemptyWrapper.kernel_single_task<class NEK1>([]() [[sycl::reqd_work_group_size(1)]] {});
3597
// expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}}
@@ -50,4 +112,67 @@ int main() {
50112
NonemptyWrapper.kernel_single_task<class NEK9>([]() [[sycl::device_has(sycl::aspect::cpu)]] {});
51113
// expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}}
52114
NonemptyWrapper.kernel_single_task<class NEK10>([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {});
115+
116+
// expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}}
117+
NonemptyWrapperWithImplicit.kernel_single_task<class NEKWI1>([]() [[sycl::reqd_work_group_size(1)]] {});
118+
// expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}}
119+
NonemptyWrapperWithImplicit.kernel_single_task<class NEKWI2>([]() [[sycl::reqd_work_group_size(1,2)]] {});
120+
// expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}}
121+
NonemptyWrapperWithImplicit.kernel_single_task<class NEKWI3>([]() [[sycl::reqd_work_group_size(1,2,3)]] {});
122+
// expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}}
123+
NonemptyWrapperWithImplicit.kernel_single_task<class NEKWI4>([]() [[sycl::work_group_size_hint(1)]] {});
124+
// expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}}
125+
NonemptyWrapperWithImplicit.kernel_single_task<class NEKWI5>([]() [[sycl::work_group_size_hint(1,2)]] {});
126+
// expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}}
127+
NonemptyWrapperWithImplicit.kernel_single_task<class NEKWI6>([]() [[sycl::work_group_size_hint(1,2,3)]] {});
128+
// expected-warning@+1 {{kernel has both attribute 'reqd_sub_group_size' and kernel properties; conflicting properties are ignored}}
129+
NonemptyWrapperWithImplicit.kernel_single_task<class NEKWI7>([]() [[sycl::reqd_sub_group_size(1)]] {});
130+
// expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}}
131+
NonemptyWrapperWithImplicit.kernel_single_task<class NEKWI8>([]() [[sycl::device_has()]] {});
132+
// expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}}
133+
NonemptyWrapperWithImplicit.kernel_single_task<class NEKWI9>([]() [[sycl::device_has(sycl::aspect::cpu)]] {});
134+
// expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}}
135+
NonemptyWrapperWithImplicit.kernel_single_task<class NEKWI10>([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {});
136+
137+
// expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}}
138+
NonemptyWrapperWithFilter.kernel_single_task<class NEKWF1>([]() [[sycl::reqd_work_group_size(1)]] {});
139+
// expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}}
140+
NonemptyWrapperWithFilter.kernel_single_task<class NEKWF2>([]() [[sycl::reqd_work_group_size(1,2)]] {});
141+
// expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}}
142+
NonemptyWrapperWithFilter.kernel_single_task<class NEKWF3>([]() [[sycl::reqd_work_group_size(1,2,3)]] {});
143+
// expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}}
144+
NonemptyWrapperWithFilter.kernel_single_task<class NEKWF4>([]() [[sycl::work_group_size_hint(1)]] {});
145+
// expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}}
146+
NonemptyWrapperWithFilter.kernel_single_task<class NEKWF5>([]() [[sycl::work_group_size_hint(1,2)]] {});
147+
// expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}}
148+
NonemptyWrapperWithFilter.kernel_single_task<class NEKWF6>([]() [[sycl::work_group_size_hint(1,2,3)]] {});
149+
// expected-warning@+1 {{kernel has both attribute 'reqd_sub_group_size' and kernel properties; conflicting properties are ignored}}
150+
NonemptyWrapperWithFilter.kernel_single_task<class NEKWF7>([]() [[sycl::reqd_sub_group_size(1)]] {});
151+
// expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}}
152+
NonemptyWrapperWithFilter.kernel_single_task<class NEKWF8>([]() [[sycl::device_has()]] {});
153+
// expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}}
154+
NonemptyWrapperWithFilter.kernel_single_task<class NEKWF9>([]() [[sycl::device_has(sycl::aspect::cpu)]] {});
155+
// expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}}
156+
NonemptyWrapperWithFilter.kernel_single_task<class NEKWF10>([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {});
157+
158+
// expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}}
159+
NonemptyWrapperWithImplicitAndFilter.kernel_single_task<class NEKWIF1>([]() [[sycl::reqd_work_group_size(1)]] {});
160+
// expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}}
161+
NonemptyWrapperWithImplicitAndFilter.kernel_single_task<class NEKWIF2>([]() [[sycl::reqd_work_group_size(1,2)]] {});
162+
// expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}}
163+
NonemptyWrapperWithImplicitAndFilter.kernel_single_task<class NEKWIF3>([]() [[sycl::reqd_work_group_size(1,2,3)]] {});
164+
// expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}}
165+
NonemptyWrapperWithImplicitAndFilter.kernel_single_task<class NEKWIF4>([]() [[sycl::work_group_size_hint(1)]] {});
166+
// expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}}
167+
NonemptyWrapperWithImplicitAndFilter.kernel_single_task<class NEKWIF5>([]() [[sycl::work_group_size_hint(1,2)]] {});
168+
// expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}}
169+
NonemptyWrapperWithImplicitAndFilter.kernel_single_task<class NEKWIF6>([]() [[sycl::work_group_size_hint(1,2,3)]] {});
170+
// expected-warning@+1 {{kernel has both attribute 'reqd_sub_group_size' and kernel properties; conflicting properties are ignored}}
171+
NonemptyWrapperWithImplicitAndFilter.kernel_single_task<class NEKWIF7>([]() [[sycl::reqd_sub_group_size(1)]] {});
172+
// expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}}
173+
NonemptyWrapperWithImplicitAndFilter.kernel_single_task<class NEKWIF8>([]() [[sycl::device_has()]] {});
174+
// expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}}
175+
NonemptyWrapperWithImplicitAndFilter.kernel_single_task<class NEKWIF9>([]() [[sycl::device_has(sycl::aspect::cpu)]] {});
176+
// expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}}
177+
NonemptyWrapperWithImplicitAndFilter.kernel_single_task<class NEKWIF10>([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {});
53178
}

‎llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -19,6 +19,7 @@
1919
#include "llvm/IR/IntrinsicInst.h"
2020
#include "llvm/IR/Module.h"
2121
#include "llvm/IR/Operator.h"
22+
#include "llvm/TargetParser/Triple.h"
2223

2324
using namespace llvm;
2425

@@ -209,6 +210,19 @@ attributeToExecModeMetadata(Module &M, const Attribute &Attr) {
209210
MDNode::get(Ctx, MD));
210211
}
211212

213+
// The sycl-single-task attribute currently only has an effect when targeting
214+
// SPIR FPGAs, in which case it will generate a "max_global_work_dim" MD node
215+
// with a 0 value, similar to applying [[intel::max_global_work_dim(0)]] to
216+
// a SYCL single_target kernel.
217+
if (AttrKindStr == "sycl-single-task" &&
218+
Triple(M.getTargetTriple()).getSubArch() == Triple::SPIRSubArch_fpga) {
219+
IntegerType *Ty = Type::getInt32Ty(Ctx);
220+
Metadata *MDVal = ConstantAsMetadata::get(Constant::getNullValue(Ty));
221+
SmallVector<Metadata *, 1> MD{MDVal};
222+
return std::pair<std::string, MDNode *>("max_global_work_dim",
223+
MDNode::get(Ctx, MD));
224+
}
225+
212226
auto getIpInterface = [](const char *Name, LLVMContext &Ctx,
213227
const Attribute &Attr) {
214228
// generate either:

‎llvm/test/SYCLLowerIR/CompileTimePropertiesPass/fpga-single-task-property.ll

Lines changed: 15 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,15 @@
1+
; RUN: opt -passes=compile-time-properties --mtriple=spir64_fpga-unknown-unknown %s -S | FileCheck %s --check-prefix CHECK-FPGA-IR
2+
; RUN: opt -passes=compile-time-properties %s -S | FileCheck %s --check-prefix CHECK-DEFAULT-IR
3+
4+
; CHECK-DEFAULT-IR-NOT: !max_global_work_dim
5+
6+
; CHECK-FPGA-IR-DAG: @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE9TheKernel0"() #0 {{.*}}!max_global_work_dim ![[MaxGlobWorkDim:[0-9]+]]
7+
; Function Attrs: convergent norecurse
8+
define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE9TheKernel0"() #0 {
9+
entry:
10+
ret void
11+
}
12+
13+
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" }
14+
15+
; CHECK-FPGA-IR-DAG: ![[MaxGlobWorkDim]] = !{i32 0}

‎sycl/include/sycl/handler.hpp

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1157,7 +1157,9 @@ class __SYCL_EXPORT handler {
11571157
template <typename KernelName, typename KernelType, typename... Props>
11581158
#ifdef __SYCL_DEVICE_ONLY__
11591159
[[__sycl_detail__::add_ir_attributes_function(
1160+
"sycl-single-task",
11601161
ext::oneapi::experimental::detail::PropertyMetaInfo<Props>::name...,
1162+
nullptr,
11611163
ext::oneapi::experimental::detail::PropertyMetaInfo<Props>::value...)]]
11621164
#endif
11631165
__SYCL_KERNEL_ATTR__ void
@@ -1174,7 +1176,9 @@ class __SYCL_EXPORT handler {
11741176
template <typename KernelName, typename KernelType, typename... Props>
11751177
#ifdef __SYCL_DEVICE_ONLY__
11761178
[[__sycl_detail__::add_ir_attributes_function(
1179+
"sycl-single-task",
11771180
ext::oneapi::experimental::detail::PropertyMetaInfo<Props>::name...,
1181+
nullptr,
11781182
ext::oneapi::experimental::detail::PropertyMetaInfo<Props>::value...)]]
11791183
#endif
11801184
__SYCL_KERNEL_ATTR__ void

0 commit comments

Comments
 (0)