Skip to content

[SYCL] Fail on kernel lambda size mismatch #6374

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 3 commits into from
Jul 8, 2022
Merged

[SYCL] Fail on kernel lambda size mismatch #6374

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
e8adf0a
[SYCL] Fail on kernel lambda size mismatch
steffenlarsen Jun 23, 2022
8166302
Fix formatting
steffenlarsen Jun 29, 2022
8a1c97c
Make return type int64_t
steffenlarsen Jun 29, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
12 changes: 12 additions & 0 deletions clang/test/SemaSYCL/Inputs/CL/sycl/detail/kernel_desc.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -35,6 +35,17 @@ __SYCL_INLINE_NAMESPACE(cl) {
int offset;
};

template <bool Cond, typename TrueT, typename FalseT>
struct conditional {
using type = TrueT;
};
template <typename TrueT, typename FalseT>
struct conditional<false, TrueT, FalseT> {
using type = FalseT;
};

using int64_t = conditional<sizeof(long) == 8, long, long long>::type;
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Ideally we should check that sizeof(long long) is 8, if sizeof(long) != 8.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I agree, but this mimics the decision made by SemaSYCL when generating the specialization, so it should be fine for our testing.


template <class KernelNameType> struct KernelInfo {
static constexpr unsigned getNumParams() { return 0; }
static const kernel_param_desc_t &getParamDesc(int) {
Expand All @@ -43,6 +54,7 @@ __SYCL_INLINE_NAMESPACE(cl) {
}
static constexpr const char *getName() { return ""; }
static constexpr bool isESIMD() { return 0; }
static constexpr int64_t getKernelSize() { return 0; }
};
} // namespace detail
} // namespace sycl
Expand Down
7 changes: 7 additions & 0 deletions sycl/include/CL/sycl/detail/kernel_desc.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,6 +13,8 @@
#include <CL/sycl/detail/defines_elementary.hpp>
#include <CL/sycl/detail/export.hpp>

#include <cstdint>
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Nice, perhaps this will also fix one of internal trackers reporting problems with undefined int8_t type slipped into integration-header file.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This inclusion is giving oneapi/dpl a stomach ache.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

because of this: https://github.com/steffenlarsen/llvm/blob/8a1c97c5705cd904795ce0e5a8856ac0e9ac7010/sycl/include/CL/sycl/detail/kernel_desc.hpp#L11


__SYCL_INLINE_NAMESPACE(cl) {
namespace sycl {
namespace detail {
Expand Down Expand Up @@ -79,6 +81,7 @@ template <class KernelNameType> struct KernelInfo {
static constexpr const char *getFunctionName() { return ""; }
static constexpr unsigned getLineNumber() { return 0; }
static constexpr unsigned getColumnNumber() { return 0; }
static constexpr int64_t getKernelSize() { return 0; }
};
#else
template <char...> struct KernelInfoData {
Expand All @@ -93,6 +96,7 @@ template <char...> struct KernelInfoData {
static constexpr const char *getFunctionName() { return ""; }
static constexpr unsigned getLineNumber() { return 0; }
static constexpr unsigned getColumnNumber() { return 0; }
static constexpr int64_t getKernelSize() { return 0; }
};

// C++14 like index_sequence and make_index_sequence
Expand Down Expand Up @@ -135,6 +139,9 @@ template <class KernelNameType> struct KernelInfo {
static constexpr const char *getFunctionName() { return ""; }
static constexpr unsigned getLineNumber() { return 0; }
static constexpr unsigned getColumnNumber() { return 0; }
static constexpr int64_t getKernelSize() {
return SubKernelInfo::getKernelSize();
}
};
#endif //__SYCL_UNNAMED_LAMBDA__

Expand Down
14 changes: 13 additions & 1 deletion sycl/include/CL/sycl/handler.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -717,13 +717,25 @@ class __SYCL_EXPORT handler {
"kernel_handler is not yet supported by host device.",
PI_ERROR_INVALID_OPERATION);
}

KernelType *KernelPtr =
ResetHostKernel<KernelType, LambdaArgType, Dims>(KernelFunc);

using KI = sycl::detail::KernelInfo<KernelName>;
constexpr bool KernelHasName =
KI::getName() != nullptr && KI::getName()[0] != '\0';

// Some host compilers may have different captures from Clang. Currently
// there is no stable way of handling this when extracting the captures, so
// a static assert is made to fail for incompatible kernel lambdas.
static_assert(!KernelHasName || sizeof(KernelFunc) == KI::getKernelSize(),
"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.");

// Empty name indicates that the compilation happens without integration
// header, so don't perform things that require it.
if (KI::getName() != nullptr && KI::getName()[0] != '\0') {
if (KernelHasName) {
// TODO support ESIMD in no-integration-header case too.
MArgs.clear();
extractArgsAndReqsFromLambda(reinterpret_cast<char *>(KernelPtr),
Expand Down
19 changes: 19 additions & 0 deletions sycl/test/basic_tests/kernel_size_mismatch.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,19 @@
// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning -o - %s

// Tests for static assertion failure when kernel lambda mismatches between host
// and device.

#include <CL/sycl.hpp>

int main() {
sycl::queue Q;
int A = 1;
Q.single_task([=]() {
#ifdef __SYCL_DEVICE_ONLY__
(void)A;
// expected-no-diagnostics
#else
// 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."}}
#endif
}).wait();
}
3 changes: 2 additions & 1 deletion sycl/unittests/SYCL2020/GetNativeOpenCL.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -127,7 +127,8 @@ TEST(GetNative, GetNativeHandle) {
sycl::buffer<int, 1> Buffer(&Data[0], sycl::range<1>(1));
Queue.submit([&](sycl::handler &cgh) {
auto Acc = Buffer.get_access<sycl::access::mode::read_write>(cgh);
cgh.single_task<TestKernel>([=]() { (void)Acc; });
constexpr size_t KS = sizeof(decltype(Acc));
cgh.single_task<TestKernel<KS>>([=]() { (void)Acc; });
});

get_native<backend::opencl>(Context);
Expand Down
2 changes: 2 additions & 0 deletions sycl/unittests/SYCL2020/KernelBundle.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -31,6 +31,7 @@ template <> struct KernelInfo<TestKernel> {
static constexpr bool isESIMD() { return false; }
static constexpr bool callsThisItem() { return false; }
static constexpr bool callsAnyThisFreeFunction() { return false; }
static constexpr int64_t getKernelSize() { return 1; }
};

template <> struct KernelInfo<TestKernelExeOnly> {
Expand All @@ -43,6 +44,7 @@ template <> struct KernelInfo<TestKernelExeOnly> {
static constexpr bool isESIMD() { return false; }
static constexpr bool callsThisItem() { return false; }
static constexpr bool callsAnyThisFreeFunction() { return false; }
static constexpr int64_t getKernelSize() { return 1; }
};

} // namespace detail
Expand Down
4 changes: 4 additions & 0 deletions sycl/unittests/SYCL2020/KernelID.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -32,6 +32,7 @@ template <> struct KernelInfo<TestKernel1> {
static constexpr bool isESIMD() { return false; }
static constexpr bool callsThisItem() { return false; }
static constexpr bool callsAnyThisFreeFunction() { return false; }
static constexpr int64_t getKernelSize() { return 1; }
};

template <> struct KernelInfo<TestKernel2> {
Expand All @@ -44,6 +45,7 @@ template <> struct KernelInfo<TestKernel2> {
static constexpr bool isESIMD() { return false; }
static constexpr bool callsThisItem() { return false; }
static constexpr bool callsAnyThisFreeFunction() { return false; }
static constexpr int64_t getKernelSize() { return 1; }
};

template <> struct KernelInfo<TestKernel3> {
Expand All @@ -56,6 +58,7 @@ template <> struct KernelInfo<TestKernel3> {
static constexpr bool isESIMD() { return false; }
static constexpr bool callsThisItem() { return false; }
static constexpr bool callsAnyThisFreeFunction() { return false; }
static constexpr int64_t getKernelSize() { return 1; }
};

template <> struct KernelInfo<ServiceKernel1> {
Expand All @@ -70,6 +73,7 @@ template <> struct KernelInfo<ServiceKernel1> {
static constexpr bool isESIMD() { return false; }
static constexpr bool callsThisItem() { return false; }
static constexpr bool callsAnyThisFreeFunction() { return false; }
static constexpr int64_t getKernelSize() { return 1; }
};
} // namespace detail
} // namespace sycl
Expand Down
1 change: 1 addition & 0 deletions sycl/unittests/SYCL2020/SpecializationConstant.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -35,6 +35,7 @@ template <> struct KernelInfo<TestKernel> {
static constexpr bool isESIMD() { return false; }
static constexpr bool callsThisItem() { return false; }
static constexpr bool callsAnyThisFreeFunction() { return false; }
static constexpr int64_t getKernelSize() { return 1; }
};

template <> const char *get_spec_constant_symbolic_ID<SpecConst1>() {
Expand Down
6 changes: 6 additions & 0 deletions sycl/unittests/assert/assert.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -50,6 +50,7 @@ template <> struct KernelInfo<TestKernel> {
static constexpr bool isESIMD() { return false; }
static constexpr bool callsThisItem() { return false; }
static constexpr bool callsAnyThisFreeFunction() { return false; }
static constexpr int64_t getKernelSize() { return 1; }
};

static constexpr const kernel_param_desc_t Signatures[] = {
Expand All @@ -68,6 +69,11 @@ struct KernelInfo<::sycl::detail::__sycl_service_kernel__::AssertInfoCopier> {
static constexpr bool isESIMD() { return 0; }
static constexpr bool callsThisItem() { return 0; }
static constexpr bool callsAnyThisFreeFunction() { return 0; }
static constexpr int64_t getKernelSize() {
// The AssertInfoCopier service kernel lambda captures an accessor.
return sizeof(sycl::accessor<sycl::detail::AssertHappened, 1,
sycl::access::mode::write>);
}
};
} // namespace detail
} // namespace sycl
Expand Down
9 changes: 6 additions & 3 deletions sycl/unittests/buffer/BufferLocation.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -116,7 +116,8 @@ TEST_F(BufferTest, BufferLocationOnly) {
cl::sycl::ext::oneapi::accessor_property_list<
cl::sycl::ext::intel::property::buffer_location::instance<2>>>
Acc{Buf, cgh, sycl::read_write, PL};
cgh.single_task<TestKernel>([=]() { Acc[0] = 4; });
constexpr size_t KS = sizeof(decltype(Acc));
cgh.single_task<TestKernel<KS>>([=]() { Acc[0] = 4; });
})
.wait();
EXPECT_EQ(PassedLocation, (uint64_t)2);
Expand Down Expand Up @@ -149,7 +150,8 @@ TEST_F(BufferTest, BufferLocationWithAnotherProp) {
cl::sycl::ext::intel::property::buffer_location::instance<5>>>
Acc{Buf, cgh, sycl::write_only, PL};

cgh.single_task<TestKernel>([=]() { Acc[0] = 4; });
constexpr size_t KS = sizeof(decltype(Acc));
cgh.single_task<TestKernel<KS>>([=]() { Acc[0] = 4; });
})
.wait();
EXPECT_EQ(PassedLocation, (uint64_t)5);
Expand Down Expand Up @@ -209,7 +211,8 @@ TEST_F(BufferTest, WOBufferLocation) {
cl::sycl::access::placeholder::false_t,
cl::sycl::ext::oneapi::accessor_property_list<>>
Acc{Buf, cgh, sycl::read_write};
cgh.single_task<TestKernel>([=]() { Acc[0] = 4; });
constexpr size_t KS = sizeof(decltype(Acc));
cgh.single_task<TestKernel<KS>>([=]() { Acc[0] = 4; });
})
.wait();
EXPECT_EQ(PassedLocation, DEFAULT_VALUE);
Expand Down
24 changes: 12 additions & 12 deletions sycl/unittests/event/EventDestruction.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -68,11 +68,11 @@ TEST_F(EventDestructionTest, EventDestruction) {

{
sycl::event E0 = Queue.submit([&](cl::sycl::handler &cgh) {
cgh.single_task<TestKernel>([]() {});
cgh.single_task<TestKernel<>>([]() {});
});
E1 = Queue.submit([&](cl::sycl::handler &cgh) {
cgh.depends_on(E0);
cgh.single_task<TestKernel>([]() {});
cgh.single_task<TestKernel<>>([]() {});
});
E1.wait();
}
Expand All @@ -85,15 +85,15 @@ TEST_F(EventDestructionTest, EventDestruction) {

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

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

Queue.submit([&](cl::sycl::handler &cgh) {
auto Acc = Buf.get_access<sycl::access::mode::read_write>(cgh);
cgh.single_task<TestKernel>([=]() {});
cgh.single_task<TestKernel<>>([=]() {});
});
sycl::event E1 = Queue.submit([&](cl::sycl::handler &cgh) {
auto Acc = Buf.get_access<sycl::access::mode::read_write>(cgh);
cgh.single_task<TestKernel>([=]() {});
cgh.single_task<TestKernel<>>([=]() {});
});
sycl::event E2 = Queue.submit([&](cl::sycl::handler &cgh) {
auto Acc = Buf.get_access<sycl::access::mode::read_write>(cgh);
cgh.single_task<TestKernel>([=]() {});
cgh.single_task<TestKernel<>>([=]() {});
});
E2.wait();
// Dependencies are deleted through one level of dependencies. When
Expand Down Expand Up @@ -172,11 +172,11 @@ TEST_F(EventDestructionTest, GetWaitList) {

{
sycl::event E0 = Queue.submit([&](cl::sycl::handler &cgh) {
cgh.single_task<TestKernel>([]() {});
cgh.single_task<TestKernel<>>([]() {});
});
E1 = Queue.submit([&](cl::sycl::handler &cgh) {
cgh.depends_on(E0);
cgh.single_task<TestKernel>([]() {});
cgh.single_task<TestKernel<>>([]() {});
});
E1.wait();
auto wait_list = E1.get_wait_list();
Expand All @@ -190,13 +190,13 @@ TEST_F(EventDestructionTest, GetWaitList) {

sycl::event E2 = Queue.submit([&](cl::sycl::handler &cgh) {
cgh.depends_on(E1);
cgh.single_task<TestKernel>([]() {});
cgh.single_task<TestKernel<>>([]() {});
});
E2.wait();

sycl::event E3 = Queue.submit([&](cl::sycl::handler &cgh) {
cgh.depends_on({E1, E2});
cgh.single_task<TestKernel>([]() {});
cgh.single_task<TestKernel<>>([]() {});
});
E3.wait();

Expand Down
5 changes: 3 additions & 2 deletions sycl/unittests/helpers/TestKernel.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,12 +10,12 @@

#include "PiImage.hpp"

class TestKernel;
template <size_t KernelSize = 1> class TestKernel;

__SYCL_INLINE_NAMESPACE(cl) {
namespace sycl {
namespace detail {
template <> struct KernelInfo<TestKernel> {
template <size_t KernelSize> struct KernelInfo<TestKernel<KernelSize>> {
static constexpr unsigned getNumParams() { return 0; }
static const kernel_param_desc_t &getParamDesc(int) {
static kernel_param_desc_t Dummy;
Expand All @@ -25,6 +25,7 @@ template <> struct KernelInfo<TestKernel> {
static constexpr bool isESIMD() { return false; }
static constexpr bool callsThisItem() { return false; }
static constexpr bool callsAnyThisFreeFunction() { return false; }
static constexpr int64_t getKernelSize() { return KernelSize; }
};

} // namespace detail
Expand Down
1 change: 1 addition & 0 deletions sycl/unittests/kernel-and-program/Cache.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -46,6 +46,7 @@ struct MockKernelInfo {
static constexpr bool isESIMD() { return false; }
static constexpr bool callsThisItem() { return false; }
static constexpr bool callsAnyThisFreeFunction() { return false; }
static constexpr int64_t getKernelSize() { return 1; }
};

template <> struct KernelInfo<TestKernel> : public MockKernelInfo {
Expand Down
7 changes: 4 additions & 3 deletions sycl/unittests/kernel-and-program/MultipleDevsCache.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -151,11 +151,12 @@ TEST_F(MultipleDeviceCacheTest, ProgramRetain) {

auto Bundle = cl::sycl::get_kernel_bundle<sycl::bundle_state::input>(
Queue.get_context());
Queue.submit(
[&](cl::sycl::handler &cgh) { cgh.single_task<TestKernel>([]() {}); });
Queue.submit([&](cl::sycl::handler &cgh) {
cgh.single_task<TestKernel<>>([]() {});
});

auto BundleObject = cl::sycl::build(Bundle, Bundle.get_devices());
auto KernelID = cl::sycl::get_kernel_id<TestKernel>();
auto KernelID = cl::sycl::get_kernel_id<TestKernel<>>();
auto Kernel = BundleObject.get_kernel(KernelID);

// Because of emulating 2 devices program is retained for each one in
Expand Down
1 change: 1 addition & 0 deletions sycl/unittests/misc/KernelBuildOptions.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -33,6 +33,7 @@ template <> struct KernelInfo<TestKernel> {
static constexpr bool isESIMD() { return true; }
static constexpr bool callsThisItem() { return false; }
static constexpr bool callsAnyThisFreeFunction() { return false; }
static constexpr int64_t getKernelSize() { return 1; }
};

} // namespace detail
Expand Down
6 changes: 4 additions & 2 deletions sycl/unittests/program_manager/EliminatedArgMask.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -20,8 +20,8 @@

class EAMTestKernel;
class EAMTestKernel2;
const char EAMTestKernelName[] = "EAMTestKernel";
const char EAMTestKernel2Name[] = "EAMTestKernel2";
constexpr const char EAMTestKernelName[] = "EAMTestKernel";
constexpr const char EAMTestKernel2Name[] = "EAMTestKernel2";
constexpr unsigned EAMTestKernelNumArgs = 4;

__SYCL_INLINE_NAMESPACE(cl) {
Expand All @@ -37,6 +37,7 @@ template <> struct KernelInfo<EAMTestKernel> {
static constexpr bool isESIMD() { return false; }
static constexpr bool callsThisItem() { return false; }
static constexpr bool callsAnyThisFreeFunction() { return false; }
static constexpr int64_t getKernelSize() { return 1; }
};

template <> struct KernelInfo<EAMTestKernel2> {
Expand All @@ -49,6 +50,7 @@ template <> struct KernelInfo<EAMTestKernel2> {
static constexpr bool isESIMD() { return false; }
static constexpr bool callsThisItem() { return false; }
static constexpr bool callsAnyThisFreeFunction() { return false; }
static constexpr int64_t getKernelSize() { return 1; }
};

} // namespace detail
Expand Down
Loading