[SYCL][Matrix] Move joint_matrix to experimental::matrix namespace (#3584)

Author: yubingex007-a11y

sycl/include/CL/sycl.hpp

@@ -13,7 +13,7 @@
 #include <CL/sycl/ONEAPI/filter_selector.hpp>
 #include <CL/sycl/ONEAPI/function_pointer.hpp>
 #include <CL/sycl/ONEAPI/group_algorithm.hpp>
-#include <CL/sycl/ONEAPI/intel_matrix/matrix.hpp>
+#include <CL/sycl/ONEAPI/matrix/matrix.hpp>
 #include <CL/sycl/ONEAPI/reduction.hpp>
 #include <CL/sycl/ONEAPI/sub_group.hpp>
 #include <CL/sycl/accessor.hpp>

sycl/include/CL/sycl/ONEAPI/intel_matrix/matrix-amx.hpp renamed to sycl/include/CL/sycl/ONEAPI/matrix/matrix-amx.hpp

@@ -46,6 +46,7 @@ template <typename T> class submatrix {
   short rows, cols;
 };
 
+// TODO: we are adding it this way until sycl::dynamic_extent gets implemented.
 constexpr size_t dynamic_extent = std::numeric_limits<size_t>::max();
 
 template <typename T> struct elems_per_dword {
@@ -62,7 +63,7 @@ ELEMS_PER_DWORD(unsigned short, 2)
 
 } // namespace detail
 
-namespace matrix {
+namespace experimental::matrix {
 using namespace cl::sycl;
 using namespace cl::sycl::ONEAPI;
 
@@ -196,10 +197,12 @@ struct joint_matrix<
   joint_matrix(Group sg) {}
 };
 
-} // namespace matrix
+} // namespace experimental::matrix
 
 namespace detail {
 
+using namespace experimental;
+
 template <typename Group, typename T, size_t NumRows, size_t NumCols,
           matrix::matrix_layout Layout>
 inline __SYCL_ALWAYS_INLINE static
@@ -231,7 +234,8 @@ inline __SYCL_ALWAYS_INLINE static
                    uint32_t row, uint32_t col, size_t stride,
                    matrix::matrix_layout layout, bool shouldreload) {
   if (shouldreload) {
-    // Force sub_m.tile's shape to be matrix::tile_size * matrix::tile_size * 4
+    // Force sub_m.tile's shape to be matrix::tile_size *
+    // matrix::tile_size * 4
     int8_t NewjmC[matrix::tile_size * matrix::tile_size * 4];
     matrix::tilestored64_internal(NumRows, NumCols * sizeof(T),
                                   reinterpret_cast<char *>(NewjmC),
@@ -309,7 +313,7 @@ inline __SYCL_ALWAYS_INLINE static
 
 } // namespace detail
 
-namespace matrix {
+namespace experimental::matrix {
 
 // This handles cases where matrix can't be accommodated by a tile
 template <typename Group, typename T, size_t NumRows, size_t NumCols,
@@ -439,7 +443,7 @@ joint_matrix_mad(Group sg,
   return;
 }
 
-} // namespace matrix
+} // namespace experimental::matrix
 } // namespace intel
 } // namespace ext
 } // namespace sycl

sycl/include/CL/sycl/ONEAPI/intel_matrix/matrix.hpp renamed to sycl/include/CL/sycl/ONEAPI/matrix/matrix.hpp

@@ -18,6 +18,6 @@
 
 #if (SYCL_EXT_ONEAPI_MATRIX == 1)
 #if defined(__AMXTILE__) && defined(__AMXINT8__) && defined(__AMXBF16__)
-#include <CL/sycl/ONEAPI/intel_matrix/matrix-amx.hpp>
+#include <CL/sycl/ONEAPI/matrix/matrix-amx.hpp>
 #endif
 #endif

sycl/test/on-device/extensions/matrix-amx-bf16-test.cpp renamed to sycl/test/matrix/matrix-amx-bf16-test.cpp

@@ -3,9 +3,8 @@
 #if (SYCL_EXT_ONEAPI_MATRIX == 1)
 #include <iostream>
 
-using namespace cl::sycl;
-using namespace cl::sycl::intel;
-using namespace cl::sycl::ext::intel::matrix;
+using namespace sycl::intel;
+using namespace sycl::ext::intel::experimental::matrix;
 
 #define TILE_SZ 16
 #define TM (3 * TILE_SZ-1)

sycl/test/on-device/extensions/matrix-amx-int8-test.cpp renamed to sycl/test/matrix/matrix-amx-int8-test.cpp

@@ -3,9 +3,8 @@
 #if (SYCL_EXT_ONEAPI_MATRIX == 1)
 #include <iostream>
 
-using namespace cl::sycl;
-using namespace cl::sycl::intel;
-using namespace cl::sycl::ext::intel::matrix;
+using namespace sycl::intel;
+using namespace sycl::ext::intel::experimental::matrix;
 
 #define TILE_SZ 16
 #define TM (4 * TILE_SZ-4)