From 826c56930e408155d2e633c8a81ed0234882c0e1 Mon Sep 17 00:00:00 2001
From: dongkyunahn-intel <dongkyun.ahn@intel.com>
Date: Tue, 23 Nov 2021 10:05:22 -0800
Subject: [PATCH] [SYCL] Fix vec class alignment on windows platform (#4953)

Currently the sycl::vec type can be copied in the way which doesn't
preserve the default alignment on windows. This can causes crashes
since the sycl:;vec code expects the vector to be aligned and uses
vector instructions. We used default alignment because we cannot
set correct alignment in all cases. The patch adds alignment of
vector types, if alignment required is larger than 64, it is limited to 64.
---
 sycl/include/CL/sycl/types.hpp          | 21 ++++++++++++---------
 sycl/test/basic_tests/stdcpp_compat.cpp |  3 +--
 2 files changed, 13 insertions(+), 11 deletions(-)

diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp
index d6dacff4dfa23..b7f3eb2316249 100644
--- a/sycl/include/CL/sycl/types.hpp
+++ b/sycl/include/CL/sycl/types.hpp
@@ -542,12 +542,13 @@ using vec_data_t = typename detail::vec_helper<T>::RetType;
 // 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)
+Requested alignment applied, limited at 64.")
+#define __SYCL_ALIGNED_VAR(type, x, var)                                       \
+  type __declspec(align((x < 64) ? x : 64)) var
 #else
-#define __SYCL_ALIGNAS(N) alignas(N)
+#define __SYCL_ALIGNED_VAR(type, x, var) alignas(x) type var
 #endif
 
 /// Provides a cross-patform vector class template that works efficiently on
@@ -1363,12 +1364,14 @@ template <typename Type, int NumElements> class vec {
   }
 
   // fields
-  // Used "__SYCL_ALIGNAS" instead "alignas" to handle MSVC compiler.
+  // Used "__SYCL_ALIGNED_VAR" instead "alignas" to handle MSVC compiler.
   // For MSVC compiler max alignment is 64, e.g. vec<double, 16> required
   // alignment of 128 and MSVC compiler cann't align a parameter with requested
-  // alignment of 128.
-  __SYCL_ALIGNAS((detail::vector_alignment<DataT, NumElements>::value))
-  DataType m_Data;
+  // alignment of 128. For alignment request larger than 64, 64-alignment
+  // is applied
+  __SYCL_ALIGNED_VAR(DataType,
+                     (detail::vector_alignment<DataT, NumElements>::value),
+                     m_Data);
 
   // friends
   template <typename T1, typename T2, typename T3, template <typename> class T4,
@@ -2497,4 +2500,4 @@ struct CheckDeviceCopyable<
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)
 
-#undef __SYCL_ALIGNAS
+#undef __SYCL_ALIGNED_VAR
diff --git a/sycl/test/basic_tests/stdcpp_compat.cpp b/sycl/test/basic_tests/stdcpp_compat.cpp
index a51501204fc8a..512a5cb9d9136 100644
--- a/sycl/test/basic_tests/stdcpp_compat.cpp
+++ b/sycl/test/basic_tests/stdcpp_compat.cpp
@@ -16,8 +16,7 @@
 // warning_extension-warning@* 0-1 {{#warning is a language extension}}
 //
 // The next warning is emitted for windows only
-// expected-warning@* 0-1 {{Alignment of class vec is not in accordance with SYCL specification requirements, a limitation of the MSVC compiler(Error C2719).Applied default alignment.}}
-
+// expected-warning@* 0-1 {{Alignment of class vec is not in accordance with SYCL specification requirements, a limitation of the MSVC compiler(Error C2719).Requested alignment applied, limited at 64.}}
 
 class KernelName1;