Beispiel #1
0
int ff_opencl_transform(AVFilterContext *ctx,
                        int width, int height, int cw, int ch,
                        const float *matrix_y, const float *matrix_uv,
                        enum InterpolateMethod interpolate,
                        enum FillMethod fill, AVFrame *in, AVFrame *out)
{
    int ret = 0;
    const size_t global_work_size = width * height + 2 * ch * cw;
    cl_int status;
    DeshakeContext *deshake = ctx->priv;
    FFOpenclParam opencl_param = {0};

    opencl_param.ctx = ctx;
    opencl_param.kernel = deshake->opencl_ctx.kernel;
    ret = av_opencl_buffer_write(deshake->opencl_ctx.cl_matrix_y, (uint8_t *)matrix_y, deshake->opencl_ctx.matrix_size * sizeof(cl_float));
    if (ret < 0)
        return ret;
    ret = av_opencl_buffer_write(deshake->opencl_ctx.cl_matrix_uv, (uint8_t *)matrix_uv, deshake->opencl_ctx.matrix_size * sizeof(cl_float));
    if (ret < 0)
        return ret;

    if ((unsigned int)interpolate > INTERPOLATE_BIQUADRATIC) {
        av_log(ctx, AV_LOG_ERROR, "Selected interpolate method is invalid\n");
        return AVERROR(EINVAL);
    }
    ret = ff_opencl_set_parameter(&opencl_param,
                                  FF_OPENCL_PARAM_INFO(deshake->opencl_ctx.cl_inbuf),
                                  FF_OPENCL_PARAM_INFO(deshake->opencl_ctx.cl_outbuf),
                                  FF_OPENCL_PARAM_INFO(deshake->opencl_ctx.cl_matrix_y),
                                  FF_OPENCL_PARAM_INFO(deshake->opencl_ctx.cl_matrix_uv),
                                  FF_OPENCL_PARAM_INFO(interpolate),
                                  FF_OPENCL_PARAM_INFO(fill),
                                  FF_OPENCL_PARAM_INFO(in->linesize[0]),
                                  FF_OPENCL_PARAM_INFO(out->linesize[0]),
                                  FF_OPENCL_PARAM_INFO(in->linesize[1]),
                                  FF_OPENCL_PARAM_INFO(out->linesize[1]),
                                  FF_OPENCL_PARAM_INFO(height),
                                  FF_OPENCL_PARAM_INFO(width),
                                  FF_OPENCL_PARAM_INFO(ch),
                                  FF_OPENCL_PARAM_INFO(cw),
                                  NULL);
    if (ret < 0)
        return ret;
    status = clEnqueueNDRangeKernel(deshake->opencl_ctx.command_queue,
                                    deshake->opencl_ctx.kernel, 1, NULL,
                                    &global_work_size, NULL, 0, NULL, NULL);
    if (status != CL_SUCCESS) {
        av_log(ctx, AV_LOG_ERROR, "OpenCL run kernel error occurred: %s\n", av_opencl_errstr(status));
        return AVERROR_EXTERNAL;
    }
    clFinish(deshake->opencl_ctx.command_queue);
    ret = av_opencl_buffer_read_image(out->data, deshake->opencl_ctx.out_plane_size,
                                      deshake->opencl_ctx.plane_num, deshake->opencl_ctx.cl_outbuf,
                                      deshake->opencl_ctx.cl_outbuf_size);
    if (ret < 0)
        return ret;
    return ret;
}
Beispiel #2
0
int ff_opencl_transform(AVFilterContext *ctx,
                        int width, int height, int cw, int ch,
                        const float *matrix_y, const float *matrix_uv,
                        enum InterpolateMethod interpolate,
                        enum FillMethod fill, AVFrame *in, AVFrame *out)
{
    int ret = 0;
    cl_int status;
    DeshakeContext *deshake = ctx->priv;
    float4 packed_matrix_lu = {matrix_y[0], matrix_y[1], matrix_y[2], matrix_y[5]};
    float4 packed_matrix_ch = {matrix_uv[0], matrix_uv[1], matrix_uv[2], matrix_uv[5]};
    size_t global_worksize_lu[2] = {(size_t)ROUND_TO_16(width), (size_t)ROUND_TO_16(height)};
    size_t global_worksize_ch[2] = {(size_t)ROUND_TO_16(cw), (size_t)(2*ROUND_TO_16(ch))};
    size_t local_worksize[2] = {16, 16};
    FFOpenclParam param_lu = {0};
    FFOpenclParam param_ch = {0};
    param_lu.ctx = param_ch.ctx = ctx;
    param_lu.kernel = deshake->opencl_ctx.kernel_luma;
    param_ch.kernel = deshake->opencl_ctx.kernel_chroma;

    if ((unsigned int)interpolate > INTERPOLATE_BIQUADRATIC) {
        av_log(ctx, AV_LOG_ERROR, "Selected interpolate method is invalid\n");
        return AVERROR(EINVAL);
    }
    ret = ff_opencl_set_parameter(&param_lu,
                                  FF_OPENCL_PARAM_INFO(deshake->opencl_ctx.cl_inbuf),
                                  FF_OPENCL_PARAM_INFO(deshake->opencl_ctx.cl_outbuf),
                                  FF_OPENCL_PARAM_INFO(packed_matrix_lu),
                                  FF_OPENCL_PARAM_INFO(interpolate),
                                  FF_OPENCL_PARAM_INFO(fill),
                                  FF_OPENCL_PARAM_INFO(in->linesize[0]),
                                  FF_OPENCL_PARAM_INFO(out->linesize[0]),
                                  FF_OPENCL_PARAM_INFO(height),
                                  FF_OPENCL_PARAM_INFO(width),
                                  NULL);
    if (ret < 0)
        return ret;
    ret = ff_opencl_set_parameter(&param_ch,
                                  FF_OPENCL_PARAM_INFO(deshake->opencl_ctx.cl_inbuf),
                                  FF_OPENCL_PARAM_INFO(deshake->opencl_ctx.cl_outbuf),
                                  FF_OPENCL_PARAM_INFO(packed_matrix_ch),
                                  FF_OPENCL_PARAM_INFO(interpolate),
                                  FF_OPENCL_PARAM_INFO(fill),
                                  FF_OPENCL_PARAM_INFO(in->linesize[0]),
                                  FF_OPENCL_PARAM_INFO(out->linesize[0]),
                                  FF_OPENCL_PARAM_INFO(in->linesize[1]),
                                  FF_OPENCL_PARAM_INFO(out->linesize[1]),
                                  FF_OPENCL_PARAM_INFO(height),
                                  FF_OPENCL_PARAM_INFO(width),
                                  FF_OPENCL_PARAM_INFO(ch),
                                  FF_OPENCL_PARAM_INFO(cw),
                                  NULL);
    if (ret < 0)
        return ret;
    status = clEnqueueNDRangeKernel(deshake->opencl_ctx.command_queue,
                                    deshake->opencl_ctx.kernel_luma, 2, NULL,
                                    global_worksize_lu, local_worksize, 0, NULL, NULL);
    status |= clEnqueueNDRangeKernel(deshake->opencl_ctx.command_queue,
                                    deshake->opencl_ctx.kernel_chroma, 2, NULL,
                                    global_worksize_ch, local_worksize, 0, NULL, NULL);
    if (status != CL_SUCCESS) {
        av_log(ctx, AV_LOG_ERROR, "OpenCL run kernel error occurred: %s\n", av_opencl_errstr(status));
        return AVERROR_EXTERNAL;
    }
    ret = av_opencl_buffer_read_image(out->data, deshake->opencl_ctx.out_plane_size,
                                      deshake->opencl_ctx.plane_num, deshake->opencl_ctx.cl_outbuf,
                                      deshake->opencl_ctx.cl_outbuf_size);
    if (ret < 0)
        return ret;
    return ret;
}
int ff_opencl_apply_unsharp(AVFilterContext *ctx, AVFrame *in, AVFrame *out)
{
    int ret;
    AVFilterLink *link = ctx->inputs[0];
    UnsharpContext *unsharp = ctx->priv;
    cl_int status;
    FFOpenclParam kernel1 = {0};
    FFOpenclParam kernel2 = {0};
    int width = link->w;
    int height = link->h;
    int cw = FF_CEIL_RSHIFT(link->w, unsharp->hsub);
    int ch = FF_CEIL_RSHIFT(link->h, unsharp->vsub);
    size_t globalWorkSize1d = width * height + 2 * ch * cw;
    size_t globalWorkSize2dLuma[2];
    size_t globalWorkSize2dChroma[2];
    size_t localWorkSize2d[2] = {16, 16};

    if (unsharp->opencl_ctx.use_fast_kernels) {
        globalWorkSize2dLuma[0] = (size_t)ROUND_TO_16(width);
        globalWorkSize2dLuma[1] = (size_t)ROUND_TO_16(height);
        globalWorkSize2dChroma[0] = (size_t)ROUND_TO_16(cw);
        globalWorkSize2dChroma[1] = (size_t)(2*ROUND_TO_16(ch));

        kernel1.ctx = ctx;
        kernel1.kernel = unsharp->opencl_ctx.kernel_luma;
        ret = ff_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->luma.amount),
                                      FF_OPENCL_PARAM_INFO(unsharp->luma.scalebits),
                                      FF_OPENCL_PARAM_INFO(unsharp->luma.halfscale),
                                      FF_OPENCL_PARAM_INFO(in->linesize[0]),
                                      FF_OPENCL_PARAM_INFO(out->linesize[0]),
                                      FF_OPENCL_PARAM_INFO(width),
                                      FF_OPENCL_PARAM_INFO(height),
                                      NULL);
        if (ret < 0)
            return ret;

        kernel2.ctx = ctx;
        kernel2.kernel = unsharp->opencl_ctx.kernel_chroma;
        ret = ff_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->chroma.amount),
                                      FF_OPENCL_PARAM_INFO(unsharp->chroma.scalebits),
                                      FF_OPENCL_PARAM_INFO(unsharp->chroma.halfscale),
                                      FF_OPENCL_PARAM_INFO(in->linesize[0]),
                                      FF_OPENCL_PARAM_INFO(in->linesize[1]),
                                      FF_OPENCL_PARAM_INFO(out->linesize[0]),
                                      FF_OPENCL_PARAM_INFO(out->linesize[1]),
                                      FF_OPENCL_PARAM_INFO(link->w),
                                      FF_OPENCL_PARAM_INFO(link->h),
                                      FF_OPENCL_PARAM_INFO(cw),
                                      FF_OPENCL_PARAM_INFO(ch),
                                      NULL);
        if (ret < 0)
            return ret;
        status = clEnqueueNDRangeKernel(unsharp->opencl_ctx.command_queue,
                                        unsharp->opencl_ctx.kernel_luma, 2, NULL,
                                        globalWorkSize2dLuma, localWorkSize2d, 0, NULL, NULL);
        status |=clEnqueueNDRangeKernel(unsharp->opencl_ctx.command_queue,
                                        unsharp->opencl_ctx.kernel_chroma, 2, NULL,
                                        globalWorkSize2dChroma, localWorkSize2d, 0, NULL, NULL);
        if (status != CL_SUCCESS) {
            av_log(ctx, AV_LOG_ERROR, "OpenCL run kernel error occurred: %s\n", av_opencl_errstr(status));
            return AVERROR_EXTERNAL;
        }
    } else {    /* use default kernel */
        kernel1.ctx = ctx;
        kernel1.kernel = unsharp->opencl_ctx.kernel_default;

        ret = ff_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_chroma_mask),
                                      FF_OPENCL_PARAM_INFO(unsharp->luma.amount),
                                      FF_OPENCL_PARAM_INFO(unsharp->chroma.amount),
                                      FF_OPENCL_PARAM_INFO(unsharp->luma.steps_x),
                                      FF_OPENCL_PARAM_INFO(unsharp->luma.steps_y),
                                      FF_OPENCL_PARAM_INFO(unsharp->chroma.steps_x),
                                      FF_OPENCL_PARAM_INFO(unsharp->chroma.steps_y),
                                      FF_OPENCL_PARAM_INFO(unsharp->luma.scalebits),
                                      FF_OPENCL_PARAM_INFO(unsharp->chroma.scalebits),
                                      FF_OPENCL_PARAM_INFO(unsharp->luma.halfscale),
                                      FF_OPENCL_PARAM_INFO(unsharp->chroma.halfscale),
                                      FF_OPENCL_PARAM_INFO(in->linesize[0]),
                                      FF_OPENCL_PARAM_INFO(in->linesize[1]),
                                      FF_OPENCL_PARAM_INFO(out->linesize[0]),
                                      FF_OPENCL_PARAM_INFO(out->linesize[1]),
                                      FF_OPENCL_PARAM_INFO(link->h),
                                      FF_OPENCL_PARAM_INFO(link->w),
                                      FF_OPENCL_PARAM_INFO(ch),
                                      FF_OPENCL_PARAM_INFO(cw),
                                      NULL);
        if (ret < 0)
            return ret;
        status = clEnqueueNDRangeKernel(unsharp->opencl_ctx.command_queue,
                                        unsharp->opencl_ctx.kernel_default, 1, NULL,
                                        &globalWorkSize1d, NULL, 0, NULL, NULL);
        if (status != CL_SUCCESS) {
            av_log(ctx, AV_LOG_ERROR, "OpenCL run kernel error occurred: %s\n", av_opencl_errstr(status));
            return AVERROR_EXTERNAL;
        }
    }
    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);
}