| 86 | |
| 87 | template<typename convT, typename T> |
| 88 | void padDataHelper(Param<convT> packed, Param<T> sig, Param<T> filter, |
| 89 | const int rank, AF_BATCH_KIND kind) { |
| 90 | Param<T> sig_tmp, filter_tmp; |
| 91 | calcParamSizes(sig_tmp, filter_tmp, packed, sig, filter, rank, kind); |
| 92 | |
| 93 | int filter_packed_elem = |
| 94 | filter_tmp.info.strides[3] * filter_tmp.info.dims[3]; |
| 95 | |
| 96 | int blocks = divup(filter_packed_elem, THREADS); |
| 97 | |
| 98 | // Locate features kernel sizes |
| 99 | auto local = sycl::range(THREADS); |
| 100 | auto global = sycl::range(blocks * THREADS); |
| 101 | |
| 102 | // Treat complex output as an array of scalars |
| 103 | using convScalarT = typename convT::value_type; |
| 104 | auto packed_num_elem = (*packed.data).get_range().size(); |
| 105 | auto filter_tmp_buffer = (*packed.data) |
| 106 | .template reinterpret<convScalarT>( |
| 107 | sycl::range<1>{packed_num_elem * 2}); |
| 108 | |
| 109 | getQueue().submit([&](auto &h) { |
| 110 | read_accessor<T> d_filter = {*filter.data, h, sycl::read_only}; |
| 111 | write_accessor<convScalarT> d_filter_tmp = {filter_tmp_buffer, h}; |
| 112 | h.parallel_for( |
| 113 | sycl::nd_range{global, local}, |
| 114 | fftconvolve_padCreateKernel<T, convScalarT>( |
| 115 | d_filter_tmp, filter_tmp.info, d_filter, filter.info)); |
| 116 | }); |
| 117 | |
| 118 | ONEAPI_DEBUG_FINISH(getQueue()); |
| 119 | } |
| 120 | } // namespace kernel |
| 121 | } // namespace oneapi |
| 122 | } // namespace arrayfire |
nothing calls this directly
no test coverage detected