diff options
33 files changed, 1541 insertions, 70 deletions
diff --git a/examples/multimedia/video/qmlvideofilter_opencl/README b/examples/multimedia/video/qmlvideofilter_opencl/README new file mode 100644 index 000000000..c239bed2f --- /dev/null +++ b/examples/multimedia/video/qmlvideofilter_opencl/README @@ -0,0 +1,18 @@ +This example performs some simple OpenCL operations on camera or video input which +is assumed to be provided in RGB format. The OpenCL operation is done on an +OpenGL texture using CL-GL interop, without any further readbacks or copies +(except for the initial texture upload, when necessary). + +Currently only OS X and Windows with desktop OpenGL (opengl32.dll) are supported. +On Windows you may need to edit testplugin.pro to specify the location of the OpenCL +headers and libraries. + +Note that an OpenCL implementation with GPU support is required. +The platform and device selection logic supports NVIDIA and Intel. +Porting to other platforms is probably simple, see clCreateContextFromType. +Note however that YUV formats, that are commonly used also for camera input +on some platforms, are not supported in this example. + +Pass the name of a video file to perform video playback or launch without +arguments to use the camera. + diff --git a/examples/multimedia/video/qmlvideofilter_opencl/main.cpp b/examples/multimedia/video/qmlvideofilter_opencl/main.cpp new file mode 100644 index 000000000..7e2796395 --- /dev/null +++ b/examples/multimedia/video/qmlvideofilter_opencl/main.cpp @@ -0,0 +1,481 @@ +/**************************************************************************** +** +** Copyright (C) 2015 Digia Plc and/or its subsidiary(-ies). +** Contact: http://www.qt-project.org/legal +** +** This file is part of the examples of the Qt Multimedia module. +** +** $QT_BEGIN_LICENSE:LGPL21$ +** Commercial License Usage +** Licensees holding valid commercial Qt licenses may use this file in +** accordance with the commercial license agreement provided with the +** Software or, alternatively, in accordance with the terms contained in +** a written agreement between you and Digia. For licensing terms and +** conditions see http://qt.digia.com/licensing. For further information +** use the contact form at http://qt.digia.com/contact-us. +** +** GNU Lesser General Public License Usage +** Alternatively, this file may be used under the terms of the GNU Lesser +** General Public License version 2.1 or version 3 as published by the Free +** Software Foundation and appearing in the file LICENSE.LGPLv21 and +** LICENSE.LGPLv3 included in the packaging of this file. Please review the +** following information to ensure the GNU Lesser General Public License +** requirements will be met: https://www.gnu.org/licenses/lgpl.html and +** http://www.gnu.org/licenses/old-licenses/lgpl-2.1.html. +** +** In addition, as a special exception, Digia gives you certain additional +** rights. These rights are described in the Digia Qt LGPL Exception +** version 1.1, included in the file LGPL_EXCEPTION.txt in this package. +** +** $QT_END_LICENSE$ +** +****************************************************************************/ + +#include <QGuiApplication> +#include <QQuickView> +#include <QOpenGLContext> +#include <QOpenGLFunctions> +#include <QAbstractVideoFilter> +#include <QQmlContext> +#include <QFileInfo> + +#ifdef Q_OS_OSX +#include <OpenCL/opencl.h> +#include <OpenGL/OpenGL.h> +#else +#include <CL/opencl.h> +#endif + +#include "rgbframehelper.h" + +static const char *openclSrc = + "__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;\n" + "__kernel void Emboss(__read_only image2d_t imgIn, __write_only image2d_t imgOut, float factor) {\n" + " const int2 pos = { get_global_id(0), get_global_id(1) };\n" + " float4 diff = read_imagef(imgIn, sampler, pos + (int2)(1,1)) - read_imagef(imgIn, sampler, pos - (int2)(1,1));\n" + " float color = (diff.x + diff.y + diff.z) / factor + 0.5f;\n" + " write_imagef(imgOut, pos, (float4)(color, color, color, 1.0f));\n" + "}\n"; + +class CLFilter : public QAbstractVideoFilter +{ + Q_OBJECT + Q_PROPERTY(qreal factor READ factor WRITE setFactor NOTIFY factorChanged) + +public: + CLFilter() : m_factor(1) { } + qreal factor() const { return m_factor; } + void setFactor(qreal v); + + QVideoFilterRunnable *createFilterRunnable() Q_DECL_OVERRIDE; + +signals: + void factorChanged(); + +private: + qreal m_factor; +}; + +class CLFilterRunnable : public QVideoFilterRunnable +{ +public: + CLFilterRunnable(CLFilter *filter); + ~CLFilterRunnable(); + QVideoFrame run(QVideoFrame *input, const QVideoSurfaceFormat &surfaceFormat, RunFlags flags) Q_DECL_OVERRIDE; + +private: + void releaseTextures(); + uint newTexture(); + + CLFilter *m_filter; + QSize m_size; + uint m_tempTexture; + uint m_outTexture; + uint m_lastInputTexture; + cl_context m_clContext; + cl_device_id m_clDeviceId; + cl_mem m_clImage[2]; + cl_command_queue m_clQueue; + cl_program m_clProgram; + cl_kernel m_clKernel; +}; + +QVideoFilterRunnable *CLFilter::createFilterRunnable() +{ + return new CLFilterRunnable(this); +} + +CLFilterRunnable::CLFilterRunnable(CLFilter *filter) : + m_filter(filter), + m_tempTexture(0), + m_outTexture(0), + m_lastInputTexture(0), + m_clContext(0), + m_clQueue(0), + m_clProgram(0), + m_clKernel(0) +{ + m_clImage[0] = m_clImage[1] = 0; + + // Set up OpenCL. + QOpenGLFunctions *f = QOpenGLContext::currentContext()->functions(); + cl_int err; + cl_uint n; + if (clGetPlatformIDs(0, 0, &n) != CL_SUCCESS) { + qWarning("Failed to get platform ID count"); + return; + } + if (n == 0) { + qWarning("No OpenCL platform found"); + return; + } + QVector<cl_platform_id> platformIds; + platformIds.resize(n); + if (clGetPlatformIDs(n, platformIds.data(), 0) != CL_SUCCESS) { + qWarning("Failed to get platform IDs"); + return; + } + cl_platform_id platform = platformIds[0]; + const char *vendor = (const char *) f->glGetString(GL_VENDOR); + qDebug("GL_VENDOR: %s", vendor); + const bool isNV = vendor && strstr(vendor, "NVIDIA"); + const bool isIntel = vendor && strstr(vendor, "Intel"); + qDebug("Found %u OpenCL platforms:", n); + for (cl_uint i = 0; i < n; ++i) { + QByteArray name; + name.resize(1024); + clGetPlatformInfo(platformIds[i], CL_PLATFORM_NAME, name.size(), name.data(), 0); + qDebug("Platform %p: %s", platformIds[i], name.constData()); + // Running with an OpenCL platform without GPU support is not going + // to cut it. In practice we want the platform for the GPU which we + // are using with OpenGL. + if (isNV && name.contains(QByteArrayLiteral("NVIDIA"))) + platform = platformIds[i]; + else if (isIntel && name.contains(QByteArrayLiteral("Intel"))) + platform = platformIds[i]; + } + qDebug("Using platform %p", platform); + + // Set up the context with OpenCL/OpenGL interop. +#if defined (Q_OS_OSX) + cl_context_properties contextProps[] = { CL_CONTEXT_PROPERTY_USE_CGL_SHAREGROUP_APPLE, + (cl_context_properties) CGLGetShareGroup(CGLGetCurrentContext()), + 0 }; +#elif defined(Q_OS_WIN) + cl_context_properties contextProps[] = { CL_CONTEXT_PLATFORM, (cl_context_properties) platform, + CL_GL_CONTEXT_KHR, (cl_context_properties) wglGetCurrentContext(), + CL_WGL_HDC_KHR, (cl_context_properties) wglGetCurrentDC(), + 0 }; +#endif + + m_clContext = clCreateContextFromType(contextProps, CL_DEVICE_TYPE_GPU, 0, 0, &err); + if (!m_clContext) { + qWarning("Failed to create OpenCL context: %d", err); + return; + } + + // Get the GPU device id +#if defined(Q_OS_OSX) + // On OS X, get the "online" device/GPU. This is required for OpenCL/OpenGL context sharing. + if (clGetGLContextInfoAPPLE(m_clContext, CGLGetCurrentContext(), + CL_CGL_DEVICE_FOR_CURRENT_VIRTUAL_SCREEN_APPLE, + sizeof(cl_device_id), &m_clDeviceId, NULL) != CL_SUCCESS) { + qWarning("Failed to get OpenCL device for current screen: %d", err); + return; + } +#else + if (clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &m_clDeviceId, 0) != CL_SUCCESS) { + qWarning("Failed to get OpenCL device"); + return; + } +#endif + + m_clQueue = clCreateCommandQueue(m_clContext, m_clDeviceId, 0, &err); + if (!m_clQueue) { + qWarning("Failed to create OpenCL command queue: %d", err); + return; + } + // Build the program. + m_clProgram = clCreateProgramWithSource(m_clContext, 1, &openclSrc, 0, &err); + if (!m_clProgram) { + qWarning("Failed to create OpenCL program: %d", err); + return; + } + if (clBuildProgram(m_clProgram, 1, &m_clDeviceId, 0, 0, 0) != CL_SUCCESS) { + qWarning("Failed to build OpenCL program"); + QByteArray log; + log.resize(2048); + clGetProgramBuildInfo(m_clProgram, m_clDeviceId, CL_PROGRAM_BUILD_LOG, log.size(), log.data(), 0); + qDebug("Build log: %s", log.constData()); + return; + } + m_clKernel = clCreateKernel(m_clProgram, "Emboss", &err); + if (!m_clKernel) { + qWarning("Failed to create emboss OpenCL kernel: %d", err); + return; + } +} + +CLFilterRunnable::~CLFilterRunnable() +{ + releaseTextures(); + if (m_clKernel) + clReleaseKernel(m_clKernel); + if (m_clProgram) + clReleaseProgram(m_clProgram); + if (m_clQueue) + clReleaseCommandQueue(m_clQueue); + if (m_clContext) + clReleaseContext(m_clContext); +} + +void CLFilterRunnable::releaseTextures() +{ + QOpenGLFunctions *f = QOpenGLContext::currentContext()->functions(); + if (m_tempTexture) + f->glDeleteTextures(1, &m_tempTexture); + if (m_outTexture) + f->glDeleteTextures(1, &m_outTexture); + m_tempTexture = m_outTexture = m_lastInputTexture = 0; + if (m_clImage[0]) + clReleaseMemObject(m_clImage[0]); + if (m_clImage[1]) + clReleaseMemObject(m_clImage[1]); + m_clImage[0] = m_clImage[1] = 0; +} + +uint CLFilterRunnable::newTexture() +{ + QOpenGLFunctions *f = QOpenGLContext::currentContext()->functions(); + GLuint texture; + f->glGenTextures(1, &texture); + f->glBindTexture(GL_TEXTURE_2D, texture); + f->glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_EDGE); + f->glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_EDGE); + f->glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST); + f->glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST); + f->glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA, m_size.width(), m_size.height(), + 0, GL_RGBA, GL_UNSIGNED_BYTE, 0); + return texture; +} + +QVideoFrame CLFilterRunnable::run(QVideoFrame *input, const QVideoSurfaceFormat &surfaceFormat, RunFlags flags) +{ + Q_UNUSED(surfaceFormat); + Q_UNUSED(flags); + + // This example supports RGB data only, either in system memory (typical with cameras on all + // platforms) or as an OpenGL texture (e.g. video playback on OS X or on Windows with ANGLE). + // The latter is the fast path where everything happens on GPU. THe former involves a texture upload. + + // ANGLE is not compatible with this example since we only do CL-GL interop, not D3D9/11. + if (QOpenGLContext::openGLModuleType() == QOpenGLContext::LibGLES) { + qWarning("ANGLE is not supported"); + return *input; + } + + if (!input->isValid() + || (input->handleType() != QAbstractVideoBuffer::NoHandle + && input->handleType() != QAbstractVideoBuffer::GLTextureHandle)) { + qWarning("Invalid input format"); + return *input; + } + + if (input->pixelFormat() == QVideoFrame::Format_YUV420P + || input->pixelFormat() == QVideoFrame::Format_YV12) { + qWarning("YUV data is not supported"); + return *input; + } + + if (m_size != input->size()) { + releaseTextures(); + m_size = input->size(); + } + + // Create a texture from the image data. + QOpenGLFunctions *f = QOpenGLContext::currentContext()->functions(); + GLuint texture; + if (input->handleType() == QAbstractVideoBuffer::NoHandle) { + // Upload. + if (m_tempTexture) + f->glBindTexture(GL_TEXTURE_2D, m_tempTexture); + else + m_tempTexture = newTexture(); + input->map(QAbstractVideoBuffer::ReadOnly); + // glTexImage2D only once and use TexSubImage later on. This avoids the need + // to recreate the CL image object on every frame. + f->glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, m_size.width(), m_size.height(), + GL_RGBA, GL_UNSIGNED_BYTE, input->bits()); + input->unmap(); + texture = m_tempTexture; + } else { + // Already an OpenGL texture. + texture = input->handle().toUInt(); + f->glBindTexture(GL_TEXTURE_2D, texture); + // Unlike on the other branch, the input texture may change, so m_clImage[0] may need to be recreated. + if (m_lastInputTexture && m_lastInputTexture != texture && m_clImage[0]) { + clReleaseMemObject(m_clImage[0]); + m_clImage[0] = 0; + } + m_lastInputTexture = texture; + } + + // OpenCL image objects cannot be read and written at the same time. So use + // a separate texture for the result. + if (!m_outTexture) + m_outTexture = newTexture(); + + // Create the image objects if not yet done. + cl_int err; + if (!m_clImage[0]) { + m_clImage[0] = clCreateFromGLTexture2D(m_clContext, CL_MEM_READ_ONLY, GL_TEXTURE_2D, 0, texture, &err); + if (!m_clImage[0]) { + qWarning("Failed to create OpenGL image object from OpenGL texture: %d", err); + return *input; + } + cl_image_format fmt; + if (clGetImageInfo(m_clImage[0], CL_IMAGE_FORMAT, sizeof(fmt), &fmt, 0) != CL_SUCCESS) { + qWarning("Failed to query image format"); + return *input; + } + if (fmt.image_channel_order != CL_RGBA) + qWarning("OpenCL image is not RGBA, expect errors"); + } + if (!m_clImage[1]) { + m_clImage[1] = clCreateFromGLTexture2D(m_clContext, CL_MEM_WRITE_ONLY, GL_TEXTURE_2D, 0, m_outTexture, &err); + if (!m_clImage[1]) { + qWarning("Failed to create output OpenGL image object from OpenGL texture: %d", err); + return *input; + } + } + + // We are all set. Queue acquiring the image objects. + f->glFinish(); + clEnqueueAcquireGLObjects(m_clQueue, 2, m_clImage, 0, 0, 0); + + // Set up the kernel arguments. + clSetKernelArg(m_clKernel, 0, sizeof(cl_mem), &m_clImage[0]); + clSetKernelArg(m_clKernel, 1, sizeof(cl_mem), &m_clImage[1]); + // Accessing dynamic properties on the filter element is simple: + cl_float factor = m_filter->factor(); + clSetKernelArg(m_clKernel, 2, sizeof(cl_float), &factor); + + // And queue the kernel. + const size_t workSize[] = { size_t(m_size.width()), size_t(m_size.height()) }; + err = clEnqueueNDRangeKernel(m_clQueue, m_clKernel, 2, 0, workSize, 0, 0, 0, 0); + if (err != CL_SUCCESS) + qWarning("Failed to enqueue kernel: %d", err); + + // Return the texture from our output image object. + // We return a texture even when the original video frame had pixel data in system memory. + // Qt Multimedia is smart enough to handle this. Once the data is on the GPU, it stays there. No readbacks, no copies. + clEnqueueReleaseGLObjects(m_clQueue, 2, m_clImage, 0, 0, 0); + clFinish(m_clQueue); + return frameFromTexture(m_outTexture, m_size, input->pixelFormat()); +} + +// InfoFilter will just provide some information about the video frame, to demonstrate +// passing arbitrary data to QML via its finished() signal. +class InfoFilter : public QAbstractVideoFilter +{ + Q_OBJECT + +public: + QVideoFilterRunnable *createFilterRunnable() Q_DECL_OVERRIDE; + +signals: + void finished(QObject *result); + +private: + friend class InfoFilterRunnable; +}; + +class InfoFilterRunnable : public QVideoFilterRunnable +{ +public: + InfoFilterRunnable(InfoFilter *filter) : m_filter(filter) { } + QVideoFrame run(QVideoFrame *input, const QVideoSurfaceFormat &surfaceFormat, RunFlags flags) Q_DECL_OVERRIDE; + +private: + InfoFilter *m_filter; +}; + +class InfoFilterResult : public QObject +{ + Q_OBJECT + Q_PROPERTY(QSize frameResolution READ frameResolution) + Q_PROPERTY(QString handleType READ handleType) + Q_PROPERTY(int pixelFormat READ pixelFormat) + +public: + InfoFilterResult() : m_pixelFormat(0) { } + QSize frameResolution() const { return m_frameResolution; } + QString handleType() const { return m_handleType; } + int pixelFormat() const { return m_pixelFormat; } + +private: + QSize m_frameResolution; + QString m_handleType; + int m_pixelFormat; + friend class InfoFilterRunnable; +}; + +void CLFilter::setFactor(qreal v) +{ + if (m_factor != v) { + m_factor = v; + emit factorChanged(); + } +} + +QVideoFilterRunnable *InfoFilter::createFilterRunnable() +{ + return new InfoFilterRunnable(this); +} + +QVideoFrame InfoFilterRunnable::run(QVideoFrame *input, const QVideoSurfaceFormat &surfaceFormat, RunFlags flags) +{ + Q_UNUSED(surfaceFormat); + Q_UNUSED(flags); + InfoFilterResult *result = new InfoFilterResult; + result->m_frameResolution = input->size(); + switch (input->handleType()) { + case QAbstractVideoBuffer::NoHandle: + result->m_handleType = QLatin1String("pixel data"); + result->m_pixelFormat = input->pixelFormat(); + break; + case QAbstractVideoBuffer::GLTextureHandle: + result->m_handleType = QLatin1String("OpenGL texture"); + break; + default: + result->m_handleType = QLatin1String("unknown"); + break; + } + emit m_filter->finished(result); // parent-less QObject -> ownership transferred to the JS engine + return *input; +} + +int main(int argc, char **argv) +{ + QGuiApplication app(argc, argv); + + qmlRegisterType<CLFilter>("qmlvideofilter.cl.test", 1, 0, "CLFilter"); + qmlRegisterType<InfoFilter>("qmlvideofilter.cl.test", 1, 0, "InfoFilter"); + + QQuickView view; + QString fn; + if (argc > 1) { + fn = QUrl::fromLocalFile(QFileInfo(QString::fromUtf8(argv[1])).absoluteFilePath()).toString(); + qDebug("Playing video %s", qPrintable(fn)); + } else { + qDebug("No video file specified, using camera instead."); + } + view.rootContext()->setContextProperty("videoFilename", fn); + view.setSource(QUrl("qrc:///main.qml")); + + view.show(); + + return app.exec(); +} + +#include "main.moc" diff --git a/examples/multimedia/video/qmlvideofilter_opencl/main.qml b/examples/multimedia/video/qmlvideofilter_opencl/main.qml new file mode 100644 index 000000000..636df568c --- /dev/null +++ b/examples/multimedia/video/qmlvideofilter_opencl/main.qml @@ -0,0 +1,114 @@ +/**************************************************************************** +** +** Copyright (C) 2015 Digia Plc and/or its subsidiary(-ies). +** Contact: http://www.qt-project.org/legal +** +** This file is part of the examples of the Qt Multimedia module. +** +** $QT_BEGIN_LICENSE:LGPL21$ +** Commercial License Usage +** Licensees holding valid commercial Qt licenses may use this file in +** accordance with the commercial license agreement provided with the +** Software or, alternatively, in accordance with the terms contained in +** a written agreement between you and Digia. For licensing terms and +** conditions see http://qt.digia.com/licensing. For further information +** use the contact form at http://qt.digia.com/contact-us. +** +** GNU Lesser General Public License Usage +** Alternatively, this file may be used under the terms of the GNU Lesser +** General Public License version 2.1 or version 3 as published by the Free +** Software Foundation and appearing in the file LICENSE.LGPLv21 and +** LICENSE.LGPLv3 included in the packaging of this file. Please review the +** following information to ensure the GNU Lesser General Public License +** requirements will be met: https://www.gnu.org/licenses/lgpl.html and +** http://www.gnu.org/licenses/old-licenses/lgpl-2.1.html. +** +** In addition, as a special exception, Digia gives you certain additional +** rights. These rights are described in the Digia Qt LGPL Exception +** version 1.1, included in the file LGPL_EXCEPTION.txt in this package. +** +** $QT_END_LICENSE$ +** +****************************************************************************/ + +import QtQuick 2.0 +import QtMultimedia 5.5 +import qmlvideofilter.cl.test 1.0 + +Item { + width: 1024 + height: 768 + + Camera { + id: camera + } + + MediaPlayer { + id: player + autoPlay: true + source: videoFilename + } + + VideoOutput { + id: output + source: videoFilename !== "" ? player : camera + filters: [ infofilter, clfilter ] + anchors.fill: parent + } + + CLFilter { + id: clfilter + // Animate a property which is passed to the OpenCL kernel. + SequentialAnimation on factor { + loops: Animation.Infinite + NumberAnimation { + from: 1 + to: 20 + duration: 6000 + } + NumberAnimation { + from: 20 + to: 1 + duration: 3000 + } + } + } + + InfoFilter { + // This filter does not change the image. Instead, it provides some results calculated from the frame. + id: infofilter + onFinished: { + info.res = result.frameResolution.width + "x" + result.frameResolution.height; + info.type = result.handleType; + info.fmt = result.pixelFormat; + } + } + + Column { + Text { + font.pointSize: 20 + color: "green" + text: "Transformed with OpenCL on GPU\nClick to disable and enable the emboss filter" + } + Text { + font.pointSize: 12 + color: "green" + text: "Emboss factor " + Math.round(clfilter.factor) + visible: clfilter.active + } + Text { + id: info + font.pointSize: 12 + color: "green" + property string res + property string type + property int fmt + text: "Input resolution: " + res + " Input frame type: " + type + (fmt ? " Pixel format: " + fmt : "") + } + } + + MouseArea { + anchors.fill: parent + onClicked: clfilter.active = !clfilter.active + } +} diff --git a/examples/multimedia/video/qmlvideofilter_opencl/qmlvideofilter_opencl.pro b/examples/multimedia/video/qmlvideofilter_opencl/qmlvideofilter_opencl.pro new file mode 100644 index 000000000..b391bb8d7 --- /dev/null +++ b/examples/multimedia/video/qmlvideofilter_opencl/qmlvideofilter_opencl.pro @@ -0,0 +1,22 @@ +TEMPLATE = app +TARGET = qmlvideofilter_opencl + +QT += quick multimedia + +SOURCES = main.cpp +HEADERS = rgbframehelper.h + +RESOURCES = qmlvideofilter_opencl.qrc +OTHER_FILES = main.qml + +target.path = $$[QT_INSTALL_EXAMPLES]/multimedia/video/qmlvideofilter_opencl +INSTALLS += target + +# Edit these as necessary +osx { + LIBS += -framework OpenCL +} else { + INCLUDEPATH += c:/cuda/include + LIBPATH += c:/cuda/lib/x64 + LIBS += -lopengl32 -lOpenCL +} diff --git a/examples/multimedia/video/qmlvideofilter_opencl/qmlvideofilter_opencl.qrc b/examples/multimedia/video/qmlvideofilter_opencl/qmlvideofilter_opencl.qrc new file mode 100644 index 000000000..5f6483ac3 --- /dev/null +++ b/examples/multimedia/video/qmlvideofilter_opencl/qmlvideofilter_opencl.qrc @@ -0,0 +1,5 @@ +<RCC> + <qresource prefix="/"> + <file>main.qml</file> + </qresource> +</RCC> diff --git a/examples/multimedia/video/qmlvideofilter_opencl/rgbframehelper.h b/examples/multimedia/video/qmlvideofilter_opencl/rgbframehelper.h new file mode 100644 index 000000000..080c3fe2b --- /dev/null +++ b/examples/multimedia/video/qmlvideofilter_opencl/rgbframehelper.h @@ -0,0 +1,124 @@ +/**************************************************************************** +** +** Copyright (C) 2015 Digia Plc and/or its subsidiary(-ies). +** Contact: http://www.qt-project.org/legal +** +** This file is part of the examples of the Qt Multimedia module. +** +** $QT_BEGIN_LICENSE:LGPL21$ +** Commercial License Usage +** Licensees holding valid commercial Qt licenses may use this file in +** accordance with the commercial license agreement provided with the +** Software or, alternatively, in accordance with the terms contained in +** a written agreement between you and Digia. For licensing terms and +** conditions see http://qt.digia.com/licensing. For further information +** use the contact form at http://qt.digia.com/contact-us. +** +** GNU Lesser General Public License Usage +** Alternatively, this file may be used under the terms of the GNU Lesser +** General Public License version 2.1 or version 3 as published by the Free +** Software Foundation and appearing in the file LICENSE.LGPLv21 and +** LICENSE.LGPLv3 included in the packaging of this file. Please review the +** following information to ensure the GNU Lesser General Public License +** requirements will be met: https://www.gnu.org/licenses/lgpl.html and +** http://www.gnu.org/licenses/old-licenses/lgpl-2.1.html. +** +** In addition, as a special exception, Digia gives you certain additional +** rights. These rights are described in the Digia Qt LGPL Exception +** version 1.1, included in the file LGPL_EXCEPTION.txt in this package. +** +** $QT_END_LICENSE$ +** +****************************************************************************/ + +#ifndef RGBFRAMEHELPER_H +#define RGBFRAMEHELPER_H + +#include <QImage> +#include <QAbstractVideoBuffer> +#include <QOpenGLContext> +#include <QOpenGLFunctions> +#include <QOpenGLFramebufferObject> + +/* + Returns a QImage that wraps the given video frame. + + This is suitable only for QAbstractVideoBuffer::NoHandle frames with RGB (or BGR) + data. YUV is not supported here. + + The QVideoFrame must be mapped and kept mapped as long as the wrapping QImage + exists. + + As a convenience the function also supports frames with a handle type of + QAbstractVideoBuffer::GLTextureHandle. This allows creating a system memory backed + QVideoFrame containing the image data from an OpenGL texture. However, readback is a + slow operation and may stall the GPU pipeline and should be avoided in production code. +*/ +QImage imageWrapper(const QVideoFrame &frame) +{ +#ifndef QT_NO_OPENGL + if (frame.handleType() == QAbstractVideoBuffer::GLTextureHandle) { + // Slow and inefficient path. Ideally what's on the GPU should remain on the GPU, instead of readbacks like this. + QImage img(frame.width(), frame.height(), QImage::Format_RGBA8888); + GLuint textureId = frame.handle().toUInt(); + QOpenGLContext *ctx = QOpenGLContext::currentContext(); + QOpenGLFunctions *f = ctx->functions(); + GLuint fbo; + f->glGenFramebuffers(1, &fbo); + GLuint prevFbo; + f->glGetIntegerv(GL_FRAMEBUFFER_BINDING, (GLint *) &prevFbo); + f->glBindFramebuffer(GL_FRAMEBUFFER, fbo); + f->glFramebufferTexture2D(GL_FRAMEBUFFER, GL_COLOR_ATTACHMENT0, GL_TEXTURE_2D, textureId, 0); + f->glReadPixels(0, 0, frame.width(), frame.height(), GL_RGBA, GL_UNSIGNED_BYTE, img.bits()); + f->glBindFramebuffer(GL_FRAMEBUFFER, prevFbo); + return img; + } else +#endif // QT_NO_OPENGL + { + if (!frame.isReadable()) { + qWarning("imageFromVideoFrame: No mapped image data available for read"); + return QImage(); + } + + QImage::Format fmt = QVideoFrame::imageFormatFromPixelFormat(frame.pixelFormat()); + if (fmt != QImage::Format_Invalid) + return QImage(frame.bits(), frame.width(), frame.height(), fmt); + + qWarning("imageFromVideoFrame: No matching QImage format"); + } + + return QImage(); +} + +#ifndef QT_NO_OPENGL +class TextureBuffer : public QAbstractVideoBuffer +{ +public: + TextureBuffer(uint id) : QAbstractVideoBuffer(GLTextureHandle), m_id(id) { } + MapMode mapMode() const { return NotMapped; } + uchar *map(MapMode, int *, int *) { return 0; } + void unmap() { } + QVariant handle() const { return QVariant::fromValue<uint>(m_id); } + +private: + GLuint m_id; +}; +#endif // QT_NO_OPENGL + +/* + Creates and returns a new video frame wrapping the OpenGL texture textureId. The size + must be passed in size, together with the format of the underlying image data in + format. When the texture originates from a QImage, use + QVideoFrame::imageFormatFromPixelFormat() to get a suitable format. Ownership is not + altered, the new QVideoFrame will not destroy the texture. +*/ +QVideoFrame frameFromTexture(uint textureId, const QSize &size, QVideoFrame::PixelFormat format) +{ +#ifndef QT_NO_OPENGL + return QVideoFrame(new TextureBuffer(textureId), size, format); +#else + return QVideoFrame(); +#endif // QT_NO_OPENGL +} + +#endif diff --git a/src/imports/multimedia/multimedia.cpp b/src/imports/multimedia/multimedia.cpp index 8b2287ce4..de0290756 100644 --- a/src/imports/multimedia/multimedia.cpp +++ b/src/imports/multimedia/multimedia.cpp @@ -38,6 +38,7 @@ #include "qsoundeffect.h" #include <private/qdeclarativevideooutput_p.h> +#include "qabstractvideofilter.h" #include "qdeclarativemultimediaglobal_p.h" #include "qdeclarativemediametadata_p.h" @@ -108,6 +109,7 @@ public: qmlRegisterRevision<QDeclarativeCamera, 2>(uri, 5, 5); qmlRegisterType<QDeclarativeMediaMetaData>(); + qmlRegisterType<QAbstractVideoFilter>(); } void initializeEngine(QQmlEngine *engine, const char *uri) diff --git a/src/imports/multimedia/plugins.qmltypes b/src/imports/multimedia/plugins.qmltypes index 12d7ea85d..0cc3a6a24 100644 --- a/src/imports/multimedia/plugins.qmltypes +++ b/src/imports/multimedia/plugins.qmltypes @@ -1225,6 +1225,7 @@ Module { Property { name: "orientation"; type: "int" } Property { name: "sourceRect"; type: "QRectF"; isReadonly: true } Property { name: "contentRect"; type: "QRectF"; isReadonly: true } + Property { name: "filters"; isList: true; isReadonly: true } Signal { name: "fillModeChanged" Parameter { type: "QDeclarativeVideoOutput::FillMode" } diff --git a/src/multimedia/qtmultimediaquicktools_headers/qdeclarativevideooutput_backend_p.h b/src/multimedia/qtmultimediaquicktools_headers/qdeclarativevideooutput_backend_p.h index 5996af773..1d267df82 100644 --- a/src/multimedia/qtmultimediaquicktools_headers/qdeclarativevideooutput_backend_p.h +++ b/src/multimedia/qtmultimediaquicktools_headers/qdeclarativevideooutput_backend_p.h @@ -46,6 +46,7 @@ QT_BEGIN_NAMESPACE class QAbstractVideoSurface; class QDeclarativeVideoOutput; class QMediaService; +class QAbstractVideoFilter; class Q_MULTIMEDIAQUICK_EXPORT QDeclarativeVideoBackend { @@ -70,6 +71,12 @@ public: // The viewport, adjusted for the pixel aspect ratio virtual QRectF adjustedViewport() const = 0; + virtual void appendFilter(QAbstractVideoFilter *filter) { Q_UNUSED(filter); } + virtual void clearFilters() { } + + virtual void releaseResources() { } + virtual void invalidateSceneGraph() { } + protected: QDeclarativeVideoOutput *q; QPointer<QMediaService> m_service; diff --git a/src/multimedia/qtmultimediaquicktools_headers/qdeclarativevideooutput_p.h b/src/multimedia/qtmultimediaquicktools_headers/qdeclarativevideooutput_p.h index e1dbf5646..f022c0e3e 100644 --- a/src/multimedia/qtmultimediaquicktools_headers/qdeclarativevideooutput_p.h +++ b/src/multimedia/qtmultimediaquicktools_headers/qdeclarativevideooutput_p.h @@ -40,6 +40,7 @@ #include <QtQuick/qquickitem.h> #include <QtCore/qpointer.h> #include <QtMultimedia/qcamerainfo.h> +#include <QtMultimedia/qabstractvideofilter.h> #include <private/qtmultimediaquickdefs_p.h> @@ -60,6 +61,7 @@ class Q_MULTIMEDIAQUICK_EXPORT QDeclarativeVideoOutput : public QQuickItem Q_PROPERTY(bool autoOrientation READ autoOrientation WRITE setAutoOrientation NOTIFY autoOrientationChanged REVISION 2) Q_PROPERTY(QRectF sourceRect READ sourceRect NOTIFY sourceRectChanged) Q_PROPERTY(QRectF contentRect READ contentRect NOTIFY contentRectChanged) + Q_PROPERTY(QQmlListProperty<QAbstractVideoFilter> filters READ filters); Q_ENUMS(FillMode) public: @@ -104,6 +106,8 @@ public: }; SourceType sourceType() const; + QQmlListProperty<QAbstractVideoFilter> filters(); + Q_SIGNALS: void sourceChanged(); void fillModeChanged(QDeclarativeVideoOutput::FillMode); @@ -116,6 +120,7 @@ protected: QSGNode *updatePaintNode(QSGNode *, UpdatePaintNodeData *); void itemChange(ItemChange change, const ItemChangeData &changeData); void geometryChanged(const QRectF &newGeometry, const QRectF &oldGeometry); + void releaseResources(); private Q_SLOTS: void _q_updateMediaObject(); @@ -123,10 +128,16 @@ private Q_SLOTS: void _q_updateNativeSize(); void _q_updateGeometry(); void _q_screenOrientationChanged(int); + void _q_invalidateSceneGraph(); private: bool createBackend(QMediaService *service); + static void filter_append(QQmlListProperty<QAbstractVideoFilter> *property, QAbstractVideoFilter *value); + static int filter_count(QQmlListProperty<QAbstractVideoFilter> *property); + static QAbstractVideoFilter *filter_at(QQmlListProperty<QAbstractVideoFilter> *property, int index); + static void filter_clear(QQmlListProperty<QAbstractVideoFilter> *property); + SourceType m_sourceType; QPointer<QObject> m_source; @@ -145,6 +156,8 @@ private: QVideoOutputOrientationHandler *m_screenOrientationHandler; QScopedPointer<QDeclarativeVideoBackend> m_backend; + + QList<QAbstractVideoFilter *> m_filters; }; QT_END_NAMESPACE diff --git a/src/multimedia/qtmultimediaquicktools_headers/qsgvideonode_p.h b/src/multimedia/qtmultimediaquicktools_headers/qsgvideonode_p.h index 8be77ff07..9a15f4afa 100644 --- a/src/multimedia/qtmultimediaquicktools_headers/qsgvideonode_p.h +++ b/src/multimedia/qtmultimediaquicktools_headers/qsgvideonode_p.h @@ -46,10 +46,16 @@ QT_BEGIN_NAMESPACE class Q_MULTIMEDIAQUICK_EXPORT QSGVideoNode : public QSGGeometryNode { public: + enum FrameFlag { + FrameFiltered = 0x01 + }; + Q_DECLARE_FLAGS(FrameFlags, FrameFlag) + QSGVideoNode(); - virtual void setCurrentFrame(const QVideoFrame &frame) = 0; + virtual void setCurrentFrame(const QVideoFrame &frame, FrameFlags flags) = 0; virtual QVideoFrame::PixelFormat pixelFormat() const = 0; + virtual QAbstractVideoBuffer::HandleType handleType() const = 0; void setTexturedRectGeometry(const QRectF &boundingRect, const QRectF &textureRect, int orientation); @@ -59,6 +65,8 @@ private: int m_orientation; }; +Q_DECLARE_OPERATORS_FOR_FLAGS(QSGVideoNode::FrameFlags) + class Q_MULTIMEDIAQUICK_EXPORT QSGVideoNodeFactoryInterface { public: diff --git a/src/multimedia/video/qabstractvideofilter.cpp b/src/multimedia/video/qabstractvideofilter.cpp new file mode 100644 index 000000000..e04cbf1df --- /dev/null +++ b/src/multimedia/video/qabstractvideofilter.cpp @@ -0,0 +1,319 @@ +/**************************************************************************** +** +** Copyright (C) 2015 Digia Plc and/or its subsidiary(-ies). +** Contact: http://www.qt-project.org/legal +** +** This file is part of the Qt Toolkit. +** +** $QT_BEGIN_LICENSE:LGPL21$ +** Commercial License Usage +** Licensees holding valid commercial Qt licenses may use this file in +** accordance with the commercial license agreement provided with the +** Software or, alternatively, in accordance with the terms contained in +** a written agreement between you and Digia. For licensing terms and +** conditions see http://qt.digia.com/licensing. For further information +** use the contact form at http://qt.digia.com/contact-us. +** +** GNU Lesser General Public License Usage +** Alternatively, this file may be used under the terms of the GNU Lesser +** General Public License version 2.1 or version 3 as published by the Free +** Software Foundation and appearing in the file LICENSE.LGPLv21 and +** LICENSE.LGPLv3 included in the packaging of this file. Please review the +** following information to ensure the GNU Lesser General Public License +** requirements will be met: https://www.gnu.org/licenses/lgpl.html and +** http://www.gnu.org/licenses/old-licenses/lgpl-2.1.html. +** +** In addition, as a special exception, Digia gives you certain additional +** rights. These rights are described in the Digia Qt LGPL Exception +** version 1.1, included in the file LGPL_EXCEPTION.txt in this package. +** +** $QT_END_LICENSE$ +** +****************************************************************************/ + +#include "qabstractvideofilter.h" + +QT_BEGIN_NAMESPACE + +/*! + \class QAbstractVideoFilter + \since 5.5 + \brief The QAbstractVideoFilter class represents a filter that is applied to the video frames + received by a VideoOutput type. + \inmodule QtMultimedia + + \ingroup multimedia + \ingroup multimedia_video + + QAbstractVideoFilter provides a convenient way for applications to run image + processing, computer vision algorithms or any generic transformation or + calculation on the output of a VideoOutput type, regardless of the source + (video or camera). By providing a simple interface it allows applications and + third parties to easily develop QML types that provide image processing + algorithms using popular frameworks like \l{http://opencv.org}{OpenCV}. Due to + the close integration with the final stages of the Qt Multimedia video + pipeline, accelerated and possibly zero-copy solutions are feasible too: for + instance, a plugin providing OpenCL-based algorithms can use OpenCL's OpenGL + interop to use the OpenGL textures created by a hardware accelerated video + decoder, without additional readbacks and copies. + + \note QAbstractVideoFilter is not always the best choice. To apply effects or + transformations using OpenGL shaders to the image shown on screen, the + standard Qt Quick approach of using ShaderEffect items in combination with + VideoOutput should be used. VideoFilter is not a replacement for this. It is + rather targeted for performing computations (that do not necessarily change + the image shown on screen) and computer vision algorithms provided by + external frameworks. + + QAbstractVideoFilter is meant to be subclassed. The subclasses are then registered to + the QML engine, so they can be used as a QML type. The list of filters are + assigned to a VideoOutput type via its \l{QtMultimedia::VideoOutput::filters}{filters} + property. + + A single filter represents one transformation or processing step on + a video frame. The output is a modified video frame, some arbitrary data or + both. For example, image transformations will result in a different image, + whereas an algorithm for detecting objects on an image will likely provide + a list of rectangles. + + Arbitrary data can be represented as properties on the QAbstractVideoFilter subclass + and on the QObject or QJSValue instances passed to its signals. What exactly + these properties and signals are, is up to the individual video + filters. Completion of the operations can be indicated by + signals. Computations that do not result in a modified image will pass the + input image through so that subsequent filters can be placed after them. + + Properties set on QAbstractVideoFilter serve as input to the computation, similarly + to how uniform values are specified in ShaderEffect types. The changed + property values are taken into use when the next video frame is processed. + + The typical usage is to subclass QAbstractVideoFilter and QVideoFilterRunnable: + + \badcode + class MyFilterRunnable : public QVideoFilterRunnable { + public: + QVideoFrame run(QVideoFrame *input, const QVideoSurfaceFormat &surfaceFormat, RunFlags flags) { ... } + }; + + class MyFilter : public QAbstractVideoFilter { + public: + QVideoFilterRunnable *createFilterRunnable() { return new MyFilterRunnable; } + signals: + void finished(QObject *result); + }; + + int main(int argc, char **argv) { + ... + qmlRegisterType<MyFilter>("my.uri", 1, 0, "MyFilter"); + ... + } + \endcode + + MyFilter is thus accessible from QML: + + \badcode + import my.uri 1.0 + + Camera { + id: camera + } + MyFilter { + id: filter + // set properties, they can also be animated + onFinished: console.log("results of the computation: " + result) + } + VideoOutput { + source: camera + filters: [ filter ] + anchors.fill: parent + } + \endcode + + This also allows providing filters in QML plugins, separately from the application. + + \sa VideoOutput, Camera, MediaPlayer, QVideoFilterRunnable +*/ + +/*! + \class QVideoFilterRunnable + \since 5.5 + \brief The QVideoFilterRunnable class represents the implementation of a filter + that owns all graphics and computational resources, and performs the actual filtering + or calculations. + \inmodule QtMultimedia + + \ingroup multimedia + \ingroup multimedia_video + + Video filters are split into QAbstractVideoFilter and corresponding QVideoFilterRunnable + instances, similar to QQuickItem and QSGNode. This is necessary to support + threaded rendering scenarios. When using the threaded render loop of the Qt + Quick scene graph, all rendering happens on a dedicated thread. + QVideoFilterRunnable instances always live on this thread and all its functions, + run(), the constructor, and the destructor, are guaranteed to be invoked on + that thread with the OpenGL context bound. QAbstractVideoFilter instances live on + the main (GUI) thread, like any other QObject and QQuickItem instances + created from QML. + + Once created, QVideoFilterRunnable instances are managed by Qt Multimedia and + will be automatically destroyed and recreated when necessary, for example + when the scene graph is invalidated or the QQuickWindow changes or is closed. + Creation happens via the QAbstractVideoFilter::createFilterRunnable() factory function. + + \sa QAbstractVideoFilter + */ + +/*! + \fn QVideoFrame QVideoFilterRunnable::run(QVideoFrame *input, const QVideoSurfaceFormat &surfaceFormat, RunFlags flags) + + Reimplement this function to perform filtering or computation on the \a + input video frame. Like the constructor and destructor, this function is + always called on the render thread with the OpenGL context bound. + + Implementations that do not modify the video frame can simply return \a input. + + It is safe to access properties of the associated QAbstractVideoFilter instance from + this function. + + \a input will not be mapped, it is up to this function to call QVideoFrame::map() + and QVideoFrame::unmap() as necessary. + + \a surfaceFormat provides additional information, for example it can be used + to determine which way is up in the input image as that is important for + filters to operate on multiple platforms with multiple cameras. + + \a flags contains additional information about the filter's invocation. For + example the LastInChain flag indicates that the filter is the last in a + VideoOutput's associated filter list. This can be very useful in cases where + multiple filters are chained together and the work is performed on image data + in some custom format (for example a format specific to some computer vision + framework). To avoid conversion on every filter in the chain, all + intermediate filters can return a QVideoFrame hosting data in the custom + format. Only the last, where the flag is set, returns a QVideoFrame in a + format compatible with Qt. + + Filters that want to expose the results of their computation to Javascript + code in QML can declare their own custom signals in the QAbstractVideoFilter + subclass to indicate the completion of the operation. For filters that only + calculate some results and do not modify the video frame, it is also possible + to operate asynchronously. They can queue the necessary operations using the + compute API and return from this function without emitting any signals. The + signal indicating the completion is then emitted only when the compute API + indicates that the operations were done and the results are available. Note + that it is strongly recommended to represent the filter's output data as a + separate instance of QJSValue or a QObject-derived class which is passed as a + parameter to the signal and becomes exposed to the Javascript engine. In case + of QObject the ownership of this object is controlled by the standard QML + rules: if it has no parent, ownership is transferred to the Javascript engine, + otherwise it stays with the emitter. Note that the signal connection may be + queued,for example when using the threaded render loop of Qt Quick, and so the + object must stay valid for a longer time, destroying it right after calling + this function is not safe. Using a dedicated results object is guaranteed to + be safe even when using threaded rendering. The same is not necessarily true + for properties on the QAbstractVideoFilter instance itself: properties can + safely be read in run() since the gui thread is blocked during that time but + writing may become problematic. + + \note Avoid time consuming operations in this function as they block the + entire rendering of the application. + + \note The handleType() and pixelFormat() of \a input is completely up to the + video decoding backend on the platform in use. On some platforms different + forms of input are used depending on the graphics stack. For example, when + playing back videos on Windows with the WMF backend, QVideoFrame contains + OpenGL-wrapped Direct3D textures in case of using ANGLE, but regular pixel + data when using desktop OpenGL (opengl32.dll). Similarly, the video file + format will often decide if the data is RGB or YUV, but this may also depend + on the decoder and the configuration in use. The returned video frame does + not have to be in the same format as the input, for example a filter with an + input of a QVideoFrame backed by system memory can output a QVideoFrame with + an OpenGL texture handle. + + \sa QVideoFrame, QVideoSurfaceFormat + */ + +/*! + \enum QVideoFilterRunnable::RunFlag + + \value LastInChain Indicates that the filter runnable's associated QAbstractVideoFilter + is the last in the corresponding VideoOutput type's filters list, meaning + that the returned frame is the one that is going to be presented to the scene + graph without invoking any further filters. + */ + +class QAbstractVideoFilterPrivate +{ +public: + QAbstractVideoFilterPrivate() : + active(true) + { } + + bool active; +}; + +/*! + \internal + */ +QVideoFilterRunnable::~QVideoFilterRunnable() +{ +} + +/*! + Constructs a new QAbstractVideoFilter instance. + */ +QAbstractVideoFilter::QAbstractVideoFilter(QObject *parent) : + QObject(parent), + d_ptr(new QAbstractVideoFilterPrivate) +{ +} + +/*! + \internal + */ +QAbstractVideoFilter::~QAbstractVideoFilter() +{ + delete d_ptr; +} + +/*! + \return \c true if the filter is active. + + By default filters are active. When set to \c false, the filter will be + ignored by the VideoOutput type. + */ +bool QAbstractVideoFilter::isActive() const +{ + Q_D(const QAbstractVideoFilter); + return d->active; +} + +/*! + \internal + */ +void QAbstractVideoFilter::setActive(bool v) +{ + Q_D(QAbstractVideoFilter); + if (d->active != v) { + d->active = v; + emit activeChanged(); + } +} + +/*! + \fn QVideoFilterRunnable *QAbstractVideoFilter::createFilterRunnable() + + Factory function to create a new instance of a QVideoFilterRunnable subclass + corresponding to this filter. + + This function is called on the thread on which the Qt Quick scene graph + performs rendering, with the OpenGL context bound. Ownership of the returned + instance is transferred: the returned instance will live on the render thread + and will be destroyed automatically when necessary. + + Typically, implementations of the function will simply construct a new + QVideoFilterRunnable instance, passing \c this to the constructor as the + filter runnables must know their associated QAbstractVideoFilter instance to + access dynamic properties and optionally emit signals. + */ + +QT_END_NAMESPACE diff --git a/src/multimedia/video/qabstractvideofilter.h b/src/multimedia/video/qabstractvideofilter.h new file mode 100644 index 000000000..01e39fb8a --- /dev/null +++ b/src/multimedia/video/qabstractvideofilter.h @@ -0,0 +1,85 @@ +/**************************************************************************** +** +** Copyright (C) 2015 Digia Plc and/or its subsidiary(-ies). +** Contact: http://www.qt-project.org/legal +** +** This file is part of the Qt Toolkit. +** +** $QT_BEGIN_LICENSE:LGPL21$ +** Commercial License Usage +** Licensees holding valid commercial Qt licenses may use this file in +** accordance with the commercial license agreement provided with the +** Software or, alternatively, in accordance with the terms contained in +** a written agreement between you and Digia. For licensing terms and +** conditions see http://qt.digia.com/licensing. For further information +** use the contact form at http://qt.digia.com/contact-us. +** +** GNU Lesser General Public License Usage +** Alternatively, this file may be used under the terms of the GNU Lesser +** General Public License version 2.1 or version 3 as published by the Free +** Software Foundation and appearing in the file LICENSE.LGPLv21 and +** LICENSE.LGPLv3 included in the packaging of this file. Please review the +** following information to ensure the GNU Lesser General Public License +** requirements will be met: https://www.gnu.org/licenses/lgpl.html and +** http://www.gnu.org/licenses/old-licenses/lgpl-2.1.html. +** +** In addition, as a special exception, Digia gives you certain additional +** rights. These rights are described in the Digia Qt LGPL Exception +** version 1.1, included in the file LGPL_EXCEPTION.txt in this package. +** +** $QT_END_LICENSE$ +** +****************************************************************************/ + +#ifndef QABSTRACTVIDEOFILTER_H +#define QABSTRACTVIDEOFILTER_H + +#include <QtCore/qobject.h> +#include <QtMultimedia/qvideoframe.h> +#include <QtMultimedia/qvideosurfaceformat.h> + +QT_BEGIN_NAMESPACE + +class QAbstractVideoFilterPrivate; + +class Q_MULTIMEDIA_EXPORT QVideoFilterRunnable +{ +public: + enum RunFlag { + LastInChain = 0x01 + }; + Q_DECLARE_FLAGS(RunFlags, RunFlag) + + virtual ~QVideoFilterRunnable(); + virtual QVideoFrame run(QVideoFrame *input, const QVideoSurfaceFormat &surfaceFormat, RunFlags flags) = 0; +}; + +Q_DECLARE_OPERATORS_FOR_FLAGS(QVideoFilterRunnable::RunFlags) + +class Q_MULTIMEDIA_EXPORT QAbstractVideoFilter : public QObject +{ + Q_OBJECT + Q_PROPERTY(bool active READ isActive WRITE setActive NOTIFY activeChanged) + +public: + QAbstractVideoFilter(QObject *parent = 0); + ~QAbstractVideoFilter(); + + bool isActive() const; + void setActive(bool v); + + virtual QVideoFilterRunnable *createFilterRunnable() = 0; + +Q_SIGNALS: + void activeChanged(); + +private: + Q_DECLARE_PRIVATE(QAbstractVideoFilter) + Q_DISABLE_COPY(QAbstractVideoFilter) + + QAbstractVideoFilterPrivate *d_ptr; +}; + +QT_END_NAMESPACE + +#endif // QABSTRACTVIDEOFILTER_H diff --git a/src/multimedia/video/qvideoframe.cpp b/src/multimedia/video/qvideoframe.cpp index 8dd23d2f5..869972a5c 100644 --- a/src/multimedia/video/qvideoframe.cpp +++ b/src/multimedia/video/qvideoframe.cpp @@ -369,6 +369,23 @@ QVideoFrame &QVideoFrame::operator =(const QVideoFrame &other) } /*! + \return \c true if this QVideoFrame and \a other reflect the same frame. + */ +bool QVideoFrame::operator==(const QVideoFrame &other) const +{ + // Due to explicit sharing we just compare the QSharedData which in turn compares the pointers. + return d == other.d; +} + +/*! + \return \c true if this QVideoFrame and \a other do not reflect the same frame. + */ +bool QVideoFrame::operator!=(const QVideoFrame &other) const +{ + return d != other.d; +} + +/*! Destroys a video frame. */ QVideoFrame::~QVideoFrame() diff --git a/src/multimedia/video/qvideoframe.h b/src/multimedia/video/qvideoframe.h index d44219d36..89cf2fe32 100644 --- a/src/multimedia/video/qvideoframe.h +++ b/src/multimedia/video/qvideoframe.h @@ -107,6 +107,8 @@ public: ~QVideoFrame(); QVideoFrame &operator =(const QVideoFrame &other); + bool operator==(const QVideoFrame &other) const; + bool operator!=(const QVideoFrame &other) const; bool isValid() const; diff --git a/src/multimedia/video/video.pri b/src/multimedia/video/video.pri index 161163783..3e39d9d18 100644 --- a/src/multimedia/video/video.pri +++ b/src/multimedia/video/video.pri @@ -6,7 +6,8 @@ PUBLIC_HEADERS += \ video/qabstractvideosurface.h \ video/qvideoframe.h \ video/qvideosurfaceformat.h \ - video/qvideoprobe.h + video/qvideoprobe.h \ + video/qabstractvideofilter.h PRIVATE_HEADERS += \ video/qabstractvideobuffer_p.h \ @@ -24,8 +25,5 @@ SOURCES += \ video/qvideooutputorientationhandler.cpp \ video/qvideosurfaceformat.cpp \ video/qvideosurfaceoutput.cpp \ - video/qvideoprobe.cpp - - - - + video/qvideoprobe.cpp \ + video/qabstractvideofilter.cpp diff --git a/src/plugins/android/videonode/qandroidsgvideonode.cpp b/src/plugins/android/videonode/qandroidsgvideonode.cpp index 4e1b2a89b..f094c0d8d 100644 --- a/src/plugins/android/videonode/qandroidsgvideonode.cpp +++ b/src/plugins/android/videonode/qandroidsgvideonode.cpp @@ -165,7 +165,7 @@ QAndroidSGVideoNode::~QAndroidSGVideoNode() m_frame = QVideoFrame(); } -void QAndroidSGVideoNode::setCurrentFrame(const QVideoFrame &frame) +void QAndroidSGVideoNode::setCurrentFrame(const QVideoFrame &frame, FrameFlags) { QMutexLocker lock(&m_frameMutex); m_frame = frame; diff --git a/src/plugins/android/videonode/qandroidsgvideonode.h b/src/plugins/android/videonode/qandroidsgvideonode.h index f3f838331..0c50d8cf9 100644 --- a/src/plugins/android/videonode/qandroidsgvideonode.h +++ b/src/plugins/android/videonode/qandroidsgvideonode.h @@ -47,8 +47,9 @@ public: QAndroidSGVideoNode(const QVideoSurfaceFormat &format); ~QAndroidSGVideoNode(); - void setCurrentFrame(const QVideoFrame &frame); + void setCurrentFrame(const QVideoFrame &frame, FrameFlags flags); QVideoFrame::PixelFormat pixelFormat() const { return m_format.pixelFormat(); } + QAbstractVideoBuffer::HandleType handleType() const { return QAbstractVideoBuffer::GLTextureHandle; } void preprocess(); diff --git a/src/plugins/videonode/egl/qsgvideonode_egl.cpp b/src/plugins/videonode/egl/qsgvideonode_egl.cpp index 4e63c0dba..15af8b5e7 100644 --- a/src/plugins/videonode/egl/qsgvideonode_egl.cpp +++ b/src/plugins/videonode/egl/qsgvideonode_egl.cpp @@ -187,7 +187,7 @@ QSGVideoNode_EGL::~QSGVideoNode_EGL() { } -void QSGVideoNode_EGL::setCurrentFrame(const QVideoFrame &frame) +void QSGVideoNode_EGL::setCurrentFrame(const QVideoFrame &frame, FrameFlags) { EGLImageKHR image = frame.handle().value<void *>(); m_material.setImage(image); @@ -199,6 +199,11 @@ QVideoFrame::PixelFormat QSGVideoNode_EGL::pixelFormat() const return m_pixelFormat; } +QAbstractVideoBuffer::HandleType QSGVideoNode_EGL::handleType() const +{ + return QAbstractVideoBuffer::EGLImageHandle; +} + static bool isExtensionSupported() { static const bool supported = eglGetProcAddress("glEGLImageTargetTexture2DOES"); diff --git a/src/plugins/videonode/egl/qsgvideonode_egl.h b/src/plugins/videonode/egl/qsgvideonode_egl.h index b6f70abfb..7e9cfe871 100644 --- a/src/plugins/videonode/egl/qsgvideonode_egl.h +++ b/src/plugins/videonode/egl/qsgvideonode_egl.h @@ -74,8 +74,9 @@ public: QSGVideoNode_EGL(const QVideoSurfaceFormat &format); ~QSGVideoNode_EGL(); - void setCurrentFrame(const QVideoFrame &frame); + void setCurrentFrame(const QVideoFrame &frame, FrameFlags flags); QVideoFrame::PixelFormat pixelFormat() const; + QAbstractVideoBuffer::HandleType handleType() const; private: QSGVideoMaterial_EGL m_material; diff --git a/src/plugins/videonode/imx6/qsgvivantevideomaterial.cpp b/src/plugins/videonode/imx6/qsgvivantevideomaterial.cpp index 948e9f9c8..073ffb34e 100644 --- a/src/plugins/videonode/imx6/qsgvivantevideomaterial.cpp +++ b/src/plugins/videonode/imx6/qsgvivantevideomaterial.cpp @@ -51,7 +51,9 @@ QSGVivanteVideoMaterial::QSGVivanteVideoMaterial() : mWidth(0), mHeight(0), mFormat(QVideoFrame::Format_Invalid), - mCurrentTexture(0) + mCurrentTexture(0), + mMappable(true), + mTexDirectTexture(0) { #ifdef QT_VIVANTE_VIDEO_DEBUG qDebug() << Q_FUNC_INFO; @@ -62,12 +64,7 @@ QSGVivanteVideoMaterial::QSGVivanteVideoMaterial() : QSGVivanteVideoMaterial::~QSGVivanteVideoMaterial() { - Q_FOREACH (GLuint id, mBitsToTextureMap.values()) { -#ifdef QT_VIVANTE_VIDEO_DEBUG - qDebug() << "delete texture: " << id; -#endif - glDeleteTextures(1, &id); - } + clearTextures(); } QSGMaterialType *QSGVivanteVideoMaterial::type() const { @@ -94,9 +91,11 @@ void QSGVivanteVideoMaterial::updateBlending() { setFlag(Blending, qFuzzyCompare(mOpacity, qreal(1.0)) ? false : true); } -void QSGVivanteVideoMaterial::setCurrentFrame(const QVideoFrame &frame) { +void QSGVivanteVideoMaterial::setCurrentFrame(const QVideoFrame &frame, QSGVideoNode::FrameFlags flags) +{ QMutexLocker lock(&mFrameMutex); mNextFrame = frame; + mMappable = !flags.testFlag(QSGVideoNode::FrameFiltered); #ifdef QT_VIVANTE_VIDEO_DEBUG qDebug() << Q_FUNC_INFO << " new frame: " << frame; @@ -122,6 +121,22 @@ void QSGVivanteVideoMaterial::bind() glBindTexture(GL_TEXTURE_2D, mCurrentTexture); } +void QSGVivanteVideoMaterial::clearTextures() +{ + Q_FOREACH (GLuint id, mBitsToTextureMap.values()) { +#ifdef QT_VIVANTE_VIDEO_DEBUG + qDebug() << "delete texture: " << id; +#endif + glDeleteTextures(1, &id); + } + mBitsToTextureMap.clear(); + + if (mTexDirectTexture) { + glDeleteTextures(1, &mTexDirectTexture); + mTexDirectTexture = 0; + } +} + GLuint QSGVivanteVideoMaterial::vivanteMapping(QVideoFrame vF) { QOpenGLContext *glcontext = QOpenGLContext::currentContext(); @@ -130,14 +145,16 @@ GLuint QSGVivanteVideoMaterial::vivanteMapping(QVideoFrame vF) return 0; } + static PFNGLTEXDIRECTVIVPROC glTexDirectVIV_LOCAL = 0; static PFNGLTEXDIRECTVIVMAPPROC glTexDirectVIVMap_LOCAL = 0; static PFNGLTEXDIRECTINVALIDATEVIVPROC glTexDirectInvalidateVIV_LOCAL = 0; - if (glTexDirectVIVMap_LOCAL == 0 || glTexDirectInvalidateVIV_LOCAL == 0) { + if (glTexDirectVIV_LOCAL == 0 || glTexDirectVIVMap_LOCAL == 0 || glTexDirectInvalidateVIV_LOCAL == 0) { + glTexDirectVIV_LOCAL = reinterpret_cast<PFNGLTEXDIRECTVIVPROC>(glcontext->getProcAddress("glTexDirectVIV")); glTexDirectVIVMap_LOCAL = reinterpret_cast<PFNGLTEXDIRECTVIVMAPPROC>(glcontext->getProcAddress("glTexDirectVIVMap")); glTexDirectInvalidateVIV_LOCAL = reinterpret_cast<PFNGLTEXDIRECTINVALIDATEVIVPROC>(glcontext->getProcAddress("glTexDirectInvalidateVIV")); } - if (glTexDirectVIVMap_LOCAL == 0 || glTexDirectInvalidateVIV_LOCAL == 0) { + if (glTexDirectVIV_LOCAL == 0 || glTexDirectVIVMap_LOCAL == 0 || glTexDirectInvalidateVIV_LOCAL == 0) { qWarning() << Q_FUNC_INFO << "couldn't find \"glTexDirectVIVMap\" and/or \"glTexDirectInvalidateVIV\" => do nothing and return"; return 0; } @@ -146,49 +163,80 @@ GLuint QSGVivanteVideoMaterial::vivanteMapping(QVideoFrame vF) mWidth = vF.width(); mHeight = vF.height(); mFormat = vF.pixelFormat(); - Q_FOREACH (GLuint id, mBitsToTextureMap.values()) { -#ifdef QT_VIVANTE_VIDEO_DEBUG - qDebug() << "delete texture: " << id; -#endif - glDeleteTextures(1, &id); - } - mBitsToTextureMap.clear(); + clearTextures(); } if (vF.map(QAbstractVideoBuffer::ReadOnly)) { - if (!mBitsToTextureMap.contains(vF.bits())) { - GLuint tmpTexId; - glGenTextures(1, &tmpTexId); - mBitsToTextureMap.insert(vF.bits(), tmpTexId); + if (mMappable) { + if (!mBitsToTextureMap.contains(vF.bits())) { + // Haven't yet seen this logical address: map to texture. + GLuint tmpTexId; + glGenTextures(1, &tmpTexId); + mBitsToTextureMap.insert(vF.bits(), tmpTexId); - const uchar *constBits = vF.bits(); - void *bits = (void*)constBits; + const uchar *constBits = vF.bits(); + void *bits = (void*)constBits; #ifdef QT_VIVANTE_VIDEO_DEBUG - qDebug() << Q_FUNC_INFO << "new texture, texId: " << tmpTexId << "; constBits: " << constBits; + qDebug() << Q_FUNC_INFO << "new texture, texId: " << tmpTexId << "; constBits: " << constBits; #endif - GLuint physical = ~0U; - - glBindTexture(GL_TEXTURE_2D, tmpTexId); - glTexDirectVIVMap_LOCAL(GL_TEXTURE_2D, - vF.width(), vF.height(), - QSGVivanteVideoNode::getVideoFormat2GLFormatMap().value(vF.pixelFormat()), - &bits, &physical); - - glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_LINEAR); - glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR); - glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_EDGE); - glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_EDGE); - glTexDirectInvalidateVIV_LOCAL(GL_TEXTURE_2D); - - return tmpTexId; - } - else { - glBindTexture(GL_TEXTURE_2D, mBitsToTextureMap.value(vF.bits())); + GLuint physical = ~0U; + + glBindTexture(GL_TEXTURE_2D, tmpTexId); + glTexDirectVIVMap_LOCAL(GL_TEXTURE_2D, + vF.width(), vF.height(), + QSGVivanteVideoNode::getVideoFormat2GLFormatMap().value(vF.pixelFormat()), + &bits, &physical); + + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_LINEAR); + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR); + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_EDGE); + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_EDGE); + glTexDirectInvalidateVIV_LOCAL(GL_TEXTURE_2D); + + return tmpTexId; + } else { + // Fastest path: already seen this logical address. Just + // indicate that the data belonging to the texture has changed. + glBindTexture(GL_TEXTURE_2D, mBitsToTextureMap.value(vF.bits())); + glTexDirectInvalidateVIV_LOCAL(GL_TEXTURE_2D); + return mBitsToTextureMap.value(vF.bits()); + } + } else { + // Cannot map. So copy. + if (!mTexDirectTexture) { + glGenTextures(1, &mTexDirectTexture); + glBindTexture(GL_TEXTURE_2D, mTexDirectTexture); + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_LINEAR); + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR); + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_EDGE); + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_EDGE); + glTexDirectVIV_LOCAL(GL_TEXTURE_2D, mCurrentFrame.width(), mCurrentFrame.height(), + QSGVivanteVideoNode::getVideoFormat2GLFormatMap().value(mCurrentFrame.pixelFormat()), + (GLvoid **) &mTexDirectPlanes); + } else { + glBindTexture(GL_TEXTURE_2D, mTexDirectTexture); + } + switch (mCurrentFrame.pixelFormat()) { + case QVideoFrame::Format_YUV420P: + case QVideoFrame::Format_YV12: + memcpy(mTexDirectPlanes[0], mCurrentFrame.bits(0), mCurrentFrame.height() * mCurrentFrame.bytesPerLine(0)); + memcpy(mTexDirectPlanes[1], mCurrentFrame.bits(1), mCurrentFrame.height() * mCurrentFrame.bytesPerLine(1)); + memcpy(mTexDirectPlanes[2], mCurrentFrame.bits(2), mCurrentFrame.height() * mCurrentFrame.bytesPerLine(2)); + break; + case QVideoFrame::Format_NV12: + case QVideoFrame::Format_NV21: + memcpy(mTexDirectPlanes[0], mCurrentFrame.bits(0), mCurrentFrame.height() * mCurrentFrame.bytesPerLine(0)); + memcpy(mTexDirectPlanes[1], mCurrentFrame.bits(1), mCurrentFrame.height() / 2 * mCurrentFrame.bytesPerLine(1)); + break; + default: + memcpy(mTexDirectPlanes[0], mCurrentFrame.bits(), mCurrentFrame.height() * mCurrentFrame.bytesPerLine()); + break; + } glTexDirectInvalidateVIV_LOCAL(GL_TEXTURE_2D); - return mBitsToTextureMap.value(vF.bits()); + return mTexDirectTexture; } } else { diff --git a/src/plugins/videonode/imx6/qsgvivantevideomaterial.h b/src/plugins/videonode/imx6/qsgvivantevideomaterial.h index 227bc633c..910595836 100644 --- a/src/plugins/videonode/imx6/qsgvivantevideomaterial.h +++ b/src/plugins/videonode/imx6/qsgvivantevideomaterial.h @@ -41,6 +41,7 @@ #include <QVideoFrame> #include <QMutex> +#include <private/qsgvideonode_p.h> class QSGVivanteVideoMaterial : public QSGMaterial { @@ -52,7 +53,7 @@ public: virtual QSGMaterialShader *createShader() const; virtual int compare(const QSGMaterial *other) const; void updateBlending(); - void setCurrentFrame(const QVideoFrame &frame); + void setCurrentFrame(const QVideoFrame &frame, QSGVideoNode::FrameFlags flags); void bind(); GLuint vivanteMapping(QVideoFrame texIdVideoFramePair); @@ -60,6 +61,8 @@ public: void setOpacity(float o) { mOpacity = o; } private: + void clearTextures(); + qreal mOpacity; int mWidth; @@ -69,8 +72,12 @@ private: QMap<const uchar*, GLuint> mBitsToTextureMap; QVideoFrame mCurrentFrame, mNextFrame; GLuint mCurrentTexture; + bool mMappable; QMutex mFrameMutex; + + GLuint mTexDirectTexture; + GLvoid *mTexDirectPlanes[3]; }; #endif // QSGVIDEOMATERIAL_VIVMAP_H diff --git a/src/plugins/videonode/imx6/qsgvivantevideonode.cpp b/src/plugins/videonode/imx6/qsgvivantevideonode.cpp index 1c1c1008b..e24ab3962 100644 --- a/src/plugins/videonode/imx6/qsgvivantevideonode.cpp +++ b/src/plugins/videonode/imx6/qsgvivantevideonode.cpp @@ -52,9 +52,9 @@ QSGVivanteVideoNode::~QSGVivanteVideoNode() { } -void QSGVivanteVideoNode::setCurrentFrame(const QVideoFrame &frame) +void QSGVivanteVideoNode::setCurrentFrame(const QVideoFrame &frame, FrameFlags flags) { - mMaterial->setCurrentFrame(frame); + mMaterial->setCurrentFrame(frame, flags); markDirty(DirtyMaterial); } diff --git a/src/plugins/videonode/imx6/qsgvivantevideonode.h b/src/plugins/videonode/imx6/qsgvivantevideonode.h index 5830cc3b4..79b6e9e57 100644 --- a/src/plugins/videonode/imx6/qsgvivantevideonode.h +++ b/src/plugins/videonode/imx6/qsgvivantevideonode.h @@ -44,8 +44,9 @@ public: QSGVivanteVideoNode(const QVideoSurfaceFormat &format); ~QSGVivanteVideoNode(); - virtual QVideoFrame::PixelFormat pixelFormat() const { return mFormat.pixelFormat(); } - void setCurrentFrame(const QVideoFrame &frame); + QVideoFrame::PixelFormat pixelFormat() const { return mFormat.pixelFormat(); } + QAbstractVideoBuffer::HandleType handleType() const { return QAbstractVideoBuffer::NoHandle; } + void setCurrentFrame(const QVideoFrame &frame, FrameFlags flags); static const QMap<QVideoFrame::PixelFormat, GLenum>& getVideoFormat2GLFormatMap(); diff --git a/src/qtmultimediaquicktools/qdeclarativevideooutput.cpp b/src/qtmultimediaquicktools/qdeclarativevideooutput.cpp index 9a9256868..c36fca67d 100644 --- a/src/qtmultimediaquicktools/qdeclarativevideooutput.cpp +++ b/src/qtmultimediaquicktools/qdeclarativevideooutput.cpp @@ -262,6 +262,12 @@ bool QDeclarativeVideoOutput::createBackend(QMediaService *service) m_backend->updateGeometry(); } + if (m_backend) { + m_backend->clearFilters(); + for (int i = 0; i < m_filters.count(); ++i) + m_backend->appendFilter(m_filters[i]); + } + return backendAvailable; } @@ -795,6 +801,12 @@ void QDeclarativeVideoOutput::itemChange(QQuickItem::ItemChange change, m_backend->itemChange(change, changeData); } +void QDeclarativeVideoOutput::releaseResources() +{ + if (m_backend) + m_backend->releaseResources(); +} + void QDeclarativeVideoOutput::geometryChanged(const QRectF &newGeometry, const QRectF &oldGeometry) { Q_UNUSED(newGeometry); @@ -809,4 +821,54 @@ void QDeclarativeVideoOutput::geometryChanged(const QRectF &newGeometry, const Q _q_updateGeometry(); } +/*! + \qmlproperty list<object> QtMultimedia::VideoOutput::filters + + This property holds the list of video filters that are run on the video + frames. The order of the filters in the list matches the order in which + they will be invoked on the video frames. The objects in the list must be + instances of a subclass of QAbstractVideoFilter. + + \sa QAbstractVideoFilter +*/ + +QQmlListProperty<QAbstractVideoFilter> QDeclarativeVideoOutput::filters() +{ + return QQmlListProperty<QAbstractVideoFilter>(this, 0, filter_append, filter_count, filter_at, filter_clear); +} + +void QDeclarativeVideoOutput::filter_append(QQmlListProperty<QAbstractVideoFilter> *property, QAbstractVideoFilter *value) +{ + QDeclarativeVideoOutput *self = static_cast<QDeclarativeVideoOutput *>(property->object); + self->m_filters.append(value); + if (self->m_backend) + self->m_backend->appendFilter(value); +} + +int QDeclarativeVideoOutput::filter_count(QQmlListProperty<QAbstractVideoFilter> *property) +{ + QDeclarativeVideoOutput *self = static_cast<QDeclarativeVideoOutput *>(property->object); + return self->m_filters.count(); +} + +QAbstractVideoFilter *QDeclarativeVideoOutput::filter_at(QQmlListProperty<QAbstractVideoFilter> *property, int index) +{ + QDeclarativeVideoOutput *self = static_cast<QDeclarativeVideoOutput *>(property->object); + return self->m_filters.at(index); +} + +void QDeclarativeVideoOutput::filter_clear(QQmlListProperty<QAbstractVideoFilter> *property) +{ + QDeclarativeVideoOutput *self = static_cast<QDeclarativeVideoOutput *>(property->object); + self->m_filters.clear(); + if (self->m_backend) + self->m_backend->clearFilters(); +} + +void QDeclarativeVideoOutput::_q_invalidateSceneGraph() +{ + if (m_backend) + m_backend->invalidateSceneGraph(); +} + QT_END_NAMESPACE diff --git a/src/qtmultimediaquicktools/qdeclarativevideooutput_render.cpp b/src/qtmultimediaquicktools/qdeclarativevideooutput_render.cpp index 657ef1767..bde416f44 100644 --- a/src/qtmultimediaquicktools/qdeclarativevideooutput_render.cpp +++ b/src/qtmultimediaquicktools/qdeclarativevideooutput_render.cpp @@ -34,6 +34,7 @@ #include "qdeclarativevideooutput_render_p.h" #include "qdeclarativevideooutput_p.h" +#include <QtMultimedia/qabstractvideofilter.h> #include <QtMultimedia/qvideorenderercontrol.h> #include <QtMultimedia/qmediaservice.h> #include <QtCore/qloggingcategory.h> @@ -41,6 +42,8 @@ #include <private/qsgvideonode_p.h> #include <QtGui/QOpenGLContext> +#include <QtQuick/QQuickWindow> +#include <QtCore/QRunnable> QT_BEGIN_NAMESPACE @@ -103,11 +106,79 @@ bool QDeclarativeVideoRendererBackend::init(QMediaService *service) return false; } +void QDeclarativeVideoRendererBackend::appendFilter(QAbstractVideoFilter *filter) +{ + QMutexLocker lock(&m_frameMutex); + m_filters.append(Filter(filter)); +} + +void QDeclarativeVideoRendererBackend::clearFilters() +{ + QMutexLocker lock(&m_frameMutex); + scheduleDeleteFilterResources(); + m_filters.clear(); +} + +class FilterRunnableDeleter : public QRunnable +{ +public: + FilterRunnableDeleter(const QList<QVideoFilterRunnable *> &runnables) : m_runnables(runnables) { } + void run() Q_DECL_OVERRIDE { + foreach (QVideoFilterRunnable *runnable, m_runnables) + delete runnable; + } +private: + QList<QVideoFilterRunnable *> m_runnables; +}; + +void QDeclarativeVideoRendererBackend::scheduleDeleteFilterResources() +{ + if (!q->window()) + return; + + QList<QVideoFilterRunnable *> runnables; + for (int i = 0; i < m_filters.count(); ++i) { + if (m_filters[i].runnable) { + runnables.append(m_filters[i].runnable); + m_filters[i].runnable = 0; + } + } + + if (!runnables.isEmpty()) { + // Request the scenegraph to run our cleanup job on the render thread. + // The execution of our QRunnable may happen after the QML tree including the QAbstractVideoFilter instance is + // destroyed on the main thread so no references to it must be used during cleanup. + q->window()->scheduleRenderJob(new FilterRunnableDeleter(runnables), QQuickWindow::BeforeSynchronizingStage); + } +} + +void QDeclarativeVideoRendererBackend::releaseResources() +{ + // Called on the gui thread when the window is closed or changed. + QMutexLocker lock(&m_frameMutex); + scheduleDeleteFilterResources(); +} + +void QDeclarativeVideoRendererBackend::invalidateSceneGraph() +{ + // Called on the render thread, e.g. when the context is lost. + QMutexLocker lock(&m_frameMutex); + for (int i = 0; i < m_filters.count(); ++i) { + if (m_filters[i].runnable) { + delete m_filters[i].runnable; + m_filters[i].runnable = 0; + } + } +} + void QDeclarativeVideoRendererBackend::itemChange(QQuickItem::ItemChange change, const QQuickItem::ItemChangeData &changeData) { - Q_UNUSED(change); - Q_UNUSED(changeData); + if (change == QQuickItem::ItemSceneChange) { + if (changeData.window) + QObject::connect(changeData.window, SIGNAL(sceneGraphInvalidated()), + q, SLOT(_q_invalidateSceneGraph()), Qt::DirectConnection); + } } void QDeclarativeVideoRendererBackend::releaseSource() @@ -216,8 +287,36 @@ QSGNode *QDeclarativeVideoRendererBackend::updatePaintNode(QSGNode *oldNode, } #endif + bool isFrameModified = false; if (m_frameChanged) { - if (videoNode && videoNode->pixelFormat() != m_frame.pixelFormat()) { + // Run the VideoFilter if there is one. This must be done before potentially changing the videonode below. + if (m_frame.isValid() && !m_filters.isEmpty()) { + const QVideoSurfaceFormat surfaceFormat = videoSurface()->surfaceFormat(); + for (int i = 0; i < m_filters.count(); ++i) { + QAbstractVideoFilter *filter = m_filters[i].filter; + QVideoFilterRunnable *&runnable = m_filters[i].runnable; + if (filter && filter->isActive()) { + // Create the filter runnable if not yet done. Ownership is taken and is tied to this thread, on which rendering happens. + if (!runnable) + runnable = filter->createFilterRunnable(); + if (!runnable) + continue; + + QVideoFilterRunnable::RunFlags flags = 0; + if (i == m_filters.count() - 1) + flags |= QVideoFilterRunnable::LastInChain; + + QVideoFrame newFrame = runnable->run(&m_frame, surfaceFormat, flags); + + if (newFrame.isValid() && newFrame != m_frame) { + isFrameModified = true; + m_frame = newFrame; + } + } + } + } + + if (videoNode && (videoNode->pixelFormat() != m_frame.pixelFormat() || videoNode->handleType() != m_frame.handleType())) { qCDebug(qLcVideo) << "updatePaintNode: deleting old video node because frame format changed"; delete videoNode; videoNode = 0; @@ -231,7 +330,9 @@ QSGNode *QDeclarativeVideoRendererBackend::updatePaintNode(QSGNode *oldNode, if (!videoNode) { foreach (QSGVideoNodeFactoryInterface* factory, m_videoNodeFactories) { - videoNode = factory->createNode(m_surface->surfaceFormat()); + // Get a node that supports our frame. The surface is irrelevant, our + // QSGVideoItemSurface supports (logically) anything. + videoNode = factory->createNode(QVideoSurfaceFormat(m_frame.size(), m_frame.pixelFormat(), m_frame.handleType())); if (videoNode) { qCDebug(qLcVideo) << "updatePaintNode: Video node created. Handle type:" << m_frame.handleType() << " Supported formats for the handle by this node:" @@ -252,7 +353,10 @@ QSGNode *QDeclarativeVideoRendererBackend::updatePaintNode(QSGNode *oldNode, videoNode->setTexturedRectGeometry(m_renderedRect, m_sourceTextureRect, qNormalizedOrientation(q->orientation())); if (m_frameChanged) { - videoNode->setCurrentFrame(m_frame); + QSGVideoNode::FrameFlags flags = 0; + if (isFrameModified) + flags |= QSGVideoNode::FrameFiltered; + videoNode->setCurrentFrame(m_frame, flags); //don't keep the frame for more than really necessary m_frameChanged = false; m_frame = QVideoFrame(); diff --git a/src/qtmultimediaquicktools/qdeclarativevideooutput_render_p.h b/src/qtmultimediaquicktools/qdeclarativevideooutput_render_p.h index dd8449fa8..a0fed3d00 100644 --- a/src/qtmultimediaquicktools/qdeclarativevideooutput_render_p.h +++ b/src/qtmultimediaquicktools/qdeclarativevideooutput_render_p.h @@ -48,6 +48,8 @@ QT_BEGIN_NAMESPACE class QSGVideoItemSurface; class QVideoRendererControl; class QOpenGLContext; +class QAbstractVideoFilter; +class QVideoFilterRunnable; class QDeclarativeVideoRendererBackend : public QDeclarativeVideoBackend { @@ -70,7 +72,14 @@ public: void present(const QVideoFrame &frame); void stop(); + void appendFilter(QAbstractVideoFilter *filter) Q_DECL_OVERRIDE; + void clearFilters() Q_DECL_OVERRIDE; + void releaseResources() Q_DECL_OVERRIDE; + void invalidateSceneGraph() Q_DECL_OVERRIDE; + private: + void scheduleDeleteFilterResources(); + QPointer<QVideoRendererControl> m_rendererControl; QList<QSGVideoNodeFactoryInterface*> m_videoNodeFactories; QSGVideoItemSurface *m_surface; @@ -83,6 +92,14 @@ private: QMutex m_frameMutex; QRectF m_renderedRect; // Destination pixel coordinates, clipped QRectF m_sourceTextureRect; // Source texture coordinates + + struct Filter { + Filter() : filter(0), runnable(0) { } + Filter(QAbstractVideoFilter *filter) : filter(filter), runnable(0) { } + QAbstractVideoFilter *filter; + QVideoFilterRunnable *runnable; + }; + QList<Filter> m_filters; }; class QSGVideoItemSurface : public QAbstractVideoSurface diff --git a/src/qtmultimediaquicktools/qsgvideonode_i420.cpp b/src/qtmultimediaquicktools/qsgvideonode_i420.cpp index 17b4924d0..3ff1f5bd9 100644 --- a/src/qtmultimediaquicktools/qsgvideonode_i420.cpp +++ b/src/qtmultimediaquicktools/qsgvideonode_i420.cpp @@ -311,7 +311,7 @@ QSGVideoNode_I420::~QSGVideoNode_I420() { } -void QSGVideoNode_I420::setCurrentFrame(const QVideoFrame &frame) +void QSGVideoNode_I420::setCurrentFrame(const QVideoFrame &frame, FrameFlags) { m_material->setCurrentFrame(frame); markDirty(DirtyMaterial); diff --git a/src/qtmultimediaquicktools/qsgvideonode_i420.h b/src/qtmultimediaquicktools/qsgvideonode_i420.h index 06d302d80..71dacf157 100644 --- a/src/qtmultimediaquicktools/qsgvideonode_i420.h +++ b/src/qtmultimediaquicktools/qsgvideonode_i420.h @@ -49,7 +49,10 @@ public: virtual QVideoFrame::PixelFormat pixelFormat() const { return m_format.pixelFormat(); } - void setCurrentFrame(const QVideoFrame &frame); + QAbstractVideoBuffer::HandleType handleType() const { + return QAbstractVideoBuffer::NoHandle; + } + void setCurrentFrame(const QVideoFrame &frame, FrameFlags flags); private: void bindTexture(int id, int unit, int w, int h, const uchar *bits); diff --git a/src/qtmultimediaquicktools/qsgvideonode_rgb.cpp b/src/qtmultimediaquicktools/qsgvideonode_rgb.cpp index 72cbc9614..486632914 100644 --- a/src/qtmultimediaquicktools/qsgvideonode_rgb.cpp +++ b/src/qtmultimediaquicktools/qsgvideonode_rgb.cpp @@ -278,7 +278,7 @@ QSGVideoNode_RGB::~QSGVideoNode_RGB() { } -void QSGVideoNode_RGB::setCurrentFrame(const QVideoFrame &frame) +void QSGVideoNode_RGB::setCurrentFrame(const QVideoFrame &frame, FrameFlags) { m_material->setVideoFrame(frame); markDirty(DirtyMaterial); diff --git a/src/qtmultimediaquicktools/qsgvideonode_rgb.h b/src/qtmultimediaquicktools/qsgvideonode_rgb.h index e6800ce9d..6fcca9bf4 100644 --- a/src/qtmultimediaquicktools/qsgvideonode_rgb.h +++ b/src/qtmultimediaquicktools/qsgvideonode_rgb.h @@ -50,7 +50,10 @@ public: virtual QVideoFrame::PixelFormat pixelFormat() const { return m_format.pixelFormat(); } - void setCurrentFrame(const QVideoFrame &frame); + QAbstractVideoBuffer::HandleType handleType() const { + return QAbstractVideoBuffer::NoHandle; + } + void setCurrentFrame(const QVideoFrame &frame, FrameFlags flags); private: QVideoSurfaceFormat m_format; diff --git a/src/qtmultimediaquicktools/qsgvideonode_texture.cpp b/src/qtmultimediaquicktools/qsgvideonode_texture.cpp index 3af14c0b8..b88327d0d 100644 --- a/src/qtmultimediaquicktools/qsgvideonode_texture.cpp +++ b/src/qtmultimediaquicktools/qsgvideonode_texture.cpp @@ -235,7 +235,7 @@ QSGVideoNode_Texture::~QSGVideoNode_Texture() { } -void QSGVideoNode_Texture::setCurrentFrame(const QVideoFrame &frame) +void QSGVideoNode_Texture::setCurrentFrame(const QVideoFrame &frame, FrameFlags) { m_material->setVideoFrame(frame); markDirty(DirtyMaterial); diff --git a/src/qtmultimediaquicktools/qsgvideonode_texture.h b/src/qtmultimediaquicktools/qsgvideonode_texture.h index c5b84c53d..1389d4f3d 100644 --- a/src/qtmultimediaquicktools/qsgvideonode_texture.h +++ b/src/qtmultimediaquicktools/qsgvideonode_texture.h @@ -50,7 +50,10 @@ public: virtual QVideoFrame::PixelFormat pixelFormat() const { return m_format.pixelFormat(); } - void setCurrentFrame(const QVideoFrame &frame); + QAbstractVideoBuffer::HandleType handleType() const { + return QAbstractVideoBuffer::GLTextureHandle; + } + void setCurrentFrame(const QVideoFrame &frame, FrameFlags flags); private: QVideoSurfaceFormat m_format; |