Skip to content

Commit bb9244f

Browse files
YuriPlyakhinYuriPlyakhin
authored
[SYCL][Joint Matrix] Resolve ambiguous fabs calls, refactor tests to remove some duplication. (#11406)
Main purpose of this PR is to resolve ambiguous fabs calls. While doing this, I also did some refactoring to reduce amount of duplication.
1 parent e7139b0 commit bb9244f

File tree

53 files changed

+105
-428
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

53 files changed

+105
-428
lines changed

‎sycl/test-e2e/Matrix/Legacy/XMX8/joint_matrix_bf16.cpp

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -10,8 +10,7 @@
1010
// RUN: %{build} -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=1
1111
// RUN: %{run} %t.out
1212

13-
#include <iostream>
14-
#include <sycl/sycl.hpp>
13+
#include "../../common.hpp"
1514

1615
using namespace sycl;
1716
using namespace sycl::ext::oneapi::experimental::matrix;

‎sycl/test-e2e/Matrix/Legacy/XMX8/joint_matrix_bfloat16.cpp

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -10,12 +10,10 @@
1010
// RUN: %{build} -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=1
1111
// RUN: %{run} %t.out
1212

13-
#include <iostream>
14-
#include <sycl/sycl.hpp>
13+
#include "../../common.hpp"
1514

1615
using namespace sycl;
1716
using namespace sycl::ext::oneapi::experimental::matrix;
18-
using bfloat16 = sycl::ext::oneapi::bfloat16;
1917

2018
#define SG_SZ 8
2119

‎sycl/test-e2e/Matrix/Legacy/XMX8/joint_matrix_bfloat16_32x64.cpp

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -12,12 +12,10 @@
1212

1313
// XFAIL: *
1414

15-
#include <iostream>
16-
#include <sycl/sycl.hpp>
15+
#include "../../common.hpp"
1716

1817
using namespace sycl;
1918
using namespace sycl::ext::oneapi::experimental::matrix;
20-
using bfloat16 = sycl::ext::oneapi::bfloat16;
2119

2220
#define SG_SZ 8
2321

‎sycl/test-e2e/Matrix/Legacy/joint_matrix_bf16.cpp

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -10,8 +10,7 @@
1010
// RUN: %{build} -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=1
1111
// RUN: %{run} %t.out
1212

13-
#include <iostream>
14-
#include <sycl/sycl.hpp>
13+
#include "../common.hpp"
1514

1615
using namespace sycl;
1716
using namespace sycl::ext::oneapi::experimental::matrix;

‎sycl/test-e2e/Matrix/Legacy/joint_matrix_bf16_impl.hpp

Lines changed: 1 addition & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -2,18 +2,6 @@
22
#define TN SG_SZ
33
#define TK 16
44

5-
#define BF16_EPSILON 0.00781250
6-
7-
template <typename T, size_t NUM_ROWS, size_t NUM_COLS> struct big_matrix {
8-
public:
9-
T *mat;
10-
11-
public:
12-
T *get_data() { return mat; }
13-
void set_data(T *data) { mat = data; }
14-
big_matrix(T *data) : mat(data) {}
15-
};
16-
175
template <typename T1, typename T2, size_t NUM_ROWS_A, size_t NUM_COLS_A,
186
size_t NUM_ROWS_B, size_t NUM_COLS_B, size_t NUM_ROWS_C,
197
size_t NUM_COLS_C>
@@ -153,13 +141,7 @@ int main() {
153141
matrix_multiply_ref((int32_t *)A, (int32_t *)B, (int32_t *)D, MATRIX_M,
154142
MATRIX_N, MATRIX_K / 2);
155143

156-
bool res = true;
157-
for (int i = 0; i < MATRIX_M; i++) {
158-
for (int j = 0; j < MATRIX_N; j++) {
159-
if (fabs(C[i][j] - D[i][j]) > BF16_EPSILON)
160-
res = false;
161-
}
162-
}
144+
bool res = matrix_compare(MATRIX_M, MATRIX_N, (float *)C, (float *)D);
163145
if (res)
164146
std::cout << "passed\n";
165147
else

‎sycl/test-e2e/Matrix/Legacy/joint_matrix_bfloat16.cpp

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -10,12 +10,10 @@
1010
// RUN: %{build} -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=1
1111
// RUN: %{run} %t.out
1212

13-
#include <iostream>
14-
#include <sycl/sycl.hpp>
13+
#include "../common.hpp"
1514

1615
using namespace sycl;
1716
using namespace sycl::ext::oneapi::experimental::matrix;
18-
using bfloat16 = sycl::ext::oneapi::bfloat16;
1917

2018
#define SG_SZ 16
2119

‎sycl/test-e2e/Matrix/Legacy/joint_matrix_bfloat16_32x64.cpp

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -12,12 +12,11 @@
1212

1313
// XFAIL: *
1414

15+
#include "../common.hpp"
1516
#include <iostream>
16-
#include <sycl/sycl.hpp>
1717

1818
using namespace sycl;
1919
using namespace sycl::ext::oneapi::experimental::matrix;
20-
using bfloat16 = sycl::ext::oneapi::bfloat16;
2120

2221
#define SG_SZ 16
2322

‎sycl/test-e2e/Matrix/Legacy/joint_matrix_bfloat16_32x64_impl.hpp

Lines changed: 2 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -2,18 +2,6 @@
22
#define TN 64
33
#define TK 16
44

5-
#define BF16_EPSILON 0.00781250
6-
7-
template <typename T, size_t NUM_ROWS, size_t NUM_COLS> struct big_matrix {
8-
private:
9-
T *mat;
10-
11-
public:
12-
T *get_data() { return mat; }
13-
void set_data(T *data) { mat = data; }
14-
big_matrix(T *data) : mat(data) {}
15-
};
16-
175
template <typename T1, typename T2, size_t M, size_t N, size_t K>
186
void matrix_multiply(big_matrix<T1, M, N> &C, big_matrix<T2, M, K> &A,
197
big_matrix<T2, K / 2, N * 2> &B) {
@@ -150,13 +138,8 @@ int main() {
150138
matrix_multiply_ref((int32_t *)Aref, (int32_t *)Bref, (int32_t *)D, MATRIX_M,
151139
MATRIX_N, MATRIX_K / 2);
152140

153-
bool res = true;
154-
for (int i = 0; i < MATRIX_M; i++) {
155-
for (int j = 0; j < MATRIX_N; j++) {
156-
if (fabs(C[i][j] - D[i][j]) > BF16_EPSILON)
157-
res = false;
158-
}
159-
}
141+
bool res = matrix_compare(MATRIX_M, MATRIX_N, (float *)C, (float *)D);
142+
160143
if (res)
161144
std::cout << "passed\n";
162145
else

‎sycl/test-e2e/Matrix/Legacy/joint_matrix_bfloat16_colmajorA_colmajorB.cpp

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -15,12 +15,10 @@
1515

1616
// XFAIL: gpu
1717

18-
#include <iostream>
19-
#include <sycl/sycl.hpp>
18+
#include "../common.hpp"
2019

2120
using namespace sycl;
2221
using namespace sycl::ext::oneapi::experimental::matrix;
23-
using bfloat16 = sycl::ext::oneapi::bfloat16;
2422

2523
#define SG_SZ 16
2624

‎sycl/test-e2e/Matrix/Legacy/joint_matrix_bfloat16_colmajorA_colmajorB_impl.hpp

Lines changed: 1 addition & 25 deletions
Original file line numberDiff line numberDiff line change
@@ -1,17 +1,6 @@
11
#define TM 8
22
#define TN SG_SZ
33
#define TK 16
4-
#define BF16_EPSILON 0.00781250
5-
6-
template <typename T, size_t NUM_ROWS, size_t NUM_COLS> struct big_matrix {
7-
private:
8-
T *mat;
9-
10-
public:
11-
T *get_data() { return mat; }
12-
void set_data(T *data) { mat = data; }
13-
big_matrix(T *data) : mat(data) {}
14-
};
154

165
template <typename T1, typename T2, size_t M, size_t N, size_t K>
176
void matrix_multiply(big_matrix<T1, M, N> &C, big_matrix<T2, M, K> &A,
@@ -81,13 +70,6 @@ bfloat16 B[MATRIX_N][MATRIX_K];
8170
float C[MATRIX_M][MATRIX_N];
8271
float D[MATRIX_M][MATRIX_N];
8372

84-
float make_fp32(bfloat16 x) {
85-
unsigned int y = *((int *)&x);
86-
y = y << 16;
87-
float *res = reinterpret_cast<float *>(&y);
88-
return *res;
89-
}
90-
9173
void matrix_multiply_ref(int M, int N, int K) {
9274
for (int m = 0; m < M; m++)
9375
for (int n = 0; n < N; n++) {
@@ -122,13 +104,7 @@ int main() {
122104
matrix_multiply(MC, MA, MB);
123105
matrix_multiply_ref(MATRIX_M, MATRIX_N, MATRIX_K);
124106

125-
bool res = true;
126-
for (int i = 0; i < MATRIX_M; i++) {
127-
for (int j = 0; j < MATRIX_N; j++) {
128-
if (fabs(C[i][j] - D[i][j]) > BF16_EPSILON)
129-
res = false;
130-
}
131-
}
107+
bool res = matrix_compare(MATRIX_M, MATRIX_N, (float *)C, (float *)D);
132108
std::cout << (res ? "passed" : "failed") << std::endl;
133109
return !res;
134110
}

‎sycl/test-e2e/Matrix/Legacy/joint_matrix_bfloat16_impl.hpp

Lines changed: 1 addition & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -2,18 +2,6 @@
22
#define TN SG_SZ
33
#define TK 16
44

5-
#define BF16_EPSILON 0.00781250
6-
7-
template <typename T, size_t NUM_ROWS, size_t NUM_COLS> struct big_matrix {
8-
private:
9-
T *mat;
10-
11-
public:
12-
T *get_data() { return mat; }
13-
void set_data(T *data) { mat = data; }
14-
big_matrix(T *data) : mat(data) {}
15-
};
16-
175
template <typename T1, typename T2, size_t M, size_t N, size_t K>
186
void matrix_multiply(big_matrix<T1, M, N> &C, big_matrix<T2, M, K> &A,
197
big_matrix<T2, K / 2, N * 2> &B) {
@@ -150,13 +138,7 @@ int main() {
150138
matrix_multiply_ref((int32_t *)Aref, (int32_t *)Bref, (int32_t *)D, MATRIX_M,
151139
MATRIX_N, MATRIX_K / 2);
152140

153-
bool res = true;
154-
for (int i = 0; i < MATRIX_M; i++) {
155-
for (int j = 0; j < MATRIX_N; j++) {
156-
if (fabs(C[i][j] - D[i][j]) > BF16_EPSILON)
157-
res = false;
158-
}
159-
}
141+
bool res = matrix_compare(MATRIX_M, MATRIX_N, (float *)C, (float *)D);
160142
std::cout << (res ? "passed" : "failed") << std::endl;
161143
return !res;
162144
}

‎sycl/test-e2e/Matrix/Legacy/joint_matrix_bfloat16_rowmajorA_rowmajorB.cpp

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -13,12 +13,10 @@
1313
// This tests support of row major layout for matrix B which does automatic VNNI
1414
// transform. This is currently only available on AMX and XMX of PVC
1515

16-
#include <iostream>
17-
#include <sycl/sycl.hpp>
16+
#include "../common.hpp"
1817

1918
using namespace sycl;
2019
using namespace sycl::ext::oneapi::experimental::matrix;
21-
using bfloat16 = sycl::ext::oneapi::bfloat16;
2220

2321
#define SG_SZ 16
2422

‎sycl/test-e2e/Matrix/Legacy/joint_matrix_bfloat16_rowmajorA_rowmajorB_impl.hpp

Lines changed: 1 addition & 24 deletions
Original file line numberDiff line numberDiff line change
@@ -3,16 +3,6 @@
33
#define TK 16
44
#define BF16_EPSILON 0.00781250
55

6-
template <typename T, size_t NUM_ROWS, size_t NUM_COLS> struct big_matrix {
7-
private:
8-
T *mat;
9-
10-
public:
11-
T *get_data() { return mat; }
12-
void set_data(T *data) { mat = data; }
13-
big_matrix(T *data) : mat(data) {}
14-
};
15-
166
template <typename T1, typename T2, size_t M, size_t N, size_t K>
177
void matrix_multiply(big_matrix<T1, M, N> &C, big_matrix<T2, M, K> &A,
188
big_matrix<T2, K, N> &B) {
@@ -81,13 +71,6 @@ bfloat16 B[MATRIX_K][MATRIX_N];
8171
float C[MATRIX_M][MATRIX_N];
8272
float D[MATRIX_M][MATRIX_N];
8373

84-
float make_fp32(bfloat16 x) {
85-
unsigned int y = *((int *)&x);
86-
y = y << 16;
87-
float *res = reinterpret_cast<float *>(&y);
88-
return *res;
89-
}
90-
9174
void matrix_multiply_ref(int M, int N, int K) {
9275
for (int m = 0; m < M; m++)
9376
for (int n = 0; n < N; n++) {
@@ -122,13 +105,7 @@ int main() {
122105
matrix_multiply(MC, MA, MB);
123106
matrix_multiply_ref(MATRIX_M, MATRIX_N, MATRIX_K);
124107

125-
bool res = true;
126-
for (int i = 0; i < MATRIX_M; i++) {
127-
for (int j = 0; j < MATRIX_N; j++) {
128-
if (fabs(C[i][j] - D[i][j]) > BF16_EPSILON)
129-
res = false;
130-
}
131-
}
108+
bool res = matrix_compare(MATRIX_M, MATRIX_N, (float *)C, (float *)D);
132109
std::cout << (res ? "passed" : "failed") << std::endl;
133110
return !res;
134111
}

‎sycl/test-e2e/Matrix/SG32/joint_matrix_all_sizes.cpp

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -10,12 +10,11 @@
1010
// RUN: %{build} -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4
1111
// RUN: %{run} %t.out
1212

13+
#include "../common.hpp"
1314
#include <iostream>
14-
#include <sycl/sycl.hpp>
1515

1616
using namespace sycl;
1717
using namespace sycl::ext::oneapi::experimental::matrix;
18-
using bfloat16 = sycl::ext::oneapi::bfloat16;
1918

2019
constexpr size_t SG_SZ = 32;
2120
// Sub-matrix N dimension

‎sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_fill_k_cache.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10,6 +10,7 @@
1010
// RUN: %{build} -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4
1111
// RUN: %{run} %t.out
1212

13+
#include "../common.hpp"
1314
#include <cstddef>
1415

1516
constexpr size_t SG_SZ = 32;

‎sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_fill_k_cache_init.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10,6 +10,7 @@
1010
// RUN: %{build} -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -DINIT_LIST
1111
// RUN: %{run} %t.out
1212

13+
#include "../common.hpp"
1314
#include <cstddef>
1415

1516
constexpr size_t SG_SZ = 32;

‎sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_fill_k_cache_unroll.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -16,6 +16,7 @@
1616
// -mllvm -inline-threshold added as a workaround,
1717
// since IGC doesn't support some variants of IR for Joint Matrix currently
1818

19+
#include "../common.hpp"
1920
#include <cstddef>
2021

2122
constexpr size_t SG_SZ = 32;

‎sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_fill_k_cache_unroll_init.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -13,6 +13,7 @@
1313
// -mllvm -inline-threshold added as a workaround,
1414
// since IGC doesn't support some variants of IR for Joint Matrix currently
1515

16+
#include "../common.hpp"
1617
#include <cstddef>
1718

1819
constexpr size_t SG_SZ = 32;

‎sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10,6 +10,7 @@
1010
// RUN: %{build} -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4
1111
// RUN: %{run} %t.out
1212

13+
#include "../common.hpp"
1314
#include <iostream>
1415
#include <sycl/sycl.hpp>
1516

‎sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16_32x64.cpp

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -12,12 +12,10 @@
1212

1313
// XFAIL: *
1414

15-
#include <iostream>
16-
#include <sycl/sycl.hpp>
15+
#include "../common.hpp"
1716

1817
using namespace sycl;
1918
using namespace sycl::ext::oneapi::experimental::matrix;
20-
using bfloat16 = sycl::ext::oneapi::bfloat16;
2119

2220
constexpr size_t SG_SZ = 32;
2321

‎sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16_array.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10,6 +10,7 @@
1010
// RUN: %{build} -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4
1111
// RUN: %{run} %t.out
1212

13+
#include "../common.hpp"
1314
#include <cstddef>
1415

1516
constexpr std::size_t SG_SZ = 32;

‎sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16_colmajorA_colmajorB.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -15,6 +15,7 @@
1515

1616
// XFAIL: gpu
1717

18+
#include "../common.hpp"
1819
#include <iostream>
1920
#include <sycl/sycl.hpp>
2021

‎sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16_rowmajorA_rowmajorB.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -15,6 +15,7 @@
1515

1616
// XFAIL: gpu
1717

18+
#include "../common.hpp"
1819
#include <iostream>
1920
#include <sycl/sycl.hpp>
2021

0 commit comments

Comments
 (0)