/*******************************************************************************
* Copyright 2021 Intel Corporation
*
* Licensed under the Apache License, Version 2.0 (the "License")* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*******************************************************************************/
#define USE_CUSTOM_GWS_GET_ID
// Temporary W/A for bf16 problems in HW and compiler
#undef cl_future_bf16_cvt
#include "gpu/intel/include/dispatch.h"
#include "gpu/intel/include/types_interop.h"
#include "gpu/intel/reorder/common.h"
#define GWS_GET_THREAD_ID(index) (get_global_id(index) + offset.array[index])
KERNEL_ATTR
__kernel void generic_reorder(__global SRC_DATA_T *restrict src,
__global DST_DATA_T *restrict dst, __global float *restrict src_scales,
__global int *restrict src_zps, __global float *restrict dst_scales,
__global int *restrict dst_zps, float sum_scale, int sum_zp,
int64x3_t offset) {
const int src_zp = GET_SRC_ZP(src_zps) const int dst_zp = GET_DST_ZP(dst_zps) float src_scale = 1.0f float dst_scale = 1.0f
src += SRC_OFFSET0 dst += DST_OFFSET0
#define LOOP_NEST_LEVEL 4
const uint sgId = get_sub_group_local_id() off_t d[6] off_t b[6]
= {0, 0, 0, 0, 0, 0}
d[0] = GWS_GET_D0() d[1] = GWS_GET_D1() d[2] = GWS_GET_D2() d[3] = GWS_GET_D3() d[4] = GWS_GET_D4() d[5] = GWS_GET_D5()
// Dispatcher code does not allow vectorization of dimensions that are not
// divisible by 16. As workaround, we cheat dispatcher by pretending
// vectorized dim is 16x larger than it really is. Here we adjust its size
// back to original.
d[VECT_DIM] /= RESCALE_COEFF // sg_off = offset into 'cache' local mem for given subgroup.
// Local memory will be split by subgroups so that given address in local
// can only be accessed by single subgroup. This lets us avoid barriers.
const uint cache_size_per_sg = D_BLK_SIZE_0 * D_BLK_SIZE_1 * D_BLK_SIZE_2
* D_BLK_SIZE_3 * VECT_SIZE const uint sg_off = get_sub_group_id() * cache_size_per_sg
// TODO: decide whether to store cache as SRC_DATA_T or DST_DATA_T
__local SRC_DATA_T cache[SG_PER_WG * cache_size_per_sg] uint iter[LOOP_NEST_LEVEL] = {0, 0, 0, 0}
// Loop across dimensions described in src_block.
// Example: block 2a4c2b would mean:
// for(a = 0..1) { for (c = 0..3) { for (b = 0..2) {}}}
#if S_BLK_SIZE_3 > 1
for_(iter[3] = 0#endif
#if S_BLK_SIZE_2 > 1
for_(iter[2] = 0#endif
#if S_BLK_SIZE_1 > 1
for_(iter[1] = 0#endif
#if S_BLK_SIZE_0 > 1
for_(iter[0] = 0#endif
{
// the same IDX could be in more than one loop, this makes offset calculation tricky
b[0] = 0 b[1] = 0 b[2] = 0 b[3] = 0 b[4] = 0 b[5] = 0 b[S_BLK_IDX_0] += iter[0] * S_BLK_STEP_0 b[S_BLK_IDX_1] += iter[1] * S_BLK_STEP_1 b[S_BLK_IDX_2] += iter[2] * S_BLK_STEP_2 b[S_BLK_IDX_3] += iter[3] * S_BLK_STEP_3
#if S_MOD_3 > 1
b[S_IDX_3] += S_MUL_3 * ((sgId / S_DIV_3) % S_MOD_3)#endif
#if S_MOD_2 > 1
b[S_IDX_2] += S_MUL_2 * ((sgId / S_DIV_2) % S_MOD_2)#endif
#if S_MOD_1 > 1
b[S_IDX_1] += S_MUL_1 * ((sgId / S_DIV_1) % S_MOD_1)#endif
#if S_MOD_0 > 1
b[S_IDX_0] += S_MUL_0 * ((sgId / S_DIV_0) % S_MOD_0)#endif
// In a majority of cases, src_off contains neighboring values for
// neighboring workitems. Exceptions occur when the src tensor has
// additional striding in some dimension(s).
const off_t src_off = SRC_OFF(d[0] + b[0], d[1] + b[1], d[2] + b[2],
d[3] + b[3], d[4] + b[4], d[5] + b[5])
// Data in cache (local mem) is organized as if it had 'fedcba' format
// tag. This is neither src's nor dst's format so some performance is
// wasted here, but otherwise the logic to calculate offsets would be
// too complicated.
uint cache_idx = sg_off + b[5] * CACHE_STRIDE_5 + b[4] * CACHE_STRIDE_4
+ b[3] * CACHE_STRIDE_3 + b[2] * CACHE_STRIDE_2
+ b[1] * CACHE_STRIDE_1 + b[0] * CACHE_STRIDE_0 const int pad_d0 = d[0] + b[0] >= SRC_D0 const int pad_d1 = NDIMS > 1 && d[1] + b[1] >= SRC_D1 const int pad_d2 = NDIMS > 2 && d[2] + b[2] >= SRC_D2 const int pad_d3 = NDIMS > 3 && d[3] + b[3] >= SRC_D3 const int pad_d4 = NDIMS > 4 && d[4] + b[4] >= SRC_D4 const int pad_d5 = NDIMS > 5 && d[5] + b[5] >= SRC_D5 const bool pad_sgid = sgId >= LIMIT_SSGID const int pad
= pad_d0 || pad_d1 || pad_d2 || pad_d3 || pad_d4 || pad_d5 if (!pad_sgid) {
SRC_DATA_T src_tmp = pad ? 0 : src[src_off] cache[cache_idx] = src_tmp }
}
for (uint i = 0 iter[i] = 0 }
#if D_BLK_SIZE_3 > 1
for_(iter[3] = 0#endif
#if D_BLK_SIZE_2 > 1
for_(iter[2] = 0#endif
#if D_BLK_SIZE_1 > 1
for_(iter[1] = 0#endif
#if D_BLK_SIZE_0 > 1
for_(iter[0] = 0#endif
{
// the same IDX could be in more than one loop, this makes offset calculation tricky
b[0] = 0 b[1] = 0 b[2] = 0 b[3] = 0 b[4] = 0 b[5] = 0 b[D_BLK_IDX_0] += iter[0] * D_BLK_STEP_0 b[D_BLK_IDX_1] += iter[1] * D_BLK_STEP_1 b[D_BLK_IDX_2] += iter[2] * D_BLK_STEP_2 b[D_BLK_IDX_3] += iter[3] * D_BLK_STEP_3
#if D_MOD_3 > 1
b[D_IDX_3] += D_MUL_3 * ((sgId / D_DIV_3) % D_MOD_3)#endif
#if D_MOD_2 > 1
b[D_IDX_2] += D_MUL_2 * ((sgId / D_DIV_2) % D_MOD_2)#endif
#if D_MOD_1 > 1
b[D_IDX_1] += D_MUL_1 * ((sgId / D_DIV_1) % D_MOD_1)#endif
#if D_MOD_0 > 1
b[D_IDX_0] += D_MUL_0 * ((sgId / D_DIV_0) % D_MOD_0)#endif
// In a majority of cases, dst_off contains neighboring values for
// neighboring workitems. Exceptions occur when the dst tensor has
// additional striding in some dimension(s), e.g., when reorder is used
// in reference concat.
const off_t dst_off = DST_OFF(d[0] + b[0], d[1] + b[1], d[2] + b[2],
d[3] + b[3], d[4] + b[4], d[5] + b[5])
DST_DATA_T dst_tmp uint cache_idx = sg_off + b[5] * CACHE_STRIDE_5 + b[4] * CACHE_STRIDE_4
+ b[3] * CACHE_STRIDE_3 + b[2] * CACHE_STRIDE_2
+ b[1] * CACHE_STRIDE_1 + b[0] * CACHE_STRIDE_0
const int pad_d0 = d[0] + b[0] >= DST_PD0 const int pad_d1 = NDIMS > 1 && d[1] + b[1] >= DST_PD1 const int pad_d2 = NDIMS > 2 && d[2] + b[2] >= DST_PD2 const int pad_d3 = NDIMS > 3 && d[3] + b[3] >= DST_PD3 const int pad_d4 = NDIMS > 4 && d[4] + b[4] >= DST_PD4 const int pad_d5 = NDIMS > 5 && d[5] + b[5] >= DST_PD5 const bool pad_sgid = sgId >= LIMIT_DSGID const int pad = pad_d0 || pad_d1 || pad_d2 || pad_d3 || pad_d4 || pad_d5
|| pad_sgid
if (!pad) {
SRC_DATA_T from_cache = cache[cache_idx]#if WITH_SUM_SCALE || WITH_SUM_ZPOINT
// TODO: move to separate loop to enable burst reads from dst?
dst_tmp = dst[dst_off]#endif
#if WITH_SRC_SCALE
off_t src_scale_idx = SCALE_OFF(SRC, d[0] + b[0], d[1] + b[1],
d[2] + b[2], d[3] + b[3], d[4] + b[4], d[5] + b[5]) src_scale = src_scale_idx < SRC_NUM_SCALES
? src_scales[src_scale_idx]
: 0.0f#endif
#if WITH_DST_SCALE
off_t dst_scale_idx = SCALE_OFF(DST, d[0] + b[0], d[1] + b[1],
d[2] + b[2], d[3] + b[3], d[4] + b[4], d[5] + b[5]) dst_scale = dst_scale_idx < DST_NUM_SCALES
? dst_scales[dst_scale_idx]
: 0.0f#endif
REORDER(DEFAULT_ROUND, dst_tmp, from_cache, src_scale, dst_scale,
sum_scale, src_zp, dst_zp, sum_zp) dst[dst_off] = dst_tmp }
}
}