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);

0 commit comments

Comments
 (0)