diff options
Diffstat (limited to 'examples/multimedia/video/qmlvideofilter_opencl/main.cpp')
| -rw-r--r-- | examples/multimedia/video/qmlvideofilter_opencl/main.cpp | 526 |
1 files changed, 0 insertions, 526 deletions
diff --git a/examples/multimedia/video/qmlvideofilter_opencl/main.cpp b/examples/multimedia/video/qmlvideofilter_opencl/main.cpp deleted file mode 100644 index 2ee2d8bf4..000000000 --- a/examples/multimedia/video/qmlvideofilter_opencl/main.cpp +++ /dev/null @@ -1,526 +0,0 @@ -/**************************************************************************** -** -** Copyright (C) 2016 The Qt Company Ltd. -** Contact: https://www.qt.io/licensing/ -** -** This file is part of the examples of the Qt Multimedia module. -** -** $QT_BEGIN_LICENSE:BSD$ -** 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 The Qt Company. For licensing terms -** and conditions see https://www.qt.io/terms-conditions. For further -** information use the contact form at https://www.qt.io/contact-us. -** -** BSD License Usage -** Alternatively, you may use this file under the terms of the BSD license -** as follows: -** -** "Redistribution and use in source and binary forms, with or without -** modification, are permitted provided that the following conditions are -** met: -** * Redistributions of source code must retain the above copyright -** notice, this list of conditions and the following disclaimer. -** * Redistributions in binary form must reproduce the above copyright -** notice, this list of conditions and the following disclaimer in -** the documentation and/or other materials provided with the -** distribution. -** * Neither the name of The Qt Company Ltd nor the names of its -** contributors may be used to endorse or promote products derived -** from this software without specific prior written permission. -** -** -** THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS -** "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT -** LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR -** A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT -** OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, -** SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT -** LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, -** DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY -** THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT -** (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE -** OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE." -** -** $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 - -#ifdef Q_OS_LINUX -#include <QtPlatformHeaders/QGLXNativeContext> -#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() 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) 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_uint n; - cl_int err = clGetPlatformIDs(0, 0, &n); - if (err != CL_SUCCESS) { - qWarning("Failed to get platform ID count (error %d)", err); - if (err == -1001) { - qDebug("Could not find OpenCL implementation. ICD missing?" -#ifdef Q_OS_LINUX - " Check /etc/OpenCL/vendors." -#endif - ); - } - return; - } - if (n == 0) { - qWarning("No OpenCL platform found"); - return; - } - QList<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"); - const bool isAMD = vendor && strstr(vendor, "ATI"); - 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]; - else if (isAMD && name.contains(QByteArrayLiteral("AMD"))) - 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 }; -#elif defined(Q_OS_LINUX) - QVariant nativeGLXHandle = QOpenGLContext::currentContext()->nativeHandle(); - QGLXNativeContext nativeGLXContext; - if (!nativeGLXHandle.isNull() && nativeGLXHandle.canConvert<QGLXNativeContext>()) - nativeGLXContext = nativeGLXHandle.value<QGLXNativeContext>(); - else - qWarning("Failed to get the underlying GLX context from the current QOpenGLContext"); - cl_context_properties contextProps[] = { CL_CONTEXT_PLATFORM, (cl_context_properties) platform, - CL_GL_CONTEXT_KHR, (cl_context_properties) nativeGLXContext.context(), - CL_GLX_DISPLAY_KHR, (cl_context_properties) glXGetCurrentDisplay(), - 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. - err = clGetGLContextInfoAPPLE(m_clContext, CGLGetCurrentContext(), - CL_CGL_DEVICE_FOR_CURRENT_VIRTUAL_SCREEN_APPLE, - sizeof(cl_device_id), &m_clDeviceId, 0); - if (err != CL_SUCCESS) { - qWarning("Failed to get OpenCL device for current screen: %d", err); - return; - } -#else - clGetGLContextInfoKHR_fn getGLContextInfo = (clGetGLContextInfoKHR_fn) clGetExtensionFunctionAddress("clGetGLContextInfoKHR"); - if (!getGLContextInfo || getGLContextInfo(contextProps, CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR, - sizeof(cl_device_id), &m_clDeviceId, 0) != CL_SUCCESS) { - err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &m_clDeviceId, 0); - if (err != CL_SUCCESS) { - qWarning("Failed to get OpenCL device: %d", err); - 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). - // The latter is the fast path where everything happens on GPU. THe former involves a texture upload. - - if (!input->isValid() - || (input->handleType() != QVideoFrame::NoHandle - && input->handleType() != QVideoFrame::GLTextureHandle)) { - qWarning("Invalid input format"); - return *input; - } - - if (input->pixelFormat() == QVideoSurfaceFormat::Format_YUV420P - || input->pixelFormat() == QVideoSurfaceFormat::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() == QVideoFrame::NoHandle) { - // Upload. - if (m_tempTexture) - f->glBindTexture(GL_TEXTURE_2D, m_tempTexture); - else - m_tempTexture = newTexture(); - input->map(QVideoFrame::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() 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) 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 QVideoFrame::NoHandle: - result->m_handleType = QLatin1String("pixel data"); - result->m_pixelFormat = input->pixelFormat(); - break; - case QVideoFrame::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) -{ -#ifdef Q_OS_WIN // avoid ANGLE on Windows - QCoreApplication::setAttribute(Qt::AA_UseDesktopOpenGL); -#endif - 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" |