MCPcopy Create free account
hub / github.com/arrayfire/arrayfire / padDataHelper

Function padDataHelper

src/backend/oneapi/kernel/fftconvolve_pad.hpp:88–119  ·  view source on GitHub ↗

Source from the content-addressed store, hash-verified

86
87template<typename convT, typename T>
88void 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

Callers

nothing calls this directly

Calls 3

calcParamSizesFunction · 0.70
rangeFunction · 0.70
getQueueFunction · 0.50

Tested by

no test coverage detected