| 77 | |
| 78 | template<typename T> |
| 79 | void coo2dense(Param<T> out, const Param<T> values, const Param<int> rowIdx, |
| 80 | const Param<int> colIdx) { |
| 81 | auto local = sycl::range(THREADS_PER_BLOCK, 1); |
| 82 | auto global = sycl::range( |
| 83 | divup(values.info.dims[0], local[0] * REPEAT) * THREADS_PER_BLOCK, 1); |
| 84 | |
| 85 | getQueue().submit([&](auto &h) { |
| 86 | sycl::accessor d_rowIdx{*rowIdx.data, h, sycl::read_only}; |
| 87 | sycl::accessor d_colIdx{*colIdx.data, h, sycl::read_only}; |
| 88 | sycl::accessor d_out{*out.data, h, sycl::write_only, sycl::no_init}; |
| 89 | sycl::accessor d_values{*values.data, h, sycl::write_only, |
| 90 | sycl::no_init}; |
| 91 | h.parallel_for(sycl::nd_range{global, local}, |
| 92 | coo2DenseCreateKernel<T>( |
| 93 | d_out, out.info, d_values, values.info, d_rowIdx, |
| 94 | rowIdx.info, d_colIdx, colIdx.info)); |
| 95 | }); |
| 96 | ONEAPI_DEBUG_FINISH(getQueue()); |
| 97 | } |
| 98 | |
| 99 | template<typename T, int THREADS> |
| 100 | class csr2DenseCreateKernel { |