| 304 | |
| 305 | template<typename T> |
| 306 | void swapIndex(Param<T> ovalues, Param<int> oindex, const Param<T> ivalues, |
| 307 | sycl::buffer<int> iindex, const Param<int> swapIdx) { |
| 308 | auto global = sycl::range(ovalues.info.dims[0]); |
| 309 | |
| 310 | getQueue().submit([&](auto &h) { |
| 311 | sycl::accessor d_ivalues{*ivalues.data, h, sycl::read_only}; |
| 312 | sycl::accessor d_iindex{iindex, h, sycl::read_only}; |
| 313 | sycl::accessor d_swapIdx{*swapIdx.data, h, sycl::read_only}; |
| 314 | sycl::accessor d_ovalues{*ovalues.data, h, sycl::write_only, |
| 315 | sycl::no_init}; |
| 316 | sycl::accessor d_oindex{*oindex.data, h, sycl::write_only, |
| 317 | sycl::no_init}; |
| 318 | |
| 319 | h.parallel_for(global, |
| 320 | swapIndexCreateKernel<T>( |
| 321 | d_ovalues, d_oindex, d_ivalues, d_iindex, d_swapIdx, |
| 322 | static_cast<int>(ovalues.info.dims[0]))); |
| 323 | }); |
| 324 | |
| 325 | ONEAPI_DEBUG_FINISH(getQueue()); |
| 326 | } |
| 327 | |
| 328 | template<typename T> |
| 329 | class csr2CooCreateKernel { |