From 33f405f1f23b0fc16967fd4828eaffbe7d48b332 Mon Sep 17 00:00:00 2001 From: Garima Gupta Date: Thu, 9 Apr 2020 12:20:24 -0700 Subject: [PATCH 1/8] [SYCL] Support more image_channel_types in image_accessors read/write API for half4 datatype on host device. Signed-off-by: Garima Gupta --- .../CL/sycl/detail/image_accessor_util.hpp | 55 +++++- .../image_accessor_readwrite_half.cpp | 161 ++++++++++++++++++ 2 files changed, 207 insertions(+), 9 deletions(-) create mode 100644 sycl/test/basic_tests/image_accessor_readwrite_half.cpp diff --git a/sycl/include/CL/sycl/detail/image_accessor_util.hpp b/sycl/include/CL/sycl/detail/image_accessor_util.hpp index 0d712cb73125f..7313fd40c8b0b 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: @@ -451,7 +465,7 @@ void convertReadData(const vec PixelData, "image_channel_type of the image.", PI_INVALID_VALUE); case image_channel_type::fp16: - RetData = PixelData.template convert(); + RetDataFloat = PixelData.template convert(); break; case image_channel_type::fp32: throw cl::sycl::invalid_parameter_error( @@ -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 @@ -539,10 +554,10 @@ template vec processFloatDataToPixel(vec WriteData, float MulFactor) { vec Temp = WriteData * MulFactor; - vec TempInInt = Temp.convert(); + vec TempInInt = Temp.template convert(); vec TempInIntSaturated = cl::sycl::clamp(TempInInt, min_v(), max_v()); - return TempInIntSaturated.convert(); + return TempInIntSaturated.template convert(); } template @@ -625,16 +640,35 @@ convertWriteData(const vec WriteData, } } +/* +template +vec processHalfDataToPixel(vec WriteData, + float MulFactor) { + vec Temp = WriteData * MulFactor; + vec TempInInt = Temp.convert(); + vec TempInIntSaturated = + cl::sycl::clamp(TempInInt, min_v(), max_v()); + return TempInIntSaturated.convert(); +}*/ + 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 +1028,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_half.cpp b/sycl/test/basic_tests/image_accessor_readwrite_half.cpp new file mode 100644 index 0000000000000..d7be57c8e9dac --- /dev/null +++ b/sycl/test/basic_tests/image_accessor_readwrite_half.cpp @@ -0,0 +1,161 @@ +// RUN: %clangxx -fsycl %s -o %t.out +// RUN: env SYCL_DEVICE_TYPE=HOST %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUNx: %CPU_RUN_PLACEHOLDER %t.out // OpenCL CPU backend does not support half datatype reads and writes for 1D images. + +//==--------------------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 only for dimension=1. + // create image: + char HostPtr[100]; + for (int i = 0; i < 100; i++) + HostPtr[i] = i; + + check_half4(HostPtr); +} From f7436798c76fc5278e9f1b5cfdcc5a6f3ed3c659 Mon Sep 17 00:00:00 2001 From: Garima Gupta Date: Fri, 10 Apr 2020 16:39:50 -0700 Subject: [PATCH 2/8] [SYCL] Clean up of image_accessor_readwrite test. Enabled it for GPU. Signed-off-by: Garima Gupta --- .../basic_tests/image_accessor_readwrite.cpp | 305 +++++------------- 1 file changed, 88 insertions(+), 217 deletions(-) diff --git a/sycl/test/basic_tests/image_accessor_readwrite.cpp b/sycl/test/basic_tests/image_accessor_readwrite.cpp index 9ebaa14e5e550..aad560e2b4382 100644 --- a/sycl/test/basic_tests/image_accessor_readwrite.cpp +++ b/sycl/test/basic_tests/image_accessor_readwrite.cpp @@ -1,8 +1,7 @@ // 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 //==--------------------image_accessor_readwrite.cpp ----------------------==// //==----------image_accessor read without sampler & write API test---------==// // @@ -23,102 +22,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 +66,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 +104,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 +121,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 +153,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 +186,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 +220,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 +250,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 +287,4 @@ int main() { check(HostPtr); check(HostPtr); check(HostPtr); - // check(HostPtr); } From 027d67446e41795756580296b4919edd0138f1a0 Mon Sep 17 00:00:00 2001 From: Garima Gupta Date: Fri, 10 Apr 2020 17:00:06 -0700 Subject: [PATCH 3/8] Use of git-clang-format Signed-off-by: Garima Gupta --- sycl/include/CL/sycl/detail/image_accessor_util.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/CL/sycl/detail/image_accessor_util.hpp b/sycl/include/CL/sycl/detail/image_accessor_util.hpp index 7313fd40c8b0b..19ae493b1985b 100644 --- a/sycl/include/CL/sycl/detail/image_accessor_util.hpp +++ b/sycl/include/CL/sycl/detail/image_accessor_util.hpp @@ -655,7 +655,7 @@ template vec convertWriteData(const vec WriteData, const image_channel_type ImageChannelType) { - vec WriteDataFloat = WriteData.convert(); + vec WriteDataFloat = WriteData.convert(); switch (ImageChannelType) { case image_channel_type::snorm_int8: // convert_char_sat_rte(h * 127.0f) From 1ad6104978049771ec232c7c728b484ef5f09d0d Mon Sep 17 00:00:00 2001 From: Garima Gupta Date: Fri, 10 Apr 2020 17:17:39 -0700 Subject: [PATCH 4/8] small change Signed-off-by: Garima Gupta --- sycl/include/CL/sycl/detail/image_accessor_util.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/include/CL/sycl/detail/image_accessor_util.hpp b/sycl/include/CL/sycl/detail/image_accessor_util.hpp index 19ae493b1985b..635fb2122dab5 100644 --- a/sycl/include/CL/sycl/detail/image_accessor_util.hpp +++ b/sycl/include/CL/sycl/detail/image_accessor_util.hpp @@ -554,10 +554,10 @@ template vec processFloatDataToPixel(vec WriteData, float MulFactor) { vec Temp = WriteData * MulFactor; - vec TempInInt = Temp.template convert(); + vec TempInInt = Temp.convert(); vec TempInIntSaturated = cl::sycl::clamp(TempInInt, min_v(), max_v()); - return TempInIntSaturated.template convert(); + return TempInIntSaturated.convert(); } template From 280dbc6bf064a8caaeaae39dc10b989f0858ddcf Mon Sep 17 00:00:00 2001 From: Garima Gupta Date: Mon, 13 Apr 2020 10:23:14 -0700 Subject: [PATCH 5/8] [SYCL]Correction of lit test case. Signed-off-by: Garima Gupta --- sycl/test/basic_tests/image_accessor_readwrite.cpp | 3 +++ sycl/test/basic_tests/image_accessor_readwrite_half.cpp | 5 +++-- 2 files changed, 6 insertions(+), 2 deletions(-) diff --git a/sycl/test/basic_tests/image_accessor_readwrite.cpp b/sycl/test/basic_tests/image_accessor_readwrite.cpp index aad560e2b4382..7f54dc2182b7c 100644 --- a/sycl/test/basic_tests/image_accessor_readwrite.cpp +++ b/sycl/test/basic_tests/image_accessor_readwrite.cpp @@ -2,6 +2,9 @@ // RUN: env SYCL_DEVICE_TYPE=HOST %t.out // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out + +// TODO: No CUDA image support +// XFAIL: cuda //==--------------------image_accessor_readwrite.cpp ----------------------==// //==----------image_accessor read without sampler & write API test---------==// // diff --git a/sycl/test/basic_tests/image_accessor_readwrite_half.cpp b/sycl/test/basic_tests/image_accessor_readwrite_half.cpp index d7be57c8e9dac..89b1a4e1a8fa0 100644 --- a/sycl/test/basic_tests/image_accessor_readwrite_half.cpp +++ b/sycl/test/basic_tests/image_accessor_readwrite_half.cpp @@ -3,6 +3,8 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // RUNx: %CPU_RUN_PLACEHOLDER %t.out // OpenCL CPU backend does not support half datatype reads and writes for 1D images. +// TODO: No CUDA image support +// XFAIL: cuda //==--------------------image_accessor_readwrite_half.cpp -------------------==// //==-image_accessor read (without sampler)& write API test for half datatype-==// // @@ -67,8 +69,7 @@ void check_read_data(s::cl_half4 ReadData, s::cl_half4 ExpectedColor) { check_read_data(ReadDatafloat, ExpectedColorfloat); } -template +template void write_type_order(char *HostPtr, const s::image_channel_order ImgOrder, WriteDataT Color) { From 0184595d85a5112c4a73d050bec427fd4c54339c Mon Sep 17 00:00:00 2001 From: Garima Gupta Date: Tue, 14 Apr 2020 11:58:37 -0700 Subject: [PATCH 6/8] [SYCL] Addition of comments. Signed-off-by: Garima Gupta --- .../CL/sycl/detail/image_accessor_util.hpp | 15 ++------------- .../basic_tests/image_accessor_readwrite_half.cpp | 10 +++++++++- 2 files changed, 11 insertions(+), 14 deletions(-) diff --git a/sycl/include/CL/sycl/detail/image_accessor_util.hpp b/sycl/include/CL/sycl/detail/image_accessor_util.hpp index 635fb2122dab5..b7ac161e7358a 100644 --- a/sycl/include/CL/sycl/detail/image_accessor_util.hpp +++ b/sycl/include/CL/sycl/detail/image_accessor_util.hpp @@ -465,8 +465,8 @@ void convertReadData(const vec PixelData, "image_channel_type of the image.", PI_INVALID_VALUE); case image_channel_type::fp16: - RetDataFloat = PixelData.template convert(); - break; + RetData = PixelData.template convert(); + return; case image_channel_type::fp32: throw cl::sycl::invalid_parameter_error( "Datatype to read - cl_half4 is incompatible with the " @@ -640,17 +640,6 @@ convertWriteData(const vec WriteData, } } -/* -template -vec processHalfDataToPixel(vec WriteData, - float MulFactor) { - vec Temp = WriteData * MulFactor; - vec TempInInt = Temp.convert(); - vec TempInIntSaturated = - cl::sycl::clamp(TempInInt, min_v(), max_v()); - return TempInIntSaturated.convert(); -}*/ - template vec convertWriteData(const vec WriteData, diff --git a/sycl/test/basic_tests/image_accessor_readwrite_half.cpp b/sycl/test/basic_tests/image_accessor_readwrite_half.cpp index 89b1a4e1a8fa0..6abb0d78d4b34 100644 --- a/sycl/test/basic_tests/image_accessor_readwrite_half.cpp +++ b/sycl/test/basic_tests/image_accessor_readwrite_half.cpp @@ -1,7 +1,7 @@ // 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 -// RUNx: %CPU_RUN_PLACEHOLDER %t.out // OpenCL CPU backend does not support half datatype reads and writes for 1D images. // TODO: No CUDA image support // XFAIL: cuda @@ -152,6 +152,14 @@ void check_half4(char *HostPtr) { }; 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]; From 69b3a1f09e6ffc113575a71c3c9b9bd6550e23c1 Mon Sep 17 00:00:00 2001 From: Garima Gupta Date: Tue, 14 Apr 2020 22:58:25 -0700 Subject: [PATCH 7/8] Added test support for CUDA. The test passes since it doesn't support fp16. Signed-off-by: Garima Gupta --- sycl/test/basic_tests/image_accessor_readwrite_half.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/sycl/test/basic_tests/image_accessor_readwrite_half.cpp b/sycl/test/basic_tests/image_accessor_readwrite_half.cpp index 6abb0d78d4b34..6c0cacd3e682f 100644 --- a/sycl/test/basic_tests/image_accessor_readwrite_half.cpp +++ b/sycl/test/basic_tests/image_accessor_readwrite_half.cpp @@ -3,8 +3,6 @@ // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out -// TODO: No CUDA image support -// XFAIL: cuda //==--------------------image_accessor_readwrite_half.cpp -------------------==// //==-image_accessor read (without sampler)& write API test for half datatype-==// // From 7d4811501d6bc8de227e41497531c6108c3e082a Mon Sep 17 00:00:00 2001 From: Garima Gupta Date: Wed, 15 Apr 2020 09:51:07 -0700 Subject: [PATCH 8/8] XFAIL to UNSUPPORTED. Signed-off-by: Garima Gupta --- sycl/test/basic_tests/image_accessor_readwrite.cpp | 3 +-- sycl/test/basic_tests/image_accessor_readwrite_half.cpp | 1 + 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test/basic_tests/image_accessor_readwrite.cpp b/sycl/test/basic_tests/image_accessor_readwrite.cpp index 7f54dc2182b7c..0f541e86efa4d 100644 --- a/sycl/test/basic_tests/image_accessor_readwrite.cpp +++ b/sycl/test/basic_tests/image_accessor_readwrite.cpp @@ -3,8 +3,7 @@ // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out -// TODO: No CUDA image support -// XFAIL: cuda +// UNSUPPORTED: cuda //==--------------------image_accessor_readwrite.cpp ----------------------==// //==----------image_accessor read without sampler & write API test---------==// // diff --git a/sycl/test/basic_tests/image_accessor_readwrite_half.cpp b/sycl/test/basic_tests/image_accessor_readwrite_half.cpp index 6c0cacd3e682f..e085556f9b319 100644 --- a/sycl/test/basic_tests/image_accessor_readwrite_half.cpp +++ b/sycl/test/basic_tests/image_accessor_readwrite_half.cpp @@ -3,6 +3,7 @@ // 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-==// //