mirror of
https://github.com/opencv/opencv.git
synced 2026-07-29 15:23:05 +04:00
Alexander Alekhin
1c8ee3f957
DNN: OpenCL/slice update * dnn(ocl/slice): make slice kernel VTune friendly - more unique names - inline code of copy functions * dnn(ocl/slice): prefer to spawn more work groups - even in case with 1D copy - perf improvement up to 2x of kernel time (due to changed configuration 128x1x1 => 128x32x1) * dnn(ocl/slice): cache kernel exec info
265 lines
7.3 KiB
Common Lisp
265 lines
7.3 KiB
Common Lisp
// This file is part of OpenCV project.
|
|
// It is subject to the license terms in the LICENSE file found in the top-level directory
|
|
// of this distribution and at http://opencv.org/license.html.
|
|
|
|
// Copyright (C) 2020, Intel Corporation, all rights reserved.
|
|
// Third party copyrights are property of their respective owners.
|
|
|
|
/*
|
|
Specialization constants:
|
|
- WSZ: size of OpenCL local group
|
|
- DIMS: number of working dimensions
|
|
- ELEMSIZE: element size in bytes
|
|
- DST_SZ_<i>: dst sizes
|
|
- SRC_START_<i>: src index shift (slice .start value)
|
|
- SRC_STEP_<i>: src steps (bytes)
|
|
- DST_STEP_<i>: dst steps (bytes), derived from DST_SZ_<i> and ELEMSIZE
|
|
- BLOCK_DIMS: number of dims for copy block (argmax(count(SRC_STEP_<i> != DST_STEP_<i>) <= 1))
|
|
- BLOCK_DIMS_CONTIGUOUS (<= BLOCK_DIMS): SRC_STEP_<i> == DST_STEP_<i> for i in [0, BLOCK_DIMS_CONTIGUOUS)
|
|
|
|
derived specialization constants:
|
|
- BLOCK_SIZE: ELEMSIZE * mul(DST_SZ_<i>) for i in [0, BLOCK_DIMS)
|
|
|
|
- USE_COPY_1D iff BLOCK_DIMS == BLOCK_DIMS_CONTIGUOUS
|
|
- BLOCK_COLS:
|
|
* with USE_COPY_1D: BLOCK_SIZE
|
|
* w/o USE_COPY_1D: ELEMSIZE * mul(DST_SZ_<i>) for i in [0, BLOCK_DIMS_CONTIGUOUS)
|
|
- BLOCK_ROWS:
|
|
* with USE_COPY_1D: N/A
|
|
* w/o USE_COPY_1D: ELEMSIZE * mul(DST_SZ_<i>) for i in [BLOCK_DIMS_CONTIGUOUS, BLOCK_DIMS)
|
|
- BLOCK_SRC_STRIDE:
|
|
* with USE_COPY_1D: N/A
|
|
* w/o USE_COPY_1D: ELEMSIZE * mul(SRC_STEP_<i>) for i in [0, BLOCK_DIMS_CONTIGUOUS)
|
|
|
|
Note: SZ, STEP values are in reversed order than OpenCV Mat:
|
|
- NCHW SZ: [cols, rows, channels, batch]
|
|
- NCHW STEP: [elemsize, cols * elemsize, rows * cols * elemsize, ...] (DIMS+1 value)
|
|
|
|
*/
|
|
|
|
/*
|
|
local: <WSZ, 1, 1>
|
|
global: <WSZ, number_of_copy_blocks, 1>
|
|
*/
|
|
|
|
#define CONCAT_(A, B) A##B
|
|
#define CONCAT(A, B) CONCAT_(A, B)
|
|
|
|
#define BLOCK_COLS_X4 (BLOCK_COLS / 4)
|
|
#define BLOCK_COLS_X16 (BLOCK_COLS / 16)
|
|
|
|
__attribute__((reqd_work_group_size(WSZ, 1, 1)))
|
|
__kernel void
|
|
CONCAT(slice_, SLICE_KERNEL_SUFFIX)(
|
|
__global const uchar* src0,
|
|
__global uchar* dst0
|
|
)
|
|
{
|
|
uint block_id = get_global_id(1);
|
|
uint dst_offset0 = block_id * BLOCK_SIZE;
|
|
uint src_offset0 = 0;
|
|
|
|
{ // calculate src_offset0
|
|
|
|
#define CALC_SRC_INDEX(dim) \
|
|
{ \
|
|
uint plane_sz = CONCAT(DST_STEP_, dim) / BLOCK_SIZE; \
|
|
CONCAT(idx_, dim) = block_id / plane_sz; \
|
|
block_id = block_id - CONCAT(idx_, dim) * plane_sz; \
|
|
}
|
|
#define UPDATE_SRC_OFFSET(dim) \
|
|
src_offset0 = mad24((uint)(CONCAT(idx_, dim) + CONCAT(SRC_START_, dim)), (uint)CONCAT(SRC_STEP_, dim), (uint)src_offset0);
|
|
/*
|
|
if (get_global_id(0) == 0 && get_global_id(1) == 0) \
|
|
printf("(%d, %d): @%d src_offset0=%d idx_dim=%d block_id=%d\n", \
|
|
get_global_id(0), get_global_id(1), \
|
|
dim, src_offset0, CONCAT(idx_, dim), block_id \
|
|
);
|
|
*/
|
|
|
|
#if DIMS > 5
|
|
#error "invalid configuration"
|
|
#endif
|
|
#if DIMS > 4
|
|
uint idx_4 = 0;
|
|
#if BLOCK_DIMS <= 4
|
|
CALC_SRC_INDEX(4)
|
|
#endif
|
|
UPDATE_SRC_OFFSET(4)
|
|
#endif
|
|
#if DIMS > 3
|
|
uint idx_3 = 0;
|
|
#if BLOCK_DIMS <= 3
|
|
CALC_SRC_INDEX(3)
|
|
#endif
|
|
UPDATE_SRC_OFFSET(3)
|
|
#endif
|
|
#if DIMS > 2
|
|
uint idx_2 = 0;
|
|
#if BLOCK_DIMS <= 2
|
|
CALC_SRC_INDEX(2)
|
|
#endif
|
|
UPDATE_SRC_OFFSET(2)
|
|
#endif
|
|
#if DIMS > 1
|
|
uint idx_1 = 0;
|
|
#if BLOCK_DIMS <= 1
|
|
CALC_SRC_INDEX(1)
|
|
#endif
|
|
UPDATE_SRC_OFFSET(1)
|
|
#endif
|
|
#if DIMS > 0
|
|
uint idx_0 = 0;
|
|
UPDATE_SRC_OFFSET(0)
|
|
#endif
|
|
|
|
/*
|
|
if (get_global_id(0) == 0)
|
|
printf("(%d, %d): src_offset0=%d dst_offset0=%d\n",
|
|
get_global_id(0), get_global_id(1),
|
|
src_offset0, dst_offset0
|
|
);
|
|
*/
|
|
|
|
} // calculate src_offset0
|
|
|
|
#ifdef USE_COPY_1D
|
|
{ // copy_block_1d
|
|
__global const uchar* src = src0 + src_offset0;
|
|
__global uchar* dst = dst0 + dst_offset0;
|
|
|
|
uint processed = 0;
|
|
|
|
#if BLOCK_COLS_X16 >= 4
|
|
{
|
|
// uchar16 x 4rows per iteration
|
|
uint i = get_local_id(0) * 16; // uchar16
|
|
while (i < BLOCK_COLS_X16 * 16)
|
|
{
|
|
uint4 idx0 = (uint4)i;
|
|
uint4 idx = idx0 + (uint4)(0, 16 * WSZ, 32 * WSZ, 48 * WSZ);
|
|
idx = select(idx0, idx, idx < (BLOCK_COLS_X16 * 16));
|
|
|
|
uchar16 a0 = vload16(0, src + idx.s0);
|
|
uchar16 a1 = vload16(0, src + idx.s1);
|
|
uchar16 a2 = vload16(0, src + idx.s2);
|
|
uchar16 a3 = vload16(0, src + idx.s3);
|
|
|
|
vstore16(a0, 0, dst + idx.s0);
|
|
vstore16(a1, 0, dst + idx.s1);
|
|
vstore16(a2, 0, dst + idx.s2);
|
|
vstore16(a3, 0, dst + idx.s3);
|
|
|
|
i += WSZ * 16 * 4;
|
|
}
|
|
processed = BLOCK_COLS_X16 * 16;
|
|
}
|
|
#else
|
|
#define SKIP_1D_BLOCK_COLS_X16 1
|
|
#endif
|
|
|
|
#if BLOCK_COLS_X4 > 0 && (defined(SKIP_1D_BLOCK_COLS_X16) || (BLOCK_COLS_X16 * 16 != BLOCK_COLS_X4 * 4))
|
|
{
|
|
// uchar4 x 4rows per iteration
|
|
uint i = get_local_id(0) * 4 + processed; // uchar4
|
|
while (i < BLOCK_COLS_X4 * 4)
|
|
{
|
|
uint4 idx0 = (uint4)i;
|
|
uint4 idx = idx0 + (uint4)(0, 4 * WSZ, 8 * WSZ, 12 * WSZ);
|
|
idx = select(idx0, idx, idx < (BLOCK_COLS_X4 * 4));
|
|
|
|
uchar4 a0 = vload4(0, src + idx.s0);
|
|
uchar4 a1 = vload4(0, src + idx.s1);
|
|
uchar4 a2 = vload4(0, src + idx.s2);
|
|
uchar4 a3 = vload4(0, src + idx.s3);
|
|
|
|
vstore4(a0, 0, dst + idx.s0);
|
|
vstore4(a1, 0, dst + idx.s1);
|
|
vstore4(a2, 0, dst + idx.s2);
|
|
vstore4(a3, 0, dst + idx.s3);
|
|
|
|
i += WSZ * 4 * 4;
|
|
}
|
|
processed = BLOCK_COLS_X4 * 4;
|
|
}
|
|
#else
|
|
#define SKIP_1D_BLOCK_COLS_X4 1
|
|
#endif // BLOCK_COLS_X4 > 0
|
|
|
|
#if (defined(SKIP_1D_BLOCK_COLS_X16) && defined(SKIP_1D_BLOCK_COLS_X4)) || BLOCK_COLS_X4 * 4 != BLOCK_COLS
|
|
{
|
|
uint i = get_local_id(0) + processed;
|
|
while (i < BLOCK_COLS)
|
|
{
|
|
uchar a0 = src[i];
|
|
dst[i] = a0;
|
|
|
|
i += WSZ;
|
|
}
|
|
}
|
|
#endif
|
|
} // copy_block_1d
|
|
|
|
#else
|
|
|
|
{ // copy_block_2d
|
|
__global const uchar* src = src0 + src_offset0;
|
|
__global uchar* dst = dst0 + dst_offset0;
|
|
|
|
uint i = get_local_id(0) * 4;
|
|
|
|
#define BLOCK_COLS_FILL_X4 (((BLOCK_COLS + 3) / 4) * 4)
|
|
#define BLOCK_SIZE_FILL_X4 (BLOCK_COLS_FILL_X4 * BLOCK_ROWS)
|
|
|
|
while (i < BLOCK_SIZE_FILL_X4)
|
|
{
|
|
int row = i / BLOCK_COLS_FILL_X4;
|
|
int col = i % BLOCK_COLS_FILL_X4;
|
|
|
|
uint src_offset = row * BLOCK_SRC_STRIDE + col;
|
|
#if BLOCK_COLS_FILL_X4 == BLOCK_COLS
|
|
uint dst_offset = i;
|
|
#else
|
|
uint dst_offset = row * BLOCK_COLS + col;
|
|
#endif
|
|
|
|
#if BLOCK_COLS_FILL_X4 != BLOCK_COLS
|
|
if (col <= BLOCK_COLS - 4)
|
|
#endif
|
|
{
|
|
uchar4 a = vload4(0, src + src_offset);
|
|
vstore4(a, 0, dst + dst_offset);
|
|
}
|
|
#if BLOCK_COLS_FILL_X4 != BLOCK_COLS
|
|
else
|
|
{
|
|
/* non-optimized reference code
|
|
while (col < BLOCK_COLS)
|
|
{
|
|
uchar a = src[src_offset];
|
|
dst[dst_offset] = a;
|
|
col++;
|
|
src_offset++;
|
|
dst_offset++;
|
|
}
|
|
*/
|
|
|
|
uint4 shift = (uint4)(0, 1, 2, 3);
|
|
shift = select((uint4)0, shift, col + shift < BLOCK_COLS);
|
|
|
|
dst[dst_offset + shift.s0] = src[src_offset + shift.s0];
|
|
|
|
#if BLOCK_COLS_FILL_X4 - BLOCK_COLS <= 2
|
|
dst[dst_offset + shift.s1] = src[src_offset + shift.s1];
|
|
#endif
|
|
#if BLOCK_COLS_FILL_X4 - BLOCK_COLS <= 1
|
|
dst[dst_offset + shift.s2] = src[src_offset + shift.s2];
|
|
#endif
|
|
}
|
|
#endif // BLOCK_COLS_FILL_X4 != BLOCK_COLS
|
|
i += WSZ * 4;
|
|
}
|
|
} // copy_block_2d
|
|
#endif
|
|
}
|