[SYCL] Added support of rounding modes for non-host devices

clang/lib/Sema/SPIRVBuiltins.td

@@ -711,31 +711,36 @@ foreach name = ["GenericPtrMemSemantics"] in {
 
 // 3.32.11. Conversion Instructions
 
-foreach IType = [UChar, UShort, UInt, ULong] in {
+foreach UIType = [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>;
+    def : SPVBuiltin<"ConvertFToU_R" # UIType.Name, [UIType, FType], Attr.Const>;
+    def : SPVBuiltin<"ConvertUToF_R" # FType.Name, [FType, UIType], Attr.Const>;
+    foreach rnd = ["", "_rte", "_rtz", "_rtp", "_rtn"] in {
+      def : SPVBuiltin<"ConvertFToU_R" # UIType.Name # rnd, [UIType, FType], Attr.Const>;
+    }
     foreach v = [2, 3, 4, 8, 16] in {
-      def : SPVBuiltin<"ConvertFToU_R" # IType.Name # v,
-                       [VectorType<IType, v>, VectorType<FType, v>],
+      def : SPVBuiltin<"ConvertFToU_R" # UIType.Name # v,
+                       [VectorType<UIType, v>, VectorType<FType, v>],
                        Attr.Const>;
       def : SPVBuiltin<"ConvertUToF_R" # FType.Name # v,
-                       [VectorType<FType, v>, VectorType<IType, v>],
+                       [VectorType<FType, v>, VectorType<UIType, v>],
                        Attr.Const>;
     }
   }
 }
 
-foreach IType = [Char, Short, Int, Long] in {
+foreach SIType = [TrueChar, 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>;
+    def : SPVBuiltin<"ConvertSToF_R" # FType.Name, [FType, SIType], Attr.Const>;
+    foreach rnd = ["", "_rte", "_rtz", "_rtp", "_rtn"] in {
+      def : SPVBuiltin<"ConvertFToS_R" # SIType.Name # rnd, [SIType, FType], Attr.Const>;
+    }
     foreach v = [2, 3, 4, 8, 16] in {
-      def : SPVBuiltin<"ConvertFToS_R" # IType.Name # v,
-                       [VectorType<IType, v>, VectorType<FType, v>],
+      def : SPVBuiltin<"ConvertFToS_R" # SIType.Name # v,
+                       [VectorType<SIType, v>, VectorType<FType, v>],
                        Attr.Const>;
       def : SPVBuiltin<"ConvertSToF_R" # FType.Name # v,
-                       [VectorType<FType, v>, VectorType<IType, v>],
+                       [VectorType<FType, v>, VectorType<SIType, v>],
                        Attr.Const>;
     }
   }

sycl/include/CL/sycl/types.hpp

@@ -219,6 +219,7 @@ detail::enable_if_t<std::is_same<T, R>::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 <typename T, typename R, rounding_mode roundingMode>
@@ -234,7 +235,6 @@ convertImpl(T Value) {
 // float to int
 template <typename T, typename R, rounding_mode roundingMode>
 detail::enable_if_t<is_float_to_int<T, R>::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,12 +264,65 @@ detail::enable_if_t<is_float_to_int<T, R>::value, R> convertImpl(T Value) {
     assert(!"Unsupported rounding mode!");
     return static_cast<R>(Value);
   };
+}
 #else
-  // TODO implement device side conversion.
+
+template <typename T, typename R, rounding_mode roundingMode>
+detail::enable_if_t<!std::is_same<T, R>::value &&
+                        (is_int_to_int<T, R>::value &&
+                         !(is_int_to_from_uint<T, R>::value) ||
+                         is_int_to_float<T, R>::value ||
+                         is_float_to_float<T, R>::value),
+                    R>
+convertImpl(T Value) {
   return static_cast<R>(Value);
-#endif
 }
 
+template <rounding_mode Mode>
+using RteOrAutomatic = detail::bool_constant<Mode == rounding_mode::automatic || Mode == rounding_mode::rte>;
+
+template <rounding_mode Mode>
+using Rtz = detail::bool_constant<Mode == rounding_mode::rtz>;
+
+template <rounding_mode Mode>
+using Rtp = detail::bool_constant<Mode == rounding_mode::rtp>;
+
+template <rounding_mode Mode>
+using Rtn = detail::bool_constant<Mode == rounding_mode::rtn>;
+
+#define __SYCL_GENERATE_CONVERT_IMPL(SPIRVOp, DestType, RoundingMode, RoundingModeCondition) \
+template <typename T, typename R, rounding_mode roundingMode>\
+detail::enable_if_t<is_float_to_int<T, R>::value && \
+                        std::is_same<R, DestType>::value && \
+                        RoundingModeCondition<roundingMode>::value, \
+                      R> \
+convertImpl(T Value) { \
+  using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t<T>; \
+  OpenCLT OpValue = cl::sycl::detail::convertDataToType<T, OpenCLT>(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
+
 } // namespace detail
 
 #if defined(_WIN32) && (_MSC_VER)

sycl/test/basic_tests/vec_convert.cpp

@@ -1,8 +1,8 @@
 // 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.
@@ -149,6 +149,6 @@ int main() {
   test<float, half, 8, rounding_mode::rtn>(
       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;
 }