Skip to content

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

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 37 commits into from
Apr 22, 2020
Merged

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

Show file tree
Hide file tree
Changes from 36 commits
Commits
Show all changes
37 commits
Select commit Hold shift + click to select a range
916dba8
support_rounding_models_for_non-host_devices #1.1
fadeeval Mar 13, 2020
0dba70d
support_rounding_models_for_non-host_devices #1.2
fadeeval Mar 16, 2020
0443893
support_rounding_models_for_non-host_devices #1.3
fadeeval Mar 16, 2020
cf7c3d4
Raw commit
fadeeval Mar 19, 2020
b4a9f93
Alpha version
fadeeval Mar 23, 2020
c671d9b
Alppha 2
fadeeval Mar 23, 2020
6c43eca
Origin
fadeeval Mar 24, 2020
bb88eff
Operative commit
fadeeval Mar 25, 2020
cf8e720
Fix the return type
fadeeval Mar 26, 2020
bbad4df
Converting SYCL types in OpenCL types
fadeeval Mar 26, 2020
b9a171e
Formatting and assert changes
fadeeval Mar 27, 2020
f927a3e
Fix of the spirv functions declaration
fadeeval Mar 30, 2020
2a96e9b
Metaprogramming
fadeeval Apr 1, 2020
a03b462
Adding macros
fadeeval Apr 1, 2020
070fbf5
Restore spirv_ops.hpp
fadeeval Apr 2, 2020
9c13847
Macros improving
fadeeval Apr 2, 2020
543fe09
Unsign to sign convert
fadeeval Apr 2, 2020
ec8806a
Reverting
fadeeval Apr 2, 2020
870059e
Patch
fadeeval Apr 3, 2020
ea768af
Formatting
fadeeval Apr 3, 2020
238ee1a
Formatting 2
fadeeval Apr 3, 2020
c03698b
Revert
fadeeval Apr 3, 2020
e0b0566
Formatting 3
fadeeval Apr 3, 2020
c5cf20e
Coments
fadeeval Apr 8, 2020
8267f30
Coments fix
fadeeval Apr 8, 2020
33220d7
Formatting 4
fadeeval Apr 8, 2020
5683577
Test fix
fadeeval Apr 10, 2020
5cd3e6e
Test fix2
fadeeval Apr 10, 2020
843bf4a
Cuda unsupp
fadeeval Apr 10, 2020
2a1fb94
delete unsupported: cuda
fadeeval Apr 13, 2020
8921bc3
vec_convert.cpp XFail: cuda
fadeeval Apr 13, 2020
f6d6777
Unsupported on cuda buffer_dev_to_dev.cpp
fadeeval Apr 14, 2020
e875fbf
TrueChar to Char
fadeeval Apr 14, 2020
f1f2897
schar to char
fadeeval Apr 14, 2020
c417b51
fix TrueChar
fadeeval Apr 15, 2020
014c0d4
Comments fix
fadeeval Apr 15, 2020
b0eb8e0
Fix
fadeeval Apr 15, 2020
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
12 changes: 8 additions & 4 deletions clang/lib/Sema/SPIRVBuiltins.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -286,8 +286,8 @@ class ConstOCLSPVBuiltin<string _Name, list<Type> _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>;
Expand Down Expand Up @@ -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<IType, v>, VectorType<FType, v>],
Expand All @@ -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<IType, v>, VectorType<FType, v>],
Expand Down
61 changes: 56 additions & 5 deletions sycl/include/CL/sycl/types.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -231,10 +231,10 @@ convertImpl(T Value) {
return static_cast<R>(Value);
}

#ifndef __SYCL_DEVICE_ONLY__
// 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:
Expand Down Expand Up @@ -264,11 +264,62 @@ 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.
return static_cast<R>(Value);
#endif
}
#else

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>;

// Convert floating-point type to integer type
#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 // __SYCL_DEVICE_ONLY__

} // namespace detail

Expand Down
3 changes: 3 additions & 0 deletions sycl/test/basic_tests/buffer/buffer_dev_to_dev.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,6 +4,9 @@
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: %ACC_RUN_PLACEHOLDER %t.out

// Unsupported: cuda
// The test fails sporadically on cuda. See https://github.com/intel/llvm/issues/1508 for more details.

Copy link
Contributor

@bader bader Apr 15, 2020

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Please, revert changes in this file.
d9b2d52 already disabled the test.

Copy link
Contributor Author

@fadeeval fadeeval Apr 15, 2020

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Do I need to do a rebase?

Copy link
Contributor

@bader bader Apr 15, 2020

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

No. CI system does the merge with tip of the sycl branch before testing.

//==---------- buffer_dev_to_dev.cpp - SYCL buffer basic test --------------==//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
Expand Down
12 changes: 6 additions & 6 deletions sycl/test/basic_tests/vec_convert.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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.
Expand All @@ -15,8 +16,7 @@

#include <cassert>

// 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;

Expand Down Expand Up @@ -151,4 +151,4 @@ int main() {
half8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f});

return 0;
}
}