diff --git a/clang/lib/Sema/SPIRVBuiltins.td b/clang/lib/Sema/SPIRVBuiltins.td index 69c8f7b82c54e..41d5ead350414 100644 --- a/clang/lib/Sema/SPIRVBuiltins.td +++ b/clang/lib/Sema/SPIRVBuiltins.td @@ -286,8 +286,8 @@ class ConstOCLSPVBuiltin _Signature> : // OpenCL v1.0/1.2/2.0 s6.1.1: Built-in Scalar Data Types. def Bool : IntType<"bool", QualType<"BoolTy">, 1>; -def TrueChar : IntType<"char", QualType<"CharTy", 0, 1>, 8>; -def Char : IntType<"schar", QualType<"SignedCharTy", 0, 1>, 8>; +def TrueChar : IntType<"_char", QualType<"CharTy", 0, 1>, 8>; +def Char : IntType<"char", QualType<"SignedCharTy", 0, 1>, 8>; def SChar : IntType<"schar", QualType<"SignedCharTy", 0, 1>, 8>; def UChar : UIntType<"uchar", QualType<"UnsignedCharTy">, 8>; def Short : IntType<"short", QualType<"ShortTy", 0, 1>, 16>; @@ -713,8 +713,10 @@ foreach name = ["GenericPtrMemSemantics"] in { foreach IType = [UChar, UShort, UInt, ULong] in { foreach FType = [Float, Double, Half] in { - def : SPVBuiltin<"ConvertFToU_R" # IType.Name, [IType, FType], Attr.Const>; def : SPVBuiltin<"ConvertUToF_R" # FType.Name, [FType, IType], Attr.Const>; + foreach rnd = ["", "_rte", "_rtz", "_rtp", "_rtn"] in { + def : SPVBuiltin<"ConvertFToU_R" # IType.Name # rnd, [IType, FType], Attr.Const>; + } foreach v = [2, 3, 4, 8, 16] in { def : SPVBuiltin<"ConvertFToU_R" # IType.Name # v, [VectorType, VectorType], @@ -728,8 +730,10 @@ foreach IType = [UChar, UShort, UInt, ULong] in { foreach IType = [Char, Short, Int, Long] in { foreach FType = [Float, Double, Half] in { - def : SPVBuiltin<"ConvertFToS_R" # IType.Name, [IType, FType], Attr.Const>; def : SPVBuiltin<"ConvertSToF_R" # FType.Name, [FType, IType], Attr.Const>; + foreach rnd = ["", "_rte", "_rtz", "_rtp", "_rtn"] in { + def : SPVBuiltin<"ConvertFToS_R" # IType.Name # rnd, [IType, FType], Attr.Const>; + } foreach v = [2, 3, 4, 8, 16] in { def : SPVBuiltin<"ConvertFToS_R" # IType.Name # v, [VectorType, VectorType], diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp index 472ce566d377a..f9fc05e4ff6a1 100644 --- a/sycl/include/CL/sycl/types.hpp +++ b/sycl/include/CL/sycl/types.hpp @@ -231,10 +231,10 @@ convertImpl(T Value) { return static_cast(Value); } +#ifndef __SYCL_DEVICE_ONLY__ // float to int template detail::enable_if_t::value, R> convertImpl(T Value) { -#ifndef __SYCL_DEVICE_ONLY__ switch (roundingMode) { // Round to nearest even is default rounding mode for floating-point types case rounding_mode::automatic: @@ -264,11 +264,62 @@ detail::enable_if_t::value, R> convertImpl(T Value) { assert(!"Unsupported rounding mode!"); return static_cast(Value); }; -#else - // TODO implement device side conversion. - return static_cast(Value); -#endif } +#else + +template +using RteOrAutomatic = detail::bool_constant; + +template +using Rtz = detail::bool_constant; + +template +using Rtp = detail::bool_constant; + +template +using Rtn = detail::bool_constant; + +// Convert floating-point type to integer type +#define __SYCL_GENERATE_CONVERT_IMPL(SPIRVOp, DestType, RoundingMode, \ + RoundingModeCondition) \ + template \ + detail::enable_if_t::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); \ + } + +#define __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(RoundingMode, \ + RoundingModeCondition) \ + __SYCL_GENERATE_CONVERT_IMPL(FToS, int, RoundingMode, RoundingModeCondition) \ + __SYCL_GENERATE_CONVERT_IMPL(FToS, char, RoundingMode, \ + RoundingModeCondition) \ + __SYCL_GENERATE_CONVERT_IMPL(FToS, short, RoundingMode, \ + RoundingModeCondition) \ + __SYCL_GENERATE_CONVERT_IMPL(FToS, long, RoundingMode, \ + RoundingModeCondition) \ + __SYCL_GENERATE_CONVERT_IMPL(FToU, uint, RoundingMode, \ + RoundingModeCondition) \ + __SYCL_GENERATE_CONVERT_IMPL(FToU, uchar, RoundingMode, \ + RoundingModeCondition) \ + __SYCL_GENERATE_CONVERT_IMPL(FToU, ushort, RoundingMode, \ + RoundingModeCondition) \ + __SYCL_GENERATE_CONVERT_IMPL(FToU, ulong, 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 + +#endif // __SYCL_DEVICE_ONLY__ } // namespace detail diff --git a/sycl/test/basic_tests/vec_convert.cpp b/sycl/test/basic_tests/vec_convert.cpp index 9ba8cd68a5669..1e7235152e1da 100644 --- a/sycl/test/basic_tests/vec_convert.cpp +++ b/sycl/test/basic_tests/vec_convert.cpp @@ -1,8 +1,9 @@ +// XFAIL: cuda // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out // RUN: env SYCL_DEVICE_TYPE=HOST %t.out -// RUNx: %CPU_RUN_PLACEHOLDER %t.out -// RUNx: %GPU_RUN_PLACEHOLDER %t.out -// RUNx: %ACC_RUN_PLACEHOLDER %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. @@ -15,8 +16,7 @@ #include -// TODO uncomment run lines on non-host devices when the rounding modes will -// be implemented. +// TODO make the test to pass on cuda using namespace cl::sycl;