diff options
Diffstat (limited to 'examples/multimedia/video/qmlvideofilter_opencl/main.cpp')
| -rw-r--r-- | examples/multimedia/video/qmlvideofilter_opencl/main.cpp | 481 |
1 files changed, 481 insertions, 0 deletions
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" |