diff --git a/sycl/include/CL/sycl/detail/image_accessor_util.hpp b/sycl/include/CL/sycl/detail/image_accessor_util.hpp index 0d712cb73125f..b7ac161e7358a 100644 --- a/sycl/include/CL/sycl/detail/image_accessor_util.hpp +++ b/sycl/include/CL/sycl/detail/image_accessor_util.hpp @@ -423,12 +423,26 @@ template void convertReadData(const vec PixelData, const image_channel_type ImageChannelType, vec &RetData) { - + vec RetDataFloat; switch (ImageChannelType) { case image_channel_type::snorm_int8: + // max(-1.0f, (half)c / 127.0f) + RetDataFloat = (PixelData.template convert()) / 127.0f; + RetDataFloat = cl::sycl::fmax(RetDataFloat, -1); + break; case image_channel_type::snorm_int16: + // max(-1.0f, (half)c / 32767.0f) + RetDataFloat = (PixelData.template convert()) / 32767.0f; + RetDataFloat = cl::sycl::fmax(RetDataFloat, -1); + break; case image_channel_type::unorm_int8: + // (half)c / 255.0f + RetDataFloat = (PixelData.template convert()) / 255.0f; + break; case image_channel_type::unorm_int16: + // (half)c / 65535.0f + RetDataFloat = (PixelData.template convert()) / 65535.0f; + break; case image_channel_type::unorm_short_565: case image_channel_type::unorm_short_555: case image_channel_type::unorm_int_101010: @@ -452,7 +466,7 @@ void convertReadData(const vec PixelData, PI_INVALID_VALUE); case image_channel_type::fp16: RetData = PixelData.template convert(); - break; + return; case image_channel_type::fp32: throw cl::sycl::invalid_parameter_error( "Datatype to read - cl_half4 is incompatible with the " @@ -461,6 +475,7 @@ void convertReadData(const vec PixelData, default: break; } + RetData = RetDataFloat.template convert(); } // Converts data to write into appropriate datatype based on the channel of the @@ -629,12 +644,20 @@ template vec convertWriteData(const vec WriteData, const image_channel_type ImageChannelType) { - + vec WriteDataFloat = WriteData.convert(); switch (ImageChannelType) { case image_channel_type::snorm_int8: + // convert_char_sat_rte(h * 127.0f) + return processFloatDataToPixel(WriteDataFloat, 127.0f); case image_channel_type::snorm_int16: + // convert_short_sat_rte(h * 32767.0f) + return processFloatDataToPixel(WriteDataFloat, 32767.0f); case image_channel_type::unorm_int8: + // convert_uchar_sat_rte(h * 255.0f) + return processFloatDataToPixel(WriteDataFloat, 255.0f); case image_channel_type::unorm_int16: + // convert_ushort_sat_rte(h * 65535.0f) + return processFloatDataToPixel(WriteDataFloat, 65535.0f); case image_channel_type::unorm_short_565: case image_channel_type::unorm_short_555: case image_channel_type::unorm_int_101010: @@ -994,10 +1017,13 @@ DataT ReadPixelDataLinearFiltMode(const cl_int8 CoordValues, // ImgChannelType. // Convert to DataT as per conversion rules in section 8.3 in OpenCL Spec. // -// TODO: -// Extend support for Step2 and Step3 for Linear Filtering Mode. -// Extend support to find out of bounds Coordinates and return appropriate -// value based on Addressing Mode. +// TODO: Add additional check for half datatype read. +// Based on OpenCL spec 2.0: +// "The read_imageh calls that take integer coordinates must use a sampler with +// filter mode set to CLK_FILTER_NEAREST, normalized coordinates set to +// CLK_NORMALIZED_COORDS_FALSE and addressing mode set to +// CLK_ADDRESS_CLAMP_TO_EDGE, CLK_ADDRESS_CLAMP or CLK_ADDRESS_NONE; otherwise +// the values returned are undefined." template DataT imageReadSamplerHostImpl(const CoordT &Coords, const sampler &Smpl, diff --git a/sycl/test/basic_tests/image_accessor_readwrite.cpp b/sycl/test/basic_tests/image_accessor_readwrite.cpp index 9ebaa14e5e550..0f541e86efa4d 100644 --- a/sycl/test/basic_tests/image_accessor_readwrite.cpp +++ b/sycl/test/basic_tests/image_accessor_readwrite.cpp @@ -1,8 +1,9 @@ // RUN: %clangxx -fsycl %s -o %t.out // RUN: env SYCL_DEVICE_TYPE=HOST %t.out // RUN: %CPU_RUN_PLACEHOLDER %t.out -// RUNx: %GPU_RUN_PLACEHOLDER %t.out -// RUNx: %ACC_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out + +// UNSUPPORTED: cuda //==--------------------image_accessor_readwrite.cpp ----------------------==// //==----------image_accessor read without sampler & write API test---------==// // @@ -23,102 +24,42 @@ namespace s = cl::sycl; template class kernel_class; -template ::value))>::type> -void check_write_data(PixelDataType *HostDataPtr, PixelDataT ExpectedData) { -#if DEBUG_OUTPUT - { - if ((HostDataPtr[0] == (PixelDataType)ExpectedData.x()) && - (HostDataPtr[1] == (PixelDataType)ExpectedData.y()) && - (HostDataPtr[2] == (PixelDataType)ExpectedData.z()) && - (HostDataPtr[3] == (PixelDataType)ExpectedData.w())) { - std::cout << "Data written is correct: " << std::endl; - } else { - std::cout << "Data written is WRONG: " << std::endl; - } - std::cout << "HostDataPtr: \t" << (float)HostDataPtr[0] << " " - << (float)HostDataPtr[1] << " " << (float)HostDataPtr[2] << " " - << (float)HostDataPtr[3] << std::endl; - - std::cout << "ExpectedData: \t" << (float)ExpectedData.x() << " " - << (float)ExpectedData.y() << " " << (float)ExpectedData.z() - << " " << (float)ExpectedData.w() << std::endl; - } -#else - assert(HostDataPtr[0] == (PixelDataType)ExpectedData.x()); - assert(HostDataPtr[1] == (PixelDataType)ExpectedData.y()); - assert(HostDataPtr[2] == (PixelDataType)ExpectedData.z()); - assert(HostDataPtr[3] == (PixelDataType)ExpectedData.w()); -#endif -} - -void check_write_data(s::cl_half *HostDataPtr, s::cl_half4 ExpectedData) { -#if DEBUG_OUTPUT - { - if ((HostDataPtr[0] == (float)ExpectedData.x()) && - (HostDataPtr[1] == (float)ExpectedData.y()) && - (HostDataPtr[2] == (float)ExpectedData.z()) && - (HostDataPtr[3] == (float)ExpectedData.w())) { - std::cout << "Data written is correct: " << std::endl; - } else { - std::cout << "Data written is WRONG: " << std::endl; - } - std::cout << "HostDataPtr: \t" << (float)HostDataPtr[0] << " " - << (float)HostDataPtr[1] << " " << (float)HostDataPtr[2] << " " - << (float)HostDataPtr[3] << std::endl; - - std::cout << "ExpectedData: \t" << (float)ExpectedData.x() << " " - << (float)ExpectedData.y() << " " << (float)ExpectedData.z() - << " " << (float)ExpectedData.w() << std::endl; - } -#else - assert(HostDataPtr[0] == (float)ExpectedData.x()); - assert(HostDataPtr[1] == (float)ExpectedData.y()); - assert(HostDataPtr[2] == (float)ExpectedData.z()); - assert(HostDataPtr[3] == (float)ExpectedData.w()); -#endif -} - template ::value) && !(std::is_same::value))>::type> void check_read_data(ReadDataT ReadData, ReadDataT ExpectedColor) { using ReadDataType = typename s::detail::TryToGetElementType::type; + bool CorrectData = false; + if ((ReadData.x() == ExpectedColor.x()) && + (ReadData.y() == ExpectedColor.y()) && + (ReadData.z() == ExpectedColor.z()) && + (ReadData.w() == ExpectedColor.w())) + CorrectData = true; + #if DEBUG_OUTPUT - { - if (((ReadDataType)ReadData.x() == (ReadDataType)ExpectedColor.x()) && - ((ReadDataType)ReadData.y() == (ReadDataType)ExpectedColor.y()) && - ((ReadDataType)ReadData.z() == (ReadDataType)ExpectedColor.z()) && - ((ReadDataType)ReadData.w() == (ReadDataType)ExpectedColor.w())) { - std::cout << "Read Data is correct: " << std::endl; - } else { - std::cout << "Read Data is WRONG: " << std::endl; - } - std::cout << "ReadData: \t" - << std::setprecision(std::numeric_limits::digits10 + - 1) - << (ReadDataType)ReadData.x() << " " - << (ReadDataType)ReadData.y() << " " - << (ReadDataType)ReadData.z() << " " - << (ReadDataType)ReadData.w() << std::endl; - - std::cout << "ExpectedColor: \t" - << std::setprecision(std::numeric_limits::digits10 + - 1) - << (ReadDataType)ExpectedColor.x() << " " - << (ReadDataType)ExpectedColor.y() << " " - << (ReadDataType)ExpectedColor.z() << " " - << (ReadDataType)ExpectedColor.w() << std::endl; - } + if (CorrectData) + std::cout << "Read Data is correct: " << std::endl; + else + std::cout << "Read Data is WRONG: " << std::endl; + + std::cout << "ReadData: \t" + << std::setprecision(std::numeric_limits::digits10 + + 1) + << ReadData.x() << " " + << ReadData.y() << " " + << ReadData.z() << " " + << ReadData.w() << std::endl; + + std::cout << "ExpectedColor: \t" + << std::setprecision(std::numeric_limits::digits10 + + 1) + << ExpectedColor.x() << " " + << ExpectedColor.y() << " " + << ExpectedColor.z() << " " + << ExpectedColor.w() << std::endl; #else - { - assert((ReadDataType)ReadData.x() == (ReadDataType)ExpectedColor.x()); - assert((ReadDataType)ReadData.y() == (ReadDataType)ExpectedColor.y()); - assert((ReadDataType)ReadData.z() == (ReadDataType)ExpectedColor.z()); - assert((ReadDataType)ReadData.w() == (ReadDataType)ExpectedColor.w()); - } + assert(CorrectData); #endif } @@ -127,37 +68,34 @@ void check_read_data(s::cl_float4 ReadData, s::cl_float4 ExpectedColor) { s::cl_int4 PixelDataInt = ReadData.template as(); s::cl_int4 ExpectedDataInt = ExpectedColor.template as(); s::cl_int4 Diff = ExpectedDataInt - PixelDataInt; + bool CorrectData = false; + if ((Diff.x() <= 1 && Diff.x() >= -1) && + (Diff.y() <= 1 && Diff.y() >= -1) && + (Diff.z() <= 1 && Diff.z() >= -1) && + (Diff.w() <= 1 && Diff.w() >= -1)) + CorrectData = true; + #if DEBUG_OUTPUT - { - if (((s::cl_int)Diff.x() <= 1 && (s::cl_int)Diff.x() >= -1) && - ((s::cl_int)Diff.y() <= 1 && (s::cl_int)Diff.y() >= -1) && - ((s::cl_int)Diff.z() <= 1 && (s::cl_int)Diff.z() >= -1) && - ((s::cl_int)Diff.w() <= 1 && (s::cl_int)Diff.w() >= -1)) { - std::cout << "Read Data is correct within precision: " << std::endl; - } else { - std::cout << "Read Data is WRONG/ outside precision: " << std::endl; - } - std::cout << "ReadData: \t" - << std::setprecision(std::numeric_limits::digits10 + - 1) - << (float)ReadData.x() << " " << (float)ReadData.y() << " " - << (float)ReadData.z() << " " << (float)ReadData.w() - << std::endl; - - std::cout << "ExpectedColor: \t" - << std::setprecision(std::numeric_limits::digits10 + - 1) - << (float)ExpectedColor.x() << " " << (float)ExpectedColor.y() - << " " << (float)ExpectedColor.z() << " " - << (float)ExpectedColor.w() << std::endl; - } + if (CorrectData) + std::cout << "Read Data is correct within precision: " << std::endl; + else + std::cout << "Read Data is WRONG/ outside precision: " << std::endl; + + std::cout << "ReadData: \t" + << std::setprecision(std::numeric_limits::digits10 + + 1) + << ReadData.x() << " " << ReadData.y() << " " + << ReadData.z() << " " << ReadData.w() + << std::endl; + + std::cout << "ExpectedColor: \t" + << std::setprecision(std::numeric_limits::digits10 + + 1) + << ExpectedColor.x() << " " << ExpectedColor.y() + << " " << ExpectedColor.z() << " " + << ExpectedColor.w() << std::endl; #else - { - assert((s::cl_int)Diff.x() <= 1 && (s::cl_int)Diff.x() >= -1); - assert((s::cl_int)Diff.y() <= 1 && (s::cl_int)Diff.y() >= -1); - assert((s::cl_int)Diff.z() <= 1 && (s::cl_int)Diff.z() >= -1); - assert((s::cl_int)Diff.w() <= 1 && (s::cl_int)Diff.w() >= -1); - } + assert(CorrectData); #endif } @@ -168,11 +106,10 @@ void check_read_data(s::cl_half4 ReadData, s::cl_half4 ExpectedColor) { check_read_data(ReadDatafloat, ExpectedColorfloat); } -template -void check_write_type_order(char *HostPtr, - const s::image_channel_order ImgOrder, - WriteDataT Color, PixelDataType ExpectedData) { +template +void write_type_order(char *HostPtr, + const s::image_channel_order ImgOrder, + WriteDataT Color) { int Coord(2); { @@ -186,17 +123,6 @@ void check_write_type_order(char *HostPtr, }); }); } - - // Check Written Data. - using PixelElementType = - typename s::detail::TryToGetElementType::type; - int NumChannels = 4; - HostPtr = - HostPtr + (2 * s::detail::getImageElementSize(NumChannels, ImgType)); - // auto HostDataPtr = reinterpret_cast(HostPtr); - auto HostDataPtr = (PixelElementType *)(HostPtr); - - check_write_data((PixelElementType *)(HostPtr), ExpectedData); } template @@ -229,37 +155,28 @@ template void check(char *); template <> void check(char *HostPtr) { // valid channel types: // s::image_channel_type::signed_int8, - check_write_type_order( + write_type_order( HostPtr, s::image_channel_order::rgba, s::cl_int4(std::numeric_limits::max(), - std::numeric_limits::min(), 123, 0), - s::cl_char4(std::numeric_limits::max(), - std::numeric_limits::min(), 123, 0)); + std::numeric_limits::min(), 123, 0)); check_read_type_order( HostPtr, s::image_channel_order::rgba, s::cl_int4(std::numeric_limits::max(), std::numeric_limits::min(), 123, 0)); // s::image_channel_type::signed_int16, - check_write_type_order( + write_type_order( HostPtr, s::image_channel_order::rgba, s::cl_int4(std::numeric_limits::max(), - std::numeric_limits::min(), 123, 0), - s::cl_short4(std::numeric_limits::max(), - std::numeric_limits::min(), 123, 0)); + std::numeric_limits::min(), 123, 0)); check_read_type_order( HostPtr, s::image_channel_order::rgba, s::cl_int4(std::numeric_limits::max(), std::numeric_limits::min(), 123, 0)); // s::image_channel_type::signed_int32. - check_write_type_order( + write_type_order( HostPtr, s::image_channel_order::rgba, - s::cl_int4(std::numeric_limits::max(), - std::numeric_limits::min(), 123, 0), s::cl_int4(std::numeric_limits::max(), std::numeric_limits::min(), 123, 0)); check_read_type_order( @@ -271,37 +188,28 @@ template <> void check(char *HostPtr) { template <> void check(char *HostPtr) { // Calling only valid channel types with s::cl_uint4. // s::image_channel_type::signed_int8 - check_write_type_order( + write_type_order( HostPtr, s::image_channel_order::rgba, s::cl_uint4(std::numeric_limits::max(), - std::numeric_limits::min(), 123, 0), - s::cl_uchar4(std::numeric_limits::max(), - std::numeric_limits::min(), 123, 0)); + std::numeric_limits::min(), 123, 0)); check_read_type_order( HostPtr, s::image_channel_order::rgba, s::cl_uint4(std::numeric_limits::max(), std::numeric_limits::min(), 123, 0)); // s::image_channel_type::signed_int16 - check_write_type_order( + write_type_order( HostPtr, s::image_channel_order::rgba, s::cl_uint4(std::numeric_limits::max(), - std::numeric_limits::min(), 123, 0), - s::cl_ushort4(std::numeric_limits::max(), - std::numeric_limits::min(), 123, 0)); + std::numeric_limits::min(), 123, 0)); check_read_type_order( HostPtr, s::image_channel_order::rgba, s::cl_uint4(std::numeric_limits::max(), std::numeric_limits::min(), 123, 0)); // s::image_channel_type::signed_int32 - check_write_type_order( + write_type_order( HostPtr, s::image_channel_order::rgba, - s::cl_uint4(std::numeric_limits::max(), - std::numeric_limits::min(), 123, 0), s::cl_uint4(std::numeric_limits::max(), std::numeric_limits::min(), 123, 0)); check_read_type_order( @@ -314,41 +222,29 @@ template <> void check(char *HostPtr) { // Calling only valid channel types with s::cl_float4. // TODO: Correct the values below. // s::image_channel_type::snorm_int8, - check_write_type_order( - HostPtr, s::image_channel_order::rgba, s::cl_float4(2, -2, 0.375f, 0), - s::cl_char4(std::numeric_limits::max(), - std::numeric_limits::min(), 48, 0)); + write_type_order( + HostPtr, s::image_channel_order::rgba, s::cl_float4(2, -2, 0.375f, 0)); check_read_type_order( HostPtr, s::image_channel_order::rgba, s::cl_float4(1, -1, ((float)48 / 127) /*0.3779527544975280762f*/, 0)); // s::image_channel_type::snorm_int16, - check_write_type_order( - HostPtr, s::image_channel_order::rgba, s::cl_float4(2, -2, 0.375f, 0), - s::cl_short4(std::numeric_limits::max(), - std::numeric_limits::min(), 12288, 0)); + write_type_order( + HostPtr, s::image_channel_order::rgba, s::cl_float4(2, -2, 0.375f, 0)); check_read_type_order( HostPtr, s::image_channel_order::rgba, s::cl_float4(1, -1, ((float)12288 / 32767) /*0.375011444091796875f*/, 0)); // s::image_channel_type::unorm_int8, - check_write_type_order( - HostPtr, s::image_channel_order::rgba, s::cl_float4(2, -2, 0.375f, 0), - s::cl_uchar4(std::numeric_limits::max(), - std::numeric_limits::min(), 96, 0)); + write_type_order( + HostPtr, s::image_channel_order::rgba, s::cl_float4(2, -2, 0.375f, 0)); check_read_type_order( HostPtr, s::image_channel_order::rgba, s::cl_float4(1, 0, ((float)96 / 255) /*0.3764705955982208252f*/, 0)); // s::image_channel_type::unorm_int16 - check_write_type_order( - HostPtr, s::image_channel_order::rgba, s::cl_float4(2, -2, 0.375f, 0), - s::cl_ushort4(std::numeric_limits::max(), - std::numeric_limits::min(), 24576, 0)); + write_type_order( + HostPtr, s::image_channel_order::rgba, s::cl_float4(2, -2, 0.375f, 0)); check_read_type_order( HostPtr, s::image_channel_order::rgba, s::cl_float4(1, 0, ((float)24576 / 65535) /*0.3750057220458984375f*/, 0)); @@ -356,57 +252,35 @@ template <> void check(char *HostPtr) { // s::image_channel_type::unorm_short_565, order::rgbx // Currently unsupported since OpenCL has no information on this. - // TODO: Enable the below call, causing an error in scheduler + // TODO: Enable the below call, causing a runtime error in OpenCL CPU/GPU: + // OpenCL API returns: -10 (CL_IMAGE_FORMAT_NOT_SUPPORTED) -10 (CL_IMAGE_FORMAT_NOT_SUPPORTED) // s::image_channel_type::unorm_short_555, order::rgbx /* - check_write_type_order( - HostPtr, s::image_channel_order::rgbx, s::cl_float4(2, -2, 0.375f, 0), - s::cl_short4(std::numeric_limits::max(), - std::numeric_limits::min(), 3, 0)); + write_type_order( + HostPtr, s::image_channel_order::rgbx, s::cl_float4(2, -2, 0.375f, 0)); // s::image_channel_type::unorm_int_101010, order::rgbx - check_write_type_order( - HostPtr, s::image_channel_order::rgbx, s::cl_float4(2, -2, 0.375f, 0), - s::cl_uint4(std::numeric_limits::max(), - std::numeric_limits::min(), 3, 0)); + write_type_order( + HostPtr, s::image_channel_order::rgbx, s::cl_float4(2, -2, 0.375f, 0)); */ // s::image_channel_type::fp16 - check_write_type_order(HostPtr, s::image_channel_order::rgba, - s::cl_float4(2, -2, 0.375f, 0), - s::cl_half4(2, -2, 0.375, 0)); + write_type_order(HostPtr, s::image_channel_order::rgba, + s::cl_float4(2, -2, 0.375f, 0)); check_read_type_order( HostPtr, s::image_channel_order::rgba, s::cl_float4(2, -2, 0.375f, 0)); // s::image_channel_type::fp32 - check_write_type_order(HostPtr, s::image_channel_order::rgba, - s::cl_float4(2, -2, 0.375f, 0), - s::cl_float4(2, -2, 0.375f, 0)); + write_type_order(HostPtr, s::image_channel_order::rgba, + s::cl_float4(2, -2, 0.375f, 0)); check_read_type_order( HostPtr, s::image_channel_order::rgba, s::cl_float4(2, -2, 0.375f, 0)); }; -/* -template <> void check(char *HostPtr) { - - // Calling only valid channel types with s::cl_half4. - // s::image_channel_type::fp16 - // TODO: Enable the below call. Currently it doesn't work because of -s::cl_half - // Datatype explicit conversion issues on stmt 71-74 - check_write_type_order( - HostPtr, s::image_channel_order::rgba, s::cl_half4(2, -2, 0.375f, 0), - s::cl_half4(2, -2, 0.375, 0)); - check_read_type_order( - HostPtr, s::image_channel_order::rgba, s::cl_half4(2, -2, 0.375f, 0)); -};*/ int main() { // Checking only for dimension=1. // 4 datatypes possible: s::cl_uint4, s::cl_int4, s::cl_float4, s::cl_half4. + // half4 datatype is checked in a different test case. // create image: char HostPtr[100]; for (int i = 0; i < 100; i++) @@ -415,5 +289,4 @@ int main() { check(HostPtr); check(HostPtr); check(HostPtr); - // check(HostPtr); } diff --git a/sycl/test/basic_tests/image_accessor_readwrite_half.cpp b/sycl/test/basic_tests/image_accessor_readwrite_half.cpp new file mode 100644 index 0000000000000..e085556f9b319 --- /dev/null +++ b/sycl/test/basic_tests/image_accessor_readwrite_half.cpp @@ -0,0 +1,169 @@ +// RUN: %clangxx -fsycl %s -o %t.out +// RUN: env SYCL_DEVICE_TYPE=HOST %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out + +// UNSUPPORTED: cuda +//==--------------------image_accessor_readwrite_half.cpp -------------------==// +//==-image_accessor read (without sampler)& write API test for half datatype-==// +// +// 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 +#if DEBUG_OUTPUT +#include +#endif + +namespace s = cl::sycl; + +template +class kernel_class; + +void check_read_data(s::cl_float4 ReadData, s::cl_float4 ExpectedColor) { + // Maximum difference of 1.5 ULP is allowed. + s::cl_int4 PixelDataInt = ReadData.template as(); + s::cl_int4 ExpectedDataInt = ExpectedColor.template as(); + s::cl_int4 Diff = ExpectedDataInt - PixelDataInt; + bool CorrectData = false; + if (((s::cl_int)Diff.x() <= 1 && (s::cl_int)Diff.x() >= -1) && + ((s::cl_int)Diff.y() <= 1 && (s::cl_int)Diff.y() >= -1) && + ((s::cl_int)Diff.z() <= 1 && (s::cl_int)Diff.z() >= -1) && + ((s::cl_int)Diff.w() <= 1 && (s::cl_int)Diff.w() >= -1)) + CorrectData = true; + +#if DEBUG_OUTPUT + if (CorrectData) + std::cout << "Read Data is correct within precision: " << std::endl; + else + std::cout << "Read Data is WRONG/ outside precision: " << std::endl; + + std::cout << "ReadData: \t" + << std::setprecision(std::numeric_limits::digits10 + + 1) + << ReadData.x() << " " << ReadData.y() << " " + << ReadData.z() << " " << ReadData.w() + << std::endl; + + std::cout << "ExpectedColor: \t" + << std::setprecision(std::numeric_limits::digits10 + + 1) + << ExpectedColor.x() << " " << ExpectedColor.y() + << " " << ExpectedColor.z() << " " + << ExpectedColor.w() << std::endl; + +#else + assert(CorrectData); +#endif +} + +void check_read_data(s::cl_half4 ReadData, s::cl_half4 ExpectedColor) { + s::cl_float4 ReadDatafloat = ReadData.convert(); + s::cl_float4 ExpectedColorfloat = ExpectedColor.convert(); + check_read_data(ReadDatafloat, ExpectedColorfloat); +} + +template +void write_type_order(char *HostPtr, + const s::image_channel_order ImgOrder, + WriteDataT Color) { + + int Coord(2); + { + // image with dim = 1; + s::image<1> Img(HostPtr, ImgOrder, ImgType, s::range<1>{10}); + s::queue Queue; + Queue.submit([&](s::handler &cgh) { + auto WriteAcc = Img.get_access(cgh); + cgh.single_task(ImgType), 0>>([=](){ + WriteAcc.write(Coord, Color); + }); + }); + } +} + +template +void check_read_type_order(char *HostPtr, const s::image_channel_order ImgOrder, + ReadDataT ExpectedColor) { + + int Coord(2); + ReadDataT ReadData; + { + // image with dim = 1 + s::image<1> Img(HostPtr, ImgOrder, ImgType, s::range<1>{10}); + s::queue Queue; + s::buffer ReadDataBuf(&ReadData, s::range<1>(1)); + Queue.submit([&](s::handler &cgh) { + auto ReadAcc = Img.get_access(cgh); + s::accessor ReadDataBufAcc( + ReadDataBuf, cgh); + + cgh.single_task(ImgType), 1>>([=](){ + ReadDataT RetColor = ReadAcc.read(Coord); + ReadDataBufAcc[0] = RetColor; + }); + }); + } + check_read_data(ReadData, ExpectedColor); +} + +void check_half4(char *HostPtr) { + + // Calling only valid channel types with s::cl_half4. + // s::image_channel_type::snorm_int8, + write_type_order( + HostPtr, s::image_channel_order::rgba, s::cl_half4(2, -2, 0.375f, 0)); + check_read_type_order( + HostPtr, s::image_channel_order::rgba, + s::cl_half4(1, -1, ((float)48 / 127) /*0.3779527544975280762f*/, 0)); + + // s::image_channel_type::snorm_int16, + write_type_order( + HostPtr, s::image_channel_order::rgba, s::cl_half4(2, -2, 0.375f, 0)); + check_read_type_order( + HostPtr, s::image_channel_order::rgba, + s::cl_half4(1, -1, ((float)12288 / 32767) /*0.375011444091796875f*/, 0)); + + // s::image_channel_type::unorm_int8, + write_type_order( + HostPtr, s::image_channel_order::rgba, s::cl_half4(2, -2, 0.375f, 0)); + check_read_type_order( + HostPtr, s::image_channel_order::rgba, + s::cl_half4(1, 0, ((float)96 / 255) /*0.3764705955982208252f*/, 0)); + + // s::image_channel_type::unorm_int16 + write_type_order( + HostPtr, s::image_channel_order::rgba, s::cl_half4(1, -1, 0.375f, 0)); + check_read_type_order( + HostPtr, s::image_channel_order::rgba, + s::cl_half4(1, 0, ((float)24576 / 65535) /*0.3750057220458984375f*/, 0)); + + // s::image_channel_type::fp16 + write_type_order( + HostPtr, s::image_channel_order::rgba, s::cl_half4(2, -2, 0.375f, 0)); + check_read_type_order( + HostPtr, s::image_channel_order::rgba, s::cl_half4(2, -2, 0.375f, 0)); +}; + +int main() { + // Checking if default selected device supports half datatype. + // Same device will be selected in the write/read functions. + s::device Dev{s::default_selector()}; + if (!Dev.is_host() && !Dev.has_extension("cl_khr_fp16")) { + std::cout << "This device doesn't support the extension cl_khr_fp16" + << std::endl; + return 0; + } + // Checking only for dimension=1. + // create image: + char HostPtr[100]; + for (int i = 0; i < 100; i++) + HostPtr[i] = i; + + check_half4(HostPtr); +}