26 #if defined(DATA_TYPE) && defined(NUM_REVERSE_DIMS)
28 #if NUM_REVERSE_DIMS > 4
29 #error("Reversing more than 4 dimensions is not currently supported")
76 const uint x_in = get_global_id(0);
77 const uint y_in = get_global_id(1);
78 const uint z_in = get_global_id(2) % depth;
79 const uint w_in = get_global_id(2) / depth;
81 const uint4 dims = (uint4)(0, 1, 2, 3);
82 int4 to_reverse = (int4)(0, 0, 0, 0);
85 #if defined(USE_INVERTED_AXIS)
86 indices =
select((
VEC_DATA_TYPE(
int, NUM_REVERSE_DIMS)) RANK - 1, -1, indices < 0) - indices;
88 indices =
select(indices, indices + RANK, indices < 0);
91 #if NUM_REVERSE_DIMS == 1
92 to_reverse = ((uint4)indices == dims);
93 #elif NUM_REVERSE_DIMS == 2
94 to_reverse = ((uint4)indices.s0 == dims) || ((uint4)indices.s1 == dims);
95 #elif NUM_REVERSE_DIMS == 3
96 to_reverse = ((uint4)indices.s0 == dims) || ((uint4)indices.s1 == dims) || ((uint4)indices.s2 == dims);
98 to_reverse = ((uint4)indices.s0 == dims) || ((uint4)indices.s1 == dims) || ((uint4)indices.s2 == dims) || ((uint4)indices.s3 == dims);
101 const uint x_out = to_reverse.s0 ? width - x_in - 1 : x_in;
102 const uint y_out = to_reverse.s1 ? height - y_in - 1 : y_in;
103 const uint z_out = to_reverse.s2 ? depth - z_in - 1 : z_in;
104 const uint w_out = to_reverse.s3 ?
batches - w_in - 1 : w_in;
106 *((__global DATA_TYPE *)
tensor4D_offset(&
dst, x_out, y_out, z_out, w_out)) = *((__global DATA_TYPE *)
src.ptr);
108 #endif // defined(DATA_TYPE) && defined(NUM_REVERSE_DIMS)