Function packDataHelper

src/backend/oneapi/kernel/fftconvolve_pack.hpp:104–138

template<typename convT, typename T>
void packDataHelper(Param<convT> packed, Param<T> sig, Param<T> filter,
                    const int rank, AF_BATCH_KIND kind) {
    Param<T> sig_tmp, filter_tmp;
    calcParamSizes(sig_tmp, filter_tmp, packed, sig, filter, rank, kind);

    int sig_packed_elem = sig_tmp.info.strides[3] * sig_tmp.info.dims[3];

    // Number of packed complex elements in dimension 0
    int sig_half_d0     = divup(sig.info.dims[0], 2);
    int sig_half_d0_odd = sig.info.dims[0] % 2;

    int blocks = divup(sig_packed_elem, THREADS);

    // Locate features kernel sizes
    auto local  = sycl::range(THREADS);
    auto global = sycl::range(blocks * THREADS);

    // Treat complex output as an array of scalars
    using convScalarT    = typename convT::value_type;
    auto packed_num_elem = (*packed.data).get_range().size();
    auto sig_tmp_buffer  = (*packed.data)
                              .template reinterpret<convScalarT>(
                                  sycl::range<1>{packed_num_elem * 2});

    getQueue().submit([&](auto &h) {
        read_accessor<T> d_sig                = {*sig.data, h};
        write_accessor<convScalarT> d_sig_tmp = {sig_tmp_buffer, h};
        h.parallel_for(sycl::nd_range{global, local},
                       fftconvolve_packCreateKernel<T, convScalarT>(
                           d_sig_tmp, sig_tmp.info, d_sig, sig.info,
                           sig_half_d0, sig_half_d0_odd));
    });

    ONEAPI_DEBUG_FINISH(getQueue());
}

}  // namespace kernel
}  // namespace oneapi