| 30 | |
| 31 | template<typename T> |
| 32 | class rangeOp { |
| 33 | public: |
| 34 | rangeOp(write_accessor<T> out, KParam oinfo, const int dim, |
| 35 | const int blocksPerMatX, const int blocksPerMatY) |
| 36 | : out_(out) |
| 37 | , oinfo_(oinfo) |
| 38 | , dim_(dim) |
| 39 | , blocksPerMatX_(blocksPerMatX) |
| 40 | , blocksPerMatY_(blocksPerMatY) {} |
| 41 | |
| 42 | void operator()(sycl::nd_item<2> it) const { |
| 43 | const int mul0 = (dim_ == 0); |
| 44 | const int mul1 = (dim_ == 1); |
| 45 | const int mul2 = (dim_ == 2); |
| 46 | const int mul3 = (dim_ == 3); |
| 47 | |
| 48 | sycl::group g = it.get_group(); |
| 49 | const int oz = g.get_group_id(0) / blocksPerMatX_; |
| 50 | const int ow = g.get_group_id(1) / blocksPerMatY_; |
| 51 | |
| 52 | const int blockIdx_x = g.get_group_id(0) - oz * blocksPerMatX_; |
| 53 | const int blockIdx_y = g.get_group_id(1) - ow * blocksPerMatY_; |
| 54 | |
| 55 | const int xx = it.get_local_id(0) + blockIdx_x * it.get_local_range(0); |
| 56 | const int yy = it.get_local_id(1) + blockIdx_y * it.get_local_range(1); |
| 57 | |
| 58 | const size_t odx = oinfo_.dims[0]; |
| 59 | const size_t ody = oinfo_.dims[1]; |
| 60 | const size_t odz = oinfo_.dims[2]; |
| 61 | const size_t odw = oinfo_.dims[3]; |
| 62 | |
| 63 | if (xx < odx && yy < ody && oz < odz && ow < odw) { |
| 64 | const int ozw = ow * oinfo_.strides[3] + oz * oinfo_.strides[2]; |
| 65 | |
| 66 | const int incy = blocksPerMatY_ * g.get_local_range(1); |
| 67 | const int incx = blocksPerMatX_ * g.get_local_range(0); |
| 68 | |
| 69 | compute_t<T> valZW = (mul3 * ow) + (mul2 * oz); |
| 70 | |
| 71 | T* optr = out_.get_pointer(); |
| 72 | for (int oy = yy; oy < oinfo_.dims[1]; oy += incy) { |
| 73 | compute_t<T> valYZW = valZW + (mul1 * oy); |
| 74 | int oyzw = ozw + oy * oinfo_.strides[1]; |
| 75 | for (int ox = xx; ox < oinfo_.dims[0]; ox += incx) { |
| 76 | int oidx = oyzw + ox; |
| 77 | compute_t<T> val = valYZW + (mul0 * ox); |
| 78 | |
| 79 | optr[oidx] = val; |
| 80 | } |
| 81 | } |
| 82 | } |
| 83 | } |
| 84 | |
| 85 | protected: |
| 86 | write_accessor<T> out_; |
| 87 | KParam oinfo_; |
| 88 | int dim_; |
| 89 | int blocksPerMatX_, blocksPerMatY_; |