Skip to content

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

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 4 commits into from
Jan 5, 2023
Merged

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

Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
fbec0b1
Removed unnecessary code. no functional changes.
JackAKirk Jan 5, 2023
a39043e
Add back required wi_data host version.
Jan 5, 2023
4c3b67f
Update runtime errors and removed old comments.
Jan 5, 2023
d215f0e
format.
Jan 5, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
39 changes: 11 additions & 28 deletions sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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__
Expand Down Expand Up @@ -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{};
Expand Down Expand Up @@ -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__)
}

Expand All @@ -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:
Expand Down Expand Up @@ -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__)
}

Expand Down Expand Up @@ -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__)
}

Expand All @@ -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:
Expand Down Expand Up @@ -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__)
}

Expand Down Expand Up @@ -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__)
}

Expand Down