Google Git
Sign in
android / platform / external / ComputeLibrary / c05589415 / . / build / android-arm64v8a / src / core / CL / cl_kernels / sobel_filter.clembed
blob: 6796c954183d0b90f9d11d89b86b036a92a82fea [file] [log] [blame]
R"(
/*
* Copyright (c) 2016, 2017 Arm Limited.
*
* SPDX-License-Identifier: MIT
*
* Permission is hereby granted, free of charge, to any person obtaining a copy
* of this software and associated documentation files (the "Software"), to
* deal in the Software without restriction, including without limitation the
* rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
* sell copies of the Software, and to permit persons to whom the Software is
* furnished to do so, subject to the following conditions:
*
* The above copyright notice and this permission notice shall be included in all
* copies or substantial portions of the Software.
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
* SOFTWARE.
*/
/*
* Copyright (c) 2016-2020 Arm Limited.
*
* SPDX-License-Identifier: MIT
*
* Permission is hereby granted, free of charge, to any person obtaining a copy
* of this software and associated documentation files (the "Software"), to
* deal in the Software without restriction, including without limitation the
* rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
* sell copies of the Software, and to permit persons to whom the Software is
* furnished to do so, subject to the following conditions:
*
* The above copyright notice and this permission notice shall be included in all
* copies or substantial portions of the Software.
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
* SOFTWARE.
*/
#ifndef ARM_COMPUTE_HELPER_H
#define ARM_COMPUTE_HELPER_H
/*
* Copyright (c) 2020 Arm Limited.
*
* SPDX-License-Identifier: MIT
*
* Permission is hereby granted, free of charge, to any person obtaining a copy
* of this software and associated documentation files (the "Software"), to
* deal in the Software without restriction, including without limitation the
* rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
* sell copies of the Software, and to permit persons to whom the Software is
* furnished to do so, subject to the following conditions:
*
* The above copyright notice and this permission notice shall be included in all
* copies or substantial portions of the Software.
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
* SOFTWARE.
*/
/** Store the 0 to (n-1)th rows of the given variables
* @name STORE_ROW_n
*
* @param[in] N0 The width of the passed in vector. Supported: 1, 2, 3, 4, 8, 16
* @param[in] DATA_TYPE The data type of the vectors
* @param[in] BASENAME The basename of the variables
* @param[in] PTR The base pointer
* @param[in] STRIDE_Y The stride value in y-axis direction
* @param[in] Z The offset in z-axis direction
* @{
*/
#define STORE_ROW_1(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(BASENAME##0, 0, (__global DATA_TYPE *)(PTR + 0 * STRIDE_Y + Z##0));
#define STORE_ROW_2(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_1(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(BASENAME##1, 0, (__global DATA_TYPE *)(PTR + 1 * STRIDE_Y + Z##1));
#define STORE_ROW_3(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_2(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(BASENAME##2, 0, (__global DATA_TYPE *)(PTR + 2 * STRIDE_Y + Z##2));
#define STORE_ROW_4(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_3(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(BASENAME##3, 0, (__global DATA_TYPE *)(PTR + 3 * STRIDE_Y + Z##3));
#define STORE_ROW_5(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_4(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(BASENAME##4, 0, (__global DATA_TYPE *)(PTR + 4 * STRIDE_Y + Z##4));
#define STORE_ROW_6(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_5(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(BASENAME##5, 0, (__global DATA_TYPE *)(PTR + 5 * STRIDE_Y + Z##5));
#define STORE_ROW_7(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_6(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(BASENAME##6, 0, (__global DATA_TYPE *)(PTR + 6 * STRIDE_Y + Z##6));
#define STORE_ROW_8(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_7(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(BASENAME##7, 0, (__global DATA_TYPE *)(PTR + 7 * STRIDE_Y + Z##7));
#define STORE_ROW_9(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_8(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(BASENAME##8, 0, (__global DATA_TYPE *)(PTR + 8 * STRIDE_Y + Z##8));
#define STORE_ROW_10(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_9(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(BASENAME##9, 0, (__global DATA_TYPE *)(PTR + 9 * STRIDE_Y + Z##9));
#define STORE_ROW_11(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_10(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(BASENAME##A, 0, (__global DATA_TYPE *)(PTR + 10 * STRIDE_Y + Z##A));
#define STORE_ROW_12(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_11(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(BASENAME##B, 0, (__global DATA_TYPE *)(PTR + 11 * STRIDE_Y + Z##B));
#define STORE_ROW_13(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_12(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(BASENAME##C, 0, (__global DATA_TYPE *)(PTR + 12 * STRIDE_Y + Z##C));
#define STORE_ROW_14(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_13(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(BASENAME##D, 0, (__global DATA_TYPE *)(PTR + 13 * STRIDE_Y + Z##D));
#define STORE_ROW_15(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_14(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(BASENAME##E, 0, (__global DATA_TYPE *)(PTR + 14 * STRIDE_Y + Z##E));
#define STORE_ROW_16(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_15(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(BASENAME##F, 0, (__global DATA_TYPE *)(PTR + 15 * STRIDE_Y + Z##F));
/** @} */ // end of groupd STORE_ROW_n
/** Convert and store the 0th to (n-1)th rows of the given variables
* @name CONVERT_STORE_ROW_n
*
* @param[in] N0 The size of the vectors
* @param[in] DATA_TYPE The data type of the vectors
* @param[in] BASENAME The basename of the variables
* @param[in] PTR The base pointer
* @param[in] STRIDE_Y The stride value in y-axis direction
* @param[in] Z The offset in z-axis direction
* @{
*/
#define CONVERT_STORE_ROW_1(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(CONVERT_SAT((BASENAME##0), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 0 * STRIDE_Y + Z##0));
#define CONVERT_STORE_ROW_2(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
CONVERT_STORE_ROW_1(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(CONVERT_SAT((BASENAME##1), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 1 * STRIDE_Y + Z##1));
#define CONVERT_STORE_ROW_3(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
CONVERT_STORE_ROW_2(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(CONVERT_SAT((BASENAME##2), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 2 * STRIDE_Y + Z##2));
#define CONVERT_STORE_ROW_4(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
CONVERT_STORE_ROW_3(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(CONVERT_SAT((BASENAME##3), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 3 * STRIDE_Y + Z##3));
#define CONVERT_STORE_ROW_5(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
CONVERT_STORE_ROW_4(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(CONVERT_SAT((BASENAME##4), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 4 * STRIDE_Y + Z##4));
#define CONVERT_STORE_ROW_6(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
CONVERT_STORE_ROW_5(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(CONVERT_SAT((BASENAME##5), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 5 * STRIDE_Y + Z##5));
#define CONVERT_STORE_ROW_7(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
CONVERT_STORE_ROW_6(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(CONVERT_SAT((BASENAME##6), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 6 * STRIDE_Y + Z##6));
#define CONVERT_STORE_ROW_8(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
CONVERT_STORE_ROW_7(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(CONVERT_SAT((BASENAME##7), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 7 * STRIDE_Y + Z##7));
#define CONVERT_STORE_ROW_9(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
CONVERT_STORE_ROW_8(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(CONVERT_SAT((BASENAME##8), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 8 * STRIDE_Y + Z##8));
#define CONVERT_STORE_ROW_10(N0, DATA, BASENAME, PTR, STRIDE_Y, Z) \
CONVERT_STORE_ROW_9(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(CONVERT_SAT((BASENAME##9), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 9 * STRIDE_Y + Z##9));
#define CONVERT_STORE_ROW_11(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
CONVERT_STORE_ROW_10(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(CONVERT_SAT((BASENAME##A), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 10 * STRIDE_Y + Z##A));
#define CONVERT_STORE_ROW_12(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
CONVERT_STORE_ROW_11(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(CONVERT_SAT((BASENAME##B), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 11 * STRIDE_Y + Z##B));
#define CONVERT_STORE_ROW_13(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
CONVERT_STORE_ROW_12(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(CONVERT_SAT((BASENAME##C), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 12 * STRIDE_Y + Z##C));
#define CONVERT_STORE_ROW_14(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
CONVERT_STORE_ROW_13(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(CONVERT_SAT((BASENAME##D), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 13 * STRIDE_Y + Z##D));
#define CONVERT_STORE_ROW_15(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
CONVERT_STORE_ROW_14(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(CONVERT_SAT((BASENAME##E), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 14 * STRIDE_Y + Z##E));
#define CONVERT_STORE_ROW_16(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
CONVERT_STORE_ROW_15(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE(N0) \
(CONVERT_SAT((BASENAME##F), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 15 * STRIDE_Y + Z##F));
/** @} */ // end of groupd CONVERT_STORE_ROW_n
/** Store a block of the given size M0xN0
* @name STORE_BLOCK
*
* Supported cases are M0=1,2,3,...,16 and N0=2,3,4,8,16.
* The data to store is expected to have consecutive names for each row.
* E.g., for M0=3 and basename=c, the expected names are c0, c1 and c2.
* The Z offset is expected to have consecutive names.
* E.g., for M0=3 and Z=zin, the expected z offset names are zin0, zin1 and zin2.
*
* @param[in] M0 The number of rows to store
* @param[in] N0 The size of each vector
* @param[in] DATA_TYPE The data type of the vectors
* @param[in] BASENAME The basename of the variables
* @param[in] PTR The base pointer
* @param[in] STRIDE_Y The stride value in y-axis direction
* @param[in] Z The offset in z-axis direction
* @{
*/
#define STORE_BLOCK_STR(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) STORE_ROW_##M0(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)
#define STORE_BLOCK(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) STORE_BLOCK_STR(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)
/** @} */ // end of group STORE_BLOCK
/** Convert and store a block of the given size M0xN0
* @name CONVERT_STORE_BLOCK
*
* Supported cases are M0=1,2,3,...,16 and N0=2,3,4,8,16.
* The data to store is expected to have consecutive names for each row.
* E.g., for M0=3 and basename=c, the expected names are c0, c1 and c2.
* The Z offset is expected to have consecutive names.
* E.g., for M0=3 and Z=zin, the expected z offset names are zin0, zin1 and zin2.
*
* @param[in] M0 The number of rows to store
* @param[in] N0 The size of each vector
* @param[in] DATA_TYPE The data type of the vectors
* @param[in] BASENAME The basename of the variables
* @param[in] PTR The base pointer
* @param[in] STRIDE_Y The stride value in y-axis direction
* @param[in] Z The offset in z-axis direction
* @{
*/
#define CONVERT_STORE_BLOCK_STR(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) CONVERT_STORE_ROW_##M0(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)
#define CONVERT_STORE_BLOCK(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) CONVERT_STORE_BLOCK_STR(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)
/** @} */ // end of group CONVERT_STORE_BLOCK
/** Partially store the 0 to (n-1)th rows of the given variables
* @name STORE_ROW_PARTIAL_n
* Within each row, store the lower @p STORE_N0 elements of vectors of width @p N0
*
* @note in case @p STORE_N0 != 1, 2, 3, 4, 8, 16, extra vstore(s) will be invoked, thus incurring small performance penalty.
*
* @param[in] N0 The width of the passed in vector. Supported: 1, 2, 3, 4, 8, 16
* @param[in] STORE_N0 The **lower** size of the vectors to store. Supported: [1-16 and <= @p N0
* @param[in] DATA_TYPE The data type of the vectors
* @param[in] BASENAME The basename of the variables
* @param[in] PTR The base pointer
* @param[in] STRIDE_Y The stride value in y-axis direction
* @param[in] Z The offset in z-axis direction
* @{
*/
#define STORE_ROW_PARTIAL_1(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE_PARTIAL(N0, STORE_N0) \
(BASENAME##0, 0, (__global DATA_TYPE *)(PTR + 0 * STRIDE_Y + Z##0));
#define STORE_ROW_PARTIAL_2(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_PARTIAL_1(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE_PARTIAL(N0, STORE_N0) \
(BASENAME##1, 0, (__global DATA_TYPE *)(PTR + 1 * STRIDE_Y + Z##1));
#define STORE_ROW_PARTIAL_3(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_PARTIAL_2(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE_PARTIAL(N0, STORE_N0) \
(BASENAME##2, 0, (__global DATA_TYPE *)(PTR + 2 * STRIDE_Y + Z##2));
#define STORE_ROW_PARTIAL_4(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_PARTIAL_3(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE_PARTIAL(N0, STORE_N0) \
(BASENAME##3, 0, (__global DATA_TYPE *)(PTR + 3 * STRIDE_Y + Z##3));
#define STORE_ROW_PARTIAL_5(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_PARTIAL_4(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE_PARTIAL(N0, STORE_N0) \
(BASENAME##4, 0, (__global DATA_TYPE *)(PTR + 4 * STRIDE_Y + Z##4));
#define STORE_ROW_PARTIAL_6(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_PARTIAL_5(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE_PARTIAL(N0, STORE_N0) \
(BASENAME##5, 0, (__global DATA_TYPE *)(PTR + 5 * STRIDE_Y + Z##5));
#define STORE_ROW_PARTIAL_7(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_PARTIAL_6(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE_PARTIAL(N0, STORE_N0) \
(BASENAME##6, 0, (__global DATA_TYPE *)(PTR + 6 * STRIDE_Y + Z##6));
#define STORE_ROW_PARTIAL_8(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_PARTIAL_7(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE_PARTIAL(N0, STORE_N0) \
(BASENAME##7, 0, (__global DATA_TYPE *)(PTR + 7 * STRIDE_Y + Z##7));
#define STORE_ROW_PARTIAL_9(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_PARTIAL_8(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE_PARTIAL(N0, STORE_N0) \
(BASENAME##8, 0, (__global DATA_TYPE *)(PTR + 8 * STRIDE_Y + Z##8));
#define STORE_ROW_PARTIAL_10(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_PARTIAL_9(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE_PARTIAL(N0, STORE_N0) \
(BASENAME##9, 0, (__global DATA_TYPE *)(PTR + 9 * STRIDE_Y + Z##9));
#define STORE_ROW_PARTIAL_11(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_PARTIAL_10(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE_PARTIAL(N0, STORE_N0) \
(BASENAME##A, 0, (__global DATA_TYPE *)(PTR + 10 * STRIDE_Y + Z##A));
#define STORE_ROW_PARTIAL_12(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_PARTIAL_11(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE_PARTIAL(N0, STORE_N0) \
(BASENAME##B, 0, (__global DATA_TYPE *)(PTR + 11 * STRIDE_Y + Z##B));
#define STORE_ROW_PARTIAL_13(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_PARTIAL_12(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE_PARTIAL(N0, STORE_N0) \
(BASENAME##C, 0, (__global DATA_TYPE *)(PTR + 12 * STRIDE_Y + Z##C));
#define STORE_ROW_PARTIAL_14(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_PARTIAL_13(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE_PARTIAL(N0, STORE_N0) \
(BASENAME##D, 0, (__global DATA_TYPE *)(PTR + 13 * STRIDE_Y + Z##D));
#define STORE_ROW_PARTIAL_15(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_PARTIAL_14(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE_PARTIAL(N0, STORE_N0) \
(BASENAME##E, 0, (__global DATA_TYPE *)(PTR + 14 * STRIDE_Y + Z##E));
#define STORE_ROW_PARTIAL_16(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
STORE_ROW_PARTIAL_15(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \
VSTORE_PARTIAL(N0, STORE_N0) \
(BASENAME##F, 0, (__global DATA_TYPE *)(PTR + 15 * STRIDE_Y + Z##F));
/** @} */ // end of groupd STORE_ROW_PARTIAL_n
/** Partially store a block of the given size STORE_M0xSTORE_N0
* @name STORE_BLOCK_PARTIAL
*
* @note The vector width @p N0 is also required for correct partial storing behaviour.
* @note in case @p STORE_N0 != 1, 2, 3, 4, 8, 16, extra vstore(s) will be invoked, thus incurring small performance penalty.
*
* The data to store is expected to have consecutive names for each row.
* E.g., for STORE_M0=3 and basename=c, the expected names are c0, c1 and c2.
* The Z offset is expected to have consecutive names.
* E.g., for STORE_M0=3 and Z=zin, the expected z offset names are zin0, zin1 and zin2.
*
* @param[in] STORE_M0 The number of rows to store. Supported: 1-16
* @param[in] STORE_N0 The lower number of elements of vectors to store. Supported: 1-16 and <= @p N0
* @param[in] N0 The size of each vector. Supported: 1, 2, 3, 4, 8, 16
* @param[in] DATA_TYPE The data type of the vectors
* @param[in] BASENAME The basename of the variables
* @param[in] PTR The base pointer
* @param[in] STRIDE_Y The stride value in y-axis direction
* @param[in] Z The offset in z-axis direction
* @{
*/
#define STORE_BLOCK_PARTIAL_STR(STORE_M0, STORE_N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) STORE_ROW_PARTIAL_##STORE_M0(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)
#define STORE_BLOCK_PARTIAL(STORE_M0, STORE_N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) STORE_BLOCK_PARTIAL_STR(STORE_M0, STORE_N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)
/** Store a block that can be partial in both x and y dimensions
*
* @note in cases @p PARTIAL_STORE_N0 != 1, 2, 3, 4, 8, 16, extra vstore(s) will be invoked, thus incurring small performance penalty.
*
* The data to store is expected to have consecutive names for each row.
* E.g., for M0=3 and basename=c, the expected names are c0, c1 and c2.
* The Z offset is expected to have consecutive names.
* E.g., for M0=3 and Z=zin, the expected z offset names are zin0, zin1 and zin2.
*
* @param[in] M0 The number of rows to store, for non-partial blocks. Supported: 1-16
* @param[in] N0 The size of each vector, for non-partial blocks. Supported: 1, 2, 3, 4, 8, 16
* @param[in] DATA_TYPE The data type of the vectors
* @param[in] BASENAME The basename of the variables
* @param[in] PTR The base pointer
* @param[in] STRIDE_Y The stride value in y-axis direction
* @param[in] Z The offset in z-axis direction
* @param[in] PARTIAL_STORE_M0 The partial size in y, for partial blocks. Supported range: [1, @p M0)
* @param[in] PARTIAL_STORE_N0 The partial size in x, for partial blocks. Supported range: [1, @p N0)
* @param[in] PARTIAL_COND_Y Condition on the y axis to perform the partial store Y. True to use PARTIAL_STORE_M0 rather than M0.
* @param[in] PARTIAL_COND_X Condition on the x axis to perform the partial store X. True to use PARTIAL_STORE_N0 rather than N0.
*/
#define STORE_BLOCK_PARTIAL_IN_X_AND_Y(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_STORE_N0, PARTIAL_COND_Y, PARTIAL_COND_X) \
if(!(PARTIAL_COND_X) && !(PARTIAL_COND_Y)) \
{ \
STORE_BLOCK_PARTIAL(M0, N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z); \
} \
else if((PARTIAL_COND_Y) && !(PARTIAL_COND_X)) \
{ \
STORE_BLOCK_PARTIAL(PARTIAL_STORE_M0, N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z); \
} \
else if(!(PARTIAL_COND_Y) && (PARTIAL_COND_X)) \
{ \
STORE_BLOCK_PARTIAL(M0, PARTIAL_STORE_N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z); \
} \
else \
{ \
STORE_BLOCK_PARTIAL(PARTIAL_STORE_M0, PARTIAL_STORE_N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z); \
}
/** Store a block that can only be partial in x but not y.
*
* @note in case @p N0 or @p PARTIAL_STORE_N0 != 1, 2, 3, 4, 8, 16, extra vstore(s) will be invoked, thus incurring small performance penalty.
*
* The data to store is expected to have consecutive names for each row.
* E.g., for M0=3 and basename=c, the expected names are c0, c1 and c2.
* The Z offset is expected to have consecutive names.
* E.g., for M0=3 and Z=zin, the expected z offset names are zin0, zin1 and zin2.
*
* @param[in] M0 The number of rows to store, for non-partial blocks. Supported: 1-16
* @param[in] N0 The size of each vector, for non-partial blocks. Supported: 1, 2, 3, 4, 8, 16
* @param[in] DATA_TYPE The data type of the vectors
* @param[in] BASENAME The basename of the variables
* @param[in] PTR The base pointer
* @param[in] STRIDE_Y The stride value in y-axis direction
* @param[in] Z The offset in z-axis direction
* @param[in] PARTIAL_STORE_N0 The partial size in x, for partial blocks. Supported range: [1, @p N0)
* @param[in] PARTIAL_COND_X Condition on the x axis to perform the partial store X. True to use PARTIAL_STORE_N0 rather than N0.
*/
#define STORE_BLOCK_PARTIAL_IN_X(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_N0, PARTIAL_COND_X) \
if(!(PARTIAL_COND_X)) \
{ \
STORE_BLOCK_PARTIAL(M0, N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z); \
} \
else \
{ \
STORE_BLOCK_PARTIAL(M0, PARTIAL_STORE_N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z); \
}
/** Store a block that can only be partial in y but not x.
*
* @note in case @p N0 or @p PARTIAL_STORE_N0 != 1, 2, 3, 4, 8, 16, extra vstore(s) will be invoked, thus incurring small performance penalty.
*
* The data to store is expected to have consecutive names for each row.
* E.g., for M0=3 and basename=c, the expected names are c0, c1 and c2.
* The Z offset is expected to have consecutive names.
* E.g., for M0=3 and Z=zin, the expected z offset names are zin0, zin1 and zin2.
*
* @param[in] M0 The number of rows to store, for non-partial blocks. Supported: 1-16
* @param[in] N0 The size of each vector, for non-partial blocks. Supported: 1, 2, 3, 4, 8, 16
* @param[in] DATA_TYPE The data type of the vectors
* @param[in] BASENAME The basename of the variables
* @param[in] PTR The base pointer
* @param[in] STRIDE_Y The stride value in y-axis direction
* @param[in] Z The offset in z-axis direction
* @param[in] PARTIAL_STORE_M0 The partial size in y, for partial blocks. Supported range: [1, @p M0)
* @param[in] PARTIAL_COND_Y Condition on the y axis to perform the partial store Y. True to use PARTIAL_STORE_M0 rather than M0.
*/
#define STORE_BLOCK_PARTIAL_IN_Y(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_COND_Y) \
if(!(PARTIAL_COND_Y)) \
{ \
STORE_BLOCK_PARTIAL(M0, N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z); \
} \
else \
{ \
STORE_BLOCK_PARTIAL(PARTIAL_STORE_M0, N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z); \
}
/** @} */ // end of group STORE_BLOCK_PARTIAL
#if defined(PARTIAL_STORE_M0) && defined(PARTIAL_STORE_N0)
/** Boundary-aware GEMM block store
* @name STORE_BLOCK_BOUNDARY_AWARE
* This macro assumes the following schemes to achieve boundary-awareness:
* - Overlapping load in Y axis from lhs tensor. This implies lhs has no padding along y dim.
* - Non-Overlapping(normal) load from rhs tensor. This imples rhs can have paddings.
* - Overlapping load in Y axis from bias tensor. This implies rhs has no padding along y dim.
* The macro then ensures that the dst tensor can be stored without any paddings in both x and y dim.
*
* In the y dimension, we place the partial blocks **at the beginning** while in the x dimension, we place the partial
* blocks **at the end**.
* Say, the dst tensor is of shape MxN and we have M0 and N0 as the block size, this is how we define "partial blocks"/
* "boundary block" (we use the 2 terms "partial blocks" and "boundary blocks" interchangeably) and its various parameters:
*
* *--x--> x == 0 x == 1
* | |<------------------------------N-------------------------->|
* y |<--------------N0------------->|<----PARTIAL_STORE_N0----->|
* | -------------#############################################################
* * | | |...............................|...........................|
* y == 0 | PAR_..._M0 |......Boundary block in y......|.Boundary block in x and y.|
* | | |...............................|...........................|
* M --#############################################################
* | | | |...........................|
* y == 1 | M0 | Non-boundary block |....Boundary block in x....|
* | | | |...........................|
* |------------#############################################################
*
* Then @p PARTIAL_STORE_M0 = M % M0 and @p PARTIAL_STORE_N0 = N % N0
*
* @note in cases @p PARTIAL_STORE_N0 != 1, 2, 3, 4, 8, 16, extra vstore(s) will be invoked, thus incurring small performance penalty.
*
* It automatically detects if a giving M,N,M0,N0 combination can yield partial blocks in either X and Y dimension,
* and select corresponding store methods such that the boundary detection logic is only added when needed.
*
* The data to store is expected to have consecutive names for each row.
* E.g., for M0=3 and basename=c, the expected names are c0, c1 and c2.
* The Z offset is expected to have consecutive names.
* E.g., for M0=3 and Z=zin, the expected z offset names are zin0, zin1 and zin2.
*
* @param[in] M0 The number of rows to store, for non-partial blocks. Supported: 1-16
* @param[in] N0 The size of each vector, for non-partial blocks. Supported: 1, 2, 3, 4, 8, 16
* @param[in] DATA_TYPE The data type of the vectors
* @param[in] BASENAME The basename of the variables
* @param[in] PTR The base pointer
* @param[in] STRIDE_Y The stride value in y-axis direction
* @param[in] Z The offset in z-axis direction
* @param[in] PARTIAL_STORE_M0 The partial size in y, for partial blocks. Supported: [0, @p M0)
* @param[in] PARTIAL_STORE_N0 The partial size in x, for partial blocks. Supported: [0, @p N0)
* @param[in] PARTIAL_COND_Y Condition on the y axis to perform the partial store Y. True to use PARTIAL_STORE_M0 rather than M0.
* @param[in] PARTIAL_COND_X Condition on the x axis to perform the partial store X. True to use PARTIAL_STORE_N0 rather than N0.
* @{
*/
#if PARTIAL_STORE_M0 == 0 && PARTIAL_STORE_N0 == 0
// Case1: No partial blocks in either x or y
#define STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_STORE_N0, PARTIAL_COND_Y, PARTIAL_COND_X) \
STORE_BLOCK(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z)
#elif PARTIAL_STORE_M0 > 0 && PARTIAL_STORE_N0 == 0
// Case2: Partial blocks in y
#define STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_STORE_N0, PARTIAL_COND_Y, PARTIAL_COND_X) \
STORE_BLOCK_PARTIAL_IN_Y(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_COND_Y)
#elif PARTIAL_STORE_M0 == 0 && PARTIAL_STORE_N0 > 0
// Case3: Partial blocks in x
#define STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_STORE_N0, PARTIAL_COND_Y, PARTIAL_COND_X) \
STORE_BLOCK_PARTIAL_IN_X(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_N0, PARTIAL_COND_X)
#else // PARTIAL_STORE_M0 == 0 && PARTIAL_STORE_N0 == 0
// Case4: Partial blocks in both x and y
#define STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_STORE_N0, PARTIAL_COND_Y, PARTIAL_COND_X) \
STORE_BLOCK_PARTIAL_IN_X_AND_Y(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_STORE_N0, PARTIAL_COND_Y, PARTIAL_COND_X)
#endif // PARTIAL_STORE_M0 == 0 && PARTIAL_STORE_N0 == 0
#endif // defined(PARTIAL_STORE_M0) && defined(PARTIAL_STORE_N0)
/** @} */ // end of group STORE_BLOCK_BOUNDARY_AWARE
#if defined(PARTIAL_STORE_M0)
/** Compute the start m0 row (LHS, BIAS and DST) in a boundary-aware way so as to avoid padding
* @name COMPUTE_M0_START_ROW
* If there're any partial blocks in y dimension, they are placed at the beginning of the rows.
* This shift amount is added to all rows such that the partial block (at the beginning) overlaps with the subsequent
* blocks in the y dimension to avoid any padding.
* EG: M0=4, PARTIAL_STORE_M0=1:
* | Non-overlapping | +M0_ROW_SHIFT (Overlapping)
* block 0 (partial)| start row = 0 | start row = 0
* block 1 (full) | start row = 4 | start row = 1
* block 2 (full) | start row = 8 | start row = 5
*
* @param[in] y Global id of current block in y.
* @param[in] M0 The number of rows to store, for non-partial blocks. Supported: 1-16
* @param[in] PARTIAL_STORE_M0 The partial size in y, for partial blocks. Supported: [0, @p M0)
* @{
*/
#define COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0) \
((uint)(max(0, (int)(y * M0) - (int)((M0 - PARTIAL_STORE_M0) % M0))))
#else // defined(PARTIAL_STORE_M0)
#define COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0) \
((uint)(y * M0))
#endif // defined(PARTIAL_STORE_M0)
/** @} */ // end of group COMPUTE_M0_START_ROW
/** Store a vector that can only be partial in x.
*
* @note in case @p vec_size or @p leftover != 1, 2, 3, 4, 8, 16, extra vstore(s) will be invoked, thus incurring small performance penalty.
*
* The data to store is expected to end in a 0.
* E.g., for basename=c, the expected name is c0.
*
* @param[in] basename The name of the variable without trailing 0
* @param[in] data_type The data type of the vector
* @param[in] ptr The base pointer
* @param[in] vec_size The vector size if cond = false. Supported: 1, 2, 3, 4, 8, 16
* @param[in] leftover The vector size if cond = true. Supported range: [1, @p vec_size0)
* @param[in] cond Condition to select either vec_size0 or vec_size1
* @{
*/
#define STORE_VECTOR_SELECT(basename, data_type, ptr, vec_size, leftover, cond) \
STORE_BLOCK_PARTIAL_IN_X(1, vec_size, data_type, basename, ptr, 0, 0, leftover, cond)
/** @} */ // end of group STORE_VECTOR_SELECT
#if defined(ARM_COMPUTE_OPENCL_FP16_ENABLED) && defined(cl_khr_fp16)
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
#endif // defined(ARM_COMPUTE_OPENCL_FP16_ENABLED) && defined(cl_khr_fp16)
#if defined(ARM_COMPUTE_OPENCL_DOT8_ENABLED) && defined(cl_arm_integer_dot_product_int8)
#pragma OPENCL EXTENSION cl_arm_integer_dot_product_int8 : enable
#endif // defined(ARM_COMPUTE_OPENCL_DOT8_ENABLED) && defined(cl_arm_integer_dot_product_int8)
#if defined(ARM_COMPUTE_OPENCL_DOT8_ACC_ENABLED) && defined(cl_arm_integer_dot_product_accumulate_int8)
#pragma OPENCL EXTENSION cl_arm_integer_dot_product_accumulate_int8 : enable
#endif // defined(ARM_COMPUTE_OPENCL_DOT8_ACC_ENABLED) && defined(cl_arm_integer_dot_product_accumulate_int8)
#if defined(ARM_COMPUTE_DEBUG_ENABLED) && defined(cl_arm_printf)
#pragma OPENCL EXTENSION cl_arm_printf : enable
#endif // defined(ARM_COMPUTE_DEBUG_ENABLED) && defined(cl_arm_printf)
#define GPU_ARCH_MIDGARD 0x100
#define GPU_ARCH_BIFROST 0x200
/** Concatenate two inputs.
*
* @param[in] a The first input to be concatenated
* @param[in] b The second input to be concatenated
*
* @return The concatenated output
*/
#define CONCAT(a, b) a##b
/** Expand the given vector
*
* @param[in] x The vector to be expanded
*
* @return The expanded output
*/
#define EXPAND(x) x
/** Clamp the given value between an upper and lower bound.
*
* @param[in] x The value to be clamped
* @param[in] min_val The lower bound
* @param[in] max_val The upper bound
*
* @return The clamped value.
*/
#define CLAMP(x, min_val, max_val) min(max(x, min_val), max_val)
/** REVn reverses the given vector whose size is n.
* @name REVn
*
* @param[in] x The vector to be reversed
*
* @return The reversed vector
* @{
*/
#define REV1(x) ((x))
#define REV2(x) ((x).s10)
#define REV3(x) ((x).s210)
#define REV4(x) ((x).s3210)
#define REV8(x) ((x).s76543210)
#define REV16(x) ((x).sFEDCBA9876543210)
/** @} */ // end of group REVn
/** Reverse the given vector.
* @name REVERSE
*
* @param[in] x The vector to be reversed
* @param[in] s The size of the vector
*
* @return The reversed vector
* @{
*/
#define REVERSE_STR(x, s) REV##s((x))
#define REVERSE(x, s) REVERSE_STR(x, s)
/** @} */ // end of group REVERSE
/** Circular-right-shift (rotate-right) the vector of size s by the amount of n.
* @name ROTs_n
*
* @param[in] x The vector to be shifted
*
* @return The shifted vector
* @{
*/
#define ROT1_0(x) ((x))
#define ROT2_0(x) ((x))
#define ROT2_1(x) ((x).s10)
#define ROT3_0(x) ((x))
#define ROT3_1(x) ((x).s201)
#define ROT3_2(x) ((x).s120)
#define ROT4_0(x) ((x))
#define ROT4_1(x) ((x).s3012)
#define ROT4_2(x) ((x).s2301)
#define ROT4_3(x) ((x).s1230)
#define ROT8_0(x) ((x))
#define ROT8_1(x) ((x).s70123456)
#define ROT8_2(x) ((x).s67012345)
#define ROT8_3(x) ((x).s56701234)
#define ROT8_4(x) ((x).s45670123)
#define ROT8_5(x) ((x).s34567012)
#define ROT8_6(x) ((x).s23456701)
#define ROT8_7(x) ((x).s12345670)
#define ROT16_0(x) ((x))
#define ROT16_1(x) ((x).sF0123456789ABCDE)
#define ROT16_2(x) ((x).sEF0123456789ABCD)
#define ROT16_3(x) ((x).sDEF0123456789ABC)
#define ROT16_4(x) ((x).sCDEF0123456789AB)
#define ROT16_5(x) ((x).sBCDEF0123456789A)
#define ROT16_6(x) ((x).sABCDEF0123456789)
#define ROT16_7(x) ((x).s9ABCDEF012345678)
#define ROT16_8(x) ((x).s89ABCDEF01234567)
#define ROT16_9(x) ((x).s789ABCDEF0123456)
#define ROT16_10(x) ((x).s6789ABCDEF012345)
#define ROT16_11(x) ((x).s56789ABCDEF01234)
#define ROT16_12(x) ((x).s456789ABCDEF0123)
#define ROT16_13(x) ((x).s3456789ABCDEF012)
#define ROT16_14(x) ((x).s23456789ABCDEF01)
#define ROT16_15(x) ((x).s123456789ABCDEF0)
/** @} */ // end of group ROTs_n
/** Circular-right-shift (rotate-right) the given vector by the given amount.
* @name ROTATE
*
* @param[in] x The vector to be shifted
* @param[in] s The size of the vector
* @param[in] n The amount to be shifted
*
* @return The shifted vector
* @{
*/
#define ROTATE_STR(x, s, n) ROT##s##_##n(x)
#define ROTATE(x, s, n) ROTATE_STR(x, s, n)
/** @} */ // end of group ROTATE
/** Creates a vector of size n filled with offset values corresponding to the location of each element.
* @name V_OFFSn
*
* @param[in] dt The data type of the output vector
*
* @return The vector filled with offset values
* @{
*/
#define V_OFFS1(dt) (dt##1)(0)
#define V_OFFS2(dt) (dt##2)(0, 1)
#define V_OFFS3(dt) (dt##3)(0, 1, 2)
#define V_OFFS4(dt) (dt##4)(0, 1, 2, 3)
#define V_OFFS8(dt) (dt##8)(0, 1, 2, 3, 4, 5, 6, 7)
#define V_OFFS16(dt) (dt##16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15)
/** @} */ // end of group V_OFFSn
/** Create a vector filled with offset values corresponding to the location of each element.
* @name VEC_OFFS
*
* @param[in] dt The data type of the output vector
* @param[in] s The size of the output vector
*
* @return The vector filled with offset values
* @{
*/
#define VEC_OFFS_STR(dt, s) V_OFFS##s(dt)
#define VEC_OFFS(dt, s) VEC_OFFS_STR(dt, s)
/** @} */ // end of group VEC_OFFS
#define VLOAD_STR(size) vload##size
#define VLOAD(size) VLOAD_STR(size)
#define PIXEL_UNIT4 1
#define PIXEL_UNIT8 2
#define PIXEL_UNIT16 4
/** Utility macro to convert a vector size in pixel unit.
*
* @name CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT
*
* @param[in] vec_size Vector size. Only 4,8 and 16 is supported
*
* @return The pixel unit (number of pixels)
* @{
*/
#define CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT_STR(vec_size) PIXEL_UNIT##vec_size
#define CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT(vec_size) CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT_STR(vec_size)
/** @} */ // end of group CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT
#define read_image2d_floatx1(img, x_coord, y_coord) (float4)(read_imagef(img, (int2)(x_coord, y_coord)));
#define read_image2d_floatx2(img, x_coord, y_coord) (float8)(read_imagef(img, (int2)(x_coord, y_coord)), read_imagef(img, (int2)(x_coord + 1, y_coord)));
#define read_image2d_floatx4(img, x_coord, y_coord) (float16)(read_imagef(img, (int2)(x_coord, y_coord)), read_imagef(img, (int2)(x_coord + 1, y_coord)), read_imagef(img, (int2)(x_coord + 2, y_coord)), read_imagef(img, (int2)(x_coord + 3, y_coord)));
#if defined(ARM_COMPUTE_OPENCL_FP16_ENABLED) && defined(cl_khr_fp16)
#define read_image2d_halfx1(img, x_coord, y_coord) (half4)(read_imageh(img, (int2)(x_coord, y_coord)));
#define read_image2d_halfx2(img, x_coord, y_coord) (half8)(read_imageh(img, (int2)(x_coord, y_coord)), read_imageh(img, (int2)(x_coord + 1, y_coord)));
#define read_image2d_halfx4(img, x_coord, y_coord) (half16)(read_imageh(img, (int2)(x_coord, y_coord)), read_imageh(img, (int2)(x_coord + 1, y_coord)), read_imageh(img, (int2)(x_coord + 2, y_coord)), read_imageh(img, (int2)(x_coord + 3, y_coord)));
#endif // defined(ARM_COMPUTE_OPENCL_FP16_ENABLED) && defined(cl_khr_fp16)
/** Utility macro to read a 2D OpenCL image object.
*
* @note Coordinates are not normalized
*
* @param[in] data_type Data type
* @param[in] n0 Number of pixel to read. Only 1,2 and 4 is supported
* @param[in] img OpenCL image object
* @param[in] x_coord The x coordinate for the top-left pixel
* @param[in] y_coord The y coordinate for the top-left pixel
*
* @return Pixels from the 2D OpenCL image object
* @{
*/
#define READ_IMAGE2D_STR(data_type, n0, img, x_coord, y_coord) read_image2d_##data_type##x##n0(img, x_coord, y_coord)
#define READ_IMAGE2D(data_type, n0, img, x_coord, y_coord) READ_IMAGE2D_STR(data_type, n0, img, x_coord, y_coord)
#define VSTORE_STR(size) vstore##size
#define VSTORE(size) VSTORE_STR(size)
#define float1 float
#define half1 half
#define char1 char
#define uchar1 uchar
#define short1 short
#define ushort1 ushort
#define int1 int
#define uint1 uint
#define long1 long
#define ulong1 ulong
#define double1 double
#define vload1(OFFSET, PTR) *(OFFSET + PTR)
#define vstore1(DATA, OFFSET, PTR) *(OFFSET + PTR) = DATA
/** Extended partial vstore that correctly handles scalar values as well.
* Store the **lower** 0 to (n-1)th elements of the given vector while minimising the amount of vstore ops
* @name VSTORE_PARTIAL
*
* @note With this macro, the passed data can be both a vector and a scalar
* @note @p store_size needs to be <= @p size
* eg 1: Valid
* VSTORE_PARTIAL(16, 15) ...;
* eg 2: Invalid
* VSTORE_PARTIAL(4, 7) ...;
*
* @param[in] size The width of @p DATA. Supported values: 1(scalar), 2, 3, 4, 8, 16
* @param[in] store_size The number of lower elements to store. Supported values: 1-16, but has to be <= @p size
* @{
*/
#define VSTORE_PARTIAL_STR(size, store_size) vstore_partial_##size##_##store_size
#define VSTORE_PARTIAL(size, store_size) VSTORE_PARTIAL_STR(size, store_size)
#define NO_STORE(data, offs, ptr) \
{ \
}
// Size == 1 (scalar)
#define vstore_partial_1_0 NO_STORE
#define vstore_partial_1_1 vstore1
#define vstore_partial_1_2 NO_STORE
#define vstore_partial_1_3 NO_STORE
#define vstore_partial_1_4 NO_STORE
#define vstore_partial_1_5 NO_STORE
#define vstore_partial_1_6 NO_STORE
#define vstore_partial_1_7 NO_STORE
#define vstore_partial_1_8 NO_STORE
#define vstore_partial_1_9 NO_STORE
#define vstore_partial_1_10 NO_STORE
#define vstore_partial_1_11 NO_STORE
#define vstore_partial_1_12 NO_STORE
#define vstore_partial_1_13 NO_STORE
#define vstore_partial_1_14 NO_STORE
#define vstore_partial_1_15 NO_STORE
#define vstore_partial_1_16 NO_STORE
// Size == 2
#define vstore_partial_2_0 NO_STORE
#define vstore_partial_2_1 vstore_partial_1
#define vstore_partial_2_2 vstore_partial_2
#define vstore_partial_2_3 NO_STORE
#define vstore_partial_2_4 NO_STORE
#define vstore_partial_2_5 NO_STORE
#define vstore_partial_2_6 NO_STORE
#define vstore_partial_2_7 NO_STORE
#define vstore_partial_2_8 NO_STORE
#define vstore_partial_2_9 NO_STORE
#define vstore_partial_2_10 NO_STORE
#define vstore_partial_2_11 NO_STORE
#define vstore_partial_2_12 NO_STORE
#define vstore_partial_2_13 NO_STORE
#define vstore_partial_2_14 NO_STORE
#define vstore_partial_2_15 NO_STORE
#define vstore_partial_2_16 NO_STORE
// Size == 3
#define vstore_partial_3_0 NO_STORE
#define vstore_partial_3_1 vstore_partial_1
#define vstore_partial_3_2 vstore_partial_2
#define vstore_partial_3_3 vstore_partial_3
#define vstore_partial_3_4 NO_STORE
#define vstore_partial_3_5 NO_STORE
#define vstore_partial_3_6 NO_STORE
#define vstore_partial_3_7 NO_STORE
#define vstore_partial_3_8 NO_STORE
#define vstore_partial_3_9 NO_STORE
#define vstore_partial_3_10 NO_STORE
#define vstore_partial_3_11 NO_STORE
#define vstore_partial_3_12 NO_STORE
#define vstore_partial_3_13 NO_STORE
#define vstore_partial_3_14 NO_STORE
#define vstore_partial_3_15 NO_STORE
#define vstore_partial_3_16 NO_STORE
// Size == 4
#define vstore_partial_4_0 NO_STORE
#define vstore_partial_4_1 vstore_partial_1
#define vstore_partial_4_2 vstore_partial_2
#define vstore_partial_4_3 vstore_partial_3
#define vstore_partial_4_4 vstore_partial_4
#define vstore_partial_4_5 NO_STORE
#define vstore_partial_4_6 NO_STORE
#define vstore_partial_4_7 NO_STORE
#define vstore_partial_4_8 NO_STORE
#define vstore_partial_4_9 NO_STORE
#define vstore_partial_4_10 NO_STORE
#define vstore_partial_4_11 NO_STORE
#define vstore_partial_4_12 NO_STORE
#define vstore_partial_4_13 NO_STORE
#define vstore_partial_4_14 NO_STORE
#define vstore_partial_4_15 NO_STORE
#define vstore_partial_4_16 NO_STORE
// Size == 8
#define vstore_partial_8_0 NO_STORE
#define vstore_partial_8_1 vstore_partial_1
#define vstore_partial_8_2 vstore_partial_2
#define vstore_partial_8_3 vstore_partial_3
#define vstore_partial_8_4 vstore_partial_4
#define vstore_partial_8_5 vstore_partial_5
#define vstore_partial_8_6 vstore_partial_6
#define vstore_partial_8_7 vstore_partial_7
#define vstore_partial_8_8 vstore_partial_8
#define vstore_partial_8_9 NO_STORE
#define vstore_partial_8_10 NO_STORE
#define vstore_partial_8_11 NO_STORE
#define vstore_partial_8_12 NO_STORE
#define vstore_partial_8_13 NO_STORE
#define vstore_partial_8_14 NO_STORE
#define vstore_partial_8_15 NO_STORE
#define vstore_partial_8_16 NO_STORE
// Size == 16
#define vstore_partial_16_0 NO_STORE
#define vstore_partial_16_1 vstore_partial_1
#define vstore_partial_16_2 vstore_partial_2
#define vstore_partial_16_3 vstore_partial_3
#define vstore_partial_16_4 vstore_partial_4
#define vstore_partial_16_5 vstore_partial_5
#define vstore_partial_16_6 vstore_partial_6
#define vstore_partial_16_7 vstore_partial_7
#define vstore_partial_16_8 vstore_partial_8
#define vstore_partial_16_9 vstore_partial_9
#define vstore_partial_16_10 vstore_partial_10
#define vstore_partial_16_11 vstore_partial_11
#define vstore_partial_16_12 vstore_partial_12
#define vstore_partial_16_13 vstore_partial_13
#define vstore_partial_16_14 vstore_partial_14
#define vstore_partial_16_15 vstore_partial_15
#define vstore_partial_16_16 vstore_partial_16
/** Partial vstore. Store the **lower** 0 to (n-1)th elements of the given vector while minimising the amount of vstore ops
* @name vstore_partial_n
*
* @note @p DATA needs to be a vector not a scalar
* @note n needs to be <= the vector width of the input variable @p DATA
* eg 1: Valid
* vstore_partial_15(var:float16, 0, 0xabcd);
* eg 2: Invalid
* vstore_partial_7(var:float4, 0, 0xabcd);
*
* @note in cases n == 1, 2, 3, 4, 8, 16, no extra vstore is invoked, thus there's no performance penalty.
*
* @param[in] DATA The name of the variable
* @param[in] OFFSET Offset in n
* @param[in] PTR The base pointer
* @{
*/
#define vstore_partial_1(DATA, OFFSET, PTR) \
vstore1(DATA.s0, OFFSET, PTR);
#define vstore_partial_2(DATA, OFFSET, PTR) \
vstore2(DATA.s01, OFFSET, PTR);
#define vstore_partial_3(DATA, OFFSET, PTR) \
vstore3(DATA.s012, OFFSET, PTR);
#define vstore_partial_4(DATA, OFFSET, PTR) \
vstore4(DATA.s0123, OFFSET, PTR);
#define vstore_partial_5(DATA, OFFSET, PTR) \
vstore_partial_4(DATA.s0123, OFFSET, PTR); \
vstore1(DATA.s4, OFFSET, PTR + 4);
#define vstore_partial_6(DATA, OFFSET, PTR) \
vstore_partial_4(DATA.s0123, OFFSET, PTR); \
vstore_partial_2(DATA.s45, OFFSET, PTR + 4);
#define vstore_partial_7(DATA, OFFSET, PTR) \
vstore_partial_4(DATA.s0123, OFFSET, PTR); \
vstore_partial_3(DATA.s456, OFFSET, PTR + 4);
#define vstore_partial_8(DATA, OFFSET, PTR) \
vstore8(DATA.s01234567, OFFSET, PTR);
#define vstore_partial_9(DATA, OFFSET, PTR) \
vstore_partial_8(DATA.s01234567, OFFSET, PTR); \
vstore1(DATA.s8, OFFSET, PTR + 8);
#define vstore_partial_10(DATA, OFFSET, PTR) \
vstore_partial_8(DATA.s01234567, OFFSET, PTR); \
vstore_partial_2(DATA.s89, OFFSET, PTR + 8);
#define vstore_partial_11(DATA, OFFSET, PTR) \
vstore_partial_8(DATA.s01234567, OFFSET, PTR); \
vstore_partial_3(DATA.s89a, OFFSET, PTR + 8);
#define vstore_partial_12(DATA, OFFSET, PTR) \
vstore_partial_8(DATA.s01234567, OFFSET, PTR); \
vstore_partial_4(DATA.s89ab, OFFSET, PTR + 8);
#define vstore_partial_13(DATA, OFFSET, PTR) \
vstore_partial_8(DATA.s01234567, OFFSET, PTR); \
vstore_partial_5(DATA.s89abcdef, OFFSET, PTR + 8);
#define vstore_partial_14(DATA, OFFSET, PTR) \
vstore_partial_8(DATA.s01234567, OFFSET, PTR); \
vstore_partial_6(DATA.s89abcdef, OFFSET, PTR + 8);
#define vstore_partial_15(DATA, OFFSET, PTR) \
vstore_partial_8(DATA.s01234567, OFFSET, PTR); \
vstore_partial_7(DATA.s89abcdef, OFFSET, PTR + 8);
#define vstore_partial_16(DATA, OFFSET, PTR) \
vstore16(DATA, OFFSET, PTR);
/** @} */ // end of groupd vstore_partial_n
/** @} */ // end of groupd VSTORE_PARTIAL
// Convert built-in functions with _sat modifier are not supported in floating point so we create defines
// without _sat to overcome this issue
#define convert_float_sat convert_float
#define convert_float1_sat convert_float
#define convert_float2_sat convert_float2
#define convert_float3_sat convert_float3
#define convert_float4_sat convert_float4
#define convert_float8_sat convert_float8
#define convert_float16_sat convert_float16
#define convert_half_sat convert_float
#define convert_half1_sat convert_half
#define convert_half2_sat convert_half2
#define convert_half3_sat convert_half3
#define convert_half4_sat convert_half4
#define convert_half8_sat convert_half8
#define convert_half16_sat convert_half16
#define convert_float1 convert_float
#define convert_half1 convert_half
#define convert_char1 convert_char
#define convert_uchar1 convert_uchar
#define convert_short1 convert_short
#define convert_ushort1 convert_ushort
#define convert_int1 convert_int
#define convert_uint1 convert_uint
#define convert_long1 convert_long
#define convert_ulong1 convert_ulong
#define convert_double1 convert_double
#define convert_char1_sat convert_char_sat
#define convert_uchar1_sat convert_uchar_sat
#define convert_short1_sat convert_short_sat
#define convert_ushort1_sat convert_ushort_sat
#define convert_int1_sat convert_int_sat
#define convert_uint1_sat convert_uint_sat
#define convert_long1_sat convert_long_sat
#define convert_ulong1_sat convert_ulong_sat
#define convert_double1_sat convert_double_sat
#define VEC_DATA_TYPE_STR(type, size) type##size
#define VEC_DATA_TYPE(type, size) VEC_DATA_TYPE_STR(type, size)
#define CONVERT_STR(x, type) (convert_##type((x)))
#define CONVERT(x, type) CONVERT_STR(x, type)
#define CONVERT_SAT_STR(x, type) (convert_##type##_sat((x)))
#define CONVERT_SAT(x, type) CONVERT_SAT_STR(x, type)
#define CONVERT_SAT_ROUND_STR(x, type, round) (convert_##type##_sat_##round((x)))
#define CONVERT_SAT_ROUND(x, type, round) CONVERT_SAT_ROUND_STR(x, type, round)
#define select_vec_dt_uchar(size) uchar##size
#define select_vec_dt_char(size) char##size
#define select_vec_dt_ushort(size) ushort##size
#define select_vec_dt_short(size) short##size
#define select_vec_dt_half(size) short##size
#define select_vec_dt_uint(size) uint##size
#define select_vec_dt_int(size) int##size
#define select_vec_dt_float(size) int##size
#define select_vec_dt_ulong(size) ulong##size
#define select_vec_dt_long(size) long##size
#define SELECT_VEC_DATA_TYPE_STR(type, size) select_vec_dt_##type(size)
#define SELECT_VEC_DATA_TYPE(type, size) SELECT_VEC_DATA_TYPE_STR(type, size)
#define SELECT_DATA_TYPE(type) SELECT_VEC_DATA_TYPE_STR(type, 1)
#define sum_reduce_1(x) (x)
#define sum_reduce_2(x) ((x).s0) + ((x).s1)
#define sum_reduce_3(x) sum_reduce_2((x).s01) + ((x).s2)
#define sum_reduce_4(x) sum_reduce_2((x).s01) + sum_reduce_2((x).s23)
#define sum_reduce_8(x) sum_reduce_4((x).s0123) + sum_reduce_4((x).s4567)
#define sum_reduce_16(x) sum_reduce_8((x).s01234567) + sum_reduce_8((x).s89ABCDEF)
#define SUM_REDUCE_STR(x, size) sum_reduce_##size(x)
#define SUM_REDUCE(x, size) SUM_REDUCE_STR(x, size)
#define max_reduce_1(x) (x)
#define max_reduce_2(x) max(((x).s0), ((x).s1))
#define max_reduce_3(x) max(max_reduce_2((x).s01), ((x).s2))
#define max_reduce_4(x) max(max_reduce_2((x).s01), max_reduce_2((x).s23))
#define max_reduce_8(x) max(max_reduce_4((x).s0123), max_reduce_4((x).s4567))
#define max_reduce_16(x) max(max_reduce_8((x).s01234567), max_reduce_8((x).s89ABCDEF))
#define MAX_REDUCE_STR(x, size) max_reduce_##size(x)
#define MAX_REDUCE(x, size) MAX_REDUCE_STR(x, size)
#define VECTOR_DECLARATION(name) \
__global uchar *name##_ptr, \
uint name##_stride_x, \
uint name##_step_x, \
uint name##_offset_first_element_in_bytes
#define IMAGE_DECLARATION(name) \
__global uchar *name##_ptr, \
uint name##_stride_x, \
uint name##_step_x, \
uint name##_stride_y, \
uint name##_step_y, \
uint name##_offset_first_element_in_bytes
#define TENSOR3D_DECLARATION(name) \
__global uchar *name##_ptr, \
uint name##_stride_x, \
uint name##_step_x, \
uint name##_stride_y, \
uint name##_step_y, \
uint name##_stride_z, \
uint name##_step_z, \
uint name##_offset_first_element_in_bytes
#define TENSOR4D_DECLARATION(name) \
__global uchar *name##_ptr, \
uint name##_stride_x, \
uint name##_step_x, \
uint name##_stride_y, \
uint name##_step_y, \
uint name##_stride_z, \
uint name##_step_z, \
uint name##_stride_w, \
uint name##_step_w, \
uint name##_offset_first_element_in_bytes
#define CONVERT_TO_VECTOR_STRUCT(name) \
update_vector_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x)
#define CONVERT_TO_VECTOR_STRUCT_NO_STEP(name) \
update_vector_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0)
#define CONVERT_TO_IMAGE_STRUCT(name) \
update_image_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y)
#define CONVERT_TO_IMAGE_STRUCT_NO_STEP(name) \
update_image_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0, name##_stride_y, 0)
#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name) \
update_image_from_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y, name##_stride_z, name##_step_z)
#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT_NO_STEP(name) \
update_image_from_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0, name##_stride_y, 0, name##_stride_z, name##_step_z)
#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name) \
update_image_from_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y, name##_stride_z, name##_step_z)
#define CONVERT_TO_TENSOR3D_STRUCT(name) \
update_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y, \
name##_stride_z, name##_step_z)
#define CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(name) \
update_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0, name##_stride_y, 0, name##_stride_z, 0)
#define CONVERT_TO_TENSOR4D_STRUCT(name, mod_size) \
update_tensor4D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y, \
name##_stride_z, name##_step_z, name##_stride_w, name##_step_w, mod_size)
#define CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(name, mod_size) \
update_tensor4D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0, name##_stride_y, 0, name##_stride_z, 0, name##_stride_w, 0, mod_size)
#define CONVERT_TO_TENSOR3D_STRUCT_NO_UPDATE_PTR(name) \
tensor3D_ptr_no_update(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y, \
name##_stride_z, name##_step_z)
/** Structure to hold Vector information */
typedef struct Vector
{
__global uchar *ptr; /**< Pointer to the starting postion of the buffer */
int offset_first_element_in_bytes; /**< The offset of the first element in the source image */
int stride_x; /**< Stride of the image in X dimension (in bytes) */
} Vector;
/** Structure to hold Image information */
typedef struct Image
{
__global uchar *ptr; /**< Pointer to the starting postion of the buffer */
int offset_first_element_in_bytes; /**< The offset of the first element in the source image */
int stride_x; /**< Stride of the image in X dimension (in bytes) */
int stride_y; /**< Stride of the image in Y dimension (in bytes) */
} Image;
/** Structure to hold 3D tensor information */
typedef struct Tensor3D
{
__global uchar *ptr; /**< Pointer to the starting postion of the buffer */
int offset_first_element_in_bytes; /**< The offset of the first element in the source image */
int stride_x; /**< Stride of the image in X dimension (in bytes) */
int stride_y; /**< Stride of the image in Y dimension (in bytes) */
int stride_z; /**< Stride of the image in Z dimension (in bytes) */
} Tensor3D;
/** Structure to hold 4D tensor information */
typedef struct Tensor4D
{
__global uchar *ptr; /**< Pointer to the starting postion of the buffer */
int offset_first_element_in_bytes; /**< The offset of the first element in the source image */
int stride_x; /**< Stride of the image in X dimension (in bytes) */
int stride_y; /**< Stride of the image in Y dimension (in bytes) */
int stride_z; /**< Stride of the image in Z dimension (in bytes) */
int stride_w; /**< Stride of the image in W dimension (in bytes) */
} Tensor4D;
/** Wrap vector information into an Vector structure, and make the pointer point at this workitem's data.
*
* @param[in] ptr Pointer to the starting postion of the buffer
* @param[in] offset_first_element_in_bytes The offset of the first element in the source vector
* @param[in] stride_x Stride of the vector in X dimension (in bytes)
* @param[in] step_x stride_x * number of elements along X processed per workitem(in bytes)
*
* @return An image object
*/
inline Vector update_vector_workitem_ptr(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x)
{
Vector vector =
{
.ptr = ptr,
.offset_first_element_in_bytes = offset_first_element_in_bytes,
.stride_x = stride_x,
};
vector.ptr += vector.offset_first_element_in_bytes + get_global_id(0) * step_x;
return vector;
}
/** Wrap image information into an Image structure, and make the pointer point at this workitem's data.
*
* @param[in] ptr Pointer to the starting postion of the buffer
* @param[in] offset_first_element_in_bytes The offset of the first element in the source image
* @param[in] stride_x Stride of the image in X dimension (in bytes)
* @param[in] step_x stride_x * number of elements along X processed per workitem(in bytes)
* @param[in] stride_y Stride of the image in Y dimension (in bytes)
* @param[in] step_y stride_y * number of elements along Y processed per workitem(in bytes)
*
* @return An image object
*/
inline Image update_image_workitem_ptr(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x, uint stride_y, uint step_y)
{
Image img =
{
.ptr = ptr,
.offset_first_element_in_bytes = offset_first_element_in_bytes,
.stride_x = stride_x,
.stride_y = stride_y
};
img.ptr += img.offset_first_element_in_bytes + get_global_id(0) * step_x + get_global_id(1) * step_y;
return img;
}
/** Wrap 3D tensor information into an image structure, and make the pointer point at this workitem's data.
*
* @param[in] ptr Pointer to the starting postion of the buffer
* @param[in] offset_first_element_in_bytes The offset of the first element in the source image
* @param[in] stride_x Stride of the image in X dimension (in bytes)
* @param[in] step_x stride_x * number of elements along X processed per workitem(in bytes)
* @param[in] stride_y Stride of the image in Y dimension (in bytes)
* @param[in] step_y stride_y * number of elements along Y processed per workitem(in bytes)
* @param[in] stride_z Stride of the image in Z dimension (in bytes)
* @param[in] step_z stride_z * number of elements along Z processed per workitem(in bytes)
*
* @return A 3D tensor object
*/
inline Image update_image_from_tensor3D_workitem_ptr(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x, uint stride_y, uint step_y, uint stride_z, uint step_z)
{
Image img =
{
.ptr = ptr,
.offset_first_element_in_bytes = offset_first_element_in_bytes,
.stride_x = stride_x,
.stride_y = stride_y
};
img.ptr += img.offset_first_element_in_bytes + get_global_id(0) * step_x + get_global_id(1) * step_y + get_global_id(2) * step_z;
return img;
}
/** Wrap 3D tensor information into an tensor structure, and make the pointer point at this workitem's data.
*
* @param[in] ptr Pointer to the starting postion of the buffer
* @param[in] offset_first_element_in_bytes The offset of the first element in the source image
* @param[in] stride_x Stride of the image in X dimension (in bytes)
* @param[in] step_x stride_x * number of elements along X processed per workitem(in bytes)
* @param[in] stride_y Stride of the image in Y dimension (in bytes)
* @param[in] step_y stride_y * number of elements along Y processed per workitem(in bytes)
* @param[in] stride_z Stride of the image in Z dimension (in bytes)
* @param[in] step_z stride_z * number of elements along Z processed per workitem(in bytes)
*
* @return A 3D tensor object
*/
inline Tensor3D update_tensor3D_workitem_ptr(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x, uint stride_y, uint step_y, uint stride_z, uint step_z)
{
Tensor3D tensor =
{
.ptr = ptr,
.offset_first_element_in_bytes = offset_first_element_in_bytes,
.stride_x = stride_x,
.stride_y = stride_y,
.stride_z = stride_z
};
tensor.ptr += tensor.offset_first_element_in_bytes + get_global_id(0) * step_x + get_global_id(1) * step_y + get_global_id(2) * step_z;
return tensor;
}
/** Wrap 3D tensor information into an tensor structure.
*
* @param[in] ptr Pointer to the starting postion of the buffer
* @param[in] offset_first_element_in_bytes The offset of the first element in the source image
* @param[in] stride_x Stride of the image in X dimension (in bytes)
* @param[in] step_x stride_x * number of elements along X processed per workitem(in bytes)
* @param[in] stride_y Stride of the image in Y dimension (in bytes)
* @param[in] step_y stride_y * number of elements along Y processed per workitem(in bytes)
* @param[in] stride_z Stride of the image in Z dimension (in bytes)
* @param[in] step_z stride_z * number of elements along Z processed per workitem(in bytes)
*
* @return A 3D tensor object
*/
inline Tensor3D tensor3D_ptr_no_update(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x, uint stride_y, uint step_y, uint stride_z, uint step_z)
{
Tensor3D tensor =
{
.ptr = ptr,
.offset_first_element_in_bytes = offset_first_element_in_bytes,
.stride_x = stride_x,
.stride_y = stride_y,
.stride_z = stride_z
};
return tensor;
}
inline Tensor4D update_tensor4D_workitem_ptr(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x, uint stride_y, uint step_y, uint stride_z, uint step_z, uint stride_w,
uint step_w,
uint mod_size)
{
Tensor4D tensor =
{
.ptr = ptr,
.offset_first_element_in_bytes = offset_first_element_in_bytes,
.stride_x = stride_x,
.stride_y = stride_y,
.stride_z = stride_z,
.stride_w = stride_w
};
tensor.ptr += tensor.offset_first_element_in_bytes + get_global_id(0) * step_x + get_global_id(1) * step_y + (get_global_id(2) % mod_size) * step_z + (get_global_id(2) / mod_size) * step_w;
return tensor;
}
/** Get the pointer position of a Vector
*
* @param[in] vec Pointer to the starting position of the buffer
* @param[in] x Relative X position
*/
inline __global const uchar *vector_offset(const Vector *vec, int x)
{
return vec->ptr + x * vec->stride_x;
}
/** Get the pointer position of a Image
*
* @param[in] img Pointer to the starting position of the buffer
* @param[in] x Relative X position
* @param[in] y Relative Y position
*/
inline __global uchar *offset(const Image *img, int x, int y)
{
return img->ptr + x * img->stride_x + y * img->stride_y;
}
/** Get the pointer position of a Tensor3D
*
* @param[in] tensor Pointer to the starting position of the buffer
* @param[in] x Relative X position
* @param[in] y Relative Y position
* @param[in] z Relative Z position
*/
inline __global const uchar *tensor3D_offset(const Tensor3D *tensor, int x, int y, int z)
{
return tensor->ptr + x * tensor->stride_x + y * tensor->stride_y + z * tensor->stride_z;
}
/** Get the pointer position of a Tensor4D
*
* @param[in] tensor Pointer to the starting position of the buffer
* @param[in] x Relative X position
* @param[in] y Relative Y position
* @param[in] z Relative Z position
* @param[in] w Relative W position
*/
inline __global const uchar *tensor4D_offset(const Tensor4D *tensor, int x, int y, int z, int w)
{
return tensor->ptr + x * tensor->stride_x + y * tensor->stride_y + z * tensor->stride_z + w * tensor->stride_w;
}
/** Get the offset for a given linear index of a Tensor3D
*
* @param[in] tensor Pointer to the starting position of the buffer
* @param[in] width Width of the input tensor
* @param[in] height Height of the input tensor
* @param[in] depth Depth of the input tensor
* @param[in] index Linear index
*/
inline __global const uchar *tensor3D_index2ptr(const Tensor3D *tensor, uint width, uint height, uint depth, uint index)
{
uint num_elements = width * height;
const uint z = index / num_elements;
index %= num_elements;
const uint y = index / width;
index %= width;
const uint x = index;
return tensor->ptr + x * tensor->stride_x + y * tensor->stride_y + z * tensor->stride_z + tensor->offset_first_element_in_bytes;
}
#endif // _HELPER_H
/***********************************************/
/* Begin implementation of Sobel3x3 filter */
/***********************************************/
/** This OpenCL kernel that computes a Sobel3x3 filter.
*
* @attention To enable computation of the X gradient -DGRAD_X must be passed at compile time, while computation of the Y gradient
* is performed when -DGRAD_Y is used. You can use both when computation of both gradients is required.
*
* @param[in] src_ptr Pointer to the source image. Supported data types: U8
* @param[in] src_stride_x Stride of the source image in X dimension (in bytes)
* @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
* @param[in] src_stride_y Stride of the source image in Y dimension (in bytes)
* @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
* @param[in] src_offset_first_element_in_bytes The offset of the first element in the source image
* @param[out] dst_gx_ptr Pointer to the destination image. Supported data types: S16
* @param[in] dst_gx_stride_x Stride of the destination image in X dimension (in bytes)
* @param[in] dst_gx_step_x dst_gx_stride_x * number of elements along X processed per workitem(in bytes)
* @param[in] dst_gx_stride_y Stride of the destination image in Y dimension (in bytes)
* @param[in] dst_gx_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes)
* @param[in] dst_gx_offset_first_element_in_bytes The offset of the first element in the destination image
* @param[out] dst_gy_ptr Pointer to the destination image. Supported data types: S16
* @param[in] dst_gy_stride_x Stride of the destination image in X dimension (in bytes)
* @param[in] dst_gy_step_x dst_gy_stride_x * number of elements along X processed per workitem(in bytes)
* @param[in] dst_gy_stride_y Stride of the destination image in Y dimension (in bytes)
* @param[in] dst_gy_step_y dst_gy_stride_y * number of elements along Y processed per workitem(in bytes)
* @param[in] dst_gy_offset_first_element_in_bytes The offset of the first element in the destination image
*/
__kernel void sobel3x3(
IMAGE_DECLARATION(src)
#ifdef GRAD_X
,
IMAGE_DECLARATION(dst_gx)
#endif /* GRAD_X */
#ifdef GRAD_Y
,
IMAGE_DECLARATION(dst_gy)
#endif /* GRAD_Y */
)
{
Image src = CONVERT_TO_IMAGE_STRUCT(src);
#ifdef GRAD_X
Image dst_gx = CONVERT_TO_IMAGE_STRUCT(dst_gx);
#endif /* GRAD_X */
#ifdef GRAD_Y
Image dst_gy = CONVERT_TO_IMAGE_STRUCT(dst_gy);
#endif /* GRAD_Y */
// Output pixels
#ifdef GRAD_X
short8 gx = (short8)0;
#endif /* GRAD_X */
#ifdef GRAD_Y
short8 gy = (short8)0;
#endif /* GRAD_Y */
// Row0
uchar16 temp = vload16(0, offset(&src, -1, -1));
short8 left = convert_short8(temp.s01234567);
short8 middle = convert_short8(temp.s12345678);
short8 right = convert_short8(temp.s23456789);
#ifdef GRAD_X
gx += left * (short8)(-1);
gx += right * (short8)(+1);
#endif /* GRAD_X */
#ifdef GRAD_Y
gy += left * (short8)(-1);
gy += middle * (short8)(-2);
gy += right * (short8)(-1);
#endif /* GRAD_Y */
// Row1
temp = vload16(0, offset(&src, -1, 0));
left = convert_short8(temp.s01234567);
right = convert_short8(temp.s23456789);
#ifdef GRAD_X
gx += left * (short8)(-2);
gx += right * (short8)(+2);
#endif /* GRAD_X */
// Row2
temp = vload16(0, offset(&src, -1, 1));
left = convert_short8(temp.s01234567);
middle = convert_short8(temp.s12345678);
right = convert_short8(temp.s23456789);
#ifdef GRAD_X
gx += left * (short8)(-1);
gx += right * (short8)(+1);
#endif /* GRAD_X */
#ifdef GRAD_Y
gy += left * (short8)(+1);
gy += middle * (short8)(+2);
gy += right * (short8)(+1);
#endif /* GRAD_Y */
// Store results
#ifdef GRAD_X
vstore8(gx, 0, ((__global short *)dst_gx.ptr));
#endif /* GRAD_X */
#ifdef GRAD_Y
vstore8(gy, 0, ((__global short *)dst_gy.ptr));
#endif /* GRAD_Y */
}
/**********************************************/
/* End implementation of Sobel3x3 filter */
/**********************************************/
/***********************************************/
/* Begin implementation of Sobel5x5 filter */
/***********************************************/
/** Compute a 1D horizontal sobel filter 1x5 for 8 bytes assuming the input is made of 1 channel of 1 byte (i.e 8 pixels).
*
* @param[in] src Pointer to source image.
* @param[in] left1_coeff_gx Weight of the most left pixel for gx
* @param[in] left2_coeff_gx Weight of the left pixel for gx
* @param[in] middle_coeff_gx Weight of the middle pixel for gx
* @param[in] right1_coeff_gx Weight of the right pixel for gx
* @param[in] right2_coeff_gx Weight of the most right pixel for gx
* @param[in] left1_coeff_gy Weight of the most left pixel for gy
* @param[in] left2_coeff_gy Weight of the left pixel for gy
* @param[in] middle_coeff_gy Weight of the middle pixel for gy
* @param[in] right1_coeff_gy Weight of the right pixel for gy
* @param[in] right2_coeff_gy Weight of the most right pixel for gy
*
* @return a short16 containing short8 gx and short8 gy values.
*/
short16 sobel1x5(
Image *src,
const short left1_coeff_gx,
const short left2_coeff_gx,
const short middle_coeff_gx,
const short right1_coeff_gx,
const short right2_coeff_gx,
const short left1_coeff_gy,
const short left2_coeff_gy,
const short middle_coeff_gy,
const short right1_coeff_gy,
const short right2_coeff_gy)
{
uchar16 temp = vload16(0, offset(src, -2, 0));
short8 gx = 0;
short8 gy = 0;
short8 val;
val = convert_short8(temp.s01234567);
gx += val * (short8)left1_coeff_gx;
gy += val * (short8)left1_coeff_gy;
val = convert_short8(temp.s12345678);
gx += val * (short8)left2_coeff_gx;
gy += val * (short8)left2_coeff_gy;
val = convert_short8(temp.s23456789);
gx += val * (short8)middle_coeff_gx;
gy += val * (short8)middle_coeff_gy;
val = convert_short8(temp.s3456789a);
gx += val * (short8)right1_coeff_gx;
gy += val * (short8)right1_coeff_gy;
val = convert_short8(temp.s456789ab);
gx += val * (short8)right2_coeff_gx;
gy += val * (short8)right2_coeff_gy;
return (short16)(gx, gy);
}
/** Compute a 1D vertical sobel filter 5x1 for 8 bytes assuming the input is made of 1 channel of 1 byte (i.e 8 pixels).
*
* @param[in] src Pointer to source image.
* @param[in] up1_coeff Weight of the most up pixel
* @param[in] up2_coeff Weight of the up pixel
* @param[in] middle_coeff Weight of the middle pixel
* @param[in] down1_coeff Weight of the down pixel
* @param[in] down2_coeff Weight of the most down pixel
*
* @return a short8 containing 8 convoluted values.
*/
short8 sobel5x1(
Image *src,
const short up1_coeff,
const short up2_coeff,
const short middle_coeff,
const short down1_coeff,
const short down2_coeff)
{
short8 val;
short8 out = (short8)0;
val = vload8(0, (__global short *)offset(src, 0, -2));
out += val * (short8)up1_coeff;
val = vload8(0, (__global short *)offset(src, 0, -1));
out += val * (short8)up2_coeff;
val = vload8(0, (__global short *)offset(src, 0, 0));
out += val * (short8)middle_coeff;
val = vload8(0, (__global short *)offset(src, 0, 1));
out += val * (short8)down1_coeff;
val = vload8(0, (__global short *)offset(src, 0, 2));
out += val * (short8)down2_coeff;
return (short8)(out);
}
/** Apply a 1x5 sobel matrix to a single channel U8 input image and output two temporary channel S16 images.
*
* @attention To enable computation of the X gradient -DGRAD_X must be passed at compile time, while computation of the Y gradient
* is performed when -DGRAD_Y is used. You can use both when computation of both gradients is required.
*
* @param[in] src_ptr Pointer to the source image.. Supported data types: U8
* @param[in] src_stride_x Stride of the source image in X dimension (in bytes)
* @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
* @param[in] src_stride_y Stride of the source image in Y dimension (in bytes)
* @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
* @param[in] src_offset_first_element_in_bytes The offset of the first element in the source image
* @param[out] dst_gx_ptr Pointer to the destination image.. Supported data types: S16
* @param[in] dst_gx_stride_x Stride of the destination image in X dimension (in bytes)
* @param[in] dst_gx_step_x dst_gx_stride_x * number of elements along X processed per workitem(in bytes)
* @param[in] dst_gx_stride_y Stride of the destination image in Y dimension (in bytes)
* @param[in] dst_gx_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes)
* @param[in] dst_gx_offset_first_element_in_bytes The offset of the first element in the destination image
* @param[out] dst_gy_ptr Pointer to the destination image. Supported data types: S16
* @param[in] dst_gy_stride_x Stride of the destination image in X dimension (in bytes)
* @param[in] dst_gy_step_x dst_gy_stride_x * number of elements along X processed per workitem(in bytes)
* @param[in] dst_gy_stride_y Stride of the destination image in Y dimension (in bytes)
* @param[in] dst_gy_step_y dst_gy_stride_y * number of elements along Y processed per workitem(in bytes)
* @param[in] dst_gy_offset_first_element_in_bytes The offset of the first element in the destination image
*/
__kernel void sobel_separable1x5(
IMAGE_DECLARATION(src)
#ifdef GRAD_X
,
IMAGE_DECLARATION(dst_gx)
#endif /* GRAD_X */
#ifdef GRAD_Y
,
IMAGE_DECLARATION(dst_gy)
#endif /* GRAD_Y */
)
{
Image src = CONVERT_TO_IMAGE_STRUCT(src);
#ifdef GRAD_X
Image dst_gx = CONVERT_TO_IMAGE_STRUCT(dst_gx);
#endif /* GRAD_X */
#ifdef GRAD_Y
Image dst_gy = CONVERT_TO_IMAGE_STRUCT(dst_gy);
#endif /* GRAD_Y */
// Output pixels
short16 gx_gy = sobel1x5(&src,
-1, -2, 0, 2, 1,
1, 4, 6, 4, 1);
// Store result in dst
#ifdef GRAD_X
vstore8(gx_gy.s01234567, 0, ((__global short *)dst_gx.ptr));
#endif /* GRAD_X */
#ifdef GRAD_Y
vstore8(gx_gy.s89ABCDEF, 0, ((__global short *)dst_gy.ptr));
#endif /* GRAD_Y */
}
/** Apply a 5x1 convolution matrix to two single channel S16 input temporary images
* and output two single channel S16 images.
*
* @attention To enable computation of the X gradient -DGRAD_X must be passed at compile time, while computation of the Y gradient
* is performed when -DGRAD_Y is used. You can use both when computation of both gradients is required.
*
* @param[in] src_x_ptr Pointer to the source image.. Supported data types: S16
* @param[in] src_x_stride_x Stride of the source image in X dimension (in bytes)
* @param[in] src_x_step_x src_x_stride_x * number of elements along X processed per workitem(in bytes)
* @param[in] src_x_stride_y Stride of the source image in Y dimension (in bytes)
* @param[in] src_x_step_y src_x_stride_y * number of elements along Y processed per workitem(in bytes)
* @param[in] src_x_offset_first_element_in_bytes The offset of the first element in the source image
* @param[out] dst_gx_ptr Pointer to the destination image. Supported data types: S16
* @param[in] dst_gx_stride_x Stride of the destination image in X dimension (in bytes)
* @param[in] dst_gx_step_x dst_gx_stride_x * number of elements along X processed per workitem(in bytes)
* @param[in] dst_gx_stride_y Stride of the destination image in Y dimension (in bytes)
* @param[in] dst_gx_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes)
* @param[in] dst_gx_offset_first_element_in_bytes The offset of the first element in the destination image
* @param[in] src_y_ptr Pointer to the source image. Supported data types: S16
* @param[in] src_y_stride_x Stride of the source image in X dimension (in bytes)
* @param[in] src_y_step_x src_y_stride_x * number of elements along X processed per workitem(in bytes)
* @param[in] src_y_stride_y Stride of the source image in Y dimension (in bytes)
* @param[in] src_y_step_y src_y_stride_y * number of elements along Y processed per workitem(in bytes)
* @param[in] src_y_offset_first_element_in_bytes The offset of the first element in the source image
* @param[out] dst_gy_ptr Pointer to the destination image. Supported data types: S16
* @param[in] dst_gy_stride_x Stride of the destination image in X dimension (in bytes)
* @param[in] dst_gy_step_x dst_gy_stride_x * number of elements along X processed per workitem(in bytes)
* @param[in] dst_gy_stride_y Stride of the destination image in Y dimension (in bytes)
* @param[in] dst_gy_step_y dst_gy_stride_y * number of elements along Y processed per workitem(in bytes)
* @param[in] dst_gy_offset_first_element_in_bytes The offset of the first element in the destination image
* @param[in] dummy Dummy parameter to easy conditional inclusion
*/
__kernel void sobel_separable5x1(
#ifdef GRAD_X
IMAGE_DECLARATION(src_x),
IMAGE_DECLARATION(dst_gx),
#endif /* GRAD_X */
#ifdef GRAD_Y
IMAGE_DECLARATION(src_y),
IMAGE_DECLARATION(dst_gy),
#endif /* GRAD_Y */
int dummy)
{
#ifdef GRAD_X
Image src_x = CONVERT_TO_IMAGE_STRUCT(src_x);
Image dst_gx = CONVERT_TO_IMAGE_STRUCT(dst_gx);
#endif /* GRAD_X */
#ifdef GRAD_Y
Image src_y = CONVERT_TO_IMAGE_STRUCT(src_y);
Image dst_gy = CONVERT_TO_IMAGE_STRUCT(dst_gy);
#endif /* GRAD_Y */
#ifdef GRAD_X
short8 gx = sobel5x1(&src_x,
1, 4, 6, 4, 1);
vstore8(gx, 0, ((__global short *)dst_gx.ptr));
#endif /* GRAD_X */
#ifdef GRAD_Y
short8 gy = sobel5x1(&src_y,
-1, -2, 0, 2, 1);
vstore8(gy, 0, ((__global short *)dst_gy.ptr));
#endif /* GRAD_Y */
}
/**********************************************/
/* End implementation of Sobel5x5 filter */
/**********************************************/
/***********************************************/
/* Begin implementation of Sobel7x7 filter */
/***********************************************/
/* Sobel 1x7 horizontal X / 7x1 vertical Y coefficients */
#define X0 -1
#define X1 -4
#define X2 -5
#define X3 0
#define X4 5
#define X5 4
#define X6 1
/* Sobel 1x7 vertical X / 7x1 horizontal Y coefficients */
#define Y0 1
#define Y1 6
#define Y2 15
#define Y3 20
#define Y4 15
#define Y5 6
#define Y6 1
/* Calculates single horizontal iteration. */
#define SOBEL1x1_HOR(src, gx, gy, idx) \
{ \
int8 val = convert_int8(vload8(0, offset(src, idx - 3, 0))); \
gx += val * X##idx; \
gy += val * Y##idx; \
}
/* Calculates single vertical iteration. */
#define SOBEL1x1_VERT(src, g, direction, idx) \
{ \
int8 val = vload8(0, (__global int *)offset(src, 0, idx - 3)); \
g += val * (int8)direction##idx; \
}
/* Calculates a 1x7 horizontal iteration. */
#define SOBEL1x7(ptr, gx, gy) \
SOBEL1x1_HOR(ptr, gx, gy, 0) \
SOBEL1x1_HOR(ptr, gx, gy, 1) \
SOBEL1x1_HOR(ptr, gx, gy, 2) \
SOBEL1x1_HOR(ptr, gx, gy, 3) \
SOBEL1x1_HOR(ptr, gx, gy, 4) \
SOBEL1x1_HOR(ptr, gx, gy, 5) \
SOBEL1x1_HOR(ptr, gx, gy, 6)
/* Calculates a 7x1 vertical iteration. */
#define SOBEL7x1(ptr, g, direction) \
SOBEL1x1_VERT(ptr, g, direction, 0) \
SOBEL1x1_VERT(ptr, g, direction, 1) \
SOBEL1x1_VERT(ptr, g, direction, 2) \
SOBEL1x1_VERT(ptr, g, direction, 3) \
SOBEL1x1_VERT(ptr, g, direction, 4) \
SOBEL1x1_VERT(ptr, g, direction, 5) \
SOBEL1x1_VERT(ptr, g, direction, 6)
/** Apply a 1x7 sobel matrix to a single channel U8 input image and output two temporary channel S16 images and leave the borders undefined.
*
* @attention To enable computation of the X gradient -DGRAD_X must be passed at compile time, while computation of the Y gradient
* is performed when -DGRAD_Y is used. You can use both when computation of both gradients is required.
*
* @param[in] src_ptr Pointer to the source image. Supported data types: U8
* @param[in] src_stride_x Stride of the source image in X dimension (in bytes)
* @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
* @param[in] src_stride_y Stride of the source image in Y dimension (in bytes)
* @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
* @param[in] src_offset_first_element_in_bytes The offset of the first element in the source image
* @param[out] dst_gx_ptr Pointer to the destination image. Supported data types: S32
* @param[in] dst_gx_stride_x Stride of the destination image in X dimension (in bytes)
* @param[in] dst_gx_step_x dst_gx_stride_x * number of elements along X processed per workitem(in bytes)
* @param[in] dst_gx_stride_y Stride of the destination image in Y dimension (in bytes)
* @param[in] dst_gx_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes)
* @param[in] dst_gx_offset_first_element_in_bytes The offset of the first element in the destination image
* @param[out] dst_gy_ptr Pointer to the destination image. Supported data types: S32
* @param[in] dst_gy_stride_x Stride of the destination image in X dimension (in bytes)
* @param[in] dst_gy_step_x dst_gy_stride_x * number of elements along X processed per workitem(in bytes)
* @param[in] dst_gy_stride_y Stride of the destination image in Y dimension (in bytes)
* @param[in] dst_gy_step_y dst_gy_stride_y * number of elements along Y processed per workitem(in bytes)
* @param[in] dst_gy_offset_first_element_in_bytes The offset of the first element in the destination image
*/
__kernel void sobel_separable1x7(
IMAGE_DECLARATION(src)
#ifdef GRAD_X
,
IMAGE_DECLARATION(dst_gx)
#endif /* GRAD_X */
#ifdef GRAD_Y
,
IMAGE_DECLARATION(dst_gy)
#endif /* GRAD_Y */
)
{
Image src = CONVERT_TO_IMAGE_STRUCT(src);
#ifdef GRAD_X
Image dst_gx = CONVERT_TO_IMAGE_STRUCT(dst_gx);
#endif /* GRAD_X */
#ifdef GRAD_Y
Image dst_gy = CONVERT_TO_IMAGE_STRUCT(dst_gy);
#endif /* GRAD_Y */
int8 gx = (int8)0;
int8 gy = (int8)0;
SOBEL1x7(&src, gx, gy);
// Store result in dst
#ifdef GRAD_X
vstore8(gx, 0, ((__global int *)dst_gx.ptr));
#endif /* GRAD_X */
#ifdef GRAD_Y
vstore8(gy, 0, ((__global int *)dst_gy.ptr));
#endif /* GRAD_Y */
}
/** Apply a 7x1 convolution matrix to two single channel S16 input temporary images and output two single channel S16 images and leave the borders undefined.
*
* @attention To enable computation of the X gradient -DGRAD_X must be passed at compile time, while computation of the Y gradient
* is performed when -DGRAD_Y is used. You can use both when computation of both gradients is required.
*
* @param[in] src_x_ptr Pointer to the source image. Supported data types: S32
* @param[in] src_x_stride_x Stride of the source image in X dimension (in bytes)
* @param[in] src_x_step_x src_x_stride_x * number of elements along X processed per workitem(in bytes)
* @param[in] src_x_stride_y Stride of the source image in Y dimension (in bytes)
* @param[in] src_x_step_y src_x_stride_y * number of elements along Y processed per workitem(in bytes)
* @param[in] src_x_offset_first_element_in_bytes The offset of the first element in the source image
* @param[out] dst_gx_ptr Pointer to the destination image. Supported data types: S16
* @param[in] dst_gx_stride_x Stride of the destination image in X dimension (in bytes)
* @param[in] dst_gx_step_x dst_gx_stride_x * number of elements along X processed per workitem(in bytes)
* @param[in] dst_gx_stride_y Stride of the destination image in Y dimension (in bytes)
* @param[in] dst_gx_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes)
* @param[in] dst_gx_offset_first_element_in_bytes The offset of the first element in the destination image
* @param[in] src_y_ptr Pointer to the source image. Supported data types: S32
* @param[in] src_y_stride_x Stride of the source image in X dimension (in bytes)
* @param[in] src_y_step_x src_y_stride_x * number of elements along X processed per workitem(in bytes)
* @param[in] src_y_stride_y Stride of the source image in Y dimension (in bytes)
* @param[in] src_y_step_y src_y_stride_y * number of elements along Y processed per workitem(in bytes)
* @param[in] src_y_offset_first_element_in_bytes The offset of the first element in the source image
* @param[out] dst_gy_ptr Pointer to the destination image. Supported data types: S16
* @param[in] dst_gy_stride_x Stride of the destination image in X dimension (in bytes)
* @param[in] dst_gy_step_x dst_gy_stride_x * number of elements along X processed per workitem(in bytes)
* @param[in] dst_gy_stride_y Stride of the destination image in Y dimension (in bytes)
* @param[in] dst_gy_step_y dst_gy_stride_y * number of elements along Y processed per workitem(in bytes)
* @param[in] dst_gy_offset_first_element_in_bytes The offset of the first element in the destination image
* @param[in] dummy Dummy parameter to easy conditional inclusion
*/
__kernel void sobel_separable7x1(
#ifdef GRAD_X
IMAGE_DECLARATION(src_x),
IMAGE_DECLARATION(dst_gx),
#endif /* GRAD_X */
#ifdef GRAD_Y
IMAGE_DECLARATION(src_y),
IMAGE_DECLARATION(dst_gy),
#endif /* GRAD_Y */
int dummy)
{
#ifdef GRAD_X
Image src_x = CONVERT_TO_IMAGE_STRUCT(src_x);
Image dst_gx = CONVERT_TO_IMAGE_STRUCT(dst_gx);
#endif /* GRAD_X */
#ifdef GRAD_Y
Image src_y = CONVERT_TO_IMAGE_STRUCT(src_y);
Image dst_gy = CONVERT_TO_IMAGE_STRUCT(dst_gy);
#endif /* GRAD_Y */
// Output pixels
#ifdef GRAD_X
int8 gx = 0;
SOBEL7x1(&src_x, gx, Y);
vstore8(gx, 0, (__global int *)dst_gx.ptr);
#endif /* GRAD_X */
#ifdef GRAD_Y
int8 gy = 0;
SOBEL7x1(&src_y, gy, X);
vstore8(gy, 0, (__global int *)dst_gy.ptr);
#endif /* GRAD_Y */
}
/**********************************************/
/* End implementation of Sobel7x7 filter */
/**********************************************/
)"