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

[SYCL][Matrix]Add test for all sizes #1628

Draft
wants to merge 10 commits into
base: intel
Choose a base branch
from
Draft

[SYCL][Matrix]Add test for all sizes #1628

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
10 commits
Select commit Hold shift + click to select a range
ba6c267
[SYCL][Matrix]Add test for odd sizes
dkhaldi Jan 18, 2023
399939f
set to xfail on GPU
dkhaldi Jan 18, 2023
c5e5e6f
Uncomment the rest of the calls
dkhaldi Jan 24, 2023
642d1c4
correct the verification code
dkhaldi Feb 28, 2023
b581424
fixed M size of the big matrix to be multiplier of M size of joint ma…
YuriPlyakhin Mar 2, 2023
a91186f
Merge branch 'intel:intel' into all_sizes_tst
YuriPlyakhin Mar 4, 2023
c6d6053
[SYCL][Matrix][NFC] updated formatting
YuriPlyakhin Mar 9, 2023
302efe4
[SYCL][Matrix] Updated tests to pass on GPU
YuriPlyakhin Mar 10, 2023
c0bcf7d
Merge branch 'intel:intel' into all_sizes_tst
YuriPlyakhin Mar 11, 2023
163ebf9
[SYCL][Matrix] Fixed size N - it should be different depending on arc…
YuriPlyakhin Mar 16, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
62 changes: 62 additions & 0 deletions SYCL/Matrix/XMX8/joint_matrix_all_sizes_bfloat16.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,62 @@
//==-------- joint_matrix_all_sizes.cpp - DPC++ joint_matrix---------------==//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
// REQUIRES: matrix-xmx8

// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -DJOINT_MATRIX_TEST_ALL_BF16
// RUN: %CPU_RUN_PLACEHOLDER %t.out

// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -DJOINT_MATRIX_TEST_BF16_1
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=0 %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=1 %GPU_RUN_PLACEHOLDER %t.out

// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -DJOINT_MATRIX_TEST_BF16_2
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=0 %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=1 %GPU_RUN_PLACEHOLDER %t.out

// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -DJOINT_MATRIX_TEST_BF16_3
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=0 %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=1 %GPU_RUN_PLACEHOLDER %t.out

// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -DJOINT_MATRIX_TEST_BF16_4
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=0 %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=1 %GPU_RUN_PLACEHOLDER %t.out

// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -DJOINT_MATRIX_TEST_BF16_5
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=0 %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=1 %GPU_RUN_PLACEHOLDER %t.out

// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -DJOINT_MATRIX_TEST_BF16_6
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=0 %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=1 %GPU_RUN_PLACEHOLDER %t.out

// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -DJOINT_MATRIX_TEST_BF16_7
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=0 %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=1 %GPU_RUN_PLACEHOLDER %t.out

// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -DJOINT_MATRIX_TEST_BF16_8
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=0 %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=1 %GPU_RUN_PLACEHOLDER %t.out

#include <iostream>
#include <sycl/sycl.hpp>

using namespace sycl;
using namespace sycl::ext::oneapi::experimental::matrix;
using bfloat16 = sycl::ext::oneapi::bfloat16;

#define SG_SZ 8

#include "../joint_matrix_all_sizes_impl.hpp"
62 changes: 62 additions & 0 deletions SYCL/Matrix/XMX8/joint_matrix_all_sizes_int8.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,62 @@
//==-------- joint_matrix_all_sizes.cpp - DPC++ joint_matrix---------------==//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
// REQUIRES: matrix-xmx8

// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -DJOINT_MATRIX_TEST_ALL_INT8
// RUN: %CPU_RUN_PLACEHOLDER %t.out

// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -DJOINT_MATRIX_TEST_INT8_1
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=0 %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=1 %GPU_RUN_PLACEHOLDER %t.out

// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -DJOINT_MATRIX_TEST_INT8_2
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=0 %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=1 %GPU_RUN_PLACEHOLDER %t.out

// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -DJOINT_MATRIX_TEST_INT8_3
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=0 %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=1 %GPU_RUN_PLACEHOLDER %t.out

// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -DJOINT_MATRIX_TEST_INT8_4
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=0 %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=1 %GPU_RUN_PLACEHOLDER %t.out

// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -DJOINT_MATRIX_TEST_INT8_5
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=0 %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=1 %GPU_RUN_PLACEHOLDER %t.out

// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -DJOINT_MATRIX_TEST_INT8_6
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=0 %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=1 %GPU_RUN_PLACEHOLDER %t.out

// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -DJOINT_MATRIX_TEST_INT8_7
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=0 %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=1 %GPU_RUN_PLACEHOLDER %t.out

// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -DJOINT_MATRIX_TEST_INT8_8
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=0 %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=1 %GPU_RUN_PLACEHOLDER %t.out

#include <iostream>
#include <sycl/sycl.hpp>

using namespace sycl;
using namespace sycl::ext::oneapi::experimental::matrix;
using bfloat16 = sycl::ext::oneapi::bfloat16;

#define SG_SZ 8

#include "../joint_matrix_all_sizes_impl.hpp"
62 changes: 62 additions & 0 deletions SYCL/Matrix/joint_matrix_all_sizes_bfloat16.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,62 @@
//==-------- joint_matrix_all_sizes.cpp - DPC++ joint_matrix---------------==//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
// REQUIRES: matrix

// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -DJOINT_MATRIX_TEST_ALL_BF16
// RUN: %CPU_RUN_PLACEHOLDER %t.out

// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -DJOINT_MATRIX_TEST_BF16_1
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=0 %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=1 %GPU_RUN_PLACEHOLDER %t.out

// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -DJOINT_MATRIX_TEST_BF16_2
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=0 %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=1 %GPU_RUN_PLACEHOLDER %t.out

// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -DJOINT_MATRIX_TEST_BF16_3
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=0 %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=1 %GPU_RUN_PLACEHOLDER %t.out

// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -DJOINT_MATRIX_TEST_BF16_4
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=0 %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=1 %GPU_RUN_PLACEHOLDER %t.out

// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -DJOINT_MATRIX_TEST_BF16_5
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=0 %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=1 %GPU_RUN_PLACEHOLDER %t.out

// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -DJOINT_MATRIX_TEST_BF16_6
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=0 %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=1 %GPU_RUN_PLACEHOLDER %t.out

// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -DJOINT_MATRIX_TEST_BF16_7
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=0 %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=1 %GPU_RUN_PLACEHOLDER %t.out

// RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -DJOINT_MATRIX_TEST_BF16_8
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=0 %GPU_RUN_PLACEHOLDER %t.out
// RUN: env IGC_JointMatrixLoadStoreOpt=1 %GPU_RUN_PLACEHOLDER %t.out

#include <iostream>
#include <sycl/sycl.hpp>

using namespace sycl;
using namespace sycl::ext::oneapi::experimental::matrix;
using bfloat16 = sycl::ext::oneapi::bfloat16;

#define SG_SZ 16

#include "joint_matrix_all_sizes_impl.hpp"
Loading