From d6e4af1dd71f1d665c8120d659c44dfb80df8afb Mon Sep 17 00:00:00 2001 From: "Plyakhin, Yury" Date: Mon, 16 Sep 2024 17:12:39 -0700 Subject: [PATCH 1/3] [SYCL][Joint Matrix][E2E] Auto-detect Joint Matrix features in LIT Remove dependency of Matrix tests on LIT's params like `gpu-intel-pvc=True`, `matrix=1`, etc. intel/llvm mostly got rid of parameters like that with auto-detection of architecture and usage of aspects. It would be more maintainable and less error-prone if all of features could be auto-detected. --- sycl/test-e2e/Matrix/SG32/element_wise_all_ops_tf32.cpp | 2 +- sycl/test-e2e/Matrix/SG32/joint_matrix_tf32.cpp | 5 +---- .../SG32/element_wise_all_ops_tf32.cpp | 4 +--- .../SPVCooperativeMatrix/SG32/joint_matrix_tf32.cpp | 7 +------ .../SPVCooperativeMatrix/element_wise_all_ops_tf32.cpp | 4 +--- .../Matrix/SPVCooperativeMatrix/joint_matrix_tf32.cpp | 4 +--- sycl/test-e2e/Matrix/element_wise_all_ops_tf32.cpp | 2 +- sycl/test-e2e/Matrix/joint_matrix_tf32.cpp | 2 +- sycl/test-e2e/lit.cfg.py | 9 --------- 9 files changed, 8 insertions(+), 31 deletions(-) diff --git a/sycl/test-e2e/Matrix/SG32/element_wise_all_ops_tf32.cpp b/sycl/test-e2e/Matrix/SG32/element_wise_all_ops_tf32.cpp index 06d459a2a3ce5..743b747e4ff45 100644 --- a/sycl/test-e2e/Matrix/SG32/element_wise_all_ops_tf32.cpp +++ b/sycl/test-e2e/Matrix/SG32/element_wise_all_ops_tf32.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix-tf32 +// REQUIRES: aspect-ext_intel_matrix, arch-intel_gpu_pvc // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -o %t.out diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_tf32.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_tf32.cpp index 18da250bc808d..62cd8b402517f 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_tf32.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_tf32.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix-tf32 +// REQUIRES: aspect-ext_intel_matrix, arch-intel_gpu_pvc // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -o %t.out @@ -13,9 +13,6 @@ #include "../common.hpp" -using namespace sycl; -using namespace sycl::ext::oneapi::experimental::matrix; - #define SG_SZ 32 constexpr size_t TN = 16; diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_tf32.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_tf32.cpp index 613bcd1f7650e..30ab8923a3ea9 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_tf32.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_tf32.cpp @@ -5,14 +5,12 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix-tf32 +// REQUIRES: aspect-ext_intel_matrix, arch-intel_gpu_pvc // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out // RUN: %{run} %t.out -// XFAIL: cpu - #include "../../common.hpp" #define SG_SZ 32 diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_tf32.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_tf32.cpp index e2ae342ed4598..41d6712e577ef 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_tf32.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_tf32.cpp @@ -5,19 +5,14 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix-tf32 +// REQUIRES: aspect-ext_intel_matrix, arch-intel_gpu_pvc // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out // RUN: %{run} %t.out -// XFAIL: cpu - #include "../../common.hpp" -using namespace sycl; -using namespace sycl::ext::oneapi::experimental::matrix; - #define SG_SZ 32 constexpr size_t TN = 16; diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_tf32.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_tf32.cpp index cc26efe30d8b8..d804228eb23df 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_tf32.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_tf32.cpp @@ -5,14 +5,12 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix-tf32 +// REQUIRES: aspect-ext_intel_matrix, arch-intel_gpu_pvc // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out // RUN: %{run} %t.out -// XFAIL: cpu - #include "../common.hpp" constexpr size_t TN = 16; diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_tf32.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_tf32.cpp index 3033631282058..f1714223abea9 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_tf32.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_tf32.cpp @@ -5,14 +5,12 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix-tf32 +// REQUIRES: aspect-ext_intel_matrix, arch-intel_gpu_pvc // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out // RUN: %{run} %t.out -// XFAIL: cpu - #include "../common.hpp" constexpr size_t TN = 16; diff --git a/sycl/test-e2e/Matrix/element_wise_all_ops_tf32.cpp b/sycl/test-e2e/Matrix/element_wise_all_ops_tf32.cpp index 6e2f8dcff6384..fe16d31455663 100644 --- a/sycl/test-e2e/Matrix/element_wise_all_ops_tf32.cpp +++ b/sycl/test-e2e/Matrix/element_wise_all_ops_tf32.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix-tf32 +// REQUIRES: aspect-ext_intel_matrix, arch-intel_gpu_pvc // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -o %t.out diff --git a/sycl/test-e2e/Matrix/joint_matrix_tf32.cpp b/sycl/test-e2e/Matrix/joint_matrix_tf32.cpp index 922b79f356e78..01ba70d8559ad 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_tf32.cpp +++ b/sycl/test-e2e/Matrix/joint_matrix_tf32.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix-tf32 +// REQUIRES: aspect-ext_intel_matrix, arch-intel_gpu_pvc // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -o %t.out diff --git a/sycl/test-e2e/lit.cfg.py b/sycl/test-e2e/lit.cfg.py index f7d56aaee8f8e..bffb4892a07e4 100644 --- a/sycl/test-e2e/lit.cfg.py +++ b/sycl/test-e2e/lit.cfg.py @@ -156,23 +156,14 @@ config.available_features.add( "matrix-fp16" ) # PVC implies the support of FP16 matrix - config.available_features.add( - "matrix-tf32" - ) # PVC implies the support of TF32 matrix if lit_config.params.get("gpu-intel-pvc-vg", False): config.available_features.add("gpu-intel-pvc-vg") config.available_features.add( "matrix-fp16" ) # PVC-VG implies the support of FP16 matrix - config.available_features.add( - "matrix-tf32" - ) # PVC-VG implies the support of TF32 matrix if lit_config.params.get("matrix", False): config.available_features.add("matrix") -if lit_config.params.get("matrix-tf32", False): - config.available_features.add("matrix-tf32") - if lit_config.params.get("matrix-xmx8", False): config.available_features.add("matrix-xmx8") config.available_features.add( From fae9196781b4f1befb0fd1715342169bcc51f3a9 Mon Sep 17 00:00:00 2001 From: "Plyakhin, Yury" Date: Tue, 17 Sep 2024 11:06:59 -0700 Subject: [PATCH 2/3] second set of changes --- .../Matrix/SG32/element_wise_all_ops_half.cpp | 2 +- .../Matrix/SG32/element_wise_all_sizes.cpp | 2 +- .../SG32/joint_matrix_annotated_ptr.cpp | 2 +- ...nt_matrix_bfloat16_colmajorA_colmajorB.cpp | 2 +- .../SG32/joint_matrix_bfloat16_packedB.cpp | 2 +- .../SG32/joint_matrix_colA_rowB_colC.cpp | 2 +- .../Matrix/SG32/joint_matrix_down_convert.cpp | 2 +- .../Matrix/SG32/joint_matrix_half.cpp | 1 - .../joint_matrix_int8_colmajorA_colmajorB.cpp | 2 +- .../SG32/element_wise_all_ops_half.cpp | 2 +- .../SG32/element_wise_all_sizes.cpp | 2 +- .../SG32/joint_matrix_annotated_ptr.cpp | 2 +- ...nt_matrix_bfloat16_colmajorA_colmajorB.cpp | 2 +- .../SG32/joint_matrix_bfloat16_packedB.cpp | 2 +- .../SG32/joint_matrix_colA_rowB_colC.cpp | 2 +- .../SG32/joint_matrix_down_convert.cpp | 2 +- .../SG32/joint_matrix_half.cpp | 1 - .../joint_matrix_int8_colmajorA_colmajorB.cpp | 2 +- .../SG32/joint_matrix_out_bounds.cpp | 2 +- .../element_wise_all_ops_half.cpp | 1 - .../joint_matrix_annotated_ptr.cpp | 2 +- .../joint_matrix_bf16_fill_k_cache_OOB.cpp | 3 --- ...nt_matrix_bfloat16_colmajorA_colmajorB.cpp | 2 +- .../joint_matrix_bfloat16_packedB.cpp | 2 +- .../joint_matrix_colA_rowB_colC.cpp | 2 +- .../joint_matrix_down_convert.cpp | 2 +- .../joint_matrix_half.cpp | 1 - .../joint_matrix_int8_colmajorA_colmajorB.cpp | 2 +- .../joint_matrix_prefetch.cpp | 2 +- .../joint_matrix_unaligned_k.cpp | 2 +- .../Matrix/element_wise_all_ops_half.cpp | 1 - .../Matrix/element_wise_all_ops_half_impl.hpp | 3 +++ .../Matrix/elemwise_irreg_size_ops_bf16.cpp | 2 +- .../Matrix/joint_matrix_annotated_ptr.cpp | 2 +- .../joint_matrix_bf16_fill_k_cache_OOB.cpp | 3 --- ...nt_matrix_bfloat16_colmajorA_colmajorB.cpp | 2 +- .../Matrix/joint_matrix_bfloat16_packedB.cpp | 2 +- .../Matrix/joint_matrix_colA_rowB_colC.cpp | 2 +- .../Matrix/joint_matrix_down_convert.cpp | 2 +- sycl/test-e2e/Matrix/joint_matrix_half.cpp | 1 - .../Matrix/joint_matrix_half_impl.hpp | 3 +++ .../joint_matrix_int8_colmajorA_colmajorB.cpp | 2 +- .../Matrix/joint_matrix_query_default.cpp | 2 +- sycl/test-e2e/lit.cfg.py | 20 ++----------------- 44 files changed, 41 insertions(+), 63 deletions(-) diff --git a/sycl/test-e2e/Matrix/SG32/element_wise_all_ops_half.cpp b/sycl/test-e2e/Matrix/SG32/element_wise_all_ops_half.cpp index 847f4a7812aa2..8468ebad1b8b3 100644 --- a/sycl/test-e2e/Matrix/SG32/element_wise_all_ops_half.cpp +++ b/sycl/test-e2e/Matrix/SG32/element_wise_all_ops_half.cpp @@ -8,7 +8,7 @@ // SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 // UNSUPPORTED: gpu-intel-dg2 // REQUIRES: aspect-fp16 -// REQUIRES: aspect-ext_intel_matrix, gpu +// REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -o %t.out diff --git a/sycl/test-e2e/Matrix/SG32/element_wise_all_sizes.cpp b/sycl/test-e2e/Matrix/SG32/element_wise_all_sizes.cpp index 4624110577ea2..55938aea88f55 100644 --- a/sycl/test-e2e/Matrix/SG32/element_wise_all_sizes.cpp +++ b/sycl/test-e2e/Matrix/SG32/element_wise_all_sizes.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -o %t.out diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_annotated_ptr.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_annotated_ptr.cpp index 7a6aca1024418..21f101e000b52 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_annotated_ptr.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_annotated_ptr.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -o %t.out diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16_colmajorA_colmajorB.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16_colmajorA_colmajorB.cpp index 6532bcfe47bff..aba19833ee581 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16_colmajorA_colmajorB.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16_colmajorA_colmajorB.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -o %t.out diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16_packedB.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16_packedB.cpp index d2ba1cdbdd38e..cd4c8993f71a0 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16_packedB.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16_packedB.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -o %t.out diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_colA_rowB_colC.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_colA_rowB_colC.cpp index 70e53441cb48f..a10810bf4549c 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_colA_rowB_colC.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_colA_rowB_colC.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -o %t.out diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_down_convert.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_down_convert.cpp index b474f846d11d5..423c225118578 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_down_convert.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_down_convert.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -o %t.out diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_half.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_half.cpp index 2b990ea3c6098..a02206ee7a805 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_half.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_half.cpp @@ -9,7 +9,6 @@ // UNSUPPORTED: gpu-intel-dg2 // REQUIRES: aspect-fp16 // REQUIRES: aspect-ext_intel_matrix -// REQUIRES: matrix-fp16 // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -o %t.out diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_int8_colmajorA_colmajorB.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_int8_colmajorA_colmajorB.cpp index c89c657c77fbc..177e4d6d861fa 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_int8_colmajorA_colmajorB.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_int8_colmajorA_colmajorB.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -o %t.out diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_half.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_half.cpp index c90aa1b824a6d..a922c241e6677 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_half.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_half.cpp @@ -8,7 +8,7 @@ // SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 // UNSUPPORTED: gpu-intel-dg2 // REQUIRES: aspect-fp16 -// REQUIRES: aspect-ext_intel_matrix, gpu +// REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_sizes.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_sizes.cpp index b473dc00a1844..bec3d15c495ef 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_sizes.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_sizes.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_annotated_ptr.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_annotated_ptr.cpp index 21b9014e7bff9..2e51654fd8c51 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_annotated_ptr.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_annotated_ptr.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_colmajorA_colmajorB.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_colmajorA_colmajorB.cpp index 80de9af24bb85..314b529ebfb7b 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_colmajorA_colmajorB.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_colmajorA_colmajorB.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_packedB.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_packedB.cpp index aa8e4bf7758ca..36d8a3317a762 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_packedB.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_packedB.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_colA_rowB_colC.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_colA_rowB_colC.cpp index 6a51e03bebd23..3ce16e94a40ba 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_colA_rowB_colC.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_colA_rowB_colC.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_down_convert.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_down_convert.cpp index efe3c2b456802..b235ba6afd89c 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_down_convert.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_down_convert.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_half.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_half.cpp index 16953caff99e0..3d82602130459 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_half.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_half.cpp @@ -9,7 +9,6 @@ // UNSUPPORTED: gpu-intel-dg2 // REQUIRES: aspect-fp16 // REQUIRES: aspect-ext_intel_matrix -// REQUIRES: matrix-fp16 // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_int8_colmajorA_colmajorB.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_int8_colmajorA_colmajorB.cpp index 375e3bb958733..a5874a5e90915 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_int8_colmajorA_colmajorB.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_int8_colmajorA_colmajorB.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_out_bounds.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_out_bounds.cpp index 6926e7951b0db..f78eb4ef6c1f9 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_out_bounds.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_out_bounds.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // UNSUPPORTED: gpu-intel-dg2, cpu diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_half.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_half.cpp index 088bda535d2de..f907360c25f5f 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_half.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_half.cpp @@ -7,7 +7,6 @@ //===----------------------------------------------------------------------===// // REQUIRES: aspect-fp16 // REQUIRES: aspect-ext_intel_matrix -// REQUIRES: matrix-fp16 // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_annotated_ptr.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_annotated_ptr.cpp index 52e680fb25070..ffca68fcfe6a2 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_annotated_ptr.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_annotated_ptr.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_OOB.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_OOB.cpp index ffa651fc98d58..af4b28e090b4c 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_OOB.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_OOB.cpp @@ -8,9 +8,6 @@ // REQUIRES: aspect-ext_intel_matrix, gpu // UNSUPPORTED: gpu-intel-dg2 -// https://github.com/intel/llvm/issues/14826 -// XFAIL: arch-intel_gpu_pvc && igc-dev - // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t_gpu_vnni.out -ffp-model=precise -DOOB -DVNNI // RUN: %{run} %t_gpu_vnni.out diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_colmajorA_colmajorB.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_colmajorA_colmajorB.cpp index ff7d2f1e05882..12b74948436e4 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_colmajorA_colmajorB.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_colmajorA_colmajorB.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_packedB.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_packedB.cpp index 1b7de0a673641..43c5d10739cfa 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_packedB.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_packedB.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27868, win: 101.5181 // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_colA_rowB_colC.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_colA_rowB_colC.cpp index 2c6323364ead9..083b7cbed8ef7 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_colA_rowB_colC.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_colA_rowB_colC.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_down_convert.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_down_convert.cpp index 9ac82ceadba6a..c17bf404cf850 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_down_convert.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_down_convert.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_half.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_half.cpp index 43b70b990b7d1..6dab4c8445baa 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_half.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_half.cpp @@ -7,7 +7,6 @@ //===----------------------------------------------------------------------===// // REQUIRES: aspect-fp16 // REQUIRES: aspect-ext_intel_matrix -// REQUIRES: matrix-fp16 // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_int8_colmajorA_colmajorB.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_int8_colmajorA_colmajorB.cpp index e0d2c5bd81769..cf462e53c86cb 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_int8_colmajorA_colmajorB.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_int8_colmajorA_colmajorB.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_prefetch.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_prefetch.cpp index 3bf4dd7608f48..cc1c848f82fba 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_prefetch.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_prefetch.cpp @@ -11,7 +11,7 @@ // XFAIL: cpu // https://github.com/intel/llvm/issues/14826 -// XFAIL: arch-intel_gpu_pvc && !dev-igc +// XFAIL: arch-intel_gpu_pvc && !igc-dev #include "../common.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_unaligned_k.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_unaligned_k.cpp index 43f3c12b70079..0f6d6fad9d673 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_unaligned_k.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_unaligned_k.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // UNSUPPORTED: gpu-intel-dg2, cpu diff --git a/sycl/test-e2e/Matrix/element_wise_all_ops_half.cpp b/sycl/test-e2e/Matrix/element_wise_all_ops_half.cpp index bb651568f9251..f97241f275bd1 100644 --- a/sycl/test-e2e/Matrix/element_wise_all_ops_half.cpp +++ b/sycl/test-e2e/Matrix/element_wise_all_ops_half.cpp @@ -7,7 +7,6 @@ //===----------------------------------------------------------------------===// // REQUIRES: aspect-fp16 // REQUIRES: aspect-ext_intel_matrix -// REQUIRES: matrix-fp16 // RUN: %{build} -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/Matrix/element_wise_all_ops_half_impl.hpp b/sycl/test-e2e/Matrix/element_wise_all_ops_half_impl.hpp index 4065c7a78a566..d2bdcbcb2d04a 100644 --- a/sycl/test-e2e/Matrix/element_wise_all_ops_half_impl.hpp +++ b/sycl/test-e2e/Matrix/element_wise_all_ops_half_impl.hpp @@ -107,6 +107,9 @@ int main() { matrix_combinations>(); for (unsigned int i = 0; i < combinations.size(); i++) { + if (combinations[i].atype != matrix_type::fp16) + continue; + if (combinations[i].nsize == 0) { // Intel AMX test(); break; diff --git a/sycl/test-e2e/Matrix/elemwise_irreg_size_ops_bf16.cpp b/sycl/test-e2e/Matrix/elemwise_irreg_size_ops_bf16.cpp index 5b4c785d46d3d..eb0a22fa30566 100644 --- a/sycl/test-e2e/Matrix/elemwise_irreg_size_ops_bf16.cpp +++ b/sycl/test-e2e/Matrix/elemwise_irreg_size_ops_bf16.cpp @@ -9,7 +9,7 @@ // SG size. This corner case only applies to AMX. Also, it tests bf16 type. // only run this on AMX // REQUIRES: cpu -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // RUN: %{build} -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/Matrix/joint_matrix_annotated_ptr.cpp b/sycl/test-e2e/Matrix/joint_matrix_annotated_ptr.cpp index e21bdaa06a328..ce19e33652011 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_annotated_ptr.cpp +++ b/sycl/test-e2e/Matrix/joint_matrix_annotated_ptr.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // RUN: %{build} -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/Matrix/joint_matrix_bf16_fill_k_cache_OOB.cpp b/sycl/test-e2e/Matrix/joint_matrix_bf16_fill_k_cache_OOB.cpp index ad0cda4824b76..0f2c8c71bd223 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_bf16_fill_k_cache_OOB.cpp +++ b/sycl/test-e2e/Matrix/joint_matrix_bf16_fill_k_cache_OOB.cpp @@ -8,9 +8,6 @@ // REQUIRES: aspect-ext_intel_matrix, gpu // UNSUPPORTED: gpu-intel-dg2 -// https://github.com/intel/llvm/issues/14826 -// XFAIL: arch-intel_gpu_pvc && igc-dev - // RUN: %{build} -o %t_gpu_vnni.out -ffp-model=precise -DOOB -DVNNI // RUN: %{run} %t_gpu_vnni.out diff --git a/sycl/test-e2e/Matrix/joint_matrix_bfloat16_colmajorA_colmajorB.cpp b/sycl/test-e2e/Matrix/joint_matrix_bfloat16_colmajorA_colmajorB.cpp index 21d5f1239cd8d..a9326e2ba312b 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_bfloat16_colmajorA_colmajorB.cpp +++ b/sycl/test-e2e/Matrix/joint_matrix_bfloat16_colmajorA_colmajorB.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // RUN: %{build} -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/Matrix/joint_matrix_bfloat16_packedB.cpp b/sycl/test-e2e/Matrix/joint_matrix_bfloat16_packedB.cpp index bf3a65d2a16d7..cf2ff9007d5b7 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_bfloat16_packedB.cpp +++ b/sycl/test-e2e/Matrix/joint_matrix_bfloat16_packedB.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27868, win: 101.5181 // RUN: %{build} -o %t.out diff --git a/sycl/test-e2e/Matrix/joint_matrix_colA_rowB_colC.cpp b/sycl/test-e2e/Matrix/joint_matrix_colA_rowB_colC.cpp index 354a71006e129..71e3cb6501bb6 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_colA_rowB_colC.cpp +++ b/sycl/test-e2e/Matrix/joint_matrix_colA_rowB_colC.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // RUN: %{build} -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/Matrix/joint_matrix_down_convert.cpp b/sycl/test-e2e/Matrix/joint_matrix_down_convert.cpp index dee504c22e7f6..5d52acd0a1d87 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_down_convert.cpp +++ b/sycl/test-e2e/Matrix/joint_matrix_down_convert.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // RUN: %{build} -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/Matrix/joint_matrix_half.cpp b/sycl/test-e2e/Matrix/joint_matrix_half.cpp index 44fa8cd86e7cb..7788b64e6d8ca 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_half.cpp +++ b/sycl/test-e2e/Matrix/joint_matrix_half.cpp @@ -7,7 +7,6 @@ //===----------------------------------------------------------------------===// // REQUIRES: aspect-fp16 // REQUIRES: aspect-ext_intel_matrix -// REQUIRES: matrix-fp16 // RUN: %{build} -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/Matrix/joint_matrix_half_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_half_impl.hpp index a9de620af4237..0ff483b051067 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_half_impl.hpp +++ b/sycl/test-e2e/Matrix/joint_matrix_half_impl.hpp @@ -112,6 +112,9 @@ int main() { matrix_combinations>(); for (unsigned int i = 0; i < combinations.size(); i++) { + if (combinations[i].atype != matrix_type::fp16) + continue; + if (combinations[i].nsize == 0) { // Intel AMX test(); break; diff --git a/sycl/test-e2e/Matrix/joint_matrix_int8_colmajorA_colmajorB.cpp b/sycl/test-e2e/Matrix/joint_matrix_int8_colmajorA_colmajorB.cpp index 37769a41f7003..f8dc8452668a1 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_int8_colmajorA_colmajorB.cpp +++ b/sycl/test-e2e/Matrix/joint_matrix_int8_colmajorA_colmajorB.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // RUN: %{build} -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/Matrix/joint_matrix_query_default.cpp b/sycl/test-e2e/Matrix/joint_matrix_query_default.cpp index f160c09525aa7..2dd752b0a9c78 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_query_default.cpp +++ b/sycl/test-e2e/Matrix/joint_matrix_query_default.cpp @@ -7,7 +7,7 @@ //===----------------------------------------------------------------------===// // Needs AMX. // REQUIRES: cpu -// REQUIRES: matrix +// REQUIRES: aspect-ext_intel_matrix // RUN: %{build} -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/lit.cfg.py b/sycl/test-e2e/lit.cfg.py index bffb4892a07e4..c2582467dd4cf 100644 --- a/sycl/test-e2e/lit.cfg.py +++ b/sycl/test-e2e/lit.cfg.py @@ -152,27 +152,11 @@ config.available_features.add("gpu-intel-dg1") if lit_config.params.get("gpu-intel-dg2", False): config.available_features.add("gpu-intel-dg2") -if lit_config.params.get("gpu-intel-pvc", False): - config.available_features.add( - "matrix-fp16" - ) # PVC implies the support of FP16 matrix if lit_config.params.get("gpu-intel-pvc-vg", False): config.available_features.add("gpu-intel-pvc-vg") - config.available_features.add( - "matrix-fp16" - ) # PVC-VG implies the support of FP16 matrix -if lit_config.params.get("matrix", False): - config.available_features.add("matrix") - -if lit_config.params.get("matrix-xmx8", False): - config.available_features.add("matrix-xmx8") - config.available_features.add( - "matrix-fp16" - ) # XMX implies the support of FP16 matrix - -if lit_config.params.get("matrix-fp16", False): - config.available_features.add("matrix-fp16") +if lit_config.params.get("igc-dev", False): + config.available_features.add("igc-dev") def check_igc_tag_and_add_feature(): if os.path.isfile(config.igc_tag_file): From 336c6d0243afe9f1ea9b112cffaf1d968450fd50 Mon Sep 17 00:00:00 2001 From: "Plyakhin, Yury" Date: Tue, 17 Sep 2024 13:17:27 -0700 Subject: [PATCH 3/3] more fixes --- sycl/test-e2e/Matrix/SG32/element_wise_all_ops_tf32.cpp | 2 +- sycl/test-e2e/Matrix/SG32/element_wise_all_sizes.cpp | 3 +++ sycl/test-e2e/Matrix/SG32/joint_matrix_down_convert.cpp | 3 +++ sycl/test-e2e/Matrix/SG32/joint_matrix_tf32.cpp | 2 +- .../SPVCooperativeMatrix/SG32/element_wise_all_ops_tf32.cpp | 2 +- .../SPVCooperativeMatrix/SG32/element_wise_all_sizes.cpp | 3 +++ .../SPVCooperativeMatrix/SG32/joint_matrix_down_convert.cpp | 3 +++ .../Matrix/SPVCooperativeMatrix/SG32/joint_matrix_tf32.cpp | 2 +- .../SPVCooperativeMatrix/element_wise_all_ops_tf32.cpp | 2 +- .../SPVCooperativeMatrix/joint_matrix_annotated_ptr.cpp | 1 + .../SPVCooperativeMatrix/joint_matrix_down_convert.cpp | 1 + .../Matrix/SPVCooperativeMatrix/joint_matrix_tf32.cpp | 2 +- sycl/test-e2e/Matrix/element_wise_all_ops_tf32.cpp | 2 +- sycl/test-e2e/Matrix/element_wise_all_ops_tf32_impl.hpp | 5 +++++ sycl/test-e2e/Matrix/joint_matrix_annotated_ptr.cpp | 1 + sycl/test-e2e/Matrix/joint_matrix_down_convert.cpp | 1 + sycl/test-e2e/Matrix/joint_matrix_tf32.cpp | 2 +- sycl/test-e2e/Matrix/joint_matrix_tf32_impl.hpp | 6 ++++++ 18 files changed, 35 insertions(+), 8 deletions(-) diff --git a/sycl/test-e2e/Matrix/SG32/element_wise_all_ops_tf32.cpp b/sycl/test-e2e/Matrix/SG32/element_wise_all_ops_tf32.cpp index 743b747e4ff45..6b6416280a6bd 100644 --- a/sycl/test-e2e/Matrix/SG32/element_wise_all_ops_tf32.cpp +++ b/sycl/test-e2e/Matrix/SG32/element_wise_all_ops_tf32.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: aspect-ext_intel_matrix, arch-intel_gpu_pvc +// REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -o %t.out diff --git a/sycl/test-e2e/Matrix/SG32/element_wise_all_sizes.cpp b/sycl/test-e2e/Matrix/SG32/element_wise_all_sizes.cpp index 55938aea88f55..df0cc30b6aae1 100644 --- a/sycl/test-e2e/Matrix/SG32/element_wise_all_sizes.cpp +++ b/sycl/test-e2e/Matrix/SG32/element_wise_all_sizes.cpp @@ -8,6 +8,9 @@ // REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 + // RUN: %{build} -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_down_convert.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_down_convert.cpp index 423c225118578..ac6a308afdd0c 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_down_convert.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_down_convert.cpp @@ -8,6 +8,9 @@ // REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 + // RUN: %{build} -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_tf32.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_tf32.cpp index 62cd8b402517f..8d9f7867e2508 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_tf32.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_tf32.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: aspect-ext_intel_matrix, arch-intel_gpu_pvc +// REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -o %t.out diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_tf32.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_tf32.cpp index 30ab8923a3ea9..b9972d74ba18d 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_tf32.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_tf32.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: aspect-ext_intel_matrix, arch-intel_gpu_pvc +// REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_sizes.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_sizes.cpp index bec3d15c495ef..0ff66f1ce7ed7 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_sizes.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_sizes.cpp @@ -8,6 +8,9 @@ // REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 + // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_down_convert.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_down_convert.cpp index b235ba6afd89c..73643d827a260 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_down_convert.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_down_convert.cpp @@ -8,6 +8,9 @@ // REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 +// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 +// UNSUPPORTED: gpu-intel-dg2 + // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_tf32.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_tf32.cpp index 41d6712e577ef..57fd883fc9a99 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_tf32.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_tf32.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: aspect-ext_intel_matrix, arch-intel_gpu_pvc +// REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_tf32.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_tf32.cpp index d804228eb23df..7a96dd76324f1 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_tf32.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_tf32.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: aspect-ext_intel_matrix, arch-intel_gpu_pvc +// REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_annotated_ptr.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_annotated_ptr.cpp index ffca68fcfe6a2..c4b748241a172 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_annotated_ptr.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_annotated_ptr.cpp @@ -6,6 +6,7 @@ // //===----------------------------------------------------------------------===// // REQUIRES: aspect-ext_intel_matrix +// UNSUPPORTED: gpu-intel-dg2 // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_down_convert.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_down_convert.cpp index c17bf404cf850..6ac1147369880 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_down_convert.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_down_convert.cpp @@ -6,6 +6,7 @@ // //===----------------------------------------------------------------------===// // REQUIRES: aspect-ext_intel_matrix +// UNSUPPORTED: gpu-intel-dg2 // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_tf32.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_tf32.cpp index f1714223abea9..322a80339de77 100644 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_tf32.cpp +++ b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_tf32.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: aspect-ext_intel_matrix, arch-intel_gpu_pvc +// REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out diff --git a/sycl/test-e2e/Matrix/element_wise_all_ops_tf32.cpp b/sycl/test-e2e/Matrix/element_wise_all_ops_tf32.cpp index fe16d31455663..e42167061de42 100644 --- a/sycl/test-e2e/Matrix/element_wise_all_ops_tf32.cpp +++ b/sycl/test-e2e/Matrix/element_wise_all_ops_tf32.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: aspect-ext_intel_matrix, arch-intel_gpu_pvc +// REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -o %t.out diff --git a/sycl/test-e2e/Matrix/element_wise_all_ops_tf32_impl.hpp b/sycl/test-e2e/Matrix/element_wise_all_ops_tf32_impl.hpp index ba63c45deea2d..2d7b3a36d8296 100644 --- a/sycl/test-e2e/Matrix/element_wise_all_ops_tf32_impl.hpp +++ b/sycl/test-e2e/Matrix/element_wise_all_ops_tf32_impl.hpp @@ -64,6 +64,11 @@ static constexpr size_t MATRIX_K = TK * 2; float A[MATRIX_M][MATRIX_K]; int main() { + queue q; + if (!is_type_supported_by_device(q, matrix_type::tf32)) { + std::cout << "Joint Matrix TF32 is not supported by this device.\n"; + return 0; + } big_matrix MA((float *)&A); diff --git a/sycl/test-e2e/Matrix/joint_matrix_annotated_ptr.cpp b/sycl/test-e2e/Matrix/joint_matrix_annotated_ptr.cpp index ce19e33652011..f29b3a46efc96 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_annotated_ptr.cpp +++ b/sycl/test-e2e/Matrix/joint_matrix_annotated_ptr.cpp @@ -6,6 +6,7 @@ // //===----------------------------------------------------------------------===// // REQUIRES: aspect-ext_intel_matrix +// UNSUPPORTED: gpu-intel-dg2 // RUN: %{build} -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/Matrix/joint_matrix_down_convert.cpp b/sycl/test-e2e/Matrix/joint_matrix_down_convert.cpp index 5d52acd0a1d87..c6772a744f43a 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_down_convert.cpp +++ b/sycl/test-e2e/Matrix/joint_matrix_down_convert.cpp @@ -6,6 +6,7 @@ // //===----------------------------------------------------------------------===// // REQUIRES: aspect-ext_intel_matrix +// UNSUPPORTED: gpu-intel-dg2 // RUN: %{build} -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/Matrix/joint_matrix_tf32.cpp b/sycl/test-e2e/Matrix/joint_matrix_tf32.cpp index 01ba70d8559ad..32072c3e01b20 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_tf32.cpp +++ b/sycl/test-e2e/Matrix/joint_matrix_tf32.cpp @@ -5,7 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: aspect-ext_intel_matrix, arch-intel_gpu_pvc +// REQUIRES: aspect-ext_intel_matrix // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 // RUN: %{build} -o %t.out diff --git a/sycl/test-e2e/Matrix/joint_matrix_tf32_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_tf32_impl.hpp index 69991884c0710..9fb40a78f8b30 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_tf32_impl.hpp +++ b/sycl/test-e2e/Matrix/joint_matrix_tf32_impl.hpp @@ -91,6 +91,12 @@ void matrix_multiply(big_matrix &C, } int main() { + queue q; + if (!is_type_supported_by_device(q, matrix_type::tf32)) { + std::cout << "Joint Matrix TF32 is not supported by this device.\n"; + return 0; + } + static constexpr size_t MATRIX_M = TM * 2; static constexpr size_t MATRIX_N = TN * 2; static constexpr size_t MATRIX_K = TK * 2;