[SYCL][Matrix] Change packed to ext_intel_packed

intel · May 31, 2023 · fed879f · fed879f
1 parent 40ac9cc
commit fed879f
Show file tree

Hide file tree

Showing 25 changed files with 42 additions and 41 deletions.
diff --git a/sycl/include/sycl/ext/oneapi/matrix/matrix-intel.hpp b/sycl/include/sycl/ext/oneapi/matrix/matrix-intel.hpp
@@ -18,7 +18,7 @@ namespace sycl {
 __SYCL_INLINE_VER_NAMESPACE(_V1) {
 namespace ext {
 namespace intel::experimental::matrix::layout {
-constexpr sycl::ext::oneapi::experimental::matrix::layout packed =
+constexpr sycl::ext::oneapi::experimental::matrix::layout ext_intel_packed =
     static_cast<sycl::ext::oneapi::experimental::matrix::layout>(2);
 }
 namespace oneapi {
@@ -36,8 +36,9 @@ template <layout Layout> struct spv_matrix_layout_traits {
 
 SPV_MATRIX_LAYOUT_TRAITS(layout::row_major, __spv::MatrixLayout::RowMajor)
 SPV_MATRIX_LAYOUT_TRAITS(layout::col_major, __spv::MatrixLayout::ColumnMajor)
-SPV_MATRIX_LAYOUT_TRAITS(sycl::ext::intel::experimental::matrix::layout::packed,
-                         __spv::MatrixLayout::Packed)
+SPV_MATRIX_LAYOUT_TRAITS(
+    sycl::ext::intel::experimental::matrix::layout::ext_intel_packed,
+    __spv::MatrixLayout::Packed)
 SPV_MATRIX_LAYOUT_TRAITS(layout::dynamic, __spv::MatrixLayout::Dynamic)
 
 template <use Use> struct spv_matrix_use_traits {

diff --git a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp
@@ -228,7 +228,7 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_load(
         Ptr, stride, __spv::MatrixLayout::ColumnMajor,
         spv_scope_traits<Group>::value);
     break;
-  case sycl::ext::intel::experimental::matrix::layout::packed:
+  case sycl::ext::intel::experimental::matrix::layout::ext_intel_packed:
     res.spvm = __spirv_JointMatrixLoadINTEL<
         DecorT, S, NumRows, NumCols,
         spv_matrix_use_traits<use::accumulator>::value,
@@ -327,7 +327,7 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_store(
         Ptr, src.spvm, stride, __spv::MatrixLayout::ColumnMajor,
         spv_scope_traits<Group>::value);
     break;
-  case sycl::ext::intel::experimental::matrix::layout::packed:
+  case sycl::ext::intel::experimental::matrix::layout::ext_intel_packed:
     __spirv_JointMatrixStoreINTEL<
         DecorT, T, NumRows, NumCols,
         spv_matrix_use_traits<use::accumulator>::value,

diff --git a/sycl/test-e2e/Matrix/element_wise_all_ops_int8_packed_impl.hpp b/sycl/test-e2e/Matrix/element_wise_all_ops_int8_packed_impl.hpp
@@ -38,7 +38,7 @@ void matrix_verify_add(queue q, big_matrix<T, M, N> &A, nd_range<2> &r,
 
            sub_group sg = spmd_item.get_sub_group();
            joint_matrix<sub_group, int8_t, use::b, TK, TN,
-                        ext::intel::experimental::matrix::layout::packed>
+                        ext::intel::experimental::matrix::layout::ext_intel_packed>
                sub_b;
 
            joint_matrix_fill(sg, sub_b, 5);
@@ -75,7 +75,7 @@ void matrix_verify_sub(queue q, big_matrix<T, M, N> &A, nd_range<2> &r,
 
            sub_group sg = spmd_item.get_sub_group();
            joint_matrix<sub_group, int8_t, use::b, TK, TN,
-                        ext::intel::experimental::matrix::layout::packed>
+                        ext::intel::experimental::matrix::layout::ext_intel_packed>
                sub_b;
 
            joint_matrix_fill(sg, sub_b, 5);
@@ -112,7 +112,7 @@ void matrix_verify_mul(queue q, big_matrix<T, M, N> &A, nd_range<2> &r,
 
            sub_group sg = spmd_item.get_sub_group();
            joint_matrix<sub_group, int8_t, use::b, TK, TN,
-                        ext::intel::experimental::matrix::layout::packed>
+                        ext::intel::experimental::matrix::layout::ext_intel_packed>
                sub_b;
 
            joint_matrix_fill(sg, sub_b, 5);
@@ -149,7 +149,7 @@ void matrix_verify_div(queue q, big_matrix<T, M, N> &A, nd_range<2> &r,
 
            sub_group sg = spmd_item.get_sub_group();
            joint_matrix<sub_group, int8_t, use::b, TK, TN,
-                        ext::intel::experimental::matrix::layout::packed>
+                        ext::intel::experimental::matrix::layout::ext_intel_packed>
                sub_b;
 
            joint_matrix_fill(sg, sub_b, 4);
@@ -186,7 +186,7 @@ void matrix_verify_logic(queue q, big_matrix<T, M, N> &A, nd_range<2> &r,
 
            sub_group sg = spmd_item.get_sub_group();
            joint_matrix<sub_group, int8_t, use::b, TK, TN,
-                        ext::intel::experimental::matrix::layout::packed>
+                        ext::intel::experimental::matrix::layout::ext_intel_packed>
                sub_b;
 
            joint_matrix_fill(sg, sub_b, 5);

diff --git a/sycl/test-e2e/Matrix/element_wise_irreg_sum_rows_impl.hpp b/sycl/test-e2e/Matrix/element_wise_irreg_sum_rows_impl.hpp
@@ -46,7 +46,7 @@ void matrix_sum_rows(queue q, big_matrix<T, M, N> &B, nd_range<2> &r) {
            ext::oneapi::sub_group sg = spmd_item.get_sub_group();
 
            joint_matrix<sub_group, T, use::b, TK, TN,
-                        ext::intel::experimental::matrix::layout::packed>
+                        ext::intel::experimental::matrix::layout::ext_intel_packed>
                sub_b;
 
            joint_matrix_load(

diff --git a/sycl/test-e2e/Matrix/element_wise_ops_impl.hpp b/sycl/test-e2e/Matrix/element_wise_ops_impl.hpp
@@ -53,7 +53,7 @@ void matrix_multiply(big_matrix<T1, NUM_ROWS_C, NUM_COLS_C> &C,
                sub_a;
            // For B, we assume B has been already VNNIed.
            joint_matrix<sub_group, int8_t, use::b, TK, TN,
-                        ext::intel::experimental::matrix::layout::packed>
+                        ext::intel::experimental::matrix::layout::ext_intel_packed>
                sub_b;
            joint_matrix<sub_group, int32_t, use::accumulator, TM, TN> sub_c;
 

diff --git a/sycl/test-e2e/Matrix/elemwise_irreg_size_ops_bf16.cpp b/sycl/test-e2e/Matrix/elemwise_irreg_size_ops_bf16.cpp
@@ -81,7 +81,7 @@ void matrix_multiply(big_matrix<T1, NUM_ROWS_C, NUM_COLS_C> &C,
                sub_a;
            // For B, we assume B has been already VNNIed.
            joint_matrix<sub_group, bfloat16, use::b, TK, TN,
-                        ext::intel::experimental::matrix::layout::packed>
+                        ext::intel::experimental::matrix::layout::ext_intel_packed>
                sub_b;
            joint_matrix<sub_group, float, use::accumulator, TM, TN> sub_c;
            joint_matrix_load(

diff --git a/sycl/test-e2e/Matrix/get_coord_bf16_gemm_impl.hpp b/sycl/test-e2e/Matrix/get_coord_bf16_gemm_impl.hpp
@@ -73,7 +73,7 @@ void matrix_multiply(big_matrix<T1, M, N> &C, big_matrix<T2, M, K> &A,
                sub_a;
            // For B, we assume B has been already VNNIed.
            joint_matrix<sub_group, bfloat16, use::b, TK, TN,
-                        ext::intel::experimental::matrix::layout::packed>
+                        ext::intel::experimental::matrix::layout::ext_intel_packed>
                sub_b;
            joint_matrix<sub_group, float, use::accumulator, TM, TN> sub_c;
 

diff --git a/sycl/test-e2e/Matrix/get_coord_bf16_matB_impl.hpp b/sycl/test-e2e/Matrix/get_coord_bf16_matB_impl.hpp
@@ -136,7 +136,7 @@ void matrix_sum_cols(queue q, big_matrix<T, M, N> &B, nd_range<2> &r) {
 
            // TK = 32, TN = 16
            joint_matrix<sub_group, int8_t, use::b, TK, TN,
-                        ext::intel::experimental::matrix::layout::packed>
+                        ext::intel::experimental::matrix::layout::ext_intel_packed>
                sub_b;
 
            joint_matrix_load(

diff --git a/sycl/test-e2e/Matrix/joint_matrix_all_sizes_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_all_sizes_impl.hpp
@@ -57,7 +57,7 @@ void matrix_multiply(big_matrix<T1, M, N> &C, big_matrix<T2, M, K> &A,
            joint_matrix<sub_group, T2, use::a, TM, TK, layout::row_major> sub_a;
            // For B, we assume B has been already VNNIed.
            joint_matrix<sub_group, T2, use::b, TK, TN,
-                        ext::intel::experimental::matrix::layout::packed>
+                        ext::intel::experimental::matrix::layout::ext_intel_packed>
                sub_b;
            joint_matrix<sub_group, T1, use::accumulator, TM, TN> sub_c;
 

diff --git a/sycl/test-e2e/Matrix/joint_matrix_bf16_fill_k_cache_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_bf16_fill_k_cache_impl.hpp
@@ -153,35 +153,35 @@ double joint_matmul(TOperand *A, TOperand *B, TResult *C, queue &q, int i) {
             ;
 
             joint_matrix<sub_group, TOperand, use::b, tK, tN,
-                         ext::intel::experimental::matrix::layout::packed>
+                         ext::intel::experimental::matrix::layout::ext_intel_packed>
                 tB[NCACHE1 / tN][KCACHE2 / KCACHE1]
 #ifdef INIT_LIST
                 =
                     {
                         joint_matrix<
                             sub_group, TOperand, use::b, tK, tN,
-                            ext::intel::experimental::matrix::layout::packed>(),
+                            ext::intel::experimental::matrix::layout::ext_intel_packed>(),
                         joint_matrix<
                             sub_group, TOperand, use::b, tK, tN,
-                            ext::intel::experimental::matrix::layout::packed>(),
+                            ext::intel::experimental::matrix::layout::ext_intel_packed>(),
                         joint_matrix<
                             sub_group, TOperand, use::b, tK, tN,
-                            ext::intel::experimental::matrix::layout::packed>(),
+                            ext::intel::experimental::matrix::layout::ext_intel_packed>(),
                         joint_matrix<
                             sub_group, TOperand, use::b, tK, tN,
-                            ext::intel::experimental::matrix::layout::packed>(),
+                            ext::intel::experimental::matrix::layout::ext_intel_packed>(),
                         joint_matrix<
                             sub_group, TOperand, use::b, tK, tN,
-                            ext::intel::experimental::matrix::layout::packed>(),
+                            ext::intel::experimental::matrix::layout::ext_intel_packed>(),
                         joint_matrix<
                             sub_group, TOperand, use::b, tK, tN,
-                            ext::intel::experimental::matrix::layout::packed>(),
+                            ext::intel::experimental::matrix::layout::ext_intel_packed>(),
                         joint_matrix<
                             sub_group, TOperand, use::b, tK, tN,
-                            ext::intel::experimental::matrix::layout::packed>(),
+                            ext::intel::experimental::matrix::layout::ext_intel_packed>(),
                         joint_matrix<
                             sub_group, TOperand, use::b, tK, tN,
-                            ext::intel::experimental::matrix::layout::packed>(),
+                            ext::intel::experimental::matrix::layout::ext_intel_packed>(),
                     }
 #endif
             ;

diff --git a/sycl/test-e2e/Matrix/joint_matrix_bfloat16_32x64_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_bfloat16_32x64_impl.hpp
@@ -47,7 +47,7 @@ void matrix_multiply(big_matrix<T1, M, N> &C, big_matrix<T2, M, K> &A,
                sub_a;
            // For B, we assume B has been already VNNIed.
            joint_matrix<sub_group, bfloat16, use::b, TK, TN,
-                        ext::intel::experimental::matrix::layout::packed>
+                        ext::intel::experimental::matrix::layout::ext_intel_packed>
                sub_b;
            joint_matrix<sub_group, float, use::accumulator, TM, TN> sub_c;
 

diff --git a/sycl/test-e2e/Matrix/joint_matrix_bfloat16_array_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_bfloat16_array_impl.hpp
@@ -61,7 +61,7 @@ void matrix_multiply(big_matrix<T1, M, N> &C, big_matrix<T2, M, K> &A,
 
            // For B, we assume B has been already VNNIed.
            joint_matrix<sub_group, bfloat16, use::b, TK, TN,
-                        ext::intel::experimental::matrix::layout::packed>
+                        ext::intel::experimental::matrix::layout::ext_intel_packed>
                sub_b;
            joint_matrix<sub_group, float, use::accumulator, TM, TN>
                sub_c[JM_ARRAY_SZ];

diff --git a/sycl/test-e2e/Matrix/joint_matrix_bfloat16_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_bfloat16_impl.hpp
@@ -47,7 +47,7 @@ void matrix_multiply(big_matrix<T1, M, N> &C, big_matrix<T2, M, K> &A,
                sub_a;
            // For B, we assume B has been already VNNIed.
            joint_matrix<sub_group, bfloat16, use::b, TK, TN,
-                        ext::intel::experimental::matrix::layout::packed>
+                        ext::intel::experimental::matrix::layout::ext_intel_packed>
                sub_b;
            joint_matrix<sub_group, float, use::accumulator, TM, TN> sub_c;
 

diff --git a/sycl/test-e2e/Matrix/joint_matrix_half_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_half_impl.hpp
@@ -52,7 +52,7 @@ void matrix_multiply(big_matrix<T1, NUM_ROWS_C, NUM_COLS_C> &C,
                    sub_a;
                // For B, we assume B has been already VNNIed.
                joint_matrix<sub_group, half, use::b, TK, TN,
-                            ext::intel::experimental::matrix::layout::packed>
+                            ext::intel::experimental::matrix::layout::ext_intel_packed>
                    sub_b;
                joint_matrix<sub_group, float, use::accumulator, TM, TN> sub_c;
 

diff --git a/sycl/test-e2e/Matrix/joint_matrix_query_default.cpp b/sycl/test-e2e/Matrix/joint_matrix_query_default.cpp
@@ -78,7 +78,7 @@ void matrix_multiply(big_matrix<T1, NUM_ROWS_C, NUM_COLS_C> &C,
 
            myparams2::joint_matrix_a<sub_group, layout::row_major> sub_a;
            myparams2::joint_matrix_b<
-               sub_group, ext::intel::experimental::matrix::layout::packed>
+               sub_group, ext::intel::experimental::matrix::layout::ext_intel_packed>
                sub_b;
            myparams2::joint_matrix_accumulator<sub_group> sub_c;
 

diff --git a/sycl/test-e2e/Matrix/joint_matrix_ss_int8_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_ss_int8_impl.hpp
@@ -53,7 +53,7 @@ void matrix_multiply(big_matrix<T1, NUM_ROWS_C, NUM_COLS_C> &C,
                sub_a;
            // For B, we assume B has been already VNNIed.
            joint_matrix<sub_group, int8_t, use::b, TK, TN,
-                        ext::intel::experimental::matrix::layout::packed>
+                        ext::intel::experimental::matrix::layout::ext_intel_packed>
                sub_b;
            joint_matrix<sub_group, int32_t, use::accumulator, TM, TN> sub_c;
 

diff --git a/sycl/test-e2e/Matrix/joint_matrix_su_int8_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_su_int8_impl.hpp
@@ -53,7 +53,7 @@ void matrix_multiply(big_matrix<T1, NUM_ROWS_C, NUM_COLS_C> &C,
                sub_a;
            // For B, we assume B has been already VNNIed.
            joint_matrix<sub_group, uint8_t, use::b, TK, TN,
-                        ext::intel::experimental::matrix::layout::packed>
+                        ext::intel::experimental::matrix::layout::ext_intel_packed>
                sub_b;
            joint_matrix<sub_group, int32_t, use::accumulator, TM, TN> sub_c;
 

diff --git a/sycl/test-e2e/Matrix/joint_matrix_us_int8_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_us_int8_impl.hpp
@@ -55,7 +55,7 @@ void matrix_multiply(big_matrix<T1, NUM_ROWS_C, NUM_COLS_C> &C,
                sub_a;
            // For B, we assume B has been already VNNIed.
            joint_matrix<sub_group, int8_t, use::b, TK, TN,
-                        ext::intel::experimental::matrix::layout::packed>
+                        ext::intel::experimental::matrix::layout::ext_intel_packed>
                sub_b;
            joint_matrix<sub_group, int32_t, use::accumulator, TM, TN> sub_c;
 

diff --git a/sycl/test-e2e/Matrix/joint_matrix_uu_int8_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_uu_int8_impl.hpp
@@ -53,7 +53,7 @@ void matrix_multiply(big_matrix<T1, NUM_ROWS_C, NUM_COLS_C> &C,
                sub_a;
            // For B, we assume B has been already VNNIed.
            joint_matrix<sub_group, uint8_t, use::b, TK, TN,
-                        ext::intel::experimental::matrix::layout::packed>
+                        ext::intel::experimental::matrix::layout::ext_intel_packed>
                sub_b;
            joint_matrix<sub_group, int32_t, use::accumulator, TM, TN> sub_c;
 

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
@@ -27,7 +27,7 @@ int main(void) {
                        layout::row_major>
               tA;
           joint_matrix<sub_group, unsigned short, use::b, 16, 16,
-                       ext::intel::experimental::matrix::layout::packed>
+                       ext::intel::experimental::matrix::layout::ext_intel_packed>
               tB;
           joint_matrix<sub_group, float, use::accumulator, 8, 16> tC;
 

diff --git a/sycl/test/matrix/matrix-bfloat16-test-coord-basicB.cpp b/sycl/test/matrix/matrix-bfloat16-test-coord-basicB.cpp
@@ -155,7 +155,7 @@ void matrix_sum_cols(queue q, big_matrix<T, M, N> &B, nd_range<2> &r) {
 
            // TK = 32, TN = 16
            joint_matrix<sub_group, int8_t, use::b, TK, TN,
-                        ext::intel::experimental::matrix::layout::packed>
+                        ext::intel::experimental::matrix::layout::ext_intel_packed>
                sub_b;
 
            joint_matrix_load(

diff --git a/sycl/test/matrix/matrix-bfloat16-test.cpp b/sycl/test/matrix/matrix-bfloat16-test.cpp
@@ -68,7 +68,7 @@ void matrix_multiply(big_matrix<T1, NUM_ROWS_C, NUM_COLS_C> &C,
            // the packed_b layout. By default, the layout is row_major and size
            // is (TK, TN).
            joint_matrix<sycl::sub_group, bfloat16, use::b, TK, TN,
-                        sycl::ext::intel::experimental::matrix::layout::packed>
+                        sycl::ext::intel::experimental::matrix::layout::ext_intel_packed>
                sub_b;
            joint_matrix<sycl::sub_group, float, use::accumulator, TM, TN> sub_c;
 

diff --git a/sycl/test/matrix/matrix-elemwise-ops.cpp b/sycl/test/matrix/matrix-elemwise-ops.cpp
@@ -69,7 +69,7 @@ void matrix_multiply(big_matrix<T1, NUM_ROWS_C, NUM_COLS_C> &C,
            // the packed_b layout. By default, the layout is row_major and size
            // is (TK, TN).
            joint_matrix<sycl::sub_group, int8_t, use::b, TK, TN,
-                        sycl::ext::intel::experimental::matrix::layout::packed>
+                        sycl::ext::intel::experimental::matrix::layout::ext_intel_packed>
                sub_b;
            joint_matrix<sycl::sub_group, int32_t, use::accumulator, TM, TN>
                sub_c;

diff --git a/sycl/test/matrix/matrix-int8-test.cpp b/sycl/test/matrix/matrix-int8-test.cpp
@@ -74,7 +74,7 @@ void matrix_multiply(big_matrix<T1, NUM_ROWS_C, NUM_COLS_C> &C,
            // the packed_b layout. By default, the layout is row_major and size
            // is (TK, TN).
            joint_matrix<sycl::sub_group, int8_t, use::b, TK, TN,
-                        sycl::ext::intel::experimental::matrix::layout::packed>
+                        sycl::ext::intel::experimental::matrix::layout::ext_intel_packed>
                sub_b;
            joint_matrix<sycl::sub_group, int32_t, use::accumulator, TM, TN>
                sub_c;

diff --git a/sycl/test/matrix/query-use.cpp b/sycl/test/matrix/query-use.cpp
@@ -64,7 +64,7 @@ void query_amx() {
           sub_group sg = spmd_item.get_sub_group();
           myparams2::joint_matrix_a<sub_group, layout::row_major> sub_a1;
           myparams2::joint_matrix_b<
-              sub_group, sycl::ext::intel::experimental::matrix::layout::packed>
+              sub_group, sycl::ext::intel::experimental::matrix::layout::ext_intel_packed>
               sub_b1;
           myparams2::joint_matrix_accumulator<sub_group> sub_c1;
 
@@ -144,7 +144,7 @@ void query_xmx8() {
           sub_group sg = spmd_item.get_sub_group();
           myparams2::joint_matrix_a<sub_group, layout::row_major> sub_a1;
           myparams2::joint_matrix_b<
-              sub_group, sycl::ext::intel::experimental::matrix::layout::packed>
+              sub_group, sycl::ext::intel::experimental::matrix::layout::ext_intel_packed>
               sub_b1;
           myparams2::joint_matrix_accumulator<sub_group> sub_c1;