| 223 | |
| 224 | template<typename T> |
| 225 | void dense2csr(Param<T> values, Param<int> rowIdx, Param<int> colIdx, |
| 226 | const Param<T> dense) { |
| 227 | int num_rows = dense.info.dims[0]; |
| 228 | int num_cols = dense.info.dims[1]; |
| 229 | |
| 230 | // sd1 contains output of scan along dim 1 of dense |
| 231 | Array<int> sd1 = createEmptyArray<int>(dim4(num_rows, num_cols)); |
| 232 | // rd1 contains output of nonzero count along dim 1 along dense |
| 233 | Array<int> rd1 = createEmptyArray<int>(num_rows); |
| 234 | |
| 235 | scan_dim<T, int, af_notzero_t, 1>(sd1, dense, true); |
| 236 | reduce_dim_default<T, int, af_notzero_t, 1>(rd1, dense, 0, 0); |
| 237 | scan_first<int, int, af_add_t>(rowIdx, rd1, false); |
| 238 | |
| 239 | const int nnz = values.info.dims[0]; |
| 240 | |
| 241 | const sycl::id<1> fillOffset(rowIdx.info.offset + |
| 242 | (rowIdx.info.dims[0] - 1)); |
| 243 | const sycl::range<1> fillRange(rowIdx.info.dims[0] - fillOffset[0]); |
| 244 | getQueue().submit([&](auto &h) { |
| 245 | sycl::accessor d_rowIdx{*rowIdx.data, h, fillRange, fillOffset}; |
| 246 | h.fill(d_rowIdx, nnz); |
| 247 | }); |
| 248 | |
| 249 | auto local = sycl::range(THREADS_X, THREADS_Y); |
| 250 | int groups_x = divup(dense.info.dims[0], local[0]); |
| 251 | int groups_y = divup(dense.info.dims[1], local[1]); |
| 252 | auto global = sycl::range(groups_x * local[0], groups_y * local[1]); |
| 253 | |
| 254 | const Param<int> sdParam = sd1; |
| 255 | |
| 256 | getQueue().submit([&](auto &h) { |
| 257 | sycl::accessor d_dense{*dense.data, h, sycl::read_only}; |
| 258 | sycl::accessor d_sdParam{*sdParam.data, h, sycl::read_only}; |
| 259 | sycl::accessor d_rowIdx{*rowIdx.data, h, sycl::read_only}; |
| 260 | sycl::accessor d_values{*values.data, h, sycl::write_only, |
| 261 | sycl::no_init}; |
| 262 | sycl::accessor d_colIdx{*colIdx.data, h, sycl::write_only, |
| 263 | sycl::no_init}; |
| 264 | h.parallel_for( |
| 265 | sycl::nd_range{global, local}, |
| 266 | dense2csrCreateKernel<T>(d_values, d_colIdx, d_dense, dense.info, |
| 267 | d_sdParam, sdParam.info, d_rowIdx)); |
| 268 | }); |
| 269 | |
| 270 | ONEAPI_DEBUG_FINISH(getQueue()); |
| 271 | } |
| 272 | |
| 273 | template<typename T> |
| 274 | class swapIndexCreateKernel { |