| 151 | |
| 152 | template<typename T, typename accType> |
| 153 | void convSep(Param<T> out, const Param<T> signal, const Param<accType> filter, |
| 154 | const int conv_dim, const bool expand) { |
| 155 | if (!(conv_dim == 0 || conv_dim == 1)) { |
| 156 | AF_ERROR( |
| 157 | "Separable convolution accepts only 0 or 1 as convolution " |
| 158 | "dimension", |
| 159 | AF_ERR_NOT_SUPPORTED); |
| 160 | } |
| 161 | constexpr int THREADS_X = 16; |
| 162 | constexpr int THREADS_Y = 16; |
| 163 | |
| 164 | const int fLen = filter.info.dims[0] * filter.info.dims[1]; |
| 165 | const size_t C0_SIZE = (THREADS_X + 2 * (fLen - 1)) * THREADS_Y; |
| 166 | const size_t C1_SIZE = (THREADS_Y + 2 * (fLen - 1)) * THREADS_X; |
| 167 | size_t locSize = (conv_dim == 0 ? C0_SIZE : C1_SIZE); |
| 168 | |
| 169 | auto local = sycl::range(THREADS_X, THREADS_Y); |
| 170 | |
| 171 | int blk_x = divup(out.info.dims[0], THREADS_X); |
| 172 | int blk_y = divup(out.info.dims[1], THREADS_Y); |
| 173 | |
| 174 | auto global = sycl::range(blk_x * signal.info.dims[2] * THREADS_X, |
| 175 | blk_y * signal.info.dims[3] * THREADS_Y); |
| 176 | |
| 177 | sycl::buffer<accType> mBuff = {sycl::range(fLen * sizeof(accType))}; |
| 178 | memcpyBuffer(mBuff, *filter.data, fLen, 0); |
| 179 | |
| 180 | getQueue().submit([&](auto &h) { |
| 181 | sycl::accessor d_signal{*signal.data, h, sycl::read_only}; |
| 182 | sycl::accessor d_out{*out.data, h, sycl::write_only, sycl::no_init}; |
| 183 | sycl::accessor d_mBuff{mBuff, h, sycl::read_only}; |
| 184 | sycl::local_accessor<T> localMem(locSize, h); |
| 185 | h.parallel_for(sycl::nd_range{global, local}, |
| 186 | convolveSeparableCreateKernel<T, accType>( |
| 187 | d_out, out.info, d_signal, signal.info, d_mBuff, |
| 188 | blk_x, blk_y, fLen, conv_dim, expand, localMem)); |
| 189 | }); |
| 190 | } |
| 191 | |
| 192 | #define INSTANTIATE(T, accT) \ |
| 193 | template void convSep<T, accT>(Param<T>, const Param<T>, \ |
nothing calls this directly
no test coverage detected