| 123 | |
| 124 | template<typename T> |
| 125 | void rotate(Param<T> out, const Param<T> in, const float theta, |
| 126 | af_interp_type method, int order) { |
| 127 | using std::string; |
| 128 | |
| 129 | using BT = typename dtype_traits<T>::base_type; |
| 130 | |
| 131 | constexpr int TX = 16; |
| 132 | constexpr int TY = 16; |
| 133 | |
| 134 | // Used for batching images |
| 135 | constexpr int TI = 4; |
| 136 | |
| 137 | const float c = cos(-theta), s = sin(-theta); |
| 138 | float tx, ty; |
| 139 | { |
| 140 | const float nx = 0.5 * (in.info.dims[0] - 1); |
| 141 | const float ny = 0.5 * (in.info.dims[1] - 1); |
| 142 | const float mx = 0.5 * (out.info.dims[0] - 1); |
| 143 | const float my = 0.5 * (out.info.dims[1] - 1); |
| 144 | const float sx = (mx * c + my * -s); |
| 145 | const float sy = (mx * s + my * c); |
| 146 | tx = -(sx - nx); |
| 147 | ty = -(sy - ny); |
| 148 | } |
| 149 | |
| 150 | // Rounding error. Anything more than 3 decimal points wont make a diff |
| 151 | tmat_t t; |
| 152 | t.tmat[0] = round(c * 1000) / 1000.0f; |
| 153 | t.tmat[1] = round(-s * 1000) / 1000.0f; |
| 154 | t.tmat[2] = round(tx * 1000) / 1000.0f; |
| 155 | t.tmat[3] = round(s * 1000) / 1000.0f; |
| 156 | t.tmat[4] = round(c * 1000) / 1000.0f; |
| 157 | t.tmat[5] = round(ty * 1000) / 1000.0f; |
| 158 | |
| 159 | auto local = sycl::range(TX, TY); |
| 160 | |
| 161 | int nimages = in.info.dims[2]; |
| 162 | int nbatches = in.info.dims[3]; |
| 163 | int global_x = local[0] * divup(out.info.dims[0], local[0]); |
| 164 | int global_y = local[1] * divup(out.info.dims[1], local[1]); |
| 165 | const int blocksXPerImage = global_x / local[0]; |
| 166 | const int blocksYPerImage = global_y / local[1]; |
| 167 | |
| 168 | if (nimages > TI) { |
| 169 | int tile_images = divup(nimages, TI); |
| 170 | nimages = TI; |
| 171 | global_x = global_x * tile_images; |
| 172 | } |
| 173 | global_y *= nbatches; |
| 174 | |
| 175 | auto global = sycl::range(global_x, global_y); |
| 176 | |
| 177 | getQueue().submit([&](auto &h) { |
| 178 | read_accessor<T> d_in{*in.data, h}; |
| 179 | write_accessor<T> d_out{*out.data, h}; |
| 180 | switch (order) { |
| 181 | case 1: |
| 182 | h.parallel_for( |