32 #define ROUND_TO_16(a) (((((a) - 1)/16)+1)*16)
37 for (i = 0; i <
len; i++) {
38 dst[i] = counter1[i] + counter2[i];
45 int counter_size =
sizeof(uint32_t) * (2 * step + 1);
46 uint32_t *temp1_counter, *temp2_counter, **counter;
62 for (i = 0; i < 2 * step + 1; i++) {
69 for (i = 0; i < 2 * step + 1; i++) {
70 memset(temp1_counter, 0, counter_size);
72 for (z = 0; z < step * 2; z += 2) {
74 memcpy(counter[z], temp1_counter, counter_size);
76 memcpy(counter[z + 1], temp2_counter, counter_size);
79 memcpy(mask, temp1_counter, counter_size);
83 for (i = 0; i < 2 * step + 1; i++) {
93 uint32_t *mask_x, *mask_y;
94 size_t size_mask_x =
sizeof(uint32_t) * (2 * step_x + 1);
95 size_t size_mask_y =
sizeof(uint32_t) * (2 * step_y + 1);
127 int i, ret = 0, step_x[2], step_y[2];
130 mask_matrix[0] = unsharp->opencl_ctx.cl_luma_mask;
131 mask_matrix[1] = unsharp->opencl_ctx.cl_chroma_mask;
132 masks[0] = unsharp->opencl_ctx.cl_luma_mask_x;
133 masks[1] = unsharp->opencl_ctx.cl_luma_mask_y;
134 masks[2] = unsharp->opencl_ctx.cl_chroma_mask_x;
135 masks[3] = unsharp->opencl_ctx.cl_chroma_mask_y;
142 if (step_x[0]>8 || step_x[1]>8 || step_y[0]>8 || step_y[1]>8)
143 unsharp->opencl_ctx.use_fast_kernels = 0;
145 unsharp->opencl_ctx.use_fast_kernels = 1;
147 if (!masks[0] || !masks[1] || !masks[2] || !masks[3]) {
151 if (!mask_matrix[0] || !mask_matrix[1]) {
155 for (i = 0; i < 2; i++) {
175 size_t globalWorkSize1d = width * height + 2 * ch * cw;
176 size_t globalWorkSize2dLuma[2];
177 size_t globalWorkSize2dChroma[2];
178 size_t localWorkSize2d[2] = {16, 16};
180 if (unsharp->opencl_ctx.use_fast_kernels) {
181 globalWorkSize2dLuma[0] = (size_t)
ROUND_TO_16(width);
182 globalWorkSize2dLuma[1] = (size_t)
ROUND_TO_16(height);
183 globalWorkSize2dChroma[0] = (size_t)
ROUND_TO_16(cw);
184 globalWorkSize2dChroma[1] = (size_t)(2*
ROUND_TO_16(ch));
187 kernel1.
kernel = unsharp->opencl_ctx.kernel_luma;
205 kernel2.
kernel = unsharp->opencl_ctx.kernel_chroma;
225 status = clEnqueueNDRangeKernel(unsharp->opencl_ctx.command_queue,
226 unsharp->opencl_ctx.kernel_luma, 2,
NULL,
227 globalWorkSize2dLuma, localWorkSize2d, 0,
NULL,
NULL);
228 status |=clEnqueueNDRangeKernel(unsharp->opencl_ctx.command_queue,
229 unsharp->opencl_ctx.kernel_chroma, 2,
NULL,
230 globalWorkSize2dChroma, localWorkSize2d, 0,
NULL,
NULL);
231 if (status != CL_SUCCESS) {
237 kernel1.
kernel = unsharp->opencl_ctx.kernel_default;
265 status = clEnqueueNDRangeKernel(unsharp->opencl_ctx.command_queue,
266 unsharp->opencl_ctx.kernel_default, 1,
NULL,
268 if (status != CL_SUCCESS) {
277 unsharp->opencl_ctx.plane_num, unsharp->opencl_ctx.cl_outbuf,
278 unsharp->opencl_ctx.cl_outbuf_size);
291 CL_MEM_READ_ONLY,
NULL);
296 CL_MEM_READ_ONLY,
NULL);
301 sizeof(uint32_t) * (2 * unsharp->
luma.
steps_x + 1),
302 CL_MEM_READ_ONLY,
NULL);
306 sizeof(uint32_t) * (2 * unsharp->
luma.
steps_y + 1),
307 CL_MEM_READ_ONLY,
NULL);
312 CL_MEM_READ_ONLY,
NULL);
317 CL_MEM_READ_ONLY,
NULL);
323 unsharp->opencl_ctx.plane_num =
PLANE_NUM;
325 if (!unsharp->opencl_ctx.command_queue) {
326 av_log(ctx,
AV_LOG_ERROR,
"Unable to get OpenCL command queue in filter 'unsharp'\n");
329 snprintf(build_opts, 96,
"-D LU_RADIUS_X=%d -D LU_RADIUS_Y=%d -D CH_RADIUS_X=%d -D CH_RADIUS_Y=%d",
332 if (!unsharp->opencl_ctx.program) {
336 if (unsharp->opencl_ctx.use_fast_kernels) {
337 if (!unsharp->opencl_ctx.kernel_luma) {
338 unsharp->opencl_ctx.kernel_luma = clCreateKernel(unsharp->opencl_ctx.program,
"unsharp_luma", &ret);
339 if (ret != CL_SUCCESS) {
344 if (!unsharp->opencl_ctx.kernel_chroma) {
345 unsharp->opencl_ctx.kernel_chroma = clCreateKernel(unsharp->opencl_ctx.program,
"unsharp_chroma", &ret);
353 if (!unsharp->opencl_ctx.kernel_default) {
354 unsharp->opencl_ctx.kernel_default = clCreateKernel(unsharp->opencl_ctx.program,
"unsharp_default", &ret);
375 clReleaseKernel(unsharp->opencl_ctx.kernel_default);
376 clReleaseKernel(unsharp->opencl_ctx.kernel_luma);
377 clReleaseKernel(unsharp->opencl_ctx.kernel_chroma);
378 clReleaseProgram(unsharp->opencl_ctx.program);
379 unsharp->opencl_ctx.command_queue =
NULL;
390 if ((!unsharp->opencl_ctx.cl_inbuf) || (!unsharp->opencl_ctx.cl_outbuf)) {
391 unsharp->opencl_ctx.in_plane_size[0] = (in->
linesize[0] * in->
height);
392 unsharp->opencl_ctx.in_plane_size[1] = (in->
linesize[1] * ch);
393 unsharp->opencl_ctx.in_plane_size[2] = (in->
linesize[2] * ch);
394 unsharp->opencl_ctx.out_plane_size[0] = (out->
linesize[0] * out->
height);
395 unsharp->opencl_ctx.out_plane_size[1] = (out->
linesize[1] * ch);
396 unsharp->opencl_ctx.out_plane_size[2] = (out->
linesize[2] * ch);
397 unsharp->opencl_ctx.cl_inbuf_size = unsharp->opencl_ctx.in_plane_size[0] +
398 unsharp->opencl_ctx.in_plane_size[1] +
399 unsharp->opencl_ctx.in_plane_size[2];
400 unsharp->opencl_ctx.cl_outbuf_size = unsharp->opencl_ctx.out_plane_size[0] +
401 unsharp->opencl_ctx.out_plane_size[1] +
402 unsharp->opencl_ctx.out_plane_size[2];
403 if (!unsharp->opencl_ctx.cl_inbuf) {
405 unsharp->opencl_ctx.cl_inbuf_size,
406 CL_MEM_READ_ONLY,
NULL);
410 if (!unsharp->opencl_ctx.cl_outbuf) {
412 unsharp->opencl_ctx.cl_outbuf_size,
413 CL_MEM_READ_WRITE,
NULL);
419 unsharp->opencl_ctx.cl_inbuf_size,
420 0, in->
data, unsharp->opencl_ctx.in_plane_size,
421 unsharp->opencl_ctx.plane_num);
static const uint16_t mask_matrix[]
This structure describes decoded (raw) audio or video data.
int h
agreed upon image height
UnsharpFilterParam luma
luma parameters (width, height, amount)
static int copy_separable_masks(cl_mem cl_mask_x, cl_mem cl_mask_y, int step_x, int step_y)
const char * av_opencl_errstr(cl_int status)
Get OpenCL error string.
AVFilterLink ** inputs
array of pointers to input links
int steps_x
horizontal step count
void ff_opencl_unsharp_uninit(AVFilterContext *ctx)
static av_cold int end(AVCodecContext *avctx)
static void add_mask_counter(uint32_t *dst, uint32_t *counter1, uint32_t *counter2, int len)
int ff_opencl_unsharp_init(AVFilterContext *ctx)
A link between two filters.
int scalebits
bits to shift pixel
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
int32_t halfscale
amount to add to pixel
static const uint16_t mask[17]
void * priv
private data for use by the filter
int avpriv_opencl_set_parameter(FFOpenclParam *opencl_param,...)
int w
agreed upon image width
cl_program av_opencl_compile(const char *program_name, const char *build_opts)
compile specific OpenCL kernel source
int av_opencl_buffer_create(cl_mem *cl_buf, size_t cl_buf_size, int flags, void *host_ptr)
Create OpenCL buffer.
int ff_opencl_unsharp_process_inout_buf(AVFilterContext *ctx, AVFrame *in, AVFrame *out)
int av_opencl_buffer_write_image(cl_mem dst_cl_buf, size_t cl_buffer_size, int dst_cl_offset, uint8_t **src_data, int *plane_size, int plane_num)
Write image data from memory to OpenCL buffer.
static int compute_mask(int step, uint32_t *mask)
int linesize[AV_NUM_DATA_POINTERS]
For video, size in bytes of each picture line.
int steps_y
vertical step count
uint8_t pi<< 24) CONV_FUNC_GROUP(AV_SAMPLE_FMT_FLT, float, AV_SAMPLE_FMT_U8, uint8_t,(*(constuint8_t *) pi-0x80)*(1.0f/(1<< 7))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_DBL, double, AV_SAMPLE_FMT_U8, uint8_t,(*(constuint8_t *) pi-0x80)*(1.0/(1<< 7))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_U8, uint8_t, AV_SAMPLE_FMT_S16, int16_t,(*(constint16_t *) pi >>8)+0x80) CONV_FUNC_GROUP(AV_SAMPLE_FMT_FLT, float, AV_SAMPLE_FMT_S16, int16_t,*(constint16_t *) pi *(1.0f/(1<< 15))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_DBL, double, AV_SAMPLE_FMT_S16, int16_t,*(constint16_t *) pi *(1.0/(1<< 15))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_U8, uint8_t, AV_SAMPLE_FMT_S32, int32_t,(*(constint32_t *) pi >>24)+0x80) CONV_FUNC_GROUP(AV_SAMPLE_FMT_FLT, float, AV_SAMPLE_FMT_S32, int32_t,*(constint32_t *) pi *(1.0f/(1U<< 31))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_DBL, double, AV_SAMPLE_FMT_S32, int32_t,*(constint32_t *) pi *(1.0/(1U<< 31))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_U8, uint8_t, AV_SAMPLE_FMT_FLT, float, av_clip_uint8(lrintf(*(constfloat *) pi *(1<< 7))+0x80)) CONV_FUNC_GROUP(AV_SAMPLE_FMT_S16, int16_t, AV_SAMPLE_FMT_FLT, float, av_clip_int16(lrintf(*(constfloat *) pi *(1<< 15)))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_S32, int32_t, AV_SAMPLE_FMT_FLT, float, av_clipl_int32(llrintf(*(constfloat *) pi *(1U<< 31)))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_U8, uint8_t, AV_SAMPLE_FMT_DBL, double, av_clip_uint8(lrint(*(constdouble *) pi *(1<< 7))+0x80)) CONV_FUNC_GROUP(AV_SAMPLE_FMT_S16, int16_t, AV_SAMPLE_FMT_DBL, double, av_clip_int16(lrint(*(constdouble *) pi *(1<< 15)))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_S32, int32_t, AV_SAMPLE_FMT_DBL, double, av_clipl_int32(llrint(*(constdouble *) pi *(1U<< 31))))#defineSET_CONV_FUNC_GROUP(ofmt, ifmt) staticvoidset_generic_function(AudioConvert *ac){}voidff_audio_convert_free(AudioConvert **ac){if(!*ac) return;ff_dither_free(&(*ac) ->dc);av_freep(ac);}AudioConvert *ff_audio_convert_alloc(AVAudioResampleContext *avr, enumAVSampleFormatout_fmt, enumAVSampleFormatin_fmt, intchannels, intsample_rate, intapply_map){AudioConvert *ac;intin_planar, out_planar;ac=av_mallocz(sizeof(*ac));if(!ac) returnNULL;ac->avr=avr;ac->out_fmt=out_fmt;ac->in_fmt=in_fmt;ac->channels=channels;ac->apply_map=apply_map;if(avr->dither_method!=AV_RESAMPLE_DITHER_NONE &&av_get_packed_sample_fmt(out_fmt)==AV_SAMPLE_FMT_S16 &&av_get_bytes_per_sample(in_fmt)>2){ac->dc=ff_dither_alloc(avr, out_fmt, in_fmt, channels, sample_rate, apply_map);if(!ac->dc){av_free(ac);returnNULL;}returnac;}in_planar=ff_sample_fmt_is_planar(in_fmt, channels);out_planar=ff_sample_fmt_is_planar(out_fmt, channels);if(in_planar==out_planar){ac->func_type=CONV_FUNC_TYPE_FLAT;ac->planes=in_planar?ac->channels:1;}elseif(in_planar) ac->func_type=CONV_FUNC_TYPE_INTERLEAVE;elseac->func_type=CONV_FUNC_TYPE_DEINTERLEAVE;set_generic_function(ac);if(ARCH_AARCH64) ff_audio_convert_init_aarch64(ac);if(ARCH_ARM) ff_audio_convert_init_arm(ac);if(ARCH_X86) ff_audio_convert_init_x86(ac);returnac;}intff_audio_convert(AudioConvert *ac, AudioData *out, AudioData *in){intuse_generic=1;intlen=in->nb_samples;intp;if(ac->dc){av_log(ac->avr, AV_LOG_TRACE,"%dsamples-audio_convert:%sto%s(dithered)\n", len, av_get_sample_fmt_name(ac->in_fmt), av_get_sample_fmt_name(ac->out_fmt));returnff_convert_dither(ac-> in
BYTE int const BYTE int int int height
UnsharpFilterParam chroma
chroma parameters (width, height, amount)
void av_opencl_buffer_release(cl_mem *cl_buf)
Release OpenCL buffer.
cl_command_queue av_opencl_get_command_queue(void)
get OpenCL command queue
uint8_t * data[AV_NUM_DATA_POINTERS]
pointer to the picture/channel planes.
common internal and external API header
int av_opencl_buffer_write(cl_mem dst_cl_buf, uint8_t *src_buf, size_t buf_size)
Write OpenCL buffer with data from src_buf.
void av_opencl_uninit(void)
Release OpenCL environment.
#define FF_OPENCL_PARAM_INFO(a)
int av_opencl_init(AVOpenCLExternalEnv *ext_opencl_env)
Initialize the run time OpenCL environment.
static void * av_mallocz_array(size_t nmemb, size_t size)
#define AVERROR_EXTERNAL
Generic error in an external library.
static int generate_mask(AVFilterContext *ctx)
void * av_mallocz(size_t size)
Allocate a block of size bytes with alignment suitable for all memory accesses (including vectors if ...
int ff_opencl_apply_unsharp(AVFilterContext *ctx, AVFrame *in, AVFrame *out)
#define AV_CEIL_RSHIFT(a, b)
int av_opencl_buffer_read_image(uint8_t **dst_data, int *plane_size, int plane_num, cl_mem src_cl_buf, size_t cl_buffer_size)
Read image data from OpenCL buffer.