diff --git a/sycl/source/stream.cpp b/sycl/source/stream.cpp index 6d000751e0354..db196c14c02ee 100644 --- a/sycl/source/stream.cpp +++ b/sycl/source/stream.cpp @@ -17,21 +17,28 @@ namespace sycl { static constexpr size_t MAX_STATEMENT_SIZE = (1 << (CHAR_BIT * detail::FLUSH_BUF_OFFSET_SIZE)) - 1; -stream::stream(size_t BufferSize, size_t MaxStatementSize, handler &CGH) - : impl(std::make_shared(BufferSize, MaxStatementSize, - CGH)), - GlobalBuf(impl->accessGlobalBuf(CGH)), - GlobalOffset(impl->accessGlobalOffset(CGH)), - // Allocate the flush buffer, which contains space for each work item - GlobalFlushBuf(impl->accessGlobalFlushBuf(CGH)), - FlushBufferSize(MaxStatementSize + detail::FLUSH_BUF_OFFSET_SIZE) { +// Checks the MaxStatementSize argument of the sycl::stream class. This is +// called on MaxStatementSize as it is passed to the constructor of the +// underlying stream_impl to make it throw before the stream buffers are +// allocated, avoiding memory leaks. +static size_t CheckMaxStatementSize(const size_t &MaxStatementSize) { if (MaxStatementSize > MAX_STATEMENT_SIZE) { throw sycl::invalid_parameter_error( "Maximum statement size exceeds limit of " + std::to_string(MAX_STATEMENT_SIZE) + " bytes.", PI_INVALID_VALUE); } + return MaxStatementSize; +} +stream::stream(size_t BufferSize, size_t MaxStatementSize, handler &CGH) + : impl(std::make_shared( + BufferSize, CheckMaxStatementSize(MaxStatementSize), CGH)), + GlobalBuf(impl->accessGlobalBuf(CGH)), + GlobalOffset(impl->accessGlobalOffset(CGH)), + // Allocate the flush buffer, which contains space for each work item + GlobalFlushBuf(impl->accessGlobalFlushBuf(CGH)), + FlushBufferSize(MaxStatementSize + detail::FLUSH_BUF_OFFSET_SIZE) { // Save stream implementation in the handler so that stream will be alive // during kernel execution CGH.addStream(impl); diff --git a/sycl/unittests/CMakeLists.txt b/sycl/unittests/CMakeLists.txt index 4bd4e70c5afb7..3efd330c112c5 100644 --- a/sycl/unittests/CMakeLists.txt +++ b/sycl/unittests/CMakeLists.txt @@ -18,6 +18,7 @@ add_subdirectory(pi) add_subdirectory(kernel-and-program) add_subdirectory(queue) add_subdirectory(scheduler) +add_subdirectory(stream) add_subdirectory(SYCL2020) add_subdirectory(thread_safety) add_subdirectory(program_manager) diff --git a/sycl/unittests/stream/CMakeLists.txt b/sycl/unittests/stream/CMakeLists.txt new file mode 100644 index 0000000000000..ae8a544262c87 --- /dev/null +++ b/sycl/unittests/stream/CMakeLists.txt @@ -0,0 +1,3 @@ +add_sycl_unittest(StreamTests OBJECT + stream.cpp +) diff --git a/sycl/unittests/stream/stream.cpp b/sycl/unittests/stream/stream.cpp new file mode 100644 index 0000000000000..557af4f6989a1 --- /dev/null +++ b/sycl/unittests/stream/stream.cpp @@ -0,0 +1,121 @@ +//==---------------- stream.cpp --- SYCL stream unit test ------------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include + +#include +#include +#include + +#include + +#include + +class TestKernel; + +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { +namespace detail { +template <> struct KernelInfo { + static constexpr unsigned getNumParams() { return 0; } + static const kernel_param_desc_t &getParamDesc(int) { + static kernel_param_desc_t Dummy; + return Dummy; + } + static constexpr const char *getName() { return "Stream_TestKernel"; } + static constexpr bool isESIMD() { return false; } + static constexpr bool callsThisItem() { return false; } + static constexpr bool callsAnyThisFreeFunction() { return false; } +}; +} // namespace detail +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) + +static sycl::unittest::PiImage generateDefaultImage() { + using namespace sycl::unittest; + + PiPropertySet PropSet; + + std::vector Bin{0, 1, 2, 3, 4, 5}; // Random data + + PiArray Entries = makeEmptyKernels({"Stream_TestKernel"}); + + PiImage Img{PI_DEVICE_BINARY_TYPE_SPIRV, // Format + __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64, // DeviceTargetSpec + "", // Compile options + "", // Link options + std::move(Bin), + std::move(Entries), + std::move(PropSet)}; + + return Img; +} + +static sycl::unittest::PiImage Img = generateDefaultImage(); +static sycl::unittest::PiImageArray<1> ImgArray{&Img}; + +size_t GBufferCreateCounter = 0; + +static pi_result +redefinedMemBufferCreate(pi_context context, pi_mem_flags flags, size_t size, + void *host_ptr, pi_mem *ret_mem, + const pi_mem_properties *properties = nullptr) { + ++GBufferCreateCounter; + *ret_mem = nullptr; + return PI_SUCCESS; +} + +TEST(Stream, TestStreamConstructorExceptionNoAllocation) { + sycl::platform Plt{sycl::default_selector()}; + if (Plt.is_host()) { + std::cout << "Not run on host - no PI buffers created in that case" + << std::endl; + return; + } + + if (Plt.get_backend() == sycl::backend::cuda) { + std::cout << "Test is not supported on CUDA platform, skipping\n"; + return; + } + + if (Plt.get_backend() == sycl::backend::hip) { + std::cout << "Test is not supported on HIP platform, skipping\n"; + return; + } + + sycl::unittest::PiMock Mock{Plt}; + setupDefaultMockAPIs(Mock); + Mock.redefine( + redefinedMemBufferCreate); + + const sycl::device Dev = Plt.get_devices()[0]; + sycl::queue Queue{Dev}; + const sycl::context Ctx = Queue.get_context(); + + sycl::kernel_bundle KernelBundle = + sycl::get_kernel_bundle(Ctx, {Dev}); + auto ExecBundle = sycl::build(KernelBundle); + + Queue.submit([&](sycl::handler &CGH) { + CGH.use_kernel_bundle(ExecBundle); + + try { + // Try to create stream with invalid workItemBufferSize parameter. + sycl::stream InvalidStream{256, std::numeric_limits::max(), CGH}; + FAIL() << "No exception was thrown."; + } catch (const sycl::invalid_parameter_error &) { + // Expected exception + } catch (...) { + FAIL() << "Unexpected exception was thrown."; + } + + CGH.single_task([=]() {}); + }); + + ASSERT_EQ(GBufferCreateCounter, 0u) << "Buffers were unexpectedly created."; +}