[SYCL][Joint Matrix] Tests query for default subgroup size (#13798)

Changes:
* Main folder now does not specify the subgroup (SG) size , it queries
it
* SG32 folder overrides the default SG size, through a `#define`
* XMX8 folder does not specify the SG size, it queries it
* minor cleanups
* `joint_matrix_opt_kernel_feature_impl.hpp` - issues a different error
code than before. I believe now that it is a named kernel it might be
taking a different code path through the runtime. Was
`errc::kernel_not_supported` now `errc::invalid`. I checked and if i
modify the tile size of `N` to be reasonable the test passes as
expected.

Author: artemrad

sycl/test-e2e/Matrix/SG32/element_wise_abc.cpp

@@ -13,7 +13,7 @@
 
 #include "../common.hpp"
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 constexpr size_t TN = 16;
 
 #include "../element_wise_abc_impl.hpp"

sycl/test-e2e/Matrix/SG32/element_wise_all_ops_half.cpp

@@ -15,7 +15,7 @@
 
 #include "../common.hpp"
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 constexpr size_t TN = 16;
 
 #include "../element_wise_all_ops_half_impl.hpp"

sycl/test-e2e/Matrix/SG32/element_wise_all_ops_int8.cpp

@@ -13,7 +13,7 @@
 
 #include "../common.hpp"
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 constexpr size_t TN = 16;
 
 #include "../element_wise_all_ops_int8_impl.hpp"

sycl/test-e2e/Matrix/SG32/element_wise_all_ops_int8_packed.cpp

@@ -15,7 +15,7 @@
 
 #include "../common.hpp"
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 constexpr size_t TN = 16;
 
 #include "../element_wise_all_ops_int8_packed_impl.hpp"

sycl/test-e2e/Matrix/SG32/element_wise_all_ops_tf32.cpp

@@ -13,7 +13,7 @@
 
 #include "../common.hpp"
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 constexpr size_t TN = 16;
 
 #include "../element_wise_all_ops_tf32_impl.hpp"

sycl/test-e2e/Matrix/SG32/element_wise_all_sizes.cpp

@@ -13,6 +13,6 @@
 
 #include "../common.hpp"
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 
 #include "../element_wise_all_sizes_impl.hpp"

sycl/test-e2e/Matrix/SG32/get_coord_float_matC.cpp

@@ -17,7 +17,7 @@
 using namespace sycl;
 using namespace sycl::ext::oneapi::experimental::matrix;
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 constexpr size_t TN = 16;
 
 #include "../get_coord_float_matC_impl.hpp"

sycl/test-e2e/Matrix/SG32/get_coord_int8_matA.cpp

@@ -17,7 +17,7 @@
 using namespace sycl;
 using namespace sycl::ext::oneapi::experimental::matrix;
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 constexpr size_t TN = 16;
 
 #include "../get_coord_int8_matA_impl.hpp"

sycl/test-e2e/Matrix/SG32/get_coord_int8_matB.cpp

@@ -18,7 +18,7 @@
 using namespace sycl;
 using namespace sycl::ext::oneapi::experimental::matrix;
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 constexpr size_t TN = 16;
 
 #include "../get_coord_int8_matB_impl.hpp"

sycl/test-e2e/Matrix/SG32/joint_matrix_all_sizes.cpp

@@ -17,7 +17,7 @@
 using namespace sycl;
 using namespace sycl::ext::oneapi::experimental::matrix;
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 // Sub-matrix N dimension
 static constexpr size_t SN = 16;
 

sycl/test-e2e/Matrix/SG32/joint_matrix_apply_bf16.cpp

@@ -13,7 +13,7 @@
 
 #include "../common.hpp"
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 constexpr size_t TN = 16;
 
 #include "../joint_matrix_apply_bf16_impl.hpp"

sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_fill_k_cache.cpp

@@ -19,7 +19,7 @@
 #include "../common.hpp"
 #include <cstddef>
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 constexpr size_t TN = 16;
 
 #include "../joint_matrix_bf16_fill_k_cache_impl.hpp"

sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_fill_k_cache_init.cpp

@@ -16,7 +16,7 @@
 #include "../common.hpp"
 #include <cstddef>
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 constexpr size_t TN = 16;
 
 #include "../joint_matrix_bf16_fill_k_cache_impl.hpp"

sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_fill_k_cache_unroll.cpp

@@ -21,7 +21,7 @@
 #include "../common.hpp"
 #include <cstddef>
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 constexpr size_t TN = 16;
 
 #include "../joint_matrix_bf16_fill_k_cache_impl.hpp"

sycl/test-e2e/Matrix/SG32/joint_matrix_bf16_fill_k_cache_unroll_init.cpp

@@ -18,7 +18,7 @@
 #include "../common.hpp"
 #include <cstddef>
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 constexpr size_t TN = 16;
 
 #include "../joint_matrix_bf16_fill_k_cache_impl.hpp"

sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16.cpp

@@ -20,7 +20,7 @@ using namespace sycl;
 using namespace sycl::ext::oneapi::experimental::matrix;
 using bfloat16 = sycl::ext::oneapi::bfloat16;
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 constexpr size_t TN = 16;
 
 #include "../joint_matrix_bfloat16_impl.hpp"

sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16_colmajorA_colmajorB.cpp

@@ -25,7 +25,7 @@ using namespace sycl;
 using namespace sycl::ext::oneapi::experimental::matrix;
 using bfloat16 = sycl::ext::oneapi::bfloat16;
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 constexpr size_t TN = 16;
 
 #include "../joint_matrix_bfloat16_colmajorA_colmajorB_impl.hpp"

sycl/test-e2e/Matrix/SG32/joint_matrix_colA_rowB_colC.cpp

@@ -15,7 +15,7 @@
 
 #include "../common.hpp"
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 constexpr size_t TN = 16;
 
 #include "../joint_matrix_colA_rowB_colC_impl.hpp"

sycl/test-e2e/Matrix/SG32/joint_matrix_down_convert.cpp

@@ -13,6 +13,6 @@
 
 #include "../common.hpp"
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 
 #include "../joint_matrix_down_convert_impl.hpp"

sycl/test-e2e/Matrix/SG32/joint_matrix_half.cpp

@@ -18,7 +18,7 @@
 using namespace sycl;
 using namespace sycl::ext::oneapi::experimental::matrix;
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 constexpr size_t TN = 16;
 
 #include "../joint_matrix_half_impl.hpp"

sycl/test-e2e/Matrix/SG32/joint_matrix_int8_colmajorA_colmajorB.cpp

@@ -21,7 +21,7 @@
 using namespace sycl;
 using namespace sycl::ext::oneapi::experimental::matrix;
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 constexpr size_t TN = 16;
 
 #include "../joint_matrix_int8_colmajorA_colmajorB_impl.hpp"

sycl/test-e2e/Matrix/SG32/joint_matrix_int8_vnni.cpp

@@ -18,7 +18,7 @@
 using namespace sycl;
 using namespace sycl::ext::oneapi::experimental::matrix;
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 constexpr size_t TN = 16;
 
 #include "../joint_matrix_int8_vnni_impl.hpp"

sycl/test-e2e/Matrix/SG32/joint_matrix_out_bounds.cpp

@@ -15,7 +15,7 @@
 
 #include "../common.hpp"
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 constexpr size_t TN = 16;
 constexpr size_t MATRIX_K = 1024 + 24;
 

sycl/test-e2e/Matrix/SG32/joint_matrix_ss_int8.cpp

@@ -16,7 +16,7 @@
 using namespace sycl;
 using namespace sycl::ext::oneapi::experimental::matrix;
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 constexpr size_t TN = 16;
 
 #include "../joint_matrix_ss_int8_impl.hpp"

sycl/test-e2e/Matrix/SG32/joint_matrix_su_int8.cpp

@@ -16,7 +16,7 @@
 using namespace sycl;
 using namespace sycl::ext::oneapi::experimental::matrix;
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 constexpr size_t TN = 16;
 
 #include "../joint_matrix_su_int8_impl.hpp"

sycl/test-e2e/Matrix/SG32/joint_matrix_tf32.cpp

@@ -16,7 +16,7 @@
 using namespace sycl;
 using namespace sycl::ext::oneapi::experimental::matrix;
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 constexpr size_t TN = 16;
 
 #include "../joint_matrix_tf32_impl.hpp"

sycl/test-e2e/Matrix/SG32/joint_matrix_transposeC.cpp

@@ -13,7 +13,7 @@
 
 #include "../common.hpp"
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 constexpr size_t TN = 16;
 
 #include "../joint_matrix_transposeC_impl.hpp"

sycl/test-e2e/Matrix/SG32/joint_matrix_unaligned_k.cpp

@@ -15,7 +15,7 @@
 
 #include "../common.hpp"
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 constexpr size_t TN = 16;
 static constexpr size_t MATRIX_K = 1024 + 14;
 

sycl/test-e2e/Matrix/SG32/joint_matrix_us_int8.cpp

@@ -16,7 +16,7 @@
 using namespace sycl;
 using namespace sycl::ext::oneapi::experimental::matrix;
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 constexpr size_t TN = 16;
 
 #include "../joint_matrix_us_int8_impl.hpp"

sycl/test-e2e/Matrix/SG32/joint_matrix_uu_int8.cpp

@@ -16,7 +16,7 @@
 using namespace sycl;
 using namespace sycl::ext::oneapi::experimental::matrix;
 
-constexpr size_t SG_SZ = 32;
+#define SG_SZ 32
 constexpr size_t TN = 16;
 
 #include "../joint_matrix_uu_int8_impl.hpp"

sycl/test-e2e/Matrix/XMX8/element_wise_abc.cpp

@@ -12,7 +12,6 @@
 
 #include "../common.hpp"
 
-#define SG_SZ 8
 constexpr size_t TN = 8;
 
 #include "../element_wise_abc_impl.hpp"

sycl/test-e2e/Matrix/XMX8/element_wise_all_ops_half.cpp

@@ -14,7 +14,6 @@
 
 #include "../common.hpp"
 
-#define SG_SZ 8
 constexpr size_t TN = 8;
 
 #include "../element_wise_all_ops_half_impl.hpp"

sycl/test-e2e/Matrix/XMX8/element_wise_all_ops_int8.cpp

@@ -12,7 +12,6 @@
 
 #include "../common.hpp"
 
-#define SG_SZ 8
 constexpr size_t TN = 8;
 
 #include "../element_wise_all_ops_int8_impl.hpp"

sycl/test-e2e/Matrix/XMX8/element_wise_all_ops_int8_packed.cpp

@@ -14,7 +14,6 @@
 
 #include "../common.hpp"
 
-#define SG_SZ 8
 constexpr size_t TN = 8;
 
 #include "../element_wise_all_ops_int8_packed_impl.hpp"

sycl/test-e2e/Matrix/XMX8/element_wise_all_sizes.cpp

@@ -13,7 +13,6 @@
 
 #include "../common.hpp"
 
-#define SG_SZ 8
 constexpr size_t TN = 8;
 
 #include "../element_wise_all_sizes_impl.hpp"

sycl/test-e2e/Matrix/XMX8/element_wise_all_sizes_no_split.cpp

@@ -16,7 +16,6 @@
 
 #include "../common.hpp"
 
-#define SG_SZ 8
 constexpr size_t TN = 8;
 
 #include "../element_wise_all_sizes_impl.hpp"

sycl/test-e2e/Matrix/XMX8/get_coord_float_matC.cpp

@@ -18,7 +18,6 @@
 using namespace sycl;
 using namespace sycl::ext::oneapi::experimental::matrix;
 
-constexpr size_t SG_SZ = 8;
 constexpr size_t TN = 8;
 
 #include "../get_coord_float_matC_impl.hpp"

sycl/test-e2e/Matrix/XMX8/get_coord_int8_matA.cpp

@@ -18,7 +18,6 @@
 using namespace sycl;
 using namespace sycl::ext::oneapi::experimental::matrix;
 
-constexpr size_t SG_SZ = 8;
 constexpr size_t TN = 8;
 
 #include "../get_coord_int8_matA_impl.hpp"

sycl/test-e2e/Matrix/XMX8/get_coord_int8_matB.cpp

@@ -17,7 +17,6 @@
 using namespace sycl;
 using namespace sycl::ext::oneapi::experimental::matrix;
 
-constexpr size_t SG_SZ = 8;
 constexpr size_t TN = 8;
 
 #include "../get_coord_int8_matB_impl.hpp"

sycl/test-e2e/Matrix/XMX8/joint_matrix_all_sizes.cpp

@@ -15,7 +15,6 @@
 using namespace sycl;
 using namespace sycl::ext::oneapi::experimental::matrix;
 
-#define SG_SZ 8
 constexpr size_t SN = 8;
 
 #include "../joint_matrix_all_sizes_impl.hpp"

sycl/test-e2e/Matrix/XMX8/joint_matrix_apply_bf16.cpp

@@ -12,7 +12,6 @@
 
 #include "../common.hpp"
 
-#define SG_SZ 8
 constexpr size_t TN = 8;
 
 #include "../joint_matrix_apply_bf16_impl.hpp"

sycl/test-e2e/Matrix/XMX8/joint_matrix_bf16_fill_k_cache.cpp

@@ -15,7 +15,6 @@
 #include "../common.hpp"
 #include <cstddef>
 
-#define SG_SZ 8
 constexpr size_t TN = 8;
 
 #include "../joint_matrix_bf16_fill_k_cache_impl.hpp"

sycl/test-e2e/Matrix/XMX8/joint_matrix_bf16_fill_k_cache_init.cpp

@@ -15,7 +15,6 @@
 #include "../common.hpp"
 #include <cstddef>
 
-#define SG_SZ 8
 constexpr size_t TN = 8;
 
 #include "../joint_matrix_bf16_fill_k_cache_impl.hpp"

sycl/test-e2e/Matrix/XMX8/joint_matrix_bf16_fill_k_cache_unroll.cpp

@@ -17,7 +17,6 @@
 #include "../common.hpp"
 #include <cstddef>
 
-#define SG_SZ 8
 constexpr size_t TN = 8;
 
 #include "../joint_matrix_bf16_fill_k_cache_impl.hpp"

sycl/test-e2e/Matrix/XMX8/joint_matrix_bf16_fill_k_cache_unroll_init.cpp

@@ -17,7 +17,6 @@
 #include "../common.hpp"
 #include <cstddef>
 
-#define SG_SZ 8
 constexpr size_t TN = 8;
 
 #include "../joint_matrix_bf16_fill_k_cache_impl.hpp"

sycl/test-e2e/Matrix/XMX8/joint_matrix_bfloat16.cpp

@@ -15,7 +15,6 @@
 using namespace sycl;
 using namespace sycl::ext::oneapi::experimental::matrix;
 
-#define SG_SZ 8
 constexpr size_t TN = 8;
 
 #include "../joint_matrix_bfloat16_impl.hpp"