Skip to content

Commit d278c67

Browse files
steffenlarsensteffenlarsen
authored
[SYCL] Fail on kernel lambda size mismatch (#6374)
* [SYCL] Fail on kernel lambda size mismatch Some host-compilers may create lambdas with captures that do not match the ones used when extracting kernel descriptors. The compiler does not currently have a stable way of handling these cases, so instead this PR adds a static assertion that the sizes match with a message informing about the limitation. Signed-off-by: Larsen, Steffen <[email protected]>
1 parent 11b3094 commit d278c67

23 files changed

+107
-31
lines changed

‎clang/test/SemaSYCL/Inputs/CL/sycl/detail/kernel_desc.hpp

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -35,6 +35,17 @@ __SYCL_INLINE_NAMESPACE(cl) {
3535
int offset;
3636
};
3737

38+
template <bool Cond, typename TrueT, typename FalseT>
39+
struct conditional {
40+
using type = TrueT;
41+
};
42+
template <typename TrueT, typename FalseT>
43+
struct conditional<false, TrueT, FalseT> {
44+
using type = FalseT;
45+
};
46+
47+
using int64_t = conditional<sizeof(long) == 8, long, long long>::type;
48+
3849
template <class KernelNameType> struct KernelInfo {
3950
static constexpr unsigned getNumParams() { return 0; }
4051
static const kernel_param_desc_t &getParamDesc(int) {
@@ -43,6 +54,7 @@ __SYCL_INLINE_NAMESPACE(cl) {
4354
}
4455
static constexpr const char *getName() { return ""; }
4556
static constexpr bool isESIMD() { return 0; }
57+
static constexpr int64_t getKernelSize() { return 0; }
4658
};
4759
} // namespace detail
4860
} // namespace sycl

‎sycl/include/CL/sycl/detail/kernel_desc.hpp

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -13,6 +13,8 @@
1313
#include <CL/sycl/detail/defines_elementary.hpp>
1414
#include <CL/sycl/detail/export.hpp>
1515

16+
#include <cstdint>
17+
1618
__SYCL_INLINE_NAMESPACE(cl) {
1719
namespace sycl {
1820
namespace detail {
@@ -79,6 +81,7 @@ template <class KernelNameType> struct KernelInfo {
7981
static constexpr const char *getFunctionName() { return ""; }
8082
static constexpr unsigned getLineNumber() { return 0; }
8183
static constexpr unsigned getColumnNumber() { return 0; }
84+
static constexpr int64_t getKernelSize() { return 0; }
8285
};
8386
#else
8487
template <char...> struct KernelInfoData {
@@ -93,6 +96,7 @@ template <char...> struct KernelInfoData {
9396
static constexpr const char *getFunctionName() { return ""; }
9497
static constexpr unsigned getLineNumber() { return 0; }
9598
static constexpr unsigned getColumnNumber() { return 0; }
99+
static constexpr int64_t getKernelSize() { return 0; }
96100
};
97101

98102
// C++14 like index_sequence and make_index_sequence
@@ -135,6 +139,9 @@ template <class KernelNameType> struct KernelInfo {
135139
static constexpr const char *getFunctionName() { return ""; }
136140
static constexpr unsigned getLineNumber() { return 0; }
137141
static constexpr unsigned getColumnNumber() { return 0; }
142+
static constexpr int64_t getKernelSize() {
143+
return SubKernelInfo::getKernelSize();
144+
}
138145
};
139146
#endif //__SYCL_UNNAMED_LAMBDA__
140147

‎sycl/include/CL/sycl/handler.hpp

Lines changed: 13 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -731,13 +731,25 @@ class __SYCL_EXPORT handler {
731731
"kernel_handler is not yet supported by host device.",
732732
PI_ERROR_INVALID_OPERATION);
733733
}
734+
734735
KernelType *KernelPtr =
735736
ResetHostKernel<KernelType, LambdaArgType, Dims>(KernelFunc);
736737

737738
using KI = sycl::detail::KernelInfo<KernelName>;
739+
constexpr bool KernelHasName =
740+
KI::getName() != nullptr && KI::getName()[0] != '\0';
741+
742+
// Some host compilers may have different captures from Clang. Currently
743+
// there is no stable way of handling this when extracting the captures, so
744+
// a static assert is made to fail for incompatible kernel lambdas.
745+
static_assert(!KernelHasName || sizeof(KernelFunc) == KI::getKernelSize(),
746+
"Unexpected kernel lambda size. This can be caused by an "
747+
"external host compiler producing a lambda with an "
748+
"unexpected layout. This is a limitation of the compiler.");
749+
738750
// Empty name indicates that the compilation happens without integration
739751
// header, so don't perform things that require it.
740-
if (KI::getName() != nullptr && KI::getName()[0] != '\0') {
752+
if (KernelHasName) {
741753
// TODO support ESIMD in no-integration-header case too.
742754
MArgs.clear();
743755
extractArgsAndReqsFromLambda(reinterpret_cast<char *>(KernelPtr),

‎sycl/test/basic_tests/kernel_size_mismatch.cpp

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,19 @@
1+
// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning -o - %s
2+
3+
// Tests for static assertion failure when kernel lambda mismatches between host
4+
// and device.
5+
6+
#include <CL/sycl.hpp>
7+
8+
int main() {
9+
sycl::queue Q;
10+
int A = 1;
11+
Q.single_task([=]() {
12+
#ifdef __SYCL_DEVICE_ONLY__
13+
(void)A;
14+
// expected-no-diagnostics
15+
#else
16+
// expected-error-re@CL/sycl/handler.hpp:* {{static_assert failed due to requirement '{{.*}}' "Unexpected kernel lambda size. This can be caused by an external host compiler producing a lambda with an unexpected layout. This is a limitation of the compiler."}}
17+
#endif
18+
}).wait();
19+
}

‎sycl/unittests/SYCL2020/GetNativeOpenCL.cpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -127,7 +127,8 @@ TEST(GetNative, GetNativeHandle) {
127127
sycl::buffer<int, 1> Buffer(&Data[0], sycl::range<1>(1));
128128
Queue.submit([&](sycl::handler &cgh) {
129129
auto Acc = Buffer.get_access<sycl::access::mode::read_write>(cgh);
130-
cgh.single_task<TestKernel>([=]() { (void)Acc; });
130+
constexpr size_t KS = sizeof(decltype(Acc));
131+
cgh.single_task<TestKernel<KS>>([=]() { (void)Acc; });
131132
});
132133

133134
get_native<backend::opencl>(Context);

‎sycl/unittests/SYCL2020/KernelBundle.cpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -31,6 +31,7 @@ template <> struct KernelInfo<TestKernel> {
3131
static constexpr bool isESIMD() { return false; }
3232
static constexpr bool callsThisItem() { return false; }
3333
static constexpr bool callsAnyThisFreeFunction() { return false; }
34+
static constexpr int64_t getKernelSize() { return 1; }
3435
};
3536

3637
template <> struct KernelInfo<TestKernelExeOnly> {
@@ -43,6 +44,7 @@ template <> struct KernelInfo<TestKernelExeOnly> {
4344
static constexpr bool isESIMD() { return false; }
4445
static constexpr bool callsThisItem() { return false; }
4546
static constexpr bool callsAnyThisFreeFunction() { return false; }
47+
static constexpr int64_t getKernelSize() { return 1; }
4648
};
4749

4850
} // namespace detail

‎sycl/unittests/SYCL2020/KernelID.cpp

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -32,6 +32,7 @@ template <> struct KernelInfo<TestKernel1> {
3232
static constexpr bool isESIMD() { return false; }
3333
static constexpr bool callsThisItem() { return false; }
3434
static constexpr bool callsAnyThisFreeFunction() { return false; }
35+
static constexpr int64_t getKernelSize() { return 1; }
3536
};
3637

3738
template <> struct KernelInfo<TestKernel2> {
@@ -44,6 +45,7 @@ template <> struct KernelInfo<TestKernel2> {
4445
static constexpr bool isESIMD() { return false; }
4546
static constexpr bool callsThisItem() { return false; }
4647
static constexpr bool callsAnyThisFreeFunction() { return false; }
48+
static constexpr int64_t getKernelSize() { return 1; }
4749
};
4850

4951
template <> struct KernelInfo<TestKernel3> {
@@ -56,6 +58,7 @@ template <> struct KernelInfo<TestKernel3> {
5658
static constexpr bool isESIMD() { return false; }
5759
static constexpr bool callsThisItem() { return false; }
5860
static constexpr bool callsAnyThisFreeFunction() { return false; }
61+
static constexpr int64_t getKernelSize() { return 1; }
5962
};
6063

6164
template <> struct KernelInfo<ServiceKernel1> {
@@ -70,6 +73,7 @@ template <> struct KernelInfo<ServiceKernel1> {
7073
static constexpr bool isESIMD() { return false; }
7174
static constexpr bool callsThisItem() { return false; }
7275
static constexpr bool callsAnyThisFreeFunction() { return false; }
76+
static constexpr int64_t getKernelSize() { return 1; }
7377
};
7478
} // namespace detail
7579
} // namespace sycl

‎sycl/unittests/SYCL2020/SpecializationConstant.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -35,6 +35,7 @@ template <> struct KernelInfo<TestKernel> {
3535
static constexpr bool isESIMD() { return false; }
3636
static constexpr bool callsThisItem() { return false; }
3737
static constexpr bool callsAnyThisFreeFunction() { return false; }
38+
static constexpr int64_t getKernelSize() { return 1; }
3839
};
3940

4041
template <> const char *get_spec_constant_symbolic_ID<SpecConst1>() {

‎sycl/unittests/assert/assert.cpp

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -50,6 +50,7 @@ template <> struct KernelInfo<TestKernel> {
5050
static constexpr bool isESIMD() { return false; }
5151
static constexpr bool callsThisItem() { return false; }
5252
static constexpr bool callsAnyThisFreeFunction() { return false; }
53+
static constexpr int64_t getKernelSize() { return 1; }
5354
};
5455

5556
static constexpr const kernel_param_desc_t Signatures[] = {
@@ -68,6 +69,11 @@ struct KernelInfo<::sycl::detail::__sycl_service_kernel__::AssertInfoCopier> {
6869
static constexpr bool isESIMD() { return 0; }
6970
static constexpr bool callsThisItem() { return 0; }
7071
static constexpr bool callsAnyThisFreeFunction() { return 0; }
72+
static constexpr int64_t getKernelSize() {
73+
// The AssertInfoCopier service kernel lambda captures an accessor.
74+
return sizeof(sycl::accessor<sycl::detail::AssertHappened, 1,
75+
sycl::access::mode::write>);
76+
}
7177
};
7278
} // namespace detail
7379
} // namespace sycl

‎sycl/unittests/buffer/BufferLocation.cpp

Lines changed: 6 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -116,7 +116,8 @@ TEST_F(BufferTest, BufferLocationOnly) {
116116
cl::sycl::ext::oneapi::accessor_property_list<
117117
cl::sycl::ext::intel::property::buffer_location::instance<2>>>
118118
Acc{Buf, cgh, sycl::read_write, PL};
119-
cgh.single_task<TestKernel>([=]() { Acc[0] = 4; });
119+
constexpr size_t KS = sizeof(decltype(Acc));
120+
cgh.single_task<TestKernel<KS>>([=]() { Acc[0] = 4; });
120121
})
121122
.wait();
122123
EXPECT_EQ(PassedLocation, (uint64_t)2);
@@ -149,7 +150,8 @@ TEST_F(BufferTest, BufferLocationWithAnotherProp) {
149150
cl::sycl::ext::intel::property::buffer_location::instance<5>>>
150151
Acc{Buf, cgh, sycl::write_only, PL};
151152

152-
cgh.single_task<TestKernel>([=]() { Acc[0] = 4; });
153+
constexpr size_t KS = sizeof(decltype(Acc));
154+
cgh.single_task<TestKernel<KS>>([=]() { Acc[0] = 4; });
153155
})
154156
.wait();
155157
EXPECT_EQ(PassedLocation, (uint64_t)5);
@@ -209,7 +211,8 @@ TEST_F(BufferTest, WOBufferLocation) {
209211
cl::sycl::access::placeholder::false_t,
210212
cl::sycl::ext::oneapi::accessor_property_list<>>
211213
Acc{Buf, cgh, sycl::read_write};
212-
cgh.single_task<TestKernel>([=]() { Acc[0] = 4; });
214+
constexpr size_t KS = sizeof(decltype(Acc));
215+
cgh.single_task<TestKernel<KS>>([=]() { Acc[0] = 4; });
213216
})
214217
.wait();
215218
EXPECT_EQ(PassedLocation, DEFAULT_VALUE);

‎sycl/unittests/event/EventDestruction.cpp

Lines changed: 12 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -68,11 +68,11 @@ TEST_F(EventDestructionTest, EventDestruction) {
6868

6969
{
7070
sycl::event E0 = Queue.submit([&](cl::sycl::handler &cgh) {
71-
cgh.single_task<TestKernel>([]() {});
71+
cgh.single_task<TestKernel<>>([]() {});
7272
});
7373
E1 = Queue.submit([&](cl::sycl::handler &cgh) {
7474
cgh.depends_on(E0);
75-
cgh.single_task<TestKernel>([]() {});
75+
cgh.single_task<TestKernel<>>([]() {});
7676
});
7777
E1.wait();
7878
}
@@ -85,15 +85,15 @@ TEST_F(EventDestructionTest, EventDestruction) {
8585

8686
sycl::event E2 = Queue.submit([&](cl::sycl::handler &cgh) {
8787
cgh.depends_on(E1);
88-
cgh.single_task<TestKernel>([]() {});
88+
cgh.single_task<TestKernel<>>([]() {});
8989
});
9090
E2.wait();
9191
// Dependencies of E1 should be cleared here. It depends on E0.
9292
EXPECT_EQ(ReleaseCounter, 1);
9393

9494
sycl::event E3 = Queue.submit([&](cl::sycl::handler &cgh) {
9595
cgh.depends_on({E1, E2});
96-
cgh.single_task<TestKernel>([]() {});
96+
cgh.single_task<TestKernel<>>([]() {});
9797
});
9898
E3.wait();
9999
// Dependency of E1 has already cleared. E2 depends on E1 that
@@ -107,20 +107,20 @@ TEST_F(EventDestructionTest, EventDestruction) {
107107
sycl::buffer<int, 1> Buf(&data[0], sycl::range<1>(2));
108108
Queue.submit([&](cl::sycl::handler &cgh) {
109109
auto Acc = Buf.get_access<sycl::access::mode::read_write>(cgh);
110-
cgh.single_task<TestKernel>([=]() {});
110+
cgh.single_task<TestKernel<>>([=]() {});
111111
});
112112

113113
Queue.submit([&](cl::sycl::handler &cgh) {
114114
auto Acc = Buf.get_access<sycl::access::mode::read_write>(cgh);
115-
cgh.single_task<TestKernel>([=]() {});
115+
cgh.single_task<TestKernel<>>([=]() {});
116116
});
117117
sycl::event E1 = Queue.submit([&](cl::sycl::handler &cgh) {
118118
auto Acc = Buf.get_access<sycl::access::mode::read_write>(cgh);
119-
cgh.single_task<TestKernel>([=]() {});
119+
cgh.single_task<TestKernel<>>([=]() {});
120120
});
121121
sycl::event E2 = Queue.submit([&](cl::sycl::handler &cgh) {
122122
auto Acc = Buf.get_access<sycl::access::mode::read_write>(cgh);
123-
cgh.single_task<TestKernel>([=]() {});
123+
cgh.single_task<TestKernel<>>([=]() {});
124124
});
125125
E2.wait();
126126
// Dependencies are deleted through one level of dependencies. When
@@ -172,11 +172,11 @@ TEST_F(EventDestructionTest, GetWaitList) {
172172

173173
{
174174
sycl::event E0 = Queue.submit([&](cl::sycl::handler &cgh) {
175-
cgh.single_task<TestKernel>([]() {});
175+
cgh.single_task<TestKernel<>>([]() {});
176176
});
177177
E1 = Queue.submit([&](cl::sycl::handler &cgh) {
178178
cgh.depends_on(E0);
179-
cgh.single_task<TestKernel>([]() {});
179+
cgh.single_task<TestKernel<>>([]() {});
180180
});
181181
E1.wait();
182182
auto wait_list = E1.get_wait_list();
@@ -190,13 +190,13 @@ TEST_F(EventDestructionTest, GetWaitList) {
190190

191191
sycl::event E2 = Queue.submit([&](cl::sycl::handler &cgh) {
192192
cgh.depends_on(E1);
193-
cgh.single_task<TestKernel>([]() {});
193+
cgh.single_task<TestKernel<>>([]() {});
194194
});
195195
E2.wait();
196196

197197
sycl::event E3 = Queue.submit([&](cl::sycl::handler &cgh) {
198198
cgh.depends_on({E1, E2});
199-
cgh.single_task<TestKernel>([]() {});
199+
cgh.single_task<TestKernel<>>([]() {});
200200
});
201201
E3.wait();
202202

‎sycl/unittests/helpers/TestKernel.hpp

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -10,12 +10,12 @@
1010

1111
#include "PiImage.hpp"
1212

13-
class TestKernel;
13+
template <size_t KernelSize = 1> class TestKernel;
1414

1515
__SYCL_INLINE_NAMESPACE(cl) {
1616
namespace sycl {
1717
namespace detail {
18-
template <> struct KernelInfo<TestKernel> {
18+
template <size_t KernelSize> struct KernelInfo<TestKernel<KernelSize>> {
1919
static constexpr unsigned getNumParams() { return 0; }
2020
static const kernel_param_desc_t &getParamDesc(int) {
2121
static kernel_param_desc_t Dummy;
@@ -25,6 +25,7 @@ template <> struct KernelInfo<TestKernel> {
2525
static constexpr bool isESIMD() { return false; }
2626
static constexpr bool callsThisItem() { return false; }
2727
static constexpr bool callsAnyThisFreeFunction() { return false; }
28+
static constexpr int64_t getKernelSize() { return KernelSize; }
2829
};
2930

3031
} // namespace detail

‎sycl/unittests/kernel-and-program/Cache.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -46,6 +46,7 @@ struct MockKernelInfo {
4646
static constexpr bool isESIMD() { return false; }
4747
static constexpr bool callsThisItem() { return false; }
4848
static constexpr bool callsAnyThisFreeFunction() { return false; }
49+
static constexpr int64_t getKernelSize() { return 1; }
4950
};
5051

5152
template <> struct KernelInfo<TestKernel> : public MockKernelInfo {

‎sycl/unittests/kernel-and-program/MultipleDevsCache.cpp

Lines changed: 4 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -151,11 +151,12 @@ TEST_F(MultipleDeviceCacheTest, ProgramRetain) {
151151

152152
auto Bundle = cl::sycl::get_kernel_bundle<sycl::bundle_state::input>(
153153
Queue.get_context());
154-
Queue.submit(
155-
[&](cl::sycl::handler &cgh) { cgh.single_task<TestKernel>([]() {}); });
154+
Queue.submit([&](cl::sycl::handler &cgh) {
155+
cgh.single_task<TestKernel<>>([]() {});
156+
});
156157

157158
auto BundleObject = cl::sycl::build(Bundle, Bundle.get_devices());
158-
auto KernelID = cl::sycl::get_kernel_id<TestKernel>();
159+
auto KernelID = cl::sycl::get_kernel_id<TestKernel<>>();
159160
auto Kernel = BundleObject.get_kernel(KernelID);
160161

161162
// Because of emulating 2 devices program is retained for each one in

‎sycl/unittests/misc/KernelBuildOptions.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -33,6 +33,7 @@ template <> struct KernelInfo<TestKernel> {
3333
static constexpr bool isESIMD() { return true; }
3434
static constexpr bool callsThisItem() { return false; }
3535
static constexpr bool callsAnyThisFreeFunction() { return false; }
36+
static constexpr int64_t getKernelSize() { return 1; }
3637
};
3738

3839
} // namespace detail

‎sycl/unittests/program_manager/EliminatedArgMask.cpp

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -20,8 +20,8 @@
2020

2121
class EAMTestKernel;
2222
class EAMTestKernel2;
23-
const char EAMTestKernelName[] = "EAMTestKernel";
24-
const char EAMTestKernel2Name[] = "EAMTestKernel2";
23+
constexpr const char EAMTestKernelName[] = "EAMTestKernel";
24+
constexpr const char EAMTestKernel2Name[] = "EAMTestKernel2";
2525
constexpr unsigned EAMTestKernelNumArgs = 4;
2626

2727
__SYCL_INLINE_NAMESPACE(cl) {
@@ -37,6 +37,7 @@ template <> struct KernelInfo<EAMTestKernel> {
3737
static constexpr bool isESIMD() { return false; }
3838
static constexpr bool callsThisItem() { return false; }
3939
static constexpr bool callsAnyThisFreeFunction() { return false; }
40+
static constexpr int64_t getKernelSize() { return 1; }
4041
};
4142

4243
template <> struct KernelInfo<EAMTestKernel2> {
@@ -49,6 +50,7 @@ template <> struct KernelInfo<EAMTestKernel2> {
4950
static constexpr bool isESIMD() { return false; }
5051
static constexpr bool callsThisItem() { return false; }
5152
static constexpr bool callsAnyThisFreeFunction() { return false; }
53+
static constexpr int64_t getKernelSize() { return 1; }
5254
};
5355

5456
} // namespace detail

0 commit comments

Comments
 (0)