Long live QtQuickCL!

Change-Id: I16954f50c631441189dd450980da2c3266a68c71
Reviewed-by: Andy Nichols <andy.nichols@theqtcompany.com>

7 files changed, 538 insertions, 0 deletions

diff --git a/examples/quickcl/histogram/histogram.cl b/examples/quickcl/histogram/histogram.cl
new file mode 100644
index 0000000..de97c2d
--- /dev/null
+++ b/examples/quickcl/histogram/histogram.cl
@@ -0,0 +1,57 @@
+// Based on https://code.google.com/p/opencl-book-samples/source/browse/trunk/src/Chapter_14/histogram
+
+#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable
+
+constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
+
+kernel void histogram(image2d_t img, int num_pixels_per_workitem, global uint *buf)
+{
+    int local_size = get_local_size(0) * get_local_size(1);
+    int item_offset = get_local_id(0) + get_local_id(1) * get_local_size(0);
+    local uint tmp[256];
+
+    int i = 0, j = 256;
+    do {
+        if (item_offset < j)
+            tmp[item_offset + i] = 0;
+        j -= local_size;
+        i += local_size;
+    } while (j > 0);
+
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    int x, image_width = get_image_width(img), image_height = get_image_height(img);
+    for (i = 0, x = get_global_id(0); i < num_pixels_per_workitem; ++i, x += get_global_size(0)) {
+        if (x < image_width && get_global_id(1) < image_height) {
+            float4 clr = read_imagef(img, sampler, (int2)(x, get_global_id(1)));
+            if (clr.w > 0.99f) {
+                float v = (clr.x + clr.y + clr.z) / 3.0f;
+                atom_inc(&tmp[convert_ushort_sat(min(v, 1.0f) * 255.0f)]);
+            }
+        }
+    }
+
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    int group_offset = (get_group_id(0) + get_group_id(1) * get_num_groups(0)) * 256;
+    i = 0;
+    j = 256;
+    do {
+        if (item_offset < j)
+            buf[group_offset + item_offset + i] = tmp[item_offset + i];
+        j -= local_size;
+        i += local_size;
+    } while (j > 0);
+}
+
+kernel void sum_histogram(global uint *buf, int num_groups, global uint *result)
+{
+    int idx = get_global_id(0) + get_group_id(0);
+    uint v = 0;
+    int group_offset = 0, n = num_groups;
+    while (--n >= 0) {
+        v += buf[group_offset + idx];
+        group_offset += 256;
+    }
+    result[idx] = v;
+}
diff --git a/examples/quickcl/histogram/histogram.cpp b/examples/quickcl/histogram/histogram.cpp
new file mode 100644
index 0000000..6ba0b65
--- /dev/null
+++ b/examples/quickcl/histogram/histogram.cpp
@@ -0,0 +1,318 @@
+/****************************************************************************
+**
+** Copyright (C) 2015 The Qt Company Ltd.
+** Contact: http://www.qt.io/licensing/
+**
+** This file is part of the examples of the Qt Quick CL module
+**
+** $QT_BEGIN_LICENSE:BSD$
+** 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$
+**
+****************************************************************************/
+
+// Demonstrates the texture-to-data use case where the results of the
+// computation are exposed to QML and visualized using Qt Quick items.
+
+#include <QGuiApplication>
+#include <QQuickView>
+#include <QQmlEngine>
+#include <QQuickCLItem>
+#include <QQuickCLImageRunnable>
+#include <QAbstractListModel>
+
+static bool profile = false;
+
+class HistogramModel : public QAbstractListModel
+{
+    Q_OBJECT
+
+public:
+    Q_INVOKABLE uint get(int idx) const { return h[idx]; }
+
+    void setHistogram(const QVector<uint> &histogram) {
+        emit beginResetModel();
+        h = histogram;
+        emit endResetModel();
+    }
+
+    int rowCount(const QModelIndex &) const Q_DECL_OVERRIDE {
+        return h.count();
+    }
+
+    QVariant data(const QModelIndex &index, int role) const Q_DECL_OVERRIDE {
+        if (index.isValid() && index.row() < h.count() && role == Qt::UserRole + 1)
+            return h[index.row()];
+        return QVariant();
+    }
+
+    QHash<int, QByteArray> roleNames() const Q_DECL_OVERRIDE {
+        QHash<int, QByteArray> roles;
+        roles[Qt::UserRole + 1] = "value";
+        return roles;
+    }
+
+private:
+    QVector<uint> h;
+};
+
+class CLRunnable;
+
+class CLItem : public QQuickCLItem
+{
+    Q_OBJECT
+    Q_PROPERTY(QQuickItem *source READ source WRITE setSource NOTIFY sourceChanged)
+    Q_PROPERTY(QAbstractItemModel *result READ result NOTIFY resultChanged)
+
+public:
+    CLItem() : m_source(0), m_result(new HistogramModel) { }
+    ~CLItem() { delete m_result; }
+
+    QQuickCLRunnable *createCL() Q_DECL_OVERRIDE;
+    void eventCompleted(cl_event event) Q_DECL_OVERRIDE;
+
+    QQuickItem *source() const { return m_source; }
+    void setSource(QQuickItem *source) {
+        if (m_source != source) {
+            m_source = source;
+            emit sourceChanged();
+            update();
+        }
+    }
+
+    QAbstractItemModel *result() const { return m_result; }
+    void updateResult(const QByteArray &histogram) {
+        const cl_uint *p = (const cl_uint *) histogram.constData();
+        QVector<uint> v(256);
+        for (int i = 0; i < v.count(); ++i)
+            v[i] = p[i];
+        m_result->setHistogram(v);
+        emit newResultsAvailable();
+    }
+
+signals:
+    void sourceChanged();
+    void resultChanged(); // not used as there is a single HistogramModel instance
+    void newResultsAvailable();
+
+private:
+    QQuickItem *m_source;
+    HistogramModel *m_result;
+    CLRunnable *m_runnable;
+};
+
+class CLRunnable : public QQuickCLImageRunnable
+{
+public:
+    CLRunnable(CLItem *item);
+    ~CLRunnable();
+    void runKernel(cl_mem inImage, cl_mem outImage, const QSize &size) Q_DECL_OVERRIDE;
+    QByteArray rawResult() const { return m_result; }
+    void resetPending() { m_resultPending.testAndSetOrdered(1, 0); }
+
+private:
+    CLItem *m_item;
+    cl_program m_program;
+    cl_kernel m_kernel;
+    cl_kernel m_sumKernel;
+    cl_mem m_resultBuf;
+    cl_mem m_sharedBuf;
+    cl_event m_doneEvent;
+    QByteArray m_result;
+    QAtomicInt m_resultPending;
+};
+
+CLRunnable::CLRunnable(CLItem *item)
+    : QQuickCLImageRunnable(item, NoOutputImage | (profile ? Profile : Flag(0))), // note the NoOutputImage flag
+      m_item(item),
+      m_program(0),
+      m_kernel(0),
+      m_sumKernel(0),
+      m_resultBuf(0),
+      m_sharedBuf(0),
+      m_doneEvent(0)
+{
+    QByteArray platform = m_item->platformName();
+    qDebug("Using platform %s", platform.constData());
+    m_program = m_item->buildProgramFromFile(QStringLiteral(":/histogram.cl"));
+    if (!m_program)
+        return;
+    cl_int err;
+    m_kernel = clCreateKernel(m_program, "histogram", &err);
+    if (!m_kernel) {
+        qWarning("Failed to create histogram OpenCL kernel: %d", err);
+        return;
+    }
+    m_sumKernel = clCreateKernel(m_program, "sum_histogram", &err);
+    if (!m_sumKernel) {
+        qWarning("Failed to create sum histogram OpenCL kernel: %d", err);
+        return;
+    }
+    m_resultBuf = clCreateBuffer(m_item->context(), CL_MEM_WRITE_ONLY, 256 * sizeof(cl_uint), 0, &err);
+    if (!m_resultBuf) {
+        qWarning("Failed to create OpenCL buffer: %d", err);
+        return;
+    }
+    m_result.resize(256 * sizeof(cl_uint));
+}
+
+CLRunnable::~CLRunnable()
+{
+    if (m_sharedBuf)
+        clReleaseMemObject(m_sharedBuf);
+    if (m_resultBuf)
+        clReleaseMemObject(m_resultBuf);
+    if (m_kernel)
+        clReleaseKernel(m_kernel);
+    if (m_sumKernel)
+        clReleaseKernel(m_sumKernel);
+    if (m_program)
+        clReleaseProgram(m_program);
+}
+
+class PendingGuard
+{
+public:
+    PendingGuard(CLRunnable *r) : r(r) { }
+    ~PendingGuard() { if (r) r->resetPending(); }
+    void take() { r = 0; }
+private:
+    CLRunnable *r;
+};
+
+#define DIV(a, b) ((a + b - 1) / b)
+
+void CLRunnable::runKernel(cl_mem inImage, cl_mem, const QSize &size)
+{
+    if (!m_program || !m_kernel || !m_sumKernel || !m_resultBuf)
+        return;
+
+    if (!m_resultPending.testAndSetOrdered(0, 1))
+        return;
+    PendingGuard pg(this);
+
+    if (profile)
+        qDebug("CL time: %f ms", elapsed());
+
+    const size_t group_size_x = 16;
+    const size_t group_size_y = 8;
+    const cl_int num_pixels_per_item = 32;
+    const size_t num_items_per_row = DIV(size.width(), num_pixels_per_item);
+    const size_t num_groups_x = DIV(num_items_per_row, group_size_x);
+    const size_t num_groups_y = DIV(size.height(), group_size_y);
+    const cl_int num_groups = cl_int(num_groups_x * num_groups_y);
+
+    const size_t localWorkSize[2] = { group_size_x, group_size_y };
+    const size_t globalWorkSize[2] = { localWorkSize[0] * num_groups_x, localWorkSize[1] * num_groups_y };
+
+    cl_int err;
+    if (!m_sharedBuf) {
+        const size_t sharedBufSize = num_groups * 256 * sizeof(cl_uint);
+        m_sharedBuf = clCreateBuffer(m_item->context(), CL_MEM_READ_WRITE, sharedBufSize, 0, &err);
+        if (!m_sharedBuf) {
+            qWarning("Failed to create shared buffer: %d", err);
+            return;
+        }
+    }
+
+    clSetKernelArg(m_kernel, 0, sizeof(cl_mem), &inImage);
+    clSetKernelArg(m_kernel, 1, sizeof(cl_int), &num_pixels_per_item);
+    clSetKernelArg(m_kernel, 2, sizeof(cl_mem), &m_sharedBuf);
+    err = clEnqueueNDRangeKernel(commandQueue(), m_kernel, 2, 0, globalWorkSize, localWorkSize, 0, 0, 0);
+    if (err != CL_SUCCESS) {
+        qWarning("Failed to enqueue kernel: %d", err);
+        return;
+    }
+
+    clSetKernelArg(m_sumKernel, 0, sizeof(cl_mem), &m_sharedBuf);
+    clSetKernelArg(m_sumKernel, 1, sizeof(cl_int), &num_groups);
+    clSetKernelArg(m_sumKernel, 2, sizeof(cl_mem), &m_resultBuf);
+    const size_t sumLocalWorkSize = 256;
+    const size_t sumGlobalWorkSize = 256;
+    err = clEnqueueNDRangeKernel(commandQueue(), m_sumKernel, 1, 0, &sumGlobalWorkSize, &sumLocalWorkSize, 0, 0, 0);
+    if (err != CL_SUCCESS) {
+        qWarning("Failed to enqueue sum kernel: %d", err);
+        return;
+    }
+
+    err = clEnqueueReadBuffer(commandQueue(), m_resultBuf, CL_FALSE, 0, m_result.size(), m_result.data(), 0, 0, &m_doneEvent);
+    if (err != CL_SUCCESS) {
+        qWarning("Failed to enqueue read buffer: %d", err);
+        return;
+    }
+
+    pg.take();
+    m_item->watchEvent(m_doneEvent);
+}
+
+QQuickCLRunnable *CLItem::createCL()
+{
+    m_runnable = new CLRunnable(this);
+    return m_runnable;
+}
+
+void CLItem::eventCompleted(cl_event event)
+{
+    // We are on the gui thread here.
+    updateResult(m_runnable->rawResult());
+
+    clReleaseEvent(event);
+    m_runnable->resetPending();
+
+    // Request an update, which will eventually also lead to issuing the next
+    // round of computation.
+    scheduleUpdate();
+}
+
+int main(int argc, char **argv)
+{
+#ifdef Q_OS_WIN
+    QCoreApplication::setAttribute(Qt::AA_UseDesktopOpenGL);
+#endif
+    QGuiApplication app(argc, argv);
+
+    if (app.arguments().contains(QStringLiteral("--profile")))
+        profile = true;
+
+    QQuickView view;
+    QObject::connect(view.engine(), SIGNAL(quit()), &app, SLOT(quit()));
+
+    qmlRegisterType<CLItem>("quickcl.qt.io", 1, 0, "CLItem");
+
+    view.setSource(QUrl("qrc:///qml/histogram.qml"));
+    view.setResizeMode(QQuickView::SizeRootObjectToView);
+
+    view.resize(1200, 600);
+    view.show();
+
+    return app.exec();
+}
+
+#include "histogram.moc"
diff --git a/examples/quickcl/histogram/histogram.pro b/examples/quickcl/histogram/histogram.pro
new file mode 100644
index 0000000..f759f6a
--- /dev/null
+++ b/examples/quickcl/histogram/histogram.pro
@@ -0,0 +1,15 @@
+TEMPLATE = app
+
+QT += qml quick quickcl
+
+SOURCES = histogram.cpp
+
+RESOURCES = histogram.qrc
+
+OTHER_FILES = $$PWD/qml/histogram.qml
+
+osx {
+    LIBS += -framework OpenCL
+} else {
+    LIBS += -lOpenCL
+}
diff --git a/examples/quickcl/histogram/histogram.qrc b/examples/quickcl/histogram/histogram.qrc
new file mode 100644
index 0000000..8061a8e
--- /dev/null
+++ b/examples/quickcl/histogram/histogram.qrc
@@ -0,0 +1,8 @@
+<RCC>
+    <qresource prefix="/">
+        <file>qml/histogram.qml</file>
+        <file>qml/image.png</file>
+        <file>qml/qt.png</file>
+        <file>histogram.cl</file>
+    </qresource>
+</RCC>
diff --git a/examples/quickcl/histogram/qml/histogram.qml b/examples/quickcl/histogram/qml/histogram.qml
new file mode 100644
index 0000000..f0fade9
--- /dev/null
+++ b/examples/quickcl/histogram/qml/histogram.qml
@@ -0,0 +1,140 @@
+/****************************************************************************
+**
+** Copyright (C) 2015 The Qt Company Ltd.
+** Contact: http://www.qt.io/licensing/
+**
+** This file is part of the examples of the Qt Quick CL module
+**
+** $QT_BEGIN_LICENSE:BSD$
+** 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$
+**
+****************************************************************************/
+
+import QtQuick 2.0
+import QtQuick.Controls 1.0
+import quickcl.qt.io 1.0
+
+Item {
+    Rectangle {
+        id: orig
+        width: parent.width / 2 - 4
+        height: parent.height
+        border.color: "magenta"
+        border.width: 4
+
+        Item {
+            id: src
+            anchors.fill: parent
+            layer.enabled: true
+            Image {
+                anchors.fill: parent
+                anchors.margins: 4
+                source: "image.png"
+            }
+            Image {
+                source: "qt.png"
+                width: 100
+                height: 100
+                smooth: true
+                SequentialAnimation on scale {
+                    loops: Animation.Infinite
+                    NumberAnimation { from: 1; to: 8; duration: 2000; easing.type: Easing.OutBounce }
+                    PauseAnimation { duration: 2000 }
+                    NumberAnimation { from: 8; to: 1; duration: 2000; easing.type: Easing.OutBounce }
+                    PauseAnimation { duration: 2000 }
+                }
+                NumberAnimation on rotation { from: 0; to: 360; duration: 2000; loops: Animation.Infinite }
+                anchors.centerIn: parent
+            }
+        }
+    }
+
+    Rectangle {
+        width: parent.width / 2 - 4
+        height: parent.height
+        anchors.left: orig.right
+        anchors.leftMargin: 4
+        border.color: "magenta"
+        border.width: 4
+
+        CLItem {
+            id: clItem
+            source: src
+            anchors.fill: parent
+            anchors.margins: 4
+            onNewResultsAvailable: {
+                var maxValue = rectRoot.height - topText.y - topText.height;
+                for (var i = 0; i < 256; ++i)
+                    rep.itemAt(i).height = Math.max(1, Math.min(result.get(i) / scaleFactor.value, maxValue))
+            }
+            Item {
+                anchors.fill: parent
+                Row {
+                    anchors.fill: parent
+                    id: rectRoot
+                    Repeater {
+                        id: rep
+                        // It might be tempting to use clItem.result as the
+                        // model. This would however result in a serious
+                        // slowdown due to dropping and adding 256 new
+                        // rectangles on every update.
+                        model: 256
+                        Rectangle {
+                            width: parent ? parent.width / rep.count : 0
+                            height: 0
+                            y: parent ? parent.height - height : 0
+                            color: "black"
+                        }
+                    }
+                }
+            }
+        }
+        Text {
+            id: topText
+            font.pointSize: 16
+            text: "Histogram calculated on the GPU with OpenCL (scale: " + Math.round(scaleFactor.value) + ")"
+            color: "green"
+            x: 10
+            y: 10
+        }
+        Slider {
+            id: scaleFactor
+            x: 10
+            anchors.top: topText.bottom
+            anchors.bottom: parent.bottom
+            anchors.margins: 10
+            minimumValue: 1
+            maximumValue: 1024
+            value: 8
+            orientation: Qt.Vertical
+        }
+    }
+}
diff --git a/examples/quickcl/histogram/qml/image.png b/examples/quickcl/histogram/qml/image.png
new file mode 100644
index 0000000..8339210
--- /dev/null
+++ b/examples/quickcl/histogram/qml/image.png
diff --git a/examples/quickcl/histogram/qml/qt.png b/examples/quickcl/histogram/qml/qt.png
new file mode 100644
index 0000000..7d3e97e
--- /dev/null
+++ b/examples/quickcl/histogram/qml/qt.png