Skip to content

Commit ba4f090

Browse files
baderbader
authored
[SYCL] Use compiler's implementation for default marray constructor (#20888)
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.
1 parent b3c6c0f commit ba4f090

File tree

6 files changed

+14
-32
lines changed

6 files changed

+14
-32
lines changed

‎sycl/include/sycl/ext/oneapi/matrix/matrix-hip.hpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -70,7 +70,7 @@ template <> struct to_hip_type<int8_t> {
7070
sycl::ext::oneapi::experimental::matrix::layout::row_major || \
7171
Layout == \
7272
sycl::ext::oneapi::experimental::matrix::layout::col_major>> { \
73-
sycl::marray<TYPE, SIZE> wi_marray; \
73+
sycl::marray<TYPE, SIZE> wi_marray{}; \
7474
};
7575

7676
__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)
9898
struct joint_matrix_hip< \
9999
TYPE, sycl::ext::oneapi::experimental::matrix::use::accumulator, M, N, \
100100
sycl::ext::oneapi::experimental::matrix::layout::dynamic> { \
101-
sycl::marray<TYPE, (M * N) / WAVEFRONT_SIZE> wi_marray; \
101+
sycl::marray<TYPE, (M * N) / WAVEFRONT_SIZE> wi_marray{}; \
102102
};
103103

104104
__SYCL_JOINT_MATRIX_OVERLOAD_ARR_ACC(float, 16, 16)

‎sycl/include/sycl/ext/oneapi/matrix/matrix-tensorcores.hpp‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -52,7 +52,7 @@ struct joint_matrix_cuda;
5252
sycl::ext::oneapi::experimental::matrix::layout::row_major || \
5353
Layout == \
5454
sycl::ext::oneapi::experimental::matrix::layout::col_major>> { \
55-
marray<TYPE, SIZE> wi_marray; \
55+
marray<TYPE, SIZE> wi_marray{}; \
5656
};
5757

5858
// m8n32k16
@@ -96,7 +96,7 @@ __SYCL_JOINT_MATRIX_OVERLOAD_ARR(double, b, 4, 8, 1)
9696
struct joint_matrix_cuda< \
9797
TYPE, sycl::ext::oneapi::experimental::matrix::use::accumulator, M, N, \
9898
sycl::ext::oneapi::experimental::matrix::layout::dynamic> { \
99-
marray<TYPE, SIZE> wi_marray; \
99+
marray<TYPE, SIZE> wi_marray{}; \
100100
};
101101

102102
__SYCL_JOINT_MATRIX_OVERLOAD_ARR_ACC(half, 8, 32, 8)
@@ -123,7 +123,7 @@ __SYCL_JOINT_MATRIX_OVERLOAD_ARR_ACC(double, 8, 8, 2)
123123
sycl::ext::oneapi::experimental::matrix::layout::row_major || \
124124
Layout == \
125125
sycl::ext::oneapi::experimental::matrix::layout::col_major>> { \
126-
marray<TYPE, SIZE> wi_marray; \
126+
marray<TYPE, SIZE> wi_marray{}; \
127127
};
128128
// m16n16k8 tf32 only
129129
__SYCL_JOINT_MATRIX_OVERLOAD_ARR_PRECISION(

‎sycl/include/sycl/marray.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -114,7 +114,7 @@ template <typename Type, std::size_t NumElements> class marray {
114114
: MData{Arr[Is]...} {}
115115

116116
public:
117-
constexpr marray() : MData{} {}
117+
constexpr marray() = default;
118118

119119
explicit constexpr marray(const Type &Arg) : MData{Arg} {
120120
initialize_data(Arg);

‎sycl/test/basic_tests/SYCL-2020-spec-constants.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -37,7 +37,7 @@ constexpr sycl::specialization_id<uint64_t> uint64_id(81);
3737
// constexpr sycl::specialization_id<half> half_id(9.0);
3838
constexpr sycl::specialization_id<float> float_id(710.0);
3939
constexpr sycl::specialization_id<double> double_id(11.0);
40-
constexpr sycl::marray<double, 5> ma;
40+
constexpr sycl::marray<double, 5> ma(33.0);
4141
constexpr sycl::specialization_id<sycl::marray<double, 5>> marray_id5(151.0);
4242
constexpr sycl::specialization_id<sycl::marray<double, 1>> marray_id1(116.0);
4343
constexpr sycl::specialization_id<sycl::marray<double, 5>> marray_id_def(ma);

‎sycl/test/basic_tests/marray/marray.cpp‎

Lines changed: 1 addition & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -129,16 +129,12 @@ int main() {
129129
CHECK_ALIAS(mfloat, float)
130130
CHECK_ALIAS(mdouble, double)
131131

132-
mint3 t000;
133132
mint3 t222{2};
134133
mint3 t123{1, 2, 3};
135134
mint3 tcpy{t123};
136135
mint3 t___;
137136
sycl::marray<bool, 3> b___;
138137

139-
// test default ctor
140-
assert(t000[0] == 0 && t000[1] == 0 && t000[2] == 0);
141-
142138
// test constant ctor
143139
assert(t222[0] == 2 && t222[1] == 2 && t222[2] == 2);
144140

@@ -201,7 +197,7 @@ int main() {
201197
CheckBinOps<double>();
202198

203199
// check copyability
204-
constexpr sycl::marray<double, 5> ma;
200+
constexpr sycl::marray<double, 5> ma(54.0);
205201
constexpr sycl::marray<double, 5> mb(ma);
206202
constexpr sycl::marray<double, 5> mc = ma;
207203

‎sycl/test/check_device_code/group_shuffle.cpp‎

Lines changed: 6 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -56,22 +56,15 @@ SYCL_EXTERNAL void test_shuffle1(sycl::sub_group &sg, vec<bfloat16, 4> *buf,
5656
// CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META29:![0-9]+]])
5757
// CHECK-NEXT: call void @llvm.lifetime.start.p0(ptr nonnull [[AGG_TMP14_I]]), !noalias [[META26]]
5858
// CHECK-NEXT: store i64 [[TMP0]], ptr [[AGG_TMP14_I]], align 8, !noalias [[META26]]
59-
// CHECK-NEXT: br label [[ARRAYINIT_BODY_I_I_I:%.*]]
60-
// CHECK: arrayinit.body.i.i.i:
61-
// CHECK-NEXT: [[ARRAYINIT_CUR_IDX_I_I_I:%.*]] = phi i64 [ 0, [[ENTRY:%.*]] ], [ [[ARRAYINIT_CUR_ADD_I_I_I:%.*]], [[ARRAYINIT_BODY_I_I_I]] ]
62-
// CHECK-NEXT: [[ARRAYINIT_CUR_PTR_I_I_I:%.*]] = getelementptr inbounds i8, ptr [[REF_TMP]], i64 [[ARRAYINIT_CUR_IDX_I_I_I]]
63-
// CHECK-NEXT: store i16 0, ptr [[ARRAYINIT_CUR_PTR_I_I_I]], align 2, !alias.scope [[META32:![0-9]+]]
64-
// CHECK-NEXT: [[ARRAYINIT_CUR_ADD_I_I_I]] = add nuw nsw i64 [[ARRAYINIT_CUR_IDX_I_I_I]], 2
65-
// CHECK-NEXT: [[ARRAYINIT_DONE_I_I_I:%.*]] = icmp eq i64 [[ARRAYINIT_CUR_ADD_I_I_I]], 8
66-
// CHECK-NEXT: br i1 [[ARRAYINIT_DONE_I_I_I]], label [[FOR_COND_I_I:%.*]], label [[ARRAYINIT_BODY_I_I_I]]
59+
// CHECK-NEXT: br label [[FOR_COND_I_I:%.*]]
6760
// CHECK: for.cond.i.i:
68-
// CHECK-NEXT: [[S_0_I_I:%.*]] = phi i32 [ [[INC_I_I:%.*]], [[FOR_BODY_I_I:%.*]] ], [ 0, [[ARRAYINIT_BODY_I_I_I]] ]
61+
// CHECK-NEXT: [[S_0_I_I:%.*]] = phi i32 [ 0, [[ENTRY:%.*]] ], [ [[INC_I_I:%.*]], [[FOR_BODY_I_I:%.*]] ]
6962
// CHECK-NEXT: [[CMP_I_I:%.*]] = icmp samesign ult i32 [[S_0_I_I]], 4
7063
// 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:%.*]]
7164
// CHECK: for.body.i.i:
7265
// CHECK-NEXT: [[CONV_I_I:%.*]] = zext nneg i32 [[S_0_I_I]] to i64
7366
// CHECK-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds %"class.sycl::_V1::ext::oneapi::bfloat16", ptr [[AGG_TMP14_I]], i64 [[CONV_I_I]]
74-
// CHECK-NEXT: [[TMP1:%.*]] = load i16, ptr [[ARRAYIDX_I_I_I]], align 2, !tbaa [[TBAA17]], !noalias [[META32]]
67+
// CHECK-NEXT: [[TMP1:%.*]] = load i16, ptr [[ARRAYIDX_I_I_I]], align 2, !tbaa [[TBAA17]], !noalias [[META32:![0-9]+]]
7568
// 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]+]]
7669
// CHECK-NEXT: [[ARRAYIDX_I13_I_I:%.*]] = getelementptr inbounds %"class.sycl::_V1::ext::oneapi::bfloat16", ptr [[REF_TMP]], i64 [[CONV_I_I]]
7770
// 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<bfloat16, 4> *buf,
10093
// 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)
10194
// CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META37]])
10295
// CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META40:![0-9]+]])
103-
// CHECK-NEXT: br label [[ARRAYINIT_BODY_I_I_I:%.*]]
104-
// CHECK: arrayinit.body.i.i.i:
105-
// CHECK-NEXT: [[ARRAYINIT_CUR_IDX_I_I_I:%.*]] = phi i64 [ 0, [[ENTRY:%.*]] ], [ [[ARRAYINIT_CUR_ADD_I_I_I:%.*]], [[ARRAYINIT_BODY_I_I_I]] ]
106-
// CHECK-NEXT: [[ARRAYINIT_CUR_PTR_I_I_I:%.*]] = getelementptr inbounds i8, ptr [[REF_TMP]], i64 [[ARRAYINIT_CUR_IDX_I_I_I]]
107-
// CHECK-NEXT: store i16 0, ptr [[ARRAYINIT_CUR_PTR_I_I_I]], align 2, !alias.scope [[META43:![0-9]+]]
108-
// CHECK-NEXT: [[ARRAYINIT_CUR_ADD_I_I_I]] = add nuw nsw i64 [[ARRAYINIT_CUR_IDX_I_I_I]], 2
109-
// CHECK-NEXT: [[ARRAYINIT_DONE_I_I_I:%.*]] = icmp eq i64 [[ARRAYINIT_CUR_ADD_I_I_I]], 10
110-
// CHECK-NEXT: br i1 [[ARRAYINIT_DONE_I_I_I]], label [[FOR_COND_I_I:%.*]], label [[ARRAYINIT_BODY_I_I_I]]
96+
// CHECK-NEXT: br label [[FOR_COND_I_I:%.*]]
11197
// CHECK: for.cond.i.i:
112-
// CHECK-NEXT: [[S_0_I_I:%.*]] = phi i32 [ [[INC_I_I:%.*]], [[FOR_BODY_I_I:%.*]] ], [ 0, [[ARRAYINIT_BODY_I_I_I]] ]
98+
// CHECK-NEXT: [[S_0_I_I:%.*]] = phi i32 [ 0, [[ENTRY:%.*]] ], [ [[INC_I_I:%.*]], [[FOR_BODY_I_I:%.*]] ]
11399
// CHECK-NEXT: [[CMP_I_I:%.*]] = icmp samesign ult i32 [[S_0_I_I]], 5
114100
// 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:%.*]]
115101
// CHECK: for.body.i.i:
116102
// CHECK-NEXT: [[CONV_I_I:%.*]] = zext nneg i32 [[S_0_I_I]] to i64
117103
// CHECK-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds %"class.sycl::_V1::ext::oneapi::bfloat16", ptr [[AGG_TMP14_I]], i64 [[CONV_I_I]]
118-
// CHECK-NEXT: [[TMP0:%.*]] = load i16, ptr [[ARRAYIDX_I_I_I]], align 2, !tbaa [[TBAA17]], !noalias [[META43]]
104+
// CHECK-NEXT: [[TMP0:%.*]] = load i16, ptr [[ARRAYIDX_I_I_I]], align 2, !tbaa [[TBAA17]], !noalias [[META43:![0-9]+]]
119105
// 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]+]]
120106
// CHECK-NEXT: [[ARRAYIDX_I13_I_I:%.*]] = getelementptr inbounds %"class.sycl::_V1::ext::oneapi::bfloat16", ptr [[REF_TMP]], i64 [[CONV_I_I]]
121107
// CHECK-NEXT: store i16 [[CALL3_I_I_I_I]], ptr [[ARRAYIDX_I13_I_I]], align 2, !tbaa [[TBAA17]], !alias.scope [[META43]]

0 commit comments

Comments
 (0)