ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
ctx->ocf.hwctx->device_id,
0, &cle);
- if (!ctx->command_queue) {
- av_log(avctx, AV_LOG_ERROR, "Failed to create OpenCL "
- "command queue: %d.\n", cle);
- err = AVERROR(EIO);
- goto fail;
- }
+ CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
+ "command queue %d.\n", cle);
// Use global kernel if mask size will be too big for the local store..
ctx->global = (ctx->luma_size_x > 17.0f ||
ctx->kernel = clCreateKernel(ctx->ocf.program,
ctx->global ? "unsharp_global"
: "unsharp_local", &cle);
- if (!ctx->kernel) {
- av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle);
- err = AVERROR(EIO);
- goto fail;
- }
+ CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle);
ctx->initialised = 1;
return 0;
CL_MEM_COPY_HOST_PTR |
CL_MEM_HOST_NO_ACCESS,
matrix_bytes, matrix, &cle);
- if (!buffer) {
- av_log(avctx, AV_LOG_ERROR, "Failed to create matrix buffer: "
- "%d.\n", cle);
- err = AVERROR(EIO);
- goto fail;
- }
+ CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create matrix buffer: "
+ "%d.\n", cle);
ctx->plane[p].matrix = buffer;
} else {
buffer = clCreateBuffer(ctx->ocf.hwctx->context,
CL_MEM_HOST_NO_ACCESS,
sizeof(ctx->plane[p].blur_x),
ctx->plane[p].blur_x, &cle);
- if (!buffer) {
- av_log(avctx, AV_LOG_ERROR, "Failed to create x-coef buffer: "
- "%d.\n", cle);
- err = AVERROR(EIO);
- goto fail;
- }
+ CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create x-coef buffer: "
+ "%d.\n", cle);
ctx->plane[p].coef_x = buffer;
buffer = clCreateBuffer(ctx->ocf.hwctx->context,
CL_MEM_HOST_NO_ACCESS,
sizeof(ctx->plane[p].blur_y),
ctx->plane[p].blur_y, &cle);
- if (!buffer) {
- av_log(avctx, AV_LOG_ERROR, "Failed to create y-coef buffer: "
- "%d.\n", cle);
- err = AVERROR(EIO);
- goto fail;
- }
+ CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create y-coef buffer: "
+ "%d.\n", cle);
ctx->plane[p].coef_y = buffer;
}
if (!dst)
break;
- cle = clSetKernelArg(ctx->kernel, 0, sizeof(cl_mem), &dst);
- if (cle != CL_SUCCESS) {
- av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
- "destination image argument: %d.\n", cle);
- goto fail;
- }
- cle = clSetKernelArg(ctx->kernel, 1, sizeof(cl_mem), &src);
- if (cle != CL_SUCCESS) {
- av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
- "source image argument: %d.\n", cle);
- goto fail;
- }
- cle = clSetKernelArg(ctx->kernel, 2, sizeof(cl_int), &ctx->plane[p].size_x);
- if (cle != CL_SUCCESS) {
- av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
- "matrix size argument: %d.\n", cle);
- goto fail;
- }
- cle = clSetKernelArg(ctx->kernel, 3, sizeof(cl_int), &ctx->plane[p].size_y);
- if (cle != CL_SUCCESS) {
- av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
- "matrix size argument: %d.\n", cle);
- goto fail;
- }
- cle = clSetKernelArg(ctx->kernel, 4, sizeof(cl_float), &ctx->plane[p].amount);
- if (cle != CL_SUCCESS) {
- av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
- "amount argument: %d.\n", cle);
- goto fail;
- }
+ CL_SET_KERNEL_ARG(ctx->kernel, 0, cl_mem, &dst);
+ CL_SET_KERNEL_ARG(ctx->kernel, 1, cl_mem, &src);
+ CL_SET_KERNEL_ARG(ctx->kernel, 2, cl_int, &ctx->plane[p].size_x);
+ CL_SET_KERNEL_ARG(ctx->kernel, 3, cl_int, &ctx->plane[p].size_y);
+ CL_SET_KERNEL_ARG(ctx->kernel, 4, cl_float, &ctx->plane[p].amount);
+
if (ctx->global) {
- cle = clSetKernelArg(ctx->kernel, 5, sizeof(cl_mem), &ctx->plane[p].matrix);
- if (cle != CL_SUCCESS) {
- av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
- "matrix argument: %d.\n", cle);
- goto fail;
- }
+ CL_SET_KERNEL_ARG(ctx->kernel, 5, cl_mem, &ctx->plane[p].matrix);
} else {
- cle = clSetKernelArg(ctx->kernel, 5, sizeof(cl_mem), &ctx->plane[p].coef_x);
- if (cle != CL_SUCCESS) {
- av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
- "x-coef argument: %d.\n", cle);
- goto fail;
- }
- cle = clSetKernelArg(ctx->kernel, 6, sizeof(cl_mem), &ctx->plane[p].coef_y);
- if (cle != CL_SUCCESS) {
- av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
- "y-coef argument: %d.\n", cle);
- goto fail;
- }
+ CL_SET_KERNEL_ARG(ctx->kernel, 5, cl_mem, &ctx->plane[p].coef_x);
+ CL_SET_KERNEL_ARG(ctx->kernel, 6, cl_mem, &ctx->plane[p].coef_y);
}
err = ff_opencl_filter_work_size_from_image(avctx, global_work, output, p,
cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
global_work, ctx->global ? NULL : local_work,
0, NULL, NULL);
- if (cle != CL_SUCCESS) {
- av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n",
- cle);
- err = AVERROR(EIO);
- goto fail;
- }
+ CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue kernel: %d.\n", cle);
}
cle = clFinish(ctx->command_queue);
- if (cle != CL_SUCCESS) {
- av_log(avctx, AV_LOG_ERROR, "Failed to finish command queue: %d.\n",
- cle);
- err = AVERROR(EIO);
- goto fail;
- }
+ CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
err = av_frame_copy_props(output, input);
if (err < 0)
{ NULL }
};
-AVFilter ff_vf_unsharp_opencl = {
+const AVFilter ff_vf_unsharp_opencl = {
.name = "unsharp_opencl",
.description = NULL_IF_CONFIG_SMALL("Apply unsharp mask to input video"),
.priv_size = sizeof(UnsharpOpenCLContext),