Skip to content

Commit 8eb588d

Browse files
fadeevalfadeeval
authored
[SYCL] Support intel::reqd_work_group_size (#1328)
Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>
1 parent 463eccd commit 8eb588d

File tree

5 files changed

+191
-7
lines changed

5 files changed

+191
-7
lines changed

‎clang/include/clang/Basic/Attr.td

Lines changed: 11 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -236,6 +236,10 @@ class DefaultIntArgument<string name, int default> : IntArgument<name, 1> {
236236
int Default = default;
237237
}
238238

239+
class DefaultUnsignedArgument<string name, int default> : UnsignedArgument<name, 1> {
240+
int Default = default;
241+
}
242+
239243
// This argument is more complex, it includes the enumerator type name,
240244
// a list of strings to accept, and a list of enumerators to map them to.
241245
class EnumArgument<string name, string type, list<string> values,
@@ -2430,13 +2434,16 @@ def NoDeref : TypeAttr {
24302434
let Documentation = [NoDerefDocs];
24312435
}
24322436

2437+
// Default arguments in ReqWorkGroupSize can be used only with
2438+
// intel::reqd_work_group_size spelling.
24332439
def ReqdWorkGroupSize : InheritableAttr {
24342440
let Spellings = [GNU<"reqd_work_group_size">,
2435-
CXX11<"cl","reqd_work_group_size">];
2436-
let Args = [UnsignedArgument<"XDim">, UnsignedArgument<"YDim">,
2437-
UnsignedArgument<"ZDim">];
2441+
CXX11<"intel","reqd_work_group_size">,
2442+
CXX11<"cl","reqd_work_group_size">];
2443+
let Args = [UnsignedArgument<"XDim">, DefaultUnsignedArgument<"YDim", 1>,
2444+
DefaultUnsignedArgument<"ZDim", 1>];
24382445
let Subjects = SubjectList<[Function], ErrorDiag>;
2439-
let Documentation = [Undocumented];
2446+
let Documentation = [ReqdWorkGroupSizeAttrDocs];
24402447
}
24412448

24422449
def WorkGroupSizeHint : InheritableAttr {

‎clang/include/clang/Basic/AttrDocs.td

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2003,6 +2003,36 @@ device kernel, the attribute is ignored and it is not propagated to a kernel.
20032003
}];
20042004
}
20052005

2006+
def ReqdWorkGroupSizeAttrDocs : Documentation {
2007+
let Category = DocCatFunction;
2008+
let Heading = "reqd_work_group_size";
2009+
let Content = [{
2010+
This attribute is documented by both OpenCL and SYCL standards
2011+
and allows to specify exact *local_work_size* which must be used as
2012+
argument to **clEnqueueNDRangeKernel** (in OpenCL) or to
2013+
**parallel_for** in SYCL. This allows the compiler to optimize the
2014+
generated code appropriately for the kernel to which attribute is applied.
2015+
2016+
While semantic of this attribute is the same between OpenCL and SYCL,
2017+
spelling is a bit different:
2018+
2019+
SYCL 1.2.1 describes ``[[cl::reqd_work_group_size(X, Y, Z)]]`` spelling: this
2020+
attribute is legal on device functions and is propagated down to any caller of
2021+
those device functions, such that the kernel attributes are the sum of all
2022+
attributes of all device functions called in this kernel.
2023+
See section 6.7 Attributes for more details.
2024+
2025+
As Intel extension, ``[[intel::reqd_work_group_size(X, Y, Z)]]`` spelling is allowed
2026+
which features optional arguments `Y` and `Z`, those simplifies its usage if
2027+
only 1- or 2-dimensional ND-range is assumed by a programmer. These arguments
2028+
defaults to ``1``.
2029+
2030+
In OpenCL C, this attribute is available in GNU spelling
2031+
(``__attribute__((reqd_work_group_size(X, Y, Z)))``), see section
2032+
6.7.2 Optional Attribute Qualifiers of OpenCL 1.2 specification for details.
2033+
}];
2034+
}
2035+
20062036
def SYCLIntelMaxWorkGroupSizeAttrDocs : Documentation {
20072037
let Category = DocCatFunction;
20082038
let Heading = "max_work_group_size (IntelFPGA)";

‎clang/lib/Sema/SemaDeclAttr.cpp

Lines changed: 11 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -2922,14 +2922,22 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) {
29222922
return;
29232923

29242924
uint32_t WGSize[3];
2925+
if (AL.getKind() == ParsedAttr::AT_ReqdWorkGroupSize &&
2926+
AL.getAttributeSpellingListIndex() ==
2927+
ReqdWorkGroupSizeAttr::CXX11_intel_reqd_work_group_size) {
2928+
WGSize[1] = ReqdWorkGroupSizeAttr::DefaultYDim;
2929+
WGSize[2] = ReqdWorkGroupSizeAttr::DefaultZDim;
2930+
} else if (!checkAttributeNumArgs(S, AL, 3))
2931+
return;
2932+
29252933
for (unsigned i = 0; i < 3; ++i) {
2926-
const Expr *E = AL.getArgAsExpr(i);
2927-
if (!checkUInt32Argument(S, AL, E, WGSize[i], i,
2934+
if (i < AL.getNumArgs() &&
2935+
!checkUInt32Argument(S, AL, AL.getArgAsExpr(i), WGSize[i], i,
29282936
/*StrictlyUnsigned=*/true))
29292937
return;
29302938
if (WGSize[i] == 0) {
29312939
S.Diag(AL.getLoc(), diag::err_attribute_argument_is_zero)
2932-
<< AL << E->getSourceRange();
2940+
<< AL << AL.getArgAsExpr(i)->getSourceRange();
29332941
return;
29342942
}
29352943
}

‎clang/test/SemaSYCL/intel-reqd-work-group-size.cpp

Lines changed: 136 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,136 @@
1+
// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -verify -DTRIGGER_ERROR %s
2+
// RUN: %clang_cc1 -fsycl -fsycl-is-device -ast-dump %s | FileCheck %s
3+
// RUN: %clang_cc1 -fsycl -fsycl-is-host -fsyntax-only -verify %s
4+
5+
#ifndef __SYCL_DEVICE_ONLY__
6+
// expected-no-diagnostics
7+
class Functor {
8+
public:
9+
[[intel::reqd_work_group_size(4)]] void operator()() {}
10+
};
11+
12+
template <typename name, typename Func>
13+
void kernel(Func kernelFunc) {
14+
kernelFunc();
15+
}
16+
17+
void bar() {
18+
Functor f;
19+
kernel<class kernel_name>(f);
20+
}
21+
#else
22+
[[intel::reqd_work_group_size(4)]] void f4x1x1() {} // expected-note {{conflicting attribute is here}}
23+
// expected-note@-1 {{conflicting attribute is here}}
24+
[[intel::reqd_work_group_size(32)]] void f32x1x1() {} // expected-note {{conflicting attribute is here}}
25+
26+
[[intel::reqd_work_group_size(16)]] void f16x1x1() {} // expected-note {{conflicting attribute is here}}
27+
[[intel::reqd_work_group_size(16, 16)]] void f16x16x1() {} // expected-note {{conflicting attribute is here}}
28+
29+
[[intel::reqd_work_group_size(32, 32)]] void f32x32x1() {} // expected-note {{conflicting attribute is here}}
30+
[[intel::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {} // expected-note {{conflicting attribute is here}}
31+
32+
#ifdef TRIGGER_ERROR
33+
class Functor32 {
34+
public:
35+
[[cl::reqd_work_group_size(32)]] void operator()() {} // expected-error {{'reqd_work_group_size' attribute requires exactly 3 arguments}}
36+
};
37+
class Functor33 {
38+
public:
39+
[[intel::reqd_work_group_size(32, -4)]] void operator()() {} // expected-error {{'reqd_work_group_size' attribute requires a non-negative integral compile time constant expression}}
40+
};
41+
#endif // TRIGGER_ERROR
42+
43+
class Functor16 {
44+
public:
45+
[[intel::reqd_work_group_size(16)]] void operator()() {}
46+
};
47+
48+
class Functor64 {
49+
public:
50+
[[intel::reqd_work_group_size(64, 64)]] void operator()() {}
51+
};
52+
53+
class Functor16x16x16 {
54+
public:
55+
[[intel::reqd_work_group_size(16, 16, 16)]] void operator()() {}
56+
};
57+
58+
class Functor8 { // expected-error {{conflicting attributes applied to a SYCL kernel}}
59+
public:
60+
[[intel::reqd_work_group_size(8)]] void operator()() { // expected-note {{conflicting attribute is here}}
61+
f4x1x1();
62+
}
63+
};
64+
65+
class Functor {
66+
public:
67+
void operator()() {
68+
f4x1x1();
69+
}
70+
};
71+
72+
class FunctorAttr {
73+
public:
74+
__attribute__((reqd_work_group_size(128, 128, 128))) void operator()() {}
75+
};
76+
77+
template <typename name, typename Func>
78+
__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
79+
kernelFunc();
80+
}
81+
82+
void bar() {
83+
Functor16 f16;
84+
kernel<class kernel_name1>(f16);
85+
86+
Functor f;
87+
kernel<class kernel_name2>(f);
88+
89+
Functor16x16x16 f16x16x16;
90+
kernel<class kernel_name3>(f16x16x16);
91+
92+
FunctorAttr fattr;
93+
kernel<class kernel_name4>(fattr);
94+
95+
kernel<class kernel_name5>([]() [[intel::reqd_work_group_size(32, 32, 32)]] {
96+
f32x32x32();
97+
});
98+
99+
#ifdef TRIGGER_ERROR
100+
Functor8 f8;
101+
kernel<class kernel_name6>(f8);
102+
103+
kernel<class kernel_name7>([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}}
104+
f4x1x1();
105+
f32x1x1();
106+
});
107+
108+
kernel<class kernel_name8>([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}}
109+
f16x1x1();
110+
f16x16x1();
111+
});
112+
113+
kernel<class kernel_name9>([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}}
114+
f32x32x32();
115+
f32x32x1();
116+
});
117+
118+
// expected-error@+1 {{expected variable name or 'this' in lambda capture list}}
119+
kernel<class kernel_name10>([[intel::reqd_work_group_size(32, 32, 32)]][]() {
120+
f32x32x32();
121+
});
122+
123+
#endif // TRIGGER_ERROR
124+
}
125+
126+
// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name1
127+
// CHECK: ReqdWorkGroupSizeAttr {{.*}} 1 1 16
128+
// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name2
129+
// CHECK: ReqdWorkGroupSizeAttr {{.*}} 1 1 4
130+
// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name3
131+
// CHECK: ReqdWorkGroupSizeAttr {{.*}} 16 16 16
132+
// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name4
133+
// CHECK: ReqdWorkGroupSizeAttr {{.*}} 128 128 128
134+
// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name5
135+
// CHECK: ReqdWorkGroupSizeAttr {{.*}} 32 32 32
136+
#endif // __SYCL_DEVICE_ONLY__

‎clang/utils/TableGen/ClangAttrEmitter.cpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1312,6 +1312,9 @@ createArgument(const Record &Arg, StringRef Attr,
13121312
Ptr = std::make_unique<TypeArgument>(Arg, Attr);
13131313
else if (ArgName == "UnsignedArgument")
13141314
Ptr = std::make_unique<SimpleArgument>(Arg, Attr, "unsigned");
1315+
else if (ArgName == "DefaultUnsignedArgument")
1316+
Ptr = std::make_unique<DefaultSimpleArgument>(Arg, Attr, "unsigned",
1317+
Arg.getValueAsInt("Default"));
13151318
else if (ArgName == "VariadicUnsignedArgument")
13161319
Ptr = std::make_unique<VariadicArgument>(Arg, Attr, "unsigned");
13171320
else if (ArgName == "VariadicStringArgument")

0 commit comments

Comments
 (0)