[FFmpeg-devel] [PATCH]avfilter/unsharp_opencl

[Michael Niedermayer]

On Wed, Nov 06, 2013 at 01:01:17AM -0600, Lenny Wang wrote:
> Add optimized opencl kernels with greatly improved overall performance
> observed on various mainstream platforms.

[...]
> @@ -225,16 +297,36 @@ int ff_opencl_unsharp_init(AVFilterContext *ctx)
>          av_log(ctx, AV_LOG_ERROR, "Unable to get OpenCL command queue in filter 'unsharp'\n");
>          return AVERROR(EINVAL);
>      }
> -    unsharp->opencl_ctx.program = av_opencl_compile("unsharp", NULL);
> +    sprintf(build_opts, "-D LU_RADIUS_X=%d -D LU_RADIUS_Y=%d -D CH_RADIUS_X=%d -D CH_RADIUS_Y=%d",
> +            2*unsharp->luma.steps_x+1, 2*unsharp->luma.steps_y+1, 2*unsharp->chroma.steps_x+1, 2*unsharp->chroma.steps_y+1);

this should use snprintf(), for saftey



[...]
> @@ -32,7 +33,156 @@ inline unsigned char clip_uint8(int a)
>          return a;
>  }
>  
> -kernel void unsharp(global  unsigned char *src,
> +kernel void unsharp_luma(
> +                    global unsigned char *src,
> +                    global unsigned char *dst,
> +                    global int *mask,
> +                    int amount,
> +                    int scalebits,
> +                    int halfscale,
> +                    int src_stride,
> +                    int dst_stride,
> +                    int width,
> +                    int height)
> +{
> +    int2 threadIdx, blockIdx, globalIdx;
> +    threadIdx.x = get_local_id(0);
> +    threadIdx.y = get_local_id(1);
> +    blockIdx.x = get_group_id(0);
> +    blockIdx.y = get_group_id(1);
> +    globalIdx.x = get_global_id(0);
> +    globalIdx.y = get_global_id(1);
> +
> +    if (!amount) {
> +        if (globalIdx.x < width && globalIdx.y < height)
> +            dst[globalIdx.x + globalIdx.y*dst_stride] = src[globalIdx.x + globalIdx.y*src_stride];
> +        return;
> +    }
> +
> +    local uchar l[32][32];
> +    local int lc[LU_RADIUS_X*LU_RADIUS_Y];
> +    int indexIx, indexIy;
> +
> +    for(int i = 0; i <= 1; i++) {

the variable should be declared outside the for( to ensure maximal
compiler compatibility


[...]
-- 
Michael     GnuPG fingerprint: 9FF2128B147EF6730BADF133611EC787040B0FAB

If you think the mosad wants you dead since a long time then you are either
wrong or dead since a long time.
-------------- next part --------------
A non-text attachment was scrubbed...
Name: not available
Type: application/pgp-signature
Size: 198 bytes
Desc: Digital signature
URL: <http://ffmpeg.org/pipermail/ffmpeg-devel/attachments/20131107/261e5fc2/attachment.asc>