Skip to content
This repository was archived by the owner on Mar 28, 2023. It is now read-only.

Commit 3d696bb

Browse files
dkhaldidkhaldi
authored
Add PVC tests and move XMX8 tests to a new folder (#1347)
1 parent 99c83d3 commit 3d696bb

Some content is hidden

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

48 files changed

+3110
-2717
lines changed

‎SYCL/Matrix/XMX8/element_wise_all_ops_bf16.cpp

Lines changed: 24 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,24 @@
1+
//==----------- element_wise_all_ops_bf16.cpp - DPC++ joint_matrix---------==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: matrix-xmx8
9+
10+
// RUN: %clangxx -fsycl %s -o %t.out
11+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
12+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
13+
14+
#include <iostream>
15+
#include <random>
16+
#include <sycl/sycl.hpp>
17+
18+
using namespace sycl;
19+
using namespace sycl::ext::intel;
20+
using namespace sycl::ext::oneapi::experimental::matrix;
21+
22+
#define SG_SZ 8
23+
24+
#include "../element_wise_all_ops_bf16_impl.hpp"

‎SYCL/Matrix/XMX8/element_wise_all_ops_half.cpp

Lines changed: 25 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,25 @@
1+
//==----------- element_wise_all_ops_half.cpp - DPC++ joint_matrix---------==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: matrix-xmx8
9+
10+
// Only runs on DPAS because AMX implementation does not support half data type
11+
// yet
12+
// RUN: %clangxx -fsycl %s -o %t.out
13+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
14+
15+
#include <iostream>
16+
#include <random>
17+
#include <sycl/sycl.hpp>
18+
19+
using namespace sycl;
20+
using namespace sycl::ext::intel;
21+
using namespace sycl::ext::oneapi::experimental::matrix;
22+
23+
#define SG_SZ 8
24+
25+
#include "../element_wise_all_ops_half_impl.hpp"

‎SYCL/Matrix/XMX8/element_wise_all_ops_int8.cpp

Lines changed: 24 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,24 @@
1+
//==----------- element_wise_all_ops_int8.cpp - DPC++ joint_matrix---------==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: matrix-xmx8
9+
10+
// RUN: %clangxx -fsycl %s -o %t.out
11+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
12+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
13+
14+
#include <iostream>
15+
#include <random>
16+
#include <sycl/sycl.hpp>
17+
18+
using namespace sycl;
19+
using namespace sycl::ext::intel;
20+
using namespace sycl::ext::oneapi::experimental::matrix;
21+
22+
#define SG_SZ 8
23+
24+
#include "../element_wise_all_ops_int8_impl.hpp"

‎SYCL/Matrix/XMX8/element_wise_all_ops_int8_packed.cpp

Lines changed: 24 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,24 @@
1+
//==------ element_wise_all_ops_int8_packed.cpp - DPC++ joint_matrix-------==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: matrix-xmx8
9+
10+
// RUN: %clangxx -fsycl %s -o %t.out
11+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
12+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
13+
14+
#include <iostream>
15+
#include <random>
16+
#include <sycl/sycl.hpp>
17+
18+
using namespace sycl;
19+
using namespace sycl::ext::intel;
20+
using namespace sycl::ext::oneapi::experimental::matrix;
21+
22+
#define SG_SZ 8
23+
24+
#include "../element_wise_all_ops_int8_packed_impl.hpp"

‎SYCL/Matrix/XMX8/element_wise_irreg_sum_rows.cpp

Lines changed: 26 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,26 @@
1+
//==-------- element_wise_irreg_sum_rows.cpp - DPC++ joint_matrix----- ----==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: matrix-xmx8
9+
10+
// RUN: %clangxx -fsycl %s -o %t.out
11+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
12+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
13+
14+
// this code calculates the sum of rows into a global array of number of rows
15+
// elements. First, partial reduction is computed inside each SG, then atomic
16+
// add is used to reduce between SG leaders
17+
18+
#include <iostream>
19+
#include <sycl/sycl.hpp>
20+
21+
using namespace sycl;
22+
using namespace sycl::ext::oneapi::experimental::matrix;
23+
24+
#define SG_SZ 8
25+
26+
#include "../element_wise_irreg_sum_rows_impl.hpp"

‎SYCL/Matrix/XMX8/element_wise_ops.cpp

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
//==----------- element_wise_ops.cpp - DPC++ joint_matrix------------- ----==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: matrix-xmx8
9+
10+
// RUN: %clangxx -fsycl %s -o %t.out
11+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
12+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
13+
14+
#include <iostream>
15+
#include <sycl/sycl.hpp>
16+
17+
using namespace sycl;
18+
using namespace sycl::ext::oneapi::experimental::matrix;
19+
20+
#define SG_SZ 8
21+
22+
#include "../element_wise_ops_impl.hpp"

‎SYCL/Matrix/XMX8/joint_matrix_bf16.cpp

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
//==-------- joint_matrix_bf16.cpp - DPC++ joint_matrix--------------- ----==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: matrix-xmx8
9+
10+
// RUN: %clangxx -fsycl %s -o %t.out
11+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
12+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
13+
14+
#include <iostream>
15+
#include <sycl/sycl.hpp>
16+
17+
using namespace sycl;
18+
using namespace sycl::ext::oneapi::experimental::matrix;
19+
20+
#define SG_SZ 8
21+
22+
#include "../joint_matrix_bf16_impl.hpp"

‎SYCL/Matrix/XMX8/joint_matrix_bfloat16.cpp

Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,23 @@
1+
//==-------- joint_matrix_bfloat16.cpp - DPC++ joint_matrix----------- ----==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: matrix-xmx8
9+
10+
// RUN: %clangxx -fsycl %s -o %t.out
11+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
12+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
13+
14+
#include <iostream>
15+
#include <sycl/sycl.hpp>
16+
17+
using namespace sycl;
18+
using namespace sycl::ext::oneapi::experimental::matrix;
19+
using bfloat16 = sycl::ext::oneapi::experimental::bfloat16;
20+
21+
#define SG_SZ 8
22+
23+
#include "../joint_matrix_bfloat16_impl.hpp"

‎SYCL/Matrix/XMX8/joint_matrix_bfloat16_use.cpp

Lines changed: 24 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,24 @@
1+
//==-------- joint_matrix_bfloat16_use.cpp - DPC++ joint_matrix------------==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: matrix-xmx8
9+
10+
// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=2
11+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
12+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
13+
14+
// XFAIL: gpu
15+
16+
#include <iostream>
17+
#include <sycl/sycl.hpp>
18+
19+
using namespace sycl::ext::oneapi::experimental::matrix;
20+
using bfloat16 = sycl::ext::oneapi::experimental::bfloat16;
21+
22+
#define SG_SZ 8
23+
24+
#include "../joint_matrix_bfloat16_use_impl.hpp"

‎SYCL/Matrix/XMX8/joint_matrix_half.cpp

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
//==-------- joint_matrix_half.cpp - DPC++ joint_matrix------------ ----==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: matrix-xmx8
9+
10+
// RUN: %clangxx -fsycl %s -o %t.out
11+
// Only run on the GPU because half is not supported on AMX hardware
12+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
13+
14+
#include <iostream>
15+
#include <sycl/sycl.hpp>
16+
17+
using namespace sycl;
18+
using namespace sycl::ext::oneapi::experimental::matrix;
19+
20+
#define SG_SZ 8
21+
22+
#include "../joint_matrix_half_impl.hpp"

‎SYCL/Matrix/XMX8/joint_matrix_int8_vnni.cpp

Lines changed: 24 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,24 @@
1+
//==-------- joint_matrix_bf16_vnni.cpp - DPC++ joint_matrix---------------==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: matrix-xmx8
9+
10+
// RUN: %clangxx -fsycl %s -o %t.out
11+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
12+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
13+
14+
// XFAIL: *
15+
16+
#include <iostream>
17+
#include <sycl/sycl.hpp>
18+
19+
using namespace sycl;
20+
using namespace sycl::ext::oneapi::experimental::matrix;
21+
22+
#define SG_SZ 8
23+
24+
#include "../joint_matrix_int8_vnni_impl.hpp"

‎SYCL/Matrix/XMX8/joint_matrix_ss_int8.cpp

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
//==-------- joint_matrix_ss_int8.cpp - DPC++ joint_matrix------------ ----==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: matrix-xmx8
9+
10+
// RUN: %clangxx -fsycl %s -o %t.out
11+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
12+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
13+
14+
#include <iostream>
15+
#include <sycl/sycl.hpp>
16+
17+
using namespace sycl;
18+
using namespace sycl::ext::oneapi::experimental::matrix;
19+
20+
#define SG_SZ 8
21+
22+
#include "../joint_matrix_ss_int8_impl.hpp"

‎SYCL/Matrix/XMX8/joint_matrix_ss_int8_use.cpp

Lines changed: 24 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,24 @@
1+
//==-------- joint_matrix_ss_int8_use.cpp - DPC++ joint_matrix-------------==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: matrix-xmx8
9+
10+
// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX=2
11+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
12+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
13+
14+
// XFAIL: gpu
15+
16+
#include <iostream>
17+
#include <sycl/sycl.hpp>
18+
19+
using namespace sycl;
20+
using namespace sycl::ext::oneapi::experimental::matrix;
21+
22+
#define SG_SZ 8
23+
24+
#include "../joint_matrix_ss_int8_use_impl.hpp"

‎SYCL/Matrix/XMX8/joint_matrix_su_int8.cpp

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
//==-------- joint_matrix_su_int8.cpp - DPC++ joint_matrix------------ ----==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: matrix-xmx8
9+
10+
// RUN: %clangxx -fsycl %s -o %t.out
11+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
12+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
13+
14+
#include <iostream>
15+
#include <sycl/sycl.hpp>
16+
17+
using namespace sycl;
18+
using namespace sycl::ext::oneapi::experimental::matrix;
19+
20+
#define SG_SZ 8
21+
22+
#include "../joint_matrix_su_int8_impl.hpp"

‎SYCL/Matrix/XMX8/joint_matrix_us_int8.cpp

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
//==-------- joint_matrix_us_int8.cpp - DPC++ joint_matrix------------ ----==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: matrix-xmx8
9+
10+
// RUN: %clangxx -fsycl %s -o %t.out
11+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
12+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
13+
14+
#include <iostream>
15+
#include <sycl/sycl.hpp>
16+
17+
using namespace sycl;
18+
using namespace sycl::ext::oneapi::experimental::matrix;
19+
20+
#define SG_SZ 8
21+
22+
#include "../joint_matrix_us_int8_impl.hpp"

‎SYCL/Matrix/XMX8/joint_matrix_uu_int8.cpp

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
//==-------- joint_matrix_uu_int8.cpp - DPC++ joint_matrix------------ ----==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: matrix-xmx8
9+
10+
// RUN: %clangxx -fsycl %s -o %t.out
11+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
12+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
13+
14+
#include <iostream>
15+
#include <sycl/sycl.hpp>
16+
17+
using namespace sycl;
18+
using namespace sycl::ext::oneapi::experimental::matrix;
19+
20+
#define SG_SZ 8
21+
22+
#include "../joint_matrix_uu_int8_impl.hpp"

0 commit comments

Comments
 (0)