From d2e8641a7f6f5ce75c6492319ba0fd6128c88acc Mon Sep 17 00:00:00 2001 From: Alexey Bader Date: Fri, 12 Dec 2025 21:27:38 -0800 Subject: [PATCH 1/3] [SYCL] Use compiler's implementaton for default marray constructor Current implementation of marray default constructor uses list-initialization to initialize elements, which leads to "value-initalization" of each element instead of default-initialization required by the SYCL specification. The fix caught a few places in the tests where marray is used in the constexpr context without initialization. --- sycl/include/sycl/marray.hpp | 2 +- .../basic_tests/SYCL-2020-spec-constants.cpp | 2 +- sycl/test/basic_tests/marray/marray.cpp | 2 +- sycl/test/check_device_code/group_shuffle.cpp | 26 +++++-------------- 4 files changed, 9 insertions(+), 23 deletions(-) diff --git a/sycl/include/sycl/marray.hpp b/sycl/include/sycl/marray.hpp index e4c193b37b635..304718100143e 100644 --- a/sycl/include/sycl/marray.hpp +++ b/sycl/include/sycl/marray.hpp @@ -114,7 +114,7 @@ template class marray { : MData{Arr[Is]...} {} public: - constexpr marray() : MData{} {} + constexpr marray() = default; explicit constexpr marray(const Type &Arg) : MData{Arg} { initialize_data(Arg); diff --git a/sycl/test/basic_tests/SYCL-2020-spec-constants.cpp b/sycl/test/basic_tests/SYCL-2020-spec-constants.cpp index 136043dba5fd6..52a432b1e9697 100644 --- a/sycl/test/basic_tests/SYCL-2020-spec-constants.cpp +++ b/sycl/test/basic_tests/SYCL-2020-spec-constants.cpp @@ -37,7 +37,7 @@ constexpr sycl::specialization_id uint64_id(81); // constexpr sycl::specialization_id half_id(9.0); constexpr sycl::specialization_id float_id(710.0); constexpr sycl::specialization_id double_id(11.0); -constexpr sycl::marray ma; +constexpr sycl::marray ma(33.0); constexpr sycl::specialization_id> marray_id5(151.0); constexpr sycl::specialization_id> marray_id1(116.0); constexpr sycl::specialization_id> marray_id_def(ma); diff --git a/sycl/test/basic_tests/marray/marray.cpp b/sycl/test/basic_tests/marray/marray.cpp index b48fe70f17520..52703e3cbfe02 100644 --- a/sycl/test/basic_tests/marray/marray.cpp +++ b/sycl/test/basic_tests/marray/marray.cpp @@ -201,7 +201,7 @@ int main() { CheckBinOps(); // check copyability - constexpr sycl::marray ma; + constexpr sycl::marray ma(54.0); constexpr sycl::marray mb(ma); constexpr sycl::marray mc = ma; diff --git a/sycl/test/check_device_code/group_shuffle.cpp b/sycl/test/check_device_code/group_shuffle.cpp index 84ca9d63ff5b1..16eadf395a02b 100644 --- a/sycl/test/check_device_code/group_shuffle.cpp +++ b/sycl/test/check_device_code/group_shuffle.cpp @@ -56,22 +56,15 @@ SYCL_EXTERNAL void test_shuffle1(sycl::sub_group &sg, vec *buf, // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META29:![0-9]+]]) // CHECK-NEXT: call void @llvm.lifetime.start.p0(ptr nonnull [[AGG_TMP14_I]]), !noalias [[META26]] // CHECK-NEXT: store i64 [[TMP0]], ptr [[AGG_TMP14_I]], align 8, !noalias [[META26]] -// CHECK-NEXT: br label [[ARRAYINIT_BODY_I_I_I:%.*]] -// CHECK: arrayinit.body.i.i.i: -// CHECK-NEXT: [[ARRAYINIT_CUR_IDX_I_I_I:%.*]] = phi i64 [ 0, [[ENTRY:%.*]] ], [ [[ARRAYINIT_CUR_ADD_I_I_I:%.*]], [[ARRAYINIT_BODY_I_I_I]] ] -// CHECK-NEXT: [[ARRAYINIT_CUR_PTR_I_I_I:%.*]] = getelementptr inbounds i8, ptr [[REF_TMP]], i64 [[ARRAYINIT_CUR_IDX_I_I_I]] -// CHECK-NEXT: store i16 0, ptr [[ARRAYINIT_CUR_PTR_I_I_I]], align 2, !alias.scope [[META32:![0-9]+]] -// CHECK-NEXT: [[ARRAYINIT_CUR_ADD_I_I_I]] = add nuw nsw i64 [[ARRAYINIT_CUR_IDX_I_I_I]], 2 -// CHECK-NEXT: [[ARRAYINIT_DONE_I_I_I:%.*]] = icmp eq i64 [[ARRAYINIT_CUR_ADD_I_I_I]], 8 -// CHECK-NEXT: br i1 [[ARRAYINIT_DONE_I_I_I]], label [[FOR_COND_I_I:%.*]], label [[ARRAYINIT_BODY_I_I_I]] +// CHECK-NEXT: br label [[FOR_COND_I_I:%.*]] // CHECK: for.cond.i.i: -// CHECK-NEXT: [[S_0_I_I:%.*]] = phi i32 [ [[INC_I_I:%.*]], [[FOR_BODY_I_I:%.*]] ], [ 0, [[ARRAYINIT_BODY_I_I_I]] ] +// CHECK-NEXT: [[S_0_I_I:%.*]] = phi i32 [ 0, [[ENTRY:%.*]] ], [ [[INC_I_I:%.*]], [[FOR_BODY_I_I:%.*]] ] // CHECK-NEXT: [[CMP_I_I:%.*]] = icmp samesign ult i32 [[S_0_I_I]], 4 // CHECK-NEXT: br i1 [[CMP_I_I]], label [[FOR_BODY_I_I]], label [[_ZN4SYCL3_V117SELECT_FROM_GROUPINS0_9SUB_GROUPENS0_6MARRAYINS0_3EXT6ONEAPI8BFLOAT16ELM4EEEEENST9ENABLE_IFIXAAOOSR3STDE9IS_SAME_VINST5DECAYIT_E4TYPEES2_ESR4SYCL3EXT6ONEAPI12EXPERIMENTALE27IS_USER_CONSTRUCTED_GROUP_VISC_EOOSR3STDE23IS_TRIVIALLY_COPYABLE_VIT0_ESR6DETAIL6IS_VECISD_EE5VALUEESD_E4TYPEESA_SD_NSA_7ID_TYPEE_EXIT:%.*]] // CHECK: for.body.i.i: // CHECK-NEXT: [[CONV_I_I:%.*]] = zext nneg i32 [[S_0_I_I]] to i64 // CHECK-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds %"class.sycl::_V1::ext::oneapi::bfloat16", ptr [[AGG_TMP14_I]], i64 [[CONV_I_I]] -// CHECK-NEXT: [[TMP1:%.*]] = load i16, ptr [[ARRAYIDX_I_I_I]], align 2, !tbaa [[TBAA17]], !noalias [[META32]] +// CHECK-NEXT: [[TMP1:%.*]] = load i16, ptr [[ARRAYIDX_I_I_I]], align 2, !tbaa [[TBAA17]], !noalias [[META32:![0-9]+]] // CHECK-NEXT: [[CALL3_I_I_I_I:%.*]] = tail call spir_func noundef zeroext i16 @_Z28__spirv_SubgroupShuffleINTELtj(i16 noundef zeroext [[TMP1]], i32 noundef 1) #[[ATTR6]], !noalias [[META33:![0-9]+]] // CHECK-NEXT: [[ARRAYIDX_I13_I_I:%.*]] = getelementptr inbounds %"class.sycl::_V1::ext::oneapi::bfloat16", ptr [[REF_TMP]], i64 [[CONV_I_I]] // CHECK-NEXT: store i16 [[CALL3_I_I_I_I]], ptr [[ARRAYIDX_I13_I_I]], align 2, !tbaa [[TBAA17]], !alias.scope [[META32]] @@ -100,22 +93,15 @@ SYCL_EXTERNAL void test_shuffle2(sycl::sub_group &sg, marray *buf, // CHECK-NEXT: call void @llvm.memcpy.p0.p4.i64(ptr noundef nonnull align 8 dereferenceable(10) [[AGG_TMP14_I]], ptr addrspace(4) noundef align 2 dereferenceable(10) [[ARRAYIDX]], i64 10, i1 false) // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META37]]) // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META40:![0-9]+]]) -// CHECK-NEXT: br label [[ARRAYINIT_BODY_I_I_I:%.*]] -// CHECK: arrayinit.body.i.i.i: -// CHECK-NEXT: [[ARRAYINIT_CUR_IDX_I_I_I:%.*]] = phi i64 [ 0, [[ENTRY:%.*]] ], [ [[ARRAYINIT_CUR_ADD_I_I_I:%.*]], [[ARRAYINIT_BODY_I_I_I]] ] -// CHECK-NEXT: [[ARRAYINIT_CUR_PTR_I_I_I:%.*]] = getelementptr inbounds i8, ptr [[REF_TMP]], i64 [[ARRAYINIT_CUR_IDX_I_I_I]] -// CHECK-NEXT: store i16 0, ptr [[ARRAYINIT_CUR_PTR_I_I_I]], align 2, !alias.scope [[META43:![0-9]+]] -// CHECK-NEXT: [[ARRAYINIT_CUR_ADD_I_I_I]] = add nuw nsw i64 [[ARRAYINIT_CUR_IDX_I_I_I]], 2 -// CHECK-NEXT: [[ARRAYINIT_DONE_I_I_I:%.*]] = icmp eq i64 [[ARRAYINIT_CUR_ADD_I_I_I]], 10 -// CHECK-NEXT: br i1 [[ARRAYINIT_DONE_I_I_I]], label [[FOR_COND_I_I:%.*]], label [[ARRAYINIT_BODY_I_I_I]] +// CHECK-NEXT: br label [[FOR_COND_I_I:%.*]] // CHECK: for.cond.i.i: -// CHECK-NEXT: [[S_0_I_I:%.*]] = phi i32 [ [[INC_I_I:%.*]], [[FOR_BODY_I_I:%.*]] ], [ 0, [[ARRAYINIT_BODY_I_I_I]] ] +// CHECK-NEXT: [[S_0_I_I:%.*]] = phi i32 [ 0, [[ENTRY:%.*]] ], [ [[INC_I_I:%.*]], [[FOR_BODY_I_I:%.*]] ] // CHECK-NEXT: [[CMP_I_I:%.*]] = icmp samesign ult i32 [[S_0_I_I]], 5 // CHECK-NEXT: br i1 [[CMP_I_I]], label [[FOR_BODY_I_I]], label [[_ZN4SYCL3_V117SELECT_FROM_GROUPINS0_9SUB_GROUPENS0_6MARRAYINS0_3EXT6ONEAPI8BFLOAT16ELM5EEEEENST9ENABLE_IFIXAAOOSR3STDE9IS_SAME_VINST5DECAYIT_E4TYPEES2_ESR4SYCL3EXT6ONEAPI12EXPERIMENTALE27IS_USER_CONSTRUCTED_GROUP_VISC_EOOSR3STDE23IS_TRIVIALLY_COPYABLE_VIT0_ESR6DETAIL6IS_VECISD_EE5VALUEESD_E4TYPEESA_SD_NSA_7ID_TYPEE_EXIT:%.*]] // CHECK: for.body.i.i: // CHECK-NEXT: [[CONV_I_I:%.*]] = zext nneg i32 [[S_0_I_I]] to i64 // CHECK-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds %"class.sycl::_V1::ext::oneapi::bfloat16", ptr [[AGG_TMP14_I]], i64 [[CONV_I_I]] -// CHECK-NEXT: [[TMP0:%.*]] = load i16, ptr [[ARRAYIDX_I_I_I]], align 2, !tbaa [[TBAA17]], !noalias [[META43]] +// CHECK-NEXT: [[TMP0:%.*]] = load i16, ptr [[ARRAYIDX_I_I_I]], align 2, !tbaa [[TBAA17]], !noalias [[META43:![0-9]+]] // CHECK-NEXT: [[CALL3_I_I_I_I:%.*]] = tail call spir_func noundef zeroext i16 @_Z28__spirv_SubgroupShuffleINTELtj(i16 noundef zeroext [[TMP0]], i32 noundef 1) #[[ATTR6]], !noalias [[META44:![0-9]+]] // CHECK-NEXT: [[ARRAYIDX_I13_I_I:%.*]] = getelementptr inbounds %"class.sycl::_V1::ext::oneapi::bfloat16", ptr [[REF_TMP]], i64 [[CONV_I_I]] // CHECK-NEXT: store i16 [[CALL3_I_I_I_I]], ptr [[ARRAYIDX_I13_I_I]], align 2, !tbaa [[TBAA17]], !alias.scope [[META43]] From a30f4acd1a7ac193caea0386694453511a00a8c1 Mon Sep 17 00:00:00 2001 From: Alexey Bader Date: Fri, 12 Dec 2025 22:09:16 -0800 Subject: [PATCH 2/3] Remove invalid test. Default marray constructor does not zero initializes int32_t type elements. --- sycl/test/basic_tests/marray/marray.cpp | 4 ---- 1 file changed, 4 deletions(-) diff --git a/sycl/test/basic_tests/marray/marray.cpp b/sycl/test/basic_tests/marray/marray.cpp index 52703e3cbfe02..7a616ddfed9c8 100644 --- a/sycl/test/basic_tests/marray/marray.cpp +++ b/sycl/test/basic_tests/marray/marray.cpp @@ -129,16 +129,12 @@ int main() { CHECK_ALIAS(mfloat, float) CHECK_ALIAS(mdouble, double) - mint3 t000; mint3 t222{2}; mint3 t123{1, 2, 3}; mint3 tcpy{t123}; mint3 t___; sycl::marray b___; - // test default ctor - assert(t000[0] == 0 && t000[1] == 0 && t000[2] == 0); - // test constant ctor assert(t222[0] == 2 && t222[1] == 2 && t222[2] == 2); From 1e363e9400caf6f10599dacc65dde0fc7a128e07 Mon Sep 17 00:00:00 2001 From: Alexey Bader Date: Fri, 12 Dec 2025 22:22:03 -0800 Subject: [PATCH 3/3] Initialize joint_matrix elements with initialization list. --- sycl/include/sycl/ext/oneapi/matrix/matrix-hip.hpp | 4 ++-- sycl/include/sycl/ext/oneapi/matrix/matrix-tensorcores.hpp | 6 +++--- 2 files changed, 5 insertions(+), 5 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/matrix/matrix-hip.hpp b/sycl/include/sycl/ext/oneapi/matrix/matrix-hip.hpp index bf224e9e78711..6f55ae3f1d6cf 100644 --- a/sycl/include/sycl/ext/oneapi/matrix/matrix-hip.hpp +++ b/sycl/include/sycl/ext/oneapi/matrix/matrix-hip.hpp @@ -70,7 +70,7 @@ template <> struct to_hip_type { sycl::ext::oneapi::experimental::matrix::layout::row_major || \ Layout == \ sycl::ext::oneapi::experimental::matrix::layout::col_major>> { \ - sycl::marray wi_marray; \ + sycl::marray wi_marray{}; \ }; __SYCL_JOINT_MATRIX_OVERLOAD_ARR(bfloat16, a, 16, 16, 4) @@ -98,7 +98,7 @@ __SYCL_JOINT_MATRIX_OVERLOAD_ARR(int8_t, b, 16, 16, 4) struct joint_matrix_hip< \ TYPE, sycl::ext::oneapi::experimental::matrix::use::accumulator, M, N, \ sycl::ext::oneapi::experimental::matrix::layout::dynamic> { \ - sycl::marray wi_marray; \ + sycl::marray wi_marray{}; \ }; __SYCL_JOINT_MATRIX_OVERLOAD_ARR_ACC(float, 16, 16) diff --git a/sycl/include/sycl/ext/oneapi/matrix/matrix-tensorcores.hpp b/sycl/include/sycl/ext/oneapi/matrix/matrix-tensorcores.hpp index 27e0160e5c7bb..71ff9aab08bf0 100644 --- a/sycl/include/sycl/ext/oneapi/matrix/matrix-tensorcores.hpp +++ b/sycl/include/sycl/ext/oneapi/matrix/matrix-tensorcores.hpp @@ -52,7 +52,7 @@ struct joint_matrix_cuda; sycl::ext::oneapi::experimental::matrix::layout::row_major || \ Layout == \ sycl::ext::oneapi::experimental::matrix::layout::col_major>> { \ - marray wi_marray; \ + marray wi_marray{}; \ }; // m8n32k16 @@ -96,7 +96,7 @@ __SYCL_JOINT_MATRIX_OVERLOAD_ARR(double, b, 4, 8, 1) struct joint_matrix_cuda< \ TYPE, sycl::ext::oneapi::experimental::matrix::use::accumulator, M, N, \ sycl::ext::oneapi::experimental::matrix::layout::dynamic> { \ - marray wi_marray; \ + marray wi_marray{}; \ }; __SYCL_JOINT_MATRIX_OVERLOAD_ARR_ACC(half, 8, 32, 8) @@ -123,7 +123,7 @@ __SYCL_JOINT_MATRIX_OVERLOAD_ARR_ACC(double, 8, 8, 2) sycl::ext::oneapi::experimental::matrix::layout::row_major || \ Layout == \ sycl::ext::oneapi::experimental::matrix::layout::col_major>> { \ - marray wi_marray; \ + marray wi_marray{}; \ }; // m16n16k8 tf32 only __SYCL_JOINT_MATRIX_OVERLOAD_ARR_PRECISION(