| 22 | namespace oneapi { |
| 23 | template<typename T> |
| 24 | Array<T> iir(const Array<T> &b, const Array<T> &a, const Array<T> &x) { |
| 25 | AF_BATCH_KIND type = x.ndims() == 1 ? AF_BATCH_NONE : AF_BATCH_SAME; |
| 26 | if (x.ndims() != b.ndims()) { |
| 27 | type = (x.ndims() < b.ndims()) ? AF_BATCH_RHS : AF_BATCH_LHS; |
| 28 | } |
| 29 | |
| 30 | // Extract the first N elements |
| 31 | Array<T> c = convolve<T, T>(x, b, type, 1, true); |
| 32 | dim4 cdims = c.dims(); |
| 33 | cdims[0] = x.dims()[0]; |
| 34 | c.resetDims(cdims); |
| 35 | |
| 36 | int num_a = a.dims()[0]; |
| 37 | |
| 38 | if (num_a == 1) { return c; } |
| 39 | |
| 40 | size_t local_bytes_req = (num_a * 2 + 1) * sizeof(T); |
| 41 | if (local_bytes_req > |
| 42 | getDevice().get_info<sycl::info::device::local_mem_size>()) { |
| 43 | char errMessage[256]; |
| 44 | snprintf(errMessage, sizeof(errMessage), |
| 45 | "\ncurrent OneAPI device does not have sufficient local " |
| 46 | "memory,\n" |
| 47 | "for iir kernel, %zu(required) > %zu(available)\n", |
| 48 | local_bytes_req, |
| 49 | getDevice().get_info<sycl::info::device::local_mem_size>()); |
| 50 | AF_ERROR(errMessage, AF_ERR_RUNTIME); |
| 51 | } |
| 52 | |
| 53 | dim4 ydims = c.dims(); |
| 54 | Array<T> y = createEmptyArray<T>(ydims); |
| 55 | |
| 56 | if (a.ndims() > 1) { |
| 57 | kernel::iir<T, true>(y, c, a); |
| 58 | } else { |
| 59 | kernel::iir<T, false>(y, c, a); |
| 60 | } |
| 61 | return y; |
| 62 | } |
| 63 | |
| 64 | #define INSTANTIATE(T) \ |
| 65 | template Array<T> iir(const Array<T> &b, const Array<T> &a, \ |