Skip to content

Commit 1c8123d

Browse files
dkhaldidkhaldi
authored
Merge pull request intel#1499 from dkhaldi/matrixfixes
[SYCL][Matrix] cherry pick fix query test and bfloat16 namespace
2 parents 05615aa + f5a7466 commit 1c8123d

11 files changed

+25
-192
lines changed

‎SYCL/Matrix/Legacy/XMX8/joint_matrix_bfloat16.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -16,7 +16,7 @@
1616

1717
using namespace sycl;
1818
using namespace sycl::ext::oneapi::experimental::matrix;
19-
using bfloat16 = sycl::ext::oneapi::experimental::bfloat16;
19+
using bfloat16 = sycl::ext::oneapi::bfloat16;
2020

2121
#define SG_SZ 8
2222

‎SYCL/Matrix/Legacy/XMX8/joint_matrix_bfloat16_32x64.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -18,7 +18,7 @@
1818

1919
using namespace sycl;
2020
using namespace sycl::ext::oneapi::experimental::matrix;
21-
using bfloat16 = sycl::ext::oneapi::experimental::bfloat16;
21+
using bfloat16 = sycl::ext::oneapi::bfloat16;
2222

2323
#define SG_SZ 8
2424

‎SYCL/Matrix/Legacy/joint_matrix_bfloat16.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -16,7 +16,7 @@
1616

1717
using namespace sycl;
1818
using namespace sycl::ext::oneapi::experimental::matrix;
19-
using bfloat16 = sycl::ext::oneapi::experimental::bfloat16;
19+
using bfloat16 = sycl::ext::oneapi::bfloat16;
2020

2121
#define SG_SZ 16
2222

‎SYCL/Matrix/Legacy/joint_matrix_bfloat16_32x64.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -18,7 +18,7 @@
1818

1919
using namespace sycl;
2020
using namespace sycl::ext::oneapi::experimental::matrix;
21-
using bfloat16 = sycl::ext::oneapi::experimental::bfloat16;
21+
using bfloat16 = sycl::ext::oneapi::bfloat16;
2222

2323
#define SG_SZ 16
2424

‎SYCL/Matrix/joint_matrix_query_default.cpp

Lines changed: 17 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
//==-------- joint_matrix_query.cpp - DPC++ joint_matrix------------ ----==//
1+
//==-------- joint_matrix_query_default.cpp - DPC++ joint_matrix-----------==//
22
//
33
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
44
// See https://llvm.org/LICENSE.txt for license information.
@@ -7,7 +7,7 @@
77
//===----------------------------------------------------------------------===//
88
// REQUIRES: matrix
99

10-
// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=1
10+
// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4
1111
// RUN: %CPU_RUN_PLACEHOLDER %t.out
1212

1313
#include <iostream>
@@ -38,9 +38,9 @@ void matrix_multiply(big_matrix<T1, NUM_ROWS_C, NUM_COLS_C> &C,
3838
assert(NUM_ROWS_C == NUM_ROWS_A && NUM_COLS_A == NUM_ROWS_B * 4);
3939

4040
using myparams2 = tpu_params<tpu::amx, int8_t, int8_t, int>;
41-
constexpr int TM = myparams2::defaultM;
42-
constexpr int TN = myparams2::defaultN;
43-
constexpr int TK = myparams2::defaultK;
41+
constexpr int TM = myparams2::M;
42+
constexpr int TN = myparams2::N;
43+
constexpr int TK = myparams2::K;
4444

4545
std::cout << "AMX query sizes are: M " << TM << " N " << TN << " K " << TK
4646
<< std::endl;
@@ -60,7 +60,7 @@ void matrix_multiply(big_matrix<T1, NUM_ROWS_C, NUM_COLS_C> &C,
6060

6161
cgh.parallel_for<class imatrix>(
6262
nd_range<2>({NDRangeM, NDRangeN * SG_SZ}, {1, 1 * SG_SZ}),
63-
[ accA, accB, accC, M, N, K ](nd_item<2> spmd_item)
63+
[accA, accB, accC, M, N, K](nd_item<2> spmd_item)
6464
[[intel::reqd_sub_group_size(SG_SZ)]]
6565

6666
{
@@ -74,29 +74,31 @@ void matrix_multiply(big_matrix<T1, NUM_ROWS_C, NUM_COLS_C> &C,
7474

7575
ext::oneapi::sub_group sg = spmd_item.get_sub_group();
7676

77-
myparams2::joint_matrix_a<sub_group> sub_a(sg);
78-
myparams2::joint_matrix_b<sub_group> sub_b(sg);
79-
myparams2::joint_matrix_c<sub_group> sub_c(sg);
77+
myparams2::joint_matrix_a<sub_group, layout::row_major> sub_a;
78+
myparams2::joint_matrix_b<
79+
sub_group, ext::intel::experimental::matrix::layout::packed>
80+
sub_b;
81+
myparams2::joint_matrix_accumulator<sub_group> sub_c;
8082

8183
joint_matrix_load(sg, sub_c,
8284
accC.get_pointer() + (sg_startx * TM) * N +
8385
sg_starty / SG_SZ * TN,
84-
N, matrix_layout::row_major);
86+
N, layout::row_major);
8587
for (int k = 0; k < K / TK; k += 1) {
8688
joint_matrix_load(
8789
sg, sub_a, accA.get_pointer() + (sg_startx * TM) * K + k * TK,
88-
K, matrix_layout::row_major);
90+
K);
8991
// Assuming B data is already in VNNI format.
9092
joint_matrix_load(sg, sub_b,
9193
accB.get_pointer() + (k * TK / 4) * (N * 4) +
9294
sg_starty / SG_SZ * TN * 4,
93-
N * 4, matrix_layout::packed_b);
95+
N * 4);
9496
sub_c = joint_matrix_mad(sg, sub_a, sub_b, sub_c);
9597
}
9698
joint_matrix_store(sg, sub_c,
9799
accC.get_pointer() + (sg_startx * TM) * N +
98100
sg_starty / SG_SZ * TN,
99-
N, matrix_layout::row_major);
101+
N, layout::row_major);
100102
}); // parallel for
101103
}).wait();
102104
}
@@ -159,8 +161,6 @@ int main() {
159161
res = false;
160162
}
161163
}
162-
if (res)
163-
std::cout << "passed\n";
164-
else
165-
std::cout << "failed\n";
164+
std::cout << (res ? "passed" : "failed") << std::endl;
165+
return !res;
166166
}

‎SYCL/Matrix/joint_matrix_query_use_default.cpp

Lines changed: 0 additions & 167 deletions
This file was deleted.

‎config_sycl/matrix_joint_matrix_query_use_default.info

Lines changed: 0 additions & 1 deletion
This file was deleted.

‎config_sycl/userdefinedreductions_user_defined_reductions_wg_size_larger_than_data_size.info

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1 @@
1+
SYCL/UserDefinedReductions/user_defined_reductions_wg_size_larger_than_data_size.cpp

‎config_sycl/matrix_joint_matrix_query_use_default.xml renamed to ‎config_sycl/userdefinedreductions_user_defined_reductions_wg_size_larger_than_data_size.xml

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
<?xml version="1.0" encoding="UTF-8" ?>
2-
<test name="matrix_joint_matrix_query_use_default" driverID="llvm_test_suite_sycl">
2+
<test name="userdefinedreductions_user_defined_reductions_wg_size_larger_than_data_size" driverID="llvm_test_suite_sycl">
33
<description>WARNING: DON'T UPDATE THIS FILE MANUALLY!!!
44
This config file auto-generated by suite_generator_sycl.pl.</description>
55
</test>

‎llvm_test_suite_sycl.xml

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -751,7 +751,6 @@ Sources repo https://github.com/intel-innersource/applications.compilers.tests.l
751751
<test configFile="config_sycl/matrix_joint_matrix_half.xml" splitGroup="matrix" testName="matrix_joint_matrix_half" />
752752
<test configFile="config_sycl/matrix_joint_matrix_int8_vnni.xml" splitGroup="matrix" testName="matrix_joint_matrix_int8_vnni" />
753753
<test configFile="config_sycl/matrix_joint_matrix_query_default.xml" splitGroup="matrix" testName="matrix_joint_matrix_query_default" />
754-
<test configFile="config_sycl/matrix_joint_matrix_query_use_default.xml" splitGroup="matrix" testName="matrix_joint_matrix_query_use_default" />
755754
<test configFile="config_sycl/matrix_joint_matrix_ss_int8.xml" splitGroup="matrix" testName="matrix_joint_matrix_ss_int8" />
756755
<test configFile="config_sycl/matrix_joint_matrix_su_int8.xml" splitGroup="matrix" testName="matrix_joint_matrix_su_int8" />
757756
<test configFile="config_sycl/matrix_joint_matrix_tensorcores.xml" splitGroup="matrix" testName="matrix_joint_matrix_tensorcores" />
@@ -1035,6 +1034,7 @@ Sources repo https://github.com/intel-innersource/applications.compilers.tests.l
10351034
<test configFile="config_sycl/tracing_image_printers.xml" splitGroup="tracing" testName="tracing_image_printers" />
10361035
<test configFile="config_sycl/tracing_pi_tracing_test.xml" splitGroup="tracing" testName="tracing_pi_tracing_test" />
10371036
<test configFile="config_sycl/userdefinedreductions_user_defined_reductions.xml" splitGroup="userdefinedreductions" testName="userdefinedreductions_user_defined_reductions" />
1037+
<test configFile="config_sycl/userdefinedreductions_user_defined_reductions_wg_size_larger_than_data_size.xml" splitGroup="userdefinedreductions" testName="userdefinedreductions_user_defined_reductions_wg_size_larger_than_data_size" />
10381038
<test configFile="config_sycl/usm_alloc_functions.xml" splitGroup="usm" testName="usm_alloc_functions" />
10391039
<test configFile="config_sycl/usm_allocator_container.xml" splitGroup="usm" testName="usm_allocator_container" />
10401040
<test configFile="config_sycl/usm_allocator_equal.xml" splitGroup="usm" testName="usm_allocator_equal" />

‎llvm_test_suite_sycl_valgrind.xml

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -751,7 +751,6 @@ Sources repo https://github.com/intel-innersource/applications.compilers.tests.l
751751
<test configFile="config_sycl/matrix_joint_matrix_half.xml" splitGroup="matrix" testName="matrix_joint_matrix_half" />
752752
<test configFile="config_sycl/matrix_joint_matrix_int8_vnni.xml" splitGroup="matrix" testName="matrix_joint_matrix_int8_vnni" />
753753
<test configFile="config_sycl/matrix_joint_matrix_query_default.xml" splitGroup="matrix" testName="matrix_joint_matrix_query_default" />
754-
<test configFile="config_sycl/matrix_joint_matrix_query_use_default.xml" splitGroup="matrix" testName="matrix_joint_matrix_query_use_default" />
755754
<test configFile="config_sycl/matrix_joint_matrix_ss_int8.xml" splitGroup="matrix" testName="matrix_joint_matrix_ss_int8" />
756755
<test configFile="config_sycl/matrix_joint_matrix_su_int8.xml" splitGroup="matrix" testName="matrix_joint_matrix_su_int8" />
757756
<test configFile="config_sycl/matrix_joint_matrix_tensorcores.xml" splitGroup="matrix" testName="matrix_joint_matrix_tensorcores" />
@@ -1035,6 +1034,7 @@ Sources repo https://github.com/intel-innersource/applications.compilers.tests.l
10351034
<test configFile="config_sycl/tracing_image_printers.xml" splitGroup="tracing" testName="tracing_image_printers" />
10361035
<test configFile="config_sycl/tracing_pi_tracing_test.xml" splitGroup="tracing" testName="tracing_pi_tracing_test" />
10371036
<test configFile="config_sycl/userdefinedreductions_user_defined_reductions.xml" splitGroup="userdefinedreductions" testName="userdefinedreductions_user_defined_reductions" />
1037+
<test configFile="config_sycl/userdefinedreductions_user_defined_reductions_wg_size_larger_than_data_size.xml" splitGroup="userdefinedreductions" testName="userdefinedreductions_user_defined_reductions_wg_size_larger_than_data_size" />
10381038
<test configFile="config_sycl/usm_alloc_functions.xml" splitGroup="usm" testName="usm_alloc_functions" />
10391039
<test configFile="config_sycl/usm_allocator_container.xml" splitGroup="usm" testName="usm_allocator_container" />
10401040
<test configFile="config_sycl/usm_allocator_equal.xml" splitGroup="usm" testName="usm_allocator_equal" />

0 commit comments

Comments
 (0)