From 5f7c343853919ad7bab6819060f4d247e572c135 Mon Sep 17 00:00:00 2001 From: "Sidorov, Dmitry" Date: Mon, 11 Nov 2024 08:17:45 -0800 Subject: [PATCH 1/2] [SYCL][NFCI] Finalize switch to SPV_KHR_cooperative_matrix Signed-off-by: Sidorov, Dmitry --- clang/lib/CodeGen/CodeGenTypes.cpp | 75 ------- clang/lib/CodeGen/CodeGenTypes.h | 8 - clang/test/CodeGenSYCL/joint_matrix.cpp | 41 ---- sycl/include/sycl/__spirv/spirv_ops.hpp | 150 -------------- sycl/include/sycl/__spirv/spirv_types.hpp | 10 - .../sycl/ext/oneapi/matrix/matrix-intel.hpp | 191 ------------------ .../oneapi/matrix/matrix-unified-utils.hpp | 2 - .../sycl/ext/oneapi/matrix/matrix-unified.hpp | 73 ------- .../SG32/element_wise_abc.cpp | 22 -- .../SG32/element_wise_all_ops.cpp | 22 -- .../SG32/element_wise_all_ops_half.cpp | 25 --- .../SG32/element_wise_all_ops_int8.cpp | 22 -- .../SG32/element_wise_all_ops_int8_packed.cpp | 24 --- .../SG32/element_wise_all_ops_tf32.cpp | 19 -- .../SG32/element_wise_all_sizes.cpp | 23 --- .../SG32/element_wise_ops.cpp | 22 -- .../SG32/get_coord_float_matC.cpp | 22 -- .../SG32/get_coord_int8_matA.cpp | 22 -- .../SG32/get_coord_int8_matB.cpp | 21 -- .../SG32/joint_matrix_all_sizes.cpp | 22 -- .../SG32/joint_matrix_annotated_ptr.cpp | 24 --- .../SG32/joint_matrix_apply_bf16.cpp | 22 -- .../SG32/joint_matrix_apply_two_matrices.cpp | 23 --- .../SG32/joint_matrix_bf16_fill_k_cache.cpp | 27 --- .../joint_matrix_bf16_fill_k_cache_SLM.cpp | 23 --- .../joint_matrix_bf16_fill_k_cache_init.cpp | 22 -- .../joint_matrix_bf16_fill_k_cache_unroll.cpp | 26 --- ...t_matrix_bf16_fill_k_cache_unroll_init.cpp | 24 --- .../SG32/joint_matrix_bfloat16.cpp | 22 -- .../SG32/joint_matrix_bfloat16_array.cpp | 22 -- ...nt_matrix_bfloat16_colmajorA_colmajorB.cpp | 31 --- .../SG32/joint_matrix_bfloat16_packedB.cpp | 23 --- .../SG32/joint_matrix_colA_rowB_colC.cpp | 21 -- .../SG32/joint_matrix_down_convert.cpp | 23 --- .../SG32/joint_matrix_half.cpp | 25 --- .../joint_matrix_int8_colmajorA_colmajorB.cpp | 27 --- .../joint_matrix_int8_rowmajorA_rowmajorB.cpp | 22 -- .../SG32/joint_matrix_out_bounds.cpp | 24 --- .../SG32/joint_matrix_prefetch.cpp | 23 --- .../SG32/joint_matrix_rowmajorA_rowmajorB.cpp | 29 --- .../SG32/joint_matrix_ss_int8.cpp | 22 -- .../SG32/joint_matrix_su_int8.cpp | 22 -- .../SG32/joint_matrix_tf32.cpp | 19 -- .../SG32/joint_matrix_transposeC.cpp | 20 -- .../SG32/joint_matrix_unaligned_k.cpp | 24 --- .../SG32/joint_matrix_us_int8.cpp | 22 -- .../SG32/joint_matrix_uu_int8.cpp | 22 -- .../SPVCooperativeMatrix/element_wise_abc.cpp | 16 -- .../element_wise_all_ops.cpp | 16 -- .../element_wise_all_ops_1d.cpp | 14 -- .../element_wise_all_ops_1d_cont.cpp | 14 -- .../element_wise_all_ops_half.cpp | 20 -- .../element_wise_all_ops_int8.cpp | 16 -- .../element_wise_all_ops_int8_packed.cpp | 18 -- .../element_wise_all_ops_scalar.cpp | 16 -- .../element_wise_all_ops_tf32.cpp | 18 -- .../element_wise_all_sizes.cpp | 22 -- .../SPVCooperativeMatrix/element_wise_ops.cpp | 16 -- .../get_coord_float_matC.cpp | 16 -- .../get_coord_int8_matA.cpp | 16 -- .../get_coord_int8_matB.cpp | 15 -- .../joint_matrix_all_sizes.cpp | 16 -- .../joint_matrix_annotated_ptr.cpp | 22 -- .../joint_matrix_apply_bf16.cpp | 16 -- .../joint_matrix_apply_two_matrices.cpp | 17 -- .../joint_matrix_bf16_fill_k_cache.cpp | 21 -- .../joint_matrix_bf16_fill_k_cache_OOB.cpp | 20 -- .../joint_matrix_bf16_fill_k_cache_SLM.cpp | 19 -- .../joint_matrix_bf16_fill_k_cache_init.cpp | 16 -- .../joint_matrix_bf16_fill_k_cache_unroll.cpp | 20 -- ...t_matrix_bf16_fill_k_cache_unroll_init.cpp | 18 -- .../joint_matrix_bfloat16.cpp | 16 -- .../joint_matrix_bfloat16_array.cpp | 16 -- ...nt_matrix_bfloat16_colmajorA_colmajorB.cpp | 22 -- .../joint_matrix_bfloat16_packedB.cpp | 20 -- .../joint_matrix_colA_rowB_colC.cpp | 19 -- .../joint_matrix_down_convert.cpp | 17 -- .../joint_matrix_half.cpp | 20 -- .../joint_matrix_int8_colmajorA_colmajorB.cpp | 22 -- .../joint_matrix_int8_rowmajorA_rowmajorB.cpp | 19 -- .../joint_matrix_opt_kernel_feature.cpp | 18 -- .../joint_matrix_out_bounds.cpp | 20 -- .../joint_matrix_prefetch.cpp | 18 -- .../joint_matrix_rowmajorA_rowmajorB.cpp | 24 --- .../joint_matrix_ss_int8.cpp | 16 -- .../joint_matrix_su_int8.cpp | 16 -- .../joint_matrix_tf32.cpp | 18 -- .../joint_matrix_transposeC.cpp | 14 -- .../joint_matrix_unaligned_k.cpp | 20 -- .../joint_matrix_us_int8.cpp | 16 -- .../joint_matrix_uu_int8.cpp | 16 -- .../matrix/matrix-int8-test.cpp | 6 +- .../matrix/matrix_load_store_as.cpp | 8 +- .../no-unsupported-without-info.cpp | 41 ---- .../no-xfail-without-tracker.cpp | 65 ------ 95 files changed, 7 insertions(+), 2348 deletions(-) delete mode 100644 clang/test/CodeGenSYCL/joint_matrix.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_abc.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_half.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_int8.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_int8_packed.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_tf32.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_sizes.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_ops.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/get_coord_float_matC.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/get_coord_int8_matA.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/get_coord_int8_matB.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_all_sizes.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_annotated_ptr.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_apply_bf16.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_apply_two_matrices.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_SLM.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_init.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_unroll.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_unroll_init.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_array.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_colmajorA_colmajorB.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_packedB.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_colA_rowB_colC.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_down_convert.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_half.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_int8_colmajorA_colmajorB.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_int8_rowmajorA_rowmajorB.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_out_bounds.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_prefetch.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_rowmajorA_rowmajorB.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_ss_int8.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_su_int8.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_tf32.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_transposeC.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_unaligned_k.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_us_int8.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_uu_int8.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_abc.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_1d.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_1d_cont.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_half.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_int8.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_int8_packed.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_scalar.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_tf32.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_sizes.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_ops.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/get_coord_float_matC.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/get_coord_int8_matA.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/get_coord_int8_matB.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_all_sizes.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_annotated_ptr.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_apply_bf16.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_apply_two_matrices.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_OOB.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_SLM.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_init.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_unroll.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_unroll_init.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_array.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_colmajorA_colmajorB.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_packedB.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_colA_rowB_colC.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_down_convert.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_half.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_int8_colmajorA_colmajorB.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_int8_rowmajorA_rowmajorB.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_opt_kernel_feature.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_out_bounds.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_prefetch.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_rowmajorA_rowmajorB.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_ss_int8.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_su_int8.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_tf32.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_transposeC.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_unaligned_k.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_us_int8.cpp delete mode 100644 sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_uu_int8.cpp diff --git a/clang/lib/CodeGen/CodeGenTypes.cpp b/clang/lib/CodeGen/CodeGenTypes.cpp index 82de188e783b3..c0123286da9e9 100644 --- a/clang/lib/CodeGen/CodeGenTypes.cpp +++ b/clang/lib/CodeGen/CodeGenTypes.cpp @@ -349,34 +349,6 @@ llvm::Type *CodeGenTypes::ConvertFunctionTypeInternal(QualType QFT) { return ResultType; } -template -llvm::Type *getJointMatrixINTELExtType(llvm::Type *CompTy, - ArrayRef TemplateArgs, - const unsigned Val = 0) { - // TODO: we should actually have exactly 5 template parameters: 1 for - // type and 4 for type parameters. But in previous version of the SPIR-V - // spec we have Layout matrix type parameter, that was later removed. - // Once we update to the newest version of the spec - this should be updated. - assert((TemplateArgs.size() == 5 || TemplateArgs.size() == 6) && - "Wrong JointMatrixINTEL template parameters number"); - // This is required to represent optional 'Component Type Interpretation' - // parameter - std::vector Params; - for (size_t I = 1; I != TemplateArgs.size(); ++I) { - assert(TemplateArgs[I].getKind() == TemplateArgument::Integral && - "Wrong JointMatrixINTEL template parameter"); - Params.push_back(TemplateArgs[I].getAsIntegral().getExtValue()); - } - // Don't add type interpretation for legacy matrices. - // Legacy matrices has 5 template parameters, while new representation - // has 6. - if (NeedTypeInterpret && TemplateArgs.size() != 5) - Params.push_back(Val); - - return llvm::TargetExtType::get(CompTy->getContext(), - "spirv.JointMatrixINTEL", {CompTy}, Params); -} - llvm::Type * getCooperativeMatrixKHRExtType(llvm::Type *CompTy, ArrayRef TemplateArgs) { @@ -393,49 +365,6 @@ getCooperativeMatrixKHRExtType(llvm::Type *CompTy, CompTy->getContext(), "spirv.CooperativeMatrixKHR", {CompTy}, Params); } -/// ConvertSYCLJointMatrixINTELType - Convert SYCL joint_matrix type -/// which is represented as a pointer to a structure to LLVM extension type -/// with the parameters that follow SPIR-V JointMatrixINTEL type. -/// The expected representation is: -/// target("spirv.JointMatrixINTEL", %element_type, %rows%, %cols%, %scope%, -/// %use%, (optional) %element_type_interpretation%) -llvm::Type *CodeGenTypes::ConvertSYCLJointMatrixINTELType(RecordDecl *RD) { - auto *TemplateDecl = cast(RD); - ArrayRef TemplateArgs = - TemplateDecl->getTemplateArgs().asArray(); - assert(TemplateArgs[0].getKind() == TemplateArgument::Type && - "1st JointMatrixINTEL template parameter must be type"); - llvm::Type *CompTy = ConvertType(TemplateArgs[0].getAsType()); - - // Per JointMatrixINTEL spec the type can have an optional - // 'Component Type Interpretation' parameter. We should emit it in case - // if on SYCL level joint matrix accepts 'bfloat16' or 'tf32' objects as - // matrix's components. Yet 'bfloat16' should be represented as 'int16' and - // 'tf32' as 'float' types. - if (CompTy->isStructTy()) { - StringRef LlvmTyName = CompTy->getStructName(); - // Emit half/int16/float for sycl[::*]::{half,bfloat16,tf32} - if (LlvmTyName.starts_with("class.sycl::") || - LlvmTyName.starts_with("class.__sycl_internal::")) - LlvmTyName = LlvmTyName.rsplit("::").second; - if (LlvmTyName == "half") { - CompTy = llvm::Type::getHalfTy(getLLVMContext()); - return getJointMatrixINTELExtType(CompTy, TemplateArgs); - } else if (LlvmTyName == "tf32") { - CompTy = llvm::Type::getFloatTy(getLLVMContext()); - // 'tf32' interpretation is mapped to '0' - return getJointMatrixINTELExtType(CompTy, TemplateArgs, 0); - } else if (LlvmTyName == "bfloat16") { - CompTy = llvm::Type::getInt16Ty(getLLVMContext()); - // 'bfloat16' interpretation is mapped to '1' - return getJointMatrixINTELExtType(CompTy, TemplateArgs, 1); - } else { - llvm_unreachable("Wrong matrix base type!"); - } - } - return getJointMatrixINTELExtType(CompTy, TemplateArgs); -} - /// ConvertSPVCooperativeMatrixType - Convert SYCL joint_matrix type /// which is represented as a pointer to a structure to LLVM extension type /// with the parameters that follow SPIR-V CooperativeMatrixKHR type. @@ -732,10 +661,6 @@ llvm::Type *CodeGenTypes::ConvertType(QualType T) { if (ClangETy && ClangETy->isStructureOrClassType()) { RecordDecl *RD = ClangETy->getAsCXXRecordDecl(); if (RD && RD->getQualifiedNameAsString() == - "__spv::__spirv_JointMatrixINTEL") { - ResultType = ConvertSYCLJointMatrixINTELType(RD); - break; - } else if (RD && RD->getQualifiedNameAsString() == "__spv::__spirv_CooperativeMatrixKHR") { ResultType = ConvertSPVCooperativeMatrixType(RD); break; diff --git a/clang/lib/CodeGen/CodeGenTypes.h b/clang/lib/CodeGen/CodeGenTypes.h index cef2892676100..c05e9f963d8e4 100644 --- a/clang/lib/CodeGen/CodeGenTypes.h +++ b/clang/lib/CodeGen/CodeGenTypes.h @@ -145,14 +145,6 @@ class CodeGenTypes { /// load/store type are the same. llvm::Type *convertTypeForLoadStore(QualType T, llvm::Type *LLVMTy = nullptr); - /// ConvertSYCLJointMatrixINTELType - Convert SYCL joint_matrix type - /// which is represented as a pointer to a structure to LLVM extension type - /// with the parameters that follow SPIR-V JointMatrixINTEL type. - /// The expected representation is: - /// target("spirv.JointMatrixINTEL", %element_type, %rows%, %cols%, %scope%, - /// %use%, (optional) %element_type_interpretation%) - llvm::Type *ConvertSYCLJointMatrixINTELType(RecordDecl *RD); - /// ConvertSPVCooperativeMatrixType - Convert SYCL joint_matrix type /// which is represented as a pointer to a structure to LLVM extension type /// with the parameters that follow SPIR-V CooperativeMatrixKHR type. diff --git a/clang/test/CodeGenSYCL/joint_matrix.cpp b/clang/test/CodeGenSYCL/joint_matrix.cpp deleted file mode 100644 index 3ba63e89f8919..0000000000000 --- a/clang/test/CodeGenSYCL/joint_matrix.cpp +++ /dev/null @@ -1,41 +0,0 @@ -// RUN: %clang_cc1 -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s -// Test that SPIR-V codegen generates the expected LLVM struct name for the -// JointMatrixINTEL type. -#include -#include - -namespace __spv { - template - struct __spirv_JointMatrixINTEL; -} - -// CHECK: @_Z2f1{{.*}}(target("spirv.JointMatrixINTEL", float, 5, 10, 0, 1, 0) -void f1(__spv::__spirv_JointMatrixINTEL *matrix) {} - -// CHECK: @_Z2f2{{.*}}(target("spirv.JointMatrixINTEL", i64, 10, 2, 0, 0, 0) -void f2(__spv::__spirv_JointMatrixINTEL *matrix) {} - -// CHECK: @_Z2f3{{.*}}(target("spirv.JointMatrixINTEL", i8, 10, 2, 0, 0, 0) -void f3(__spv::__spirv_JointMatrixINTEL *matrix) {} - -namespace sycl { - class half {}; - class bfloat16 {}; - class tf32 {}; -} -typedef sycl::half my_half; - -// CHECK: @_Z2f4{{.*}}(target("spirv.JointMatrixINTEL", half, 10, 2, 0, 0, 0) -void f4(__spv::__spirv_JointMatrixINTEL *matrix) {} - -// CHECK: @_Z2f5{{.*}}(target("spirv.JointMatrixINTEL", i16, 10, 2, 0, 0, 0, 1) -void f5(__spv::__spirv_JointMatrixINTEL *matrix) {} - -// CHECK: @_Z2f6{{.*}}(target("spirv.JointMatrixINTEL", i128, 10, 2, 0, 0, 0) -void f6(__spv::__spirv_JointMatrixINTEL<_BitInt(128), 10, 2, 0, 0, 0> *matrix) {} - -// CHECK: @_Z2f7{{.*}}(target("spirv.JointMatrixINTEL", float, 10, 2, 0, 0, 0, 0) -void f7(__spv::__spirv_JointMatrixINTEL *matrix) {} - -// CHECK: @_Z2f8{{.*}}(target("spirv.JointMatrixINTEL", double, 5, 10, 0, 1, 0) -void f8(__spv::__spirv_JointMatrixINTEL *matrix) {} diff --git a/sycl/include/sycl/__spirv/spirv_ops.hpp b/sycl/include/sycl/__spirv/spirv_ops.hpp index 29cb4ffa6aa2f..b95cfa7c1d7f3 100644 --- a/sycl/include/sycl/__spirv/spirv_ops.hpp +++ b/sycl/include/sycl/__spirv/spirv_ops.hpp @@ -27,155 +27,6 @@ extern __DPCPP_SYCL_EXTERNAL float __spirv_RoundFToTF32INTEL(float a); -#ifndef __SPIRV_USE_COOPERATIVE_MATRIX -template -extern __DPCPP_SYCL_EXTERNAL - __spv::__spirv_JointMatrixINTEL * - __spirv_JointMatrixLoadINTEL(T *Ptr, std::size_t Stride, - __spv::MatrixLayout Layout = L, - __spv::Scope::Flag Sc = S, int MemOperand = 0); - -template -extern __DPCPP_SYCL_EXTERNAL void __spirv_JointMatrixStoreINTEL( - T *Ptr, __spv::__spirv_JointMatrixINTEL *Object, - std::size_t Stride, __spv::MatrixLayout Layout = L, - __spv::Scope::Flag Sc = S, int MemOperand = 0); - -template -extern __DPCPP_SYCL_EXTERNAL - __spv::__spirv_JointMatrixINTEL * - __spirv_CooperativeMatrixConstructCheckedINTEL(int32_t CoordX, - int32_t CoordY, - uint32_t Height, - uint32_t Width, - const T Value); - -template -extern __DPCPP_SYCL_EXTERNAL - __spv::__spirv_JointMatrixINTEL * - __spirv_CooperativeMatrixLoadCheckedINTEL( - T *Ptr, int32_t CoordX, int32_t CoordY, __spv::MatrixLayout Layout = L, - uint32_t Height = 0, uint32_t Width = 0, std::size_t Stride = 0, - int MemOperand = 0); - -template -extern __DPCPP_SYCL_EXTERNAL void __spirv_CooperativeMatrixStoreCheckedINTEL( - T *Ptr, int32_t CoordX, int32_t CoordY, - __spv::__spirv_JointMatrixINTEL *Object, - __spv::MatrixLayout Layout = L, uint32_t Height = 0, uint32_t Width = 0, - std::size_t Stride = 0, int MemOperand = 0); - -template -extern __DPCPP_SYCL_EXTERNAL - __spv::__spirv_JointMatrixINTEL * - __spirv_JointMatrixMadINTEL( - __spv::__spirv_JointMatrixINTEL *A, - __spv::__spirv_JointMatrixINTEL *B, - __spv::__spirv_JointMatrixINTEL *C, - __spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup); - -template -extern __DPCPP_SYCL_EXTERNAL - __spv::__spirv_JointMatrixINTEL * - __spirv_JointMatrixUUMadINTEL( - __spv::__spirv_JointMatrixINTEL *A, - __spv::__spirv_JointMatrixINTEL *B, - __spv::__spirv_JointMatrixINTEL *C, - __spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup); - -template -extern __DPCPP_SYCL_EXTERNAL - __spv::__spirv_JointMatrixINTEL * - __spirv_JointMatrixUSMadINTEL( - __spv::__spirv_JointMatrixINTEL *A, - __spv::__spirv_JointMatrixINTEL *B, - __spv::__spirv_JointMatrixINTEL *C, - __spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup); - -template -extern __DPCPP_SYCL_EXTERNAL - __spv::__spirv_JointMatrixINTEL * - __spirv_JointMatrixSUMadINTEL( - __spv::__spirv_JointMatrixINTEL *A, - __spv::__spirv_JointMatrixINTEL *B, - __spv::__spirv_JointMatrixINTEL *C, - __spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup); - -template -extern __DPCPP_SYCL_EXTERNAL - __spv::__spirv_JointMatrixINTEL * - __spirv_CompositeConstruct(const T v); - -template -extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t -__spirv_JointMatrixGetElementCoordINTEL( - __spv::__spirv_JointMatrixINTEL *, size_t i); - -template -extern __DPCPP_SYCL_EXTERNAL size_t __spirv_JointMatrixWorkItemLengthINTEL( - __spv::__spirv_JointMatrixINTEL *); - -template -extern __DPCPP_SYCL_EXTERNAL Ts __spirv_VectorExtractDynamic( - __spv::__spirv_JointMatrixINTEL *, size_t i); - -template -extern __DPCPP_SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL * -__spirv_VectorInsertDynamic(__spv::__spirv_JointMatrixINTEL *, - Ts val, size_t i); -#else // __SPIRV_USE_COOPERATIVE_MATRIX template *Object, __spv::MatrixLayout Layout = L, uint32_t Height = 0, uint32_t Width = 0, std::size_t Stride = 0, int MemOperand = 0); -#endif // __SPIRV_USE_COOPERATIVE_MATRIX template extern __DPCPP_SYCL_EXTERNAL void __spirv_CooperativeMatrixPrefetchINTEL( diff --git a/sycl/include/sycl/__spirv/spirv_types.hpp b/sycl/include/sycl/__spirv/spirv_types.hpp index 21563c3b8a67a..3a082b3fdf3aa 100644 --- a/sycl/include/sycl/__spirv/spirv_types.hpp +++ b/sycl/include/sycl/__spirv/spirv_types.hpp @@ -118,7 +118,6 @@ enum class MatrixLayout : uint32_t { enum class MatrixUse : uint32_t { MatrixA = 0, MatrixB = 1, Accumulator = 2 }; -#ifdef __SPIRV_USE_COOPERATIVE_MATRIX enum class MatrixOperands : uint32_t { // SPV_KHR_cooperative_matrix operands NoneKHR = 0, @@ -133,19 +132,10 @@ enum class MatrixOperands : uint32_t { MatrixCBFloat16ComponentsINTEL = 0x80, MatrixResultBFloat16ComponentsINTEL = 0x100 }; -#endif // __SPIRV_USE_COOPERATIVE_MATRIX -#ifndef __SPIRV_USE_COOPERATIVE_MATRIX - -template -struct __spirv_JointMatrixINTEL; -#else template struct __spirv_CooperativeMatrixKHR; -#endif // __SPIRV_USE_COOPERATIVE_MATRIX struct __spirv_TaskSequenceINTEL; diff --git a/sycl/include/sycl/ext/oneapi/matrix/matrix-intel.hpp b/sycl/include/sycl/ext/oneapi/matrix/matrix-intel.hpp index ebd2ad91de359..b8484a077c5fc 100644 --- a/sycl/include/sycl/ext/oneapi/matrix/matrix-intel.hpp +++ b/sycl/include/sycl/ext/oneapi/matrix/matrix-intel.hpp @@ -115,13 +115,8 @@ class wi_element { inline __SYCL_ALWAYS_INLINE std::tuple get_coord() { #if defined(__SYCL_DEVICE_ONLY__) -#ifndef __SPIRV_USE_COOPERATIVE_MATRIX __ocl_vec_t coord = __spirv_JointMatrixGetElementCoordINTEL(M.spvm, idx); -#else - __ocl_vec_t coord = - __spirv_JointMatrixGetElementCoordINTEL(M.spvm, idx); -#endif // __SPIRV_USE_COOPERATIVE_MATRIX const size_t row = coord[0]; const size_t col = coord[1]; return std::make_tuple(row, col); @@ -133,20 +128,11 @@ class wi_element { operator storage_element_type() { #ifdef __SYCL_DEVICE_ONLY__ -#ifndef __SPIRV_USE_COOPERATIVE_MATRIX - storage_element_type elem = - __spirv_VectorExtractDynamic::value, - spv_matrix_layout_traits::value, - spv_scope_traits::value>(M.spvm, - idx); -#else storage_element_type *ExtractP = __spirv_AccessChain::value, spv_scope_traits::value>(&M.spvm, idx); storage_element_type elem = *ExtractP; -#endif // __SPIRV_USE_COOPERATIVE_MATRIX return elem; #else throw exception(make_error_code(errc::runtime), @@ -156,20 +142,11 @@ class wi_element { explicit operator bool() { #ifdef __SYCL_DEVICE_ONLY__ -#ifndef __SPIRV_USE_COOPERATIVE_MATRIX - return __spirv_VectorExtractDynamic::value, - spv_matrix_layout_traits::value, - spv_scope_traits::value>( - M.spvm, idx) != static_cast(0); -#else storage_element_type *ExtractP = __spirv_AccessChain::value, spv_scope_traits::value>(&M.spvm, idx); return *ExtractP != static_cast(0); -#endif // __SPIRV_USE_COOPERATIVE_MATRIX #else throw exception(make_error_code(errc::runtime), "joint matrix is not supported on host."); @@ -178,16 +155,11 @@ class wi_element { template wi_element &operator=(const T2 &rhs) { #ifdef __SYCL_DEVICE_ONLY__ -#ifndef __SPIRV_USE_COOPERATIVE_MATRIX - M.spvm = __spirv_VectorInsertDynamic( - M.spvm, static_cast(rhs), idx); -#else storage_element_type *InsertP = __spirv_AccessChain::value, spv_scope_traits::value>(&M.spvm, idx); *InsertP = static_cast(rhs); -#endif // __SPIRV_USE_COOPERATIVE_MATRIX return *this; #else (void)rhs; @@ -199,16 +171,6 @@ class wi_element { wi_element & operator=(const wi_element &rhs) { #ifdef __SYCL_DEVICE_ONLY__ -#ifndef __SPIRV_USE_COOPERATIVE_MATRIX - M.spvm = __spirv_VectorInsertDynamic( - M.spvm, - __spirv_VectorExtractDynamic::value, - spv_matrix_layout_traits::value, - spv_scope_traits::value>(rhs.M.spvm, - rhs.idx), - idx); -#else storage_element_type *ExtractP = __spirv_AccessChain::value, @@ -219,7 +181,6 @@ class wi_element { spv_matrix_use_traits::value, spv_scope_traits::value>(&M.spvm, idx); *InsertP = *ExtractP; -#endif // __SPIRV_USE_COOPERATIVE_MATRIX return *this; #else (void)rhs; @@ -229,22 +190,6 @@ class wi_element { } #if __SYCL_DEVICE_ONLY__ -#ifndef __SPIRV_USE_COOPERATIVE_MATRIX -#define OP(op) \ - template wi_element &operator op##=(const T2 & rhs) { \ - M.spvm = __spirv_VectorInsertDynamic( \ - M.spvm, \ - static_cast( \ - __spirv_VectorExtractDynamic< \ - storage_element_type, T, NumRows, NumCols, \ - spv_matrix_use_traits::value, \ - spv_matrix_layout_traits::value, \ - spv_scope_traits::value>(M.spvm, idx) \ - op static_cast(rhs)), \ - idx); \ - return *this; \ - } -#else // __SPIRV_USE_COOPERATIVE_MATRIX #define OP(op) \ template wi_element &operator op##=(const T2 & rhs) { \ storage_element_type *ExtractP = \ @@ -259,7 +204,6 @@ class wi_element { *InsertP = *ExtractP op static_cast(rhs); \ return *this; \ } -#endif // __SPIRV_USE_COOPERATIVE_MATRIX #else // __SYCL_DEVICE_ONLY__ #define OP(op) \ template wi_element &operator op##=(const T2 & rhs) { \ @@ -294,13 +238,8 @@ class wi_element get_coord() { #if defined(__SYCL_DEVICE_ONLY__) -#ifndef __SPIRV_USE_COOPERATIVE_MATRIX - __ocl_vec_t coord = - __spirv_JointMatrixGetElementCoordINTEL(M.spvm, idx); -#else __ocl_vec_t coord = __spirv_JointMatrixGetElementCoordINTEL(M.spvm, idx); -#endif // __SPIRV_USE_COOPERATIVE_MATRIX const uint32_t row = coord[0]; const uint32_t col = coord[1]; return std::make_tuple(row, col); @@ -312,20 +251,12 @@ class wi_element::value, - spv_matrix_layout_traits::value, - spv_scope_traits::value>(M.spvm, idx); -#else sycl::ext::oneapi::bfloat16 *ExtractP = __spirv_AccessChain::value, spv_scope_traits::value>(&M.spvm, idx); return *ExtractP; -#endif // __SPIRV_USE_COOPERATIVE_MATRIX #else throw exception(make_error_code(errc::runtime), "joint matrix is not supported on host."); @@ -334,15 +265,6 @@ class wi_element( - __spirv_VectorExtractDynamic< - sycl::ext::oneapi::bfloat16, sycl::ext::oneapi::bfloat16, - NumRows, NumCols, spv_matrix_use_traits::value, - spv_matrix_layout_traits::value, - spv_scope_traits::value>(M.spvm, idx))) >= - std::numeric_limits::epsilon(); -#else sycl::ext::oneapi::bfloat16 *ExtractP = __spirv_AccessChain(Elem)) >= std::numeric_limits::epsilon(); -#endif // __SPIRV_USE_COOPERATIVE_MATRIX #else throw exception(make_error_code(errc::runtime), "joint matrix is not supported on host."); @@ -360,16 +281,12 @@ class wi_element::value, spv_scope_traits::value>(&M.spvm, idx); *InsertP = rhs; -#endif // __SPIRV_USE_COOPERATIVE_MATRIX return *this; #else (void)rhs; @@ -381,18 +298,6 @@ class wi_element &rhs) { #ifdef __SYCL_DEVICE_ONLY__ -#ifndef __SPIRV_USE_COOPERATIVE_MATRIX - M.spvm = __spirv_VectorInsertDynamic( - M.spvm, - __spirv_VectorExtractDynamic::value, - spv_matrix_layout_traits::value, - spv_scope_traits::value>(rhs.M.spvm, - rhs.idx), - idx); - return *this; -#else sycl::ext::oneapi::bfloat16 *ExtractP = __spirv_AccessChain::value>(&M.spvm, idx); *InsertP = *ExtractP; return *this; -#endif // __SPIRV_USE_COOPERATIVE_MATRIX #else (void)rhs; throw exception(make_error_code(errc::runtime), @@ -415,20 +319,6 @@ class wi_element::value, \ - spv_matrix_layout_traits::value, \ - spv_scope_traits::value>(M.spvm, idx) op rhs, \ - idx); \ - return *this; \ - } -#else #define OP(opassign, op) \ wi_element &operator opassign(const sycl::ext::oneapi::bfloat16 & rhs) { \ sycl::ext::oneapi::bfloat16 *ExtractP = \ @@ -444,7 +334,6 @@ class wi_element &lhs, \ - const sycl::ext::oneapi::bfloat16 &rhs) { \ - return __spirv_VectorExtractDynamic< \ - sycl::ext::oneapi::bfloat16, sycl::ext::oneapi::bfloat16, NumRows, \ - NumCols, spv_matrix_use_traits::value, \ - spv_matrix_layout_traits::value, \ - spv_scope_traits::value>(lhs.M.spvm, lhs.idx) op rhs; \ - } \ - friend type operator op( \ - const sycl::ext::oneapi::bfloat16 &lhs, \ - const wi_element &rhs) { \ - return __spirv_VectorExtractDynamic< \ - sycl::ext::oneapi::bfloat16, sycl::ext::oneapi::bfloat16, NumRows, \ - NumCols, spv_matrix_use_traits::value, \ - spv_matrix_layout_traits::value, \ - spv_scope_traits::value>(rhs.M.spvm, rhs.idx) op lhs; \ - } -#else #define OP(type, op) \ friend type operator op( \ const wi_element &lhs, \ - const sycl::ext::oneapi::bfloat16 &rhs) { \ - return type{static_cast( \ - __spirv_VectorExtractDynamic< \ - sycl::ext::oneapi::bfloat16, sycl::ext::oneapi::bfloat16, NumRows, \ - NumCols, spv_matrix_use_traits::value, \ - spv_matrix_layout_traits::value, \ - spv_scope_traits::value>(lhs.M.spvm, lhs.idx)) \ - op static_cast(rhs)}; \ - } \ - friend type operator op( \ - const sycl::ext::oneapi::bfloat16 &lhs, \ - const wi_element &rhs) { \ - return type{static_cast( \ - __spirv_VectorExtractDynamic< \ - sycl::ext::oneapi::bfloat16, sycl::ext::oneapi::bfloat16, NumRows, \ - NumCols, spv_matrix_use_traits::value, \ - spv_matrix_layout_traits::value, \ - spv_scope_traits::value>(rhs.M.spvm, rhs.idx)) \ - op static_cast(lhs)}; \ - } -#else #define OP(type, op) \ friend type operator op( \ const wi_element(*ExtractP) op static_cast(lhs)}; \ } -#endif // __SPIRV_USE_COOPERATIVE_MATRIX OP(bool, ==) OP(bool, !=) OP(bool, <) @@ -631,11 +468,7 @@ class wi_data { public: size_t length() { #if __SYCL_DEVICE_ONLY__ -#ifndef __SPIRV_USE_COOPERATIVE_MATRIX - return __spirv_JointMatrixWorkItemLengthINTEL(jm.spvm); -#else return __spirv_CooperativeMatrixLengthKHR(jm.spvm); -#endif // __SPIRV_USE_COOPERATIVE_MATRIX #else throw exception(make_error_code(errc::runtime), "joint matrix is not supported on host."); @@ -690,17 +523,6 @@ joint_matrix_store(Group, // intel's impl using DecorT = typename sycl::detail::DecoratedType::type; DecorT *Ptr = sycl::detail::getDecorated(dst); -#ifndef __SPIRV_USE_COOPERATIVE_MATRIX - __spirv_JointMatrixStoreINTEL::value, - sycl::ext::oneapi::experimental::matrix:: - spv_matrix_layout_traits::value>( - Ptr, src.spvm, stride, - sycl::ext::oneapi::experimental::matrix::spv_matrix_layout_traits< - Layout>::value, - sycl::ext::oneapi::experimental::matrix::spv_scope_traits::value); -#else __spirv_CooperativeMatrixStoreKHR< DecorT, Tp, NumRows, NumCols, sycl::ext::oneapi::experimental::matrix::spv_matrix_use_traits< @@ -711,7 +533,6 @@ joint_matrix_store(Group, sycl::ext::oneapi::experimental::matrix::spv_matrix_layout_traits< Layout>::value, stride); -#endif // __SPIRV_USE_COOPERATIVE_MATRIX #endif // defined(__NVPTX__) #else std::ignore = src; @@ -748,17 +569,6 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_store( #else // intel's impl T *Ptr = dst.get(); -#ifndef __SPIRV_USE_COOPERATIVE_MATRIX - __spirv_JointMatrixStoreINTEL::value, - sycl::ext::oneapi::experimental::matrix:: - spv_matrix_layout_traits::value>( - Ptr, src.spvm, stride, - sycl::ext::oneapi::experimental::matrix::spv_matrix_layout_traits< - Layout>::value, - sycl::ext::oneapi::experimental::matrix::spv_scope_traits::value); -#else __spirv_CooperativeMatrixStoreKHR< T, Tp, NumRows, NumCols, sycl::ext::oneapi::experimental::matrix::spv_matrix_use_traits< @@ -769,7 +579,6 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_store( sycl::ext::oneapi::experimental::matrix::spv_matrix_layout_traits< Layout>::value, stride); -#endif // __SPIRV_USE_COOPERATIVE_MATRIX #endif // defined(__NVPTX__) #else std::ignore = src; diff --git a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified-utils.hpp b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified-utils.hpp index 9d2759bdd3ad5..ec14cf6da1931 100644 --- a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified-utils.hpp +++ b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified-utils.hpp @@ -82,7 +82,6 @@ inline __SYCL_ALWAYS_INLINE __spv::MatrixLayout joint_matrix_layout_to_spv( } } -#ifdef __SPIRV_USE_COOPERATIVE_MATRIX template constexpr uint32_t CalculateMatrixOperand() { if constexpr (std::is_same::value && @@ -104,7 +103,6 @@ constexpr uint32_t CalculateMatrixOperand() { } return 0; } -#endif // __SPIRV_USE_COOPERATIVE_MATRIX } // namespace detail } // namespace _V1 diff --git a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp index c8d2918b6b105..2cf2eebc3bac5 100644 --- a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp +++ b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp @@ -50,15 +50,9 @@ struct joint_matrix { sycl::ext::oneapi::detail::joint_matrix_hip matrix_impl; #elif defined(__SPIR__) || defined(__SPIRV__) -#ifndef __SPIRV_USE_COOPERATIVE_MATRIX - __spv::__spirv_JointMatrixINTEL< - T, Rows, Cols, spv_matrix_layout_traits::value, - spv_scope_traits::value, spv_matrix_use_traits::value> *spvm; -#else __spv::__spirv_CooperativeMatrixKHR::value, Rows, Cols, spv_matrix_use_traits::value> *spvm; -#endif // __SPIRV_USE_COOPERATIVE_MATRIX #else static_assert(false, "The joint_matrix API is only supported by the Intel, " "CUDA and HIP (GFX90A) backends"); @@ -206,20 +200,11 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_load( std::ignore = sg; using DecorT = typename sycl::detail::DecoratedType::type; DecorT *Ptr = sycl::detail::getDecorated(src); -#ifndef __SPIRV_USE_COOPERATIVE_MATRIX - res.spvm = __spirv_JointMatrixLoadINTEL< - DecorT, S, NumRows, NumCols, - spv_matrix_use_traits::value, - spv_matrix_layout_traits::value>( - Ptr, stride, sycl::detail::joint_matrix_layout_to_spv(Layout), - spv_scope_traits::value); -#else res.spvm = __spirv_CooperativeMatrixLoadKHR< DecorT, S, NumRows, NumCols, spv_matrix_use_traits::value, spv_matrix_layout_traits::value>( Ptr, sycl::detail::joint_matrix_layout_to_spv(Layout), stride); -#endif // __SPIRV_USE_COOPERATIVE_MATRIX #endif // defined(__NVPTX__) #else std::ignore = sg; @@ -260,20 +245,11 @@ joint_matrix_load(Group sg, std::ignore = sg; using DecorT = typename sycl::detail::DecoratedType::type; DecorT *Ptr = sycl::detail::getDecorated(src); -#ifndef __SPIRV_USE_COOPERATIVE_MATRIX - res.spvm = - __spirv_JointMatrixLoadINTEL::value, - spv_matrix_layout_traits::value>( - Ptr, stride, spv_matrix_layout_traits::value, - spv_scope_traits::value); -#else res.spvm = __spirv_CooperativeMatrixLoadKHR::value, spv_matrix_layout_traits::value>( Ptr, spv_matrix_layout_traits::value, stride); -#endif // __SPIRV_USE_COOPERATIVE_MATRIX #endif // defined(__NVPTX__) #else std::ignore = sg; @@ -306,18 +282,10 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_load( #else std::ignore = sg; T *Ptr = src.get(); -#ifndef __SPIRV_USE_COOPERATIVE_MATRIX - res.spvm = __spirv_JointMatrixLoadINTEL< - T, S, NumRows, NumCols, spv_matrix_use_traits::value, - spv_matrix_layout_traits::value>( - Ptr, stride, sycl::detail::joint_matrix_layout_to_spv(Layout), - spv_scope_traits::value); -#else res.spvm = __spirv_CooperativeMatrixLoadKHR< T, S, NumRows, NumCols, spv_matrix_use_traits::value, spv_matrix_layout_traits::value>( Ptr, sycl::detail::joint_matrix_layout_to_spv(Layout), stride); -#endif // __SPIRV_USE_COOPERATIVE_MATRIX #endif // defined(__NVPTX__) #else std::ignore = sg; @@ -352,20 +320,11 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_load( #else std::ignore = sg; T *Ptr = src.get(); -#ifndef __SPIRV_USE_COOPERATIVE_MATRIX - res.spvm = - __spirv_JointMatrixLoadINTEL::value, - spv_matrix_layout_traits::value>( - Ptr, stride, spv_matrix_layout_traits::value, - spv_scope_traits::value); -#else res.spvm = __spirv_CooperativeMatrixLoadKHR::value, spv_matrix_layout_traits::value>( Ptr, spv_matrix_layout_traits::value, stride); -#endif // __SPIRV_USE_COOPERATIVE_MATRIX #endif // defined(__NVPTX__) #else std::ignore = sg; @@ -402,20 +361,11 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_store( std::ignore = sg; using DecorT = typename sycl::detail::DecoratedType::type; DecorT *Ptr = sycl::detail::getDecorated(dst); -#ifndef __SPIRV_USE_COOPERATIVE_MATRIX - __spirv_JointMatrixStoreINTEL< - DecorT, T, NumRows, NumCols, - spv_matrix_use_traits::value, - spv_matrix_layout_traits::value>( - Ptr, src.spvm, stride, sycl::detail::joint_matrix_layout_to_spv(Layout), - spv_scope_traits::value); -#else __spirv_CooperativeMatrixStoreKHR< DecorT, T, NumRows, NumCols, spv_matrix_use_traits::value, spv_matrix_layout_traits::value>( Ptr, src.spvm, sycl::detail::joint_matrix_layout_to_spv(Layout), stride); -#endif // __SPIRV_USE_COOPERATIVE_MATRIX #endif // defined(__NVPTX__) #else std::ignore = sg; @@ -448,18 +398,10 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_store( #else std::ignore = sg; T *Ptr = dst.get(); -#ifndef __SPIRV_USE_COOPERATIVE_MATRIX - __spirv_JointMatrixStoreINTEL< - T, T, NumRows, NumCols, spv_matrix_use_traits::value, - spv_matrix_layout_traits::value>( - Ptr, src.spvm, stride, sycl::detail::joint_matrix_layout_to_spv(Layout), - spv_scope_traits::value); -#else __spirv_CooperativeMatrixStoreKHR< T, T, NumRows, NumCols, spv_matrix_use_traits::value, spv_matrix_layout_traits::value>( Ptr, src.spvm, sycl::detail::joint_matrix_layout_to_spv(Layout), stride); -#endif // __SPIRV_USE_COOPERATIVE_MATRIX #endif // defined(__NVPTX__) #else std::ignore = sg; @@ -515,26 +457,11 @@ joint_matrix_mad( assert(false && "Ta != Tb : In the HIP backend joint_matrix_mad " "requires that joint_matrix data types Ta and Tb match"); } -#else -#ifndef __SPIRV_USE_COOPERATIVE_MATRIX - if constexpr (std::is_same::value && - std::is_same::value && - std::is_same::value) - D.spvm = __spirv_JointMatrixMadINTEL(A.spvm, B.spvm, C.spvm); - else if constexpr (std::is_unsigned::value && std::is_unsigned::value) - D.spvm = __spirv_JointMatrixUUMadINTEL(A.spvm, B.spvm, C.spvm); - else if constexpr (std::is_signed::value && std::is_unsigned::value) - D.spvm = __spirv_JointMatrixSUMadINTEL(A.spvm, B.spvm, C.spvm); - else if constexpr (std::is_unsigned::value && std::is_signed::value) - D.spvm = __spirv_JointMatrixUSMadINTEL(A.spvm, B.spvm, C.spvm); - else - D.spvm = __spirv_JointMatrixMadINTEL(A.spvm, B.spvm, C.spvm); #else constexpr uint32_t MatrixOperand = sycl::detail::CalculateMatrixOperand(); D.spvm = __spirv_CooperativeMatrixMulAddKHR(A.spvm, B.spvm, C.spvm, MatrixOperand); -#endif // __SPIRV_USE_COOPERATIVE_MATRIX #endif // defined(__NVPTX__) #else std::ignore = A; diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_abc.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_abc.cpp deleted file mode 100644 index 16e1951eab262..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_abc.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==----------- element_wise_abc.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 -// -//===----------------------------------------------------------------------===// -// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 -// UNSUPPORTED: gpu-intel-dg2 -// REQUIRES: aspect-ext_intel_matrix -// 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 - -#include "../../element_wise_abc_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops.cpp deleted file mode 100644 index 6c80692109ca8..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==------------ element_wise_all_ops.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: 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 - -// XFAIL: cpu - -#include "../../common.hpp" - -#define SG_SZ 32 - -#include "../../element_wise_all_ops_impl.hpp" 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 deleted file mode 100644 index 4acd374a74211..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_half.cpp +++ /dev/null @@ -1,25 +0,0 @@ -//==----------- element_wise_all_ops_half.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 -// -//===----------------------------------------------------------------------===// -// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 -// SYCL Joint Matrix fp16 operations are not supported on SPR -// UNSUPPORTED: gpu-intel-dg2, arch-intel_cpu_spr - -// REQUIRES: aspect-fp16 -// REQUIRES: aspect-ext_intel_matrix -// 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 - -#include "../../element_wise_all_ops_half_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_int8.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_int8.cpp deleted file mode 100644 index 4d7fdb9285023..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_int8.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==----------- element_wise_all_ops_int8.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 -// -//===----------------------------------------------------------------------===// -// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 -// UNSUPPORTED: gpu-intel-dg2 -// REQUIRES: aspect-ext_intel_matrix -// 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 - -#include "../../element_wise_all_ops_int8_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_int8_packed.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_int8_packed.cpp deleted file mode 100644 index 87ede89ab00c8..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_int8_packed.cpp +++ /dev/null @@ -1,24 +0,0 @@ -//==------ element_wise_all_ops_int8_packed.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 -// -//===----------------------------------------------------------------------===// -// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 -// UNSUPPORTED: gpu-intel-dg2 -// REQUIRES: aspect-ext_intel_matrix -// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// XFAIL: cpu - -// This test stores the matrix B that is VNNIed (packed). - -#include "../../common.hpp" - -#define SG_SZ 32 - -#include "../../element_wise_all_ops_int8_packed_impl.hpp" 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 deleted file mode 100644 index b9972d74ba18d..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_tf32.cpp +++ /dev/null @@ -1,19 +0,0 @@ -//==----------- element_wise_all_ops_tf32.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: aspect-ext_intel_matrix -// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -#include "../../common.hpp" - -#define SG_SZ 32 -constexpr size_t TN = 16; - -#include "../../element_wise_all_ops_tf32_impl.hpp" 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 deleted file mode 100644 index 0ff66f1ce7ed7..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_all_sizes.cpp +++ /dev/null @@ -1,23 +0,0 @@ -//==----------- element_wise_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: 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 - -// XFAIL: cpu - -#include "../../common.hpp" - -#define SG_SZ 32 - -#include "../../element_wise_all_sizes_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_ops.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_ops.cpp deleted file mode 100644 index 611c369b99011..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/element_wise_ops.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==----------- element_wise_ops.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: 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 - -// XFAIL: cpu - -#include "../../common.hpp" - -#define SG_SZ 32 - -#include "../../element_wise_ops_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/get_coord_float_matC.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/get_coord_float_matC.cpp deleted file mode 100644 index 8ef78f76b3509..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/get_coord_float_matC.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==----------- get_coord_float_matC.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 -// -//===----------------------------------------------------------------------===// -// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 -// UNSUPPORTED: gpu-intel-dg2 -// REQUIRES: aspect-ext_intel_matrix -// 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 - -#include "../../get_coord_float_matC_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/get_coord_int8_matA.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/get_coord_int8_matA.cpp deleted file mode 100644 index 9d3e62726720c..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/get_coord_int8_matA.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==----------- get_coord_int8_matA.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 -// -//===----------------------------------------------------------------------===// -// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 -// UNSUPPORTED: gpu-intel-dg2 -// REQUIRES: aspect-ext_intel_matrix -// 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 - -#include "../../get_coord_int8_matA_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/get_coord_int8_matB.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/get_coord_int8_matB.cpp deleted file mode 100644 index 0b6dac6047681..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/get_coord_int8_matB.cpp +++ /dev/null @@ -1,21 +0,0 @@ -//==----------- get_coord_int8_matB.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 -// -//===----------------------------------------------------------------------===// -// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 -// UNSUPPORTED: gpu-intel-dg2 -// REQUIRES: aspect-ext_intel_matrix -// 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 - -#include "../../get_coord_int8_matB_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_all_sizes.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_all_sizes.cpp deleted file mode 100644 index cddc7659167a2..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_all_sizes.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==-------- 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 -// -//===----------------------------------------------------------------------===// -// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 -// UNSUPPORTED: gpu-intel-dg2 -// REQUIRES: aspect-ext_intel_matrix -// 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 - -#include "../../joint_matrix_all_sizes_impl.hpp" 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 deleted file mode 100644 index 2e51654fd8c51..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_annotated_ptr.cpp +++ /dev/null @@ -1,24 +0,0 @@ -//==-------- joint_matrix_annotated_ptr.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: aspect-ext_intel_matrix -// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// XFAIL: cpu - -// Currently row major B fails when annotated_ptr is used -// XFAIL: gpu - -#include "../../common.hpp" - -#define SG_SZ 32 -constexpr size_t TN = 16; - -#include "../../joint_matrix_annotated_ptr_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_apply_bf16.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_apply_bf16.cpp deleted file mode 100644 index b5fb8bf2c6dfe..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_apply_bf16.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==----------- joint_matrix_apply_bf16.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 -// -//===----------------------------------------------------------------------===// -// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 -// UNSUPPORTED: gpu-intel-dg2 -// REQUIRES: aspect-ext_intel_matrix -// 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 - -#include "../../joint_matrix_apply_bf16_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_apply_two_matrices.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_apply_two_matrices.cpp deleted file mode 100644 index 84a2bc54d791f..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_apply_two_matrices.cpp +++ /dev/null @@ -1,23 +0,0 @@ -//==------ joint_matrix_apply_two_matrices.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: aspect-ext_intel_matrix - -// 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 %fp-model-precise -o %t.out -// RUN: %{run} %t.out - -// XFAIL: cpu -// XFAIL: gpu - -#include "../../common.hpp" - -#define SG_SZ 32 - -#include "../../joint_matrix_apply_two_matrices_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache.cpp deleted file mode 100644 index 8420948046337..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache.cpp +++ /dev/null @@ -1,27 +0,0 @@ -//==--- joint_matrix_bf16_fill_k_cache.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 -// -//===----------------------------------------------------------------------===// -// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 -// UNSUPPORTED: gpu-intel-dg2 -// REQUIRES: aspect-ext_intel_matrix -// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t_vnni.out -DVNNI %fp-model-precise -// RUN: %{run} %t_vnni.out - -// TODO: add row major compilation and run once Sub-group size 32 -// support becomes available in IGC for row major - -// XFAIL: cpu - -// -ffp-model=precise is added to not depend on compiler defaults. - -#include "../../common.hpp" - -#define SG_SZ 32 - -#include "../../joint_matrix_bf16_fill_k_cache_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_SLM.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_SLM.cpp deleted file mode 100644 index fb6eeed328995..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_SLM.cpp +++ /dev/null @@ -1,23 +0,0 @@ -//==--- joint_matrix_bf16_fill_k_cache_SLM.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 -// -//===----------------------------------------------------------------------===// -// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 -// UNSUPPORTED: gpu-intel-dg2 -// REQUIRES: aspect-ext_intel_matrix, gpu - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t_gpu_vnni.out %fp-model-precise -DSLM -DVNNI -// RUN: %{run} %t_gpu_vnni.out - -// TODO: add row major compilation and run once Sub-group size 32 -// support becomes available in IGC for row major - -// -ffp-model=precise is added to not depend on compiler defaults. - -#include "../../common.hpp" -#define SG_SZ 32 - -#include "../../joint_matrix_bf16_fill_k_cache_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_init.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_init.cpp deleted file mode 100644 index 6bcdcbcb79e17..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_init.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==---joint_matrix_bf16_fill_k_cache_init.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 -// -//===----------------------------------------------------------------------===// -// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 -// UNSUPPORTED: gpu-intel-dg2 -// REQUIRES: aspect-ext_intel_matrix, gpu -// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -DINIT_LIST -DVNNI %fp-model-precise -// RUN: %{run} %t.out - -// -ffp-model=precise is added to not depend on compiler defaults. - -#include "../../common.hpp" - -#define SG_SZ 32 - -#include "../../joint_matrix_bf16_fill_k_cache_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_unroll.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_unroll.cpp deleted file mode 100644 index c55158d5717e8..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_unroll.cpp +++ /dev/null @@ -1,26 +0,0 @@ -//==---joint_matrix_bf16_fill_k_cache_unroll.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 -// -//===----------------------------------------------------------------------===// -// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 -// UNSUPPORTED: gpu-intel-dg2 -// REQUIRES: aspect-ext_intel_matrix -// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -mllvm -inline-threshold=5000 %fp-model-precise -o %t.out -DMANUAL_UNROLL -DVNNI -// RUN: %{run} %t.out - -// XFAIL: cpu - -// -mllvm -inline-threshold added as a workaround, -// since IGC doesn't support some variants of IR for Joint Matrix currently -// -ffp-model=precise is added to not depend on compiler defaults. - -#include "../../common.hpp" - -#define SG_SZ 32 - -#include "../../joint_matrix_bf16_fill_k_cache_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_unroll_init.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_unroll_init.cpp deleted file mode 100644 index b428a505f0e90..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_unroll_init.cpp +++ /dev/null @@ -1,24 +0,0 @@ -//==--joint_matrix_bf16_fill_k_cache_unroll_init.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 -// -//===----------------------------------------------------------------------===// -// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 -// UNSUPPORTED: gpu-intel-dg2 -// REQUIRES: aspect-ext_intel_matrix, gpu -// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -mllvm -inline-threshold=5000 %fp-model-precise -o %t_gpu.out -DINIT_LIST -DMANUAL_UNROLL -DVNNI -// RUN: %{run} %t_gpu.out - -// -mllvm -inline-threshold added as a workaround, -// since IGC doesn't support some variants of IR for Joint Matrix currently -// -ffp-model=precise is added to not depend on compiler defaults. - -#include "../../common.hpp" - -#define SG_SZ 32 - -#include "../../joint_matrix_bf16_fill_k_cache_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16.cpp deleted file mode 100644 index 637e9ebcd858c..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==-------- joint_matrix_bfloat16.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 -// -//===----------------------------------------------------------------------===// -// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 -// UNSUPPORTED: gpu-intel-dg2 -// REQUIRES: aspect-ext_intel_matrix -// 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 - -#include "../../joint_matrix_bfloat16_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_array.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_array.cpp deleted file mode 100644 index eee85175d678d..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_array.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==-------- joint_matrix_bfloat16_array.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 -// -//===----------------------------------------------------------------------===// -// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 -// UNSUPPORTED: gpu-intel-dg2 -// REQUIRES: aspect-ext_intel_matrix -// 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 - -#include "../../joint_matrix_bfloat16_array_impl.hpp" 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 deleted file mode 100644 index 314b529ebfb7b..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_colmajorA_colmajorB.cpp +++ /dev/null @@ -1,31 +0,0 @@ -//==-- joint_matrix_bfloat16_colmajorA_colmajorB.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: aspect-ext_intel_matrix -// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// This tests support of col major layout for matrix B which does transpose and -// then VNNI transform. This is currently only available on AMX - -// XFAIL: gpu - -#include "../../common.hpp" -#include -#include -#include - -using namespace sycl; -using namespace sycl::ext::oneapi::experimental::matrix; -using bfloat16 = sycl::ext::oneapi::bfloat16; - -#define SG_SZ 32 -constexpr size_t TN = 16; - -#include "../../joint_matrix_bfloat16_colmajorA_colmajorB_impl.hpp" 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 deleted file mode 100644 index 60a3d2eb75aee..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_packedB.cpp +++ /dev/null @@ -1,23 +0,0 @@ -//==----- joint_matrix_bfloat16_packedB.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: 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 - -// XFAIL: gpu -// XFAIL: cpu - -#include "../../common.hpp" - -#define SG_SZ 32 -#include "../../joint_matrix_bfloat16_packedB_impl.hpp" 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 deleted file mode 100644 index 3ce16e94a40ba..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_colA_rowB_colC.cpp +++ /dev/null @@ -1,21 +0,0 @@ -//==---------- joint_matrix_colA_rowB_colC.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: aspect-ext_intel_matrix -// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// XFAIL:* - -#include "../../common.hpp" - -#define SG_SZ 32 -constexpr size_t TN = 16; - -#include "../../joint_matrix_colA_rowB_colC_impl.hpp" 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 deleted file mode 100644 index 73643d827a260..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_down_convert.cpp +++ /dev/null @@ -1,23 +0,0 @@ -//==-------- joint_matrix_down_convert.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: 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 - -// XFAIL: cpu - -#include "../../common.hpp" - -#define SG_SZ 32 - -#include "../../joint_matrix_down_convert_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_half.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_half.cpp deleted file mode 100644 index a0acb935699de..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_half.cpp +++ /dev/null @@ -1,25 +0,0 @@ -//==-------- joint_matrix_half.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 -// -//===----------------------------------------------------------------------===// -// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 -// SYCL Joint Matrix fp16 operations are not supported on SPR -// UNSUPPORTED: gpu-intel-dg2, arch-intel_cpu_spr - -// REQUIRES: aspect-fp16 -// REQUIRES: aspect-ext_intel_matrix -// 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 - -#include "../../joint_matrix_half_impl.hpp" 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 deleted file mode 100644 index a5874a5e90915..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_int8_colmajorA_colmajorB.cpp +++ /dev/null @@ -1,27 +0,0 @@ -//==----- joint_matrix_int8_colmajorA_colmajorB.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: aspect-ext_intel_matrix -// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// This tests support of col major layout for matrix B which does transpose and -// then VNNI transform. This is currently only available on AMX - -// XFAIL: gpu - -#include "../../common.hpp" - -using namespace sycl; -using namespace sycl::ext::oneapi::experimental::matrix; - -#define SG_SZ 32 -constexpr size_t TN = 16; - -#include "../../joint_matrix_int8_colmajorA_colmajorB_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_int8_rowmajorA_rowmajorB.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_int8_rowmajorA_rowmajorB.cpp deleted file mode 100644 index 654bdfe695116..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_int8_rowmajorA_rowmajorB.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==--- joint_matrix_int8_rowmajorA_rowmajorB.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 -// -//===----------------------------------------------------------------------===// -// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 -// UNSUPPORTED: gpu-intel-dg2 -// REQUIRES: aspect-ext_intel_matrix -// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// XFAIL: gpu - -#include "../../common.hpp" - -#define SG_SZ 32 - -#include "../../joint_matrix_int8_rowmajorA_rowmajorB_impl.hpp" 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 deleted file mode 100644 index f78eb4ef6c1f9..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_out_bounds.cpp +++ /dev/null @@ -1,24 +0,0 @@ -//==-------- joint_matrix_out_bounds.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: aspect-ext_intel_matrix -// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 - -// UNSUPPORTED: gpu-intel-dg2, cpu - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// XFAIL:gpu - -#include "../../common.hpp" - -#define SG_SZ 32 -constexpr size_t TN = 16; -constexpr size_t MATRIX_K = 1024 + 24; - -#include "../../joint_matrix_out_bounds_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_prefetch.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_prefetch.cpp deleted file mode 100644 index 7073feae64ac4..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_prefetch.cpp +++ /dev/null @@ -1,23 +0,0 @@ -//==-------- joint_matrix_prefetch.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-INTEL-DRIVER: lin: 27501, win: 101.4943 -// REQUIRES: aspect-ext_intel_matrix -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// XFAIL: cpu -// XFAIL: gpu - -// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 -// UNSUPPORTED: gpu-intel-dg2 - -#include "../../common.hpp" - -#define SG_SZ 32 -constexpr size_t TN = 16; -#include "../../joint_matrix_prefetch_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_rowmajorA_rowmajorB.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_rowmajorA_rowmajorB.cpp deleted file mode 100644 index 610b2b5bf6e5c..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_rowmajorA_rowmajorB.cpp +++ /dev/null @@ -1,29 +0,0 @@ -//==--------joint_matrix_rowmajorA_rowmajorB.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 -// -//===----------------------------------------------------------------------===// -// This tests support of row major layout for matrix B which does automatic VNNI -// REQUIRES: aspect-ext_intel_matrix -// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 -// VNNI transform and sub-group size 32 are not supported yet on DG2 by IGC -// UNSUPPORTED: gpu-intel-dg2 - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// XFAIL: cpu - -// Sub-group size 32 support for this test is not currently available in IGC -// XFAIL: gpu - -#include "../../common.hpp" - -using namespace sycl; -using namespace sycl::ext::oneapi::experimental::matrix; - -#define SG_SZ 32 - -#include "../../joint_matrix_rowmajorA_rowmajorB_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_ss_int8.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_ss_int8.cpp deleted file mode 100644 index fce90d5a42fa3..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_ss_int8.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==-------- joint_matrix_ss_int8.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 -// -//===----------------------------------------------------------------------===// -// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 -// UNSUPPORTED: gpu-intel-dg2 -// REQUIRES: aspect-ext_intel_matrix -// 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 - -#include "../../joint_matrix_ss_int8_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_su_int8.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_su_int8.cpp deleted file mode 100644 index c3d0302a3f187..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_su_int8.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==-------- joint_matrix_su_int8.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 -// -//===----------------------------------------------------------------------===// -// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 -// UNSUPPORTED: gpu-intel-dg2 -// REQUIRES: aspect-ext_intel_matrix -// 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 - -#include "../../joint_matrix_su_int8_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_tf32.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_tf32.cpp deleted file mode 100644 index 57fd883fc9a99..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_tf32.cpp +++ /dev/null @@ -1,19 +0,0 @@ -//==---------------- joint_matrix_tf32.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: aspect-ext_intel_matrix -// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -#include "../../common.hpp" - -#define SG_SZ 32 -constexpr size_t TN = 16; - -#include "../../joint_matrix_tf32_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_transposeC.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_transposeC.cpp deleted file mode 100644 index d0fd090b2e371..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_transposeC.cpp +++ /dev/null @@ -1,20 +0,0 @@ -//==----------- joint_matrix_transposeC.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 -// -//===----------------------------------------------------------------------===// -// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 -// UNSUPPORTED: gpu-intel-dg2 -// REQUIRES: aspect-ext_intel_matrix -// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -#include "../../common.hpp" - -#define SG_SZ 32 - -#include "../../joint_matrix_transposeC_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_unaligned_k.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_unaligned_k.cpp deleted file mode 100644 index d0ac32d7661e6..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_unaligned_k.cpp +++ /dev/null @@ -1,24 +0,0 @@ -//==-------- joint_matrix_unaligned_k.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: aspect-ext_intel_matrix -// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 - -// UNSUPPORTED: gpu-intel-dg2, cpu - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// XFAIL: gpu - -#include "../../common.hpp" - -#define SG_SZ 32 -constexpr size_t TN = 16; -static constexpr size_t MATRIX_K = 1024 + 14; - -#include "../../joint_matrix_out_bounds_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_us_int8.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_us_int8.cpp deleted file mode 100644 index 56ae3f112bb85..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_us_int8.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==-------- joint_matrix_us_int8.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 -// -//===----------------------------------------------------------------------===// -// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 -// UNSUPPORTED: gpu-intel-dg2 -// REQUIRES: aspect-ext_intel_matrix -// 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 - -#include "../../joint_matrix_us_int8_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_uu_int8.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_uu_int8.cpp deleted file mode 100644 index daf87d386d3c8..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/SG32/joint_matrix_uu_int8.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==-------- joint_matrix_uu_int8.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 -// -//===----------------------------------------------------------------------===// -// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 -// UNSUPPORTED: gpu-intel-dg2 -// REQUIRES: aspect-ext_intel_matrix -// 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 - -#include "../../joint_matrix_uu_int8_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_abc.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_abc.cpp deleted file mode 100644 index 9cc96a4396066..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_abc.cpp +++ /dev/null @@ -1,16 +0,0 @@ -//==----------- element_wise_abc.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: aspect-ext_intel_matrix - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// XFAIL: cpu - -#include "../common.hpp" -#include "../element_wise_abc_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops.cpp deleted file mode 100644 index edf0db4193add..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops.cpp +++ /dev/null @@ -1,16 +0,0 @@ -//==------------ element_wise_all_ops.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: aspect-ext_intel_matrix - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// XFAIL: cpu - -#include "../common.hpp" -#include "../element_wise_all_ops_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_1d.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_1d.cpp deleted file mode 100644 index 0058fb19af747..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_1d.cpp +++ /dev/null @@ -1,14 +0,0 @@ -//==-------- element_wise_all_ops_1d.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: aspect-ext_intel_matrix, gpu - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: env IGC_JointMatrixLoadStoreOpt=1 %{run} %t.out - -#include "../common.hpp" -#include "../element_wise_all_ops_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_1d_cont.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_1d_cont.cpp deleted file mode 100644 index 4a79a2832e424..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_1d_cont.cpp +++ /dev/null @@ -1,14 +0,0 @@ -//==-------- element_wise_all_ops_1d_cont.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: aspect-ext_intel_matrix, gpu - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: env IGC_JointMatrixLoadStoreOpt=2 %{run} %t.out - -#include "../common.hpp" -#include "../element_wise_all_ops_impl.hpp" 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 deleted file mode 100644 index f1873f848ac72..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_half.cpp +++ /dev/null @@ -1,20 +0,0 @@ -//==----------- element_wise_all_ops_half.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: aspect-fp16 -// REQUIRES: aspect-ext_intel_matrix - -// SYCL Joint Matrix fp16 operations are not supported on SPR -// UNSUPPORTED: arch-intel_cpu_spr - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// XFAIL: cpu - -#include "../common.hpp" -#include "../element_wise_all_ops_half_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_int8.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_int8.cpp deleted file mode 100644 index aa07934f6da8c..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_int8.cpp +++ /dev/null @@ -1,16 +0,0 @@ -//==----------- element_wise_all_ops_int8.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: aspect-ext_intel_matrix - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// XFAIL: cpu - -#include "../common.hpp" -#include "../element_wise_all_ops_int8_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_int8_packed.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_int8_packed.cpp deleted file mode 100644 index bd11f8a14313c..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_int8_packed.cpp +++ /dev/null @@ -1,18 +0,0 @@ -//==------ element_wise_all_ops_int8_packed.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: aspect-ext_intel_matrix - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// XFAIL: cpu - -// This test stores the matrix B that is VNNIed (packed). - -#include "../common.hpp" -#include "../element_wise_all_ops_int8_packed_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_scalar.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_scalar.cpp deleted file mode 100644 index 98a316f7e02d2..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_scalar.cpp +++ /dev/null @@ -1,16 +0,0 @@ -//==---- element_wise_all_ops_scalar.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: aspect-ext_intel_matrix, gpu - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: env IGC_JointMatrixLoadStoreOpt=0 %{run} %t.out - -// XFAIL: cpu - -#include "../common.hpp" -#include "../element_wise_all_ops_impl.hpp" 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 deleted file mode 100644 index 7a96dd76324f1..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_ops_tf32.cpp +++ /dev/null @@ -1,18 +0,0 @@ -//==----------- element_wise_all_ops_tf32.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: aspect-ext_intel_matrix -// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -#include "../common.hpp" - -constexpr size_t TN = 16; - -#include "../element_wise_all_ops_tf32_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_sizes.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_sizes.cpp deleted file mode 100644 index a5578c4caaded..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_all_sizes.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==----------- element_wise_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: aspect-ext_intel_matrix -// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// This is a version of the test with disabled device code -// split to test against fixed bug in IGC -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -fsycl-device-code-split=off -o %t_split.out -// RUN: %if gpu-intel-dg2 %{ %{run} %t_split.out %} - -// XFAIL: cpu - -#include "../common.hpp" -#include "../element_wise_all_sizes_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_ops.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_ops.cpp deleted file mode 100644 index 46a4171264063..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/element_wise_ops.cpp +++ /dev/null @@ -1,16 +0,0 @@ -//==----------- element_wise_ops.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: aspect-ext_intel_matrix - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// XFAIL: cpu - -#include "../common.hpp" -#include "../element_wise_ops_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/get_coord_float_matC.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/get_coord_float_matC.cpp deleted file mode 100644 index 29a97e665bc19..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/get_coord_float_matC.cpp +++ /dev/null @@ -1,16 +0,0 @@ -//==----------- get_coord_float_matC.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: aspect-ext_intel_matrix - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// XFAIL: cpu - -#include "../common.hpp" -#include "../get_coord_float_matC_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/get_coord_int8_matA.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/get_coord_int8_matA.cpp deleted file mode 100644 index 6b2c772267081..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/get_coord_int8_matA.cpp +++ /dev/null @@ -1,16 +0,0 @@ -//==----------- get_coord_int8_matA.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: aspect-ext_intel_matrix - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// XFAIL: cpu - -#include "../common.hpp" -#include "../get_coord_int8_matA_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/get_coord_int8_matB.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/get_coord_int8_matB.cpp deleted file mode 100644 index fdc80946352c3..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/get_coord_int8_matB.cpp +++ /dev/null @@ -1,15 +0,0 @@ -//==----------- get_coord_int8_matB.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: aspect-ext_intel_matrix - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out -// XFAIL: cpu - -#include "../common.hpp" -#include "../get_coord_int8_matB_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_all_sizes.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_all_sizes.cpp deleted file mode 100644 index 10d6d9ee62d56..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_all_sizes.cpp +++ /dev/null @@ -1,16 +0,0 @@ -//==-------- 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: aspect-ext_intel_matrix - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// XFAIL: cpu - -#include "../common.hpp" -#include "../joint_matrix_all_sizes_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_annotated_ptr.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_annotated_ptr.cpp deleted file mode 100644 index c4b748241a172..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_annotated_ptr.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==-------- joint_matrix_annotated_ptr.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: aspect-ext_intel_matrix -// UNSUPPORTED: gpu-intel-dg2 - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out -// RUN: %if gpu %{ env IGC_JointMatrixLoadStoreOpt=0 %{run} %t.out %} -// RUN: %if gpu %{ env IGC_JointMatrixLoadStoreOpt=1 %{run} %t.out %} - -// XFAIL: cpu - -#include "../common.hpp" - -constexpr size_t TN = 16; - -#include "../joint_matrix_annotated_ptr_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_apply_bf16.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_apply_bf16.cpp deleted file mode 100644 index ff4d2a6c716d6..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_apply_bf16.cpp +++ /dev/null @@ -1,16 +0,0 @@ -//==----------- joint_matrix_apply_bf16.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: aspect-ext_intel_matrix - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// XFAIL: cpu - -#include "../common.hpp" -#include "../joint_matrix_apply_bf16_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_apply_two_matrices.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_apply_two_matrices.cpp deleted file mode 100644 index 0a266a7792843..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_apply_two_matrices.cpp +++ /dev/null @@ -1,17 +0,0 @@ -//==------ joint_matrix_apply_two_matrices.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: aspect-ext_intel_matrix - -// RUN: %{build} %fp-model-precise -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// XFAIL: cpu -// XFAIL: gpu && !gpu-intel-dg2 - -#include "../common.hpp" -#include "../joint_matrix_apply_two_matrices_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache.cpp deleted file mode 100644 index 7a3631fb2f05b..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache.cpp +++ /dev/null @@ -1,21 +0,0 @@ -//==--- joint_matrix_bf16_fill_k_cache.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: aspect-ext_intel_matrix - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t_vnni.out -DVNNI %fp-model-precise -// RUN: %{run} %t_vnni.out - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out %fp-model-precise -// RUN: %{run} %t.out - -// XFAIL: cpu - -// -ffp-model=precise is added to not depend on compiler defaults. - -#include "../common.hpp" -#include "../joint_matrix_bf16_fill_k_cache_impl.hpp" 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 deleted file mode 100644 index e657f16ec2664..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_OOB.cpp +++ /dev/null @@ -1,20 +0,0 @@ -//==--- joint_matrix_bf16_fill_k_cache_OOB.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: aspect-ext_intel_matrix, gpu -// UNSUPPORTED: gpu-intel-dg2 - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t_gpu_vnni.out %fp-model-precise -DOOB -DVNNI -// RUN: %{run} %t_gpu_vnni.out - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t_gpu.out %fp-model-precise -DOOB -// RUN: %{run} %t_gpu.out - -// -ffp-model=precise is added to not depend on compiler defaults. - -#include "../common.hpp" -#include "../joint_matrix_bf16_fill_k_cache_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_SLM.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_SLM.cpp deleted file mode 100644 index c6c33cc3c4632..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_SLM.cpp +++ /dev/null @@ -1,19 +0,0 @@ -//==--- joint_matrix_bf16_fill_k_cache_SLM.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: aspect-ext_intel_matrix, gpu - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t_gpu_vnni.out %fp-model-precise -DSLM -DVNNI -// RUN: %{run} %t_gpu_vnni.out - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t_gpu.out %fp-model-precise -DSLM -// RUN: %{run} %t_gpu.out - -// -ffp-model=precise is added to not depend on compiler defaults. - -#include "../common.hpp" -#include "../joint_matrix_bf16_fill_k_cache_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_init.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_init.cpp deleted file mode 100644 index 60df4c0a6192e..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_init.cpp +++ /dev/null @@ -1,16 +0,0 @@ -//==---joint_matrix_bf16_fill_k_cache_init.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: aspect-ext_intel_matrix, gpu - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -DINIT_LIST -DVNNI %fp-model-precise -// RUN: %{run} %t.out - -// -ffp-model=precise is added to not depend on compiler defaults. - -#include "../common.hpp" -#include "../joint_matrix_bf16_fill_k_cache_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_unroll.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_unroll.cpp deleted file mode 100644 index ac4bbaec52169..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_unroll.cpp +++ /dev/null @@ -1,20 +0,0 @@ -//==---joint_matrix_bf16_fill_k_cache_unroll.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: aspect-ext_intel_matrix - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -mllvm -inline-threshold=2000 %fp-model-precise -o %t.out -DMANUAL_UNROLL -DVNNI -// RUN: %{run} %t.out - -// XFAIL: cpu - -// -mllvm -inline-threshold=2000 added as a workaround, -// since IGC doesn't support some variants of IR for Joint Matrix currently -// -ffp-model=precise is added to not depend on compiler defaults. - -#include "../common.hpp" -#include "../joint_matrix_bf16_fill_k_cache_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_unroll_init.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_unroll_init.cpp deleted file mode 100644 index ded7492812912..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_unroll_init.cpp +++ /dev/null @@ -1,18 +0,0 @@ -//==--joint_matrix_bf16_fill_k_cache_unroll_init.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: aspect-ext_intel_matrix, gpu - -// RUN: %{build} -mllvm -inline-threshold=2000 %fp-model-precise -o %t_gpu.out -DINIT_LIST -DMANUAL_UNROLL -DVNNI -// RUN: %{run} %t_gpu.out - -// -mllvm -inline-threshold=2000 added as a workaround, -// since IGC doesn't support some variants of IR for Joint Matrix currently -// -ffp-model=precise is added to not depend on compiler defaults. - -#include "../common.hpp" -#include "../joint_matrix_bf16_fill_k_cache_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16.cpp deleted file mode 100644 index c6f1db06e625f..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16.cpp +++ /dev/null @@ -1,16 +0,0 @@ -//==-------- joint_matrix_bfloat16.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: aspect-ext_intel_matrix - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// XFAIL: cpu - -#include "../common.hpp" -#include "../joint_matrix_bfloat16_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_array.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_array.cpp deleted file mode 100644 index 69f9aa8553bd3..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_array.cpp +++ /dev/null @@ -1,16 +0,0 @@ -//==-------- joint_matrix_bfloat16_array.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: aspect-ext_intel_matrix - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// XFAIL: cpu - -#include "../common.hpp" -#include "../joint_matrix_bfloat16_array_impl.hpp" 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 deleted file mode 100644 index 12b74948436e4..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_colmajorA_colmajorB.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==-- joint_matrix_bfloat16_colmajorA_colmajorB.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: aspect-ext_intel_matrix - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// This tests support of col major layout for matrix B which does transpose and -// then VNNI transform. This is currently only available on AMX - -// XFAIL: gpu - -#include "../common.hpp" - -constexpr size_t TN = 16; - -#include "../joint_matrix_bfloat16_colmajorA_colmajorB_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_packedB.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_packedB.cpp deleted file mode 100644 index 43c5d10739cfa..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_packedB.cpp +++ /dev/null @@ -1,20 +0,0 @@ -//==----- joint_matrix_bfloat16_packedB.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: aspect-ext_intel_matrix -// REQUIRES-INTEL-DRIVER: lin: 27868, win: 101.5181 - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out -// RUN: %if gpu %{ env IGC_JointMatrixLoadStoreOpt=2 %{run} %t.out %} -// RUN: %if gpu %{ env IGC_JointMatrixLoadStoreOpt=1 %{run} %t.out %} -// RUN: %if gpu %{ env IGC_JointMatrixLoadStoreOpt=0 %{run} %t.out %} - -// XFAIL: cpu - -#include "../common.hpp" -#include "../joint_matrix_bfloat16_packedB_impl.hpp" 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 deleted file mode 100644 index 083b7cbed8ef7..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_colA_rowB_colC.cpp +++ /dev/null @@ -1,19 +0,0 @@ -//==---------- joint_matrix_colA_rowB_colC.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: aspect-ext_intel_matrix - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// XFAIL:* - -#include "../common.hpp" - -constexpr size_t TN = 16; - -#include "../joint_matrix_colA_rowB_colC_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_down_convert.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_down_convert.cpp deleted file mode 100644 index 6ac1147369880..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_down_convert.cpp +++ /dev/null @@ -1,17 +0,0 @@ -//==-------- joint_matrix_down_convert.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: aspect-ext_intel_matrix -// UNSUPPORTED: gpu-intel-dg2 - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// XFAIL: cpu - -#include "../common.hpp" -#include "../joint_matrix_down_convert_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_half.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_half.cpp deleted file mode 100644 index 52a6f6165a355..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_half.cpp +++ /dev/null @@ -1,20 +0,0 @@ -//==-------- joint_matrix_half.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: aspect-fp16 -// REQUIRES: aspect-ext_intel_matrix - -// SYCL Joint Matrix fp16 operations are not supported on SPR -// UNSUPPORTED: arch-intel_cpu_spr - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// XFAIL: cpu - -#include "../common.hpp" -#include "../joint_matrix_half_impl.hpp" 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 deleted file mode 100644 index cf462e53c86cb..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_int8_colmajorA_colmajorB.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==----- joint_matrix_int8_colmajorA_colmajorB.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: aspect-ext_intel_matrix - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// This tests support of col major layout for matrix B which does transpose and -// then VNNI transform. This is currently only available on AMX - -// XFAIL: gpu - -#include "../common.hpp" - -constexpr size_t TN = 16; - -#include "../joint_matrix_int8_colmajorA_colmajorB_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_int8_rowmajorA_rowmajorB.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_int8_rowmajorA_rowmajorB.cpp deleted file mode 100644 index 31bc890fdd197..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_int8_rowmajorA_rowmajorB.cpp +++ /dev/null @@ -1,19 +0,0 @@ -//==----- joint_matrix_int8_rowmajorA_rowmajorB.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: aspect-ext_intel_matrix - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// Run these 2 tests on PVC only for now. Check can be updated to "gpu", -// when newer IGC is used in intel/llvm pre-checkin testing on Intel Arc -// RUN: %if arch-intel_gpu_pvc %{ env IGC_JointMatrixLoadStoreOpt=0 %{run} %t.out %} -// RUN: %if arch-intel_gpu_pvc %{ env IGC_JointMatrixLoadStoreOpt=1 %{run} %t.out %} - -#include "../common.hpp" -#include "../joint_matrix_int8_rowmajorA_rowmajorB_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_opt_kernel_feature.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_opt_kernel_feature.cpp deleted file mode 100644 index 2ef42e3b499e9..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_opt_kernel_feature.cpp +++ /dev/null @@ -1,18 +0,0 @@ -//===---joint_matrix_opt_kernel_feature.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: aspect-ext_intel_matrix - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// Test checks that exception will be thrown in case matrix parameters are -// incompatible on the current device - -#include "../common.hpp" -#include "../joint_matrix_opt_kernel_feature_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_out_bounds.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_out_bounds.cpp deleted file mode 100644 index 138f6738155f9..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_out_bounds.cpp +++ /dev/null @@ -1,20 +0,0 @@ -//==-------- joint_matrix_out_bounds.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: aspect-ext_intel_matrix - -// UNSUPPORTED: gpu-intel-dg2, cpu -// -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -#include "../common.hpp" - -constexpr size_t TN = 16; -constexpr size_t MATRIX_K = 1024 + 24; - -#include "../joint_matrix_out_bounds_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_prefetch.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_prefetch.cpp deleted file mode 100644 index daeadb6c8f658..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_prefetch.cpp +++ /dev/null @@ -1,18 +0,0 @@ -//==-------- joint_matrix_prefetch.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: aspect-ext_intel_matrix -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// XFAIL: cpu -// XFAIL-TRACKER: CMPLRLLVM-62790 - -#include "../common.hpp" - -constexpr size_t TN = 16; -#include "../joint_matrix_prefetch_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_rowmajorA_rowmajorB.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_rowmajorA_rowmajorB.cpp deleted file mode 100644 index 6d213ba4ed870..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_rowmajorA_rowmajorB.cpp +++ /dev/null @@ -1,24 +0,0 @@ -//==-------joint_matrix_rowmajorA_rowmajorB.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: aspect-ext_intel_matrix -// VNNI transform is not supported yet by IGC on DG2 -// UNSUPPORTED: gpu-intel-dg2 - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out -// RUN: %if gpu %{ env IGC_JointMatrixLoadStoreOpt=2 %{run} %t.out %} -// RUN: %if gpu %{ env IGC_JointMatrixLoadStoreOpt=1 %{run} %t.out %} -// RUN: %if gpu %{ env IGC_JointMatrixLoadStoreOpt=0 %{run} %t.out %} - -// XFAIL: cpu - -// This tests support of row major layout for matrix B which does automatic VNNI -// transform. This is currently only available on AMX and XMX of PVC - -#include "../common.hpp" -#include "../joint_matrix_rowmajorA_rowmajorB_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_ss_int8.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_ss_int8.cpp deleted file mode 100644 index d43e10ffed568..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_ss_int8.cpp +++ /dev/null @@ -1,16 +0,0 @@ -//==-------- joint_matrix_ss_int8.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: aspect-ext_intel_matrix - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// XFAIL: cpu - -#include "../common.hpp" -#include "../joint_matrix_ss_int8_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_su_int8.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_su_int8.cpp deleted file mode 100644 index 1432abf115508..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_su_int8.cpp +++ /dev/null @@ -1,16 +0,0 @@ -//==-------- joint_matrix_su_int8.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: aspect-ext_intel_matrix - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// XFAIL: cpu - -#include "../common.hpp" -#include "../joint_matrix_su_int8_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_tf32.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_tf32.cpp deleted file mode 100644 index 322a80339de77..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_tf32.cpp +++ /dev/null @@ -1,18 +0,0 @@ -//==---------------- joint_matrix_tf32.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: aspect-ext_intel_matrix -// REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943 - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -#include "../common.hpp" - -constexpr size_t TN = 16; - -#include "../joint_matrix_tf32_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_transposeC.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_transposeC.cpp deleted file mode 100644 index 2a2ee8eccaf66..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_transposeC.cpp +++ /dev/null @@ -1,14 +0,0 @@ -//==----------- joint_matrix_transposeC.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: aspect-ext_intel_matrix - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -#include "../common.hpp" -#include "../joint_matrix_transposeC_impl.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 deleted file mode 100644 index 0f6d6fad9d673..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_unaligned_k.cpp +++ /dev/null @@ -1,20 +0,0 @@ -//==-------- joint_matrix_unaligned_k.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: aspect-ext_intel_matrix - -// UNSUPPORTED: gpu-intel-dg2, cpu - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -#include "../common.hpp" - -constexpr size_t TN = 16; -static constexpr size_t MATRIX_K = 1024 + 14; - -#include "../joint_matrix_out_bounds_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_us_int8.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_us_int8.cpp deleted file mode 100644 index 93050559c1dfe..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_us_int8.cpp +++ /dev/null @@ -1,16 +0,0 @@ -//==-------- joint_matrix_us_int8.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: aspect-ext_intel_matrix - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// XFAIL: cpu - -#include "../common.hpp" -#include "../joint_matrix_us_int8_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_uu_int8.cpp b/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_uu_int8.cpp deleted file mode 100644 index 2b988bdf92bcc..0000000000000 --- a/sycl/test-e2e/Matrix/SPVCooperativeMatrix/joint_matrix_uu_int8.cpp +++ /dev/null @@ -1,16 +0,0 @@ -//==-------- joint_matrix_uu_int8.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: aspect-ext_intel_matrix - -// RUN: %{build} -D__SPIRV_USE_COOPERATIVE_MATRIX -o %t.out -// RUN: %{run} %t.out - -// XFAIL: cpu - -#include "../common.hpp" -#include "../joint_matrix_uu_int8_impl.hpp" diff --git a/sycl/test/check_device_code/matrix/matrix-int8-test.cpp b/sycl/test/check_device_code/matrix/matrix-int8-test.cpp index ee56a0c73fe61..690450238f2a9 100644 --- a/sycl/test/check_device_code/matrix/matrix-int8-test.cpp +++ b/sycl/test/check_device_code/matrix/matrix-int8-test.cpp @@ -1,8 +1,8 @@ // RUN: %clangxx -fsycl -fsycl-device-only -O2 -S -emit-llvm -o - %s | FileCheck %s -// CHECK-DAG: target("spirv.JointMatrixINTEL", i8, 12, 48, 0, 3, 0) -// CHECK-DAG: target("spirv.JointMatrixINTEL", i32, 12, 12, 3, 3, 2) -// CHECK-DAG: target("spirv.JointMatrixINTEL", i8, 48, 12, 2, 3, 1) +// CHECK-DAG: target("spirv.CooperativeMatrixKHR", i8, 3, 12, 48, 0) +// CHECK-DAG: target("spirv.CooperativeMatrixKHR", i32, 3, 12, 12, 2) +// CHECK-DAG: target("spirv.CooperativeMatrixKHR", i8, 3, 48, 12, 1) // CHECK: !{!"matrix_type::sint32,use::accumulator,12,12;matrix_type::sint8,use::a,12,48;matrix_type::sint8,use::b,48,12"} // CHECK: !{!"matrix_type::sint8,matrix_type::sint8,matrix_type::sint32,matrix_type::sint32,12,48,12"} diff --git a/sycl/test/check_device_code/matrix/matrix_load_store_as.cpp b/sycl/test/check_device_code/matrix/matrix_load_store_as.cpp index 771235690ac8d..c35cdc0bd9186 100644 --- a/sycl/test/check_device_code/matrix/matrix_load_store_as.cpp +++ b/sycl/test/check_device_code/matrix/matrix_load_store_as.cpp @@ -1,7 +1,7 @@ // RUN: %clangxx -fsycl-device-only -S -emit-llvm -o - %s | FileCheck %s // Check that SROA and mem2reg won't leave alloca of matrix type in IR -// CHECK-NOT: alloca target("spirv.JointMatrixINTEL" +// CHECK-NOT: alloca target("spirv.CooperativeMatrixKHR" // check that correct address spaces are used to load from and store to #include @@ -27,14 +27,14 @@ SYCL_EXTERNAL [[sycl::reqd_sub_group_size(16)]] void matrix_store_as( it.barrier(access::fence_space::local_space); // A should load from local address space - // CHECK: %{{.*}} = tail call spir_func noundef target("spirv.JointMatrixINTEL", i16, 8, 16, 0, 3, 0) @_Z[[#]]__spirv_JointMatrixLoadINTEL{{.*}}(ptr addrspace(3) noundef %{{.*}}, i64 noundef 16, i32 noundef 0, i32 noundef 3, i32 noundef 0) #{{.*}} + // CHECK: %{{.*}} = tail call spir_func noundef target("spirv.CooperativeMatrixKHR", i16, 3, 8, 16, 0) @_Z[[#]]__spirv_CooperativeMatrixLoadKHR{{.*}}(ptr addrspace(3) noundef %{{.*}}, i32 noundef 0, i64 noundef 16, i32 noundef 0) #{{.*}} joint_matrix_load( sg, tA, tileA.template get_multi_ptr(), 16); // B should load from global address space - // CHECK: %{{.*}} = tail call spir_func noundef target("spirv.JointMatrixINTEL", i16, 16, 16, 2, 3, 1) @_Z[[#]]__spirv_JointMatrixLoadINTEL{{.*}}(ptr addrspace(1) noundef %{{.*}}, i64 noundef 32, i32 noundef 2, i32 noundef 3, i32 noundef 0) #{{.*}} + // CHECK: %{{.*}} = tail call spir_func noundef target("spirv.CooperativeMatrixKHR", i16, 3, 16, 16, 1) @_Z[[#]]__spirv_CooperativeMatrixLoadKHR{{.*}}(ptr addrspace(1) noundef %{{.*}}, i32 noundef 2, i64 noundef 32, i32 noundef 0) #{{.*}} joint_matrix_load(sg, tB, pB, 32); joint_matrix_mad(sg, tC, tA, tB, tC); // C should store to global address space - // CHECK: tail call spir_func void @_Z[[#]]__spirv_JointMatrixStoreINTEL{{.*}}(ptr addrspace(1) noundef %{{.*}}, target("spirv.JointMatrixINTEL", float, 8, 16, 3, 3, 2) noundef %{{.*}}, i64 noundef 16, i32 noundef 0, i32 noundef 3, i32 noundef 0) #{{.*}} + // CHECK: tail call spir_func void @_Z[[#]]__spirv_CooperativeMatrixStoreKHR{{.*}}(ptr addrspace(1) noundef %{{.*}}, target("spirv.CooperativeMatrixKHR", float, 3, 8, 16, 2) noundef %{{.*}}, i32 noundef 0, i64 noundef 16, i32 noundef 0) #{{.*}} joint_matrix_store(sg, tC, pC, 16, layout::row_major); } diff --git a/sycl/test/e2e_test_requirements/no-unsupported-without-info.cpp b/sycl/test/e2e_test_requirements/no-unsupported-without-info.cpp index 3602d7d01ae95..0201373566102 100644 --- a/sycl/test/e2e_test_requirements/no-unsupported-without-info.cpp +++ b/sycl/test/e2e_test_requirements/no-unsupported-without-info.cpp @@ -302,47 +302,6 @@ // CHECK-NEXT: Matrix/SG32/joint_matrix_unaligned_k.cpp // CHECK-NEXT: Matrix/SG32/joint_matrix_us_int8.cpp // CHECK-NEXT: Matrix/SG32/joint_matrix_uu_int8.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/element_wise_abc.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_half.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_int8.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_int8_packed.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/element_wise_all_sizes.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/element_wise_ops.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/get_coord_float_matC.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/get_coord_int8_matA.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/get_coord_int8_matB.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_all_sizes.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_apply_bf16.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_apply_two_matrices.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_SLM.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_init.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_unroll.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_unroll_init.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_array.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_packedB.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_down_convert.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_half.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_int8_rowmajorA_rowmajorB.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_out_bounds.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_prefetch.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_rowmajorA_rowmajorB.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_ss_int8.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_su_int8.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_transposeC.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_unaligned_k.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_us_int8.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_uu_int8.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/element_wise_all_ops_half.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/joint_matrix_annotated_ptr.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_OOB.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/joint_matrix_down_convert.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/joint_matrix_half.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/joint_matrix_out_bounds.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/joint_matrix_rowmajorA_rowmajorB.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/joint_matrix_unaligned_k.cpp // CHECK-NEXT: Matrix/joint_matrix_annotated_ptr.cpp // CHECK-NEXT: Matrix/joint_matrix_bf16_fill_k_cache_OOB.cpp // CHECK-NEXT: Matrix/joint_matrix_bf16_fill_k_cache_prefetch.cpp diff --git a/sycl/test/e2e_test_requirements/no-xfail-without-tracker.cpp b/sycl/test/e2e_test_requirements/no-xfail-without-tracker.cpp index 7274f59992c72..82cf0ceb75ab8 100644 --- a/sycl/test/e2e_test_requirements/no-xfail-without-tracker.cpp +++ b/sycl/test/e2e_test_requirements/no-xfail-without-tracker.cpp @@ -96,71 +96,6 @@ // CHECK-NEXT: Matrix/SG32/joint_matrix_prefetch.cpp // CHECK-NEXT: Matrix/SG32/joint_matrix_rowmajorA_rowmajorB.cpp // CHECK-NEXT: Matrix/SG32/joint_matrix_unaligned_k.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/element_wise_abc.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_half.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_int8.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/element_wise_all_ops_int8_packed.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/element_wise_all_sizes.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/element_wise_ops.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/get_coord_float_matC.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/get_coord_int8_matA.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/get_coord_int8_matB.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_all_sizes.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_annotated_ptr.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_annotated_ptr.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_apply_bf16.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_apply_two_matrices.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bf16_fill_k_cache_unroll.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_array.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_colmajorA_colmajorB.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_bfloat16_packedB.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_colA_rowB_colC.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_down_convert.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_half.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_int8_colmajorA_colmajorB.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_int8_rowmajorA_rowmajorB.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_out_bounds.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_prefetch.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_rowmajorA_rowmajorB.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_rowmajorA_rowmajorB.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_ss_int8.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_su_int8.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_unaligned_k.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_us_int8.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/SG32/joint_matrix_uu_int8.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/element_wise_abc.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/element_wise_all_ops.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/element_wise_all_ops_half.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/element_wise_all_ops_int8.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/element_wise_all_ops_int8_packed.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/element_wise_all_ops_scalar.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/element_wise_all_sizes.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/element_wise_ops.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/get_coord_float_matC.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/get_coord_int8_matA.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/get_coord_int8_matB.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/joint_matrix_all_sizes.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/joint_matrix_annotated_ptr.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/joint_matrix_apply_bf16.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/joint_matrix_apply_two_matrices.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/joint_matrix_bf16_fill_k_cache_unroll.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_array.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_colmajorA_colmajorB.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/joint_matrix_bfloat16_packedB.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/joint_matrix_colA_rowB_colC.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/joint_matrix_down_convert.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/joint_matrix_half.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/joint_matrix_int8_colmajorA_colmajorB.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/joint_matrix_rowmajorA_rowmajorB.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/joint_matrix_ss_int8.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/joint_matrix_su_int8.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/joint_matrix_us_int8.cpp -// CHECK-NEXT: Matrix/SPVCooperativeMatrix/joint_matrix_uu_int8.cpp // CHECK-NEXT: Matrix/joint_matrix_bfloat16_colmajorA_colmajorB.cpp // CHECK-NEXT: Matrix/joint_matrix_colA_rowB_colC.cpp // CHECK-NEXT: Matrix/joint_matrix_int8_colmajorA_colmajorB.cpp From d3866012568acbe24a3b4c332f165df2eec7af6a Mon Sep 17 00:00:00 2001 From: "Sidorov, Dmitry" Date: Mon, 11 Nov 2024 14:36:38 -0800 Subject: [PATCH 2/2] format and test Signed-off-by: Sidorov, Dmitry --- clang/lib/CodeGen/CodeGenTypes.cpp | 2 +- sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp | 4 ++-- .../e2e_test_requirements/no-unsupported-without-info.cpp | 2 +- sycl/test/e2e_test_requirements/no-xfail-without-tracker.cpp | 2 +- 4 files changed, 5 insertions(+), 5 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenTypes.cpp b/clang/lib/CodeGen/CodeGenTypes.cpp index c0123286da9e9..7da27900ac498 100644 --- a/clang/lib/CodeGen/CodeGenTypes.cpp +++ b/clang/lib/CodeGen/CodeGenTypes.cpp @@ -661,7 +661,7 @@ llvm::Type *CodeGenTypes::ConvertType(QualType T) { if (ClangETy && ClangETy->isStructureOrClassType()) { RecordDecl *RD = ClangETy->getAsCXXRecordDecl(); if (RD && RD->getQualifiedNameAsString() == - "__spv::__spirv_CooperativeMatrixKHR") { + "__spv::__spirv_CooperativeMatrixKHR") { ResultType = ConvertSPVCooperativeMatrixType(RD); break; } else if (RD && RD->getQualifiedNameAsString() == diff --git a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp index 2cf2eebc3bac5..18404f99184a9 100644 --- a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp +++ b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp @@ -460,8 +460,8 @@ joint_matrix_mad( #else constexpr uint32_t MatrixOperand = sycl::detail::CalculateMatrixOperand(); - D.spvm = __spirv_CooperativeMatrixMulAddKHR(A.spvm, B.spvm, C.spvm, - MatrixOperand); + D.spvm = + __spirv_CooperativeMatrixMulAddKHR(A.spvm, B.spvm, C.spvm, MatrixOperand); #endif // defined(__NVPTX__) #else std::ignore = A; diff --git a/sycl/test/e2e_test_requirements/no-unsupported-without-info.cpp b/sycl/test/e2e_test_requirements/no-unsupported-without-info.cpp index 0201373566102..6b7337b3b0769 100644 --- a/sycl/test/e2e_test_requirements/no-unsupported-without-info.cpp +++ b/sycl/test/e2e_test_requirements/no-unsupported-without-info.cpp @@ -54,7 +54,7 @@ // tests to match the required format and in that case you should just update // (i.e. reduce) the number and the list below. // -// NUMBER-OF-UNSUPPORTED-WITHOUT-INFO: 477 +// NUMBER-OF-UNSUPPORTED-WITHOUT-INFO: 436 // // List of improperly UNSUPPORTED tests. // Remove the CHECK once the test has been properly UNSUPPORTED. diff --git a/sycl/test/e2e_test_requirements/no-xfail-without-tracker.cpp b/sycl/test/e2e_test_requirements/no-xfail-without-tracker.cpp index 82cf0ceb75ab8..c867ce1f4f420 100644 --- a/sycl/test/e2e_test_requirements/no-xfail-without-tracker.cpp +++ b/sycl/test/e2e_test_requirements/no-xfail-without-tracker.cpp @@ -51,7 +51,7 @@ // tests to match the required format and in that case you should just update // (i.e. reduce) the number and the list below. // -// NUMBER-OF-XFAIL-WITHOUT-TRACKER: 142 +// NUMBER-OF-XFAIL-WITHOUT-TRACKER: 77 // // List of improperly XFAIL-ed tests. // Remove the CHECK once the test has been properly XFAIL-ed.