| 138 | |
| 139 | template<typename T> |
| 140 | void csr2dense(Param<T> output, const Param<T> values, const Param<int> rowIdx, |
| 141 | const Param<int> colIdx) { |
| 142 | constexpr int MAX_GROUPS = 4096; |
| 143 | // FIXME: This needs to be based non nonzeros per row |
| 144 | constexpr int threads = 64; |
| 145 | |
| 146 | const int M = rowIdx.info.dims[0] - 1; |
| 147 | |
| 148 | auto local = sycl::range(threads, 1); |
| 149 | int groups_x = std::min((int)(divup(M, local[0])), MAX_GROUPS); |
| 150 | auto global = sycl::range(local[0] * groups_x, 1); |
| 151 | |
| 152 | getQueue().submit([&](auto &h) { |
| 153 | sycl::accessor d_values{*values.data, h, sycl::read_only}; |
| 154 | sycl::accessor d_rowIdx{*rowIdx.data, h, sycl::read_only}; |
| 155 | sycl::accessor d_colIdx{*colIdx.data, h, sycl::read_only}; |
| 156 | sycl::accessor d_output{*output.data, h, sycl::write_only, |
| 157 | sycl::no_init}; |
| 158 | h.parallel_for(sycl::nd_range{global, local}, |
| 159 | csr2DenseCreateKernel<T, threads>( |
| 160 | d_output, d_values, d_rowIdx, d_colIdx, M, |
| 161 | static_cast<int>(values.info.offset), |
| 162 | static_cast<int>(rowIdx.info.offset), |
| 163 | static_cast<int>(colIdx.info.offset))); |
| 164 | }); |
| 165 | |
| 166 | ONEAPI_DEBUG_FINISH(getQueue()); |
| 167 | } |
| 168 | |
| 169 | template<typename T> |
| 170 | class dense2csrCreateKernel { |