- 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->dims[p]);
- 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_mem), &ctx->matrix[p]);
- if (cle != CL_SUCCESS) {
- av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
- "matrix argument: %d.\n", cle);
- goto fail;
- }
- cle = clSetKernelArg(ctx->kernel, 4, sizeof(cl_float), &ctx->rdivs[p]);
- if (cle != CL_SUCCESS) {
- av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
- "rdiv argument: %d.\n", cle);
- goto fail;
- }
- cle = clSetKernelArg(ctx->kernel, 5, sizeof(cl_float), &ctx->biases[p]);
- if (cle != CL_SUCCESS) {
- av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
- "bias argument: %d.\n", cle);
- goto fail;
- }
-
-
- err = ff_opencl_filter_work_size_from_image(avctx, global_work, output, p, 0);
- if (err < 0)
- goto fail;
-
- av_log(avctx, AV_LOG_DEBUG, "Run kernel on plane %d "
- "(%"SIZE_SPECIFIER"x%"SIZE_SPECIFIER").\n",
- p, global_work[0], global_work[1]);
-
- cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
- global_work, NULL,
- 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;
+ if (!strcmp(avctx->filter->name, "convolution_opencl")) {
+ 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->dims[p]);
+ CL_SET_KERNEL_ARG(ctx->kernel, 3, cl_mem, &ctx->matrix[p]);
+ CL_SET_KERNEL_ARG(ctx->kernel, 4, cl_float, &ctx->rdivs[p]);
+ CL_SET_KERNEL_ARG(ctx->kernel, 5, cl_float, &ctx->biases[p]);
+
+ err = ff_opencl_filter_work_size_from_image(avctx, global_work, output, p, 0);
+ if (err < 0)
+ goto fail;
+
+ av_log(avctx, AV_LOG_DEBUG, "Run kernel on plane %d "
+ "(%"SIZE_SPECIFIER"x%"SIZE_SPECIFIER").\n",
+ p, global_work[0], global_work[1]);
+
+ cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
+ global_work, NULL,
+ 0, NULL, NULL);
+ CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue "
+ "kernel: %d.\n", cle);
+ } else {
+ if (!(ctx->planes & (1 << p))) {
+ err = ff_opencl_filter_work_size_from_image(avctx, region, output, p, 0);
+ if (err < 0)
+ goto fail;
+
+ cle = clEnqueueCopyImage(ctx->command_queue, src, dst,
+ origin, origin, region, 0, NULL, NULL);
+ CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to copy plane %d: %d.\n",
+ p, cle);
+ } else {
+ 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_float, &ctx->scale);
+ CL_SET_KERNEL_ARG(ctx->kernel, 3, cl_float, &ctx->delta);
+
+ err = ff_opencl_filter_work_size_from_image(avctx, global_work, output, p, 0);
+ if (err < 0)
+ goto fail;
+
+ av_log(avctx, AV_LOG_DEBUG, "Run kernel on plane %d "
+ "(%"SIZE_SPECIFIER"x%"SIZE_SPECIFIER").\n",
+ p, global_work[0], global_work[1]);
+
+ cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
+ global_work, NULL,
+ 0, NULL, NULL);
+ CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue "
+ "kernel: %d.\n", cle);
+ }