| 27 | |
| 28 | template<typename T, typename Tr> |
| 29 | void svdInPlace(Array<Tr> &s, Array<T> &u, Array<T> &vt, Array<T> &in) { |
| 30 | dim4 iDims = in.dims(); |
| 31 | int64_t M = iDims[0]; |
| 32 | int64_t N = iDims[1]; |
| 33 | |
| 34 | dim4 iStrides = in.strides(); |
| 35 | dim4 uStrides = u.strides(); |
| 36 | dim4 vStrides = vt.strides(); |
| 37 | int64_t LDA = iStrides[1]; |
| 38 | int64_t LDU = uStrides[1]; |
| 39 | int64_t LDVt = vStrides[1]; |
| 40 | |
| 41 | int64_t scratch_size = |
| 42 | ::oneapi::mkl::lapack::gesvd_scratchpad_size<compute_t<T>>( |
| 43 | getQueue(), ::oneapi::mkl::jobsvd::vectors, |
| 44 | ::oneapi::mkl::jobsvd::vectors, M, N, LDA, LDU, LDVt); |
| 45 | |
| 46 | auto scratchpad = memAlloc<compute_t<T>>(scratch_size); |
| 47 | |
| 48 | sycl::buffer<compute_t<T>> in_buffer = |
| 49 | in.template getBufferWithOffset<compute_t<T>>(); |
| 50 | |
| 51 | sycl::buffer<compute_t<Tr>> sBuf = |
| 52 | s.template getBufferWithOffset<compute_t<Tr>>(); |
| 53 | sycl::buffer<compute_t<T>> uBuf = |
| 54 | u.template getBufferWithOffset<compute_t<T>>(); |
| 55 | sycl::buffer<compute_t<T>> vtBuf = |
| 56 | vt.template getBufferWithOffset<compute_t<T>>(); |
| 57 | |
| 58 | ::oneapi::mkl::lapack::gesvd(getQueue(), ::oneapi::mkl::jobsvd::vectors, |
| 59 | ::oneapi::mkl::jobsvd::vectors, M, N, |
| 60 | in_buffer, LDA, sBuf, uBuf, LDU, vtBuf, LDVt, |
| 61 | *scratchpad, scratchpad->size()); |
| 62 | } |
| 63 | |
| 64 | template<typename T, typename Tr> |
| 65 | void svd(Array<Tr> &s, Array<T> &u, Array<T> &vt, const Array<T> &in) { |