From bf662640ee536fc3d2e489783eef9302ebb95191 Mon Sep 17 00:00:00 2001 From: AidanBeltonS Date: Tue, 16 Sep 2025 17:10:14 +0100 Subject: [PATCH] SWDEV-539805, SWDEV-553860 - Resolve GCC clang ABI mismatch and check vector alignment (#909) * SWDEV-539805 - Add checks for vector alignment and size * SWDEV-553860 - Alter alignment for gcc * SWDEV-553860 - Align fallback method * SWDEV-553860 - Alter alignment requirement --- .../hip/amd_detail/amd_hip_vector_types.h | 20 +- .../catch/unit/vector_types/vector_types.cc | 189 +++++++++++++++++- 2 files changed, 196 insertions(+), 13 deletions(-) diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_vector_types.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_vector_types.h index 7d0f327b8f..b96bcc1e16 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_vector_types.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_vector_types.h @@ -41,7 +41,7 @@ THE SOFTWARE. #define __HIP_USE_NATIVE_VECTOR__ 1 #define __NATIVE_VECTOR__(n, T) T __attribute__((ext_vector_type(n))) #else -#define __NATIVE_VECTOR__(n, T) alignas(__hip_vec_align_v()) T[n] +#define __NATIVE_VECTOR__(n, T) alignas(n * sizeof(T)) T[n] #endif #if defined(__cplusplus) @@ -51,10 +51,6 @@ THE SOFTWARE. #include #endif // defined(__HIPCC_RTC__) -template constexpr __hip_internal::size_t __hip_vec_align_v() { - return (_hip_N == 4 && alignof(T) == 8) ? 16 : _hip_N * alignof(T); -} - template struct HIP_vector_base; template struct HIP_vector_type; @@ -63,8 +59,9 @@ template __attribute__((always_inline)) __HOST_DEVI typename HIP_vector_base::Native_vec_* get_native_pointer(HIP_vector_base& base_vec) { static_assert(sizeof(base_vec) == sizeof(typename HIP_vector_base::Native_vec_)); - static_assert(__hip_internal::alignment_of>::value == - __hip_internal::alignment_of::Native_vec_>::value); + static_assert( + (__hip_internal::alignment_of>::value % + __hip_internal::alignment_of::Native_vec_>::value) == 0); return reinterpret_cast::Native_vec_*>(&base_vec); }; @@ -72,8 +69,9 @@ template __attribute__((always_inline)) __HOST_DEVICE__ const typename HIP_vector_base::Native_vec_* get_native_pointer(const HIP_vector_base& base_vec) { static_assert(sizeof(base_vec) == sizeof(typename HIP_vector_base::Native_vec_)); - static_assert(__hip_internal::alignment_of>::value == - __hip_internal::alignment_of::Native_vec_>::value); + static_assert( + (__hip_internal::alignment_of>::value % + __hip_internal::alignment_of::Native_vec_>::value) == 0); return reinterpret_cast::Native_vec_*>(&base_vec); }; } // Namespace hip_impl. @@ -111,7 +109,7 @@ template struct HIP_vector_base { HIP_vector_base& operator=(const HIP_vector_base&) = default; }; -template struct alignas(alignof(__NATIVE_VECTOR__(2, T))) HIP_vector_base { +template struct alignas(2 * sizeof(T)) HIP_vector_base { using Native_vec_ = __NATIVE_VECTOR__(2, T); T x, y; @@ -268,7 +266,7 @@ template struct HIP_vector_base { HIP_vector_base& operator=(HIP_vector_base&&) = default; }; -template struct alignas(__NATIVE_VECTOR__(4, T)) HIP_vector_base { +template struct alignas(4 * sizeof(T)) HIP_vector_base { using Native_vec_ = __NATIVE_VECTOR__(4, T); T x, y, z, w; diff --git a/projects/hip-tests/catch/unit/vector_types/vector_types.cc b/projects/hip-tests/catch/unit/vector_types/vector_types.cc index d3393de400..cb3c2f1fc2 100644 --- a/projects/hip-tests/catch/unit/vector_types/vector_types.cc +++ b/projects/hip-tests/catch/unit/vector_types/vector_types.cc @@ -298,7 +298,7 @@ TEMPLATE_TEST_CASE("Unit_VectorAndValueTypeOperations_SanityCheck_Basic_Device", /** * Test Description * ------------------------ - * - Checks that vectors can be used with structured bindigns + * - Checks that vectors can be used with structured bindings * - Tests from the host side * Test source * ------------------------ @@ -328,7 +328,7 @@ __global__ void generate_my_kernel() { static_assert(func()); } /** * Test Description * ------------------------ - * - Checks that vectors can be used with structured bindigns + * - Checks that vectors can work with constexpr * - Tests from the host and device side * Test source * ------------------------ @@ -341,6 +341,191 @@ TEST_CASE("Unit_VectorConstexpr_SanityCheck_Basic_host_device", "") { generate_my_kernel<<<1, 1>>>(); static_assert(func()); } + +struct padded_struct { + int2 data1; + float3 data2; +}; + +__host__ __device__ void check_alignment() { + // char/uchar + static_assert(std::alignment_of_v == 1); + static_assert(std::alignment_of_v == 2); + static_assert(std::alignment_of_v == 1); + static_assert(std::alignment_of_v == 4); + static_assert(std::alignment_of_v == 1); + static_assert(std::alignment_of_v == 2); + static_assert(std::alignment_of_v == 1); + static_assert(std::alignment_of_v == 4); + + // short/ushort + static_assert(std::alignment_of_v == 2); + static_assert(std::alignment_of_v == 4); + static_assert(std::alignment_of_v == 2); + static_assert(std::alignment_of_v == 8); + static_assert(std::alignment_of_v == 2); + static_assert(std::alignment_of_v == 4); + static_assert(std::alignment_of_v == 2); + static_assert(std::alignment_of_v == 8); + + // int/uint + static_assert(std::alignment_of_v == 4); + static_assert(std::alignment_of_v == 8); + static_assert(std::alignment_of_v == 4); + static_assert(std::alignment_of_v == 16); + static_assert(std::alignment_of_v == 4); + static_assert(std::alignment_of_v == 8); + static_assert(std::alignment_of_v == 4); + static_assert(std::alignment_of_v == 16); + + // long/ulong + constexpr size_t long_size = sizeof(long); // needed to handle MSVC long defintion + static_assert(std::alignment_of_v == long_size); + static_assert(std::alignment_of_v == 2 * long_size); + static_assert(std::alignment_of_v == long_size); + static_assert(std::alignment_of_v == 4 * long_size); + static_assert(std::alignment_of_v == long_size); + static_assert(std::alignment_of_v == 2 * long_size); + static_assert(std::alignment_of_v == long_size); + static_assert(std::alignment_of_v == 4 * long_size); + + // longlong/ulonglong + static_assert(std::alignment_of_v == 8); + static_assert(std::alignment_of_v == 16); + static_assert(std::alignment_of_v == 8); + static_assert(std::alignment_of_v == 32); + static_assert(std::alignment_of_v == 8); + static_assert(std::alignment_of_v == 16); + static_assert(std::alignment_of_v == 8); + static_assert(std::alignment_of_v == 32); + + // float + static_assert(std::alignment_of_v == 4); + static_assert(std::alignment_of_v == 8); + static_assert(std::alignment_of_v == 4); + static_assert(std::alignment_of_v == 16); + + // double + static_assert(std::alignment_of_v == 8); + static_assert(std::alignment_of_v == 16); + static_assert(std::alignment_of_v == 8); + static_assert(std::alignment_of_v == 32); + + // padded struct + static_assert(std::alignment_of_v == 8); +} + +__global__ void check_alignment_device() { check_alignment(); } + +/** + * Test Description + * ------------------------ + * - Compile-time test checking vector type alignement + * - Tests from the host and device side + * Test source + * ------------------------ + * - unit/vector_types/vector_types.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Vector_alignment_check", "") { + check_alignment_device<<<1, 1>>>(); + check_alignment(); +} + +__host__ __device__ void check_size() { + // char/uchar + constexpr size_t char_size = sizeof(char); + static_assert(sizeof(char1) == 1 * char_size); + static_assert(sizeof(char2) == 2 * char_size); + static_assert(sizeof(char3) == 3 * char_size); + static_assert(sizeof(char4) == 4 * char_size); + static_assert(sizeof(uchar1) == 1 * char_size); + static_assert(sizeof(uchar2) == 2 * char_size); + static_assert(sizeof(uchar3) == 3 * char_size); + static_assert(sizeof(uchar4) == 4 * char_size); + + // short/ushort + constexpr size_t short_size = sizeof(short); + static_assert(sizeof(short1) == 1 * short_size); + static_assert(sizeof(short2) == 2 * short_size); + static_assert(sizeof(short3) == 3 * short_size); + static_assert(sizeof(short4) == 4 * short_size); + static_assert(sizeof(ushort1) == 1 * short_size); + static_assert(sizeof(ushort2) == 2 * short_size); + static_assert(sizeof(ushort3) == 3 * short_size); + static_assert(sizeof(ushort4) == 4 * short_size); + + // int/uint + constexpr size_t int_size = sizeof(int); + static_assert(sizeof(int1) == 1 * int_size); + static_assert(sizeof(int2) == 2 * int_size); + static_assert(sizeof(int3) == 3 * int_size); + static_assert(sizeof(int4) == 4 * int_size); + static_assert(sizeof(uint1) == 1 * int_size); + static_assert(sizeof(uint2) == 2 * int_size); + static_assert(sizeof(uint3) == 3 * int_size); + static_assert(sizeof(uint4) == 4 * int_size); + + // long/ulong + constexpr size_t long_size = sizeof(long); + static_assert(sizeof(long1) == 1 * long_size); + static_assert(sizeof(long2) == 2 * long_size); + static_assert(sizeof(long3) == 3 * long_size); + static_assert(sizeof(long4) == 4 * long_size); + static_assert(sizeof(ulong1) == 1 * long_size); + static_assert(sizeof(ulong2) == 2 * long_size); + static_assert(sizeof(ulong3) == 3 * long_size); + static_assert(sizeof(ulong4) == 4 * long_size); + + // longlong/ulonglong + constexpr size_t longlong_size = sizeof(long long); + static_assert(sizeof(longlong1) == 1 * longlong_size); + static_assert(sizeof(longlong2) == 2 * longlong_size); + static_assert(sizeof(longlong3) == 3 * longlong_size); + static_assert(sizeof(longlong4) == 4 * longlong_size); + static_assert(sizeof(ulonglong1) == 1 * longlong_size); + static_assert(sizeof(ulonglong2) == 2 * longlong_size); + static_assert(sizeof(ulonglong3) == 3 * longlong_size); + static_assert(sizeof(ulonglong4) == 4 * longlong_size); + + // float + constexpr size_t float_size = sizeof(float); + static_assert(sizeof(float1) == 1 * float_size); + static_assert(sizeof(float2) == 2 * float_size); + static_assert(sizeof(float3) == 3 * float_size); + static_assert(sizeof(float4) == 4 * float_size); + + // double + constexpr size_t double_size = sizeof(double); + static_assert(sizeof(double1) == 1 * double_size); + static_assert(sizeof(double2) == 2 * double_size); + static_assert(sizeof(double3) == 3 * double_size); + static_assert(sizeof(double4) == 4 * double_size); + + // padded struct + static_assert(sizeof(padded_struct) == 24); +} + +__global__ void check_size_device() { check_size(); } + +/** + * Test Description + * ------------------------ + * - Compile-time test checking vector type size + * - Tests from the host and device side + * Test source + * ------------------------ + * - unit/vector_types/vector_types.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Vector_size_check", "") { + check_size_device<<<1, 1>>>(); + check_size(); +} #endif // HT_AMD /**