26 #if defined(DATA_TYPE) && defined(AXIS)
64 const int px = get_global_id(0);
65 const int py = get_global_id(1);
66 const int pz = get_global_id(2) % OUTPUT_DIM_Z;
67 const int pw = (get_global_id(2) / OUTPUT_DIM_Z );
75 const uint index = *(__global
const uint *)
tensor4D_offset(&indices, px, 0, 0, 0);
76 const uint safe_index =
select((uint)0, index, index < INDEX_LIMIT);
78 #elif INDICES_DIMS == 2
79 const uint index = *(__global
const uint *)
tensor4D_offset(&indices, px, py, 0, 0);
80 const uint safe_index =
select((uint)0, index, index < INDEX_LIMIT);
82 #elif INDICES_DIMS == 3
83 const uint index = *(__global
const uint *)
tensor4D_offset(&indices, px, py, pz, 0);
84 const uint safe_index =
select((uint)0, index, index < INDEX_LIMIT);
86 #elif INDICES_DIMS == 4
87 const uint index = *(__global
const uint *)
tensor4D_offset(&indices, px, py, pz, pw);
88 const uint safe_index =
select((uint)0, index, index < INDEX_LIMIT);
94 const uint index = *(__global
const uint *)
tensor4D_offset(&indices, py, 0, 0, 0);
95 const uint safe_index =
select((uint)0, index, index < INDEX_LIMIT);
97 #elif INDICES_DIMS == 2
98 const uint index = *(__global
const uint *)
tensor4D_offset(&indices, py, pz, 0, 0);
99 const uint safe_index =
select((uint)0, index, index < INDEX_LIMIT);
101 #elif INDICES_DIMS == 3
102 const uint index = *(__global
const uint *)
tensor4D_offset(&indices, py, pz, pw, 0);
103 const uint safe_index =
select((uint)0, index, index < INDEX_LIMIT);
105 #endif //INDICES_DIMS
108 #if INDICES_DIMS == 1
109 const uint index = *(__global
const uint *)
tensor4D_offset(&indices, pz, 0, 0, 0);
110 const uint safe_index =
select((uint)0, index, index < INDEX_LIMIT);
112 #elif INDICES_DIMS == 2
113 const uint index = *(__global
const uint *)
tensor4D_offset(&indices, pz, pw, 0, 0);
114 const uint safe_index =
select((uint)0, index, index < INDEX_LIMIT);
116 #endif //INDICES_DIMS
119 #if INDICES_DIMS == 1
120 const uint index = *(__global
const uint *)
tensor4D_offset(&indices, pw, 0, 0, 0);
121 const uint safe_index =
select((uint)0, index, index < INDEX_LIMIT);
123 #endif //INDICES_DIMS
127 *(__global DATA_TYPE *)output.
ptr =
select((DATA_TYPE)0, *((__global
const DATA_TYPE *)input_addr), (DATA_TYPE)(index < INDEX_LIMIT));
130 #endif //defined(DATA_TYPE) && defined(AXIS)