diff options
author | Alexey Titov <alex.ti.771@gmail.com> | 2015-02-10 23:21:05 +0300 |
---|---|---|
committer | Michael Niedermayer <michaelni@gmx.at> | 2015-02-11 15:51:58 +0300 |
commit | a05a7373166c3c1b2eed3db66e143e9d227e0354 (patch) | |
tree | 6299b33a7c50ca1818ab145c4ac284d97fb0de49 /libavfilter/unsharp_opencl_kernel.h | |
parent | ac494e5a66507c8124e8ba399b3f0685a51d7a82 (diff) |
avfilter/unsharp: OpenCL unsharpen filter optimization: substitute N^2 filter computation with 2N+C
i7-4770K luma 21% faster, chroma 18% faster A10-7850K luma 42% faster, chroma 37% faster on 1920x1080 res
Signed-off-by: Michael Niedermayer <michaelni@gmx.at>
Diffstat (limited to 'libavfilter/unsharp_opencl_kernel.h')
-rw-r--r-- | libavfilter/unsharp_opencl_kernel.h | 122 |
1 files changed, 89 insertions, 33 deletions
diff --git a/libavfilter/unsharp_opencl_kernel.h b/libavfilter/unsharp_opencl_kernel.h index 9c4fd65031..307d0f1814 100644 --- a/libavfilter/unsharp_opencl_kernel.h +++ b/libavfilter/unsharp_opencl_kernel.h @@ -36,7 +36,8 @@ inline unsigned char clip_uint8(int a) kernel void unsharp_luma( global unsigned char *src, global unsigned char *dst, - global int *mask, + global int *mask_x, + global int *mask_y, int amount, int scalebits, int halfscale, @@ -59,10 +60,12 @@ kernel void unsharp_luma( return; } - local uchar l[32][32]; - local int lc[LU_RADIUS_X*LU_RADIUS_Y]; + local unsigned int l[32][32]; + local unsigned int lcx[LU_RADIUS_X]; + local unsigned int lcy[LU_RADIUS_Y]; int indexIx, indexIy, i, j; + //load up tile: actual workspace + halo of 8 points in x and y \n for(i = 0; i <= 1; i++) { indexIy = -8 + (blockIdx.y + i) * 16 + threadIdx.y; indexIy = indexIy < 0 ? 0 : indexIy; @@ -76,27 +79,54 @@ kernel void unsharp_luma( } int indexL = threadIdx.y*16 + threadIdx.x; - if (indexL < LU_RADIUS_X*LU_RADIUS_Y) - lc[indexL] = mask[indexL]; + if (indexL < LU_RADIUS_X) + lcx[indexL] = mask_x[indexL]; + if (indexL < LU_RADIUS_Y) + lcy[indexL] = mask_y[indexL]; barrier(CLK_LOCAL_MEM_FENCE); + //needed for unsharp mask application in the end \n + int orig_value = (int)l[threadIdx.y + 8][threadIdx.x + 8]; + int idx, idy, maskIndex; - int sum = 0; - int steps_x = LU_RADIUS_X/2; - int steps_y = LU_RADIUS_Y/2; + int temp[2] = {0}; + int steps_x = (LU_RADIUS_X-1)/2; + int steps_y = (LU_RADIUS_Y-1)/2; - \n#pragma unroll\n - for (i = -steps_y; i <= steps_y; i++) { - idy = 8 + i + threadIdx.y; - \n#pragma unroll\n - for (j = -steps_x; j <= steps_x; j++) { - idx = 8 + j + threadIdx.x; - maskIndex = (i + steps_y)*LU_RADIUS_X + j + steps_x; - sum += (int)l[idy][idx] * lc[maskIndex]; + // compute the actual workspace + left&right halos \n + \n#pragma unroll\n + for (j = 0; j <=1; j++) { + //extra work to cover left and right halos \n + idx = 16*j + threadIdx.x; + \n#pragma unroll\n + for (i = -steps_y; i <= steps_y; i++) { + idy = 8 + i + threadIdx.y; + maskIndex = (i + steps_y); + temp[j] += (int)l[idy][idx] * lcy[maskIndex]; } } - int temp = (int)l[threadIdx.y + 8][threadIdx.x + 8]; - int res = temp + (((temp - (int)((sum + halfscale) >> scalebits)) * amount) >> 16); + barrier(CLK_LOCAL_MEM_FENCE); + //save results from the vertical filter in local memory \n + idy = 8 + threadIdx.y; + \n#pragma unroll\n + for (j = 0; j <=1; j++) { + idx = 16*j + threadIdx.x; + l[idy][idx] = temp[j]; + } + barrier(CLK_LOCAL_MEM_FENCE); + + //compute results with the horizontal filter \n + int sum = 0; + idy = 8 + threadIdx.y; + \n#pragma unroll\n + for (j = -steps_x; j <= steps_x; j++) { + idx = 8 + j + threadIdx.x; + maskIndex = j + steps_x; + sum += (int)l[idy][idx] * lcx[maskIndex]; + } + + int res = orig_value + (((orig_value - (int)((sum + halfscale) >> scalebits)) * amount) >> 16); + if (globalIdx.x < width && globalIdx.y < height) dst[globalIdx.x + globalIdx.y*dst_stride] = clip_uint8(res); } @@ -104,7 +134,8 @@ kernel void unsharp_luma( kernel void unsharp_chroma( global unsigned char *src_y, global unsigned char *dst_y, - global int *mask, + global int *mask_x, + global int *mask_y, int amount, int scalebits, int halfscale, @@ -141,8 +172,9 @@ kernel void unsharp_chroma( return; } - local uchar l[32][32]; - local int lc[CH_RADIUS_X*CH_RADIUS_Y]; + local unsigned int l[32][32]; + local unsigned int lcx[CH_RADIUS_X]; + local unsigned int lcy[CH_RADIUS_Y]; int indexIx, indexIy, i, j; for(i = 0; i <= 1; i++) { indexIy = -8 + (blockIdx.y + i) * 16 + threadIdx.y; @@ -157,27 +189,51 @@ kernel void unsharp_chroma( } int indexL = threadIdx.y*16 + threadIdx.x; - if (indexL < CH_RADIUS_X*CH_RADIUS_Y) - lc[indexL] = mask[indexL]; + if (indexL < CH_RADIUS_X) + lcx[indexL] = mask_x[indexL]; + if (indexL < CH_RADIUS_Y) + lcy[indexL] = mask_y[indexL]; barrier(CLK_LOCAL_MEM_FENCE); + int orig_value = (int)l[threadIdx.y + 8][threadIdx.x + 8]; + int idx, idy, maskIndex; - int sum = 0; int steps_x = CH_RADIUS_X/2; int steps_y = CH_RADIUS_Y/2; + int temp[2] = {0,0}; \n#pragma unroll\n - for (i = -steps_y; i <= steps_y; i++) { - idy = 8 + i + threadIdx.y; + for (j = 0; j <= 1; j++) { + idx = 16*j + threadIdx.x; \n#pragma unroll\n - for (j = -steps_x; j <= steps_x; j++) { - idx = 8 + j + threadIdx.x; - maskIndex = (i + steps_y)*CH_RADIUS_X + j + steps_x; - sum += (int)l[idy][idx] * lc[maskIndex]; - } + for (i = -steps_y; i <= steps_y; i++) { + idy = 8 + i + threadIdx.y; + maskIndex = i + steps_y; + temp[j] += (int)l[idy][idx] * lcy[maskIndex]; + } + } + + barrier(CLK_LOCAL_MEM_FENCE); + idy = 8 + threadIdx.y; + \n#pragma unroll\n + for (j = 0; j <= 1; j++) { + idx = 16*j + threadIdx.x; + l[idy][idx] = temp[j]; } - int temp = (int)l[threadIdx.y + 8][threadIdx.x + 8]; - int res = temp + (((temp - (int)((sum + halfscale) >> scalebits)) * amount) >> 16); + barrier(CLK_LOCAL_MEM_FENCE); + + //compute results with the horizontal filter \n + int sum = 0; + idy = 8 + threadIdx.y; + \n#pragma unroll\n + for (j = -steps_x; j <= steps_x; j++) { + idx = 8 + j + threadIdx.x; + maskIndex = j + steps_x; + sum += (int)l[idy][idx] * lcx[maskIndex]; + } + + int res = orig_value + (((orig_value - (int)((sum + halfscale) >> scalebits)) * amount) >> 16); + if (globalIdx.x < cw && globalIdx.y < ch) dst[globalIdx.x + globalIdx.y*dst_stride_ch] = clip_uint8(res); } |