[SYCL] Fix for memory layout and alignment of the class vec.

alexeyvoronov-intel · bader · commit 32f0cd57c898 · 2019-09-01T10:06:09.000+03:00
Now vectors with a length of 3 will store 4 elements of dataT.

Fixed boolean lit test on debug mode.
Fixed the alignment of vec&lt;half, N&gt; variants.

Now the vec class conform the requirements of item
"4.10.2.6 Memory layout and alignment" of the SYCL specification:
'''
The elements of an instance of the SYCL vec class template are stored
in memory sequentially and contiguously and are aligned to the size
of the element type in bytes multiplied by the number of elements:
  sizeof(dataT) * numElements
The exception to this is when the number of element is three in which
case the SYCL vec is aligned to the size of the element type in bytes
multiplied by four:
  sizeof(dataT) * 4.
'''

The lit test type.cpp was splited into two: aliases.cpp and types.cpp.
This is done to make the tests smaller and make the tests consistent
with what they're testing.

Signed-off-by: Alexey Voronov &lt;alexey.voronov@intel.com&gt;
diff --git a/sycl/include/CL/sycl/detail/boolean.hpp b/sycl/include/CL/sycl/detail/boolean.hpp
@@ -43,18 +43,18 @@ template <> struct Assigner<0> {
   }
 };
 
-template <int N> struct alignas(N == 3 ? 4 : N) Boolean {
+template <int N> struct alignas(VectorAlignment<bool, N>::value) Boolean {
   static_assert(((N == 2) || (N == 3) || (N == 4) || (N == 8) || (N == 16)),
                 "Invalid size");
 
   using element_type = bool;
 
 #ifdef __SYCL_DEVICE_ONLY__
   using DataType =
-      element_type __attribute__((ext_vector_type(N)));
+      element_type __attribute__((ext_vector_type(VectorLength<N>::value)));
   using vector_t = DataType;
 #else
-  using DataType = element_type[N];
+  using DataType = element_type[VectorLength<N>::value];
 #endif
 
   Boolean() : value{false} {}
diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp
@@ -306,15 +306,24 @@ detail::enable_if_t<is_float_to_int<T, R>::value, R> convertImpl(T Value) {
 #endif
 }
 
+// 4.10.2.6 Memory layout and alignment
+template <int N> struct VectorLength { constexpr static int value = N; };
+
+template <> struct VectorLength<3> { constexpr static int value = 4; };
+
+template <typename T, int N> struct VectorAlignment {
+  constexpr static int value = sizeof(T) * VectorLength<N>::value;
+};
+
 } // namespace detail
 
 template <typename Type, int NumElements> class vec {
   using DataT = Type;
 
   // This represent type of underlying value. There should be only one field
   // in the class, so vec<float, 16> should be equal to float16 in memory.
-  using DataType =
-      typename detail::BaseCLTypeConverter<DataT, NumElements>::DataType;
+  using DataType = typename detail::BaseCLTypeConverter<
+      DataT, detail::VectorLength<NumElements>::value>::DataType;
 
   template <bool B, class T, class F>
   using conditional_t = typename std::conditional<B, T, F>::type;
@@ -1792,16 +1801,18 @@ using cl_schar16 = cl_char16;
 // As a result half values will be converted to the integer and passed as a
 // kernel argument which is expected to be floating point number.
 #ifndef __SYCL_DEVICE_ONLY__
-template <int NumElements, typename CLType> struct alignas(CLType) half_vec {
-  std::array<half, NumElements> s;
+template <int NumElements>
+struct alignas(
+    cl::sycl::detail::VectorAlignment<half, NumElements>::value) half_vec {
+  std::array<half, cl::sycl::detail::VectorLength<NumElements>::value> s;
 };
 
-typedef half __half_t;
-typedef half_vec<2, cl_half2> __half2_vec_t;
-typedef half_vec<4, cl_half3> __half3_vec_t;
-typedef half_vec<4, cl_half4> __half4_vec_t;
-typedef half_vec<8, cl_half8> __half8_vec_t;
-typedef half_vec<16, cl_half16> __half16_vec_t;
+using __half_t = half;
+using __half2_vec_t = half_vec<2>;
+using __half3_vec_t = half_vec<4>;
+using __half4_vec_t = half_vec<4>;
+using __half8_vec_t = half_vec<8>;
+using __half16_vec_t = half_vec<16>;
 #endif
 
 #define GET_CL_HALF_TYPE(target, num) __##target##num##_vec_t
diff --git a/sycl/test/basic_tests/aliases.cpp b/sycl/test/basic_tests/aliases.cpp
@@ -0,0 +1,108 @@
+// RUN: %clangxx -fsycl %s -o %t.out -lOpenCL
+//==------------ aliases.cpp - SYCL type aliases test ----------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+#include <CL/sycl.hpp>
+#include <CL/sycl/detail/common.hpp>
+#include <cassert>
+#include <iostream>
+#include <type_traits>
+
+using namespace std;
+
+using cl_schar = cl_char;
+using cl_schar4 = cl_char4;
+
+namespace s = cl::sycl;
+
+#define CHECK_TYPE(TYPE)                                                       \
+  static_assert(sizeof(cl_##TYPE) == sizeof(s::cl_##TYPE), "")
+
+#define CHECK_SIZE(TYPE, SIZE) static_assert(sizeof(TYPE) == SIZE, "");
+
+#define CHECK_SIZE_VEC_N(TYPE, N)                                              \
+  static_assert(N * sizeof(TYPE) == sizeof(s::vec<TYPE, N>), "");
+
+#define CHECK_SIZE_VEC_N3(TYPE)                                                \
+  static_assert(sizeof(s::vec<TYPE, 3>) == sizeof(s::vec<TYPE, 4>), "");
+
+#define CHECK_SIZE_VEC(TYPE)                                                   \
+  CHECK_SIZE_VEC_N(TYPE, 2);                                                   \
+  CHECK_SIZE_VEC_N3(TYPE);                                                     \
+  CHECK_SIZE_VEC_N(TYPE, 4);                                                   \
+  CHECK_SIZE_VEC_N(TYPE, 8);                                                   \
+  CHECK_SIZE_VEC_N(TYPE, 16);
+
+#define CHECK_SIZE_TYPE_I(TYPE, SIZE)                                          \
+  CHECK_SIZE(TYPE, SIZE)                                                       \
+  static_assert(std::is_signed<TYPE>::value, "");
+
+#define CHECK_SIZE_TYPE_UI(TYPE, SIZE)                                         \
+  CHECK_SIZE(TYPE, SIZE)                                                       \
+  static_assert(std::is_unsigned<TYPE>::value, "");
+
+#define CHECK_SIZE_TYPE_F(TYPE, SIZE)                                          \
+  CHECK_SIZE(TYPE, SIZE)                                                       \
+  static_assert(std::numeric_limits<TYPE>::is_iec559, "");
+
+int main() {
+  CHECK_TYPE(bool);
+  CHECK_TYPE(char);
+  CHECK_TYPE(schar);
+  CHECK_TYPE(uchar);
+  CHECK_TYPE(short);
+  CHECK_TYPE(ushort);
+  CHECK_TYPE(half);
+  CHECK_TYPE(int);
+  CHECK_TYPE(uint);
+  CHECK_TYPE(long);
+  CHECK_TYPE(ulong);
+  CHECK_TYPE(float);
+  CHECK_TYPE(double);
+  CHECK_TYPE(char2);
+  CHECK_TYPE(uchar3);
+  CHECK_TYPE(short4);
+  CHECK_TYPE(ushort8);
+  CHECK_TYPE(half16);
+  CHECK_TYPE(int2);
+  CHECK_TYPE(uint3);
+  CHECK_TYPE(long4);
+  CHECK_TYPE(schar4);
+  CHECK_TYPE(ulong8);
+  CHECK_TYPE(float16);
+  CHECK_TYPE(double2);
+
+  // Table 4.93: Scalar data type aliases supported by SYCL
+  CHECK_SIZE_TYPE_UI(s::byte, 1);
+
+  CHECK_SIZE_TYPE_I(s::cl_char, 1);
+  CHECK_SIZE_TYPE_I(s::cl_short, 2);
+  CHECK_SIZE_TYPE_I(s::cl_int, 4);
+  CHECK_SIZE_TYPE_I(s::cl_long, 8);
+
+  CHECK_SIZE_TYPE_UI(s::cl_uchar, 1);
+  CHECK_SIZE_TYPE_UI(s::cl_ushort, 2);
+  CHECK_SIZE_TYPE_UI(s::cl_uint, 4);
+  CHECK_SIZE_TYPE_UI(s::cl_ulong, 8);
+
+  CHECK_SIZE_TYPE_F(s::cl_float, 4);
+  CHECK_SIZE_TYPE_F(s::cl_double, 8);
+  CHECK_SIZE(s::cl_half, 2);
+
+  CHECK_SIZE_VEC(s::cl_char);
+  CHECK_SIZE_VEC(s::cl_schar);
+  CHECK_SIZE_VEC(s::cl_uchar);
+  CHECK_SIZE_VEC(s::cl_short);
+  CHECK_SIZE_VEC(s::cl_ushort);
+  CHECK_SIZE_VEC(s::cl_half);
+  CHECK_SIZE_VEC(s::cl_int);
+  CHECK_SIZE_VEC(s::cl_uint);
+  CHECK_SIZE_VEC(s::cl_long);
+  CHECK_SIZE_VEC(s::cl_ulong);
+  CHECK_SIZE_VEC(s::cl_float);
+  CHECK_SIZE_VEC(s::cl_double);
+}
diff --git a/sycl/test/basic_tests/types.cpp b/sycl/test/basic_tests/types.cpp
@@ -6,119 +6,56 @@
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
-#include <CL/sycl.hpp>
-#include <CL/sycl/detail/common.hpp>
-#include <cassert>
-#include <iostream>
-#include <type_traits>
-
-using namespace std;
-
-using cl_schar = cl_char;
-using cl_schar4 = cl_char4;
-
-namespace s = cl::sycl;
-
-#define CHECK_TYPE(type)                                                       \
-  static_assert(sizeof(cl_##type) == sizeof(cl::sycl::cl_##type), "Wrong "     \
-                                                                  "size")
 
-#define CHECK_SIZE(T, S) static_assert(sizeof(T) == S, "Wrong size of type");
-
-#define CHECK_SIZE_VEC_N(T, n)                                \
-  static_assert(n * sizeof(T) == sizeof(cl::sycl::vec<T, n>), \
-                    "Wrong size of vec<type>");
-
-#define CHECK_SIZE_VEC_N3(T)                                  \
-  static_assert(sizeof(cl::sycl::vec<T, 3>) == sizeof(cl::sycl::vec<T, 4>), \
-                    "Wrong size of vec<cl_type3>");
-
-#define CHECK_SIZE_VEC(T) \
-  CHECK_SIZE_VEC_N(T, 2); \
-  CHECK_SIZE_VEC_N3(T);   \
-  CHECK_SIZE_VEC_N(T, 4); \
-  CHECK_SIZE_VEC_N(T, 8); \
-  CHECK_SIZE_VEC_N(T, 16);
+#include <CL/sycl.hpp>
 
-#define CHECK_SIZE_TYPE_I(T, S)                                                \
-  CHECK_SIZE(T, S)                                                             \
-  static_assert(std::is_signed<T>::value, "Expected signed type");
+template <typename T, int N> inline void checkVectorSizeAndAlignment() {
+  using VectorT = cl::sycl::vec<T, N>;
+  constexpr auto RealLength = (N != 3 ? N : 4);
+  static_assert(sizeof(VectorT) == (sizeof(T) * RealLength),
+                "Wrong size of vec<T, N>");
+  static_assert(alignof(VectorT) == (alignof(T) * RealLength),
+                "Wrong alignment of vec<T, N>");
+}
 
-#define CHECK_SIZE_TYPE_UI(T, S)                                               \
-  CHECK_SIZE(T, S)                                                             \
-  static_assert(std::is_unsigned<T>::value, "Expected unsigned type");
+template <typename T> inline void checkVectorsWithN() {
+  checkVectorSizeAndAlignment<T, 1>();
+  checkVectorSizeAndAlignment<T, 2>();
+  checkVectorSizeAndAlignment<T, 3>();
+  checkVectorSizeAndAlignment<T, 4>();
+  checkVectorSizeAndAlignment<T, 8>();
+  checkVectorSizeAndAlignment<T, 16>();
+}
 
-#define CHECK_SIZE_TYPE_F(T, S)                                                \
-  CHECK_SIZE(T, S)                                                             \
-  static_assert(std::numeric_limits<T>::is_iec559,                             \
-                "Expected type conformed to the IEEE 754");
+inline void checkVectors() {
+  checkVectorsWithN<half>();
+  checkVectorsWithN<float>();
+  checkVectorsWithN<double>();
+  checkVectorsWithN<char>();
+  checkVectorsWithN<signed char>();
+  checkVectorsWithN<unsigned char>();
+  checkVectorsWithN<signed short>();
+  checkVectorsWithN<unsigned short>();
+  checkVectorsWithN<signed int>();
+  checkVectorsWithN<unsigned int>();
+  checkVectorsWithN<signed long>();
+  checkVectorsWithN<unsigned long>();
+  checkVectorsWithN<signed long long>();
+  checkVectorsWithN<unsigned long long>();
+  checkVectorsWithN<::cl_char>();
+  checkVectorsWithN<::cl_uchar>();
+  checkVectorsWithN<::cl_short>();
+  checkVectorsWithN<::cl_ushort>();
+  checkVectorsWithN<::cl_int>();
+  checkVectorsWithN<::cl_uint>();
+  checkVectorsWithN<::cl_long>();
+  checkVectorsWithN<::cl_ulong>();
+  checkVectorsWithN<::cl_half>();
+  checkVectorsWithN<::cl_float>();
+  checkVectorsWithN<::cl_double>();
+}
 
 int main() {
-  CHECK_TYPE(bool);
-  CHECK_TYPE(char);
-  CHECK_TYPE(schar);
-  CHECK_TYPE(uchar);
-  CHECK_TYPE(short);
-  CHECK_TYPE(ushort);
-  CHECK_TYPE(half);
-  CHECK_TYPE(int);
-  CHECK_TYPE(uint);
-  CHECK_TYPE(long);
-  CHECK_TYPE(ulong);
-  CHECK_TYPE(float);
-  CHECK_TYPE(double);
-  CHECK_TYPE(char2);
-  CHECK_TYPE(uchar3);
-  CHECK_TYPE(short4);
-  CHECK_TYPE(ushort8);
-  CHECK_TYPE(half16);
-  CHECK_TYPE(int2);
-  CHECK_TYPE(uint3);
-  CHECK_TYPE(long4);
-  CHECK_TYPE(schar4);
-  CHECK_TYPE(ulong8);
-  CHECK_TYPE(float16);
-  CHECK_TYPE(double2);
-
-  // Table 4.93: Scalar data type aliases supported by SYCL
-  CHECK_SIZE_TYPE_UI(cl::sycl::byte, 1);
-
-  CHECK_SIZE_TYPE_I(cl::sycl::cl_char, 1);
-  CHECK_SIZE_TYPE_I(cl::sycl::cl_short, 2);
-  CHECK_SIZE_TYPE_I(cl::sycl::cl_int, 4);
-  CHECK_SIZE_TYPE_I(cl::sycl::cl_long, 8);
-
-  CHECK_SIZE_TYPE_UI(cl::sycl::cl_uchar, 1);
-  CHECK_SIZE_TYPE_UI(cl::sycl::cl_ushort, 2);
-  CHECK_SIZE_TYPE_UI(cl::sycl::cl_uint, 4);
-  CHECK_SIZE_TYPE_UI(cl::sycl::cl_ulong, 8);
-
-  CHECK_SIZE_TYPE_F(cl::sycl::cl_float, 4);
-  CHECK_SIZE_TYPE_F(cl::sycl::cl_double, 8);
-  // CHECK_SIZE_TYPE_F(cl::sycl::cl_half, 2);
-
-  CHECK_SIZE_VEC(char);
-  CHECK_SIZE_VEC(short);
-  CHECK_SIZE_VEC(unsigned short);
-  CHECK_SIZE_VEC(int);
-  CHECK_SIZE_VEC(unsigned int);
-  CHECK_SIZE_VEC(long);
-  CHECK_SIZE_VEC(unsigned long);
-  CHECK_SIZE_VEC(long long);
-  CHECK_SIZE_VEC(unsigned long long);
-  CHECK_SIZE_VEC(float);
-  CHECK_SIZE_VEC(double);
-
-  CHECK_SIZE_VEC(s::cl_char);
-  CHECK_SIZE_VEC(s::cl_schar);
-  CHECK_SIZE_VEC(s::cl_uchar);
-  CHECK_SIZE_VEC(s::cl_short);
-  CHECK_SIZE_VEC(s::cl_ushort);
-  CHECK_SIZE_VEC(s::cl_half);
-  CHECK_SIZE_VEC(s::cl_int);
-  CHECK_SIZE_VEC(s::cl_uint);
-  CHECK_SIZE_VEC(s::cl_long);
-  CHECK_SIZE_VEC(s::cl_ulong);
-  CHECK_SIZE_VEC(s::cl_float);
-  CHECK_SIZE_VEC(s::cl_double);
+  checkVectors();
+  return 0;
 }
