Skip to content

Commit fad9855

Browse files
mylermyler
committed
resolve conflicts
1 parent bff4365 commit fad9855

25 files changed

+1
-531
lines changed

‎CMakeLists.txt

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -6,16 +6,13 @@ option(TEST_SUITE_FORTRAN "Enable Fortran test suite" OFF)
66

77
project(test-suite C CXX)
88

9-
<<<<<<< HEAD
10-
=======
119
if("Fortran" IN_LIST TEST_SUITE_SUBDIRS)
1210
set(TEST_SUITE_FORTRAN_default ON)
1311
else()
1412
set(TEST_SUITE_FORTRAN_default OFF)
1513
endif()
1614
option(TEST_SUITE_FORTRAN "Enable Fortran test suite" ${TEST_SUITE_FORTRAN_default})
1715

18-
>>>>>>> temp_main
1916
if(TEST_SUITE_FORTRAN)
2017
enable_language(Fortran)
2118
endif()

‎SYCL/InlineAsm/asm_multiple_instructions.cpp

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -58,12 +58,8 @@ struct KernelFunctor : WithInputBuffers<T, 3>, WithOutputBuffer<T> {
5858
D[wiID] = output;
5959
#else
6060
#if defined(__SYCL_DEVICE_ONLY__)
61-
<<<<<<< HEAD
62-
asm volatile("{\n"
63-
=======
6461
asm volatile(
6562
"{\n"
66-
>>>>>>> 417b60f35 ([SYCL] Use Subgroup size 16 as default for InlineASM tests (#1476))
6763
"add (M1, 16) %1(0, 0)<1> %1(0, 0)<1;1,0> %2(0, 0)<1;1,0>\n"
6864
"add (M1, 16) %1(0, 0)<1> %1(0, 0)<1;1,0> %3(0, 0)<1;1,0>\n"
6965
"mov (M1, 16) %0(0, 0)<1> %1(0, 0)<1;1,0>\n"

‎SYCL/InlineAsm/asm_no_operands.cpp

Lines changed: 0 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -25,15 +25,9 @@ int main() {
2525
// Submitting command group(work) to queue
2626
Queue.submit([&](sycl::handler &cgh) {
2727
// Executing kernel
28-
<<<<<<< HEAD
29-
cgh.parallel_for<no_operands_kernel>(NumOfWorkItems,
30-
[=](sycl::id<1> WIid)
31-
[[intel::reqd_sub_group_size(16)]] {
32-
=======
3328
cgh.parallel_for<no_operands_kernel>(
3429
NumOfWorkItems,
3530
[=](sycl::id<1> WIid) [[intel::reqd_sub_group_size(16)]] {
36-
>>>>>>> 417b60f35 ([SYCL] Use Subgroup size 16 as default for InlineASM tests (#1476))
3731
#if defined(__SYCL_DEVICE_ONLY__)
3832
asm("barrier");
3933
#endif

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

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -7,11 +7,7 @@
77
//===----------------------------------------------------------------------===//
88
// REQUIRES: matrix-xmx8
99

10-
<<<<<<<< HEAD:SYCL/Matrix/Legacy/XMX8/joint_matrix_bf16.cpp
1110
// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=1
12-
========
13-
// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4
14-
>>>>>>>> cbbfcc6c1 ([SYCL] Add matrix tests that use the new API (unified API) (#1391)):SYCL/Matrix/XMX8/joint_matrix_bf16.cpp
1511
// RUN: %CPU_RUN_PLACEHOLDER %t.out
1612
// RUN: %GPU_RUN_PLACEHOLDER %t.out
1713

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

Lines changed: 0 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -7,15 +7,7 @@
77
//===----------------------------------------------------------------------===//
88
// REQUIRES: matrix-xmx8
99

10-
<<<<<<< HEAD
11-
<<<<<<<< HEAD:SYCL/Matrix/Legacy/XMX8/joint_matrix_bfloat16.cpp
1210
// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=1
13-
========
14-
// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4
15-
>>>>>>>> cbbfcc6c1 ([SYCL] Add matrix tests that use the new API (unified API) (#1391)):SYCL/Matrix/XMX8/joint_matrix_bfloat16_use.cpp
16-
=======
17-
// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=1
18-
>>>>>>> cbbfcc6c1 ([SYCL] Add matrix tests that use the new API (unified API) (#1391))
1911
// RUN: %CPU_RUN_PLACEHOLDER %t.out
2012
// RUN: %GPU_RUN_PLACEHOLDER %t.out
2113

@@ -24,15 +16,7 @@
2416

2517
using namespace sycl;
2618
using namespace sycl::ext::oneapi::experimental::matrix;
27-
<<<<<<< HEAD
28-
<<<<<<< HEAD
29-
using bfloat16 = sycl::ext::oneapi::bfloat16;
30-
=======
31-
using bfloat16 = sycl::ext::oneapi::experimental::bfloat16;
32-
>>>>>>> cbbfcc6c1 ([SYCL] Add matrix tests that use the new API (unified API) (#1391))
33-
=======
3419
using bfloat16 = sycl::ext::oneapi::bfloat16;
35-
>>>>>>> 87f7445c7 ([SYCL][Matrix] Fix bfloat16 namespace in the legacy tests (#1478))
3620

3721
#define SG_SZ 8
3822

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

Lines changed: 0 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -18,15 +18,7 @@
1818

1919
using namespace sycl;
2020
using namespace sycl::ext::oneapi::experimental::matrix;
21-
<<<<<<< HEAD
22-
<<<<<<< HEAD
2321
using bfloat16 = sycl::ext::oneapi::bfloat16;
24-
=======
25-
using bfloat16 = sycl::ext::oneapi::experimental::bfloat16;
26-
>>>>>>> cbbfcc6c1 ([SYCL] Add matrix tests that use the new API (unified API) (#1391))
27-
=======
28-
using bfloat16 = sycl::ext::oneapi::bfloat16;
29-
>>>>>>> 87f7445c7 ([SYCL][Matrix] Fix bfloat16 namespace in the legacy tests (#1478))
3022

3123
#define SG_SZ 8
3224

‎SYCL/Matrix/Legacy/element_wise_all_ops_bf16.cpp

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -11,10 +11,6 @@
1111
// RUN: %CPU_RUN_PLACEHOLDER %t.out
1212
// RUN: %GPU_RUN_PLACEHOLDER %t.out
1313

14-
<<<<<<< HEAD
15-
16-
=======
17-
>>>>>>> cbbfcc6c1 ([SYCL] Add matrix tests that use the new API (unified API) (#1391))
1814
#include <iostream>
1915
#include <random>
2016
#include <sycl/sycl.hpp>

‎SYCL/Matrix/Legacy/element_wise_all_ops_half_impl.hpp

Lines changed: 0 additions & 50 deletions
Original file line numberDiff line numberDiff line change
@@ -41,15 +41,7 @@ void matrix_verify_add(queue q, big_matrix<T, M, N> &A, nd_range<2> &r,
4141
ext::oneapi::sub_group sg = spmd_item.get_sub_group();
4242
joint_matrix<T, TM, TK> sub_a(sg);
4343

44-
<<<<<<< HEAD
45-
<<<<<<< HEAD
4644
joint_matrix_fill(sg, sub_a, 5);
47-
=======
48-
joint_matrix_fill(sg, sub_a, 5.0);
49-
>>>>>>> cbbfcc6c1 ([SYCL] Add matrix tests that use the new API (unified API) (#1391))
50-
=======
51-
joint_matrix_fill(sg, sub_a, 5);
52-
>>>>>>> e3a7db8e9 ([SYCL][Matrix] fix ATS-M double bug (#1496))
5345

5446
auto wi_slice_a = sub_a.get_wi_data();
5547
for (int i = 0; i < wi_slice_a.length(); i++) {
@@ -82,15 +74,7 @@ void matrix_verify_sub(queue q, big_matrix<T, M, N> &A, nd_range<2> &r,
8274
ext::oneapi::sub_group sg = spmd_item.get_sub_group();
8375
joint_matrix<T, TM, TK> sub_a(sg);
8476

85-
<<<<<<< HEAD
86-
<<<<<<< HEAD
87-
joint_matrix_fill(sg, sub_a, 5);
88-
=======
89-
joint_matrix_fill(sg, sub_a, 5.0);
90-
>>>>>>> cbbfcc6c1 ([SYCL] Add matrix tests that use the new API (unified API) (#1391))
91-
=======
9277
joint_matrix_fill(sg, sub_a, 5);
93-
>>>>>>> e3a7db8e9 ([SYCL][Matrix] fix ATS-M double bug (#1496))
9478

9579
auto wi_slice_a = sub_a.get_wi_data();
9680
for (int i = 0; i < wi_slice_a.length(); i++) {
@@ -123,15 +107,7 @@ void matrix_verify_mul(queue q, big_matrix<T, M, N> &A, nd_range<2> &r,
123107
ext::oneapi::sub_group sg = spmd_item.get_sub_group();
124108
joint_matrix<T, TM, TK> sub_a(sg);
125109

126-
<<<<<<< HEAD
127-
<<<<<<< HEAD
128110
joint_matrix_fill(sg, sub_a, 5);
129-
=======
130-
joint_matrix_fill(sg, sub_a, 5.0);
131-
>>>>>>> cbbfcc6c1 ([SYCL] Add matrix tests that use the new API (unified API) (#1391))
132-
=======
133-
joint_matrix_fill(sg, sub_a, 5);
134-
>>>>>>> e3a7db8e9 ([SYCL][Matrix] fix ATS-M double bug (#1496))
135111

136112
auto wi_slice_a = sub_a.get_wi_data();
137113
for (int i = 0; i < wi_slice_a.length(); i++) {
@@ -164,15 +140,7 @@ void matrix_verify_div(queue q, big_matrix<T, M, N> &A, nd_range<2> &r,
164140
ext::oneapi::sub_group sg = spmd_item.get_sub_group();
165141
joint_matrix<T, TM, TK> sub_a(sg);
166142

167-
<<<<<<< HEAD
168-
<<<<<<< HEAD
169-
joint_matrix_fill(sg, sub_a, 4);
170-
=======
171-
joint_matrix_fill(sg, sub_a, 4.0);
172-
>>>>>>> cbbfcc6c1 ([SYCL] Add matrix tests that use the new API (unified API) (#1391))
173-
=======
174143
joint_matrix_fill(sg, sub_a, 4);
175-
>>>>>>> e3a7db8e9 ([SYCL][Matrix] fix ATS-M double bug (#1496))
176144

177145
auto wi_slice_a = sub_a.get_wi_data();
178146
for (int i = 0; i < wi_slice_a.length(); i++) {
@@ -205,15 +173,7 @@ void matrix_verify_logic(queue q, big_matrix<T, M, N> &A, nd_range<2> &r,
205173
ext::oneapi::sub_group sg = spmd_item.get_sub_group();
206174
joint_matrix<T, TM, TK> sub_a(sg);
207175

208-
<<<<<<< HEAD
209-
<<<<<<< HEAD
210176
joint_matrix_fill(sg, sub_a, 5);
211-
=======
212-
joint_matrix_fill(sg, sub_a, 5.0);
213-
>>>>>>> cbbfcc6c1 ([SYCL] Add matrix tests that use the new API (unified API) (#1391))
214-
=======
215-
joint_matrix_fill(sg, sub_a, 5);
216-
>>>>>>> e3a7db8e9 ([SYCL][Matrix] fix ATS-M double bug (#1496))
217177

218178
auto wi_slice_a = sub_a.get_wi_data();
219179
for (int i = 0; i < wi_slice_a.length(); i++) {
@@ -229,18 +189,8 @@ void matrix_verify_logic(queue q, big_matrix<T, M, N> &A, nd_range<2> &r,
229189
val++;
230190
if (wi_slice_a[i] == static_cast<half>(2.0)) {
231191
val -= 2;
232-
<<<<<<< HEAD
233-
<<<<<<< HEAD
234-
val *= 3;
235-
val /= 2;
236-
=======
237-
val *= 3.0;
238-
val /= 2.0;
239-
>>>>>>> cbbfcc6c1 ([SYCL] Add matrix tests that use the new API (unified API) (#1391))
240-
=======
241192
val *= 3;
242193
val /= 2;
243-
>>>>>>> e3a7db8e9 ([SYCL][Matrix] fix ATS-M double bug (#1496))
244194
} else {
245195
val += 2;
246196
}

‎SYCL/Matrix/Legacy/joint_matrix_bf16.cpp

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -7,11 +7,7 @@
77
//===----------------------------------------------------------------------===//
88
// REQUIRES: matrix
99

10-
<<<<<<<< HEAD:SYCL/Matrix/Legacy/joint_matrix_bf16.cpp
1110
// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=1
12-
========
13-
// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4
14-
>>>>>>>> cbbfcc6c1 ([SYCL] Add matrix tests that use the new API (unified API) (#1391)):SYCL/Matrix/joint_matrix_bf16.cpp
1511
// RUN: %CPU_RUN_PLACEHOLDER %t.out
1612
// RUN: %GPU_RUN_PLACEHOLDER %t.out
1713

‎SYCL/Matrix/Legacy/joint_matrix_bf16_impl.hpp

Lines changed: 0 additions & 37 deletions
Original file line numberDiff line numberDiff line change
@@ -52,27 +52,6 @@ void matrix_multiply(big_matrix<T1, NUM_ROWS_C, NUM_COLS_C> &C,
5252
const auto sg_starty = global_idy - spmd_item.get_local_id(1);
5353

5454
sub_group sg = spmd_item.get_sub_group();
55-
<<<<<<< HEAD
56-
joint_matrix<sub_group, unsigned short, use::a, TM, TK,
57-
layout::row_major>
58-
sub_a;
59-
// For B, we assume B has been already VNNIed.
60-
joint_matrix<sub_group, unsigned short, use::b, TK, TN,
61-
ext::intel::experimental::matrix::layout::packed>
62-
sub_b;
63-
joint_matrix<sub_group, float, use::accumulator, TM, TN> sub_c;
64-
joint_matrix_load(sg, sub_c,
65-
accC.get_pointer() + (sg_startx * TM) * N +
66-
sg_starty / SG_SZ * TN,
67-
N, layout::row_major);
68-
for (int k = 0; k < K; k += TK) {
69-
joint_matrix_load(
70-
sg, sub_a, accA.get_pointer() + (sg_startx * TM) * K + k, K);
71-
joint_matrix_load(sg, sub_b,
72-
accB.get_pointer() + k * N +
73-
sg_starty / SG_SZ * TN * 2,
74-
N * 2);
75-
=======
7655
joint_matrix<unsigned short, TM, TK> sub_a(sg);
7756
// For B, since current implementation does not support non-packed
7857
// layout, users need to specify the packed_b layout.
@@ -93,17 +72,12 @@ void matrix_multiply(big_matrix<T1, NUM_ROWS_C, NUM_COLS_C> &C,
9372
accB.get_pointer() + (k) * (N) +
9473
sg_starty / SG_SZ * TN * 2,
9574
N * 2, matrix_layout::packed_b);
96-
>>>>>>> cbbfcc6c1 ([SYCL] Add matrix tests that use the new API (unified API) (#1391))
9775
sub_c = joint_matrix_mad(sg, sub_a, sub_b, sub_c);
9876
}
9977
joint_matrix_store(sg, sub_c,
10078
accC.get_pointer() + (sg_startx * TM) * N +
10179
sg_starty / SG_SZ * TN,
102-
<<<<<<< HEAD
103-
N, layout::row_major);
104-
=======
10580
N, matrix_layout::row_major);
106-
>>>>>>> cbbfcc6c1 ([SYCL] Add matrix tests that use the new API (unified API) (#1391))
10781
}); // parallel for
10882
}).wait();
10983
}
@@ -131,20 +105,14 @@ unsigned short make_bf16(float x) {
131105

132106
void matrix_multiply_ref(int *A_mem, int *B_mem, int *C_mem, int M, int N,
133107
int K) {
134-
<<<<<<< HEAD
135-
=======
136108
// tiling
137-
>>>>>>> cbbfcc6c1 ([SYCL] Add matrix tests that use the new API (unified API) (#1391))
138109
for (int m = 0; m < M; m++)
139110
for (int n = 0; n < N; n++) {
140111
for (int k = 0; k < K; k++) {
141112
short *va = (short *)(A_mem + m * K + k);
142113
short *vb = (short *)(B_mem + k * N + n);
143114
float acc = *((float *)(C_mem + m * N + n));
144-
<<<<<<< HEAD
145-
=======
146115
// FIXME: Should we do reduce-add in another version?
147-
>>>>>>> cbbfcc6c1 ([SYCL] Add matrix tests that use the new API (unified API) (#1391))
148116
for (int i = 0; i < 2; i++) {
149117
acc += (make_fp32(va[i]) * make_fp32(vb[i]));
150118
}
@@ -187,13 +155,8 @@ int main() {
187155
res = false;
188156
}
189157
}
190-
<<<<<<< HEAD
191-
std::cout << (res ? "passed" : "failed") << std::endl;
192-
return !res;
193-
=======
194158
if (res)
195159
std::cout << "passed\n";
196160
else
197161
std::cout << "failed\n";
198-
>>>>>>> cbbfcc6c1 ([SYCL] Add matrix tests that use the new API (unified API) (#1391))
199162
}

‎SYCL/Matrix/Legacy/joint_matrix_bfloat16.cpp

Lines changed: 0 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -16,15 +16,7 @@
1616

1717
using namespace sycl;
1818
using namespace sycl::ext::oneapi::experimental::matrix;
19-
<<<<<<< HEAD
20-
<<<<<<< HEAD
2119
using bfloat16 = sycl::ext::oneapi::bfloat16;
22-
=======
23-
using bfloat16 = sycl::ext::oneapi::experimental::bfloat16;
24-
>>>>>>> cbbfcc6c1 ([SYCL] Add matrix tests that use the new API (unified API) (#1391))
25-
=======
26-
using bfloat16 = sycl::ext::oneapi::bfloat16;
27-
>>>>>>> 87f7445c7 ([SYCL][Matrix] Fix bfloat16 namespace in the legacy tests (#1478))
2820

2921
#define SG_SZ 16
3022

‎SYCL/Matrix/Legacy/joint_matrix_bfloat16_32x64.cpp

Lines changed: 0 additions & 23 deletions
Original file line numberDiff line numberDiff line change
@@ -1,12 +1,4 @@
1-
<<<<<<< HEAD
2-
<<<<<<<< HEAD:SYCL/Matrix/Legacy/joint_matrix_bfloat16_32x64.cpp
31
//==----- joint_matrix_bfloat16_32x64.cpp - DPC++ joint_matrix-------------==//
4-
========
5-
//==-------- joint_matrix_bf16_vnni.cpp - DPC++ joint_matrix---------------==//
6-
>>>>>>>> cbbfcc6c1 ([SYCL] Add matrix tests that use the new API (unified API) (#1391)):SYCL/Matrix/Legacy/joint_matrix_int8_vnni.cpp
7-
=======
8-
//==----- joint_matrix_bfloat16_32x64.cpp - DPC++ joint_matrix-------------==//
9-
>>>>>>> cbbfcc6c1 ([SYCL] Add matrix tests that use the new API (unified API) (#1391))
102
//
113
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
124
// See https://llvm.org/LICENSE.txt for license information.
@@ -26,23 +18,8 @@
2618

2719
using namespace sycl;
2820
using namespace sycl::ext::oneapi::experimental::matrix;
29-
<<<<<<< HEAD
30-
<<<<<<< HEAD
31-
32-
#define SG_SZ 16
33-
34-
<<<<<<<< HEAD:SYCL/Matrix/Legacy/joint_matrix_bfloat16_32x64.cpp
35-
#include "joint_matrix_bfloat16_32x64_impl.hpp"
36-
========
37-
#include "joint_matrix_int8_vnni_impl.hpp"
38-
>>>>>>>> cbbfcc6c1 ([SYCL] Add matrix tests that use the new API (unified API) (#1391)):SYCL/Matrix/Legacy/joint_matrix_int8_vnni.cpp
39-
=======
40-
using bfloat16 = sycl::ext::oneapi::experimental::bfloat16;
41-
=======
4221
using bfloat16 = sycl::ext::oneapi::bfloat16;
43-
>>>>>>> 87f7445c7 ([SYCL][Matrix] Fix bfloat16 namespace in the legacy tests (#1478))
4422

4523
#define SG_SZ 16
4624

4725
#include "joint_matrix_bfloat16_32x64_impl.hpp"
48-
>>>>>>> cbbfcc6c1 ([SYCL] Add matrix tests that use the new API (unified API) (#1391))

‎SYCL/Matrix/Legacy/joint_matrix_bfloat16_32x64_impl.hpp

Lines changed: 0 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -120,29 +120,13 @@ int main() {
120120
for (int j = 0; j < MATRIX_K; j++) {
121121
// bfloat16 is created using unsigned short since conversion from float to
122122
// bfloat16 is not supported on the host side yet
123-
<<<<<<< HEAD
124-
<<<<<<< HEAD
125123
A[i][j] = bfloat16(1.0f * (i + j));
126-
=======
127-
A[i][j] = bfloat16::from_bits(make_bf16(1.0f * (i + j)));
128-
>>>>>>> cbbfcc6c1 ([SYCL] Add matrix tests that use the new API (unified API) (#1391))
129-
=======
130-
A[i][j] = bfloat16(1.0f * (i + j));
131-
>>>>>>> 2722bd134 ([SYCL][Matrix]update recent tests to use the new API and remove deprecated bfloat16::from_bits (#1494))
132124
Aref[i][j] = make_bf16(1.0f * (i + j));
133125
}
134126
}
135127
for (int i = 0; i < MATRIX_K / 2; i++) {
136128
for (int j = 0; j < MATRIX_N * 2; j++) {
137-
<<<<<<< HEAD
138-
<<<<<<< HEAD
139-
B[i][j] = bfloat16(2.0f * i + 3.0f * j);
140-
=======
141-
B[i][j] = bfloat16::from_bits((make_bf16(2.0f * i + 3.0f * j)));
142-
>>>>>>> cbbfcc6c1 ([SYCL] Add matrix tests that use the new API (unified API) (#1391))
143-
=======
144129
B[i][j] = bfloat16(2.0f * i + 3.0f * j);
145-
>>>>>>> 2722bd134 ([SYCL][Matrix]update recent tests to use the new API and remove deprecated bfloat16::from_bits (#1494))
146130
Bref[i][j] = make_bf16(2.0f * i + 3.0f * j);
147131
}
148132
}

0 commit comments

Comments
 (0)