| 443 | |
| 444 | template<typename T> |
| 445 | void coo2csr(Param<T> ovalues, Param<int> orowIdx, Param<int> ocolIdx, |
| 446 | const Param<T> ivalues, const Param<int> irowIdx, |
| 447 | const Param<int> icolIdx, Param<int> index, Param<int> rowCopy, |
| 448 | const int M) { |
| 449 | // Now we need to sort this into column major |
| 450 | kernel::sort0ByKeyIterative<int, int>(rowCopy, index, true); |
| 451 | |
| 452 | // Now use index to sort values and rows |
| 453 | kernel::swapIndex<T>(ovalues, ocolIdx, ivalues, *icolIdx.data, index); |
| 454 | |
| 455 | ONEAPI_DEBUG_FINISH(getQueue()); |
| 456 | |
| 457 | auto global = sycl::range(irowIdx.info.dims[0]); |
| 458 | |
| 459 | getQueue().submit([&](auto &h) { |
| 460 | sycl::accessor d_orowIdx{*orowIdx.data, h, sycl::write_only}; |
| 461 | sycl::accessor d_rowCopy{*rowCopy.data, h, sycl::read_only}; |
| 462 | h.parallel_for( |
| 463 | sycl::range{global}, |
| 464 | csrReduceKernel<T>(d_orowIdx, d_rowCopy, M, |
| 465 | static_cast<int>(ovalues.info.dims[0]))); |
| 466 | }); |
| 467 | ONEAPI_DEBUG_FINISH(getQueue()); |
| 468 | } |
| 469 | |
| 470 | } // namespace kernel |
| 471 | } // namespace oneapi |