diff --git a/sycl/test-e2e/Matrix/common.hpp b/sycl/test-e2e/Matrix/Inputs/common.hpp similarity index 100% rename from sycl/test-e2e/Matrix/common.hpp rename to sycl/test-e2e/Matrix/Inputs/common.hpp diff --git a/sycl/test-e2e/Matrix/element_wise_abc_impl.hpp b/sycl/test-e2e/Matrix/Inputs/element_wise_abc_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/element_wise_abc_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/element_wise_abc_impl.hpp diff --git a/sycl/test-e2e/Matrix/element_wise_all_ops_half_impl.hpp b/sycl/test-e2e/Matrix/Inputs/element_wise_all_ops_half_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/element_wise_all_ops_half_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/element_wise_all_ops_half_impl.hpp diff --git a/sycl/test-e2e/Matrix/element_wise_all_ops_impl.hpp b/sycl/test-e2e/Matrix/Inputs/element_wise_all_ops_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/element_wise_all_ops_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/element_wise_all_ops_impl.hpp diff --git a/sycl/test-e2e/Matrix/element_wise_all_ops_int8_impl.hpp b/sycl/test-e2e/Matrix/Inputs/element_wise_all_ops_int8_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/element_wise_all_ops_int8_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/element_wise_all_ops_int8_impl.hpp diff --git a/sycl/test-e2e/Matrix/element_wise_all_ops_int8_packed_impl.hpp b/sycl/test-e2e/Matrix/Inputs/element_wise_all_ops_int8_packed_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/element_wise_all_ops_int8_packed_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/element_wise_all_ops_int8_packed_impl.hpp diff --git a/sycl/test-e2e/Matrix/element_wise_all_ops_tf32_impl.hpp b/sycl/test-e2e/Matrix/Inputs/element_wise_all_ops_tf32_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/element_wise_all_ops_tf32_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/element_wise_all_ops_tf32_impl.hpp diff --git a/sycl/test-e2e/Matrix/element_wise_all_sizes_impl.hpp b/sycl/test-e2e/Matrix/Inputs/element_wise_all_sizes_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/element_wise_all_sizes_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/element_wise_all_sizes_impl.hpp diff --git a/sycl/test-e2e/Matrix/element_wise_ops_impl.hpp b/sycl/test-e2e/Matrix/Inputs/element_wise_ops_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/element_wise_ops_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/element_wise_ops_impl.hpp diff --git a/sycl/test-e2e/Matrix/get_coordinate_ops_impl.hpp b/sycl/test-e2e/Matrix/Inputs/get_coordinate_ops_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/get_coordinate_ops_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/get_coordinate_ops_impl.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_all_sizes_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_all_sizes_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_all_sizes_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_all_sizes_impl.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_annotated_ptr_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_annotated_ptr_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_annotated_ptr_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_annotated_ptr_impl.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_apply_bf16_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_apply_bf16_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_apply_bf16_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_apply_bf16_impl.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_apply_cuda.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_apply_cuda.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_apply_cuda.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_apply_cuda.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_apply_two_matrices_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_apply_two_matrices_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_apply_two_matrices_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_apply_two_matrices_impl.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_bf16_fill_k_cache_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_bf16_fill_k_cache_impl.hpp similarity index 94% rename from sycl/test-e2e/Matrix/joint_matrix_bf16_fill_k_cache_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_bf16_fill_k_cache_impl.hpp index 097d7e42bfd96..ac611281168ad 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_bf16_fill_k_cache_impl.hpp +++ b/sycl/test-e2e/Matrix/Inputs/joint_matrix_bf16_fill_k_cache_impl.hpp @@ -38,16 +38,17 @@ template class MatMul; template < #if !defined(ARG_DIM) && !defined(RUNTIME_DIM) - size_t rowsA, size_t colsA, size_t rowsB, size_t colsB, + size_t rowsA, size_t colsA, size_t rowsB, size_t colsB, #endif // ARG_DIM, RUNTIME_DIM - size_t vnniFactor, typename TOperand, typename TResult, size_t TM, - size_t TN, size_t TK, size_t MCache1, size_t NCache1, size_t KCache1, - size_t MCache2, size_t NCache2, size_t KCache2> + size_t vnniFactor, typename TOperand, typename TResult, size_t TM, + size_t TN, size_t TK, size_t MCache1, size_t NCache1, size_t KCache1, + size_t MCache2, size_t NCache2, size_t KCache2> double joint_matmul(TOperand *A, TOperand *B, TResult *C, queue &q, int i #if defined(ARG_DIM) || defined(RUNTIME_DIM) - , size_t rowsA, size_t colsA, size_t rowsB, size_t colsB + , + size_t rowsA, size_t colsA, size_t rowsB, size_t colsB #endif // ARG_DIM, RUNTIME_DIM - ) { +) { size_t sgSize = get_sg_size>(q); range<2> global{rowsA / MCache1, (colsB / NCache1) * sgSize}; @@ -118,12 +119,12 @@ double joint_matmul(TOperand *A, TOperand *B, TResult *C, queue &q, int i // pm1B and pn1B are used to identify the distribution of subgroups // along the workgroup prefetch for B matrix. For A matrix, sgId is // enough. - size_t pm1B = sgId / 16; // prefetch m1 (sgId/16) - size_t pn1B = sgId & 0xF; // prefetch n1 (sgId%16) -#else // VNNI + size_t pm1B = sgId / 16; // prefetch m1 (sgId/16) + size_t pn1B = sgId & 0xF; // prefetch n1 (sgId%16) +#else // VNNI size_t pm1B = sgId / 8; // prefetch m1 (sgId/8) size_t pn1B = sgId & 0x7; // prefetch n1 (sgId%8) -#endif // VNNI +#endif // VNNI constexpr size_t prefDistance = 3; for (int p = 0; p < prefDistance; p++) joint_matrix_prefetch( @@ -306,8 +307,8 @@ double joint_matmul(TOperand *A, TOperand *B, TResult *C, queue &q, int i pm1B * prefRow) * (colsB)*vnniFactor + (n2 * NCache2 * vnniFactor + pn1B * prefCol); - if ((prefetch_offsetB + (prefRow * colsB * vnniFactor) + - prefCol) < (rowsB * colsB)) + if ((prefetch_offsetB + (prefRow * colsB * vnniFactor) + prefCol) < + (rowsB * colsB)) joint_matrix_prefetch( sg, B + prefetch_offsetB, colsB * vnniFactor, layout::row_major, @@ -395,18 +396,17 @@ void test(size_t matrix_size_input) { // run testIterations time, aggregate and calculate average run time double totalDuration = 0; for (unsigned int i = 0; i < testIterations; i++) { - double duration = - joint_matmul< + double duration = joint_matmul< #if !defined(ARG_DIM) && !defined(RUNTIME_DIM) - matrix_size, matrix_size, matrix_size, matrix_size, + matrix_size, matrix_size, matrix_size, matrix_size, #endif // ARG_DIM, RUNTIME_DIM - vnniFactor, T, TResult, TM, TN, TK, MCache1, NCache1, - KCache1, MCache2, NCache2, KCache2> - (A, B, C, q, i + vnniFactor, T, TResult, TM, TN, TK, MCache1, NCache1, KCache1, MCache2, + NCache2, KCache2>(A, B, C, q, i #if defined(ARG_DIM) || defined(RUNTIME_DIM) - , matrix_size, matrix_size, matrix_size, matrix_size + , + matrix_size, matrix_size, matrix_size, matrix_size #endif // ARG_DIM, RUNTIME_DIM - ); + ); if (i >= recordThresh) { totalDuration += duration; @@ -431,11 +431,11 @@ void test(size_t matrix_size_input) { int main( #ifdef RUNTIME_DIM - int argc, char *argv[] -#endif //RUNTIME_DIM - ) { + int argc, char *argv[] +#endif // RUNTIME_DIM +) { -size_t matrix_size = -1; + size_t matrix_size = -1; #ifdef RUNTIME_DIM if (argc == 2) { matrix_size = std::stoul(argv[1]); @@ -443,7 +443,7 @@ size_t matrix_size = -1; std::cerr << "Usage: ./program matrix_size\n"; return 1; // Error if no argument } -#endif //RUNTIME_DIM +#endif // RUNTIME_DIM queue q; std::vector combinations = diff --git a/sycl/test-e2e/Matrix/joint_matrix_bf16_rowmajorB_load_store_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_bf16_rowmajorB_load_store_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_bf16_rowmajorB_load_store_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_bf16_rowmajorB_load_store_impl.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_bf16_rowmajorB_pair_load_store_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_bf16_rowmajorB_pair_load_store_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_bf16_rowmajorB_pair_load_store_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_bf16_rowmajorB_pair_load_store_impl.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_bfloat16_array_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_bfloat16_array_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_bfloat16_array_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_bfloat16_array_impl.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_bfloat16_colmajorA_colmajorB_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_bfloat16_colmajorA_colmajorB_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_bfloat16_colmajorA_colmajorB_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_bfloat16_colmajorA_colmajorB_impl.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_bfloat16_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_bfloat16_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_bfloat16_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_bfloat16_impl.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_bfloat16_packedB_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_bfloat16_packedB_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_bfloat16_packedB_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_bfloat16_packedB_impl.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_colA_rowB_colC_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_colA_rowB_colC_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_colA_rowB_colC_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_colA_rowB_colC_impl.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_down_convert_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_down_convert_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_down_convert_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_down_convert_impl.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_gemm_cuda.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_gemm_cuda.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_gemm_cuda.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_gemm_cuda.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_half_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_half_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_half_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_half_impl.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_hip_apply.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_hip_apply.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_hip_apply.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_hip_apply.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_hip_copy.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_hip_copy.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_hip_copy.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_hip_copy.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_hip_fill.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_hip_fill.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_hip_fill.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_hip_fill.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_hip_mfma.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_hip_mfma.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_hip_mfma.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_hip_mfma.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_int8_colmajorA_colmajorB_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_int8_colmajorA_colmajorB_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_int8_colmajorA_colmajorB_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_int8_colmajorA_colmajorB_impl.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_int8_rowmajorA_rowmajorB_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_int8_rowmajorA_rowmajorB_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_int8_rowmajorA_rowmajorB_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_int8_rowmajorA_rowmajorB_impl.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_opt_kernel_feature_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_opt_kernel_feature_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_opt_kernel_feature_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_opt_kernel_feature_impl.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_out_bounds_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_out_bounds_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_out_bounds_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_out_bounds_impl.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_prefetch_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_prefetch_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_prefetch_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_prefetch_impl.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_rowmajorA_rowmajorB_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_rowmajorA_rowmajorB_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_rowmajorA_rowmajorB_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_rowmajorA_rowmajorB_impl.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_ss_int8_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_ss_int8_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_ss_int8_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_ss_int8_impl.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_su_int8_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_su_int8_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_su_int8_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_su_int8_impl.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_tf32_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_tf32_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_tf32_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_tf32_impl.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_transposeC_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_transposeC_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_transposeC_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_transposeC_impl.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_us_int8_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_us_int8_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_us_int8_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_us_int8_impl.hpp diff --git a/sycl/test-e2e/Matrix/joint_matrix_uu_int8_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_uu_int8_impl.hpp similarity index 100% rename from sycl/test-e2e/Matrix/joint_matrix_uu_int8_impl.hpp rename to sycl/test-e2e/Matrix/Inputs/joint_matrix_uu_int8_impl.hpp diff --git a/sycl/test-e2e/Matrix/slm_utils.hpp b/sycl/test-e2e/Matrix/Inputs/slm_utils.hpp similarity index 100% rename from sycl/test-e2e/Matrix/slm_utils.hpp rename to sycl/test-e2e/Matrix/Inputs/slm_utils.hpp diff --git a/sycl/test-e2e/Matrix/SG32/element_wise_abc.cpp b/sycl/test-e2e/Matrix/SG32/element_wise_abc.cpp index 1fdd989ae091f..e549f7a1ae57a 100644 --- a/sycl/test-e2e/Matrix/SG32/element_wise_abc.cpp +++ b/sycl/test-e2e/Matrix/SG32/element_wise_abc.cpp @@ -13,8 +13,8 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 -#include "../element_wise_abc_impl.hpp" +#include "element_wise_abc_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/element_wise_all_ops.cpp b/sycl/test-e2e/Matrix/SG32/element_wise_all_ops.cpp index fdcc7d2aef9f7..f39fc09ccf531 100644 --- a/sycl/test-e2e/Matrix/SG32/element_wise_all_ops.cpp +++ b/sycl/test-e2e/Matrix/SG32/element_wise_all_ops.cpp @@ -13,8 +13,8 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 -#include "../element_wise_all_ops_impl.hpp" +#include "element_wise_all_ops_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/element_wise_all_ops_half.cpp b/sycl/test-e2e/Matrix/SG32/element_wise_all_ops_half.cpp index 8468ebad1b8b3..156c180b9fcc1 100644 --- a/sycl/test-e2e/Matrix/SG32/element_wise_all_ops_half.cpp +++ b/sycl/test-e2e/Matrix/SG32/element_wise_all_ops_half.cpp @@ -14,8 +14,8 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 -#include "../element_wise_all_ops_half_impl.hpp" +#include "element_wise_all_ops_half_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/element_wise_all_ops_int8.cpp b/sycl/test-e2e/Matrix/SG32/element_wise_all_ops_int8.cpp index 984ff9a9b082f..5831c869e18bd 100644 --- a/sycl/test-e2e/Matrix/SG32/element_wise_all_ops_int8.cpp +++ b/sycl/test-e2e/Matrix/SG32/element_wise_all_ops_int8.cpp @@ -13,8 +13,8 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 -#include "../element_wise_all_ops_int8_impl.hpp" +#include "element_wise_all_ops_int8_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/element_wise_all_ops_int8_packed.cpp b/sycl/test-e2e/Matrix/SG32/element_wise_all_ops_int8_packed.cpp index af2f4df82b648..3ee62e2454b82 100644 --- a/sycl/test-e2e/Matrix/SG32/element_wise_all_ops_int8_packed.cpp +++ b/sycl/test-e2e/Matrix/SG32/element_wise_all_ops_int8_packed.cpp @@ -15,8 +15,8 @@ // This test stores the matrix B that is VNNIed (packed). -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 -#include "../element_wise_all_ops_int8_packed_impl.hpp" +#include "element_wise_all_ops_int8_packed_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/element_wise_all_ops_tf32.cpp b/sycl/test-e2e/Matrix/SG32/element_wise_all_ops_tf32.cpp index 6b6416280a6bd..9ba12f04856df 100644 --- a/sycl/test-e2e/Matrix/SG32/element_wise_all_ops_tf32.cpp +++ b/sycl/test-e2e/Matrix/SG32/element_wise_all_ops_tf32.cpp @@ -11,9 +11,9 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 constexpr size_t TN = 16; -#include "../element_wise_all_ops_tf32_impl.hpp" +#include "element_wise_all_ops_tf32_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/element_wise_all_sizes.cpp b/sycl/test-e2e/Matrix/SG32/element_wise_all_sizes.cpp index df0cc30b6aae1..f1221c668886f 100644 --- a/sycl/test-e2e/Matrix/SG32/element_wise_all_sizes.cpp +++ b/sycl/test-e2e/Matrix/SG32/element_wise_all_sizes.cpp @@ -14,8 +14,8 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 -#include "../element_wise_all_sizes_impl.hpp" +#include "element_wise_all_sizes_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/element_wise_ops.cpp b/sycl/test-e2e/Matrix/SG32/element_wise_ops.cpp index 1a09518e65ffb..c458c3c24c456 100644 --- a/sycl/test-e2e/Matrix/SG32/element_wise_ops.cpp +++ b/sycl/test-e2e/Matrix/SG32/element_wise_ops.cpp @@ -13,8 +13,8 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 -#include "../element_wise_ops_impl.hpp" +#include "element_wise_ops_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/get_coordinate_ops.cpp b/sycl/test-e2e/Matrix/SG32/get_coordinate_ops.cpp index 0402a56e5c15d..491b2410ac174 100644 --- a/sycl/test-e2e/Matrix/SG32/get_coordinate_ops.cpp +++ b/sycl/test-e2e/Matrix/SG32/get_coordinate_ops.cpp @@ -19,8 +19,8 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 -#include "../get_coordinate_ops_impl.hpp" +#include "get_coordinate_ops_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_all_sizes.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_all_sizes.cpp index 3fe1f33343b17..664044dd242f6 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_all_sizes.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_all_sizes.cpp @@ -13,8 +13,8 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 -#include "../joint_matrix_all_sizes_impl.hpp" +#include "joint_matrix_all_sizes_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_annotated_ptr.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_annotated_ptr.cpp index 13656b8e3f4ec..dc5378808439a 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_annotated_ptr.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_annotated_ptr.cpp @@ -15,9 +15,9 @@ // XFAIL: gpu // XFAIL-TRACKER: GSD-4181 -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 constexpr size_t TN = 16; -#include "../joint_matrix_annotated_ptr_impl.hpp" +#include "joint_matrix_annotated_ptr_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_apply_bf16.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_apply_bf16.cpp index 88d5ba29922de..12434f940e3f8 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_apply_bf16.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_apply_bf16.cpp @@ -13,8 +13,8 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 -#include "../joint_matrix_apply_bf16_impl.hpp" +#include "joint_matrix_apply_bf16_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_apply_two_matrices.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_apply_two_matrices.cpp index e0aa84e460731..95e37bf1cca57 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_apply_two_matrices.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_apply_two_matrices.cpp @@ -13,8 +13,8 @@ // RUN: %{build} %fp-model-precise -o %t.out // RUN: %{run} %t.out -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 -#include "../joint_matrix_apply_two_matrices_impl.hpp" +#include "joint_matrix_apply_two_matrices_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_fill_k_cache.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_fill_k_cache.cpp index 10e01af93d239..a46c4388be56f 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_fill_k_cache.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_fill_k_cache.cpp @@ -18,8 +18,8 @@ // -ffp-model=precise is added to not depend on compiler defaults. -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 -#include "../joint_matrix_bf16_fill_k_cache_impl.hpp" +#include "joint_matrix_bf16_fill_k_cache_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_fill_k_cache_SLM.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_fill_k_cache_SLM.cpp index 40f5576c0042b..360fea486323e 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_fill_k_cache_SLM.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_fill_k_cache_SLM.cpp @@ -17,7 +17,7 @@ // -ffp-model=precise is added to not depend on compiler defaults. -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 -#include "../joint_matrix_bf16_fill_k_cache_impl.hpp" +#include "joint_matrix_bf16_fill_k_cache_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_fill_k_cache_init.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_fill_k_cache_init.cpp index db358793a39f7..e8d0caa96d376 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_fill_k_cache_init.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_fill_k_cache_init.cpp @@ -15,8 +15,8 @@ // -ffp-model=precise is added to not depend on compiler defaults. -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 -#include "../joint_matrix_bf16_fill_k_cache_impl.hpp" +#include "joint_matrix_bf16_fill_k_cache_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_fill_k_cache_unroll.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_fill_k_cache_unroll.cpp index 79b50f2d88a5d..4d0cf849c5fcc 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_fill_k_cache_unroll.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_fill_k_cache_unroll.cpp @@ -17,8 +17,8 @@ // 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 "common.hpp" #define SG_SZ 32 -#include "../joint_matrix_bf16_fill_k_cache_impl.hpp" +#include "joint_matrix_bf16_fill_k_cache_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_fill_k_cache_unroll_init.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_fill_k_cache_unroll_init.cpp index 42c04e480b9ae..a01a050446f8e 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_fill_k_cache_unroll_init.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_fill_k_cache_unroll_init.cpp @@ -17,8 +17,8 @@ // 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 "common.hpp" #define SG_SZ 32 -#include "../joint_matrix_bf16_fill_k_cache_impl.hpp" +#include "joint_matrix_bf16_fill_k_cache_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_rowmajorB_load_store.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_rowmajorB_load_store.cpp index 731ae6a053bec..4c5a977bf496e 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_rowmajorB_load_store.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_rowmajorB_load_store.cpp @@ -10,8 +10,8 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 -#include "../joint_matrix_bf16_rowmajorB_load_store_impl.hpp" +#include "joint_matrix_bf16_rowmajorB_load_store_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_rowmajorB_pair_load_store.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_rowmajorB_pair_load_store.cpp index 890938ae59c43..37618ee0e085b 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_rowmajorB_pair_load_store.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_rowmajorB_pair_load_store.cpp @@ -10,8 +10,8 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 -#include "../joint_matrix_bf16_rowmajorB_pair_load_store_impl.hpp" +#include "joint_matrix_bf16_rowmajorB_pair_load_store_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16.cpp index 4857ded37a011..6f51a3f64e6e8 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16.cpp @@ -13,8 +13,8 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 -#include "../joint_matrix_bfloat16_impl.hpp" +#include "joint_matrix_bfloat16_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16_array.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16_array.cpp index 87fd837446618..e4e3190235bcc 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16_array.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16_array.cpp @@ -13,8 +13,8 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 -#include "../joint_matrix_bfloat16_array_impl.hpp" +#include "joint_matrix_bfloat16_array_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16_colmajorA_colmajorB.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16_colmajorA_colmajorB.cpp index 504da33d936f1..f4f2e1719dd6c 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16_colmajorA_colmajorB.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16_colmajorA_colmajorB.cpp @@ -17,6 +17,6 @@ // XFAIL: gpu // XFAIL-TRACKER: GSD-5768 -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 -#include "../joint_matrix_bfloat16_colmajorA_colmajorB_impl.hpp" +#include "joint_matrix_bfloat16_colmajorA_colmajorB_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16_packedB.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16_packedB.cpp index 1944a331bfe3a..aa2745bd05406 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16_packedB.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16_packedB.cpp @@ -17,7 +17,7 @@ // XFAIL: gpu // XFAIL-TRACKER: GSD-4181 -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 -#include "../joint_matrix_bfloat16_packedB_impl.hpp" +#include "joint_matrix_bfloat16_packedB_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_colA_rowB_colC.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_colA_rowB_colC.cpp index 98ded99791115..83a19b55436e8 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_colA_rowB_colC.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_colA_rowB_colC.cpp @@ -14,6 +14,6 @@ // XFAIL: gpu && run-mode // XFAIL-TRACKER: GSD-5768 -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 -#include "../joint_matrix_colA_rowB_colC_impl.hpp" +#include "joint_matrix_colA_rowB_colC_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_down_convert.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_down_convert.cpp index ac6a308afdd0c..bbc02f8588f8c 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_down_convert.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_down_convert.cpp @@ -14,8 +14,8 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 -#include "../joint_matrix_down_convert_impl.hpp" +#include "joint_matrix_down_convert_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_half.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_half.cpp index a02206ee7a805..ef732b5ac0c8f 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_half.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_half.cpp @@ -14,8 +14,8 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 -#include "../joint_matrix_half_impl.hpp" +#include "joint_matrix_half_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_int8_colmajorA_colmajorB.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_int8_colmajorA_colmajorB.cpp index 7db1b0a618823..554a1ceb0d0bf 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_int8_colmajorA_colmajorB.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_int8_colmajorA_colmajorB.cpp @@ -17,7 +17,7 @@ // XFAIL: gpu // XFAIL-TRACKER: GSD-5768 -#include "../common.hpp" +#include "common.hpp" using namespace sycl; using namespace sycl::ext::oneapi::experimental::matrix; @@ -25,4 +25,4 @@ using namespace sycl::ext::oneapi::experimental::matrix; #define SG_SZ 32 constexpr size_t TN = 16; -#include "../joint_matrix_int8_colmajorA_colmajorB_impl.hpp" +#include "joint_matrix_int8_colmajorA_colmajorB_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_int8_rowmajorA_rowmajorB.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_int8_rowmajorA_rowmajorB.cpp index 3068b6bcca684..2b262f835c5fa 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_int8_rowmajorA_rowmajorB.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_int8_rowmajorA_rowmajorB.cpp @@ -16,8 +16,8 @@ // XFAIL: gpu // XFAIL-TRACKER: GSD-4181 -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 -#include "../joint_matrix_int8_rowmajorA_rowmajorB_impl.hpp" +#include "joint_matrix_int8_rowmajorA_rowmajorB_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_out_bounds.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_out_bounds.cpp index 97a6d17b4e1ee..732ddc84243cd 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_out_bounds.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_out_bounds.cpp @@ -16,10 +16,10 @@ // XFAIL:gpu // XFAIL-TRACKER: GSD-4181 -#include "../common.hpp" +#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" +#include "joint_matrix_out_bounds_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_prefetch.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_prefetch.cpp index 7e1520e95ff1c..96520227ea4eb 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_prefetch.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_prefetch.cpp @@ -16,8 +16,8 @@ // SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2 // UNSUPPORTED: gpu-intel-dg2 -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 constexpr size_t TN = 16; -#include "../joint_matrix_prefetch_impl.hpp" +#include "joint_matrix_prefetch_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_rowmajorA_rowmajorB.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_rowmajorA_rowmajorB.cpp index ca7310f6f5e15..cc931d4e52653 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_rowmajorA_rowmajorB.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_rowmajorA_rowmajorB.cpp @@ -18,11 +18,11 @@ // XFAIL: gpu // XFAIL-TRACKER: GSD-4181 -#include "../common.hpp" +#include "common.hpp" using namespace sycl; using namespace sycl::ext::oneapi::experimental::matrix; #define SG_SZ 32 -#include "../joint_matrix_rowmajorA_rowmajorB_impl.hpp" +#include "joint_matrix_rowmajorA_rowmajorB_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_ss_int8.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_ss_int8.cpp index 08f44d2a00090..72e2aeb66c061 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_ss_int8.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_ss_int8.cpp @@ -13,8 +13,8 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 -#include "../joint_matrix_ss_int8_impl.hpp" +#include "joint_matrix_ss_int8_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_su_int8.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_su_int8.cpp index c260cb7558ed5..74df83b3b2800 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_su_int8.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_su_int8.cpp @@ -13,8 +13,8 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 -#include "../joint_matrix_su_int8_impl.hpp" +#include "joint_matrix_su_int8_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_tf32.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_tf32.cpp index 8d9f7867e2508..f4b19ceff6859 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_tf32.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_tf32.cpp @@ -11,9 +11,9 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 constexpr size_t TN = 16; -#include "../joint_matrix_tf32_impl.hpp" +#include "joint_matrix_tf32_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_transposeC.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_transposeC.cpp index 6cea5a248e0b2..fd28fff805529 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_transposeC.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_transposeC.cpp @@ -13,8 +13,8 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 -#include "../joint_matrix_transposeC_impl.hpp" +#include "joint_matrix_transposeC_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_unaligned_k.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_unaligned_k.cpp index 342648810fea4..fa88c72729a40 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_unaligned_k.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_unaligned_k.cpp @@ -16,10 +16,10 @@ // XFAIL:gpu // XFAIL-TRACKER: GSD-4181 -#include "../common.hpp" +#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" +#include "joint_matrix_out_bounds_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_us_int8.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_us_int8.cpp index a0ed38373a43a..f2bfcefe34427 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_us_int8.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_us_int8.cpp @@ -13,8 +13,8 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 -#include "../joint_matrix_us_int8_impl.hpp" +#include "joint_matrix_us_int8_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_uu_int8.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_uu_int8.cpp index d5c5281f67c9f..94eb8a9efa1e8 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_uu_int8.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_uu_int8.cpp @@ -13,8 +13,8 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include "../common.hpp" +#include "common.hpp" #define SG_SZ 32 -#include "../joint_matrix_uu_int8_impl.hpp" +#include "joint_matrix_uu_int8_impl.hpp" diff --git a/sycl/test-e2e/Matrix/lit.local.cfg b/sycl/test-e2e/Matrix/lit.local.cfg index fe942ed141799..6f297c3012e11 100644 --- a/sycl/test-e2e/Matrix/lit.local.cfg +++ b/sycl/test-e2e/Matrix/lit.local.cfg @@ -2,3 +2,11 @@ # At the moment support of the feature depends on SPIR-V Backend & run-time # drivers version, so we temporarily mark it as unsupported (CMPLRLLVM-64705). config.unsupported_features += ['spirv-backend'] + +config.substitutions.append(("%helper-includes", "-I {}/Inputs".format(os.path.dirname(os.path.abspath(__file__))))) +original_clangxx="" +for substitution in config.substitutions: + if substitution[0] == "%clangxx": + original_clangxx=substitution[1] +config.substitutions.insert(0, + ("%clangxx", original_clangxx + ' %helper-includes '))