Function csrCalcOutNNZ

src/backend/oneapi/kernel/sparse_arith.hpp:418–453

};

static void csrCalcOutNNZ(Param<int> outRowIdx, unsigned &nnzC, const uint M,
                          const uint N, uint nnzA, const Param<int> lrowIdx,
                          const Param<int> lcolIdx, uint nnzB,
                          const Param<int> rrowIdx, const Param<int> rcolIdx) {
    UNUSED(N);
    UNUSED(nnzA);
    UNUSED(nnzB);

    auto local  = sycl::range(256);
    auto global = sycl::range(divup(M, local[0]) * local[0]);

    Array<unsigned> out = createValueArray<unsigned>(1, 0);
    auto out_get = out.get();

    getQueue().submit([&](auto &h) {
        sycl::accessor d_out{*out_get, h, sycl::write_only};
        sycl::accessor d_outRowIdx{*outRowIdx.data, h, sycl::write_only};
        sycl::accessor d_lRowIdx{*lrowIdx.data, h, sycl::read_only};
        sycl::accessor d_lColIdx{*lcolIdx.data, h, sycl::read_only};
        sycl::accessor d_rRowIdx{*rrowIdx.data, h, sycl::read_only};
        sycl::accessor d_rColIdx{*rcolIdx.data, h, sycl::read_only};

        auto blkNNZ = sycl::local_accessor<unsigned, 1>(local[0], h);
        h.parallel_for(
            sycl::nd_range{global, local},
            csrCalcOutNNZKernel(d_out, d_outRowIdx, M, d_lRowIdx, d_lColIdx,
                                d_rRowIdx, d_rColIdx, blkNNZ));
    });

    {
        sycl::host_accessor nnz_acc{*out.get(), sycl::read_only};
        nnzC = nnz_acc[0];
    }

    ONEAPI_DEBUG_FINISH(getQueue());
}

template<typename T, af_op_t op>
class ssarithCSRKernel {