[FFmpeg-devel] [PATCH 1/2] libavfilter/unsharp opencl optimization
Wei Gao
highgod0401 at gmail.com
Mon Jan 5 01:56:12 CET 2015
2015-01-04 7:34 GMT+08:00 Alexey Titov <alex.ti.771 at gmail.com>:
> From: atitov <alexey.titov at amd.com>
>
> ---
> libavfilter/unsharp.h | 4 ++
> libavfilter/unsharp_opencl.c | 74 +++++++++++++++-------
> libavfilter/unsharp_opencl_kernel.h | 122
> ++++++++++++++++++++++++++----------
> 3 files changed, 145 insertions(+), 55 deletions(-)
>
> diff --git a/libavfilter/unsharp.h b/libavfilter/unsharp.h
> index c2aed64..fc651c0 100644
> --- a/libavfilter/unsharp.h
> +++ b/libavfilter/unsharp.h
> @@ -41,6 +41,10 @@ typedef struct {
> cl_kernel kernel_chroma;
> cl_mem cl_luma_mask;
> cl_mem cl_chroma_mask;
> + cl_mem cl_luma_mask_x;
> + cl_mem cl_chroma_mask_x;
> + cl_mem cl_luma_mask_y;
> + cl_mem cl_chroma_mask_y;
> int in_plane_size[8];
> int out_plane_size[8];
> int plane_num;
> diff --git a/libavfilter/unsharp_opencl.c b/libavfilter/unsharp_opencl.c
> index 5c6b5ef..1923cb3 100644
> --- a/libavfilter/unsharp_opencl.c
> +++ b/libavfilter/unsharp_opencl.c
> @@ -87,11 +87,12 @@ end:
> return ret;
> }
>
> -static int compute_mask_matrix(cl_mem cl_mask_matrix, int step_x, int
> step_y)
> +static int copy_separable_masks(cl_mem cl_mask_x, cl_mem cl_mask_y, int
> step_x, int step_y)
> {
> - int i, j, ret = 0;
> - uint32_t *mask_matrix, *mask_x, *mask_y;
> - size_t size_matrix = sizeof(uint32_t) * (2 * step_x + 1) * (2 *
> step_y + 1);
> + int ret = 0;
> + uint32_t *mask_x, *mask_y;
> + size_t size_mask_x = sizeof(uint32_t) * (2 * step_x + 1);
> + size_t size_mask_y = sizeof(uint32_t) * (2 * step_y + 1);
> mask_x = av_mallocz_array(2 * step_x + 1, sizeof(uint32_t));
> if (!mask_x) {
> ret = AVERROR(ENOMEM);
> @@ -102,37 +103,33 @@ static int compute_mask_matrix(cl_mem
> cl_mask_matrix, int step_x, int step_y)
> ret = AVERROR(ENOMEM);
> goto end;
> }
> - mask_matrix = av_mallocz(size_matrix);
> - if (!mask_matrix) {
> - ret = AVERROR(ENOMEM);
> - goto end;
> - }
> ret = compute_mask(step_x, mask_x);
> if (ret < 0)
> goto end;
> ret = compute_mask(step_y, mask_y);
> if (ret < 0)
> goto end;
> - for (j = 0; j < 2 * step_y + 1; j++) {
> - for (i = 0; i < 2 * step_x + 1; i++) {
> - mask_matrix[i + j * (2 * step_x + 1)] = mask_y[j] * mask_x[i];
> - }
> - }
> - ret = av_opencl_buffer_write(cl_mask_matrix, (uint8_t *)mask_matrix,
> size_matrix);
> + ret = av_opencl_buffer_write(cl_mask_x, (uint8_t *)mask_x,
> size_mask_x);
> + ret = av_opencl_buffer_write(cl_mask_y, (uint8_t *)mask_y,
> size_mask_y);
> end:
> av_freep(&mask_x);
> av_freep(&mask_y);
> - av_freep(&mask_matrix);
> return ret;
> }
>
> static int generate_mask(AVFilterContext *ctx)
> {
> - UnsharpContext *unsharp = ctx->priv;
> - int i, ret = 0, step_x[2], step_y[2];
> + cl_mem masks[4];
> cl_mem mask_matrix[2];
> + int i, ret = 0, step_x[2], step_y[2];
> +
> + UnsharpContext *unsharp = ctx->priv;
> mask_matrix[0] = unsharp->opencl_ctx.cl_luma_mask;
> mask_matrix[1] = unsharp->opencl_ctx.cl_chroma_mask;
> + masks[0] = unsharp->opencl_ctx.cl_luma_mask_x;
> + masks[1] = unsharp->opencl_ctx.cl_luma_mask_y;
> + masks[2] = unsharp->opencl_ctx.cl_chroma_mask_x;
> + masks[3] = unsharp->opencl_ctx.cl_chroma_mask_y;
> step_x[0] = unsharp->luma.steps_x;
> step_x[1] = unsharp->chroma.steps_x;
> step_y[0] = unsharp->luma.steps_y;
> @@ -144,12 +141,16 @@ static int generate_mask(AVFilterContext *ctx)
> else
> unsharp->opencl_ctx.use_fast_kernels = 1;
>
> + if (!masks[0] || !masks[1] || !masks[2] || !masks[3]) {
> + av_log(ctx, AV_LOG_ERROR, "Luma mask and chroma mask should not
> be NULL\n");
> + return AVERROR(EINVAL);
> + }
> if (!mask_matrix[0] || !mask_matrix[1]) {
> av_log(ctx, AV_LOG_ERROR, "Luma mask and chroma mask should not
> be NULL\n");
> return AVERROR(EINVAL);
> }
> for (i = 0; i < 2; i++) {
> - ret = compute_mask_matrix(mask_matrix[i], step_x[i], step_y[i]);
> + ret = copy_separable_masks(masks[2*i], masks[2*i+1], step_x[i],
> step_y[i]);
> if (ret < 0)
> return ret;
> }
> @@ -184,7 +185,8 @@ int ff_opencl_apply_unsharp(AVFilterContext *ctx,
> AVFrame *in, AVFrame *out)
> ret = avpriv_opencl_set_parameter(&kernel1,
>
> FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_inbuf),
>
> FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_outbuf),
> -
> FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_luma_mask),
> +
> FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_luma_mask_x),
> +
> FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_luma_mask_y),
>
> FF_OPENCL_PARAM_INFO(unsharp->luma.amount),
>
> FF_OPENCL_PARAM_INFO(unsharp->luma.scalebits),
>
> FF_OPENCL_PARAM_INFO(unsharp->luma.halfscale),
> @@ -201,7 +203,8 @@ int ff_opencl_apply_unsharp(AVFilterContext *ctx,
> AVFrame *in, AVFrame *out)
> ret = avpriv_opencl_set_parameter(&kernel2,
>
> FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_inbuf),
>
> FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_outbuf),
> -
> FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_chroma_mask),
> +
> FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_chroma_mask_x),
> +
> FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_chroma_mask_y),
>
> FF_OPENCL_PARAM_INFO(unsharp->chroma.amount),
>
> FF_OPENCL_PARAM_INFO(unsharp->chroma.scalebits),
>
> FF_OPENCL_PARAM_INFO(unsharp->chroma.halfscale),
> @@ -264,7 +267,9 @@ int ff_opencl_apply_unsharp(AVFilterContext *ctx,
> AVFrame *in, AVFrame *out)
> return AVERROR_EXTERNAL;
> }
> }
> - clFinish(unsharp->opencl_ctx.command_queue);
> + //blocking map is suffficient, no need for clFinish
> + //clFinish(unsharp->opencl_ctx.command_queue);
> +
> return av_opencl_buffer_read_image(out->data,
> unsharp->opencl_ctx.out_plane_size,
> unsharp->opencl_ctx.plane_num,
> unsharp->opencl_ctx.cl_outbuf,
>
> unsharp->opencl_ctx.cl_outbuf_size);
> @@ -286,6 +291,27 @@ int ff_opencl_unsharp_init(AVFilterContext *ctx)
> ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_chroma_mask,
> sizeof(uint32_t) * (2 *
> unsharp->chroma.steps_x + 1) * (2 * unsharp->chroma.steps_y + 1),
> CL_MEM_READ_ONLY, NULL);
> + // separable filters
> + if (ret < 0)
> + return ret;
> + ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_luma_mask_x,
> + sizeof(uint32_t) * (2 *
> unsharp->luma.steps_x + 1),
> + CL_MEM_READ_ONLY, NULL);
> + if (ret < 0)
> + return ret;
> + ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_luma_mask_y,
> + sizeof(uint32_t) * (2 *
> unsharp->luma.steps_y + 1),
> + CL_MEM_READ_ONLY, NULL);
> + if (ret < 0)
> + return ret;
> + ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_chroma_mask_x,
> + sizeof(uint32_t) * (2 *
> unsharp->chroma.steps_x + 1),
> + CL_MEM_READ_ONLY, NULL);
> + if (ret < 0)
> + return ret;
> + ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_chroma_mask_y,
> + sizeof(uint32_t) * (2 *
> unsharp->chroma.steps_y + 1),
> + CL_MEM_READ_ONLY, NULL);
> if (ret < 0)
> return ret;
> ret = generate_mask(ctx);
> @@ -339,6 +365,10 @@ void ff_opencl_unsharp_uninit(AVFilterContext *ctx)
> av_opencl_buffer_release(&unsharp->opencl_ctx.cl_outbuf);
> av_opencl_buffer_release(&unsharp->opencl_ctx.cl_luma_mask);
> av_opencl_buffer_release(&unsharp->opencl_ctx.cl_chroma_mask);
> + av_opencl_buffer_release(&unsharp->opencl_ctx.cl_luma_mask_x);
> + av_opencl_buffer_release(&unsharp->opencl_ctx.cl_chroma_mask_x);
> + av_opencl_buffer_release(&unsharp->opencl_ctx.cl_luma_mask_y);
> + av_opencl_buffer_release(&unsharp->opencl_ctx.cl_chroma_mask_y);
> clReleaseKernel(unsharp->opencl_ctx.kernel_default);
> clReleaseKernel(unsharp->opencl_ctx.kernel_luma);
> clReleaseKernel(unsharp->opencl_ctx.kernel_chroma);
> diff --git a/libavfilter/unsharp_opencl_kernel.h
> b/libavfilter/unsharp_opencl_kernel.h
> index 9c4fd65..0fc802e 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 short l[32][32];
> + local int lcx[LU_RADIUS_X];
> + local 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 ushort l[32][32];
> + local int lcx[CH_RADIUS_X];
> + local 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);
> }
> --
> 1.8.4.msysgit.0
>
> Hi
This is the first part of the whole patch right? Could you send the second
part?
Thanks
Best regads
> _______________________________________________
> ffmpeg-devel mailing list
> ffmpeg-devel at ffmpeg.org
> http://ffmpeg.org/mailman/listinfo/ffmpeg-devel
>
More information about the ffmpeg-devel
mailing list