diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp index f9fc05e4ff6a1..b94c183e01e41 100644 --- a/sycl/include/CL/sycl/types.hpp +++ b/sycl/include/CL/sycl/types.hpp @@ -199,6 +199,32 @@ using is_int_to_int = std::integral_constant::value && std::is_integral::value>; +template +using is_sint_to_sint = + std::integral_constant::value && + is_sigeninteger::value>; + +template +using is_uint_to_uint = + std::integral_constant::value && + is_sugeninteger::value>; + +template +using is_sint_to_from_uint = std::integral_constant< + bool, is_sugeninteger::value && is_sigeninteger::value || + is_sigeninteger::value && is_sugeninteger::value>; + +template +using is_sint_to_float = + std::integral_constant::value && + !(std::is_unsigned::value) && + detail::is_floating_point::value>; + +template +using is_uint_to_float = + std::integral_constant::value && + detail::is_floating_point::value>; + template using is_int_to_float = std::integral_constant::value && @@ -213,15 +239,23 @@ template using is_float_to_float = std::integral_constant::value && detail::is_floating_point::value>; +template +using is_standard_type = std::integral_constant< + bool, detail::is_sgentype::value && !std::is_same::value && + !std::is_same::value>; -template +template detail::enable_if_t::value, R> convertImpl(T Value) { return Value; } +#ifndef __SYCL_DEVICE_ONLY__ + // Note for float to half conversions, static_cast calls the conversion operator // implemented for host that takes care of the precision requirements. -template +template detail::enable_if_t::value && (is_int_to_int::value || is_int_to_float::value || @@ -231,9 +265,9 @@ convertImpl(T Value) { return static_cast(Value); } -#ifndef __SYCL_DEVICE_ONLY__ // float to int -template +template detail::enable_if_t::value, R> convertImpl(T Value) { switch (roundingMode) { // Round to nearest even is default rounding mode for floating-point types @@ -280,16 +314,145 @@ using Rtp = detail::bool_constant; template using Rtn = detail::bool_constant; -// Convert floating-point type to integer type +// convert signed and unsigned types with an equal size and diff names +template +detail::enable_if_t::value && + (is_sint_to_sint::value || + is_uint_to_uint::value) && + std::is_same::value, + R> +convertImpl(T Value) { + return static_cast(Value); +} + +// signed to signed +#define __SYCL_GENERATE_CONVERT_IMPL(DestType) \ + template \ + detail::enable_if_t::value && \ + is_sint_to_sint::value && \ + (std::is_same::value || \ + std::is_same::value && \ + std::is_same::value) && \ + !std::is_same::value, \ + R> \ + convertImpl(T Value) { \ + OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); \ + return __spirv_SConvert##_R##DestType(OpValue); \ + } + +__SYCL_GENERATE_CONVERT_IMPL(char) +__SYCL_GENERATE_CONVERT_IMPL(short) +__SYCL_GENERATE_CONVERT_IMPL(int) +__SYCL_GENERATE_CONVERT_IMPL(long) +__SYCL_GENERATE_CONVERT_IMPL(longlong) + +#undef __SYCL_GENERATE_CONVERT_IMPL + +// unsigned to unsigned +#define __SYCL_GENERATE_CONVERT_IMPL(DestType) \ + template \ + detail::enable_if_t::value && \ + is_uint_to_uint::value && \ + std::is_same::value && \ + !std::is_same::value, \ + R> \ + convertImpl(T Value) { \ + OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); \ + return __spirv_UConvert##_R##DestType(OpValue); \ + } + +__SYCL_GENERATE_CONVERT_IMPL(uchar) +__SYCL_GENERATE_CONVERT_IMPL(ushort) +__SYCL_GENERATE_CONVERT_IMPL(uint) +__SYCL_GENERATE_CONVERT_IMPL(ulong) + +#undef __SYCL_GENERATE_CONVERT_IMPL + +// unsigned to (from) signed +template +detail::enable_if_t::value, R> convertImpl(T Value) { + return static_cast(Value); +} + +// sint to float +#define __SYCL_GENERATE_CONVERT_IMPL(SPIRVOp, DestType) \ + template \ + detail::enable_if_t< \ + is_sint_to_float::value && std::is_same::value, R> \ + convertImpl(T Value) { \ + OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); \ + return __spirv_Convert##SPIRVOp##_R##DestType(OpValue); \ + } + +__SYCL_GENERATE_CONVERT_IMPL(SToF, half) +__SYCL_GENERATE_CONVERT_IMPL(SToF, float) +__SYCL_GENERATE_CONVERT_IMPL(SToF, double) + +#undef __SYCL_GENERATE_CONVERT_IMPL + +// uint to float +#define __SYCL_GENERATE_CONVERT_IMPL(SPIRVOp, DestType) \ + template \ + detail::enable_if_t< \ + is_uint_to_float::value && std::is_same::value, R> \ + convertImpl(T Value) { \ + OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); \ + return __spirv_Convert##SPIRVOp##_R##DestType(OpValue); \ + } + +__SYCL_GENERATE_CONVERT_IMPL(UToF, half) +__SYCL_GENERATE_CONVERT_IMPL(UToF, float) +__SYCL_GENERATE_CONVERT_IMPL(UToF, double) + +#undef __SYCL_GENERATE_CONVERT_IMPL + +// float to float +#define __SYCL_GENERATE_CONVERT_IMPL(DestType, RoundingMode, \ + RoundingModeCondition) \ + template \ + detail::enable_if_t::value && \ + is_float_to_float::value && \ + std::is_same::value && \ + RoundingModeCondition::value, \ + R> \ + convertImpl(T Value) { \ + OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); \ + return __spirv_FConvert##_R##DestType##_##RoundingMode(OpValue); \ + } + +#define __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(RoundingMode, \ + RoundingModeCondition) \ + __SYCL_GENERATE_CONVERT_IMPL(double, RoundingMode, RoundingModeCondition) \ + __SYCL_GENERATE_CONVERT_IMPL(float, RoundingMode, RoundingModeCondition) \ + __SYCL_GENERATE_CONVERT_IMPL(half, RoundingMode, RoundingModeCondition) + +__SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rte, RteOrAutomatic) +__SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtz, Rtz) +__SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtp, Rtp) +__SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtn, Rtn) + +#undef __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE +#undef __SYCL_GENERATE_CONVERT_IMPL + +// float to int #define __SYCL_GENERATE_CONVERT_IMPL(SPIRVOp, DestType, RoundingMode, \ RoundingModeCondition) \ - template \ + template \ detail::enable_if_t::value && \ - std::is_same::value && \ + (std::is_same::value || \ + std::is_same::value && \ + std::is_same::value) && \ RoundingModeCondition::value, \ R> \ convertImpl(T Value) { \ - using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; \ OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); \ return __spirv_Convert##SPIRVOp##_R##DestType##_##RoundingMode(OpValue); \ } @@ -319,6 +482,18 @@ __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtn, Rtn) #undef __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE #undef __SYCL_GENERATE_CONVERT_IMPL +// Back up +template +detail::enable_if_t< + (!is_standard_type::value && !is_standard_type::value || + !is_standard_type::value && !is_standard_type::value) && + !std::is_same::value, + R> +convertImpl(T Value) { + return static_cast(Value); +} + #endif // __SYCL_DEVICE_ONLY__ } // namespace detail @@ -627,9 +802,13 @@ template class vec { detail::is_floating_point::value, "Unsupported convertT"); vec Result; + using OpenCLT = detail::ConvertToOpenCLType_t; + using OpenCLR = detail::ConvertToOpenCLType_t; for (size_t I = 0; I < NumElements; ++I) { Result.setValue( - I, detail::convertImpl(getValue(I))); + I, + detail::convertImpl( + getValue(I))); } return Result; } diff --git a/sycl/test/basic_tests/vec_convert.cpp b/sycl/test/basic_tests/vec_convert.cpp deleted file mode 100644 index 1e7235152e1da..0000000000000 --- a/sycl/test/basic_tests/vec_convert.cpp +++ /dev/null @@ -1,154 +0,0 @@ -// XFAIL: cuda -// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out -// RUN: env SYCL_DEVICE_TYPE=HOST %t.out -// RUN: %CPU_RUN_PLACEHOLDER %t.out -// RUN: %GPU_RUN_PLACEHOLDER %t.out -// RUN: %ACC_RUN_PLACEHOLDER %t.out -//==------------ vec_convert.cpp - SYCL vec class convert method test ------==// -// -// 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 - -// TODO make the test to pass on cuda - -using namespace cl::sycl; - -template class kernel_name; - -template struct helper; - -template <> struct helper<0> { - template - static void compare(const vec &x, - const vec &y) { - const T xs = x.template swizzle<0>(); - const T ys = y.template swizzle<0>(); - assert(xs == ys); - } -}; - -template struct helper { - template - static void compare(const vec &x, - const vec &y) { - const T xs = x.template swizzle(); - const T ys = y.template swizzle(); - helper::compare(x, y); - assert(xs == ys); - } -}; - -template -void test(const vec &ToConvert, - const vec &Expected) { - vec Converted{0}; - { - buffer, 1> Buffer{&Converted, range<1>{1}}; - queue Queue; - Queue.submit([&](handler &CGH) { - accessor, 1, access::mode::write> Accessor( - Buffer, CGH); - CGH.single_task(roundingMode)>>([=]() { - Accessor[0] = ToConvert.template convert(); - }); - }); - } - helper::compare(Converted, Expected); -} - -int main() { - // automatic - test( - int8{2, 3, 3, -2, -3, -3, 0, 0}, - int8{2, 3, 3, -2, -3, -3, 0, 0}); - test( - float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}, - int8{2, 2, 3, -2, -2, -3, 0, 0}); - test( - int8{2, 3, 3, -2, -3, -3, 0, 0}, - float8{2.f, 3.f, 3.f, -2.f, -3.f, -3.f, 0.f, 0.f}); - test( - float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}, - float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}); - test( - float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}, - half8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}); - - // rte - test( - int8{2, 3, 3, -2, -3, -3, 0, 0}, - int8{2, 3, 3, -2, -3, -3, 0, 0}); - test( - float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}, - int8{2, 2, 3, -2, -2, -3, 0, 0}); - test( - int8{2, 3, 3, -2, -3, -3, 0, 0}, - float8{2.f, 3.f, 3.f, -2.f, -3.f, -3.f, 0.f, 0.f}); - test( - float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}, - float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}); - test( - float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}, - half8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}); - - // rtz - test( - int8{2, 3, 3, -2, -3, -3, 0, 0}, - int8{2, 3, 3, -2, -3, -3, 0, 0}); - test( - float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}, - int8{2, 2, 2, -2, -2, -2, 0, 0}); - test( - int8{2, 3, 3, -2, -3, -3, 0, 0}, - float8{2.f, 3.f, 3.f, -2.f, -3.f, -3.f, 0.f, 0.f}); - test( - float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}, - float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}); - test( - float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}, - half8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}); - - // rtp - test( - int8{2, 3, 3, -2, -3, -3, 0, 0}, - int8{2, 3, 3, -2, -3, -3, 0, 0}); - test( - float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}, - int8{3, 3, 3, -2, -2, -2, 0, 0}); - test( - int8{2, 3, 3, -2, -3, -3, 0, 0}, - float8{2.f, 3.f, 3.f, -2.f, -3.f, -3.f, 0.f, 0.f}); - test( - float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}, - float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}); - test( - float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}, - half8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}); - - // rtn - test( - int8{2, 3, 3, -2, -3, -3, 0, 0}, - int8{2, 3, 3, -2, -3, -3, 0, 0}); - test( - float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}, - int8{2, 2, 2, -3, -3, -3, 0, 0}); - test( - int8{2, 3, 3, -2, -3, -3, 0, 0}, - float8{2.f, 3.f, 3.f, -2.f, -3.f, -3.f, 0.f, 0.f}); - test( - float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}, - float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}); - test( - float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}, - half8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}); - - return 0; -} diff --git a/sycl/test/basic_tests/vec_convert.hpp b/sycl/test/basic_tests/vec_convert.hpp new file mode 100644 index 0000000000000..981b69a189ce9 --- /dev/null +++ b/sycl/test/basic_tests/vec_convert.hpp @@ -0,0 +1,53 @@ +#include + +#include + +using namespace cl::sycl; + +template +class kernel_name; + +template +struct helper; + +template <> +struct helper<0> { + template + static void compare(const vec &x, + const vec &y) { + const T xs = x.template swizzle<0>(); + const T ys = y.template swizzle<0>(); + assert(xs == ys); + } +}; + +template +struct helper { + template + static void compare(const vec &x, + const vec &y) { + const T xs = x.template swizzle(); + const T ys = y.template swizzle(); + helper::compare(x, y); + assert(xs == ys); + } +}; + +template +void test(const vec &ToConvert, + const vec &Expected) { + vec Converted{0}; + { + buffer, 1> Buffer{&Converted, range<1>{1}}; + queue Queue; + Queue.submit([&](handler &CGH) { + accessor, 1, access::mode::write> Accessor( + Buffer, CGH); + CGH.single_task(roundingMode)>>([=]() { + Accessor[0] = ToConvert.template convert(); + }); + }); + } + helper::compare(Converted, Expected); +} diff --git a/sycl/test/basic_tests/vec_convert_f_to_f.cpp b/sycl/test/basic_tests/vec_convert_f_to_f.cpp new file mode 100644 index 0000000000000..803385f90da00 --- /dev/null +++ b/sycl/test/basic_tests/vec_convert_f_to_f.cpp @@ -0,0 +1,61 @@ +// XFAIL: cuda +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out +//==------------ vec_convert.cpp - SYCL vec class convert method test ------==// +// +// 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 "vec_convert.hpp" + +// TODO make the convertion on CPU and HOST identical +// TODO make the test to pass on cuda + +int main() { + // automatic + test( + double8{1234567890.0, 987654304.0, 100.0, -50.0, 111111.111, 625.625, 50625.0009765625, -2500000.875}, + float8{1234567936.0f, 987654272.0f, 100.0f, -50.0f, 111111.109375f, 625.625f, 50625.0f, -2500001.0f}); + test( + float8{1234567936.0f, 987654272.0f, 100.0f, -50.0f, 111111.109375f, 625.625f, 50625.0f, -2500001.0f}, + double8{1234567936.0, 987654272.0, 100.0, -50.0, 111111.109375, 625.625, 50625.0, -2500001.0}); + + // rte + test( + double8{1234567890.0, 987654304.0, 100.0, -50.0, 111111.111, 625.625, 50625.0009765625, -2500000.875}, + float8{1234567936.0f, 987654272.0f, 100.0f, -50.0f, 111111.109375f, 625.625f, 50625.0f, -2500001.0f}); + test( + float8{1234567936.0f, 987654272.0f, 100.0f, -50.0f, 111111.109375f, 625.625f, 50625.0f, -2500001.0f}, + double8{1234567936.0, 987654272.0, 100.0, -50.0, 111111.109375, 625.625, 50625.0, -2500001.0}); + + // rtp + test( + double8{1234567890.0, 987654304.0, 100.0, -50.0, 111111.111, 625.625, 50625.0009765625, -2500000.875}, + float8{1234567936.0f, 987654336.0f, 100.0f, -50.0f, 111111.1171875f, 625.625f, 50625.00390625f, -2500000.75f}); + test( + float8{1234567936.0f, 987654272.0f, 100.0f, -50.0f, 111111.109375f, 625.625f, 50625.0f, -2500001.0f}, + double8{1234567936.0, 987654272.0, 100.0, -50.0, 111111.109375, 625.625, 50625.0, -2500001.0}); + + // rtn + test( + double8{1234567890.0, 987654304.0, 100.0, -50.0, 111111.111, 625.625, 50625.0009765625, -2500000.875}, + float8{1234567808.0f, 987654272.0f, 100.0f, -50.0f, 111111.109375f, 625.625f, 50625.0f, -2500001.0f}); + test( + float8{1234567936.0f, 987654272.0f, 100.0f, -50.0f, 111111.109375f, 625.625f, 50625.0f, -2500001.0f}, + double8{1234567936.0, 987654272.0, 100.0, -50.0, 111111.109375, 625.625, 50625.0, -2500001.0}); + + // rtz + test( + double8{1234567890.0, 987654304.0, 100.0, -50.0, 111111.111, 625.625, 50625.0009765625, -2500000.875}, + float8{1234567808.0f, 987654272.0f, 100.0f, -50.0f, 111111.109375f, 625.625f, 50625.0f, -2500000.75f}); + test( + float8{1234567936.0f, 987654272.0f, 100.0f, -50.0f, 111111.109375f, 625.625f, 50625.0f, -2500001.0f}, + double8{1234567936.0, 987654272.0, 100.0, -50.0, 111111.109375, 625.625, 50625.0, -2500001.0}); + + return 0; +} diff --git a/sycl/test/basic_tests/vec_convert_f_to_i.cpp b/sycl/test/basic_tests/vec_convert_f_to_i.cpp new file mode 100644 index 0000000000000..65c8a67e21a01 --- /dev/null +++ b/sycl/test/basic_tests/vec_convert_f_to_i.cpp @@ -0,0 +1,60 @@ +// XFAIL: cuda +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out +// RUN: env SYCL_DEVICE_TYPE=HOST %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out +//==------------ vec_convert.cpp - SYCL vec class convert method test ------==// +// +// 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 "vec_convert.hpp" + +// TODO make the test to pass on cuda + +int main() { + // automatic + test( + float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}, + int8{2, 2, 3, -2, -2, -3, 0, 0}); + test( + int8{2, 3, 3, -2, -3, -3, 0, 0}, + float8{2.f, 3.f, 3.f, -2.f, -3.f, -3.f, 0.f, 0.f}); + + // rte + test( + float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}, + int8{2, 2, 3, -2, -2, -3, 0, 0}); + test( + int8{2, 3, 3, -2, -3, -3, 0, 0}, + float8{2.f, 3.f, 3.f, -2.f, -3.f, -3.f, 0.f, 0.f}); + + // rtz + test( + float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}, + int8{2, 2, 2, -2, -2, -2, 0, 0}); + test( + int8{2, 3, 3, -2, -3, -3, 0, 0}, + float8{2.f, 3.f, 3.f, -2.f, -3.f, -3.f, 0.f, 0.f}); + + // rtp + test( + float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}, + int8{3, 3, 3, -2, -2, -2, 0, 0}); + test( + int8{2, 3, 3, -2, -3, -3, 0, 0}, + float8{2.f, 3.f, 3.f, -2.f, -3.f, -3.f, 0.f, 0.f}); + + // rtn + test( + float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}, + int8{2, 2, 2, -3, -3, -3, 0, 0}); + test( + int8{2, 3, 3, -2, -3, -3, 0, 0}, + float8{2.f, 3.f, 3.f, -2.f, -3.f, -3.f, 0.f, 0.f}); + return 0; +} diff --git a/sycl/test/basic_tests/vec_convert_half.cpp b/sycl/test/basic_tests/vec_convert_half.cpp new file mode 100644 index 0000000000000..cecd812d76932 --- /dev/null +++ b/sycl/test/basic_tests/vec_convert_half.cpp @@ -0,0 +1,102 @@ +// XFAIL: cuda +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out +//==------------ vec_convert_half.cpp - SYCL vec class convert method test ------==// +// +// 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 + +// TODO make the convertion on CPU and HOST identical +// TODO make the test to pass on cuda + +using namespace cl::sycl; + +template +class kernel_name; + +template +struct helper; + +template <> +struct helper<0> { + template + static void compare(const vec &x, + const vec &y) { + const T xs = x.template swizzle<0>(); + const T ys = y.template swizzle<0>(); + assert(xs == ys); + } +}; + +template +struct helper { + template + static void compare(const vec &x, + const vec &y) { + const T xs = x.template swizzle(); + const T ys = y.template swizzle(); + helper::compare(x, y); + assert(xs == ys); + } +}; + +template +void test(const vec &ToConvert, + const vec &Expected) { + vec Converted{0}; + { + buffer, 1> Buffer{&Converted, range<1>{1}}; + queue Queue; + + cl::sycl::device D = Queue.get_device(); + if (!D.has_extension("cl_khr_fp16")) + exit(0); + + Queue.submit([&](handler &CGH) { + accessor, 1, access::mode::write> Accessor( + Buffer, CGH); + CGH.single_task(roundingMode)>>([=]() { + Accessor[0] = ToConvert.template convert(); + }); + }); + } + helper::compare(Converted, Expected); +} + +int main() { + //automatic + test( + double4{12345.0, 100.0, -50.0, 11111.111}, + half4{12344.0f, 100.0, -50.0, 11112}); + + //rte + test( + double4{12345.0, 100.0, -50.0, 11111.111}, + half4{12344.0f, 100.0, -50.0, 11112}); + //rtp + test( + double4{12345.0, 100.0, -50.0, 11111.111}, + half4{12352.0f, 100.0, -50.0, 11112}); + + //rtn + test( + double4{12345.0, 100.0, -50.0, 11111.111}, + half4{12344.0f, 100.0, -50.0, 11104}); + + //rtz + test( + double4{12345.0, 100.0, -50.0, 11111.111}, + half4{12344.0f, 100.0, -50.0, 11104}); + + return 0; +} diff --git a/sycl/test/basic_tests/vec_convert_i_to_i.cpp b/sycl/test/basic_tests/vec_convert_i_to_i.cpp new file mode 100644 index 0000000000000..7228b37777665 --- /dev/null +++ b/sycl/test/basic_tests/vec_convert_i_to_i.cpp @@ -0,0 +1,51 @@ +// XFAIL: cuda +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out +// RUN: env SYCL_DEVICE_TYPE=HOST %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out +//==------------ vec_convert_i_to_i.cpp - SYCL vec class convert method test ------==// +// +// 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 "vec_convert.hpp" + +// TODO make the test to pass on cuda + +int main() { + + test( + short8{300, -300, 100, -50, 128, -129, 0, 1}, + char8{44, -44, 100, -50, -128, 127, 0, 1}); + test( + int8{100000, -100000, 100, -50, 32768, -32769, 0, 1}, + short8{-31072, 31072, 100, -50, -32768, 32767, 0, 1}); + test( + long8{3000000000, -3000000000, 100, -50, 2147483648, -2147483649, 0, 1}, + int8{-1294967296, 1294967296, 100, -50, -2147483648, 2147483647, 0, 1}); + + test( + ushort8{300, 255, 100, 150, 128, 256, 0, 1}, + uchar8{44, 255, 100, 150, 128, 0, 0, 1}); + test( + uint8{100000, 65535, 100, 150, 32768, 65536, 0, 1}, + ushort8{34464, 65535, 100, 150, 32768, 0, 0, 1}); + test( + ulong8{10000000000, 4294967295, 100, 150, 2147483648, 4294967296, 0, 1}, + uint8{1410065408, 4294967295, 100, 150, 2147483648, 0, 0, 1}); + + test( + int8{2147483647, -1, 100, 150, -100, -2147483648, 0, 1}, + uint8{2147483647, 4294967295, 100, 150, 4294967196, 2147483648, 0, 1}); + test( + short8{32767, -1, 100, 150, -100, -32768, 0, 1}, + uint8{32767, 4294967295, 100, 150, 4294967196, 4294934528, 0, 1}); + test( + ulong8{3000000000, 2147483647, 100, 150, 2147483648, 1000, 0, 1}, + int8{-1294967296, 2147483647, 100, 150, -2147483648, 1000, 0, 1}); + return 0; +}