Hi, I've fixed the patch according to the reviewers suggestions. Attached are the two patches (original and fixes)
Regards, Alexey >From e94ea7add32791a1c9ac04dbd68b78cf799ca92d Mon Sep 17 00:00:00 2001 From: atitov <alexey.ti...@amd.com> Date: Wed, 17 Dec 2014 16:02:17 -0800 Subject: [PATCH 1/2] libavfilter: unsharpen opencl filter optimizations --- libavfilter/unsharp.h | 4 ++ libavfilter/unsharp_opencl.c | 76 +++++++++++++++------- libavfilter/unsharp_opencl_kernel.h | 122 ++++++++++++++++++++++++++---------- libavutil/opencl.c | 19 +++++- 4 files changed, 165 insertions(+), 56 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..a99fc5b 100644 --- a/libavfilter/unsharp_opencl.c +++ b/libavfilter/unsharp_opencl.c @@ -87,42 +87,36 @@ 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); - mask_x = av_mallocz_array(2 * step_x + 1, sizeof(uint32_t)); + 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(size_mask_x); if (!mask_x) { ret = AVERROR(ENOMEM); goto end; } - mask_y = av_mallocz_array(2 * step_y + 1, sizeof(uint32_t)); + mask_y = av_mallocz_array(size_mask_y); if (!mask_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; } @@ -133,6 +127,11 @@ static int generate_mask(AVFilterContext *ctx) cl_mem mask_matrix[2]; mask_matrix[0] = unsharp->opencl_ctx.cl_luma_mask; mask_matrix[1] = unsharp->opencl_ctx.cl_chroma_mask; + cl_mem masks[4]; + 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 +143,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 +187,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 +205,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 +269,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 +293,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 +367,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..e013e2f 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,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); } diff --git a/libavutil/opencl.c b/libavutil/opencl.c index 36cb6fe..738d0db 100644 --- a/libavutil/opencl.c +++ b/libavutil/opencl.c @@ -450,7 +450,24 @@ cl_program av_opencl_compile(const char *program_name, const char *build_opts) status = clBuildProgram(program, 1, &(opencl_ctx.device_id), build_opts, NULL, NULL); if (status != CL_SUCCESS) { av_log(&opencl_ctx, AV_LOG_ERROR, - "Compilation failed with OpenCL program: %s\n", program_name); + "Compilation failed with OpenCL program: '%s' with error %d \n", program_name, status); + + // Determine the size of the log + size_t log_size; + clGetProgramBuildInfo(program, &(opencl_ctx.device_id), CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size); + + // Allocate memory for the log + char *log = (char *) malloc(log_size+1); + + // Get the log + clGetProgramBuildInfo(program, &(opencl_ctx.device_id), CL_PROGRAM_BUILD_LOG, log_size, log, NULL); + log[log_size] = '\0'; + printf("--- Build log ---\n "); + // Print the log + printf("%s\n", log); + printf("--- End Build log ---\n "); + + free(log); program = NULL; goto end; } -- 1.8.4.msysgit.0 >From 72f4e48f3adf06fd859abc4aecede4e15a36e9d9 Mon Sep 17 00:00:00 2001 From: atitov <alexey.ti...@amd.com> Date: Tue, 23 Dec 2014 13:16:35 -0800 Subject: [PATCH 2/2] Fixed the patch according to the suggestions build log is now capped and is using av_log, other issues are fixed --- libavfilter/unsharp_opencl.c | 11 ++++++----- libavfilter/unsharp_opencl_kernel.h | 2 +- libavutil/opencl.c | 29 +++++++++++++++++------------ libavutil/opencl.h | 2 ++ 4 files changed, 26 insertions(+), 18 deletions(-) diff --git a/libavfilter/unsharp_opencl.c b/libavfilter/unsharp_opencl.c index a99fc5b..4adad63 100644 --- a/libavfilter/unsharp_opencl.c +++ b/libavfilter/unsharp_opencl.c @@ -93,12 +93,12 @@ static int copy_separable_masks(cl_mem cl_mask_x, cl_mem cl_mask_y, int step_x, 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(size_mask_x); + mask_x = av_mallocz_array((2 * step_x + 1), sizeof(uint32_t)); if (!mask_x) { ret = AVERROR(ENOMEM); goto end; } - mask_y = av_mallocz_array(size_mask_y); + mask_y = av_mallocz_array((2 * step_y + 1), sizeof(uint32_t)); if (!mask_y) { ret = AVERROR(ENOMEM); goto end; @@ -122,12 +122,13 @@ end: 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; - cl_mem masks[4]; 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; diff --git a/libavfilter/unsharp_opencl_kernel.h b/libavfilter/unsharp_opencl_kernel.h index e013e2f..0fc802e 100644 --- a/libavfilter/unsharp_opencl_kernel.h +++ b/libavfilter/unsharp_opencl_kernel.h @@ -89,7 +89,7 @@ kernel void unsharp_luma( int orig_value = (int)l[threadIdx.y + 8][threadIdx.x + 8]; int idx, idy, maskIndex; - int temp[2] = {0,0}; + int temp[2] = {0}; int steps_x = (LU_RADIUS_X-1)/2; int steps_y = (LU_RADIUS_Y-1)/2; diff --git a/libavutil/opencl.c b/libavutil/opencl.c index 738d0db..3cd2bf8 100644 --- a/libavutil/opencl.c +++ b/libavutil/opencl.c @@ -449,27 +449,32 @@ cl_program av_opencl_compile(const char *program_name, const char *build_opts) } status = clBuildProgram(program, 1, &(opencl_ctx.device_id), build_opts, NULL, NULL); if (status != CL_SUCCESS) { + size_t log_size; + char *log; av_log(&opencl_ctx, AV_LOG_ERROR, "Compilation failed with OpenCL program: '%s' with error %d \n", program_name, status); // Determine the size of the log - size_t log_size; - clGetProgramBuildInfo(program, &(opencl_ctx.device_id), CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size); + clGetProgramBuildInfo(program, opencl_ctx.device_id, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size); - // Allocate memory for the log - char *log = (char *) malloc(log_size+1); + if( log_size <= AV_OPENCL_MAX_PROG_BUILD_LOG_SIZE) { - // Get the log - clGetProgramBuildInfo(program, &(opencl_ctx.device_id), CL_PROGRAM_BUILD_LOG, log_size, log, NULL); - log[log_size] = '\0'; - printf("--- Build log ---\n "); - // Print the log - printf("%s\n", log); - printf("--- End Build log ---\n "); + // Allocate memory for the log + log = av_malloc(log_size+1); - free(log); + // Get the log + clGetProgramBuildInfo(program, opencl_ctx.device_id, CL_PROGRAM_BUILD_LOG, log_size, log, NULL); + log[log_size] = '\0'; + // Print the log + av_log(NULL, AV_LOG_ERROR, "--- Build log ---\n"); + av_log(NULL, AV_LOG_ERROR, "%s\n", log); + av_log(NULL, AV_LOG_ERROR, "--- End Build log ---\n "); + + av_free(log); + } program = NULL; goto end; + } opencl_ctx.kernel_code[kernel_code_idx].is_compiled = 1; diff --git a/libavutil/opencl.h b/libavutil/opencl.h index 4655cba..c2c327e 100644 --- a/libavutil/opencl.h +++ b/libavutil/opencl.h @@ -51,6 +51,8 @@ #define AV_OPENCL_MAX_PLATFORM_NAME_SIZE 100 +#define AV_OPENCL_MAX_PROG_BUILD_LOG_SIZE 1000 + typedef struct { int device_type; char device_name[AV_OPENCL_MAX_DEVICE_NAME_SIZE]; -- 1.8.4.msysgit.0 _______________________________________________ ffmpeg-devel mailing list ffmpeg-devel@ffmpeg.org http://ffmpeg.org/mailman/listinfo/ffmpeg-devel