Function range

src/backend/oneapi/kernel/range.hpp:93–114

template<typename T>
void range(Param<T> out, const int dim) {
    constexpr int RANGE_TX    = 32;
    constexpr int RANGE_TY    = 8;
    constexpr int RANGE_TILEX = 512;
    constexpr int RANGE_TILEY = 32;

    sycl::range<2> local(RANGE_TX, RANGE_TY);

    int blocksPerMatX = divup(out.info.dims[0], RANGE_TILEX);
    int blocksPerMatY = divup(out.info.dims[1], RANGE_TILEY);
    sycl::range<2> global(local[0] * blocksPerMatX * out.info.dims[2],
                          local[1] * blocksPerMatY * out.info.dims[3]);
    sycl::nd_range<2> ndrange(global, local);

    getQueue().submit([&](sycl::handler& h) {
        write_accessor<T> out_acc{*out.data, h};

        h.parallel_for(ndrange, rangeOp<T>(out_acc, out.info, dim,
                                           blocksPerMatX, blocksPerMatY));
    });
    ONEAPI_DEBUG_FINISH(getQueue());
}

}  // namespace kernel
}  // namespace oneapi