From 568df1b6a4bc9eac1c62b179cb55909c956ca496 Mon Sep 17 00:00:00 2001 From: Xianzhi Yu Date: Thu, 25 May 2023 17:35:23 +0800 Subject: [PATCH 1/4] delete unused file delete unused file --- .../src/gpu/mali/cl/lut_yuvnv21_preprocess.cl | 48 ------------------- 1 file changed, 48 deletions(-) delete mode 100644 compute/image/src/gpu/mali/cl/lut_yuvnv21_preprocess.cl diff --git a/compute/image/src/gpu/mali/cl/lut_yuvnv21_preprocess.cl b/compute/image/src/gpu/mali/cl/lut_yuvnv21_preprocess.cl deleted file mode 100644 index 4de4dc93..00000000 --- a/compute/image/src/gpu/mali/cl/lut_yuvnv21_preprocess.cl +++ /dev/null @@ -1,48 +0,0 @@ -// Copyright (C) 2019. Huawei Technologies Co., Ltd. All rights reserved. - -// 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. - -#include "kernel_def.h" - -__kernel void KERNEL_NAME(__global const uchar *_srcptr, - __global T *_dstptr, - const int src_chw, - const int src_row, - const int src_step, - const int dst_n, - const int dst_h, - const int dst_w, - const int dst_chw, - const int dst_hw) -{ - int x = get_global_id(0); - int y = get_global_id(1); - int z = get_global_id(2); - if (x >= dst_w || y >= dst_h || z >= dst_n) { - return; - } - - __global const uchar *srcptr = _srcptr + z * src_chw; - __global T *dstptr = _dstptr + z * dst_chw; - - __global const uchar *ysrc = srcptr + mad24(y << 1, src_step, (x << 1)); - __global const uchar *usrc = srcptr + mad24(src_row + y, src_step, (x << 1)); - - T3 r = (T3)(ysrc[1], usrc[0], usrc[1]); - T3 scale = (T3)0.003921; - r *= scale; - - int index = mad24(y, dst_w, x); - dstptr[index] = r.s0; - dstptr[index + dst_hw] = r.s1; - dstptr[index + (dst_hw << 1)] = r.s2; -} From c27780033fbd0cf225aae6a31de538f638e2de37 Mon Sep 17 00:00:00 2001 From: Xianzhi Yu Date: Thu, 25 May 2023 17:36:25 +0800 Subject: [PATCH 2/4] delete unused file delete unused file --- .../src/gpu/mali/cl/lut_yuvnv21_trilinear.cl | 138 ------------------ 1 file changed, 138 deletions(-) delete mode 100644 compute/image/src/gpu/mali/cl/lut_yuvnv21_trilinear.cl diff --git a/compute/image/src/gpu/mali/cl/lut_yuvnv21_trilinear.cl b/compute/image/src/gpu/mali/cl/lut_yuvnv21_trilinear.cl deleted file mode 100644 index 5ad6ee41..00000000 --- a/compute/image/src/gpu/mali/cl/lut_yuvnv21_trilinear.cl +++ /dev/null @@ -1,138 +0,0 @@ -// Copyright (C) 2019. Huawei Technologies Co., Ltd. All rights reserved. - -// 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. - -#include "kernel_def.h" - -inline int4 get_id1(int x, int3 y) -{ - int4 v; - v.s0 = x + y.s2; - v.s1 = v.s0 + y.s0; - v.s2 = v.s0 + y.s1; - v.s3 = v.s2 + y.s0; - return v; -} - -inline T8 get_lut(__global const T *lut, int4 id, int offset) -{ - int4 tmp_id = id + offset; - T8 v; - v.s01 = vload2(0, lut + tmp_id.s0); - v.s23 = vload2(0, lut + tmp_id.s1); - v.s45 = vload2(0, lut + tmp_id.s2); - v.s67 = vload2(0, lut + tmp_id.s3); - return v; -} - -inline T4 get_ww1(float g_d, float b_d) -{ - T4 ww; - T2 tmp_gd = (T2)(1 - g_d, g_d); - ww.s01 = tmp_gd * (T2)(1 - b_d); - ww.s23 = tmp_gd * (T2)b_d; - return ww; -} - -inline T8 get_ww2(float r_d, T4 ww1) -{ - T8 ww; - ww.s0246 = (T4)(1 - r_d) * ww1; - ww.s1357 = (T4)r_d * ww1; - return ww; -} - -__kernel void KERNEL_NAME(__global const uchar *_image, - __global const T *lut, - __global uchar *_output, - const int dim, - const int shift, - const float binsize, - const int src_step, - const int rows, - const int src_n, - const int src_chw, - const int dst_chw) -{ - int x = get_global_id(0); - int y = get_global_id(1); - int z = get_global_id(2); - int x_2 = x << 1; - int y_2 = y << 1; - if (x_2 >= src_step || y_2 >= rows || z >= src_n) { - return; - } - __global const uchar *image = _image + z * src_chw; - __global uchar *output = _output + z * dst_chw; - - int index_y = mad24(y_2, src_step, x_2); - int index_u = mad24(rows + y, src_step, x_2); - __global const uchar *ysrc = image + index_y; - __global const uchar *usrc = image + index_u; - __global uchar *ydst = output + index_y; - __global uchar *udst = output + index_u; - - float4 ydata; - float2 uvdata; - ydata.s01 = (float2)(ysrc[0], ysrc[1]) / 255.0f; - ydata.s23 = (float2)(ysrc[src_step], ysrc[src_step + 1]) / 255.0f; - uvdata = (float2)(usrc[0], usrc[1]) / 255.0f; - - float4 y_gid = ydata / binsize; - float2 uv_gid = uvdata / binsize; - - float4 y_fid = floor(y_gid); - float2 uv_fid = floor(uv_gid); - - float4 y_dd = y_gid - y_fid; - float2 uv_dd = uv_gid - uv_fid; - - int3 tmp_id; - tmp_id.s0 = dim; - tmp_id.s1 = dim * dim; - tmp_id.s2 = tmp_id.s0 * uv_fid.s0 + tmp_id.s1 * uv_fid.s1; - - int4 id1 = get_id1(y_fid.s0, tmp_id); - int4 id2 = get_id1(y_fid.s1, tmp_id); - int4 id3 = get_id1(y_fid.s2, tmp_id); - int4 id4 = get_id1(y_fid.s3, tmp_id); - - T4 ww1 = get_ww1(uv_dd.s0, uv_dd.s1); - T8 tmp_ww1 = get_ww2(y_dd.s0, ww1); - T8 tmp_ww2 = get_ww2(y_dd.s1, ww1); - T8 tmp_ww3 = get_ww2(y_dd.s2, ww1); - T8 tmp_ww4 = get_ww2(y_dd.s3, ww1); - - T8 y_lut1 = get_lut(lut, id1, 0); - T8 y_lut2 = get_lut(lut, id2, 0); - T8 y_lut3 = get_lut(lut, id3, 0); - T8 y_lut4 = get_lut(lut, id4, 0); - T8 u_lut = get_lut(lut, id1, shift); - T8 v_lut = get_lut(lut, id1, shift * 2); - - float4 tmp_y; - float2 tmp_uv; - tmp_y.s0 = dot(tmp_ww1.s0123, y_lut1.s0123) + dot(tmp_ww1.s4567, y_lut1.s4567); - tmp_y.s1 = dot(tmp_ww2.s0123, y_lut2.s0123) + dot(tmp_ww2.s4567, y_lut2.s4567); - tmp_y.s2 = dot(tmp_ww3.s0123, y_lut3.s0123) + dot(tmp_ww3.s4567, y_lut3.s4567); - tmp_y.s3 = dot(tmp_ww4.s0123, y_lut4.s0123) + dot(tmp_ww4.s4567, y_lut4.s4567); - - tmp_uv.s0 = dot(tmp_ww1.s0123, u_lut.s0123) + dot(tmp_ww1.s4567, u_lut.s4567); - tmp_uv.s1 = dot(tmp_ww1.s0123, v_lut.s0123) + dot(tmp_ww1.s4567, v_lut.s4567); - - tmp_y *= 255.0f; - tmp_uv *= 255.0f; - - vstore2((uchar2)(tmp_y.s0, tmp_y.s1), 0, ydst); - vstore2((uchar2)(tmp_y.s2, tmp_y.s3), 0, ydst + src_step); - vstore2((uchar2)(tmp_uv.s0, tmp_uv.s1), 0, udst); -} From a19a21064341741d6051850388689fae1f0dc27f Mon Sep 17 00:00:00 2001 From: Xianzhi Yu Date: Thu, 25 May 2023 17:38:10 +0800 Subject: [PATCH 3/4] delete unused file delete unused file --- .../gpu/mali/cl/bilateral_slice_apply_c12.cl | 77 +------------------ 1 file changed, 3 insertions(+), 74 deletions(-) diff --git a/compute/tensor/src/gpu/mali/cl/bilateral_slice_apply_c12.cl b/compute/tensor/src/gpu/mali/cl/bilateral_slice_apply_c12.cl index 43e43ce3..709f10ff 100644 --- a/compute/tensor/src/gpu/mali/cl/bilateral_slice_apply_c12.cl +++ b/compute/tensor/src/gpu/mali/cl/bilateral_slice_apply_c12.cl @@ -16,80 +16,9 @@ /*these parameters are belong to matrix mult/add and conv*/ /*they are extract from HDR model*/ /*they may be changful for different model*/ -inline T guide_cal0(T3 v) +inline T guide_cal(T3 v) { - T3 tmp; - tmp.x = v.x * (T)0.900616 - v.y * (T)0.1006 - v.z * (T)0.058384 + (T)0.072721; - tmp.y = -v.x * (T)0.079311 + v.y * (T)0.91976 - v.z * (T)0.037624 + (T)0.124359; - tmp.z = -v.x * (T)0.068347 - v.y * (T)0.069032 + v.z * (T)0.975032 + (T)0.129721; - tmp = max(tmp, (T)0); - tmp.x = tmp.x * (T)0.003211 * 16; - tmp.y = tmp.y * (T)0.007948 * 16; - tmp.z = tmp.z * (T)0.046259 * 16; - T g = tmp.x * (T)0.249512 + tmp.y * (T)0.274577 + tmp.z * (T)0.324276 + (T)0.078941; - return g; -} - -inline T guide_cal1(T3 v) -{ - T4 a = {v.x, v.y, v.z, 1}; - - T4 wx = {0.9266905188560486, -0.07651382684707642, -0.11796596646308899, 0.03732128441333771}; - T4 wy = {0.016965966671705246, 1.0332931280136108, 0.09558156877756119, 0.049296945333480835}; - T4 wz = {-0.060142070055007935, -0.0184615608304739, 0.9641872048377991, 0.03588166460394859}; - - T x = dot(a, wx); - T y = dot(a, wy); - T z = dot(a, wz); - - T16 sx = {-0.04031608998775482, 0.203898087143898, 0.21509018540382385, 0.2156994342803955, - 0.22189579904079437, 0.2710961699485779, 0.33060845732688904, 0.3510134816169739, - 0.3799624741077423, 0.4165642559528351, 0.5429311394691467, 0.6519719958305359, - 0.7579551339149475, 0.8117461800575256, 0.8115477561950684, 0.811525821685791}; - - T16 sy = {-0.04493796080350876, 0.2501078248023987, 0.24961410462856293, 0.24829524755477905, - 0.25029096007347107, 0.25275537371635437, 0.2535839378833771, 0.25915712118148804, - 0.992545485496521, 0.869307279586792, 0.8143411874771118, 0.8268355131149292, - 0.849763810634613, 0.8641695380210876, 0.8749480843544006, 0.9124495387077332}; - - T16 sz = {-0.0450710691511631, 0.17914339900016785, 0.20727036893367767, 0.21128158271312714, - 0.785589873790741, 0.40014126896858215, 0.39716723561286926, 0.4003089666366577, - 0.5749346613883972, 0.6277766227722168, 0.7884474992752075, 0.788446307182312, - 0.789533257484436, 0.7905913591384888, 0.7964500188827515, 0.7964839339256287}; - - T16 rx = max(x - sx, (T)0); - T16 ry = max(x - sy, (T)0); - T16 rz = max(z - sz, (T)0); - - T16 mx = {0.9483454823493958, -0.02504969760775566, -0.0731356292963028, -0.08960649371147156, - -0.0989985391497612, -0.0911787822842598, -0.07849951088428497, -0.07431424409151077, - -0.05982533469796181, -0.027073463425040245, 0.09377846121788025, 0.07562971860170364, - -0.05076618492603302, 0.2615104913711548, 0.42631882429122925, 0.6887183785438538}; - - T16 my = {0.9732255339622498, -0.03841959312558174, -0.07476486265659332, -0.08849595487117767, - -0.10008298605680466, -0.10915014147758484, -0.1108635663986206, -0.09364574402570724, - -0.04355158284306526, -0.015994733199477196, -0.025348246097564697, -0.051913388073444366, - -0.07183714956045151, -0.0823502317070961, -0.09460879862308502, -0.13453315198421478}; - - T16 mz = {0.951180636882782, -0.014929438941180706, -0.022745108231902122, -0.042111292481422424, - 0.061638616025447845, -0.04308458790183067, -0.050973013043403625, -0.045611534267663956, - 0.037990815937519073, 0.04962018504738808, 0.15617141127586365, 0.13662904500961304, - 0.16109246015548706, 0.160025492310524, 0.12079561501741409, 0.15001150965690613}; - - x = dot(rx.s0123, mx.s0123) + dot(rx.s4567, mx.s4567) + dot(rx.s89ab, mx.s89ab) + - dot(rx.scdef, mx.scdef); - y = dot(ry.s0123, my.s0123) + dot(ry.s4567, my.s4567) + dot(ry.s89ab, my.s89ab) + - dot(ry.scdef, my.scdef); - z = dot(rz.s0123, mz.s0123) + dot(rz.s4567, mz.s4567) + dot(rz.s89ab, mz.s89ab) + - dot(rz.scdef, mz.scdef); - - T4 t1 = {x, y, z, 1}; - T4 w = {0.28540247678756714, 0.31782254576683044, 0.28381019830703735, 0.06326253712177277}; - T g = dot(t1, w); - - g = min(max((T)0., g), (T)1.); - - return g; + return v.x; } __kernel void KERNEL_NAME @@ -135,7 +64,7 @@ __kernel void KERNEL_NAME T gx = (x + (T)0.5) * (T)scale_x; #if defined(CONV) - T gz = guide_cal1(in_val); + T gz = guide_cal(in_val); #else T gz = guide[in_off]; #endif From 0bc7832a834dd18ae3e699e9113410e7f9f84492 Mon Sep 17 00:00:00 2001 From: Xianzhi Yu Date: Fri, 9 Jun 2023 17:25:48 +0800 Subject: [PATCH 4/4] Update CHANGELOG.md --- docs/CHANGELOG.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/docs/CHANGELOG.md b/docs/CHANGELOG.md index 3ab47fd9..fd37cb11 100644 --- a/docs/CHANGELOG.md +++ b/docs/CHANGELOG.md @@ -16,7 +16,7 @@ and this project adheres to [Semantic Versioning](