Skip to content

[SYCL] Use compiler's implementation for default marray constructor #20888

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

Open
bader wants to merge 3 commits into
base: sycl
Choose a base branch
from
Open

[SYCL] Use compiler's implementation for default marray constructor #20888

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
d2e8641
[SYCL] Use compiler's implementaton for default marray constructor
bader Dec 13, 2025
a30f4ac
Remove invalid test.
bader Dec 13, 2025
1e363e9
Initialize joint_matrix elements with initialization list.
bader Dec 13, 2025
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
4 changes: 2 additions & 2 deletions sycl/include/sycl/ext/oneapi/matrix/matrix-hip.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -70,7 +70,7 @@ template <> struct to_hip_type<int8_t> {
sycl::ext::oneapi::experimental::matrix::layout::row_major || \
Layout == \
sycl::ext::oneapi::experimental::matrix::layout::col_major>> { \
sycl::marray<TYPE, SIZE> wi_marray; \
sycl::marray<TYPE, SIZE> wi_marray{}; \
};

__SYCL_JOINT_MATRIX_OVERLOAD_ARR(bfloat16, a, 16, 16, 4)
Expand Down Expand Up @@ -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<TYPE, (M * N) / WAVEFRONT_SIZE> wi_marray; \
sycl::marray<TYPE, (M * N) / WAVEFRONT_SIZE> wi_marray{}; \
};

__SYCL_JOINT_MATRIX_OVERLOAD_ARR_ACC(float, 16, 16)
Expand Down
6 changes: 3 additions & 3 deletions sycl/include/sycl/ext/oneapi/matrix/matrix-tensorcores.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<TYPE, SIZE> wi_marray; \
marray<TYPE, SIZE> wi_marray{}; \
};

// m8n32k16
Expand Down Expand Up @@ -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<TYPE, SIZE> wi_marray; \
marray<TYPE, SIZE> wi_marray{}; \
};

__SYCL_JOINT_MATRIX_OVERLOAD_ARR_ACC(half, 8, 32, 8)
Expand All @@ -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<TYPE, SIZE> wi_marray; \
marray<TYPE, SIZE> wi_marray{}; \
};
// m16n16k8 tf32 only
__SYCL_JOINT_MATRIX_OVERLOAD_ARR_PRECISION(
Expand Down
2 changes: 1 addition & 1 deletion sycl/include/sycl/marray.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -114,7 +114,7 @@ template <typename Type, std::size_t NumElements> class marray {
: MData{Arr[Is]...} {}

public:
constexpr marray() : MData{} {}
constexpr marray() = default;

explicit constexpr marray(const Type &Arg) : MData{Arg} {
initialize_data(Arg);
Expand Down
2 changes: 1 addition & 1 deletion sycl/test/basic_tests/SYCL-2020-spec-constants.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -37,7 +37,7 @@ constexpr sycl::specialization_id<uint64_t> uint64_id(81);
// constexpr sycl::specialization_id<half> half_id(9.0);
constexpr sycl::specialization_id<float> float_id(710.0);
constexpr sycl::specialization_id<double> double_id(11.0);
constexpr sycl::marray<double, 5> ma;
constexpr sycl::marray<double, 5> ma(33.0);
constexpr sycl::specialization_id<sycl::marray<double, 5>> marray_id5(151.0);
constexpr sycl::specialization_id<sycl::marray<double, 1>> marray_id1(116.0);
constexpr sycl::specialization_id<sycl::marray<double, 5>> marray_id_def(ma);
Expand Down
6 changes: 1 addition & 5 deletions sycl/test/basic_tests/marray/marray.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<bool, 3> 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);

Expand Down Expand Up @@ -201,7 +197,7 @@ int main() {
CheckBinOps<double>();

// check copyability
constexpr sycl::marray<double, 5> ma;
constexpr sycl::marray<double, 5> ma(54.0);
constexpr sycl::marray<double, 5> mb(ma);
constexpr sycl::marray<double, 5> mc = ma;

Expand Down
26 changes: 6 additions & 20 deletions sycl/test/check_device_code/group_shuffle.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -56,22 +56,15 @@ SYCL_EXTERNAL void test_shuffle1(sycl::sub_group &sg, vec<bfloat16, 4> *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]]
Expand Down Expand Up @@ -100,22 +93,15 @@ SYCL_EXTERNAL void test_shuffle2(sycl::sub_group &sg, marray<bfloat16, 4> *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]]
Expand Down
Loading