From 50b081be0dedeb923d067b9f194fd004c0cb00d5 Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Tue, 10 Mar 2020 14:31:56 +0300 Subject: [PATCH 1/7] [SYCL] Rewrite image_api LIT as Google Unit test Avoids usage of private APIs in end-to-end tests, places a scheduler test alongside its siblings. Signed-off-by: Alexander Batashev --- sycl/include/CL/sycl/handler.hpp | 4 +- sycl/test/basic_tests/image_api.cpp | 164 ----------------------- sycl/unittests/scheduler/ImageApi.cpp | 72 ++++++++++ sycl/unittests/scheduler/MockHandler.hpp | 24 ++++ 4 files changed, 98 insertions(+), 166 deletions(-) delete mode 100644 sycl/test/basic_tests/image_api.cpp create mode 100644 sycl/unittests/scheduler/ImageApi.cpp create mode 100644 sycl/unittests/scheduler/MockHandler.hpp diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp index a7c0583d5ae23..2e7e4f8981d21 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/sycl/include/CL/sycl/handler.hpp @@ -135,7 +135,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. @@ -1240,7 +1240,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 4e7976311416d..0000000000000 --- a/sycl/test/basic_tests/image_api.cpp +++ /dev/null @@ -1,164 +0,0 @@ -// 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 -// 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/ImageApi.cpp b/sycl/unittests/scheduler/ImageApi.cpp new file mode 100644 index 0000000000000..690c354c1e4bc --- /dev/null +++ b/sycl/unittests/scheduler/ImageApi.cpp @@ -0,0 +1,72 @@ +//==----------------- 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(Scheduler, ImageApi) { + + size_t Size = 256; + 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 int Dimensions = 2; + 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 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; + } +}; From 765e5c6894369ddd300ff1bb09a6b1a387d0ac69 Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Thu, 12 Mar 2020 11:44:44 +0300 Subject: [PATCH 2/7] Fix tests Signed-off-by: Alexander Batashev --- sycl/unittests/scheduler/CMakeLists.txt | 1 + sycl/unittests/scheduler/ImageApi.cpp | 4 ++-- 2 files changed, 3 insertions(+), 2 deletions(-) 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 index 690c354c1e4bc..c82862ce148fc 100644 --- a/sycl/unittests/scheduler/ImageApi.cpp +++ b/sycl/unittests/scheduler/ImageApi.cpp @@ -13,9 +13,9 @@ #include -TEST_F(Scheduler, ImageApi) { +TEST_F(SchedulerTest, ImageApi) { - size_t Size = 256; + constexpr size_t Size = 256; std::array Src; std::array Dest; std::fill(Src.begin(), Src.end(), sycl::float4{1.0f, 2.0f, 3.0f, 4.0f}); From 1670fac5b6aa5e1416baa49540e3fb21f651f749 Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Thu, 12 Mar 2020 14:35:59 +0300 Subject: [PATCH 3/7] Fix warnings revealed by the test Signed-off-by: Alexander Batashev --- sycl/include/CL/sycl/detail/cg.hpp | 2 +- .../CL/sycl/detail/image_accessor_util.hpp | 20 ++++++++++++++++--- sycl/include/CL/sycl/handler.hpp | 3 ++- 3 files changed, 20 insertions(+), 5 deletions(-) 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..ccc82cc5beacb 100644 --- a/sycl/include/CL/sycl/detail/image_accessor_util.hpp +++ b/sycl/include/CL/sycl/detail/image_accessor_util.hpp @@ -530,13 +530,27 @@ 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()); + cl::sycl::clamp(TempInInt, min_v::type>(), + max_v::type>()); return TempInIntSaturated.convert(); } @@ -616,7 +630,7 @@ 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 +673,7 @@ 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); } } diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp index f885934e99a06..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 @@ -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)); From 2abcd3613ae37b752c8813acf263f3b7a57d8df6 Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Tue, 17 Mar 2020 11:23:56 +0300 Subject: [PATCH 4/7] Reorder constants Signed-off-by: Alexander Batashev --- sycl/unittests/scheduler/ImageApi.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/unittests/scheduler/ImageApi.cpp b/sycl/unittests/scheduler/ImageApi.cpp index c82862ce148fc..1e4d70b1f3b33 100644 --- a/sycl/unittests/scheduler/ImageApi.cpp +++ b/sycl/unittests/scheduler/ImageApi.cpp @@ -16,12 +16,13 @@ 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 int Dimensions = 2; constexpr sycl::image_channel_order ChannelOrder = sycl::image_channel_order::rgba; constexpr sycl::image_channel_type ChannelType = From fe762cd4fb63f568b6c3bc161f8c831bb2f997bd Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Tue, 17 Mar 2020 13:37:37 +0300 Subject: [PATCH 5/7] Fix test on windows Signed-off-by: Alexander Batashev --- sycl/unittests/scheduler/ImageApi.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/unittests/scheduler/ImageApi.cpp b/sycl/unittests/scheduler/ImageApi.cpp index 1e4d70b1f3b33..1145f2df49f1c 100644 --- a/sycl/unittests/scheduler/ImageApi.cpp +++ b/sycl/unittests/scheduler/ImageApi.cpp @@ -45,7 +45,7 @@ TEST_F(SchedulerTest, ImageApi) { sycl::addressing_mode::clamp, sycl::filtering_mode::nearest); - CGH.parallel_for(Range, [=](sycl::id ID) { + 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; From 8c86418604f2d23c71d8598af33425ef7dfc74fc Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Tue, 17 Mar 2020 13:38:07 +0300 Subject: [PATCH 6/7] Apply clang-format Signed-off-by: Alexander Batashev --- .../CL/sycl/detail/image_accessor_util.hpp | 25 ++++++++----------- 1 file changed, 10 insertions(+), 15 deletions(-) diff --git a/sycl/include/CL/sycl/detail/image_accessor_util.hpp b/sycl/include/CL/sycl/detail/image_accessor_util.hpp index ccc82cc5beacb..d7856cf5aadb7 100644 --- a/sycl/include/CL/sycl/detail/image_accessor_util.hpp +++ b/sycl/include/CL/sycl/detail/image_accessor_util.hpp @@ -530,27 +530,20 @@ convertWriteData(const vec WriteData, } } -template -struct pixel_type_helper { - using type = T; -}; +template struct pixel_type_helper { using type = T; }; -template <> struct pixel_type_helper { - using type = int32_t; -}; +template <> struct pixel_type_helper { using type = int32_t; }; -template <> struct pixel_type_helper { - using type = int64_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::type>(), - max_v::type>()); + vec TempInIntSaturated = cl::sycl::clamp( + TempInInt, min_v::type>(), + max_v::type>()); return TempInIntSaturated.convert(); } @@ -630,7 +623,8 @@ convertWriteData(const vec WriteData, case image_channel_type::fp32: return WriteData.convert(); default: - throw cl::sycl::invalid_parameter_error("Unsupported data type", PI_INVALID_VALUE); + throw cl::sycl::invalid_parameter_error("Unsupported data type", + PI_INVALID_VALUE); } } @@ -673,7 +667,8 @@ convertWriteData(const vec WriteData, "image_channel_type of the image.", PI_INVALID_VALUE); default: - throw cl::sycl::invalid_parameter_error("Unsupported data type", PI_INVALID_VALUE); + throw cl::sycl::invalid_parameter_error("Unsupported data type", + PI_INVALID_VALUE); } } From 1930dac69e7bf893944ead26ceaca0ada804d134 Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Tue, 17 Mar 2020 15:10:42 +0300 Subject: [PATCH 7/7] Fix more warnings Signed-off-by: Alexander Batashev --- .../CL/sycl/detail/image_accessor_util.hpp | 36 +++++++++---------- 1 file changed, 17 insertions(+), 19 deletions(-) diff --git a/sycl/include/CL/sycl/detail/image_accessor_util.hpp b/sycl/include/CL/sycl/detail/image_accessor_util.hpp index d7856cf5aadb7..cea5845775e97 100644 --- a/sycl/include/CL/sycl/detail/image_accessor_util.hpp +++ b/sycl/include/CL/sycl/detail/image_accessor_util.hpp @@ -552,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) @@ -581,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. @@ -791,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, @@ -917,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();