From 916dba8e32c24808057ad62d46ff6c312f5ffb32 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Fri, 13 Mar 2020 15:13:13 +0300 Subject: [PATCH 01/37] support_rounding_models_for_non-host_devices #1.1 Signed-off-by: Aleksander Fadeev --- sycl/include/CL/sycl/types.hpp | 22 +++++++++++++++++++++- 1 file changed, 21 insertions(+), 1 deletion(-) diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp index 472ce566d377a..40590626fcba1 100644 --- a/sycl/include/CL/sycl/types.hpp +++ b/sycl/include/CL/sycl/types.hpp @@ -266,7 +266,27 @@ detail::enable_if_t::value, R> convertImpl(T Value) { }; #else // TODO implement device side conversion. - return static_cast(Value); + switch (roundingMode) { + // Round to nearest even is default rounding mode for floating-point types + case rounding_mode::automatic: + // Round to nearest even. + case rounding_mode::rte: { + R Result = std::rint(Value); + return Result; + } + // Round toward zero. + case rounding_mode::rtz: + return std::trunc(Value); + // Round toward positive infinity. + case rounding_mode::rtp: + return std::ceil(Value); + // Round toward negative infinity. + case rounding_mode::rtn: + return std::floor(Value); + default: + assert(!"Unsupported rounding mode!"); + return static_cast(Value); + }; #endif } From 0dba70d0708b954205fde0efe0a398cd03604616 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Mon, 16 Mar 2020 12:54:07 +0300 Subject: [PATCH 02/37] support_rounding_models_for_non-host_devices #1.2 Signed-off-by: Aleksander Fadeev --- sycl/include/CL/sycl/types.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp index 40590626fcba1..132b5098e443e 100644 --- a/sycl/include/CL/sycl/types.hpp +++ b/sycl/include/CL/sycl/types.hpp @@ -271,7 +271,7 @@ detail::enable_if_t::value, R> convertImpl(T Value) { case rounding_mode::automatic: // Round to nearest even. case rounding_mode::rte: { - R Result = std::rint(Value); + R Result = __spirv_OpConvertFToS_RScalarType_rte; return Result; } // Round toward zero. From 0443893dc8db62dfeafe7afbe7eff7e887de0467 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Mon, 16 Mar 2020 16:57:38 +0300 Subject: [PATCH 03/37] support_rounding_models_for_non-host_devices #1.3 Signed-off-by: Aleksander Fadeev --- sycl/include/CL/sycl/types.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp index 132b5098e443e..6e52c9dba4c96 100644 --- a/sycl/include/CL/sycl/types.hpp +++ b/sycl/include/CL/sycl/types.hpp @@ -271,7 +271,7 @@ detail::enable_if_t::value, R> convertImpl(T Value) { case rounding_mode::automatic: // Round to nearest even. case rounding_mode::rte: { - R Result = __spirv_OpConvertFToS_RScalarType_rte; + R Result = __spirv_ConvertFToS_Rint_rte(Value); return Result; } // Round toward zero. From cf7c3d4be4569e2a27039e7bfbcbbaa4a98d9aef Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Thu, 19 Mar 2020 13:47:56 +0300 Subject: [PATCH 04/37] Raw commit Signed-off-by: Aleksander Fadeev --- sycl/include/CL/sycl/types.hpp | 12 +++++------- 1 file changed, 5 insertions(+), 7 deletions(-) diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp index 6e52c9dba4c96..5bc0f177dfde5 100644 --- a/sycl/include/CL/sycl/types.hpp +++ b/sycl/include/CL/sycl/types.hpp @@ -270,19 +270,17 @@ detail::enable_if_t::value, R> convertImpl(T Value) { // Round to nearest even is default rounding mode for floating-point types case rounding_mode::automatic: // Round to nearest even. - case rounding_mode::rte: { - R Result = __spirv_ConvertFToS_Rint_rte(Value); - return Result; - } + case rounding_mode::rte: + return __spirv_ConvertFToS_Rint_rte(Value); // Round toward zero. case rounding_mode::rtz: - return std::trunc(Value); + return __spirv_ConvertFToS_Rint_rtz(Value); // Round toward positive infinity. case rounding_mode::rtp: - return std::ceil(Value); + return __spirv_ConvertFToS_Rint_rtp(Value); // Round toward negative infinity. case rounding_mode::rtn: - return std::floor(Value); + return __spirv_ConvertFToS_Rint_rtn(Value); default: assert(!"Unsupported rounding mode!"); return static_cast(Value); From b4a9f93f028bf32c301b31e229239a786288aabf Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Mon, 23 Mar 2020 15:13:17 +0300 Subject: [PATCH 05/37] Alpha version Signed-off-by: Aleksander Fadeev --- .../SPIRV/runtime/OpenCL/inc/spirv_convert.h | 2 +- sycl/include/CL/sycl/types.hpp | 22 +++---------------- sycl/test/basic_tests/vec_convert.cpp | 10 +++++++-- 3 files changed, 12 insertions(+), 22 deletions(-) diff --git a/llvm-spirv/lib/SPIRV/runtime/OpenCL/inc/spirv_convert.h b/llvm-spirv/lib/SPIRV/runtime/OpenCL/inc/spirv_convert.h index facd047dbee0b..1a7eed187b37f 100644 --- a/llvm-spirv/lib/SPIRV/runtime/OpenCL/inc/spirv_convert.h +++ b/llvm-spirv/lib/SPIRV/runtime/OpenCL/inc/spirv_convert.h @@ -1096,7 +1096,7 @@ __attribute__((overloadable)) int __spirv_ConvertFToS_Rint_sat_rtp(half); __attribute__((overloadable)) int __spirv_ConvertFToS_Rint_sat_rtn(half); __attribute__((overloadable)) int __spirv_ConvertFToS_Rint(float); __attribute__((overloadable)) int __spirv_ConvertFToS_Rint_rtz(float); -__attribute__((overloadable)) int __spirv_ConvertFToS_Rint_rte(float); +SYCL_EXTERNAL int __spirv_ConvertFToS_Rint_rte(float); __attribute__((overloadable)) int __spirv_ConvertFToS_Rint_rtp(float); __attribute__((overloadable)) int __spirv_ConvertFToS_Rint_rtn(float); __attribute__((overloadable)) int __spirv_ConvertFToS_Rint_sat(float); diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp index 5bc0f177dfde5..b3f09d5d1d43a 100644 --- a/sycl/include/CL/sycl/types.hpp +++ b/sycl/include/CL/sycl/types.hpp @@ -266,25 +266,9 @@ detail::enable_if_t::value, R> convertImpl(T Value) { }; #else // TODO implement device side conversion. - switch (roundingMode) { - // Round to nearest even is default rounding mode for floating-point types - case rounding_mode::automatic: - // Round to nearest even. - case rounding_mode::rte: - return __spirv_ConvertFToS_Rint_rte(Value); - // Round toward zero. - case rounding_mode::rtz: - return __spirv_ConvertFToS_Rint_rtz(Value); - // Round toward positive infinity. - case rounding_mode::rtp: - return __spirv_ConvertFToS_Rint_rtp(Value); - // Round toward negative infinity. - case rounding_mode::rtn: - return __spirv_ConvertFToS_Rint_rtn(Value); - default: - assert(!"Unsupported rounding mode!"); - return static_cast(Value); - }; +#include "/localdisk2/icl/fadeeval/sycl_workspace/llvm/llvm-spirv/lib/SPIRV/runtime/OpenCL/inc/spirv_convert.h" + R Result = __spirv_ConvertFToS_Rint_rte(Value); + return 1; #endif } diff --git a/sycl/test/basic_tests/vec_convert.cpp b/sycl/test/basic_tests/vec_convert.cpp index 9ba8cd68a5669..705845d8632da 100644 --- a/sycl/test/basic_tests/vec_convert.cpp +++ b/sycl/test/basic_tests/vec_convert.cpp @@ -30,7 +30,10 @@ template <> struct helper<0> { const vec &y) { const T xs = x.template swizzle<0>(); const T ys = y.template swizzle<0>(); - assert(xs == ys); + if (xs != ys) { + std::cerr << "sometihng failed"; + exit(1); + } } }; @@ -41,7 +44,10 @@ template struct helper { const T xs = x.template swizzle(); const T ys = y.template swizzle(); helper::compare(x, y); - assert(xs == ys); + if (xs != ys) { + std::cerr << "sometihng failed"; + exit(1); + } } }; From c671d9b008c37963573d1220447a85c87e70a8b0 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Mon, 23 Mar 2020 16:31:40 +0300 Subject: [PATCH 06/37] Alppha 2 Signed-off-by: Aleksander Fadeev --- sycl/include/CL/sycl/types.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp index b3f09d5d1d43a..5616e4683cda5 100644 --- a/sycl/include/CL/sycl/types.hpp +++ b/sycl/include/CL/sycl/types.hpp @@ -268,7 +268,7 @@ detail::enable_if_t::value, R> convertImpl(T Value) { // TODO implement device side conversion. #include "/localdisk2/icl/fadeeval/sycl_workspace/llvm/llvm-spirv/lib/SPIRV/runtime/OpenCL/inc/spirv_convert.h" R Result = __spirv_ConvertFToS_Rint_rte(Value); - return 1; + return Result; #endif } From 6c43ecacb6d14a2452f3dbb54b3f7c30e7e9b5f7 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Tue, 24 Mar 2020 13:16:33 +0300 Subject: [PATCH 07/37] Origin Signed-off-by: Aleksander Fadeev --- .../SPIRV/runtime/OpenCL/inc/spirv_convert.h | 6 +++--- sycl/include/CL/sycl/types.hpp | 21 +++++++++++++++++-- sycl/test/basic_tests/vec_convert.cpp | 2 +- 3 files changed, 23 insertions(+), 6 deletions(-) diff --git a/llvm-spirv/lib/SPIRV/runtime/OpenCL/inc/spirv_convert.h b/llvm-spirv/lib/SPIRV/runtime/OpenCL/inc/spirv_convert.h index 1a7eed187b37f..34d3a24d754f3 100644 --- a/llvm-spirv/lib/SPIRV/runtime/OpenCL/inc/spirv_convert.h +++ b/llvm-spirv/lib/SPIRV/runtime/OpenCL/inc/spirv_convert.h @@ -1095,10 +1095,10 @@ __attribute__((overloadable)) int __spirv_ConvertFToS_Rint_sat_rte(half); __attribute__((overloadable)) int __spirv_ConvertFToS_Rint_sat_rtp(half); __attribute__((overloadable)) int __spirv_ConvertFToS_Rint_sat_rtn(half); __attribute__((overloadable)) int __spirv_ConvertFToS_Rint(float); -__attribute__((overloadable)) int __spirv_ConvertFToS_Rint_rtz(float); +SYCL_EXTERNAL int __spirv_ConvertFToS_Rint_rtz(float); SYCL_EXTERNAL int __spirv_ConvertFToS_Rint_rte(float); -__attribute__((overloadable)) int __spirv_ConvertFToS_Rint_rtp(float); -__attribute__((overloadable)) int __spirv_ConvertFToS_Rint_rtn(float); +SYCL_EXTERNAL int __spirv_ConvertFToS_Rint_rtp(float); +SYCL_EXTERNAL int __spirv_ConvertFToS_Rint_rtn(float); __attribute__((overloadable)) int __spirv_ConvertFToS_Rint_sat(float); __attribute__((overloadable)) int __spirv_ConvertFToS_Rint_sat_rtz(float); __attribute__((overloadable)) int __spirv_ConvertFToS_Rint_sat_rte(float); diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp index 5616e4683cda5..af7c15fa3350f 100644 --- a/sycl/include/CL/sycl/types.hpp +++ b/sycl/include/CL/sycl/types.hpp @@ -267,8 +267,25 @@ detail::enable_if_t::value, R> convertImpl(T Value) { #else // TODO implement device side conversion. #include "/localdisk2/icl/fadeeval/sycl_workspace/llvm/llvm-spirv/lib/SPIRV/runtime/OpenCL/inc/spirv_convert.h" - R Result = __spirv_ConvertFToS_Rint_rte(Value); - return Result; + switch (roundingMode) { + // Round to nearest even is default rounding mode for floating-point types + case rounding_mode::automatic: + // Round to nearest even. + case rounding_mode::rte: + return __spirv_ConvertFToS_Rint_rte(Value); + // Round toward zero. + case rounding_mode::rtz: + return __spirv_ConvertFToS_Rint_rtz(Value); + // Round toward positive infinity. + case rounding_mode::rtp: + return __spirv_ConvertFToS_Rint_rtp(Value); + // Round toward negative infinity. + case rounding_mode::rtn: + return __spirv_ConvertFToS_Rint_rtn(Value); + default: + assert(!"Unsupported rounding mode!"); + return static_cast(Value); + }; #endif } diff --git a/sycl/test/basic_tests/vec_convert.cpp b/sycl/test/basic_tests/vec_convert.cpp index 705845d8632da..abd069a738eb9 100644 --- a/sycl/test/basic_tests/vec_convert.cpp +++ b/sycl/test/basic_tests/vec_convert.cpp @@ -1,5 +1,5 @@ // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out -// RUN: env SYCL_DEVICE_TYPE=HOST %t.out +// RUN: env SYCL_DEVICE_TYPE=CPU %t.out // RUNx: %CPU_RUN_PLACEHOLDER %t.out // RUNx: %GPU_RUN_PLACEHOLDER %t.out // RUNx: %ACC_RUN_PLACEHOLDER %t.out From bb88eff963ae2e22645b9bbaf93de2b57c6761d3 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Wed, 25 Mar 2020 15:18:03 +0300 Subject: [PATCH 08/37] Operative commit Signed-off-by: Aleksander Fadeev --- .../lib/SPIRV/runtime/OpenCL/inc/spirv_convert.h | 8 ++++---- sycl/include/CL/__spirv/spirv_ops.hpp | 11 +++++++++++ sycl/include/CL/sycl/types.hpp | 2 +- sycl/test/basic_tests/vec_convert.cpp | 8 ++++---- 4 files changed, 20 insertions(+), 9 deletions(-) diff --git a/llvm-spirv/lib/SPIRV/runtime/OpenCL/inc/spirv_convert.h b/llvm-spirv/lib/SPIRV/runtime/OpenCL/inc/spirv_convert.h index 34d3a24d754f3..facd047dbee0b 100644 --- a/llvm-spirv/lib/SPIRV/runtime/OpenCL/inc/spirv_convert.h +++ b/llvm-spirv/lib/SPIRV/runtime/OpenCL/inc/spirv_convert.h @@ -1095,10 +1095,10 @@ __attribute__((overloadable)) int __spirv_ConvertFToS_Rint_sat_rte(half); __attribute__((overloadable)) int __spirv_ConvertFToS_Rint_sat_rtp(half); __attribute__((overloadable)) int __spirv_ConvertFToS_Rint_sat_rtn(half); __attribute__((overloadable)) int __spirv_ConvertFToS_Rint(float); -SYCL_EXTERNAL int __spirv_ConvertFToS_Rint_rtz(float); -SYCL_EXTERNAL int __spirv_ConvertFToS_Rint_rte(float); -SYCL_EXTERNAL int __spirv_ConvertFToS_Rint_rtp(float); -SYCL_EXTERNAL int __spirv_ConvertFToS_Rint_rtn(float); +__attribute__((overloadable)) int __spirv_ConvertFToS_Rint_rtz(float); +__attribute__((overloadable)) int __spirv_ConvertFToS_Rint_rte(float); +__attribute__((overloadable)) int __spirv_ConvertFToS_Rint_rtp(float); +__attribute__((overloadable)) int __spirv_ConvertFToS_Rint_rtn(float); __attribute__((overloadable)) int __spirv_ConvertFToS_Rint_sat(float); __attribute__((overloadable)) int __spirv_ConvertFToS_Rint_sat_rtz(float); __attribute__((overloadable)) int __spirv_ConvertFToS_Rint_sat_rte(float); diff --git a/sycl/include/CL/__spirv/spirv_ops.hpp b/sycl/include/CL/__spirv/spirv_ops.hpp index 1d758901d05f7..3e02252a63764 100644 --- a/sycl/include/CL/__spirv/spirv_ops.hpp +++ b/sycl/include/CL/__spirv/spirv_ops.hpp @@ -49,6 +49,17 @@ template extern SYCL_EXTERNAL TempRetT __spirv_ImageSampleExplicitLod(SampledType, TempArgT, int, float); +template + extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rint_rte(TValue Value); + +template + extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rint_rtz(TValue Value); + +template + extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rint_rtp(TValue Value); + +template + extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rint_rtn(TValue Value); #define OpGroupAsyncCopyGlobalToLocal __spirv_GroupAsyncCopy #define OpGroupAsyncCopyLocalToGlobal __spirv_GroupAsyncCopy diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp index af7c15fa3350f..6303085f3f5f4 100644 --- a/sycl/include/CL/sycl/types.hpp +++ b/sycl/include/CL/sycl/types.hpp @@ -51,6 +51,7 @@ #include #include #include +#include "CL/__spirv/spirv_ops.hpp" #include #include @@ -266,7 +267,6 @@ detail::enable_if_t::value, R> convertImpl(T Value) { }; #else // TODO implement device side conversion. -#include "/localdisk2/icl/fadeeval/sycl_workspace/llvm/llvm-spirv/lib/SPIRV/runtime/OpenCL/inc/spirv_convert.h" switch (roundingMode) { // Round to nearest even is default rounding mode for floating-point types case rounding_mode::automatic: diff --git a/sycl/test/basic_tests/vec_convert.cpp b/sycl/test/basic_tests/vec_convert.cpp index abd069a738eb9..db689b488551b 100644 --- a/sycl/test/basic_tests/vec_convert.cpp +++ b/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=CPU %t.out -// RUNx: %CPU_RUN_PLACEHOLDER %t.out -// RUNx: %GPU_RUN_PLACEHOLDER %t.out -// RUNx: %ACC_RUN_PLACEHOLDER %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. From cf8e720d91ec9d16f07ad7335c8ce716e2c9d701 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Thu, 26 Mar 2020 14:09:38 +0300 Subject: [PATCH 09/37] Fix the return type Signed-off-by: Aleksander Fadeev --- sycl/include/CL/__spirv/spirv_ops.hpp | 45 ++++++++++++++++++++++++--- sycl/include/CL/sycl/types.hpp | 31 ++++++++++++++++-- 2 files changed, 69 insertions(+), 7 deletions(-) diff --git a/sycl/include/CL/__spirv/spirv_ops.hpp b/sycl/include/CL/__spirv/spirv_ops.hpp index 3e02252a63764..2a12b44729910 100644 --- a/sycl/include/CL/__spirv/spirv_ops.hpp +++ b/sycl/include/CL/__spirv/spirv_ops.hpp @@ -49,17 +49,54 @@ template extern SYCL_EXTERNAL TempRetT __spirv_ImageSampleExplicitLod(SampledType, TempArgT, int, float); + +template +extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rint_rte(TValue Value); + +template +extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rchar_rte(TValue Value); + +template +extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rshort_rte(TValue Value); + +template +extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rlong_rte(TValue Value); + +template +extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rint_rtz(TValue Value); + +template +extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rchar_rtz(TValue Value); + +template +extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rshort_rtz(TValue Value); + +template +extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rlong_rtz(TValue Value); + +template +extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rint_rtp(TValue Value); + +template +extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rchar_rtp(TValue Value); + +template +extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rshort_rtp(TValue Value); + +template +extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rlong_rtp(TValue Value); + template - extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rint_rte(TValue Value); +extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rint_rtn(TValue Value); template - extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rint_rtz(TValue Value); +extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rchar_rtn(TValue Value); template - extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rint_rtp(TValue Value); +extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rshort_rtn(TValue Value); template - extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rint_rtn(TValue Value); +extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rlong_rtn(TValue Value); #define OpGroupAsyncCopyGlobalToLocal __spirv_GroupAsyncCopy #define OpGroupAsyncCopyLocalToGlobal __spirv_GroupAsyncCopy diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp index 6303085f3f5f4..b2fe53cb6d269 100644 --- a/sycl/include/CL/sycl/types.hpp +++ b/sycl/include/CL/sycl/types.hpp @@ -272,19 +272,44 @@ detail::enable_if_t::value, R> convertImpl(T Value) { case rounding_mode::automatic: // Round to nearest even. case rounding_mode::rte: + if (std::is_same::value) return __spirv_ConvertFToS_Rint_rte(Value); + else if (std::is_same::value) + return __spirv_ConvertFToS_Rchar_rte(Value); + else if (std::is_same::value) + return __spirv_ConvertFToS_Rshort_rte(Value); + else if (std::is_same::value) + return __spirv_ConvertFToS_Rlong_rte(Value); // Round toward zero. case rounding_mode::rtz: + if (std::is_same::value) return __spirv_ConvertFToS_Rint_rtz(Value); + else if (std::is_same::value) + return __spirv_ConvertFToS_Rchar_rtz(Value); + else if (std::is_same::value) + return __spirv_ConvertFToS_Rshort_rtz(Value); + else if (std::is_same::value) + return __spirv_ConvertFToS_Rlong_rtz(Value); // Round toward positive infinity. case rounding_mode::rtp: + if (std::is_same::value) return __spirv_ConvertFToS_Rint_rtp(Value); + else if (std::is_same::value) + return __spirv_ConvertFToS_Rchar_rtp(Value); + else if (std::is_same::value) + return __spirv_ConvertFToS_Rshort_rtp(Value); + else if (std::is_same::value) + return __spirv_ConvertFToS_Rlong_rtp(Value); // Round toward negative infinity. case rounding_mode::rtn: + if (std::is_same::value) return __spirv_ConvertFToS_Rint_rtn(Value); - default: - assert(!"Unsupported rounding mode!"); - return static_cast(Value); + else if (std::is_same::value) + return __spirv_ConvertFToS_Rchar_rtn(Value); + else if (std::is_same::value) + return __spirv_ConvertFToS_Rshort_rtn(Value); + else if (std::is_same::value) + return __spirv_ConvertFToS_Rlong_rtn(Value); }; #endif } From bbad4dfb734acce3acca7f29ca47ce3705474ffd Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Thu, 26 Mar 2020 16:32:46 +0300 Subject: [PATCH 10/37] Converting SYCL types in OpenCL types Signed-off-by: Aleksander Fadeev --- sycl/include/CL/sycl/types.hpp | 44 ++++++++++++++++++---------------- 1 file changed, 23 insertions(+), 21 deletions(-) diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp index b2fe53cb6d269..905aaae219e42 100644 --- a/sycl/include/CL/sycl/types.hpp +++ b/sycl/include/CL/sycl/types.hpp @@ -267,49 +267,51 @@ detail::enable_if_t::value, R> convertImpl(T Value) { }; #else // TODO implement device side conversion. + using OpenCLT= cl::sycl::detail::ConvertToOpenCLType_t; + OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); switch (roundingMode) { // Round to nearest even is default rounding mode for floating-point types case rounding_mode::automatic: // Round to nearest even. - case rounding_mode::rte: - if (std::is_same::value) - return __spirv_ConvertFToS_Rint_rte(Value); + case rounding_mode::rte: + if (std::is_same::value) + return __spirv_ConvertFToS_Rint_rte(OpValue); else if (std::is_same::value) - return __spirv_ConvertFToS_Rchar_rte(Value); + return __spirv_ConvertFToS_Rchar_rte(OpValue); else if (std::is_same::value) - return __spirv_ConvertFToS_Rshort_rte(Value); + return __spirv_ConvertFToS_Rshort_rte(OpValue); else if (std::is_same::value) - return __spirv_ConvertFToS_Rlong_rte(Value); + return __spirv_ConvertFToS_Rlong_rte(OpValue); // Round toward zero. case rounding_mode::rtz: - if (std::is_same::value) - return __spirv_ConvertFToS_Rint_rtz(Value); + if (std::is_same::value) + return __spirv_ConvertFToS_Rint_rtz(OpValue); else if (std::is_same::value) - return __spirv_ConvertFToS_Rchar_rtz(Value); + return __spirv_ConvertFToS_Rchar_rtz(OpValue); else if (std::is_same::value) - return __spirv_ConvertFToS_Rshort_rtz(Value); + return __spirv_ConvertFToS_Rshort_rtz(OpValue); else if (std::is_same::value) - return __spirv_ConvertFToS_Rlong_rtz(Value); + return __spirv_ConvertFToS_Rlong_rtz(OpValue); // Round toward positive infinity. case rounding_mode::rtp: - if (std::is_same::value) - return __spirv_ConvertFToS_Rint_rtp(Value); + if (std::is_same::value) + return __spirv_ConvertFToS_Rint_rtp(OpValue); else if (std::is_same::value) - return __spirv_ConvertFToS_Rchar_rtp(Value); + return __spirv_ConvertFToS_Rchar_rtp(OpValue); else if (std::is_same::value) - return __spirv_ConvertFToS_Rshort_rtp(Value); + return __spirv_ConvertFToS_Rshort_rtp(OpValue); else if (std::is_same::value) - return __spirv_ConvertFToS_Rlong_rtp(Value); + return __spirv_ConvertFToS_Rlong_rtp(OpValue); // Round toward negative infinity. case rounding_mode::rtn: - if (std::is_same::value) - return __spirv_ConvertFToS_Rint_rtn(Value); + if (std::is_same::value) + return __spirv_ConvertFToS_Rint_rtn(OpValue); else if (std::is_same::value) - return __spirv_ConvertFToS_Rchar_rtn(Value); + return __spirv_ConvertFToS_Rchar_rtn(OpValue); else if (std::is_same::value) - return __spirv_ConvertFToS_Rshort_rtn(Value); + return __spirv_ConvertFToS_Rshort_rtn(OpValue); else if (std::is_same::value) - return __spirv_ConvertFToS_Rlong_rtn(Value); + return __spirv_ConvertFToS_Rlong_rtn(OpValue); }; #endif } From b9a171ed533ac5c9b1f8a1ece3eaf2aa7ff03574 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Fri, 27 Mar 2020 13:39:12 +0300 Subject: [PATCH 11/37] Formatting and assert changes Signed-off-by: Aleksander Fadeev --- sycl/include/CL/sycl/types.hpp | 4 ++-- sycl/test/basic_tests/vec_convert.cpp | 10 ++-------- 2 files changed, 4 insertions(+), 10 deletions(-) diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp index 905aaae219e42..19ef619b2d2cd 100644 --- a/sycl/include/CL/sycl/types.hpp +++ b/sycl/include/CL/sycl/types.hpp @@ -267,8 +267,8 @@ detail::enable_if_t::value, R> convertImpl(T Value) { }; #else // TODO implement device side conversion. - using OpenCLT= cl::sycl::detail::ConvertToOpenCLType_t; - OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); + using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; + OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); switch (roundingMode) { // Round to nearest even is default rounding mode for floating-point types case rounding_mode::automatic: diff --git a/sycl/test/basic_tests/vec_convert.cpp b/sycl/test/basic_tests/vec_convert.cpp index db689b488551b..21e6249b1adc5 100644 --- a/sycl/test/basic_tests/vec_convert.cpp +++ b/sycl/test/basic_tests/vec_convert.cpp @@ -30,10 +30,7 @@ template <> struct helper<0> { const vec &y) { const T xs = x.template swizzle<0>(); const T ys = y.template swizzle<0>(); - if (xs != ys) { - std::cerr << "sometihng failed"; - exit(1); - } + assert(xs == ys); } }; @@ -44,10 +41,7 @@ template struct helper { const T xs = x.template swizzle(); const T ys = y.template swizzle(); helper::compare(x, y); - if (xs != ys) { - std::cerr << "sometihng failed"; - exit(1); - } + assert(xs == ys); } }; From f927a3e0b7a014a586824d9df9dcef362559fda7 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Mon, 30 Mar 2020 14:42:09 +0300 Subject: [PATCH 12/37] Fix of the spirv functions declaration Signed-off-by: Aleksander Fadeev --- clang/lib/Sema/SPIRVBuiltins.td | 9 +++- sycl/include/CL/__spirv/spirv_ops.hpp | 71 --------------------------- sycl/include/CL/sycl/types.hpp | 1 - 3 files changed, 7 insertions(+), 74 deletions(-) diff --git a/clang/lib/Sema/SPIRVBuiltins.td b/clang/lib/Sema/SPIRVBuiltins.td index 69c8f7b82c54e..e592413705479 100644 --- a/clang/lib/Sema/SPIRVBuiltins.td +++ b/clang/lib/Sema/SPIRVBuiltins.td @@ -715,6 +715,9 @@ 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], @@ -726,10 +729,12 @@ foreach IType = [UChar, UShort, UInt, ULong] in { } } -foreach IType = [Char, Short, Int, Long] in { +foreach IType = [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>; + 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/__spirv/spirv_ops.hpp b/sycl/include/CL/__spirv/spirv_ops.hpp index 2a12b44729910..3173741b12c53 100644 --- a/sycl/include/CL/__spirv/spirv_ops.hpp +++ b/sycl/include/CL/__spirv/spirv_ops.hpp @@ -27,77 +27,6 @@ "SPIR-V built-ins are not available. Please set -fdeclare-spirv-builtins flag." #endif -template -extern SYCL_EXTERNAL RetT __spirv_ImageQueryFormat(ImageT); - -template -extern SYCL_EXTERNAL RetT __spirv_ImageQueryOrder(ImageT); - -template -extern SYCL_EXTERNAL RetT __spirv_ImageQuerySize(ImageT); - -template -extern SYCL_EXTERNAL void __spirv_ImageWrite(ImageT, CoordT, ValT); - -template -extern SYCL_EXTERNAL RetT __spirv_ImageRead(ImageT, TempArgT); - -template -extern SYCL_EXTERNAL SampledType __spirv_SampledImage(ImageT, __ocl_sampler_t); - -template -extern SYCL_EXTERNAL TempRetT __spirv_ImageSampleExplicitLod(SampledType, - TempArgT, int, - float); - -template -extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rint_rte(TValue Value); - -template -extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rchar_rte(TValue Value); - -template -extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rshort_rte(TValue Value); - -template -extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rlong_rte(TValue Value); - -template -extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rint_rtz(TValue Value); - -template -extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rchar_rtz(TValue Value); - -template -extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rshort_rtz(TValue Value); - -template -extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rlong_rtz(TValue Value); - -template -extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rint_rtp(TValue Value); - -template -extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rchar_rtp(TValue Value); - -template -extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rshort_rtp(TValue Value); - -template -extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rlong_rtp(TValue Value); - -template -extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rint_rtn(TValue Value); - -template -extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rchar_rtn(TValue Value); - -template -extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rshort_rtn(TValue Value); - -template -extern SYCL_EXTERNAL int __spirv_ConvertFToS_Rlong_rtn(TValue Value); - #define OpGroupAsyncCopyGlobalToLocal __spirv_GroupAsyncCopy #define OpGroupAsyncCopyLocalToGlobal __spirv_GroupAsyncCopy diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp index 19ef619b2d2cd..856878c9a6ed0 100644 --- a/sycl/include/CL/sycl/types.hpp +++ b/sycl/include/CL/sycl/types.hpp @@ -51,7 +51,6 @@ #include #include #include -#include "CL/__spirv/spirv_ops.hpp" #include #include From 2a96e9b11ec239841a6f6f0769e46daf10fcef1a Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Wed, 1 Apr 2020 11:37:24 +0300 Subject: [PATCH 13/37] Metaprogramming Signed-off-by: Aleksander Fadeev --- sycl/include/CL/sycl/types.hpp | 227 ++++++++++++++++++++++++++------- 1 file changed, 180 insertions(+), 47 deletions(-) diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp index 856878c9a6ed0..ef4baa221d80e 100644 --- a/sycl/include/CL/sycl/types.hpp +++ b/sycl/include/CL/sycl/types.hpp @@ -232,9 +232,9 @@ convertImpl(T Value) { } // float to int +#ifndef __SYCL_DEVICE_ONLY__ 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,57 +264,190 @@ detail::enable_if_t::value, R> convertImpl(T Value) { assert(!"Unsupported rounding mode!"); return static_cast(Value); }; +} #else - // TODO implement device side conversion. + +template +detail::enable_if_t::value && + std::is_same::value && + (roundingMode == rounding_mode::automatic || + roundingMode == rounding_mode::rte), + R> +convertImpl(T Value) { using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); - switch (roundingMode) { - // Round to nearest even is default rounding mode for floating-point types - case rounding_mode::automatic: - // Round to nearest even. - case rounding_mode::rte: - if (std::is_same::value) - return __spirv_ConvertFToS_Rint_rte(OpValue); - else if (std::is_same::value) - return __spirv_ConvertFToS_Rchar_rte(OpValue); - else if (std::is_same::value) - return __spirv_ConvertFToS_Rshort_rte(OpValue); - else if (std::is_same::value) - return __spirv_ConvertFToS_Rlong_rte(OpValue); - // Round toward zero. - case rounding_mode::rtz: - if (std::is_same::value) - return __spirv_ConvertFToS_Rint_rtz(OpValue); - else if (std::is_same::value) - return __spirv_ConvertFToS_Rchar_rtz(OpValue); - else if (std::is_same::value) - return __spirv_ConvertFToS_Rshort_rtz(OpValue); - else if (std::is_same::value) - return __spirv_ConvertFToS_Rlong_rtz(OpValue); - // Round toward positive infinity. - case rounding_mode::rtp: - if (std::is_same::value) - return __spirv_ConvertFToS_Rint_rtp(OpValue); - else if (std::is_same::value) - return __spirv_ConvertFToS_Rchar_rtp(OpValue); - else if (std::is_same::value) - return __spirv_ConvertFToS_Rshort_rtp(OpValue); - else if (std::is_same::value) - return __spirv_ConvertFToS_Rlong_rtp(OpValue); - // Round toward negative infinity. - case rounding_mode::rtn: - if (std::is_same::value) - return __spirv_ConvertFToS_Rint_rtn(OpValue); - else if (std::is_same::value) - return __spirv_ConvertFToS_Rchar_rtn(OpValue); - else if (std::is_same::value) - return __spirv_ConvertFToS_Rshort_rtn(OpValue); - else if (std::is_same::value) - return __spirv_ConvertFToS_Rlong_rtn(OpValue); - }; -#endif + return __spirv_ConvertFToS_Rint_rte(OpValue); +} + +template +detail::enable_if_t::value && + std::is_same::value && + roundingMode == rounding_mode::rtz, + R> +convertImpl(T Value) { + using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; + OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); + return __spirv_ConvertFToS_Rint_rtz(OpValue); } +template +detail::enable_if_t::value && + std::is_same::value && + roundingMode == rounding_mode::rtp, + R> +convertImpl(T Value) { + using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; + OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); + return __spirv_ConvertFToS_Rint_rtp(OpValue); +} + +template +detail::enable_if_t::value && + std::is_same::value && + roundingMode == rounding_mode::rtn, + R> +convertImpl(T Value) { + using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; + OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); + return __spirv_ConvertFToS_Rint_rtn(OpValue); +} + +template +detail::enable_if_t::value && + std::is_same::value && + (roundingMode == rounding_mode::automatic || + roundingMode == rounding_mode::rte), + R> +convertImpl(T Value) { + using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; + OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); + return __spirv_ConvertFToS_Rchar_rte(OpValue); +} + +template +detail::enable_if_t::value && + std::is_same::value && + roundingMode == rounding_mode::rtz, + R> +convertImpl(T Value) { + using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; + OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); + return __spirv_ConvertFToS_Rchar_rtz(OpValue); +} + +template +detail::enable_if_t::value && + std::is_same::value && + roundingMode == rounding_mode::rtp, + R> +convertImpl(T Value) { + using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; + OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); + return __spirv_ConvertFToS_Rchar_rtp(OpValue); +} + +template +detail::enable_if_t::value && + std::is_same::value && + roundingMode == rounding_mode::rtn, + R> +convertImpl(T Value) { + using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; + OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); + return __spirv_ConvertFToS_Rchar_rtn(OpValue); +} + +template +detail::enable_if_t::value && + std::is_same::value && + (roundingMode == rounding_mode::automatic || + roundingMode == rounding_mode::rte), + R> +convertImpl(T Value) { + using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; + OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); + return __spirv_ConvertFToS_Rshort_rte(OpValue); +} + +template +detail::enable_if_t::value && + std::is_same::value && + roundingMode == rounding_mode::rtz, + R> +convertImpl(T Value) { + using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; + OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); + return __spirv_ConvertFToS_Rshort_rtz(OpValue); +} + +template +detail::enable_if_t::value && + std::is_same::value && + roundingMode == rounding_mode::rtp, + R> +convertImpl(T Value) { + using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; + OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); + return __spirv_ConvertFToS_Rshort_rtp(OpValue); +} + +template +detail::enable_if_t::value && + std::is_same::value && + roundingMode == rounding_mode::rtn, + R> +convertImpl(T Value) { + using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; + OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); + return __spirv_ConvertFToS_Rshort_rtn(OpValue); +} + +template +detail::enable_if_t::value && + std::is_same::value && + (roundingMode == rounding_mode::automatic || + roundingMode == rounding_mode::rte), + R> +convertImpl(T Value) { + using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; + OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); + return __spirv_ConvertFToS_Rint_rte(OpValue); +} + +template +detail::enable_if_t::value && + std::is_same::value && + roundingMode == rounding_mode::rtz, + R> +convertImpl(T Value) { + using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; + OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); + return __spirv_ConvertFToS_Rlong_rtz(OpValue); +} + +template +detail::enable_if_t::value && + std::is_same::value && + roundingMode == rounding_mode::rtp, + R> +convertImpl(T Value) { + using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; + OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); + return __spirv_ConvertFToS_Rlong_rtp(OpValue); +} + +template +detail::enable_if_t::value && + std::is_same::value && + roundingMode == rounding_mode::rtn, + R> +convertImpl(T Value) { + using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; + OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); + return __spirv_ConvertFToS_Rlong_rtn(OpValue); +} +#endif + } // namespace detail #if defined(_WIN32) && (_MSC_VER) From a03b462c6eb89c4da01669461787c6109ecbf3e6 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Wed, 1 Apr 2020 18:34:34 +0300 Subject: [PATCH 14/37] Adding macros Signed-off-by: Aleksander Fadeev --- sycl/include/CL/sycl/types.hpp | 209 +++++++-------------------------- 1 file changed, 45 insertions(+), 164 deletions(-) diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp index ef4baa221d80e..98295c9ab09b6 100644 --- a/sycl/include/CL/sycl/types.hpp +++ b/sycl/include/CL/sycl/types.hpp @@ -267,185 +267,66 @@ detail::enable_if_t::value, R> convertImpl(T Value) { } #else -template -detail::enable_if_t::value && - std::is_same::value && - (roundingMode == rounding_mode::automatic || - roundingMode == rounding_mode::rte), - R> -convertImpl(T Value) { - using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; - OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); - return __spirv_ConvertFToS_Rint_rte(OpValue); -} -template -detail::enable_if_t::value && - std::is_same::value && - roundingMode == rounding_mode::rtz, - R> -convertImpl(T Value) { - using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; - OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); - return __spirv_ConvertFToS_Rint_rtz(OpValue); -} +template +using RteOrAutomatic = detail::bool_constant; -template -detail::enable_if_t::value && - std::is_same::value && - roundingMode == rounding_mode::rtp, - R> -convertImpl(T Value) { - using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; - OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); - return __spirv_ConvertFToS_Rint_rtp(OpValue); -} +template +using Rtz = detail::bool_constant; -template -detail::enable_if_t::value && - std::is_same::value && - roundingMode == rounding_mode::rtn, - R> -convertImpl(T Value) { - using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; - OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); - return __spirv_ConvertFToS_Rint_rtn(OpValue); -} +template +using Rtp = detail::bool_constant; -template -detail::enable_if_t::value && - std::is_same::value && - (roundingMode == rounding_mode::automatic || - roundingMode == rounding_mode::rte), - R> -convertImpl(T Value) { - using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; - OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); - return __spirv_ConvertFToS_Rchar_rte(OpValue); -} +template +using Rtn = detail::bool_constant; -template -detail::enable_if_t::value && - std::is_same::value && - roundingMode == rounding_mode::rtz, - R> -convertImpl(T Value) { - using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; - OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); - return __spirv_ConvertFToS_Rchar_rtz(OpValue); -} -template -detail::enable_if_t::value && - std::is_same::value && - roundingMode == rounding_mode::rtp, - R> -convertImpl(T Value) { - using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; - OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); - return __spirv_ConvertFToS_Rchar_rtp(OpValue); -} +#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); \ +} \ -template -detail::enable_if_t::value && - std::is_same::value && - roundingMode == rounding_mode::rtn, - R> -convertImpl(T Value) { - using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; - OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); - return __spirv_ConvertFToS_Rchar_rtn(OpValue); -} +__SYCL_GENERATE_CONVERT_IMPL(FToS, int, rte, RteOrAutomatic) -template -detail::enable_if_t::value && - std::is_same::value && - (roundingMode == rounding_mode::automatic || - roundingMode == rounding_mode::rte), - R> -convertImpl(T Value) { - using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; - OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); - return __spirv_ConvertFToS_Rshort_rte(OpValue); -} +__SYCL_GENERATE_CONVERT_IMPL(FToU, char, rte, RteOrAutomatic) -template -detail::enable_if_t::value && - std::is_same::value && - roundingMode == rounding_mode::rtz, - R> -convertImpl(T Value) { - using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; - OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); - return __spirv_ConvertFToS_Rshort_rtz(OpValue); -} +__SYCL_GENERATE_CONVERT_IMPL(FToU, short, rte, RteOrAutomatic) -template -detail::enable_if_t::value && - std::is_same::value && - roundingMode == rounding_mode::rtp, - R> -convertImpl(T Value) { - using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; - OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); - return __spirv_ConvertFToS_Rshort_rtp(OpValue); -} +__SYCL_GENERATE_CONVERT_IMPL(FToU, long, rte, RteOrAutomatic) -template -detail::enable_if_t::value && - std::is_same::value && - roundingMode == rounding_mode::rtn, - R> -convertImpl(T Value) { - using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; - OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); - return __spirv_ConvertFToS_Rshort_rtn(OpValue); -} +__SYCL_GENERATE_CONVERT_IMPL(FToS, int, rtz, Rtz) -template -detail::enable_if_t::value && - std::is_same::value && - (roundingMode == rounding_mode::automatic || - roundingMode == rounding_mode::rte), - R> -convertImpl(T Value) { - using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; - OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); - return __spirv_ConvertFToS_Rint_rte(OpValue); -} +__SYCL_GENERATE_CONVERT_IMPL(FToS, char, rtz, Rtz) -template -detail::enable_if_t::value && - std::is_same::value && - roundingMode == rounding_mode::rtz, - R> -convertImpl(T Value) { - using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; - OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); - return __spirv_ConvertFToS_Rlong_rtz(OpValue); -} +__SYCL_GENERATE_CONVERT_IMPL(FToS, short, rtz, Rtz) -template -detail::enable_if_t::value && - std::is_same::value && - roundingMode == rounding_mode::rtp, - R> -convertImpl(T Value) { - using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; - OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); - return __spirv_ConvertFToS_Rlong_rtp(OpValue); -} +__SYCL_GENERATE_CONVERT_IMPL(FToS, long, rtz, Rtz) + +__SYCL_GENERATE_CONVERT_IMPL(FToS, int, rtp, Rtp) + +__SYCL_GENERATE_CONVERT_IMPL(FToS, char, rtp, Rtp) + +__SYCL_GENERATE_CONVERT_IMPL(FToS, short, rtp, Rtp) + +__SYCL_GENERATE_CONVERT_IMPL(FToS, long, rtp, Rtp) + +__SYCL_GENERATE_CONVERT_IMPL(FToS, int, rtn, Rtn) + +__SYCL_GENERATE_CONVERT_IMPL(FToS, char, rtn, Rtn) + +__SYCL_GENERATE_CONVERT_IMPL(FToS, short, rtn, Rtn) + +__SYCL_GENERATE_CONVERT_IMPL(FToS, long, rtn, Rtn) + +#undef __SYCL_GENERATE_CONVERT_IMPL -template -detail::enable_if_t::value && - std::is_same::value && - roundingMode == rounding_mode::rtn, - R> -convertImpl(T Value) { - using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t; - OpenCLT OpValue = cl::sycl::detail::convertDataToType(Value); - return __spirv_ConvertFToS_Rlong_rtn(OpValue); -} #endif } // namespace detail From 070fbf519fab9d7205105eb7fce9b7498d812779 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Thu, 2 Apr 2020 10:29:14 +0300 Subject: [PATCH 15/37] Restore spirv_ops.hpp Signed-off-by: Aleksander Fadeev --- sycl/include/CL/__spirv/spirv_ops.hpp | 23 +++++++++++++++++++++++ 1 file changed, 23 insertions(+) diff --git a/sycl/include/CL/__spirv/spirv_ops.hpp b/sycl/include/CL/__spirv/spirv_ops.hpp index 3173741b12c53..1d758901d05f7 100644 --- a/sycl/include/CL/__spirv/spirv_ops.hpp +++ b/sycl/include/CL/__spirv/spirv_ops.hpp @@ -27,6 +27,29 @@ "SPIR-V built-ins are not available. Please set -fdeclare-spirv-builtins flag." #endif +template +extern SYCL_EXTERNAL RetT __spirv_ImageQueryFormat(ImageT); + +template +extern SYCL_EXTERNAL RetT __spirv_ImageQueryOrder(ImageT); + +template +extern SYCL_EXTERNAL RetT __spirv_ImageQuerySize(ImageT); + +template +extern SYCL_EXTERNAL void __spirv_ImageWrite(ImageT, CoordT, ValT); + +template +extern SYCL_EXTERNAL RetT __spirv_ImageRead(ImageT, TempArgT); + +template +extern SYCL_EXTERNAL SampledType __spirv_SampledImage(ImageT, __ocl_sampler_t); + +template +extern SYCL_EXTERNAL TempRetT __spirv_ImageSampleExplicitLod(SampledType, + TempArgT, int, + float); + #define OpGroupAsyncCopyGlobalToLocal __spirv_GroupAsyncCopy #define OpGroupAsyncCopyLocalToGlobal __spirv_GroupAsyncCopy From 9c1384784957ec8feac6b8fee4add2dd0a0dbf1a Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Thu, 2 Apr 2020 13:12:53 +0300 Subject: [PATCH 16/37] Macros improving Signed-off-by: Aleksander Fadeev --- sycl/include/CL/sycl/types.hpp | 49 ++++++++++++---------------------- 1 file changed, 17 insertions(+), 32 deletions(-) diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp index 98295c9ab09b6..898e386204500 100644 --- a/sycl/include/CL/sycl/types.hpp +++ b/sycl/include/CL/sycl/types.hpp @@ -293,38 +293,23 @@ convertImpl(T Value) { \ return __spirv_Convert##SPIRVOp##_R##DestType##_##RoundingMode(OpValue); \ } \ -__SYCL_GENERATE_CONVERT_IMPL(FToS, int, rte, RteOrAutomatic) - -__SYCL_GENERATE_CONVERT_IMPL(FToU, char, rte, RteOrAutomatic) - -__SYCL_GENERATE_CONVERT_IMPL(FToU, short, rte, RteOrAutomatic) - -__SYCL_GENERATE_CONVERT_IMPL(FToU, long, rte, RteOrAutomatic) - -__SYCL_GENERATE_CONVERT_IMPL(FToS, int, rtz, Rtz) - -__SYCL_GENERATE_CONVERT_IMPL(FToS, char, rtz, Rtz) - -__SYCL_GENERATE_CONVERT_IMPL(FToS, short, rtz, Rtz) - -__SYCL_GENERATE_CONVERT_IMPL(FToS, long, rtz, Rtz) - -__SYCL_GENERATE_CONVERT_IMPL(FToS, int, rtp, Rtp) - -__SYCL_GENERATE_CONVERT_IMPL(FToS, char, rtp, Rtp) - -__SYCL_GENERATE_CONVERT_IMPL(FToS, short, rtp, Rtp) - -__SYCL_GENERATE_CONVERT_IMPL(FToS, long, rtp, Rtp) - -__SYCL_GENERATE_CONVERT_IMPL(FToS, int, rtn, Rtn) - -__SYCL_GENERATE_CONVERT_IMPL(FToS, char, rtn, Rtn) - -__SYCL_GENERATE_CONVERT_IMPL(FToS, short, rtn, Rtn) - -__SYCL_GENERATE_CONVERT_IMPL(FToS, long, rtn, Rtn) - +#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 From 543fe09589890fa0d5312afabefb8ed0ddc7b059 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Thu, 2 Apr 2020 15:50:07 +0300 Subject: [PATCH 17/37] Unsign to sign convert Signed-off-by: Aleksander Fadeev --- sycl/include/CL/sycl/types.hpp | 44 +++++++++++++++++++++++++++++++++- 1 file changed, 43 insertions(+), 1 deletion(-) diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp index 898e386204500..6d7edf6d14c77 100644 --- a/sycl/include/CL/sycl/types.hpp +++ b/sycl/include/CL/sycl/types.hpp @@ -199,6 +199,12 @@ using is_int_to_int = std::integral_constant::value && std::is_integral::value>; +template +using is_int_to_from_uint = + std::integral_constant::value && + std::is_signed::value || + std::is_signed::value && std::is_unsigned::value>; + template using is_int_to_float = std::integral_constant::value && @@ -219,6 +225,7 @@ 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 @@ -232,7 +239,6 @@ convertImpl(T Value) { } // float to int -#ifndef __SYCL_DEVICE_ONLY__ template detail::enable_if_t::value, R> convertImpl(T Value) { switch (roundingMode) { @@ -267,6 +273,16 @@ detail::enable_if_t::value, R> convertImpl(T Value) { } #else +template +detail::enable_if_t::value && + (is_int_to_int::value && + !(is_int_to_from_uint::value) || + is_int_to_float::value || + is_float_to_float::value), + R> +convertImpl(T Value) { + return static_cast(Value); +} template using RteOrAutomatic = detail::bool_constant; @@ -280,6 +296,32 @@ using Rtp = detail::bool_constant; template using Rtn = detail::bool_constant; +#define __SYCL_INT_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_INT_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(RoundingMode, RoundingModeCondition) \ + __SYCL_INT_GENERATE_CONVERT_IMPL(UToS, int, RoundingMode, RoundingModeCondition) \ + __SYCL_INT_GENERATE_CONVERT_IMPL(UToS, char, RoundingMode, RoundingModeCondition) \ + __SYCL_INT_GENERATE_CONVERT_IMPL(UToS, short, RoundingMode, RoundingModeCondition) \ + __SYCL_INT_GENERATE_CONVERT_IMPL(UToS, long, RoundingMode, RoundingModeCondition) \ + __SYCL_INT_GENERATE_CONVERT_IMPL(SToU, uint, RoundingMode, RoundingModeCondition) \ + __SYCL_INT_GENERATE_CONVERT_IMPL(SToU, uchar, RoundingMode, RoundingModeCondition) \ + __SYCL_INT_GENERATE_CONVERT_IMPL(SToU, ushort, RoundingMode, RoundingModeCondition) \ + __SYCL_INT_GENERATE_CONVERT_IMPL(SToU, ulong, RoundingMode, RoundingModeCondition) \ + +__SYCL_INT_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rte, RteOrAutomatic) +__SYCL_INT_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtz, Rtz) +__SYCL_INT_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtp, Rtp) +__SYCL_INT_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtn, Rtn) #define __SYCL_GENERATE_CONVERT_IMPL(SPIRVOp, DestType, RoundingMode, RoundingModeCondition) \ template \ From ec8806a6aeea0bbff60fa1f6e3468d3840e5ecb7 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Thu, 2 Apr 2020 20:51:56 +0300 Subject: [PATCH 18/37] Reverting Signed-off-by: Aleksander Fadeev --- clang/lib/Sema/SPIRVBuiltins.td | 26 ++++++++++----------- sycl/include/CL/sycl/types.hpp | 33 --------------------------- sycl/test/basic_tests/vec_convert.cpp | 2 +- 3 files changed, 14 insertions(+), 47 deletions(-) diff --git a/clang/lib/Sema/SPIRVBuiltins.td b/clang/lib/Sema/SPIRVBuiltins.td index e592413705479..6199787bc6f38 100644 --- a/clang/lib/Sema/SPIRVBuiltins.td +++ b/clang/lib/Sema/SPIRVBuiltins.td @@ -711,36 +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" # IType.Name # rnd, [IType, FType], Attr.Const>; + 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, VectorType], + def : SPVBuiltin<"ConvertFToU_R" # UIType.Name # v, + [VectorType, VectorType], Attr.Const>; def : SPVBuiltin<"ConvertUToF_R" # FType.Name # v, - [VectorType, VectorType], + [VectorType, VectorType], Attr.Const>; } } } -foreach IType = [TrueChar, Short, Int, Long] in { +foreach SIType = [TrueChar, Short, Int, Long] in { foreach FType = [Float, Double, Half] in { - 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" # IType.Name # rnd, [IType, FType], Attr.Const>; + 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, VectorType], + def : SPVBuiltin<"ConvertFToS_R" # SIType.Name # v, + [VectorType, VectorType], Attr.Const>; def : SPVBuiltin<"ConvertSToF_R" # FType.Name # v, - [VectorType, VectorType], + [VectorType, VectorType], Attr.Const>; } } diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp index 6d7edf6d14c77..ef605714f3ceb 100644 --- a/sycl/include/CL/sycl/types.hpp +++ b/sycl/include/CL/sycl/types.hpp @@ -199,12 +199,6 @@ using is_int_to_int = std::integral_constant::value && std::is_integral::value>; -template -using is_int_to_from_uint = - std::integral_constant::value && - std::is_signed::value || - std::is_signed::value && std::is_unsigned::value>; - template using is_int_to_float = std::integral_constant::value && @@ -296,33 +290,6 @@ using Rtp = detail::bool_constant; template using Rtn = detail::bool_constant; -#define __SYCL_INT_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_INT_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(RoundingMode, RoundingModeCondition) \ - __SYCL_INT_GENERATE_CONVERT_IMPL(UToS, int, RoundingMode, RoundingModeCondition) \ - __SYCL_INT_GENERATE_CONVERT_IMPL(UToS, char, RoundingMode, RoundingModeCondition) \ - __SYCL_INT_GENERATE_CONVERT_IMPL(UToS, short, RoundingMode, RoundingModeCondition) \ - __SYCL_INT_GENERATE_CONVERT_IMPL(UToS, long, RoundingMode, RoundingModeCondition) \ - __SYCL_INT_GENERATE_CONVERT_IMPL(SToU, uint, RoundingMode, RoundingModeCondition) \ - __SYCL_INT_GENERATE_CONVERT_IMPL(SToU, uchar, RoundingMode, RoundingModeCondition) \ - __SYCL_INT_GENERATE_CONVERT_IMPL(SToU, ushort, RoundingMode, RoundingModeCondition) \ - __SYCL_INT_GENERATE_CONVERT_IMPL(SToU, ulong, RoundingMode, RoundingModeCondition) \ - -__SYCL_INT_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rte, RteOrAutomatic) -__SYCL_INT_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtz, Rtz) -__SYCL_INT_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtp, Rtp) -__SYCL_INT_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtn, Rtn) - #define __SYCL_GENERATE_CONVERT_IMPL(SPIRVOp, DestType, RoundingMode, RoundingModeCondition) \ template \ detail::enable_if_t::value && \ diff --git a/sycl/test/basic_tests/vec_convert.cpp b/sycl/test/basic_tests/vec_convert.cpp index 21e6249b1adc5..6641b80962f4b 100644 --- a/sycl/test/basic_tests/vec_convert.cpp +++ b/sycl/test/basic_tests/vec_convert.cpp @@ -149,6 +149,6 @@ int main() { 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; } From 870059e553f3956639198dfd1518f341e4eec3ff Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Fri, 3 Apr 2020 10:12:04 +0300 Subject: [PATCH 19/37] Patch Signed-off-by: Aleksander Fadeev --- sycl/include/CL/sycl/types.hpp | 14 ++------------ 1 file changed, 2 insertions(+), 12 deletions(-) diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp index ef605714f3ceb..98da745aae9ea 100644 --- a/sycl/include/CL/sycl/types.hpp +++ b/sycl/include/CL/sycl/types.hpp @@ -219,7 +219,7 @@ 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 @@ -232,6 +232,7 @@ convertImpl(T Value) { return static_cast(Value); } +#ifndef __SYCL_DEVICE_ONLY__ // float to int template detail::enable_if_t::value, R> convertImpl(T Value) { @@ -267,17 +268,6 @@ detail::enable_if_t::value, R> convertImpl(T Value) { } #else -template -detail::enable_if_t::value && - (is_int_to_int::value && - !(is_int_to_from_uint::value) || - is_int_to_float::value || - is_float_to_float::value), - R> -convertImpl(T Value) { - return static_cast(Value); -} - template using RteOrAutomatic = detail::bool_constant; From ea768af708a3a299f9d96657d3673cfa81a38553 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Fri, 3 Apr 2020 10:17:45 +0300 Subject: [PATCH 20/37] Formatting Signed-off-by: Aleksander Fadeev --- sycl/include/CL/sycl/types.hpp | 179 +++++++++++++------------- sycl/test/basic_tests/vec_convert.cpp | 14 +- 2 files changed, 100 insertions(+), 93 deletions(-) diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp index 98da745aae9ea..2acf820befa79 100644 --- a/sycl/include/CL/sycl/types.hpp +++ b/sycl/include/CL/sycl/types.hpp @@ -44,8 +44,8 @@ #error "SYCL device compiler is built without ext_vector_type support" #endif // __HAS_EXT_VECTOR_TYPE__ -#include #include +#include #include #include #include @@ -219,7 +219,6 @@ detail::enable_if_t::value, R> convertImpl(T Value) { return Value; } - // Note for float to half conversions, static_cast calls the conversion operator // implemented for host that takes care of the precision requirements. template @@ -269,7 +268,8 @@ detail::enable_if_t::value, R> convertImpl(T Value) { #else template -using RteOrAutomatic = detail::bool_constant; +using RteOrAutomatic = detail::bool_constant; template using Rtz = detail::bool_constant; @@ -280,34 +280,41 @@ using Rtp = detail::bool_constant; template using Rtn = detail::bool_constant; -#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) \ +#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(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 @@ -326,7 +333,7 @@ __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtn, Rtn) // For information on calling conventions for x64 processors, see // Calling Convention // (https://docs.microsoft.com/en-us/cpp/build/x64-calling-convention). -#pragma message ("Alignment of class vec is not in accordance with SYCL \ +#pragma message("Alignment of class vec is not in accordance with SYCL \ specification requirements, a limitation of the MSVC compiler(Error C2719).\ Applied default alignment.") #define SYCL_ALIGNAS(x) @@ -346,14 +353,15 @@ template class vec { // SizeChecker is needed for vec(const argTN &... args) ctor to validate args. template - struct SizeChecker: detail::conditional_t {}; + struct SizeChecker : detail::conditional_t { + }; template struct SizeChecker : detail::conditional_t, - std::false_type> {}; + SizeChecker, + std::false_type> {}; #define ALLOW_VECTOR_SIZES(num_elements) \ template \ @@ -473,13 +481,13 @@ template class vec { *this = Rhs.template as(); return *this; } - + #ifdef __SYCL_USE_EXT_VECTOR_TYPE__ template using EnableIfNotHostHalf = typename std::enable_if< !std::is_same::value || !std::is_same::value, + cl::sycl::detail::host_half_impl::half>::value, T>::type; template using EnableIfHostHalf = typename std::enable_if< @@ -547,9 +555,10 @@ template class vec { // Helper type to make specific constructors available only for specific // number of elements. template - using EnableIfMultipleElems = typename std::enable_if< - std::is_convertible::value && NumElements == IdxNum, - DataT>::type; + using EnableIfMultipleElems = + typename std::enable_if::value && + NumElements == IdxNum, + DataT>::type; template vec(const EnableIfMultipleElems<2, Ty> Arg0, const EnableIfNotHostHalf Arg1) @@ -673,10 +682,7 @@ template class vec { // function would activate a bug in MSVC that is fixed only in v19.20. // Until then MSVC does not recognize such constexpr functions as const and // thus does not let using them in template parameters inside swizzle.def. - template - struct Indexer { - static constexpr int value = Index; - }; + template struct Indexer { static constexpr int value = Index; }; public: #ifdef __SYCL_ACCESS_RETURN @@ -754,8 +760,8 @@ template class vec { return *this; \ } \ template \ - typename std::enable_if::type \ - operator OPASSIGN(const DataT &Rhs) { \ + typename std::enable_if::type operator OPASSIGN( \ + const DataT &Rhs) { \ *this = *this BINOP vec(Rhs); \ return *this; \ } @@ -891,8 +897,7 @@ template class vec { // Use __SYCL_DEVICE_ONLY__ macro because cast to OpenCL vector type is defined // by SYCL device compiler only. #ifdef __SYCL_DEVICE_ONLY__ - return vec{ - (typename vec::DataType)~m_Data}; + return vec{(typename vec::DataType) ~m_Data}; #else vec Ret; for (size_t I = 0; I < NumElements; ++I) { @@ -907,7 +912,7 @@ template class vec { // by SYCL device compiler only. #ifdef __SYCL_DEVICE_ONLY__ return vec{ - (typename vec::DataType)!m_Data}; + (typename vec::DataType) !m_Data}; #else vec Ret; for (size_t I = 0; I < NumElements; ++I) { @@ -1109,7 +1114,7 @@ template class vec { // alignment of 128 and MSVC compiler cann't align a parameter with requested // alignment of 128. SYCL_ALIGNAS((detail::vector_alignment::value)) - DataType m_Data; + DataType m_Data; // friends template class T4, @@ -1122,7 +1127,7 @@ template class vec { // all compilers supporting deduction guides also support fold expressions template ::value && ...)>> -vec(T, U...)->vec; +vec(T, U...) -> vec; #endif namespace detail { @@ -1271,8 +1276,7 @@ class SwizzleOp { #undef __SYCL_UOP template - typename std::enable_if::value, vec_t>::type - operator~() { + typename std::enable_if::value, vec_t>::type operator~() { vec_t Tmp = *this; return ~Tmp; } @@ -1584,8 +1588,7 @@ class SwizzleOp { // function would activate a bug in MSVC that is fixed only in v19.20. // Until then MSVC does not recognize such constexpr functions as const and // thus does not let using them in template parameters inside swizzle.def. - template - struct Indexer { + template struct Indexer { static constexpr int IDXs[sizeof...(Indexes)] = {Indexes...}; static constexpr int value = IDXs[Index >= getNumElements() ? 0 : Index]; }; @@ -1714,9 +1717,10 @@ class SwizzleOp { (std::is_fundamental::value || \ std::is_same::type, half>::value), \ vec>::type \ - operator BINOP(const T &Lhs, \ - const detail::SwizzleOp &Rhs) {\ + operator BINOP( \ + const T &Lhs, \ + const detail::SwizzleOp &Rhs) { \ vec Tmp = Rhs; \ return Lhs BINOP Tmp; \ } \ @@ -1801,33 +1805,32 @@ __SYCL_RELLOGOP(||) } // namespace sycl } // __SYCL_INLINE_NAMESPACE(cl) - #ifdef __SYCL_USE_EXT_VECTOR_TYPE__ #define DECLARE_TYPE_VIA_CL_T(type) \ -using __##type##_t = cl::sycl::cl_##type; \ -using __##type##2_vec_t = cl::sycl::cl_##type \ - __attribute__((ext_vector_type(2))); \ -using __##type##3_vec_t = cl::sycl::cl_##type \ - __attribute__((ext_vector_type(3))); \ -using __##type##4_vec_t = cl::sycl::cl_##type \ - __attribute__((ext_vector_type(4))); \ -using __##type##8_vec_t = cl::sycl::cl_##type \ - __attribute__((ext_vector_type(8))); \ -using __##type##16_vec_t = cl::sycl::cl_##type \ - __attribute__((ext_vector_type(16))); + using __##type##_t = cl::sycl::cl_##type; \ + using __##type##2_vec_t = \ + cl::sycl::cl_##type __attribute__((ext_vector_type(2))); \ + using __##type##3_vec_t = \ + cl::sycl::cl_##type __attribute__((ext_vector_type(3))); \ + using __##type##4_vec_t = \ + cl::sycl::cl_##type __attribute__((ext_vector_type(4))); \ + using __##type##8_vec_t = \ + cl::sycl::cl_##type __attribute__((ext_vector_type(8))); \ + using __##type##16_vec_t = \ + cl::sycl::cl_##type __attribute__((ext_vector_type(16))); #define DECLARE_TYPE_T(type) \ -using __##type##_t = cl::sycl::type; \ -using __##type##2_vec_t = cl::sycl::type \ - __attribute__((ext_vector_type(2))); \ -using __##type##3_vec_t = cl::sycl::type \ - __attribute__((ext_vector_type(3))); \ -using __##type##4_vec_t = cl::sycl::type \ - __attribute__((ext_vector_type(4))); \ -using __##type##8_vec_t = cl::sycl::type \ - __attribute__((ext_vector_type(8))); \ -using __##type##16_vec_t = cl::sycl::type \ - __attribute__((ext_vector_type(16))); + using __##type##_t = cl::sycl::type; \ + using __##type##2_vec_t = \ + cl::sycl::type __attribute__((ext_vector_type(2))); \ + using __##type##3_vec_t = \ + cl::sycl::type __attribute__((ext_vector_type(3))); \ + using __##type##4_vec_t = \ + cl::sycl::type __attribute__((ext_vector_type(4))); \ + using __##type##8_vec_t = \ + cl::sycl::type __attribute__((ext_vector_type(8))); \ + using __##type##16_vec_t = \ + cl::sycl::type __attribute__((ext_vector_type(16))); DECLARE_TYPE_VIA_CL_T(char); DECLARE_TYPE_T(schar); @@ -1870,10 +1873,10 @@ namespace detail { // sizeof(IN). expected to handle scalar types in IN. template using select_apply_cl_t = - conditional_t>>; -} // detail + conditional_t>>; +} // namespace detail #define DECLARE_CONVERTER(base, num) \ template <> class BaseCLTypeConverter { \ @@ -1884,24 +1887,25 @@ using select_apply_cl_t = #define DECLARE_SIGNED_INTEGRAL_CONVERTER(base, num) \ template <> class BaseCLTypeConverter { \ public: \ - using DataType = detail::select_apply_cl_t; \ + using DataType = detail::select_apply_cl_t< \ + base, GET_CL_TYPE(char, num), GET_CL_TYPE(short, num), \ + GET_CL_TYPE(int, num), GET_CL_TYPE(long, num)>; \ }; #define DECLARE_UNSIGNED_INTEGRAL_CONVERTER(base, num) \ template <> class BaseCLTypeConverter { \ public: \ - using DataType = detail::select_apply_cl_t; \ + using DataType = detail::select_apply_cl_t< \ + base, GET_CL_TYPE(uchar, num), GET_CL_TYPE(ushort, num), \ + GET_CL_TYPE(uint, num), GET_CL_TYPE(ulong, num)>; \ }; #define DECLARE_FLOAT_CONVERTER(base, num) \ template <> class BaseCLTypeConverter { \ public: \ - using DataType = detail::select_apply_cl_t; \ + using DataType = detail::select_apply_cl_t< \ + base, std::false_type, GET_CL_HALF_TYPE(half, num), \ + GET_CL_TYPE(float, num), GET_CL_TYPE(double, num)>; \ }; #define DECLARE_LONGLONG_CONVERTER(base, num) \ @@ -1976,7 +1980,6 @@ using select_apply_cl_t = DECLARE_SCALAR_CONVERTER(base) \ } // namespace detail - #define DECLARE_HALF_VECTOR_CONVERTERS(base) \ namespace detail { \ DECLARE_HALF_CONVERTER(base, 2) \ diff --git a/sycl/test/basic_tests/vec_convert.cpp b/sycl/test/basic_tests/vec_convert.cpp index 6641b80962f4b..409035315d1cc 100644 --- a/sycl/test/basic_tests/vec_convert.cpp +++ b/sycl/test/basic_tests/vec_convert.cpp @@ -20,11 +20,14 @@ using namespace cl::sycl; -template class kernel_name; +template +class kernel_name; -template struct helper; +template +struct helper; -template <> struct helper<0> { +template <> +struct helper<0> { template static void compare(const vec &x, const vec &y) { @@ -34,7 +37,8 @@ template <> struct helper<0> { } }; -template struct helper { +template +struct helper { template static void compare(const vec &x, const vec &y) { @@ -149,6 +153,6 @@ int main() { 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; } From 238ee1a3e3926871f52eed1eeea78655575a18f1 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Fri, 3 Apr 2020 10:26:04 +0300 Subject: [PATCH 21/37] Formatting 2 Signed-off-by: Aleksander Fadeev --- sycl/include/CL/sycl/types.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp index 2acf820befa79..af78583fcf844 100644 --- a/sycl/include/CL/sycl/types.hpp +++ b/sycl/include/CL/sycl/types.hpp @@ -1127,7 +1127,7 @@ template class vec { // all compilers supporting deduction guides also support fold expressions template ::value && ...)>> -vec(T, U...) -> vec; +vec(T, U...)->vec; #endif namespace detail { From c03698bf53711f234dcae8c42f994de8d78ec4f1 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Fri, 3 Apr 2020 11:25:53 +0300 Subject: [PATCH 22/37] Revert Signed-off-by: Aleksander Fadeev --- clang/lib/Sema/SPIRVBuiltins.td | 25 ++-- sycl/include/CL/sycl/types.hpp | 177 +++++++++++++------------- sycl/test/basic_tests/vec_convert.cpp | 14 +- 3 files changed, 104 insertions(+), 112 deletions(-) diff --git a/clang/lib/Sema/SPIRVBuiltins.td b/clang/lib/Sema/SPIRVBuiltins.td index 6199787bc6f38..ea07ae585e029 100644 --- a/clang/lib/Sema/SPIRVBuiltins.td +++ b/clang/lib/Sema/SPIRVBuiltins.td @@ -711,36 +711,35 @@ foreach name = ["GenericPtrMemSemantics"] in { // 3.32.11. Conversion Instructions -foreach UIType = [UChar, UShort, UInt, ULong] in { +foreach IType = [UChar, UShort, UInt, ULong] in { foreach FType = [Float, Double, Half] in { - def : SPVBuiltin<"ConvertFToU_R" # UIType.Name, [UIType, FType], Attr.Const>; - def : SPVBuiltin<"ConvertUToF_R" # FType.Name, [FType, UIType], Attr.Const>; + def : SPVBuiltin<"ConvertUToF_R" # FType.Name, [FType, IType], Attr.Const>; foreach rnd = ["", "_rte", "_rtz", "_rtp", "_rtn"] in { - def : SPVBuiltin<"ConvertFToU_R" # UIType.Name # rnd, [UIType, FType], Attr.Const>; + def : SPVBuiltin<"ConvertFToU_R" # IType.Name # rnd, [IType, FType], Attr.Const>; } foreach v = [2, 3, 4, 8, 16] in { - def : SPVBuiltin<"ConvertFToU_R" # UIType.Name # v, - [VectorType, VectorType], + def : SPVBuiltin<"ConvertFToU_R" # IType.Name # v, + [VectorType, VectorType], Attr.Const>; def : SPVBuiltin<"ConvertUToF_R" # FType.Name # v, - [VectorType, VectorType], + [VectorType, VectorType], Attr.Const>; } } } -foreach SIType = [TrueChar, Short, Int, Long] in { +foreach IType = [TrueChar, Short, Int, Long] in { foreach FType = [Float, Double, Half] in { - def : SPVBuiltin<"ConvertSToF_R" # FType.Name, [FType, SIType], Attr.Const>; + def : SPVBuiltin<"ConvertSToF_R" # FType.Name, [FType, IType], Attr.Const>; foreach rnd = ["", "_rte", "_rtz", "_rtp", "_rtn"] in { - def : SPVBuiltin<"ConvertFToS_R" # SIType.Name # rnd, [SIType, FType], Attr.Const>; + def : SPVBuiltin<"ConvertFToS_R" # IType.Name # rnd, [IType, FType], Attr.Const>; } foreach v = [2, 3, 4, 8, 16] in { - def : SPVBuiltin<"ConvertFToS_R" # SIType.Name # v, - [VectorType, VectorType], + def : SPVBuiltin<"ConvertFToS_R" # IType.Name # v, + [VectorType, VectorType], Attr.Const>; def : SPVBuiltin<"ConvertSToF_R" # FType.Name # v, - [VectorType, VectorType], + [VectorType, VectorType], Attr.Const>; } } diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp index af78583fcf844..98da745aae9ea 100644 --- a/sycl/include/CL/sycl/types.hpp +++ b/sycl/include/CL/sycl/types.hpp @@ -44,8 +44,8 @@ #error "SYCL device compiler is built without ext_vector_type support" #endif // __HAS_EXT_VECTOR_TYPE__ -#include #include +#include #include #include #include @@ -219,6 +219,7 @@ detail::enable_if_t::value, R> convertImpl(T Value) { return Value; } + // Note for float to half conversions, static_cast calls the conversion operator // implemented for host that takes care of the precision requirements. template @@ -268,8 +269,7 @@ detail::enable_if_t::value, R> convertImpl(T Value) { #else template -using RteOrAutomatic = detail::bool_constant; +using RteOrAutomatic = detail::bool_constant; template using Rtz = detail::bool_constant; @@ -280,41 +280,34 @@ using Rtp = detail::bool_constant; template using Rtn = detail::bool_constant; -#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) \ +#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(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 @@ -333,7 +326,7 @@ __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtn, Rtn) // For information on calling conventions for x64 processors, see // Calling Convention // (https://docs.microsoft.com/en-us/cpp/build/x64-calling-convention). -#pragma message("Alignment of class vec is not in accordance with SYCL \ +#pragma message ("Alignment of class vec is not in accordance with SYCL \ specification requirements, a limitation of the MSVC compiler(Error C2719).\ Applied default alignment.") #define SYCL_ALIGNAS(x) @@ -353,15 +346,14 @@ template class vec { // SizeChecker is needed for vec(const argTN &... args) ctor to validate args. template - struct SizeChecker : detail::conditional_t { - }; + struct SizeChecker: detail::conditional_t {}; template struct SizeChecker : detail::conditional_t, - std::false_type> {}; + SizeChecker, + std::false_type> {}; #define ALLOW_VECTOR_SIZES(num_elements) \ template \ @@ -481,13 +473,13 @@ template class vec { *this = Rhs.template as(); return *this; } - + #ifdef __SYCL_USE_EXT_VECTOR_TYPE__ template using EnableIfNotHostHalf = typename std::enable_if< !std::is_same::value || !std::is_same::value, + cl::sycl::detail::host_half_impl::half>::value, T>::type; template using EnableIfHostHalf = typename std::enable_if< @@ -555,10 +547,9 @@ template class vec { // Helper type to make specific constructors available only for specific // number of elements. template - using EnableIfMultipleElems = - typename std::enable_if::value && - NumElements == IdxNum, - DataT>::type; + using EnableIfMultipleElems = typename std::enable_if< + std::is_convertible::value && NumElements == IdxNum, + DataT>::type; template vec(const EnableIfMultipleElems<2, Ty> Arg0, const EnableIfNotHostHalf Arg1) @@ -682,7 +673,10 @@ template class vec { // function would activate a bug in MSVC that is fixed only in v19.20. // Until then MSVC does not recognize such constexpr functions as const and // thus does not let using them in template parameters inside swizzle.def. - template struct Indexer { static constexpr int value = Index; }; + template + struct Indexer { + static constexpr int value = Index; + }; public: #ifdef __SYCL_ACCESS_RETURN @@ -760,8 +754,8 @@ template class vec { return *this; \ } \ template \ - typename std::enable_if::type operator OPASSIGN( \ - const DataT &Rhs) { \ + typename std::enable_if::type \ + operator OPASSIGN(const DataT &Rhs) { \ *this = *this BINOP vec(Rhs); \ return *this; \ } @@ -897,7 +891,8 @@ template class vec { // Use __SYCL_DEVICE_ONLY__ macro because cast to OpenCL vector type is defined // by SYCL device compiler only. #ifdef __SYCL_DEVICE_ONLY__ - return vec{(typename vec::DataType) ~m_Data}; + return vec{ + (typename vec::DataType)~m_Data}; #else vec Ret; for (size_t I = 0; I < NumElements; ++I) { @@ -912,7 +907,7 @@ template class vec { // by SYCL device compiler only. #ifdef __SYCL_DEVICE_ONLY__ return vec{ - (typename vec::DataType) !m_Data}; + (typename vec::DataType)!m_Data}; #else vec Ret; for (size_t I = 0; I < NumElements; ++I) { @@ -1114,7 +1109,7 @@ template class vec { // alignment of 128 and MSVC compiler cann't align a parameter with requested // alignment of 128. SYCL_ALIGNAS((detail::vector_alignment::value)) - DataType m_Data; + DataType m_Data; // friends template class T4, @@ -1276,7 +1271,8 @@ class SwizzleOp { #undef __SYCL_UOP template - typename std::enable_if::value, vec_t>::type operator~() { + typename std::enable_if::value, vec_t>::type + operator~() { vec_t Tmp = *this; return ~Tmp; } @@ -1588,7 +1584,8 @@ class SwizzleOp { // function would activate a bug in MSVC that is fixed only in v19.20. // Until then MSVC does not recognize such constexpr functions as const and // thus does not let using them in template parameters inside swizzle.def. - template struct Indexer { + template + struct Indexer { static constexpr int IDXs[sizeof...(Indexes)] = {Indexes...}; static constexpr int value = IDXs[Index >= getNumElements() ? 0 : Index]; }; @@ -1717,10 +1714,9 @@ class SwizzleOp { (std::is_fundamental::value || \ std::is_same::type, half>::value), \ vec>::type \ - operator BINOP( \ - const T &Lhs, \ - const detail::SwizzleOp &Rhs) { \ + operator BINOP(const T &Lhs, \ + const detail::SwizzleOp &Rhs) {\ vec Tmp = Rhs; \ return Lhs BINOP Tmp; \ } \ @@ -1805,32 +1801,33 @@ __SYCL_RELLOGOP(||) } // namespace sycl } // __SYCL_INLINE_NAMESPACE(cl) + #ifdef __SYCL_USE_EXT_VECTOR_TYPE__ #define DECLARE_TYPE_VIA_CL_T(type) \ - using __##type##_t = cl::sycl::cl_##type; \ - using __##type##2_vec_t = \ - cl::sycl::cl_##type __attribute__((ext_vector_type(2))); \ - using __##type##3_vec_t = \ - cl::sycl::cl_##type __attribute__((ext_vector_type(3))); \ - using __##type##4_vec_t = \ - cl::sycl::cl_##type __attribute__((ext_vector_type(4))); \ - using __##type##8_vec_t = \ - cl::sycl::cl_##type __attribute__((ext_vector_type(8))); \ - using __##type##16_vec_t = \ - cl::sycl::cl_##type __attribute__((ext_vector_type(16))); +using __##type##_t = cl::sycl::cl_##type; \ +using __##type##2_vec_t = cl::sycl::cl_##type \ + __attribute__((ext_vector_type(2))); \ +using __##type##3_vec_t = cl::sycl::cl_##type \ + __attribute__((ext_vector_type(3))); \ +using __##type##4_vec_t = cl::sycl::cl_##type \ + __attribute__((ext_vector_type(4))); \ +using __##type##8_vec_t = cl::sycl::cl_##type \ + __attribute__((ext_vector_type(8))); \ +using __##type##16_vec_t = cl::sycl::cl_##type \ + __attribute__((ext_vector_type(16))); #define DECLARE_TYPE_T(type) \ - using __##type##_t = cl::sycl::type; \ - using __##type##2_vec_t = \ - cl::sycl::type __attribute__((ext_vector_type(2))); \ - using __##type##3_vec_t = \ - cl::sycl::type __attribute__((ext_vector_type(3))); \ - using __##type##4_vec_t = \ - cl::sycl::type __attribute__((ext_vector_type(4))); \ - using __##type##8_vec_t = \ - cl::sycl::type __attribute__((ext_vector_type(8))); \ - using __##type##16_vec_t = \ - cl::sycl::type __attribute__((ext_vector_type(16))); +using __##type##_t = cl::sycl::type; \ +using __##type##2_vec_t = cl::sycl::type \ + __attribute__((ext_vector_type(2))); \ +using __##type##3_vec_t = cl::sycl::type \ + __attribute__((ext_vector_type(3))); \ +using __##type##4_vec_t = cl::sycl::type \ + __attribute__((ext_vector_type(4))); \ +using __##type##8_vec_t = cl::sycl::type \ + __attribute__((ext_vector_type(8))); \ +using __##type##16_vec_t = cl::sycl::type \ + __attribute__((ext_vector_type(16))); DECLARE_TYPE_VIA_CL_T(char); DECLARE_TYPE_T(schar); @@ -1873,10 +1870,10 @@ namespace detail { // sizeof(IN). expected to handle scalar types in IN. template using select_apply_cl_t = - conditional_t>>; -} // namespace detail + conditional_t>>; +} // detail #define DECLARE_CONVERTER(base, num) \ template <> class BaseCLTypeConverter { \ @@ -1887,25 +1884,24 @@ using select_apply_cl_t = #define DECLARE_SIGNED_INTEGRAL_CONVERTER(base, num) \ template <> class BaseCLTypeConverter { \ public: \ - using DataType = detail::select_apply_cl_t< \ - base, GET_CL_TYPE(char, num), GET_CL_TYPE(short, num), \ - GET_CL_TYPE(int, num), GET_CL_TYPE(long, num)>; \ + using DataType = detail::select_apply_cl_t; \ }; #define DECLARE_UNSIGNED_INTEGRAL_CONVERTER(base, num) \ template <> class BaseCLTypeConverter { \ public: \ - using DataType = detail::select_apply_cl_t< \ - base, GET_CL_TYPE(uchar, num), GET_CL_TYPE(ushort, num), \ - GET_CL_TYPE(uint, num), GET_CL_TYPE(ulong, num)>; \ + using DataType = detail::select_apply_cl_t; \ }; #define DECLARE_FLOAT_CONVERTER(base, num) \ template <> class BaseCLTypeConverter { \ public: \ - using DataType = detail::select_apply_cl_t< \ - base, std::false_type, GET_CL_HALF_TYPE(half, num), \ - GET_CL_TYPE(float, num), GET_CL_TYPE(double, num)>; \ + using DataType = detail::select_apply_cl_t; \ }; #define DECLARE_LONGLONG_CONVERTER(base, num) \ @@ -1980,6 +1976,7 @@ using select_apply_cl_t = DECLARE_SCALAR_CONVERTER(base) \ } // namespace detail + #define DECLARE_HALF_VECTOR_CONVERTERS(base) \ namespace detail { \ DECLARE_HALF_CONVERTER(base, 2) \ diff --git a/sycl/test/basic_tests/vec_convert.cpp b/sycl/test/basic_tests/vec_convert.cpp index 409035315d1cc..6641b80962f4b 100644 --- a/sycl/test/basic_tests/vec_convert.cpp +++ b/sycl/test/basic_tests/vec_convert.cpp @@ -20,14 +20,11 @@ using namespace cl::sycl; -template -class kernel_name; +template class kernel_name; -template -struct helper; +template struct helper; -template <> -struct helper<0> { +template <> struct helper<0> { template static void compare(const vec &x, const vec &y) { @@ -37,8 +34,7 @@ struct helper<0> { } }; -template -struct helper { +template struct helper { template static void compare(const vec &x, const vec &y) { @@ -153,6 +149,6 @@ int main() { 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; } From e0b05669653660d6aba8117435a59ac0ae7114d4 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Fri, 3 Apr 2020 13:14:37 +0300 Subject: [PATCH 23/37] Formatting 3 Signed-off-by: Aleksander Fadeev --- sycl/include/CL/sycl/types.hpp | 53 +++++++++++++++------------ sycl/test/basic_tests/vec_convert.cpp | 3 +- 2 files changed, 31 insertions(+), 25 deletions(-) diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp index 98da745aae9ea..2d3c519e4d011 100644 --- a/sycl/include/CL/sycl/types.hpp +++ b/sycl/include/CL/sycl/types.hpp @@ -219,7 +219,6 @@ detail::enable_if_t::value, R> convertImpl(T Value) { return Value; } - // Note for float to half conversions, static_cast calls the conversion operator // implemented for host that takes care of the precision requirements. template @@ -269,7 +268,8 @@ detail::enable_if_t::value, R> convertImpl(T Value) { #else template -using RteOrAutomatic = detail::bool_constant; +using RteOrAutomatic = detail::bool_constant; template using Rtz = detail::bool_constant; @@ -280,34 +280,41 @@ using Rtp = detail::bool_constant; template using Rtn = detail::bool_constant; -#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) \ +#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(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 diff --git a/sycl/test/basic_tests/vec_convert.cpp b/sycl/test/basic_tests/vec_convert.cpp index 6641b80962f4b..07d0307bea847 100644 --- a/sycl/test/basic_tests/vec_convert.cpp +++ b/sycl/test/basic_tests/vec_convert.cpp @@ -149,6 +149,5 @@ int main() { 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; -} +} \ No newline at end of file From c5cf20ef557ed844278857eb73786c10dbeb0545 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Wed, 8 Apr 2020 15:52:52 +0300 Subject: [PATCH 24/37] Coments Signed-off-by: Aleksander Fadeev --- sycl/include/CL/sycl/types.hpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp index 2d3c519e4d011..c0d02296267ab 100644 --- a/sycl/include/CL/sycl/types.hpp +++ b/sycl/include/CL/sycl/types.hpp @@ -280,6 +280,7 @@ using Rtp = detail::bool_constant; template using Rtn = detail::bool_constant; +// Convert floating-point type to integer #define __SYCL_GENERATE_CONVERT_IMPL(SPIRVOp, DestType, RoundingMode, \ RoundingModeCondition) \ template \ @@ -318,7 +319,7 @@ __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtn, Rtn) #undef __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE #undef __SYCL_GENERATE_CONVERT_IMPL -#endif +#endif // __SYCL_DEVICE_ONLY__ } // namespace detail From 8267f30a6880efb52ad78b8dbf726cf3a6246eb7 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Wed, 8 Apr 2020 15:55:34 +0300 Subject: [PATCH 25/37] Coments fix Signed-off-by: Aleksander Fadeev --- sycl/include/CL/sycl/types.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp index c0d02296267ab..c89ac1e729ec6 100644 --- a/sycl/include/CL/sycl/types.hpp +++ b/sycl/include/CL/sycl/types.hpp @@ -280,7 +280,7 @@ using Rtp = detail::bool_constant; template using Rtn = detail::bool_constant; -// Convert floating-point type to integer +// Convert floating-point type to integer type #define __SYCL_GENERATE_CONVERT_IMPL(SPIRVOp, DestType, RoundingMode, \ RoundingModeCondition) \ template \ From 33220d7334c3ee515494071a383b69655cf0a763 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Wed, 8 Apr 2020 16:00:44 +0300 Subject: [PATCH 26/37] Formatting 4 Signed-off-by: Aleksander Fadeev --- sycl/include/CL/sycl/types.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp index c89ac1e729ec6..f9fc05e4ff6a1 100644 --- a/sycl/include/CL/sycl/types.hpp +++ b/sycl/include/CL/sycl/types.hpp @@ -280,7 +280,7 @@ using Rtp = detail::bool_constant; template using Rtn = detail::bool_constant; -// Convert floating-point type to integer type +// Convert floating-point type to integer type #define __SYCL_GENERATE_CONVERT_IMPL(SPIRVOp, DestType, RoundingMode, \ RoundingModeCondition) \ template \ From 56835779ec1b038d57711409346162e53c432e4e Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Fri, 10 Apr 2020 14:06:39 +0300 Subject: [PATCH 27/37] Test fix Signed-off-by: Aleksander Fadeev --- sycl/test/basic_tests/vec_convert.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/sycl/test/basic_tests/vec_convert.cpp b/sycl/test/basic_tests/vec_convert.cpp index 07d0307bea847..7684530ba3b9e 100644 --- a/sycl/test/basic_tests/vec_convert.cpp +++ b/sycl/test/basic_tests/vec_convert.cpp @@ -1,3 +1,4 @@ +// UNSUPPORTED: 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 From 5cd3e6e9f875fbec9c005a78b5348633a9055257 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Fri, 10 Apr 2020 14:11:20 +0300 Subject: [PATCH 28/37] Test fix2 Signed-off-by: Aleksander Fadeev --- sycl/test/basic_tests/handler/handler_mem_op.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/sycl/test/basic_tests/handler/handler_mem_op.cpp b/sycl/test/basic_tests/handler/handler_mem_op.cpp index 124c57e62d82c..b95c40c20631c 100644 --- a/sycl/test/basic_tests/handler/handler_mem_op.cpp +++ b/sycl/test/basic_tests/handler/handler_mem_op.cpp @@ -1,3 +1,4 @@ +// UNSUPPORTED: 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 From 843bf4a8f5f2947927dc8fdb0f78bbd3aec07cf0 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Fri, 10 Apr 2020 14:51:54 +0300 Subject: [PATCH 29/37] Cuda unsupp Signed-off-by: Aleksander Fadeev --- sycl/test/basic_tests/buffer/buffer_dev_to_dev.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/sycl/test/basic_tests/buffer/buffer_dev_to_dev.cpp b/sycl/test/basic_tests/buffer/buffer_dev_to_dev.cpp index c84a06bbbace4..05cabafbaefdf 100644 --- a/sycl/test/basic_tests/buffer/buffer_dev_to_dev.cpp +++ b/sycl/test/basic_tests/buffer/buffer_dev_to_dev.cpp @@ -1,3 +1,4 @@ +// UNSUPPORTED: 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 From 2a1fb9483665fe0dc9de11e0d3b926bd03d6f31a Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Mon, 13 Apr 2020 13:39:24 +0300 Subject: [PATCH 30/37] delete unsupported: cuda Signed-off-by: Aleksander Fadeev --- sycl/test/basic_tests/buffer/buffer_dev_to_dev.cpp | 1 - sycl/test/basic_tests/handler/handler_mem_op.cpp | 1 - 2 files changed, 2 deletions(-) diff --git a/sycl/test/basic_tests/buffer/buffer_dev_to_dev.cpp b/sycl/test/basic_tests/buffer/buffer_dev_to_dev.cpp index 05cabafbaefdf..c84a06bbbace4 100644 --- a/sycl/test/basic_tests/buffer/buffer_dev_to_dev.cpp +++ b/sycl/test/basic_tests/buffer/buffer_dev_to_dev.cpp @@ -1,4 +1,3 @@ -// UNSUPPORTED: 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 diff --git a/sycl/test/basic_tests/handler/handler_mem_op.cpp b/sycl/test/basic_tests/handler/handler_mem_op.cpp index b95c40c20631c..124c57e62d82c 100644 --- a/sycl/test/basic_tests/handler/handler_mem_op.cpp +++ b/sycl/test/basic_tests/handler/handler_mem_op.cpp @@ -1,4 +1,3 @@ -// UNSUPPORTED: 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 From 8921bc3f64adc77d09dbcfbd403dcceeec5230cb Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Mon, 13 Apr 2020 14:22:37 +0300 Subject: [PATCH 31/37] vec_convert.cpp XFail: cuda Signed-off-by: Aleksander Fadeev --- sycl/test/basic_tests/vec_convert.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/test/basic_tests/vec_convert.cpp b/sycl/test/basic_tests/vec_convert.cpp index 7684530ba3b9e..7683f9635ec1a 100644 --- a/sycl/test/basic_tests/vec_convert.cpp +++ b/sycl/test/basic_tests/vec_convert.cpp @@ -1,4 +1,4 @@ -// UNSUPPORTED: cuda +// 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 @@ -18,6 +18,7 @@ // 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; From f6d6777e2aaf40184059566cd9bd76c0f794d351 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Tue, 14 Apr 2020 10:25:40 +0300 Subject: [PATCH 32/37] Unsupported on cuda buffer_dev_to_dev.cpp Signed-off-by: Aleksander Fadeev --- sycl/test/basic_tests/buffer/buffer_dev_to_dev.cpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/sycl/test/basic_tests/buffer/buffer_dev_to_dev.cpp b/sycl/test/basic_tests/buffer/buffer_dev_to_dev.cpp index c84a06bbbace4..d00ba5a0db099 100644 --- a/sycl/test/basic_tests/buffer/buffer_dev_to_dev.cpp +++ b/sycl/test/basic_tests/buffer/buffer_dev_to_dev.cpp @@ -4,6 +4,9 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // RUN: %ACC_RUN_PLACEHOLDER %t.out +// Unsupported: cuda +// The test fails sporadically on cuda. + //==---------- buffer_dev_to_dev.cpp - SYCL buffer basic test --------------==// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. From e875fbf334bd10612637d5d04574b369485d2a53 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Tue, 14 Apr 2020 13:11:53 +0300 Subject: [PATCH 33/37] TrueChar to Char Signed-off-by: Aleksander Fadeev --- clang/lib/Sema/SPIRVBuiltins.td | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/Sema/SPIRVBuiltins.td b/clang/lib/Sema/SPIRVBuiltins.td index ea07ae585e029..c798481452c62 100644 --- a/clang/lib/Sema/SPIRVBuiltins.td +++ b/clang/lib/Sema/SPIRVBuiltins.td @@ -728,7 +728,7 @@ foreach IType = [UChar, UShort, UInt, ULong] in { } } -foreach IType = [TrueChar, Short, Int, Long] in { +foreach IType = [Char, Short, Int, Long] in { foreach FType = [Float, Double, Half] in { def : SPVBuiltin<"ConvertSToF_R" # FType.Name, [FType, IType], Attr.Const>; foreach rnd = ["", "_rte", "_rtz", "_rtp", "_rtn"] in { From f1f289720159d2b7947e7955fb15347dd2f01ec0 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Tue, 14 Apr 2020 15:16:55 +0300 Subject: [PATCH 34/37] schar to char Signed-off-by: Aleksander Fadeev --- clang/lib/Sema/SPIRVBuiltins.td | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/Sema/SPIRVBuiltins.td b/clang/lib/Sema/SPIRVBuiltins.td index c798481452c62..8e0eefa10d73b 100644 --- a/clang/lib/Sema/SPIRVBuiltins.td +++ b/clang/lib/Sema/SPIRVBuiltins.td @@ -287,7 +287,7 @@ 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 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>; From c417b51e8d7c408eea99c3e6442b6364f0f3361e Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Wed, 15 Apr 2020 10:53:34 +0300 Subject: [PATCH 35/37] fix TrueChar Signed-off-by: Aleksander Fadeev --- clang/lib/Sema/SPIRVBuiltins.td | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/Sema/SPIRVBuiltins.td b/clang/lib/Sema/SPIRVBuiltins.td index 8e0eefa10d73b..41d5ead350414 100644 --- a/clang/lib/Sema/SPIRVBuiltins.td +++ b/clang/lib/Sema/SPIRVBuiltins.td @@ -286,7 +286,7 @@ 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 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>; From 014c0d4e96e5cfaff18240fca5f5ce558e7acfb2 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Wed, 15 Apr 2020 14:49:48 +0300 Subject: [PATCH 36/37] Comments fix Signed-off-by: Aleksander Fadeev --- sycl/test/basic_tests/buffer/buffer_dev_to_dev.cpp | 2 +- sycl/test/basic_tests/vec_convert.cpp | 3 +-- 2 files changed, 2 insertions(+), 3 deletions(-) diff --git a/sycl/test/basic_tests/buffer/buffer_dev_to_dev.cpp b/sycl/test/basic_tests/buffer/buffer_dev_to_dev.cpp index d00ba5a0db099..91f02653b2e0f 100644 --- a/sycl/test/basic_tests/buffer/buffer_dev_to_dev.cpp +++ b/sycl/test/basic_tests/buffer/buffer_dev_to_dev.cpp @@ -5,7 +5,7 @@ // RUN: %ACC_RUN_PLACEHOLDER %t.out // Unsupported: cuda -// The test fails sporadically on cuda. +// The test fails sporadically on cuda. See https://github.com/intel/llvm/issues/1508 for more details. //==---------- buffer_dev_to_dev.cpp - SYCL buffer basic test --------------==// // diff --git a/sycl/test/basic_tests/vec_convert.cpp b/sycl/test/basic_tests/vec_convert.cpp index 7683f9635ec1a..de186178f1cb0 100644 --- a/sycl/test/basic_tests/vec_convert.cpp +++ b/sycl/test/basic_tests/vec_convert.cpp @@ -16,8 +16,6 @@ #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; @@ -151,5 +149,6 @@ int main() { 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; } \ No newline at end of file From b0eb8e026e1c8913f590021f5eec49692c9c8c52 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Wed, 15 Apr 2020 19:00:22 +0300 Subject: [PATCH 37/37] Fix Signed-off-by: Aleksander Fadeev --- sycl/test/basic_tests/buffer/buffer_dev_to_dev.cpp | 3 --- sycl/test/basic_tests/vec_convert.cpp | 2 +- 2 files changed, 1 insertion(+), 4 deletions(-) diff --git a/sycl/test/basic_tests/buffer/buffer_dev_to_dev.cpp b/sycl/test/basic_tests/buffer/buffer_dev_to_dev.cpp index 91f02653b2e0f..c84a06bbbace4 100644 --- a/sycl/test/basic_tests/buffer/buffer_dev_to_dev.cpp +++ b/sycl/test/basic_tests/buffer/buffer_dev_to_dev.cpp @@ -4,9 +4,6 @@ // 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. - //==---------- buffer_dev_to_dev.cpp - SYCL buffer basic test --------------==// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. diff --git a/sycl/test/basic_tests/vec_convert.cpp b/sycl/test/basic_tests/vec_convert.cpp index de186178f1cb0..1e7235152e1da 100644 --- a/sycl/test/basic_tests/vec_convert.cpp +++ b/sycl/test/basic_tests/vec_convert.cpp @@ -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; -} \ No newline at end of file +}