Long live QtQuickCL!

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

1 files changed, 318 insertions, 0 deletions

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"