diff --git a/sycl/include/CL/sycl/detail/cg.hpp b/sycl/include/CL/sycl/detail/cg.hpp index a87daa3e8e154..4b63d9c3289ee 100644 --- a/sycl/include/CL/sycl/detail/cg.hpp +++ b/sycl/include/CL/sycl/detail/cg.hpp @@ -208,7 +208,7 @@ class HostKernel : public HostKernelBase { // If local size for host is not set explicitly, let's adjust it to 1, // so nd_range_error for zero local size is not thrown. if (AdjustedRange.LocalSize[0] == 0) - for (int I = 0; I < AdjustedRange.Dims; ++I) + for (size_t I = 0; I < AdjustedRange.Dims; ++I) AdjustedRange.LocalSize[I] = 1; if (HPI) HPI->start(); diff --git a/sycl/include/CL/sycl/detail/image_accessor_util.hpp b/sycl/include/CL/sycl/detail/image_accessor_util.hpp index f0a4cefbe43b0..cea5845775e97 100644 --- a/sycl/include/CL/sycl/detail/image_accessor_util.hpp +++ b/sycl/include/CL/sycl/detail/image_accessor_util.hpp @@ -530,13 +530,20 @@ convertWriteData(const vec WriteData, } } +template struct pixel_type_helper { using type = T; }; + +template <> struct pixel_type_helper { using type = int32_t; }; + +template <> struct pixel_type_helper { using type = int64_t; }; + template vec processFloatDataToPixel(vec WriteData, float MulFactor) { vec Temp = WriteData * MulFactor; vec TempInInt = Temp.convert(); - vec TempInIntSaturated = - cl::sycl::clamp(TempInInt, min_v(), max_v()); + vec TempInIntSaturated = cl::sycl::clamp( + TempInInt, min_v::type>(), + max_v::type>()); return TempInIntSaturated.convert(); } @@ -545,8 +552,6 @@ vec convertWriteData(const vec WriteData, const image_channel_type ImageChannelType) { - vec PixelData; - switch (ImageChannelType) { case image_channel_type::snorm_int8: // convert_char_sat_rte(f * 127.0f) @@ -574,7 +579,8 @@ convertWriteData(const vec WriteData, { vec PixelData = processFloatDataToPixel(WriteData, 32.0f); - PixelData = cl::sycl::min(PixelData, static_cast(0x1f)); + PixelData = + cl::sycl::min(PixelData, static_cast(0x1f)); // Compressing the data into the first element of PixelData. // This is needed so that the data can be directly stored into the pixel // location from the first element. @@ -616,7 +622,8 @@ convertWriteData(const vec WriteData, case image_channel_type::fp32: return WriteData.convert(); default: - break; + throw cl::sycl::invalid_parameter_error("Unsupported data type", + PI_INVALID_VALUE); } } @@ -659,7 +666,8 @@ convertWriteData(const vec WriteData, "image_channel_type of the image.", PI_INVALID_VALUE); default: - break; + throw cl::sycl::invalid_parameter_error("Unsupported data type", + PI_INVALID_VALUE); } } @@ -782,8 +790,8 @@ void imageWriteHostImpl(const CoordT &Coords, const WriteDataT &Color, template DataT ReadPixelData(const cl_int4 PixelCoord, const id<3> ImgPitch, const image_channel_type ImageChannelType, - const image_channel_order ImageChannelOrder, - void *BasePtr, const uint8_t ElementSize) { + const image_channel_order ImageChannelOrder, void *BasePtr, + const uint8_t ElementSize) { DataT Color(0); auto Ptr = static_cast(BasePtr) + getImageOffset(PixelCoord, ImgPitch, @@ -908,29 +916,28 @@ DataT ReadPixelDataLinearFiltMode(const cl_int8 CoordValues, cl_int i0 = CoordValues.s0(), j0 = CoordValues.s1(), k0 = CoordValues.s2(), i1 = CoordValues.s4(), j1 = CoordValues.s5(), k1 = CoordValues.s6(); - auto getColorInFloat = - [&](cl_int4 V) { - DataT Res = getColor(V, SmplAddrMode, - ImgRange, ImgPitch, ImgChannelType, - ImgChannelOrder, BasePtr, ElementSize); - return Res.template convert(); - }; + auto getColorInFloat = [&](cl_int4 V) { + DataT Res = + getColor(V, SmplAddrMode, ImgRange, ImgPitch, ImgChannelType, + ImgChannelOrder, BasePtr, ElementSize); + return Res.template convert(); + }; // Get Color Values at each Coordinate. cl_float4 Ci0j0k0 = getColorInFloat(cl_int4{i0, j0, k0, 0}); - + cl_float4 Ci1j0k0 = getColorInFloat(cl_int4{i1, j0, k0, 0}); - + cl_float4 Ci0j1k0 = getColorInFloat(cl_int4{i0, j1, k0, 0}); - + cl_float4 Ci1j1k0 = getColorInFloat(cl_int4{i1, j1, k0, 0}); - + cl_float4 Ci0j0k1 = getColorInFloat(cl_int4{i0, j0, k1, 0}); - + cl_float4 Ci1j0k1 = getColorInFloat(cl_int4{i1, j0, k1, 0}); - + cl_float4 Ci0j1k1 = getColorInFloat(cl_int4{i0, j1, k1, 0}); - + cl_float4 Ci1j1k1 = getColorInFloat(cl_int4{i1, j1, k1, 0}); cl_float a = abc.x(); diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp index 2ea4e4e83db66..482dd6920726d 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/sycl/include/CL/sycl/handler.hpp @@ -23,6 +23,7 @@ #include #include +#include #include #include #include @@ -135,7 +136,7 @@ device getDeviceFromHandler(handler &); /// "finalization" it constructs CG object, that represents specific operation, /// passing fields that are required only. class handler { -private: +protected: /// Constructs SYCL handler from queue. /// /// \param Queue is a SYCL queue. @@ -337,7 +338,7 @@ class handler { using KI = sycl::detail::KernelInfo; // Empty name indicates that the compilation happens without integration // header, so don't perform things that require it. - if (KI::getName() != "") { + if (std::strlen(KI::getName()) > 0) { MArgs.clear(); extractArgsAndReqsFromLambda(MHostKernel->getPtr(), KI::getNumParams(), &KI::getParamDesc(0)); @@ -1241,7 +1242,7 @@ class handler { MCGType = detail::CG::PREFETCH_USM; } -private: +protected: shared_ptr_class MQueue; /// The storage for the arguments passed. /// We need to store a copy of values that are passed explicitly through diff --git a/sycl/test/basic_tests/image_api.cpp b/sycl/test/basic_tests/image_api.cpp deleted file mode 100644 index 1a5abc258af79..0000000000000 --- a/sycl/test/basic_tests/image_api.cpp +++ /dev/null @@ -1,165 +0,0 @@ -// REQUIRES: opencl - -// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -I %sycl_source_dir %s -o %t1.out -// RUN: %clangxx -I %sycl_source_dir %s -o %t3.out -lsycl -I %sycl_include -// RUN: env SYCL_DEVICE_TYPE=HOST %t1.out -// RUN: env SYCL_DEVICE_TYPE=HOST %t3.out -// RUN: %CPU_RUN_PLACEHOLDER %t1.out -// RUN: %GPU_RUN_PLACEHOLDER %t1.out -// RUN: %ACC_RUN_PLACEHOLDER %t1.out - -#include -// FIXME do not use internal methods in tests. -#include -#include -#include - -#include -#include -#include -#include -#include -#include - -namespace s = cl::sycl; -namespace d = cl::sycl::detail; - -struct FakeHandler { - d::NDRDescT MNDRDesc; - s::unique_ptr_class MHostKernel; - s::shared_ptr_class MSyclKernel; - s::vector_class> MArgsStorage; - s::vector_class MAccStorage; - s::vector_class MRequirements; - s::vector_class MArgs; - s::vector_class> MSharedPtrStorage; - s::string_class MKernelName; - d::OSModuleHandle MOSModuleHandle; - s::vector_class> MStreamStorage; -}; - -int main() { - constexpr size_t Size = 256; - std::array Src; - std::array Dest; - std::fill(Src.begin(), Src.end(), s::float4{1.0f, 2.0f, 3.0f, 4.0f}); - std::fill(Dest.begin(), Dest.end(), s::float4{0.0f, 0.0f, 0.0f, 0.0}); - - try { - constexpr int Dimensions = 2; - constexpr s::image_channel_order ChannelOrder = - s::image_channel_order::rgba; - constexpr s::image_channel_type ChannelType = s::image_channel_type::fp32; - const s::range Range{16, 16}; - s::image SrcImage(Src.data(), ChannelOrder, ChannelType, Range); - s::image DestImage(Dest.data(), ChannelOrder, ChannelType, - Range); - - s::queue Queue; - s::device Device = Queue.get_device(); - if(!Device.get_info()) { - std::cout << "Images are not supported. The result can not be checked." - << std::endl; - return 0; - } - - FakeHandler Handler; - Handler.MNDRDesc.set(Range); - if (Queue.is_host()) { - auto KernelFunc = [&](s::id ID) { - s::int2 Coords(ID[0], ID[1]); - int linearId = ID[1] + ID[0] * Range[0]; - s::float4 Color = Src[linearId]; - Color *= 10.0f; - Dest[linearId] = Color; - }; - - Handler.MHostKernel.reset( - new d::HostKernel, - Dimensions>(KernelFunc)); - } else { - s::context Context = Queue.get_context(); - s::program Program(Context); - Program.build_with_source( - "__kernel void ImageTest(__read_only image2d_t Input, __write_only " - "image2d_t Output, sampler_t sampler) {" - " int2 Coords = (int2)(get_global_id(0), get_global_id(1));" - " float4 Color = read_imagef(Input, sampler, Coords);" - " Color *= 10.0f;" - " write_imagef(Output, Coords, Color);" - "}\n"); - s::kernel Kernel = Program.get_kernel("ImageTest"); - Handler.MSyclKernel = d::getSyclObjImpl(Kernel); - Handler.MKernelName = - Handler.MSyclKernel->get_info(); - Handler.MOSModuleHandle = - d::OSUtil::getOSModuleHandle(Handler.MKernelName.c_str()); - } - - auto addFakeImageAccessor = [&Handler, Dimensions](s::image Image, - s::access::mode Mode, int Index) { - const s::id<3> Offset{0, 0, 0}; - const s::range<3> AccessRange{Image.get_range()[0], Image.get_range()[1], 1}; - const s::range<3> MemoryRange{Image.get_range()[0], Image.get_range()[1], 1}; - d::SYCLMemObjI *SYCLMemObject = static_cast(d::getSyclObjImpl(Image).get()); - const int ElemSize = d::getSyclObjImpl(Image)->getElementSize(); - - d::AccessorImplPtr AccImpl = std::make_shared(Offset, - AccessRange,MemoryRange, Mode, SYCLMemObject, Dimensions, ElemSize); - - d::Requirement *Req = AccImpl.get(); - Handler.MRequirements.push_back(Req); - Handler.MAccStorage.push_back(AccImpl); - Handler.MArgs.emplace_back(d::kernel_param_kind_t::kind_accessor, Req, - static_cast(s::access::target::image), - Index); - }; - - addFakeImageAccessor(SrcImage, s::access::mode::read, 0); - addFakeImageAccessor(DestImage, s::access::mode::write, 1); - - s::sampler Sampler(s::coordinate_normalization_mode::unnormalized, - s::addressing_mode::clamp, s::filtering_mode::nearest); - Handler.MArgsStorage.emplace_back(sizeof(s::sampler)); - s::sampler *SamplerPtr = - reinterpret_cast(Handler.MArgsStorage.back().data()); - *SamplerPtr = Sampler; - Handler.MArgs.emplace_back(d::kernel_param_kind_t::kind_sampler, - static_cast(SamplerPtr), - sizeof(s::sampler), 2); - - s::unique_ptr_class CommandGroup; - CommandGroup.reset(new d::CGExecKernel( - std::move(Handler.MNDRDesc), std::move(Handler.MHostKernel), - std::move(Handler.MSyclKernel), std::move(Handler.MArgsStorage), - std::move(Handler.MAccStorage), std::move(Handler.MSharedPtrStorage), - std::move(Handler.MRequirements), /*DepsEvents*/ {}, - std::move(Handler.MArgs), std::move(Handler.MKernelName), - std::move(Handler.MOSModuleHandle), std::move(Handler.MStreamStorage), - d::CG::KERNEL)); - - d::EventImplPtr Event = d::Scheduler::getInstance().addCG( - std::move(CommandGroup), d::getSyclObjImpl(Queue)); - - s::event EventRet = d::createSyclObjFromImpl(Event); - EventRet.wait(); - } catch (const s::exception &E) { - std::cout << "SYCL exception caught: " << E.what() << std::endl; - } - - s::float4 Expected{10.f, 20.f, 30.f, 40.f}; - - bool Result = std::all_of(Dest.cbegin(), Dest.cend(), - [Expected](const s::float4 &Value) -> bool { - return s::all(s::isequal(Value, Expected)); - }); - - if (Result) { - std::cout << "The result is correct." << std::endl; - } else { - std::cout << "The result is incorrect." << std::endl; - assert(Result); - } - return 0; -} diff --git a/sycl/unittests/scheduler/CMakeLists.txt b/sycl/unittests/scheduler/CMakeLists.txt index 3c630071a2398..03c971a1a3f0d 100644 --- a/sycl/unittests/scheduler/CMakeLists.txt +++ b/sycl/unittests/scheduler/CMakeLists.txt @@ -11,6 +11,7 @@ add_sycl_unittest(SchedulerTests FinishedCmdCleanup.cpp LeafLimit.cpp MemObjCommandCleanup.cpp + ImageApi.cpp utils.cpp ) diff --git a/sycl/unittests/scheduler/ImageApi.cpp b/sycl/unittests/scheduler/ImageApi.cpp new file mode 100644 index 0000000000000..1145f2df49f1c --- /dev/null +++ b/sycl/unittests/scheduler/ImageApi.cpp @@ -0,0 +1,73 @@ +//==----------------- ImageApi.cpp --- Scheduler unit tests ----------------==// +// +// 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 "MockHandler.hpp" +#include "SchedulerTest.hpp" + +#include + +#include + +TEST_F(SchedulerTest, ImageApi) { + + constexpr size_t Size = 256; + constexpr int Dimensions = 2; + + std::array Src; + std::array Dest; + std::fill(Src.begin(), Src.end(), sycl::float4{1.0f, 2.0f, 3.0f, 4.0f}); + std::fill(Dest.begin(), Dest.end(), sycl::float4{0.0f, 0.0f, 0.0f, 0.0}); + + constexpr sycl::image_channel_order ChannelOrder = + sycl::image_channel_order::rgba; + constexpr sycl::image_channel_type ChannelType = + sycl::image_channel_type::fp32; + const sycl::range Range{16, 16}; + + sycl::image SrcImg(Src.data(), ChannelOrder, ChannelType, Range); + sycl::image DstImg(Dest.data(), ChannelOrder, ChannelType, Range); + + auto CGHLambda = [&](MockHandler &CGH) { + auto SrcAcc = + SrcImg.template get_access(CGH); + auto DstAcc = + DstImg.template get_access( + CGH); + + EXPECT_EQ(CGH.getKernelAccessors().size(), 2UL); + + sycl::sampler Sampler(sycl::coordinate_normalization_mode::unnormalized, + sycl::addressing_mode::clamp, + sycl::filtering_mode::nearest); + + CGH.parallel_for(Range, [=](sycl::id<2> ID) { + sycl::int2 Coords{ID[1], ID[0]}; + sycl::float4 Color = SrcAcc.read(Coords, Sampler); + Color *= 10.0f; + DstAcc.write(Coords, Color); + }); + }; + + MockHandler MockCGH(sycl::detail::getSyclObjImpl(MQueue)); + + CGHLambda(MockCGH); + + EXPECT_EQ(MockCGH.getKernelArgs().size(), 2UL); + + auto Event = MockCGH.mockFinalize(); + Event.wait(); + + sycl::float4 Expected{10.f, 20.f, 30.f, 40.f}; + + bool Result = std::all_of(Dest.cbegin(), Dest.cend(), + [Expected](const sycl::float4 &Value) -> bool { + return sycl::all(sycl::isequal(Value, Expected)); + }); + + EXPECT_TRUE(Result); +} diff --git a/sycl/unittests/scheduler/MockHandler.hpp b/sycl/unittests/scheduler/MockHandler.hpp new file mode 100644 index 0000000000000..136e056d84bef --- /dev/null +++ b/sycl/unittests/scheduler/MockHandler.hpp @@ -0,0 +1,24 @@ +//==------------ MockHandler.hpp --- Scheduler unit tests ------------------==// +// +// 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 +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include +#include + +class MockHandler : public sycl::handler { +public: + MockHandler(sycl::shared_ptr_class Queue) + : sycl::handler(std::move(Queue), true) {} + sycl::event mockFinalize() { return finalize(); } + + sycl::vector_class &getKernelArgs() { return MArgs; } + sycl::vector_class &getKernelAccessors() { + return MAssociatedAccesors; + } +};