[SYCL][CUDA] Implementation of matrix ext using new "unified" interface

[JackAKirk]

CUDA backend implementation using the "unified" matrix extension interface. The same interface will be used for a future Intel backend implementation of the matrix extension.

• New "unified" interface uses SYCL_EXT_ONEAPI_MATRIX_VERSION=4
• joint_matrix_load, joint_matrix_store, joint_matrix_mad and joint_matrix interfaces match the new spec from [SYCL][Spec] Update the matrix spec based on new use argument #6662
• Separated joint_matrix_* functions into new header matrix-unified.hpp: Intel backend implementations can be called from the same functions in the future.
• C++17 everywhere in line with [SYCL] Emit an error on attempt to compile in less than C++17 mode #6678
• Updated device code tests to use new interfaces
• Completely removed uint16 implementations that are replaced by bfloat16 that is being moved out of the experimental namespace
• Updated all CUDA runtime matrix tests here: [SYCL][CUDA] Unified matrix interface updated tests llvm-test-suite#1183

[JackAKirk]

/verify with intel/llvm-test-suite#1183

[dkhaldi]

LGTM

[JackAKirk]

@bader Can this be merged now that it has the two approvals?

AMD failures are unrelated.

Thanks

[JackAKirk]

/verify with intel/llvm-test-suite#1334

[bader]

GitHub says we need one more approval from @intel/llvm-reviewers-runtime team.

[JackAKirk]

GitHub says we need one more approval from @intel/llvm-reviewers-runtime team.

OK. It would be great if this can get a review quite quickly. I will be on holiday after tomorrow and we wanted to have this before the 2023.1 code freeze also. It would mean that we could publish the joint_matrix optimizations for SYCL-DNN and SYCL-BLAS for the 2023.1 release too.

[steffenlarsen]

Looks okay from a high-level perspective and @dkhaldi and @yubingex007-a11y have been thorough. 👍

[JackAKirk]

Looks okay from a high-level perspective and @dkhaldi and @yubingex007-a11y have been thorough. +1

Thanks for the review!

[yubingex007-a11y]

LGTM

[JackAKirk]

I just had to merge the sycl branch to resolve the conflict with c103a6a where the now unnecessary c++17 checks were removed.
This makes no difference to the patch.

[steffenlarsen]

/verify with intel/llvm-test-suite#1334

[JackAKirk]

/verify with intel/llvm-test-suite#1334

[steffenlarsen]

Verification failures in Windows CI are unrelated and have been reported:

• code_location_e2e in intel/llvm-test-suite times out on Windows #7740
• HostInteropTask in intel/llvm-test-suite time out on Windows #7741
• multiple-kernel-linking.cpp in intel/llvm-test-suite fails on Windows #7742
• Multiple intel/llvm-test-suite reduction tests failing with L0 on Windows #7743
• Multiple intel/llvm-test-suite DeviceLib tests failing with L0 on Windows #7744
• alloc_pinned_host_memory and group_async_copy intel/llvm-test-suite tests failing with L0 on Windows #7745

[yubingex007-a11y]

@AerialMantis @JackAKirk @dkhaldi
since we can't provide wi_data in both cuda&intel's header, i will make wi_data unified again and provide wi_data of host version, so the return type should be "decltype(auto)".

[yubingex007-a11y]

@AerialMantis @JackAKirk
sorry, i remember previously it is:

#if defined(__SYCL_DEVICE_ONLY__)
#if defined(__NVPTX__)
  sycl::ext::oneapi::detail::joint_matrix_cuda<T, Use, Rows, Cols, Layout>
      cuda_impl;
#else
  __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(__SYCL_DEVICE_ONLY__)
#endif

in intel side, we can't let host compilation use sycl::ext::oneapi::detail::joint_matrix_cuda. so i go back to the previous code and i can still get passed in cuda's testcases.