|
| 1 | +//==----------- KernelArgMemObj.cpp ---- Scheduler unit tests ---------- ---==// |
| 2 | +// |
| 3 | +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. |
| 4 | +// See https://llvm.org/LICENSE.txt for license information. |
| 5 | +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception |
| 6 | +// |
| 7 | +//===----------------------------------------------------------------------===// |
| 8 | + |
| 9 | +#include <gtest/gtest.h> |
| 10 | + |
| 11 | +#include <helpers/MockKernelInfo.hpp> |
| 12 | +#include <helpers/PiImage.hpp> |
| 13 | +#include <helpers/PiMock.hpp> |
| 14 | + |
| 15 | +class TestKernelWithMemObj; |
| 16 | + |
| 17 | +namespace sycl { |
| 18 | +__SYCL_INLINE_VER_NAMESPACE(_V1) { |
| 19 | +namespace detail { |
| 20 | +template <> |
| 21 | +struct KernelInfo<TestKernelWithMemObj> : public unittest::MockKernelInfoBase { |
| 22 | + static constexpr const char *getName() { return "TestKernelWithMemObj"; } |
| 23 | + static constexpr unsigned getNumParams() { return 1; } |
| 24 | + static const detail::kernel_param_desc_t &getParamDesc(int) { |
| 25 | + static detail::kernel_param_desc_t desc{ |
| 26 | + detail::kernel_param_kind_t::kind_accessor, |
| 27 | + int(access::target::device) /*info*/, 0 /*offset*/}; |
| 28 | + return desc; |
| 29 | + } |
| 30 | + static constexpr uint32_t getKernelSize() { return 32; } |
| 31 | +}; |
| 32 | +} // namespace detail |
| 33 | +} // __SYCL_INLINE_VER_NAMESPACE(_V1) |
| 34 | +} // namespace sycl |
| 35 | + |
| 36 | +static sycl::unittest::PiImage generateImage() { |
| 37 | + using namespace sycl::unittest; |
| 38 | + |
| 39 | + PiPropertySet PropSet; |
| 40 | + |
| 41 | + std::vector<unsigned char> Bin{0, 1, 2, 3, 4, 5}; // Random data |
| 42 | + |
| 43 | + PiArray<PiOffloadEntry> Entries = makeEmptyKernels({"TestKernelWithMemObj"}); |
| 44 | + |
| 45 | + PiImage Img{PI_DEVICE_BINARY_TYPE_SPIRV, // Format |
| 46 | + __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64, // DeviceTargetSpec |
| 47 | + "", // Compile options |
| 48 | + "", // Link options |
| 49 | + std::move(Bin), |
| 50 | + std::move(Entries), |
| 51 | + std::move(PropSet)}; |
| 52 | + |
| 53 | + return Img; |
| 54 | +} |
| 55 | + |
| 56 | +static sycl::unittest::PiImage Img = generateImage(); |
| 57 | +static sycl::unittest::PiImageArray<1> ImgArray{&Img}; |
| 58 | + |
| 59 | +using namespace sycl; |
| 60 | + |
| 61 | +bool PropertyPresent = false; |
| 62 | +pi_mem_obj_property PropsCopy{}; |
| 63 | + |
| 64 | +pi_result redefinedKernelSetArgMemObj(pi_kernel kernel, pi_uint32 arg_index, |
| 65 | + const pi_mem_obj_property *arg_properties, |
| 66 | + const pi_mem *arg_value) { |
| 67 | + PropertyPresent = arg_properties != nullptr; |
| 68 | + if (PropertyPresent) |
| 69 | + PropsCopy = *arg_properties; |
| 70 | + return PI_SUCCESS; |
| 71 | +} |
| 72 | + |
| 73 | +class BufferTestPiArgs : public ::testing::Test { |
| 74 | +public: |
| 75 | + BufferTestPiArgs() |
| 76 | + : Mock(sycl::backend::ext_oneapi_level_zero), Plt{Mock.getPlatform()} {} |
| 77 | + |
| 78 | +protected: |
| 79 | + void SetUp() override { |
| 80 | + PropertyPresent = false; |
| 81 | + PropsCopy = {}; |
| 82 | + Mock.redefineBefore<detail::PiApiKind::piextKernelSetArgMemObj>( |
| 83 | + redefinedKernelSetArgMemObj); |
| 84 | + } |
| 85 | + |
| 86 | + template <sycl::access::mode AccessMode> |
| 87 | + void TestFunc(pi_mem_obj_access ExpectedAccessMode) { |
| 88 | + queue Queue(context(Plt), default_selector_v); |
| 89 | + sycl::buffer<int, 1> Buf(3); |
| 90 | + Queue |
| 91 | + .submit([&](sycl::handler &cgh) { |
| 92 | + auto acc = Buf.get_access<AccessMode>(cgh); |
| 93 | + cgh.single_task<TestKernelWithMemObj>([=]() { |
| 94 | + if constexpr (AccessMode != sycl::access::mode::read) |
| 95 | + acc[0] = 4; |
| 96 | + else |
| 97 | + std::ignore = acc[0]; |
| 98 | + }); |
| 99 | + }) |
| 100 | + .wait(); |
| 101 | + ASSERT_TRUE(PropertyPresent); |
| 102 | + EXPECT_EQ(PropsCopy.type, PI_KERNEL_ARG_MEM_OBJ_ACCESS); |
| 103 | + EXPECT_EQ(PropsCopy.mem_access, ExpectedAccessMode); |
| 104 | + } |
| 105 | + |
| 106 | +protected: |
| 107 | + sycl::unittest::PiMock Mock; |
| 108 | + sycl::platform Plt; |
| 109 | +}; |
| 110 | + |
| 111 | +TEST_F(BufferTestPiArgs, KernelSetArgMemObjReadWrite) { |
| 112 | + TestFunc<sycl::access::mode::read_write>(PI_ACCESS_READ_WRITE); |
| 113 | +} |
| 114 | + |
| 115 | +TEST_F(BufferTestPiArgs, KernelSetArgMemObjDiscardReadWrite) { |
| 116 | + TestFunc<sycl::access::mode::discard_read_write>(PI_ACCESS_READ_WRITE); |
| 117 | +} |
| 118 | + |
| 119 | +TEST_F(BufferTestPiArgs, KernelSetArgMemObjRead) { |
| 120 | + TestFunc<sycl::access::mode::read>(PI_ACCESS_READ_ONLY); |
| 121 | +} |
| 122 | + |
| 123 | +TEST_F(BufferTestPiArgs, KernelSetArgMemObjWrite) { |
| 124 | + TestFunc<sycl::access::mode::write>(PI_ACCESS_WRITE_ONLY); |
| 125 | +} |
| 126 | + |
| 127 | +TEST_F(BufferTestPiArgs, KernelSetArgMemObjDiscardWrite) { |
| 128 | + TestFunc<sycl::access::mode::discard_write>(PI_ACCESS_WRITE_ONLY); |
| 129 | +} |
0 commit comments