[SYCL][NFC] Cleanup wi_data/joint_matrix code (#7929)

There doesn't appear to be any reason to have an #ifdef in get_wi_data,
so I used the template inference in both cases. I've updated the runtime
errors to be accurate, and removed some of my old out of date comments.

Signed-off-by: JackAKirk <[email protected]>

Author: JackAKirk

sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp

@@ -28,8 +28,8 @@ struct joint_matrix {
   __spv::__spirv_JointMatrixINTEL<
       T, Rows, Cols, spv_matrix_layout_traits<Layout>::value,
       spv_scope_traits<Group>::value, spv_matrix_use_traits<Use>::value> *spvm;
+#endif // defined(__NVPTX__)
 #endif // defined(__SYCL_DEVICE_ONLY__)
-#endif
 
   joint_matrix() {
 #ifndef __SYCL_DEVICE_ONLY__
@@ -93,12 +93,8 @@ template <typename Group, typename T, use Use, size_t Rows, size_t Cols,
 inline __SYCL_ALWAYS_INLINE decltype(auto)
 get_wi_data(Group sg, joint_matrix<Group, T, Use, Rows, Cols, Layout> &jm) {
 #if defined(__SYCL_DEVICE_ONLY__)
-#if defined(__NVPTX__)
   std::ignore = sg;
   return wi_data(jm);
-#else
-  return wi_data<Group, T, Use, Rows, Cols, Layout>(jm);
-#endif // defined(__NVPTX__)
 #else
   if constexpr (std::is_same_v<T, precision::tf32>) {
     marray<float, 1> unused{};
@@ -131,10 +127,8 @@ joint_matrix_fill(Group sg,
   std::ignore = sg;
   std::ignore = res;
   std::ignore = v;
-  throw runtime_error(
-      "This version of the matrix extension is only currently supported on "
-      "Nvidia devices",
-      PI_ERROR_INVALID_DEVICE);
+  throw runtime_error("joint matrix is not supported on host device.",
+                      PI_ERROR_INVALID_DEVICE);
 #endif // defined(__SYCL_DEVICE_ONLY__)
 }
 
@@ -155,8 +149,6 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_load(
   sycl::ext::oneapi::detail::load_accumulator_cuda(res.cuda_impl, src, stride,
                                                    Layout);
 #else
-  // intel's impl
-  // matL is determined by matrix.use?
   T *Ptr = src.get();
   switch (Layout) {
   default:
@@ -189,10 +181,8 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_load(
   std::ignore = res;
   std::ignore = src;
   std::ignore = stride;
-  throw runtime_error(
-      "This version of the matrix extension is only currently supported on "
-      "Nvidia devices",
-      PI_ERROR_INVALID_DEVICE);
+  throw runtime_error("joint matrix is not supported on host device.",
+                      PI_ERROR_INVALID_DEVICE);
 #endif // defined(__SYCL_DEVICE_ONLY__)
 }
 
@@ -228,10 +218,8 @@ joint_matrix_load(Group sg,
   std::ignore = res;
   std::ignore = src;
   std::ignore = stride;
-  throw runtime_error(
-      "This version of the matrix extension is only currently supported on "
-      "Nvidia devices",
-      PI_ERROR_INVALID_DEVICE);
+  throw runtime_error("joint matrix is not supported on host device.",
+                      PI_ERROR_INVALID_DEVICE);
 #endif // defined(__SYCL_DEVICE_ONLY__)
 }
 
@@ -250,7 +238,6 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_store(
                                                      Space>(src.cuda_impl, dst,
                                                             stride, Layout);
 #else
-  // intel's impl
   T *Ptr = dst.get();
   switch (Layout) {
   default:
@@ -283,10 +270,8 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_store(
   std::ignore = src;
   std::ignore = dst;
   std::ignore = stride;
-  throw runtime_error(
-      "This version of the matrix extension is only currently supported on "
-      "Nvidia devices",
-      PI_ERROR_INVALID_DEVICE);
+  throw runtime_error("joint matrix is not supported on host device.",
+                      PI_ERROR_INVALID_DEVICE);
 #endif // defined(__SYCL_DEVICE_ONLY__)
 }
 
@@ -337,10 +322,8 @@ inline __SYCL_ALWAYS_INLINE
   std::ignore = A;
   std::ignore = B;
   std::ignore = C;
-  throw runtime_error(
-      "This version of the matrix extension is only currently supported on "
-      "Nvidia devices",
-      PI_ERROR_INVALID_DEVICE);
+  throw runtime_error("joint matrix is not supported on host device.",
+                      PI_ERROR_INVALID_DEVICE);
 #endif // defined(__SYCL_DEVICE_ONLY__)
 }