| 363 | |
| 364 | template<typename T> |
| 365 | void csr2coo(Param<T> ovalues, Param<int> orowIdx, Param<int> ocolIdx, |
| 366 | const Param<T> ivalues, const Param<int> irowIdx, |
| 367 | const Param<int> icolIdx, Param<int> index) { |
| 368 | const int MAX_GROUPS = 4096; |
| 369 | int M = irowIdx.info.dims[0] - 1; |
| 370 | // FIXME: This needs to be based non nonzeros per row |
| 371 | int threads = 64; |
| 372 | |
| 373 | auto scratch = memAlloc<int>(orowIdx.info.dims[0]); |
| 374 | |
| 375 | auto local = sycl::range(threads, 1); |
| 376 | int groups_x = std::min((int)(divup(M, local[0])), MAX_GROUPS); |
| 377 | auto global = sycl::range(local[0] * groups_x, 1); |
| 378 | |
| 379 | getQueue().submit([&](auto &h) { |
| 380 | sycl::accessor d_irowIdx{*irowIdx.data, h, sycl::read_only}; |
| 381 | sycl::accessor d_icolIdx{*icolIdx.data, h, sycl::read_only}; |
| 382 | sycl::accessor d_scratch{*scratch, h, sycl::write_only, sycl::no_init}; |
| 383 | sycl::accessor d_ocolIdx{*ocolIdx.data, h, sycl::write_only, |
| 384 | sycl::no_init}; |
| 385 | h.parallel_for(sycl::nd_range{global, local}, |
| 386 | csr2CooCreateKernel<T>(d_scratch, d_ocolIdx, d_irowIdx, |
| 387 | d_icolIdx, M)); |
| 388 | }); |
| 389 | |
| 390 | // Now we need to sort this into column major |
| 391 | kernel::sort0ByKeyIterative<int, int>(ocolIdx, index, true); |
| 392 | |
| 393 | // Now use index to sort values and rows |
| 394 | kernel::swapIndex<T>(ovalues, orowIdx, ivalues, *scratch, index); |
| 395 | |
| 396 | ONEAPI_DEBUG_FINISH(getQueue()); |
| 397 | } |
| 398 | |
| 399 | template<typename T> |
| 400 | class csrReduceKernel { |