| 97 | |
| 98 | template<typename T> |
| 99 | void reorder(Param<T> out, const Param<T> in, const dim_t* rdims) { |
| 100 | constexpr int TX = 32; |
| 101 | constexpr int TY = 8; |
| 102 | constexpr int TILEX = 512; |
| 103 | constexpr int TILEY = 32; |
| 104 | |
| 105 | auto local = sycl::range(TX, TY); |
| 106 | |
| 107 | int blocksPerMatX = divup(out.info.dims[0], TILEX); |
| 108 | int blocksPerMatY = divup(out.info.dims[1], TILEY); |
| 109 | auto global = sycl::range(local[0] * blocksPerMatX * out.info.dims[2], |
| 110 | local[1] * blocksPerMatY * out.info.dims[3]); |
| 111 | |
| 112 | getQueue().submit([&](auto& h) { |
| 113 | read_accessor<T> d_in{*in.data, h}; |
| 114 | write_accessor<T> d_out{*out.data, h}; |
| 115 | h.parallel_for( |
| 116 | sycl::nd_range{global, local}, |
| 117 | reorderCreateKernel<T>( |
| 118 | d_out, d_in, out.info, in.info, static_cast<int>(rdims[0]), |
| 119 | static_cast<int>(rdims[1]), static_cast<int>(rdims[2]), |
| 120 | static_cast<int>(rdims[3]), blocksPerMatX, blocksPerMatY)); |
| 121 | }); |
| 122 | |
| 123 | ONEAPI_DEBUG_FINISH(getQueue()); |
| 124 | } |
| 125 | } // namespace kernel |
| 126 | } // namespace oneapi |
| 127 | } // namespace arrayfire |