Skip to content

[SYCL] Fix vec class alignment on windows platform #4953

New issue

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

Sign up for GitHub

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

Already on GitHub? Sign in to your account

Merged
merged 3 commits into from
Nov 23, 2021
Merged

[SYCL] Fix vec class alignment on windows platform #4953

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
20263f2
[SYCL] Data alignment is applied for vector type under Windows
dongkyunahn-intel Nov 13, 2021
a3de672
clang-format fix
dongkyunahn-intel Nov 13, 2021
889f0ad
Warning message fix for stdcpp_compat.cpp
dongkyunahn-intel Nov 17, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
21 changes: 12 additions & 9 deletions sycl/include/CL/sycl/types.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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.")
Copy link
Contributor

Choose a reason for hiding this comment

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

@alexbatashev
If I remember correctly you found that such a change is an ABI breaking one, could you please confirm?

Copy link
Contributor

Choose a reason for hiding this comment

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

It is

Copy link
Contributor

Choose a reason for hiding this comment

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

Could you please remind why?

Copy link
Contributor

Choose a reason for hiding this comment

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

Type alignment is part of ABI: assembled code has no notion of structure layout, it operates raw bytes.

Copy link
Contributor

Choose a reason for hiding this comment

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

In this particular change I do not see how setting alignment can break anything:

  1. The vector type is not passed directly to functions, but as part of the struct, so it's either passed by pointer or thru a register.
  2. The structure has only one field which is already considered to be aligned by the clang as it's a vector type(generates vector instructions)
  3. We set the alignment of vector type being equal to the size of the field itself, so we should not change the size. of the structure to trigger it to be passed differently.

#define __SYCL_ALIGNED_VAR(type, x, var) \
type __declspec(align((x < 64) ? x : 64)) var
Copy link
Contributor

Choose a reason for hiding this comment

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

Can't we continue using alignas specifier?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

alignas failed as follows during kernel compilation for 'test_handler' from CTS.

In file included from C:\work\dongkyun\github\opensource\llvm_build\include\sycl\CL/sycl.hpp:11:
In file included from C:\work\dongkyun\github\opensource\llvm_build\include\sycl\CL/sycl/accessor.hpp:13:
In file included from C:\work\dongkyun\github\opensource\llvm_build\include\sycl\CL/sycl/buffer.hpp:11:
In file included from C:\work\dongkyun\github\opensource\llvm_build\include\sycl\CL/sycl/detail/buffer_impl.hpp:20:
C:\work\dongkyun\github\opensource\llvm_build\include\sycl\CL/sycl/types.hpp(1370,3): error: requested alignment is less than minimum alignment of 128 for type 'sycl::vec<double, 16>::DataType' (aka 'double __attribute__((ext_vector_type(16)))')
  __SYCL_ALIGNAS((detail::vector_alignment<DataT, NumElements>::value))
  ^
C:\work\dongkyun\github\opensource\llvm_build\include\sycl\CL/sycl/types.hpp(548,27): note: expanded from macro '__SYCL_ALIGNAS'
#define __SYCL_ALIGNAS(x) alignas((x<64)?x:64)
                          ^

#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
Expand Down Expand Up @@ -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,
Expand Down Expand Up @@ -2494,4 +2497,4 @@ struct CheckDeviceCopyable<
} // namespace sycl
} // __SYCL_INLINE_NAMESPACE(cl)

#undef __SYCL_ALIGNAS
#undef __SYCL_ALIGNED_VAR
3 changes: 1 addition & 2 deletions sycl/test/basic_tests/stdcpp_compat.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;

Expand Down