rhi: Add compute api and implement for Vulkan and Metal

D3D11 and GL (4.3+, ES 3.1+) will come separately at a
later time.

Change-Id: If30f2f3d062fa27e57e9912674669225b82a7b93
Reviewed-by: Lars Knoll <lars.knoll@qt.io>

src/gui/rhi/qrhi.cpp

@@ -266,6 +266,18 @@ QT_BEGIN_NAMESPACE
     transitions. Such synchronization is done implicitly by the backends, where
     applicable (for example, Vulkan), by tracking resource usage as necessary.
 
+    \note Resources within a render or compute pass are expected to be bound to
+    a single usage during that pass. For example, a buffer can be used as
+    vertex, index, uniform, or storage buffer, but not a combination of them
+    within a single pass. However, it is perfectly fine to use a buffer as a
+    storage buffer in a compute pass, and then as a vertex buffer in a render
+    pass, for example, assuming the buffer declared both usages upon creation.
+
+    \note Textures have this rule relaxed in certain cases, because using two
+    subresources (typically two different mip levels) of the same texture for
+    different access (one for load, one for store) is supported even within the
+    same pass.
+
     \section3 Resource reuse
 
     From the user's point of view a QRhiResource is reusable immediately after
@@ -481,6 +493,8 @@ QT_BEGIN_NAMESPACE
     when running on plain OpenGL ES 2.0 implementations without the necessary
     extension. When false, only 16-bit unsigned elements are supported in the
     index buffer.
+
+    \value Compute Indicates that compute shaders are supported.
  */
 
 /*!
@@ -1131,21 +1145,22 @@ QDebug operator<<(QDebug dbg, const QRhiVertexInputLayout &v)
 #endif
 
 /*!
-    \class QRhiGraphicsShaderStage
+    \class QRhiShaderStage
     \inmodule QtRhi
-    \brief Specifies the type and the shader code for a shader stage in the graphics pipeline.
+    \brief Specifies the type and the shader code for a shader stage in the pipeline.
  */
 
 /*!
-    \enum QRhiGraphicsShaderStage::Type
+    \enum QRhiShaderStage::Type
     Specifies the type of the shader stage.
 
     \value Vertex Vertex stage
     \value Fragment Fragment (pixel) stage
+    \value Compute Compute stage (this may not always be supported at run time)
  */
 
 /*!
-    \fn QRhiGraphicsShaderStage::QRhiGraphicsShaderStage()
+    \fn QRhiShaderStage::QRhiShaderStage()
 
     Constructs a shader stage description for the vertex stage with an empty
     QShader.
@@ -1160,7 +1175,7 @@ QDebug operator<<(QDebug dbg, const QRhiVertexInputLayout &v)
     In addition, it can also contain variants of the shader with slightly
     modified code. \a v can then be used to select the desired variant.
  */
-QRhiGraphicsShaderStage::QRhiGraphicsShaderStage(Type type, const QShader &shader, QShader::Variant v)
+QRhiShaderStage::QRhiShaderStage(Type type, const QShader &shader, QShader::Variant v)
     : m_type(type),
       m_shader(shader),
       m_shaderVariant(v)
@@ -1168,12 +1183,12 @@ QRhiGraphicsShaderStage::QRhiGraphicsShaderStage(Type type, const QShader &shade
 }
 
 /*!
-    \return \c true if the values in the two QRhiGraphicsShaderStage objects
+    \return \c true if the values in the two QRhiShaderStage objects
     \a a and \a b are equal.
 
-    \relates QRhiGraphicsShaderStage
+    \relates QRhiShaderStage
  */
-bool operator==(const QRhiGraphicsShaderStage &a, const QRhiGraphicsShaderStage &b) Q_DECL_NOTHROW
+bool operator==(const QRhiShaderStage &a, const QRhiShaderStage &b) Q_DECL_NOTHROW
 {
     return a.type() == b.type()
             && a.shader() == b.shader()
@@ -1181,12 +1196,12 @@ bool operator==(const QRhiGraphicsShaderStage &a, const QRhiGraphicsShaderStage
 }
 
 /*!
-    \return \c false if the values in the two QRhiGraphicsShaderStage
+    \return \c false if the values in the two QRhiShaderStage
     objects \a a and \a b are equal; otherwise returns \c true.
 
-    \relates QRhiGraphicsShaderStage
+    \relates QRhiShaderStage
 */
-bool operator!=(const QRhiGraphicsShaderStage &a, const QRhiGraphicsShaderStage &b) Q_DECL_NOTHROW
+bool operator!=(const QRhiShaderStage &a, const QRhiShaderStage &b) Q_DECL_NOTHROW
 {
     return !(a == b);
 }
@@ -1194,18 +1209,18 @@ bool operator!=(const QRhiGraphicsShaderStage &a, const QRhiGraphicsShaderStage
 /*!
     \return the hash value for \a v, using \a seed to seed the calculation.
 
-    \relates QRhiGraphicsShaderStage
+    \relates QRhiShaderStage
  */
-uint qHash(const QRhiGraphicsShaderStage &v, uint seed) Q_DECL_NOTHROW
+uint qHash(const QRhiShaderStage &v, uint seed) Q_DECL_NOTHROW
 {
     return v.type() + qHash(v.shader(), seed) + v.shaderVariant();
 }
 
 #ifndef QT_NO_DEBUG_STREAM
-QDebug operator<<(QDebug dbg, const QRhiGraphicsShaderStage &s)
+QDebug operator<<(QDebug dbg, const QRhiShaderStage &s)
 {
     QDebugStateSaver saver(dbg);
-    dbg.nospace() << "QRhiGraphicsShaderStage(type=" << s.type()
+    dbg.nospace() << "QRhiShaderStage(type=" << s.type()
                   << " shader=" << s.shader()
                   << " variant=" << s.shaderVariant()
                   << ')';
@@ -1781,9 +1796,25 @@ quint64 QRhiResource::globalResourceId() const
     \enum QRhiBuffer::UsageFlag
     Flag values to specify how the buffer is going to be used.
 
-    \value VertexBuffer Vertex buffer
+    \value VertexBuffer Vertex buffer. This allows the QRhiBuffer to be used in
-    \value IndexBuffer Index buffer
+    \l{setVertexInput()}{QRhiCommandBuffer::setVertexInput()}.
-    \value UniformBuffer Uniform (constant) buffer
+
+    \value IndexBuffer Index buffer. This allows the QRhiBuffer to be used in
+    \l{setVertexInput()}{QRhiCommandBuffer::setVertexInput()}.
+
+    \value UniformBuffer Uniform buffer (also called constant buffer). This
+    allows the QRhiBuffer to be used in combination with
+    \l{UniformBuffer}{QRhiShaderResourceBinding::UniformBuffer}. When
+    \l{QRhi::NonDynamicUniformBuffers}{NonDynamicUniformBuffers} is reported as
+    not supported, this usage can only be combined with the type Dynamic.
+
+    \value StorageBuffer Storage buffer. This allows the QRhiBuffer to be used
+    in combination with \l{BufferLoad}{QRhiShaderResourceBinding::BufferLoad},
+    \l{BufferStore}{QRhiShaderResourceBinding::BufferStore}, or
+    \l{BufferLoadStore}{QRhiShaderResourceBinding::BufferLoadStore}. This usage
+    can only be combined with the types Immutable or Static, and is only
+    available when the \l{QRhi::Compute}{Compute feature} is reported as
+    supported.
  */
 
 /*!
@@ -1941,6 +1972,9 @@ QRhiResource::Type QRhiRenderBuffer::resourceType() const
 
      \value UsedWithGenerateMips The texture is going to be used with
      QRhiResourceUpdateBatch::generateMips().
+
+     \value UsedWithLoadStore The texture is going to be used with image
+     load/store operations, for example, in a compute shader.
  */
 
 /*!
@@ -2438,7 +2472,26 @@ bool QRhiShaderResourceBindings::isLayoutCompatible(const QRhiShaderResourceBind
     Specifies type of the shader resource bound to a binding point
 
     \value UniformBuffer Uniform buffer
+
     \value SampledTexture Combined image sampler
+
+    \value ImageLoad Image load (with GLSL this maps to doing imageLoad() on a
+    single level - and either one or all layers - of a texture exposed to the
+    shader as an image object)
+
+    \value ImageStore Image store (with GLSL this maps to doing imageStore() or
+    imageAtomic*() on a single level - and either one or all layers - of a
+    texture exposed to the shader as an image object)
+
+    \value ImageLoadStore Image load and store
+
+    \value BufferLoad Storage buffer store (with GLSL this maps to reading from
+    a shader storage buffer)
+
+    \value BufferStore Storage buffer store (with GLSL this maps to writing to
+    a shader storage buffer)
+
+    \value BufferLoadStore Storage buffer load and store
  */
 
 /*!
@@ -2447,6 +2500,7 @@ bool QRhiShaderResourceBindings::isLayoutCompatible(const QRhiShaderResourceBind
 
     \value VertexStage Vertex stage
     \value FragmentStage Fragment (pixel) stage
+    \value ComputeStage Compute stage
  */
 
 /*!
@@ -2513,6 +2567,8 @@ bool QRhiShaderResourceBinding::isLayoutCompatible(const QRhiShaderResourceBindi
 /*!
     \return a shader resource binding for the given binding number, pipeline
     stages, and buffer specified by \a binding, \a stage, and \a buf.
+
+    \note \a buf must have been created with QRhiBuffer::UniformBuffer.
  */
 QRhiShaderResourceBinding QRhiShaderResourceBinding::uniformBuffer(
         int binding, StageFlags stage, QRhiBuffer *buf)
@@ -2539,21 +2595,17 @@ QRhiShaderResourceBinding QRhiShaderResourceBinding::uniformBuffer(
     QRhi::ubufAlignment().
 
     \note \a size must be greater than 0.
+
+    \note \a buf must have been created with QRhiBuffer::UniformBuffer.
  */
 QRhiShaderResourceBinding QRhiShaderResourceBinding::uniformBuffer(
         int binding, StageFlags stage, QRhiBuffer *buf, int offset, int size)
 {
     Q_ASSERT(size > 0);
-    QRhiShaderResourceBinding b;
+    QRhiShaderResourceBinding b = uniformBuffer(binding, stage, buf);
     QRhiShaderResourceBindingPrivate *d = QRhiShaderResourceBindingPrivate::get(&b);
-    Q_ASSERT(d->ref.load() == 1);
-    d->binding = binding;
-    d->stage = stage;
-    d->type = UniformBuffer;
-    d->u.ubuf.buf = buf;
     d->u.ubuf.offset = offset;
     d->u.ubuf.maybeSize = size;
-    d->u.ubuf.hasDynamicOffset = false;
     return b;
 }
 
@@ -2565,19 +2617,14 @@ QRhiShaderResourceBinding QRhiShaderResourceBinding::uniformBuffer(
     varying offset values without creating new bindings for the buffer. The
     size of the bound region is specified by \a size. Like with non-dynamic
     offsets, \c{offset + size} cannot exceed the size of \a buf.
+
+    \note \a buf must have been created with QRhiBuffer::UniformBuffer.
  */
 QRhiShaderResourceBinding QRhiShaderResourceBinding::uniformBufferWithDynamicOffset(
         int binding, StageFlags stage, QRhiBuffer *buf, int size)
 {
-    QRhiShaderResourceBinding b;
+    QRhiShaderResourceBinding b = uniformBuffer(binding, stage, buf, 0, size);
     QRhiShaderResourceBindingPrivate *d = QRhiShaderResourceBindingPrivate::get(&b);
-    Q_ASSERT(d->ref.load() == 1);
-    d->binding = binding;
-    d->stage = stage;
-    d->type = UniformBuffer;
-    d->u.ubuf.buf = buf;
-    d->u.ubuf.offset = 0;
-    d->u.ubuf.maybeSize = size;
     d->u.ubuf.hasDynamicOffset = true;
     return b;
 }
@@ -2601,6 +2648,167 @@ QRhiShaderResourceBinding QRhiShaderResourceBinding::sampledTexture(
     return b;
 }
 
+/*!
+   \return a shader resource binding for a read-only storage image with the
+   given \a binding number and pipeline \a stage. The image load operations
+   will have access to all layers of the specified \a level. (so if the texture
+   is a cubemap, the shader must use imageCube instead of image2D)
+
+   \note \a tex must have been created with QRhiTexture::UsedWithLoadStore.
+ */
+QRhiShaderResourceBinding QRhiShaderResourceBinding::imageLoad(
+        int binding, StageFlags stage, QRhiTexture *tex, int level)
+{
+    QRhiShaderResourceBinding b;
+    QRhiShaderResourceBindingPrivate *d = QRhiShaderResourceBindingPrivate::get(&b);
+    Q_ASSERT(d->ref.load() == 1);
+    d->binding = binding;
+    d->stage = stage;
+    d->type = ImageLoad;
+    d->u.simage.tex = tex;
+    d->u.simage.level = level;
+    return b;
+}
+
+/*!
+   \return a shader resource binding for a write-only storage image with the
+   given \a binding number and pipeline \a stage. The image store operations
+   will have access to all layers of the specified \a level. (so if the texture
+   is a cubemap, the shader must use imageCube instead of image2D)
+
+   \note \a tex must have been created with QRhiTexture::UsedWithLoadStore.
+ */
+QRhiShaderResourceBinding QRhiShaderResourceBinding::imageStore(
+        int binding, StageFlags stage, QRhiTexture *tex, int level)
+{
+    QRhiShaderResourceBinding b = imageLoad(binding, stage, tex, level);
+    QRhiShaderResourceBindingPrivate *d = QRhiShaderResourceBindingPrivate::get(&b);
+    d->type = ImageStore;
+    return b;
+}
+
+/*!
+   \return a shader resource binding for a read/write storage image with the
+   given \a binding number and pipeline \a stage. The image load/store operations
+   will have access to all layers of the specified \a level. (so if the texture
+   is a cubemap, the shader must use imageCube instead of image2D)
+
+   \note \a tex must have been created with QRhiTexture::UsedWithLoadStore.
+ */
+QRhiShaderResourceBinding QRhiShaderResourceBinding::imageLoadStore(
+        int binding, StageFlags stage, QRhiTexture *tex, int level)
+{
+    QRhiShaderResourceBinding b = imageLoad(binding, stage, tex, level);
+    QRhiShaderResourceBindingPrivate *d = QRhiShaderResourceBindingPrivate::get(&b);
+    d->type = ImageLoadStore;
+    return b;
+}
+
+/*!
+    \return a shader resource binding for a read-only storage buffer with the
+    given \a binding number and pipeline \a stage.
+
+    \note \a buf must have been created with QRhiBuffer::StorageBuffer.
+ */
+QRhiShaderResourceBinding QRhiShaderResourceBinding::bufferLoad(
+        int binding, StageFlags stage, QRhiBuffer *buf)
+{
+    QRhiShaderResourceBinding b;
+    QRhiShaderResourceBindingPrivate *d = QRhiShaderResourceBindingPrivate::get(&b);
+    Q_ASSERT(d->ref.load() == 1);
+    d->binding = binding;
+    d->stage = stage;
+    d->type = BufferLoad;
+    d->u.sbuf.buf = buf;
+    d->u.sbuf.offset = 0;
+    d->u.sbuf.maybeSize = 0; // entire buffer
+    return b;
+}
+
+/*!
+    \return a shader resource binding for a read-only storage buffer with the
+    given \a binding number and pipeline \a stage. This overload binds a region
+    only, as specified by \a offset and \a size.
+
+    \note \a buf must have been created with QRhiBuffer::StorageBuffer.
+ */
+QRhiShaderResourceBinding QRhiShaderResourceBinding::bufferLoad(
+        int binding, StageFlags stage, QRhiBuffer *buf, int offset, int size)
+{
+    Q_ASSERT(size > 0);
+    QRhiShaderResourceBinding b = bufferLoad(binding, stage, buf);
+    QRhiShaderResourceBindingPrivate *d = QRhiShaderResourceBindingPrivate::get(&b);
+    d->u.sbuf.offset = offset;
+    d->u.sbuf.maybeSize = size;
+    return b;
+}
+
+/*!
+    \return a shader resource binding for a write-only storage buffer with the
+    given \a binding number and pipeline \a stage.
+
+    \note \a buf must have been created with QRhiBuffer::StorageBuffer.
+ */
+QRhiShaderResourceBinding QRhiShaderResourceBinding::bufferStore(
+        int binding, StageFlags stage, QRhiBuffer *buf)
+{
+    QRhiShaderResourceBinding b = bufferLoad(binding, stage, buf);
+    QRhiShaderResourceBindingPrivate *d = QRhiShaderResourceBindingPrivate::get(&b);
+    d->type = BufferStore;
+    return b;
+}
+
+/*!
+    \return a shader resource binding for a write-only storage buffer with the
+    given \a binding number and pipeline \a stage. This overload binds a region
+    only, as specified by \a offset and \a size.
+
+    \note \a buf must have been created with QRhiBuffer::StorageBuffer.
+ */
+QRhiShaderResourceBinding QRhiShaderResourceBinding::bufferStore(
+        int binding, StageFlags stage, QRhiBuffer *buf, int offset, int size)
+{
+    Q_ASSERT(size > 0);
+    QRhiShaderResourceBinding b = bufferStore(binding, stage, buf);
+    QRhiShaderResourceBindingPrivate *d = QRhiShaderResourceBindingPrivate::get(&b);
+    d->u.sbuf.offset = offset;
+    d->u.sbuf.maybeSize = size;
+    return b;
+}
+
+/*!
+    \return a shader resource binding for a read-write storage buffer with the
+    given \a binding number and pipeline \a stage.
+
+    \note \a buf must have been created with QRhiBuffer::StorageBuffer.
+ */
+QRhiShaderResourceBinding QRhiShaderResourceBinding::bufferLoadStore(
+        int binding, StageFlags stage, QRhiBuffer *buf)
+{
+    QRhiShaderResourceBinding b = bufferLoad(binding, stage, buf);
+    QRhiShaderResourceBindingPrivate *d = QRhiShaderResourceBindingPrivate::get(&b);
+    d->type = BufferLoadStore;
+    return b;
+}
+
+/*!
+    \return a shader resource binding for a read-write storage buffer with the
+    given \a binding number and pipeline \a stage. This overload binds a region
+    only, as specified by \a offset and \a size.
+
+    \note \a buf must have been created with QRhiBuffer::StorageBuffer.
+ */
+QRhiShaderResourceBinding QRhiShaderResourceBinding::bufferLoadStore(
+        int binding, StageFlags stage, QRhiBuffer *buf, int offset, int size)
+{
+    Q_ASSERT(size > 0);
+    QRhiShaderResourceBinding b = bufferLoadStore(binding, stage, buf);
+    QRhiShaderResourceBindingPrivate *d = QRhiShaderResourceBindingPrivate::get(&b);
+    d->u.sbuf.offset = offset;
+    d->u.sbuf.maybeSize = size;
+    return b;
+}
+
 /*!
     \return \c true if the contents of the two QRhiShaderResourceBinding
     objects \a a and \a b are equal. This includes the resources (buffer,
@@ -2639,6 +2847,29 @@ bool operator==(const QRhiShaderResourceBinding &a, const QRhiShaderResourceBind
             return false;
         }
         break;
+    case QRhiShaderResourceBinding::ImageLoad:
+        Q_FALLTHROUGH();
+    case QRhiShaderResourceBinding::ImageStore:
+        Q_FALLTHROUGH();
+    case QRhiShaderResourceBinding::ImageLoadStore:
+        if (a.d->u.simage.tex != b.d->u.simage.tex
+                || a.d->u.simage.level != b.d->u.simage.level)
+        {
+            return false;
+        }
+        break;
+    case QRhiShaderResourceBinding::BufferLoad:
+        Q_FALLTHROUGH();
+    case QRhiShaderResourceBinding::BufferStore:
+        Q_FALLTHROUGH();
+    case QRhiShaderResourceBinding::BufferLoadStore:
+        if (a.d->u.sbuf.buf != b.d->u.sbuf.buf
+                || a.d->u.sbuf.offset != b.d->u.sbuf.offset
+                || a.d->u.sbuf.maybeSize != b.d->u.sbuf.maybeSize)
+        {
+            return false;
+        }
+        break;
     default:
         Q_UNREACHABLE();
         return false;
@@ -2693,6 +2924,45 @@ QDebug operator<<(QDebug dbg, const QRhiShaderResourceBinding &b)
                       << " sampler=" << d->u.stex.sampler
                       << ')';
         break;
+    case QRhiShaderResourceBinding::ImageLoad:
+        dbg.nospace() << " ImageLoad("
+                      << "texture=" << d->u.simage.tex
+                      << " level=" << d->u.simage.level
+                      << ')';
+        break;
+    case QRhiShaderResourceBinding::ImageStore:
+        dbg.nospace() << " ImageStore("
+                      << "texture=" << d->u.simage.tex
+                      << " level=" << d->u.simage.level
+                      << ')';
+        break;
+    case QRhiShaderResourceBinding::ImageLoadStore:
+        dbg.nospace() << " ImageLoadStore("
+                      << "texture=" << d->u.simage.tex
+                      << " level=" << d->u.simage.level
+                      << ')';
+        break;
+    case QRhiShaderResourceBinding::BufferLoad:
+        dbg.nospace() << " BufferLoad("
+                      << "buffer=" << d->u.sbuf.buf
+                      << " offset=" << d->u.sbuf.offset
+                      << " maybeSize=" << d->u.sbuf.maybeSize
+                      << ')';
+        break;
+    case QRhiShaderResourceBinding::BufferStore:
+        dbg.nospace() << " BufferStore("
+                      << "buffer=" << d->u.sbuf.buf
+                      << " offset=" << d->u.sbuf.offset
+                      << " maybeSize=" << d->u.sbuf.maybeSize
+                      << ')';
+        break;
+    case QRhiShaderResourceBinding::BufferLoadStore:
+        dbg.nospace() << " BufferLoadStore("
+                      << "buffer=" << d->u.sbuf.buf
+                      << " offset=" << d->u.sbuf.offset
+                      << " maybeSize=" << d->u.sbuf.maybeSize
+                      << ')';
+        break;
     default:
         Q_UNREACHABLE();
         break;
@@ -3195,6 +3465,34 @@ QRhiResource::Type QRhiSwapChain::resourceType() const
     Regardless of the return value, calling release() is always safe.
  */
 
+/*!
+    \class QRhiComputePipeline
+    \inmodule QtRhi
+    \brief Compute pipeline state resource.
+
+    \note Setting the shader resource bindings is mandatory. The referenced
+    QRhiShaderResourceBindings must already be built by the time build() is
+    called.
+
+    \note Setting the shader is mandatory.
+ */
+
+/*!
+    \return the resource type.
+ */
+QRhiResource::Type QRhiComputePipeline::resourceType() const
+{
+    return ComputePipeline;
+}
+
+/*!
+    \internal
+ */
+QRhiComputePipeline::QRhiComputePipeline(QRhiImplementation *rhi)
+    : QRhiResource(rhi)
+{
+}
+
 /*!
     \class QRhiCommandBuffer
     \inmodule QtRhi
@@ -3982,8 +4280,8 @@ void QRhiCommandBuffer::endPass(QRhiResourceUpdateBatch *resourceUpdates)
     therefore overoptimizing to avoid calls to this function is not necessary
     on the applications' side.
 
-    \note This function can only be called inside a pass, meaning between a
+    \note This function can only be called inside a render pass, meaning
-    beginPass() end endPass() call.
+    between a beginPass() and endPass() call.
  */
 void QRhiCommandBuffer::setGraphicsPipeline(QRhiGraphicsPipeline *ps)
 {
@@ -3994,14 +4292,13 @@ void QRhiCommandBuffer::setGraphicsPipeline(QRhiGraphicsPipeline *ps)
     Records binding a set of shader resources, such as, uniform buffers or
     textures, that are made visible to one or more shader stages.
 
-    \a srb can be null in which case the current graphics pipeline's associated
+    \a srb can be null in which case the current graphics or compute pipeline's
-    QRhiGraphicsPipeline::shaderResourceBindings() is used. When \a srb is
+    associated QRhiShaderResourceBindings is used. When \a srb is non-null, it
-    non-null, it must be
+    must be
     \l{QRhiShaderResourceBindings::isLayoutCompatible()}{layout-compatible},
     meaning the layout (number of bindings, the type and binding number of each
     binding) must fully match the QRhiShaderResourceBindings that was
-    associated with the pipeline at the time of calling
+    associated with the pipeline at the time of calling the pipeline's build().
-    QRhiGraphicsPipeline::build().
 
     There are cases when a seemingly unnecessary setShaderResources() call is
     mandatory: when rebuilding a resource referenced from \a srb, for example
@@ -4029,8 +4326,9 @@ void QRhiCommandBuffer::setGraphicsPipeline(QRhiGraphicsPipeline *ps)
     the conditions described above into account), so therefore overoptimizing
     to avoid calls to this function is not necessary on the applications' side.
 
-    \note This function can only be called inside a pass, meaning between a
+    \note This function can only be called inside a render or compute pass,
-    beginPass() end endPass() call.
+    meaning between a beginPass() and endPass(), or beginComputePass() and
+    endComputePass().
  */
 void QRhiCommandBuffer::setShaderResources(QRhiShaderResourceBindings *srb,
                                            int dynamicOffsetCount,
@@ -4056,8 +4354,8 @@ void QRhiCommandBuffer::setShaderResources(QRhiShaderResourceBindings *srb,
     automatically with most backends and therefore applications do not need to
     overoptimize to avoid calls to this function.
 
-    \note This function can only be called inside a pass, meaning between a
+    \note This function can only be called inside a render pass, meaning
-    beginPass() end endPass() call.
+    between a beginPass() and endPass() call.
 
     As a simple example, take a vertex shader with two inputs:
 
@@ -4110,8 +4408,8 @@ void QRhiCommandBuffer::setVertexInput(int startBinding, int bindingCount, const
     \note QRhi assumes OpenGL-style viewport coordinates, meaning x and y are
     bottom-left.
 
-    \note This function can only be called inside a pass, meaning between a
+    \note This function can only be called inside a render pass, meaning
-    beginPass() end endPass() call.
+    between a beginPass() and endPass() call.
  */
 void QRhiCommandBuffer::setViewport(const QRhiViewport &viewport)
 {
@@ -4129,8 +4427,8 @@ void QRhiCommandBuffer::setViewport(const QRhiViewport &viewport)
     \note QRhi assumes OpenGL-style viewport coordinates, meaning x and y are
     bottom-left.
 
-    \note This function can only be called inside a pass, meaning between a
+    \note This function can only be called inside a render pass, meaning
-    beginPass() end endPass() call.
+    between a beginPass() and endPass() call.
  */
 void QRhiCommandBuffer::setScissor(const QRhiScissor &scissor)
 {
@@ -4143,8 +4441,8 @@ void QRhiCommandBuffer::setScissor(const QRhiScissor &scissor)
     This can only be called when the bound pipeline has
     QRhiGraphicsPipeline::UsesBlendConstants set.
 
-    \note This function can only be called inside a pass, meaning between a
+    \note This function can only be called inside a render pass, meaning
-    beginPass() end endPass() call.
+    between a beginPass() and endPass() call.
  */
 void QRhiCommandBuffer::setBlendConstants(const QColor &c)
 {
@@ -4157,8 +4455,8 @@ void QRhiCommandBuffer::setBlendConstants(const QColor &c)
     This can only be called when the bound pipeline has
     QRhiGraphicsPipeline::UsesStencilRef set.
 
-    \note This function can only be called inside a pass, meaning between a
+    \note This function can only be called inside a render pass, meaning between
-    beginPass() end endPass() call.
+    a beginPass() and endPass() call.
  */
 void QRhiCommandBuffer::setStencilRef(quint32 refValue)
 {
@@ -4173,8 +4471,8 @@ void QRhiCommandBuffer::setStencilRef(quint32 refValue)
     the index of the first vertex to draw. \a firstInstance is the instance ID
     of the first instance to draw.
 
-    \note This function can only be called inside a pass, meaning between a
+    \note This function can only be called inside a render pass, meaning
-    beginPass() end endPass() call.
+    between a beginPass() and endPass() call.
  */
 void QRhiCommandBuffer::draw(quint32 vertexCount,
                              quint32 instanceCount, quint32 firstVertex, quint32 firstInstance)
@@ -4200,8 +4498,8 @@ void QRhiCommandBuffer::draw(quint32 vertexCount,
 
     \a vertexOffset is added to the vertex index.
 
-    \note This function can only be called inside a pass, meaning between a
+    \note This function can only be called inside a render pass, meaning
-    beginPass() end endPass() call.
+    between a beginPass() and endPass() call.
  */
 void QRhiCommandBuffer::drawIndexed(quint32 indexCount,
                                     quint32 instanceCount, quint32 firstIndex,
@@ -4254,6 +4552,69 @@ void QRhiCommandBuffer::debugMarkMsg(const QByteArray &msg)
     m_rhi->debugMarkMsg(this, msg);
 }
 
+/*!
+    Records starting a new compute pass.
+
+    \a resourceUpdates, when not null, specifies a resource update batch that
+    is to be committed and then released.
+
+    \note Do not assume that any state or resource bindings persist between
+    passes.
+
+    \note A compute pass can record setComputePipeline(), setShaderResources(),
+    and dispatch() calls, not graphics ones. General functionality, such as,
+    debug markers and beginExternal() is available both in render and compute
+    passes.
+
+    \note Compute is only available when the \l{QRhi::Compute}{Compute} feature
+    is reported as supported.
+ */
+void QRhiCommandBuffer::beginComputePass(QRhiResourceUpdateBatch *resourceUpdates)
+{
+    m_rhi->beginComputePass(this, resourceUpdates);
+}
+
+/*!
+    Records ending the current compute pass.
+
+    \a resourceUpdates, when not null, specifies a resource update batch that
+    is to be committed and then released.
+ */
+void QRhiCommandBuffer::endComputePass(QRhiResourceUpdateBatch *resourceUpdates)
+{
+    m_rhi->endComputePass(this, resourceUpdates);
+}
+
+/*!
+    Records setting a new compute pipeline \a ps.
+
+    \note This function must be called before recording setShaderResources() or
+    dispatch() commands on the command buffer.
+
+    \note QRhi will optimize out unnecessary invocations within a pass, so
+    therefore overoptimizing to avoid calls to this function is not necessary
+    on the applications' side.
+
+    \note This function can only be called inside a compute pass, meaning
+    between a beginComputePass() and endComputePass() call.
+ */
+void QRhiCommandBuffer::setComputePipeline(QRhiComputePipeline *ps)
+{
+    m_rhi->setComputePipeline(this, ps);
+}
+
+/*!
+    Records dispatching compute work items, with \a x, \a y, and \a z
+    specifying the number of local workgroups in the corresponding dimension.
+
+    \note This function can only be called inside a compute pass, meaning
+    between a beginComputePass() and endComputePass() call.
+ */
+void QRhiCommandBuffer::dispatch(int x, int y, int z)
+{
+    m_rhi->dispatch(this, x, y, z);
+}
+
 /*!
     \return a pointer to a backend-specific QRhiNativeHandles subclass, such as
     QRhiVulkanCommandBufferNativeHandles. The returned value is null when
@@ -4479,6 +4840,19 @@ QRhiGraphicsPipeline *QRhi::newGraphicsPipeline()
     return d->createGraphicsPipeline();
 }
 
+/*!
+    \return a new compute pipeline resource.
+
+    \note Compute is only available when the \l{QRhi::Compute}{Compute} feature
+    is reported as supported.
+
+    \sa QRhiResource::release()
+ */
+QRhiComputePipeline *QRhi::newComputePipeline()
+{
+    return d->createComputePipeline();
+}
+
 /*!
     \return a new shader resource binding collection resource.
 
@@ -4493,7 +4867,8 @@ QRhiShaderResourceBindings *QRhi::newShaderResourceBindings()
     \return a new buffer with the specified \a type, \a usage, and \a size.
 
     \note Some \a usage and \a type combinations may not be supported by all
-    backends. See \l{QRhi::NonDynamicUniformBuffers}{the feature flags}.
+    backends. See \l{QRhiBuffer::UsageFlag}{UsageFlags} and
+    \l{QRhi::NonDynamicUniformBuffers}{the feature flags}.
 
     \sa QRhiResource::release()
  */
@@ -4840,32 +5215,30 @@ static inline QRhiPassResourceTracker::BufferStage earlierStage(QRhiPassResource
     return QRhiPassResourceTracker::BufferStage(qMin(int(a), int(b)));
 }
 
-void QRhiPassResourceTracker::registerBufferOnce(QRhiBuffer *buf, int slot, BufferAccess access, BufferStage stage,
+void QRhiPassResourceTracker::registerBuffer(QRhiBuffer *buf, int slot, BufferAccess *access, BufferStage *stage,
-                                                 const UsageState &stateAtPassBegin)
+                                             const UsageState &state)
 {
     auto it = std::find_if(m_buffers.begin(), m_buffers.end(), [buf](const Buffer &b) { return b.buf == buf; });
     if (it != m_buffers.end()) {
-        if (it->access != access) {
+        if (it->access != *access) {
             const QByteArray name = buf->name();
             qWarning("Buffer %p (%s) used with different accesses within the same pass, this is not allowed.",
                      buf, name.constData());
             return;
         }
-        if (it->stage != stage)
+        if (it->stage != *stage) {
-            it->stage = earlierStage(it->stage, stage);
+            it->stage = earlierStage(it->stage, *stage);
-        // Multiple registrations of the same buffer is fine as long is it is
+            *stage = it->stage;
-        // a compatible usage. stateAtPassBegin is not actually the state at
+        }
-        // pass begin in the second, third, etc. invocation but that's fine
-        // since we'll just return here.
         return;
     }
 
     Buffer b;
     b.buf = buf;
     b.slot = slot;
-    b.access = access;
+    b.access = *access;
-    b.stage = stage;
+    b.stage = *stage;
-    b.stateAtPassBegin = stateAtPassBegin;
+    b.stateAtPassBegin = state; // first use -> initial state
     m_buffers.append(b);
 }
 
@@ -4875,30 +5248,44 @@ static inline QRhiPassResourceTracker::TextureStage earlierStage(QRhiPassResourc
     return QRhiPassResourceTracker::TextureStage(qMin(int(a), int(b)));
 }
 
-void QRhiPassResourceTracker::registerTextureOnce(QRhiTexture *tex, TextureAccess access, TextureStage stage,
+static inline bool isImageLoadStore(QRhiPassResourceTracker::TextureAccess access)
-                                                  const UsageState &stateAtPassBegin)
+{
+    return access == QRhiPassResourceTracker::TexStorageLoad
+            || access == QRhiPassResourceTracker::TexStorageStore
+            || access == QRhiPassResourceTracker::TexStorageLoadStore;
+}
+
+void QRhiPassResourceTracker::registerTexture(QRhiTexture *tex, TextureAccess *access, TextureStage *stage,
+                                              const UsageState &state)
 {
     auto it = std::find_if(m_textures.begin(), m_textures.end(), [tex](const Texture &t) { return t.tex == tex; });
     if (it != m_textures.end()) {
-        if (it->access != access) {
+        if (it->access != *access) {
-            const QByteArray name = tex->name();
+            // Different subresources of a texture may be used for both load
-            qWarning("Texture %p (%s) used with different accesses within the same pass, this is not allowed.",
+            // and store in the same pass. (think reading from one mip level
-                     tex, name.constData());
+            // and writing to another one in a compute shader) This we can
+            // handle by treating the entire resource as read-write.
+            if (isImageLoadStore(it->access) && isImageLoadStore(*access)) {
+                it->access = QRhiPassResourceTracker::TexStorageLoadStore;
+                *access = it->access;
+            } else {
+                const QByteArray name = tex->name();
+                qWarning("Texture %p (%s) used with different accesses within the same pass, this is not allowed.",
+                         tex, name.constData());
+            }
+        }
+        if (it->stage != *stage) {
+            it->stage = earlierStage(it->stage, *stage);
+            *stage = it->stage;
         }
-        if (it->stage != stage)
-            it->stage = earlierStage(it->stage, stage);
-        // Multiple registrations of the same texture is fine as long is it is
-        // a compatible usage. stateAtPassBegin is not actually the state at
-        // pass begin in the second, third, etc. invocation but that's fine
-        // since we'll just return here.
         return;
     }
 
     Texture t;
     t.tex = tex;
-    t.access = access;
+    t.access = *access;
-    t.stage = stage;
+    t.stage = *stage;
-    t.stateAtPassBegin = stateAtPassBegin;
+    t.stateAtPassBegin = state; // first use -> initial state
     m_textures.append(t);
 }
 

src/gui/rhi/qrhi_p.h

@@ -259,17 +259,18 @@ Q_GUI_EXPORT uint qHash(const QRhiVertexInputLayout &v, uint seed = 0) Q_DECL_NO
 Q_GUI_EXPORT QDebug operator<<(QDebug, const QRhiVertexInputLayout &);
 #endif
 
-class Q_GUI_EXPORT QRhiGraphicsShaderStage
+class Q_GUI_EXPORT QRhiShaderStage
 {
 public:
     enum Type {
         Vertex,
-        Fragment
+        Fragment,
+        Compute
     };
 
-    QRhiGraphicsShaderStage() = default;
+    QRhiShaderStage() = default;
-    QRhiGraphicsShaderStage(Type type, const QShader &shader,
+    QRhiShaderStage(Type type, const QShader &shader,
-                            QShader::Variant v = QShader::StandardShader);
+                    QShader::Variant v = QShader::StandardShader);
 
     Type type() const { return m_type; }
     void setType(Type t) { m_type = t; }
@@ -286,26 +287,35 @@ private:
     QShader::Variant m_shaderVariant = QShader::StandardShader;
 };
 
-Q_DECLARE_TYPEINFO(QRhiGraphicsShaderStage, Q_MOVABLE_TYPE);
+Q_DECLARE_TYPEINFO(QRhiShaderStage, Q_MOVABLE_TYPE);
 
-Q_GUI_EXPORT bool operator==(const QRhiGraphicsShaderStage &a, const QRhiGraphicsShaderStage &b) Q_DECL_NOTHROW;
+Q_GUI_EXPORT bool operator==(const QRhiShaderStage &a, const QRhiShaderStage &b) Q_DECL_NOTHROW;
-Q_GUI_EXPORT bool operator!=(const QRhiGraphicsShaderStage &a, const QRhiGraphicsShaderStage &b) Q_DECL_NOTHROW;
+Q_GUI_EXPORT bool operator!=(const QRhiShaderStage &a, const QRhiShaderStage &b) Q_DECL_NOTHROW;
-Q_GUI_EXPORT uint qHash(const QRhiGraphicsShaderStage &s, uint seed = 0) Q_DECL_NOTHROW;
+Q_GUI_EXPORT uint qHash(const QRhiShaderStage &s, uint seed = 0) Q_DECL_NOTHROW;
 #ifndef QT_NO_DEBUG_STREAM
-Q_GUI_EXPORT QDebug operator<<(QDebug, const QRhiGraphicsShaderStage &);
+Q_GUI_EXPORT QDebug operator<<(QDebug, const QRhiShaderStage &);
 #endif
 
+using QRhiGraphicsShaderStage = QRhiShaderStage;
+
 class Q_GUI_EXPORT QRhiShaderResourceBinding
 {
 public:
     enum Type {
         UniformBuffer,
-        SampledTexture
+        SampledTexture,
+        ImageLoad,
+        ImageStore,
+        ImageLoadStore,
+        BufferLoad,
+        BufferStore,
+        BufferLoadStore
     };
 
     enum StageFlag {
         VertexStage = 1 << 0,
-        FragmentStage = 1 << 1
+        FragmentStage = 1 << 1,
+        ComputeStage = 1 << 2
     };
     Q_DECLARE_FLAGS(StageFlags, StageFlag)
 
@@ -320,8 +330,20 @@ public:
     static QRhiShaderResourceBinding uniformBuffer(int binding, StageFlags stage, QRhiBuffer *buf);
     static QRhiShaderResourceBinding uniformBuffer(int binding, StageFlags stage, QRhiBuffer *buf, int offset, int size);
     static QRhiShaderResourceBinding uniformBufferWithDynamicOffset(int binding, StageFlags stage, QRhiBuffer *buf, int size);
+
     static QRhiShaderResourceBinding sampledTexture(int binding, StageFlags stage, QRhiTexture *tex, QRhiSampler *sampler);
 
+    static QRhiShaderResourceBinding imageLoad(int binding, StageFlags stage, QRhiTexture *tex, int level);
+    static QRhiShaderResourceBinding imageStore(int binding, StageFlags stage, QRhiTexture *tex, int level);
+    static QRhiShaderResourceBinding imageLoadStore(int binding, StageFlags stage, QRhiTexture *tex, int level);
+
+    static QRhiShaderResourceBinding bufferLoad(int binding, StageFlags stage, QRhiBuffer *buf);
+    static QRhiShaderResourceBinding bufferLoad(int binding, StageFlags stage, QRhiBuffer *buf, int offset, int size);
+    static QRhiShaderResourceBinding bufferStore(int binding, StageFlags stage, QRhiBuffer *buf);
+    static QRhiShaderResourceBinding bufferStore(int binding, StageFlags stage, QRhiBuffer *buf, int offset, int size);
+    static QRhiShaderResourceBinding bufferLoadStore(int binding, StageFlags stage, QRhiBuffer *buf);
+    static QRhiShaderResourceBinding bufferLoadStore(int binding, StageFlags stage, QRhiBuffer *buf, int offset, int size);
+
 private:
     QRhiShaderResourceBindingPrivate *d;
     friend class QRhiShaderResourceBindingPrivate;
@@ -558,6 +580,7 @@ public:
         ShaderResourceBindings,
         GraphicsPipeline,
         SwapChain,
+        ComputePipeline,
         CommandBuffer
     };
 
@@ -594,7 +617,8 @@ public:
     enum UsageFlag {
         VertexBuffer = 1 << 0,
         IndexBuffer = 1 << 1,
-        UniformBuffer = 1 << 2
+        UniformBuffer = 1 << 2,
+        StorageBuffer = 1 << 3
     };
     Q_DECLARE_FLAGS(UsageFlags, UsageFlag)
 
@@ -629,7 +653,8 @@ public:
         MipMapped = 1 << 3,
         sRGB = 1 << 4,
         UsedAsTransferSource = 1 << 5,
-        UsedWithGenerateMips = 1 << 6
+        UsedWithGenerateMips = 1 << 6,
+        UsedWithLoadStore = 1 << 7
     };
     Q_DECLARE_FLAGS(Flags, Flag)
 
@@ -1043,8 +1068,8 @@ public:
     int sampleCount() const { return m_sampleCount; }
     void setSampleCount(int s) { m_sampleCount = s; }
 
-    QVector<QRhiGraphicsShaderStage> shaderStages() const { return m_shaderStages; }
+    QVector<QRhiShaderStage> shaderStages() const { return m_shaderStages; }
-    void setShaderStages(const QVector<QRhiGraphicsShaderStage> &stages) { m_shaderStages = stages; }
+    void setShaderStages(const QVector<QRhiShaderStage> &stages) { m_shaderStages = stages; }
 
     QRhiVertexInputLayout vertexInputLayout() const { return m_vertexInputLayout; }
     void setVertexInputLayout(const QRhiVertexInputLayout &layout) { m_vertexInputLayout = layout; }
@@ -1073,7 +1098,7 @@ protected:
     quint32 m_stencilReadMask = 0xFF;
     quint32 m_stencilWriteMask = 0xFF;
     int m_sampleCount = 1;
-    QVector<QRhiGraphicsShaderStage> m_shaderStages;
+    QVector<QRhiShaderStage> m_shaderStages;
     QRhiVertexInputLayout m_vertexInputLayout;
     QRhiShaderResourceBindings *m_shaderResourceBindings = nullptr;
     QRhiRenderPassDescriptor *m_renderPassDesc = nullptr;
@@ -1133,6 +1158,24 @@ protected:
 
 Q_DECLARE_OPERATORS_FOR_FLAGS(QRhiSwapChain::Flags)
 
+class Q_GUI_EXPORT QRhiComputePipeline : public QRhiResource
+{
+public:
+    QRhiResource::Type resourceType() const override;
+    virtual bool build() = 0;
+
+    QRhiShaderStage shaderStage() const { return m_shaderStage; }
+    void setShaderStage(const QRhiShaderStage &stage) { m_shaderStage = stage; }
+
+    QRhiShaderResourceBindings *shaderResourceBindings() const { return m_shaderResourceBindings; }
+    void setShaderResourceBindings(QRhiShaderResourceBindings *srb) { m_shaderResourceBindings = srb; }
+
+protected:
+    QRhiComputePipeline(QRhiImplementation *rhi);
+    QRhiShaderStage m_shaderStage;
+    QRhiShaderResourceBindings *m_shaderResourceBindings = nullptr;
+};
+
 class Q_GUI_EXPORT QRhiCommandBuffer : public QRhiResource
 {
 public:
@@ -1181,6 +1224,11 @@ public:
     void debugMarkEnd();
     void debugMarkMsg(const QByteArray &msg);
 
+    void beginComputePass(QRhiResourceUpdateBatch *resourceUpdates = nullptr);
+    void endComputePass(QRhiResourceUpdateBatch *resourceUpdates = nullptr);
+    void setComputePipeline(QRhiComputePipeline *ps);
+    void dispatch(int x, int y, int z);
+
     const QRhiNativeHandles *nativeHandles();
     void beginExternal();
     void endExternal();
@@ -1263,7 +1311,8 @@ public:
         NonFourAlignedEffectiveIndexBufferOffset,
         NPOTTextureRepeat,
         RedOrAlpha8IsRed,
-        ElementIndexUint
+        ElementIndexUint,
+        Compute
     };
 
     enum BeginFrameFlag {
@@ -1297,6 +1346,7 @@ public:
     void runCleanup();
 
     QRhiGraphicsPipeline *newGraphicsPipeline();
+    QRhiComputePipeline *newComputePipeline();
     QRhiShaderResourceBindings *newShaderResourceBindings();
 
     QRhiBuffer *newBuffer(QRhiBuffer::Type type,

src/gui/rhi/qrhi_p_p.h

@@ -70,6 +70,7 @@ public:
     virtual void destroy() = 0;
 
     virtual QRhiGraphicsPipeline *createGraphicsPipeline() = 0;
+    virtual QRhiComputePipeline *createComputePipeline() = 0;
     virtual QRhiShaderResourceBindings *createShaderResourceBindings() = 0;
     virtual QRhiBuffer *createBuffer(QRhiBuffer::Type type,
                                      QRhiBuffer::UsageFlags usage,
@@ -133,6 +134,11 @@ public:
     virtual void debugMarkEnd(QRhiCommandBuffer *cb) = 0;
     virtual void debugMarkMsg(QRhiCommandBuffer *cb, const QByteArray &msg) = 0;
 
+    virtual void beginComputePass(QRhiCommandBuffer *cb, QRhiResourceUpdateBatch *resourceUpdates) = 0;
+    virtual void endComputePass(QRhiCommandBuffer *cb, QRhiResourceUpdateBatch *resourceUpdates) = 0;
+    virtual void setComputePipeline(QRhiCommandBuffer *cb, QRhiComputePipeline *ps) = 0;
+    virtual void dispatch(QRhiCommandBuffer *cb, int x, int y, int z) = 0;
+
     virtual const QRhiNativeHandles *nativeHandles(QRhiCommandBuffer *cb) = 0;
     virtual void beginExternal(QRhiCommandBuffer *cb) = 0;
     virtual void endExternal(QRhiCommandBuffer *cb) = 0;
@@ -200,6 +206,7 @@ public:
 protected:
     bool debugMarkers = false;
     int currentFrameSlot = 0; // for vk, mtl, and similar. unused by gl and d3d11.
+    bool inFrame = false;
 
 private:
     QRhi::Implementation implType;
@@ -210,7 +217,6 @@ private:
     QSet<QRhiResource *> resources;
     QSet<QRhiResource *> pendingReleaseAndDestroyResources;
     QVector<QRhi::CleanupCallback> cleanupCallbacks;
-    bool inFrame = false;
 
     friend class QRhi;
     friend class QRhiResourceUpdateBatchPrivate;
@@ -393,9 +399,20 @@ public:
         QRhiTexture *tex;
         QRhiSampler *sampler;
     };
+    struct StorageImageData {
+        QRhiTexture *tex;
+        int level;
+    };
+    struct StorageBufferData {
+        QRhiBuffer *buf;
+        int offset;
+        int maybeSize;
+    };
     union {
         UniformBufferData ubuf;
         SampledTextureData stex;
+        StorageImageData simage;
+        StorageBufferData sbuf;
     } u;
 };
 
@@ -487,33 +504,41 @@ public:
     enum BufferStage {
         BufVertexInputStage,
         BufVertexStage,
-        BufFragmentStage
+        BufFragmentStage,
+        BufComputeStage
     };
 
     enum BufferAccess {
         BufVertexInput,
         BufIndexRead,
-        BufUniformRead
+        BufUniformRead,
+        BufStorageLoad,
+        BufStorageStore,
+        BufStorageLoadStore
     };
 
-    void registerBufferOnce(QRhiBuffer *buf, int slot, BufferAccess access, BufferStage stage,
+    void registerBuffer(QRhiBuffer *buf, int slot, BufferAccess *access, BufferStage *stage,
-                            const UsageState &stateAtPassBegin);
+                        const UsageState &state);
 
     enum TextureStage {
         TexVertexStage,
         TexFragmentStage,
         TexColorOutputStage,
-        TexDepthOutputStage
+        TexDepthOutputStage,
+        TexComputeStage
     };
 
     enum TextureAccess {
         TexSample,
         TexColorOutput,
-        TexDepthOutput
+        TexDepthOutput,
+        TexStorageLoad,
+        TexStorageStore,
+        TexStorageLoadStore
     };
 
-    void registerTextureOnce(QRhiTexture *tex, TextureAccess access, TextureStage stage,
+    void registerTexture(QRhiTexture *tex, TextureAccess *access, TextureStage *stage,
-                             const UsageState &stateAtPassBegin);
+                         const UsageState &state);
 
     struct Buffer {
         QRhiBuffer *buf;

src/gui/rhi/qrhid3d11.cpp

@@ -375,6 +375,8 @@ bool QRhiD3D11::isFeatureSupported(QRhi::Feature feature) const
         return true;
     case QRhi::ElementIndexUint:
         return true;
+    case QRhi::Compute:
+        return false;
     default:
         Q_UNREACHABLE();
         return false;
@@ -443,6 +445,11 @@ QRhiGraphicsPipeline *QRhiD3D11::createGraphicsPipeline()
     return new QD3D11GraphicsPipeline(this);
 }
 
+QRhiComputePipeline *QRhiD3D11::createComputePipeline()
+{
+    return new QD3D11ComputePipeline(this);
+}
+
 QRhiShaderResourceBindings *QRhiD3D11::createShaderResourceBindings()
 {
     return new QD3D11ShaderResourceBindings(this);
@@ -450,9 +457,8 @@ QRhiShaderResourceBindings *QRhiD3D11::createShaderResourceBindings()
 
 void QRhiD3D11::setGraphicsPipeline(QRhiCommandBuffer *cb, QRhiGraphicsPipeline *ps)
 {
-    Q_ASSERT(inPass);
-
     QD3D11CommandBuffer *cbD = QRHI_RES(QD3D11CommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QD3D11CommandBuffer::RenderPass);
     QD3D11GraphicsPipeline *psD = QRHI_RES(QD3D11GraphicsPipeline, ps);
     const bool pipelineChanged = cbD->currentPipeline != ps || cbD->currentPipelineGeneration != psD->generation;
 
@@ -471,9 +477,8 @@ void QRhiD3D11::setShaderResources(QRhiCommandBuffer *cb, QRhiShaderResourceBind
                                    int dynamicOffsetCount,
                                    const QRhiCommandBuffer::DynamicOffset *dynamicOffsets)
 {
-    Q_ASSERT(inPass);
-
     QD3D11CommandBuffer *cbD = QRHI_RES(QD3D11CommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QD3D11CommandBuffer::RenderPass);
     Q_ASSERT(cbD->currentPipeline);
 
     if (!srb)
@@ -568,8 +573,8 @@ void QRhiD3D11::setVertexInput(QRhiCommandBuffer *cb,
                                int startBinding, int bindingCount, const QRhiCommandBuffer::VertexInput *bindings,
                                QRhiBuffer *indexBuf, quint32 indexOffset, QRhiCommandBuffer::IndexFormat indexFormat)
 {
-    Q_ASSERT(inPass);
     QD3D11CommandBuffer *cbD = QRHI_RES(QD3D11CommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QD3D11CommandBuffer::RenderPass);
 
     bool needsBindVBuf = false;
     for (int i = 0; i < bindingCount; ++i) {
@@ -632,8 +637,8 @@ void QRhiD3D11::setVertexInput(QRhiCommandBuffer *cb,
 
 void QRhiD3D11::setViewport(QRhiCommandBuffer *cb, const QRhiViewport &viewport)
 {
-    Q_ASSERT(inPass);
     QD3D11CommandBuffer *cbD = QRHI_RES(QD3D11CommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QD3D11CommandBuffer::RenderPass);
     Q_ASSERT(cbD->currentTarget);
     const QSize outputSize = cbD->currentTarget->pixelSize();
 
@@ -656,8 +661,8 @@ void QRhiD3D11::setViewport(QRhiCommandBuffer *cb, const QRhiViewport &viewport)
 
 void QRhiD3D11::setScissor(QRhiCommandBuffer *cb, const QRhiScissor &scissor)
 {
-    Q_ASSERT(inPass);
     QD3D11CommandBuffer *cbD = QRHI_RES(QD3D11CommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QD3D11CommandBuffer::RenderPass);
     Q_ASSERT(cbD->currentTarget);
     const QSize outputSize = cbD->currentTarget->pixelSize();
 
@@ -678,8 +683,9 @@ void QRhiD3D11::setScissor(QRhiCommandBuffer *cb, const QRhiScissor &scissor)
 
 void QRhiD3D11::setBlendConstants(QRhiCommandBuffer *cb, const QColor &c)
 {
-    Q_ASSERT(inPass);
     QD3D11CommandBuffer *cbD = QRHI_RES(QD3D11CommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QD3D11CommandBuffer::RenderPass);
+
     QD3D11CommandBuffer::Command cmd;
     cmd.cmd = QD3D11CommandBuffer::Command::BlendConstants;
     cmd.args.blendConstants.ps = QRHI_RES(QD3D11GraphicsPipeline, cbD->currentPipeline);
@@ -692,8 +698,9 @@ void QRhiD3D11::setBlendConstants(QRhiCommandBuffer *cb, const QColor &c)
 
 void QRhiD3D11::setStencilRef(QRhiCommandBuffer *cb, quint32 refValue)
 {
-    Q_ASSERT(inPass);
     QD3D11CommandBuffer *cbD = QRHI_RES(QD3D11CommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QD3D11CommandBuffer::RenderPass);
+
     QD3D11CommandBuffer::Command cmd;
     cmd.cmd = QD3D11CommandBuffer::Command::StencilRef;
     cmd.args.stencilRef.ps = QRHI_RES(QD3D11GraphicsPipeline, cbD->currentPipeline);
@@ -704,8 +711,9 @@ void QRhiD3D11::setStencilRef(QRhiCommandBuffer *cb, quint32 refValue)
 void QRhiD3D11::draw(QRhiCommandBuffer *cb, quint32 vertexCount,
                      quint32 instanceCount, quint32 firstVertex, quint32 firstInstance)
 {
-    Q_ASSERT(inPass);
     QD3D11CommandBuffer *cbD = QRHI_RES(QD3D11CommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QD3D11CommandBuffer::RenderPass);
+
     QD3D11CommandBuffer::Command cmd;
     cmd.cmd = QD3D11CommandBuffer::Command::Draw;
     cmd.args.draw.ps = QRHI_RES(QD3D11GraphicsPipeline, cbD->currentPipeline);
@@ -719,8 +727,9 @@ void QRhiD3D11::draw(QRhiCommandBuffer *cb, quint32 vertexCount,
 void QRhiD3D11::drawIndexed(QRhiCommandBuffer *cb, quint32 indexCount,
                             quint32 instanceCount, quint32 firstIndex, qint32 vertexOffset, quint32 firstInstance)
 {
-    Q_ASSERT(inPass);
     QD3D11CommandBuffer *cbD = QRHI_RES(QD3D11CommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QD3D11CommandBuffer::RenderPass);
+
     QD3D11CommandBuffer::Command cmd;
     cmd.cmd = QD3D11CommandBuffer::Command::DrawIndexed;
     cmd.args.drawIndexed.ps = QRHI_RES(QD3D11GraphicsPipeline, cbD->currentPipeline);
@@ -777,28 +786,23 @@ const QRhiNativeHandles *QRhiD3D11::nativeHandles(QRhiCommandBuffer *cb)
 
 void QRhiD3D11::beginExternal(QRhiCommandBuffer *cb)
 {
-    Q_ASSERT(inPass);
     Q_UNUSED(cb);
     flushCommandBuffer();
 }
 
 void QRhiD3D11::endExternal(QRhiCommandBuffer *cb)
 {
-    Q_ASSERT(inPass);
     QD3D11CommandBuffer *cbD = QRHI_RES(QD3D11CommandBuffer, cb);
-    Q_ASSERT(cbD->currentTarget);
     Q_ASSERT(cbD->commands.isEmpty());
     cbD->resetCachedState();
-    enqueueSetRenderTarget(cbD, cbD->currentTarget);
+    if (cbD->currentTarget) // could be compute, no rendertarget then
+        enqueueSetRenderTarget(cbD, cbD->currentTarget);
 }
 
 QRhi::FrameOpResult QRhiD3D11::beginFrame(QRhiSwapChain *swapChain, QRhi::BeginFrameFlags flags)
 {
     Q_UNUSED(flags);
 
-    Q_ASSERT(!inFrame);
-    inFrame = true;
-
     QD3D11SwapChain *swapChainD = QRHI_RES(QD3D11SwapChain, swapChain);
     contextState.currentSwapChain = swapChainD;
     const int currentFrameSlot = swapChainD->currentFrameSlot;
@@ -844,9 +848,6 @@ QRhi::FrameOpResult QRhiD3D11::beginFrame(QRhiSwapChain *swapChain, QRhi::BeginF
 
 QRhi::FrameOpResult QRhiD3D11::endFrame(QRhiSwapChain *swapChain, QRhi::EndFrameFlags flags)
 {
-    Q_ASSERT(inFrame);
-    inFrame = false;
-
     QD3D11SwapChain *swapChainD = QRHI_RES(QD3D11SwapChain, swapChain);
     Q_ASSERT(contextState.currentSwapChain = swapChainD);
     const int currentFrameSlot = swapChainD->currentFrameSlot;
@@ -899,8 +900,6 @@ QRhi::FrameOpResult QRhiD3D11::endFrame(QRhiSwapChain *swapChain, QRhi::EndFrame
 
 QRhi::FrameOpResult QRhiD3D11::beginOffscreenFrame(QRhiCommandBuffer **cb)
 {
-    Q_ASSERT(!inFrame);
-    inFrame = true;
     ofr.active = true;
 
     ofr.cbWrapper.resetState();
@@ -911,8 +910,6 @@ QRhi::FrameOpResult QRhiD3D11::beginOffscreenFrame(QRhiCommandBuffer **cb)
 
 QRhi::FrameOpResult QRhiD3D11::endOffscreenFrame()
 {
-    Q_ASSERT(inFrame && ofr.active);
-    inFrame = false;
     ofr.active = false;
 
     executeCommandBuffer(&ofr.cbWrapper);
@@ -1047,8 +1044,6 @@ static inline bool isDepthTextureFormat(QRhiTexture::Format format)
 
 QRhi::FrameOpResult QRhiD3D11::finish()
 {
-    Q_ASSERT(!inPass);
-
     if (inFrame)
         flushCommandBuffer();
 
@@ -1379,7 +1374,7 @@ static inline QD3D11RenderTargetData *rtData(QRhiRenderTarget *rt)
 
 void QRhiD3D11::resourceUpdate(QRhiCommandBuffer *cb, QRhiResourceUpdateBatch *resourceUpdates)
 {
-    Q_ASSERT(inFrame && !inPass);
+    Q_ASSERT(QRHI_RES(QD3D11CommandBuffer, cb)->recordingPass == QD3D11CommandBuffer::NoPass);
 
     enqueueResourceUpdates(cb, resourceUpdates);
 }
@@ -1398,12 +1393,12 @@ void QRhiD3D11::beginPass(QRhiCommandBuffer *cb,
                           const QRhiDepthStencilClearValue &depthStencilClearValue,
                           QRhiResourceUpdateBatch *resourceUpdates)
 {
-    Q_ASSERT(inFrame && !inPass);
+    QD3D11CommandBuffer *cbD = QRHI_RES(QD3D11CommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QD3D11CommandBuffer::NoPass);
 
     if (resourceUpdates)
         enqueueResourceUpdates(cb, resourceUpdates);
 
-    QD3D11CommandBuffer *cbD = QRHI_RES(QD3D11CommandBuffer, cb);
     bool wantsColorClear = true;
     bool wantsDsClear = true;
     QD3D11RenderTargetData *rtD = rtData(rt);
@@ -1431,17 +1426,15 @@ void QRhiD3D11::beginPass(QRhiCommandBuffer *cb,
     clearCmd.args.clear.s = depthStencilClearValue.stencilClearValue();
     cbD->commands.append(clearCmd);
 
+    cbD->recordingPass = QD3D11CommandBuffer::RenderPass;
     cbD->currentTarget = rt;
-
-    inPass = true;
 }
 
 void QRhiD3D11::endPass(QRhiCommandBuffer *cb, QRhiResourceUpdateBatch *resourceUpdates)
 {
-    Q_ASSERT(inPass);
-    inPass = false;
-
     QD3D11CommandBuffer *cbD = QRHI_RES(QD3D11CommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QD3D11CommandBuffer::RenderPass);
+
     if (cbD->currentTarget->resourceType() == QRhiResource::TextureRenderTarget) {
         QD3D11TextureRenderTarget *rtTex = QRHI_RES(QD3D11TextureRenderTarget, cbD->currentTarget);
         const QVector<QRhiColorAttachment> colorAttachments = rtTex->m_desc.colorAttachments();
@@ -1491,12 +1484,49 @@ void QRhiD3D11::endPass(QRhiCommandBuffer *cb, QRhiResourceUpdateBatch *resource
         }
     }
 
+    cbD->recordingPass = QD3D11CommandBuffer::NoPass;
     cbD->currentTarget = nullptr;
 
     if (resourceUpdates)
         enqueueResourceUpdates(cb, resourceUpdates);
 }
 
+void QRhiD3D11::beginComputePass(QRhiCommandBuffer *cb, QRhiResourceUpdateBatch *resourceUpdates)
+{
+    QD3D11CommandBuffer *cbD = QRHI_RES(QD3D11CommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QD3D11CommandBuffer::NoPass);
+
+    if (resourceUpdates)
+        enqueueResourceUpdates(cb, resourceUpdates);
+
+    cbD->recordingPass = QD3D11CommandBuffer::ComputePass;
+}
+
+void QRhiD3D11::endComputePass(QRhiCommandBuffer *cb, QRhiResourceUpdateBatch *resourceUpdates)
+{
+    QD3D11CommandBuffer *cbD = QRHI_RES(QD3D11CommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QD3D11CommandBuffer::ComputePass);
+
+    cbD->recordingPass = QD3D11CommandBuffer::NoPass;
+
+    if (resourceUpdates)
+        enqueueResourceUpdates(cb, resourceUpdates);
+}
+
+void QRhiD3D11::setComputePipeline(QRhiCommandBuffer *cb, QRhiComputePipeline *ps)
+{
+    Q_UNUSED(cb);
+    Q_UNUSED(ps);
+}
+
+void QRhiD3D11::dispatch(QRhiCommandBuffer *cb, int x, int y, int z)
+{
+    Q_UNUSED(cb);
+    Q_UNUSED(x);
+    Q_UNUSED(y);
+    Q_UNUSED(z);
+}
+
 void QRhiD3D11::updateShaderResourceBindings(QD3D11ShaderResourceBindings *srbD)
 {
     srbD->vsubufs.clear();
@@ -1709,6 +1739,8 @@ void QRhiD3D11::setRenderTarget(QRhiRenderTarget *rt)
 
 void QRhiD3D11::executeCommandBuffer(QD3D11CommandBuffer *cbD, QD3D11SwapChain *timestampSwapChain)
 {
+    Q_ASSERT(cbD->recordingPass == QD3D11CommandBuffer::NoPass);
+
     quint32 stencilRef = 0;
     float blendConstants[] = { 1, 1, 1, 1 };
 
@@ -1911,6 +1943,11 @@ bool QD3D11Buffer::build()
     if (buffer)
         release();
 
+    if (m_usage.testFlag(QRhiBuffer::UniformBuffer) && m_type != Dynamic) {
+        qWarning("UniformBuffer must always be combined with Dynamic on D3D11");
+        return false;
+    }
+
     const int nonZeroSize = m_size <= 0 ? 256 : m_size;
     const int roundedSize = m_usage.testFlag(QRhiBuffer::UniformBuffer) ? aligned(nonZeroSize, 256) : nonZeroSize;
 
@@ -3008,7 +3045,7 @@ bool QD3D11GraphicsPipeline::build()
     }
 
     QByteArray vsByteCode;
-    for (const QRhiGraphicsShaderStage &shaderStage : qAsConst(m_shaderStages)) {
+    for (const QRhiShaderStage &shaderStage : qAsConst(m_shaderStages)) {
         QString error;
         QByteArray bytecode = compileHlslShaderSource(shaderStage.shader(), shaderStage.shaderVariant(), &error);
         if (bytecode.isEmpty()) {
@@ -3016,7 +3053,7 @@ bool QD3D11GraphicsPipeline::build()
             return false;
         }
         switch (shaderStage.type()) {
-        case QRhiGraphicsShaderStage::Vertex:
+        case QRhiShaderStage::Vertex:
             hr = rhiD->dev->CreateVertexShader(bytecode.constData(), bytecode.size(), nullptr, &vs);
             if (FAILED(hr)) {
                 qWarning("Failed to create vertex shader: %s", qPrintable(comErrorMessage(hr)));
@@ -3024,7 +3061,7 @@ bool QD3D11GraphicsPipeline::build()
             }
             vsByteCode = bytecode;
             break;
-        case QRhiGraphicsShaderStage::Fragment:
+        case QRhiShaderStage::Fragment:
             hr = rhiD->dev->CreatePixelShader(bytecode.constData(), bytecode.size(), nullptr, &fs);
             if (FAILED(hr)) {
                 qWarning("Failed to create pixel shader: %s", qPrintable(comErrorMessage(hr)));
@@ -3072,6 +3109,25 @@ bool QD3D11GraphicsPipeline::build()
     return true;
 }
 
+QD3D11ComputePipeline::QD3D11ComputePipeline(QRhiImplementation *rhi)
+    : QRhiComputePipeline(rhi)
+{
+}
+
+QD3D11ComputePipeline::~QD3D11ComputePipeline()
+{
+    release();
+}
+
+void QD3D11ComputePipeline::release()
+{
+}
+
+bool QD3D11ComputePipeline::build()
+{
+    return false;
+}
+
 QD3D11CommandBuffer::QD3D11CommandBuffer(QRhiImplementation *rhi)
     : QRhiCommandBuffer(rhi)
 {

src/gui/rhi/qrhid3d11_p_p.h

@@ -254,6 +254,14 @@ struct QD3D11GraphicsPipeline : public QRhiGraphicsPipeline
     friend class QRhiD3D11;
 };
 
+struct QD3D11ComputePipeline : public QRhiComputePipeline
+{
+    QD3D11ComputePipeline(QRhiImplementation *rhi);
+    ~QD3D11ComputePipeline();
+    void release() override;
+    bool build() override;
+};
+
 struct QD3D11SwapChain;
 
 struct QD3D11CommandBuffer : public QRhiCommandBuffer
@@ -387,7 +395,14 @@ struct QD3D11CommandBuffer : public QRhiCommandBuffer
         } args;
     };
 
+    enum PassType {
+        NoPass,
+        RenderPass,
+        ComputePass
+    };
+
     QVector<Command> commands;
+    PassType recordingPass;
     QRhiRenderTarget *currentTarget;
     QRhiGraphicsPipeline *currentPipeline;
     uint currentPipelineGeneration;
@@ -418,6 +433,7 @@ struct QD3D11CommandBuffer : public QRhiCommandBuffer
     }
     void resetState() {
         resetCommands();
+        recordingPass = NoPass;
         currentTarget = nullptr;
         resetCachedState();
     }
@@ -484,6 +500,7 @@ public:
     void destroy() override;
 
     QRhiGraphicsPipeline *createGraphicsPipeline() override;
+    QRhiComputePipeline *createComputePipeline() override;
     QRhiShaderResourceBindings *createShaderResourceBindings() override;
     QRhiBuffer *createBuffer(QRhiBuffer::Type type,
                              QRhiBuffer::UsageFlags usage,
@@ -548,6 +565,11 @@ public:
     void debugMarkEnd(QRhiCommandBuffer *cb) override;
     void debugMarkMsg(QRhiCommandBuffer *cb, const QByteArray &msg) override;
 
+    void beginComputePass(QRhiCommandBuffer *cb, QRhiResourceUpdateBatch *resourceUpdates) override;
+    void endComputePass(QRhiCommandBuffer *cb, QRhiResourceUpdateBatch *resourceUpdates) override;
+    void setComputePipeline(QRhiCommandBuffer *cb, QRhiComputePipeline *ps) override;
+    void dispatch(QRhiCommandBuffer *cb, int x, int y, int z) override;
+
     const QRhiNativeHandles *nativeHandles(QRhiCommandBuffer *cb) override;
     void beginExternal(QRhiCommandBuffer *cb) override;
     void endExternal(QRhiCommandBuffer *cb) override;
@@ -591,9 +613,6 @@ public:
     bool hasDxgi2 = false;
     QRhiD3D11NativeHandles nativeHandlesStruct;
 
-    bool inFrame = false;
-    bool inPass = false;
-
     struct {
         int vsHighestActiveSrvBinding = -1;
         int fsHighestActiveSrvBinding = -1;

src/gui/rhi/qrhigles2.cpp

@@ -616,6 +616,8 @@ bool QRhiGles2::isFeatureSupported(QRhi::Feature feature) const
         return caps.coreProfile;
     case QRhi::ElementIndexUint:
         return caps.elementIndexUint;
+    case QRhi::Compute:
+        return false;
     default:
         Q_UNREACHABLE();
         return false;
@@ -692,11 +694,15 @@ QRhiShaderResourceBindings *QRhiGles2::createShaderResourceBindings()
     return new QGles2ShaderResourceBindings(this);
 }
 
+QRhiComputePipeline *QRhiGles2::createComputePipeline()
+{
+    return new QGles2ComputePipeline(this);
+}
+
 void QRhiGles2::setGraphicsPipeline(QRhiCommandBuffer *cb, QRhiGraphicsPipeline *ps)
 {
-    Q_ASSERT(inPass);
-
     QGles2CommandBuffer *cbD = QRHI_RES(QGles2CommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QGles2CommandBuffer::RenderPass);
     QGles2GraphicsPipeline *psD = QRHI_RES(QGles2GraphicsPipeline, ps);
     const bool pipelineChanged = cbD->currentPipeline != ps || cbD->currentPipelineGeneration != psD->generation;
 
@@ -715,9 +721,8 @@ void QRhiGles2::setShaderResources(QRhiCommandBuffer *cb, QRhiShaderResourceBind
                                    int dynamicOffsetCount,
                                    const QRhiCommandBuffer::DynamicOffset *dynamicOffsets)
 {
-    Q_ASSERT(inPass);
-
     QGles2CommandBuffer *cbD = QRHI_RES(QGles2CommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QGles2CommandBuffer::RenderPass);
     Q_ASSERT(cbD->currentPipeline);
 
     if (!srb)
@@ -770,8 +775,8 @@ void QRhiGles2::setVertexInput(QRhiCommandBuffer *cb,
                                int startBinding, int bindingCount, const QRhiCommandBuffer::VertexInput *bindings,
                                QRhiBuffer *indexBuf, quint32 indexOffset, QRhiCommandBuffer::IndexFormat indexFormat)
 {
-    Q_ASSERT(inPass);
     QGles2CommandBuffer *cbD = QRHI_RES(QGles2CommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QGles2CommandBuffer::RenderPass);
 
     for (int i = 0; i < bindingCount; ++i) {
         QRhiBuffer *buf = bindings[i].first;
@@ -801,7 +806,9 @@ void QRhiGles2::setVertexInput(QRhiCommandBuffer *cb,
 
 void QRhiGles2::setViewport(QRhiCommandBuffer *cb, const QRhiViewport &viewport)
 {
-    Q_ASSERT(inPass);
+    QGles2CommandBuffer *cbD = QRHI_RES(QGles2CommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QGles2CommandBuffer::RenderPass);
+
     QGles2CommandBuffer::Command cmd;
     cmd.cmd = QGles2CommandBuffer::Command::Viewport;
     const std::array<float, 4> r = viewport.viewport();
@@ -811,12 +818,14 @@ void QRhiGles2::setViewport(QRhiCommandBuffer *cb, const QRhiViewport &viewport)
     cmd.args.viewport.h = r[3];
     cmd.args.viewport.d0 = viewport.minDepth();
     cmd.args.viewport.d1 = viewport.maxDepth();
-    QRHI_RES(QGles2CommandBuffer, cb)->commands.append(cmd);
+    cbD->commands.append(cmd);
 }
 
 void QRhiGles2::setScissor(QRhiCommandBuffer *cb, const QRhiScissor &scissor)
 {
-    Q_ASSERT(inPass);
+    QGles2CommandBuffer *cbD = QRHI_RES(QGles2CommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QGles2CommandBuffer::RenderPass);
+
     QGles2CommandBuffer::Command cmd;
     cmd.cmd = QGles2CommandBuffer::Command::Scissor;
     const std::array<int, 4> r = scissor.scissor();
@@ -824,25 +833,27 @@ void QRhiGles2::setScissor(QRhiCommandBuffer *cb, const QRhiScissor &scissor)
     cmd.args.scissor.y = qMax(0, r[1]);
     cmd.args.scissor.w = r[2];
     cmd.args.scissor.h = r[3];
-    QRHI_RES(QGles2CommandBuffer, cb)->commands.append(cmd);
+    cbD->commands.append(cmd);
 }
 
 void QRhiGles2::setBlendConstants(QRhiCommandBuffer *cb, const QColor &c)
 {
-    Q_ASSERT(inPass);
+    QGles2CommandBuffer *cbD = QRHI_RES(QGles2CommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QGles2CommandBuffer::RenderPass);
+
     QGles2CommandBuffer::Command cmd;
     cmd.cmd = QGles2CommandBuffer::Command::BlendConstants;
     cmd.args.blendConstants.r = c.redF();
     cmd.args.blendConstants.g = c.greenF();
     cmd.args.blendConstants.b = c.blueF();
     cmd.args.blendConstants.a = c.alphaF();
-    QRHI_RES(QGles2CommandBuffer, cb)->commands.append(cmd);
+    cbD->commands.append(cmd);
 }
 
 void QRhiGles2::setStencilRef(QRhiCommandBuffer *cb, quint32 refValue)
 {
-    Q_ASSERT(inPass);
     QGles2CommandBuffer *cbD = QRHI_RES(QGles2CommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QGles2CommandBuffer::RenderPass);
 
     QGles2CommandBuffer::Command cmd;
     cmd.cmd = QGles2CommandBuffer::Command::StencilRef;
@@ -854,10 +865,10 @@ void QRhiGles2::setStencilRef(QRhiCommandBuffer *cb, quint32 refValue)
 void QRhiGles2::draw(QRhiCommandBuffer *cb, quint32 vertexCount,
                      quint32 instanceCount, quint32 firstVertex, quint32 firstInstance)
 {
-    Q_ASSERT(inPass);
     Q_UNUSED(instanceCount); // no instancing
     Q_UNUSED(firstInstance);
     QGles2CommandBuffer *cbD = QRHI_RES(QGles2CommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QGles2CommandBuffer::RenderPass);
 
     QGles2CommandBuffer::Command cmd;
     cmd.cmd = QGles2CommandBuffer::Command::Draw;
@@ -870,11 +881,11 @@ void QRhiGles2::draw(QRhiCommandBuffer *cb, quint32 vertexCount,
 void QRhiGles2::drawIndexed(QRhiCommandBuffer *cb, quint32 indexCount,
                             quint32 instanceCount, quint32 firstIndex, qint32 vertexOffset, quint32 firstInstance)
 {
-    Q_ASSERT(inPass);
     Q_UNUSED(instanceCount); // no instancing
     Q_UNUSED(firstInstance);
     Q_UNUSED(vertexOffset); // no glDrawElementsBaseVertex
     QGles2CommandBuffer *cbD = QRHI_RES(QGles2CommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QGles2CommandBuffer::RenderPass);
 
     QGles2CommandBuffer::Command cmd;
     cmd.cmd = QGles2CommandBuffer::Command::DrawIndexed;
@@ -918,19 +929,17 @@ const QRhiNativeHandles *QRhiGles2::nativeHandles(QRhiCommandBuffer *cb)
 
 void QRhiGles2::beginExternal(QRhiCommandBuffer *cb)
 {
-    Q_ASSERT(inPass);
     Q_UNUSED(cb);
     flushCommandBuffer(); // also ensures the context is current
 }
 
 void QRhiGles2::endExternal(QRhiCommandBuffer *cb)
 {
-    Q_ASSERT(inPass);
     QGles2CommandBuffer *cbD = QRHI_RES(QGles2CommandBuffer, cb);
-    Q_ASSERT(cbD->currentTarget);
     Q_ASSERT(cbD->commands.isEmpty());
     cbD->resetCachedState();
-    enqueueBindFramebuffer(cbD->currentTarget, cbD);
+    if (cbD->currentTarget)
+        enqueueBindFramebuffer(cbD->currentTarget, cbD);
 }
 
 static void addBoundaryCommand(QGles2CommandBuffer *cb, QGles2CommandBuffer::Command::Cmd type)
@@ -943,13 +952,11 @@ static void addBoundaryCommand(QGles2CommandBuffer *cb, QGles2CommandBuffer::Com
 QRhi::FrameOpResult QRhiGles2::beginFrame(QRhiSwapChain *swapChain, QRhi::BeginFrameFlags flags)
 {
     Q_UNUSED(flags);
-    Q_ASSERT(!inFrame);
 
     QGles2SwapChain *swapChainD = QRHI_RES(QGles2SwapChain, swapChain);
     if (!ensureContext(swapChainD->surface))
         return QRhi::FrameOpError;
 
-    inFrame = true;
     currentSwapChain = swapChainD;
 
     QRhiProfilerPrivate *rhiP = profilerPrivateOrNull();
@@ -965,9 +972,6 @@ QRhi::FrameOpResult QRhiGles2::beginFrame(QRhiSwapChain *swapChain, QRhi::BeginF
 
 QRhi::FrameOpResult QRhiGles2::endFrame(QRhiSwapChain *swapChain, QRhi::EndFrameFlags flags)
 {
-    Q_ASSERT(inFrame);
-    inFrame = false;
-
     QGles2SwapChain *swapChainD = QRHI_RES(QGles2SwapChain, swapChain);
     Q_ASSERT(currentSwapChain == swapChainD);
 
@@ -996,12 +1000,9 @@ QRhi::FrameOpResult QRhiGles2::endFrame(QRhiSwapChain *swapChain, QRhi::EndFrame
 
 QRhi::FrameOpResult QRhiGles2::beginOffscreenFrame(QRhiCommandBuffer **cb)
 {
-    Q_ASSERT(!inFrame);
-
     if (!ensureContext())
         return QRhi::FrameOpError;
 
-    inFrame = true;
     ofr.active = true;
 
     executeDeferredReleases();
@@ -1015,8 +1016,7 @@ QRhi::FrameOpResult QRhiGles2::beginOffscreenFrame(QRhiCommandBuffer **cb)
 
 QRhi::FrameOpResult QRhiGles2::endOffscreenFrame()
 {
-    Q_ASSERT(inFrame && ofr.active);
+    Q_ASSERT(ofr.active);
-    inFrame = false;
     ofr.active = false;
 
     addBoundaryCommand(&ofr.cbWrapper, QGles2CommandBuffer::Command::EndFrame);
@@ -1031,7 +1031,6 @@ QRhi::FrameOpResult QRhiGles2::endOffscreenFrame()
 
 QRhi::FrameOpResult QRhiGles2::finish()
 {
-    Q_ASSERT(!inPass); // because that's what the QRhi docs say, even though not required by this backend
     return inFrame ? flushCommandBuffer() : QRhi::FrameOpSuccess;
 }
 
@@ -1501,6 +1500,8 @@ static inline GLenum toGlTextureCompareFunc(QRhiSampler::CompareOp op)
 void QRhiGles2::executeCommandBuffer(QRhiCommandBuffer *cb)
 {
     QGles2CommandBuffer *cbD = QRHI_RES(QGles2CommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QGles2CommandBuffer::NoPass);
+
     GLenum indexType = GL_UNSIGNED_SHORT;
     quint32 indexStride = sizeof(quint16);
     quint32 indexOffset = 0;
@@ -1970,7 +1971,7 @@ void QRhiGles2::bindShaderResources(QRhiGraphicsPipeline *ps, QRhiShaderResource
 
 void QRhiGles2::resourceUpdate(QRhiCommandBuffer *cb, QRhiResourceUpdateBatch *resourceUpdates)
 {
-    Q_ASSERT(inFrame && !inPass);
+    Q_ASSERT(QRHI_RES(QGles2CommandBuffer, cb)->recordingPass == QGles2CommandBuffer::NoPass);
 
     enqueueResourceUpdates(cb, resourceUpdates);
 }
@@ -2018,12 +2019,12 @@ void QRhiGles2::beginPass(QRhiCommandBuffer *cb,
                           const QRhiDepthStencilClearValue &depthStencilClearValue,
                           QRhiResourceUpdateBatch *resourceUpdates)
 {
-    Q_ASSERT(inFrame && !inPass);
+    QGles2CommandBuffer *cbD = QRHI_RES(QGles2CommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QGles2CommandBuffer::NoPass);
 
     if (resourceUpdates)
         enqueueResourceUpdates(cb, resourceUpdates);
 
-    QGles2CommandBuffer *cbD = QRHI_RES(QGles2CommandBuffer, cb);
     bool wantsColorClear, wantsDsClear;
     QGles2RenderTargetData *rtD = enqueueBindFramebuffer(rt, cbD, &wantsColorClear, &wantsDsClear);
 
@@ -2042,17 +2043,15 @@ void QRhiGles2::beginPass(QRhiCommandBuffer *cb,
     clearCmd.args.clear.s = depthStencilClearValue.stencilClearValue();
     cbD->commands.append(clearCmd);
 
+    cbD->recordingPass = QGles2CommandBuffer::RenderPass;
     cbD->currentTarget = rt;
-
-    inPass = true;
 }
 
 void QRhiGles2::endPass(QRhiCommandBuffer *cb, QRhiResourceUpdateBatch *resourceUpdates)
 {
-    Q_ASSERT(inPass);
-    inPass = false;
-
     QGles2CommandBuffer *cbD = QRHI_RES(QGles2CommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QGles2CommandBuffer::RenderPass);
+
     if (cbD->currentTarget->resourceType() == QRhiResource::TextureRenderTarget) {
         QGles2TextureRenderTarget *rtTex = QRHI_RES(QGles2TextureRenderTarget, cbD->currentTarget);
         const QVector<QRhiColorAttachment> colorAttachments = rtTex->m_desc.colorAttachments();
@@ -2083,12 +2082,49 @@ void QRhiGles2::endPass(QRhiCommandBuffer *cb, QRhiResourceUpdateBatch *resource
         }
     }
 
+    cbD->recordingPass = QGles2CommandBuffer::NoPass;
     cbD->currentTarget = nullptr;
 
     if (resourceUpdates)
         enqueueResourceUpdates(cb, resourceUpdates);
 }
 
+void QRhiGles2::beginComputePass(QRhiCommandBuffer *cb, QRhiResourceUpdateBatch *resourceUpdates)
+{
+    QGles2CommandBuffer *cbD = QRHI_RES(QGles2CommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QGles2CommandBuffer::NoPass);
+
+    if (resourceUpdates)
+        enqueueResourceUpdates(cb, resourceUpdates);
+
+    cbD->recordingPass = QGles2CommandBuffer::ComputePass;
+}
+
+void QRhiGles2::endComputePass(QRhiCommandBuffer *cb, QRhiResourceUpdateBatch *resourceUpdates)
+{
+    QGles2CommandBuffer *cbD = QRHI_RES(QGles2CommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QGles2CommandBuffer::ComputePass);
+
+    cbD->recordingPass = QGles2CommandBuffer::NoPass;
+
+    if (resourceUpdates)
+        enqueueResourceUpdates(cb, resourceUpdates);
+}
+
+void QRhiGles2::setComputePipeline(QRhiCommandBuffer *cb, QRhiComputePipeline *ps)
+{
+    Q_UNUSED(cb);
+    Q_UNUSED(ps);
+}
+
+void QRhiGles2::dispatch(QRhiCommandBuffer *cb, int x, int y, int z)
+{
+    Q_UNUSED(cb);
+    Q_UNUSED(x);
+    Q_UNUSED(y);
+    Q_UNUSED(z);
+}
+
 QGles2Buffer::QGles2Buffer(QRhiImplementation *rhi, Type type, UsageFlags usage, int size)
     : QRhiBuffer(rhi, type, usage, size)
 {
@@ -2742,9 +2778,9 @@ bool QGles2GraphicsPipeline::build()
     program = rhiD->f->glCreateProgram();
 
     int sourceVer = 0;
-    for (const QRhiGraphicsShaderStage &shaderStage : qAsConst(m_shaderStages)) {
+    for (const QRhiShaderStage &shaderStage : qAsConst(m_shaderStages)) {
-        const bool isVertex = shaderStage.type() == QRhiGraphicsShaderStage::Vertex;
+        const bool isVertex = shaderStage.type() == QRhiShaderStage::Vertex;
-        const bool isFragment = shaderStage.type() == QRhiGraphicsShaderStage::Fragment;
+        const bool isFragment = shaderStage.type() == QRhiShaderStage::Fragment;
         if (!isVertex && !isFragment)
             continue;
 
@@ -2902,6 +2938,25 @@ bool QGles2GraphicsPipeline::build()
     return true;
 }
 
+QGles2ComputePipeline::QGles2ComputePipeline(QRhiImplementation *rhi)
+    : QRhiComputePipeline(rhi)
+{
+}
+
+QGles2ComputePipeline::~QGles2ComputePipeline()
+{
+    release();
+}
+
+void QGles2ComputePipeline::release()
+{
+}
+
+bool QGles2ComputePipeline::build()
+{
+    return false;
+}
+
 QGles2CommandBuffer::QGles2CommandBuffer(QRhiImplementation *rhi)
     : QRhiCommandBuffer(rhi)
 {

src/gui/rhi/qrhigles2_p_p.h

@@ -248,6 +248,14 @@ struct QGles2GraphicsPipeline : public QRhiGraphicsPipeline
 Q_DECLARE_TYPEINFO(QGles2GraphicsPipeline::Uniform, Q_MOVABLE_TYPE);
 Q_DECLARE_TYPEINFO(QGles2GraphicsPipeline::Sampler, Q_MOVABLE_TYPE);
 
+struct QGles2ComputePipeline : public QRhiComputePipeline
+{
+    QGles2ComputePipeline(QRhiImplementation *rhi);
+    ~QGles2ComputePipeline();
+    void release() override;
+    bool build() override;
+};
+
 struct QGles2CommandBuffer : public QRhiCommandBuffer
 {
     QGles2CommandBuffer(QRhiImplementation *rhi);
@@ -426,7 +434,14 @@ struct QGles2CommandBuffer : public QRhiCommandBuffer
         } args;
     };
 
+    enum PassType {
+        NoPass,
+        RenderPass,
+        ComputePass
+    };
+
     QVector<Command> commands;
+    PassType recordingPass;
     QRhiRenderTarget *currentTarget;
     QRhiGraphicsPipeline *currentPipeline;
     uint currentPipelineGeneration;
@@ -452,6 +467,7 @@ struct QGles2CommandBuffer : public QRhiCommandBuffer
     }
     void resetState() {
         resetCommands();
+        recordingPass = NoPass;
         currentTarget = nullptr;
         resetCachedState();
     }
@@ -495,6 +511,7 @@ public:
     void destroy() override;
 
     QRhiGraphicsPipeline *createGraphicsPipeline() override;
+    QRhiComputePipeline *createComputePipeline() override;
     QRhiShaderResourceBindings *createShaderResourceBindings() override;
     QRhiBuffer *createBuffer(QRhiBuffer::Type type,
                              QRhiBuffer::UsageFlags usage,
@@ -559,6 +576,11 @@ public:
     void debugMarkEnd(QRhiCommandBuffer *cb) override;
     void debugMarkMsg(QRhiCommandBuffer *cb, const QByteArray &msg) override;
 
+    void beginComputePass(QRhiCommandBuffer *cb, QRhiResourceUpdateBatch *resourceUpdates) override;
+    void endComputePass(QRhiCommandBuffer *cb, QRhiResourceUpdateBatch *resourceUpdates) override;
+    void setComputePipeline(QRhiCommandBuffer *cb, QRhiComputePipeline *ps) override;
+    void dispatch(QRhiCommandBuffer *cb, int x, int y, int z) override;
+
     const QRhiNativeHandles *nativeHandles(QRhiCommandBuffer *cb) override;
     void beginExternal(QRhiCommandBuffer *cb) override;
     void endExternal(QRhiCommandBuffer *cb) override;
@@ -645,8 +667,6 @@ public:
         uint uniformBuffers : 1;
         uint elementIndexUint : 1;
     } caps;
-    bool inFrame = false;
-    bool inPass = false;
     QGles2SwapChain *currentSwapChain = nullptr;
     QVector<GLint> supportedCompressedFormats;
     mutable QVector<int> supportedSampleCountList;

src/gui/rhi/qrhimetal.mm

@@ -36,6 +36,7 @@
 
 #include "qrhimetal_p_p.h"
 #include "qshader_p.h"
+#include "qshaderdescription_p.h"
 #include <QGuiApplication>
 #include <QWindow>
 #include <qmath.h>
@@ -51,14 +52,12 @@ QT_BEGIN_NAMESPACE
 
 /*
     Metal backend. Double buffers and throttles to vsync. "Dynamic" buffers are
-    Shared (host visible) and duplicated (due to 2 frames in flight), "static"
+    Shared (host visible) and duplicated (to help having 2 frames in flight),
-    are Managed on macOS and Shared on iOS/tvOS, and still duplicated.
+    "static" and "immutable" are Managed on macOS and Shared on iOS/tvOS.
-    "Immutable" is like "static" but with only one native buffer underneath.
     Textures are Private (device local) and a host visible staging buffer is
     used to upload data to them. Does not rely on strong objects refs from
-    command buffers (hence uses commandBufferWithUnretainedReferences), but
+    command buffers but does rely on the automatic resource tracking of the
-    does rely on automatic dependency tracking between encoders (hence no
+    command encoders.
-    MTLResourceHazardTrackingModeUntracked atm).
 */
 
 #if __has_feature(objc_arc)
@@ -173,6 +172,7 @@ struct QRhiMetalData
             struct {
                 id<MTLTexture> texture;
                 id<MTLBuffer> stagingBuffers[QMTL_FRAMES_IN_FLIGHT];
+                id<MTLTexture> views[QRhi::MAX_LEVELS];
             } texture;
             struct {
                 id<MTLSamplerState> samplerState;
@@ -213,6 +213,7 @@ Q_DECLARE_TYPEINFO(QRhiMetalData::ActiveReadback, Q_MOVABLE_TYPE);
 struct QMetalBufferData
 {
     bool managed;
+    bool slotted;
     id<MTLBuffer> buf[QMTL_FRAMES_IN_FLIGHT];
     QVector<QRhiResourceUpdateBatchPrivate::DynamicBufferUpdate> pendingUpdates[QMTL_FRAMES_IN_FLIGHT];
 };
@@ -225,10 +226,16 @@ struct QMetalRenderBufferData
 
 struct QMetalTextureData
 {
+    QMetalTextureData(QMetalTexture *t) : q(t) { }
+
+    QMetalTexture *q;
     MTLPixelFormat format;
     id<MTLTexture> tex = nil;
     id<MTLBuffer> stagingBuf[QMTL_FRAMES_IN_FLIGHT];
     bool owns = true;
+    id<MTLTexture> perLevelViews[QRhi::MAX_LEVELS];
+
+    id<MTLTexture> viewForLevel(int level);
 };
 
 struct QMetalSamplerData
@@ -239,7 +246,8 @@ struct QMetalSamplerData
 struct QMetalCommandBufferData
 {
     id<MTLCommandBuffer> cb;
-    id<MTLRenderCommandEncoder> currentPassEncoder;
+    id<MTLRenderCommandEncoder> currentRenderPassEncoder;
+    id<MTLComputeCommandEncoder> currentComputePassEncoder;
     MTLRenderPassDescriptor *currentPassRpDesc;
     int currentFirstVertexBinding;
     QRhiBatchedBindings<id<MTLBuffer> > currentVertexInputsBuffers;
@@ -286,6 +294,14 @@ struct QMetalGraphicsPipelineData
     id<MTLFunction> fsFunc = nil;
 };
 
+struct QMetalComputePipelineData
+{
+    id<MTLComputePipelineState> ps = nil;
+    id<MTLLibrary> csLib = nil;
+    id<MTLFunction> csFunc = nil;
+    MTLSize localSize;
+};
+
 struct QMetalSwapChainData
 {
     CAMetalLayer *layer = nullptr;
@@ -505,6 +521,8 @@ bool QRhiMetal::isFeatureSupported(QRhi::Feature feature) const
         return true;
     case QRhi::ElementIndexUint:
         return true;
+    case QRhi::Compute:
+        return true;
     default:
         Q_UNREACHABLE();
         return false;
@@ -573,6 +591,11 @@ QRhiGraphicsPipeline *QRhiMetal::createGraphicsPipeline()
     return new QMetalGraphicsPipeline(this);
 }
 
+QRhiComputePipeline *QRhiMetal::createComputePipeline()
+{
+    return new QMetalComputePipeline(this);
+}
+
 QRhiShaderResourceBindings *QRhiMetal::createShaderResourceBindings()
 {
     return new QMetalShaderResourceBindings(this);
@@ -583,7 +606,7 @@ void QRhiMetal::enqueueShaderResourceBindings(QMetalShaderResourceBindings *srbD
                                               const QRhiCommandBuffer::DynamicOffset *dynamicOffsets,
                                               bool offsetOnlyChange)
 {
-    static const int KNOWN_STAGES = 2;
+    static const int KNOWN_STAGES = 3;
     struct {
         QRhiBatchedBindings<id<MTLBuffer> > buffers;
         QRhiBatchedBindings<NSUInteger> bufferOffsets;
@@ -597,7 +620,7 @@ void QRhiMetal::enqueueShaderResourceBindings(QMetalShaderResourceBindings *srbD
         case QRhiShaderResourceBinding::UniformBuffer:
         {
             QMetalBuffer *bufD = QRHI_RES(QMetalBuffer, b->u.ubuf.buf);
-            id<MTLBuffer> mtlbuf = bufD->d->buf[bufD->m_type == QRhiBuffer::Immutable ? 0 : currentFrameSlot];
+            id<MTLBuffer> mtlbuf = bufD->d->buf[bufD->d->slotted ? currentFrameSlot : 0];
             uint offset = b->u.ubuf.offset;
             for (int i = 0; i < dynamicOffsetCount; ++i) {
                 const QRhiCommandBuffer::DynamicOffset &dynOfs(dynamicOffsets[i]);
@@ -614,6 +637,10 @@ void QRhiMetal::enqueueShaderResourceBindings(QMetalShaderResourceBindings *srbD
                 res[1].buffers.feed(b->binding, mtlbuf);
                 res[1].bufferOffsets.feed(b->binding, offset);
             }
+            if (b->stage.testFlag(QRhiShaderResourceBinding::ComputeStage)) {
+                res[2].buffers.feed(b->binding, mtlbuf);
+                res[2].bufferOffsets.feed(b->binding, offset);
+            }
         }
             break;
         case QRhiShaderResourceBinding::SampledTexture:
@@ -628,6 +655,49 @@ void QRhiMetal::enqueueShaderResourceBindings(QMetalShaderResourceBindings *srbD
                 res[1].textures.feed(b->binding, texD->d->tex);
                 res[1].samplers.feed(b->binding, samplerD->d->samplerState);
             }
+            if (b->stage.testFlag(QRhiShaderResourceBinding::ComputeStage)) {
+                res[2].textures.feed(b->binding, texD->d->tex);
+                res[2].samplers.feed(b->binding, samplerD->d->samplerState);
+            }
+        }
+            break;
+        case QRhiShaderResourceBinding::ImageLoad:
+            Q_FALLTHROUGH();
+        case QRhiShaderResourceBinding::ImageStore:
+            Q_FALLTHROUGH();
+        case QRhiShaderResourceBinding::ImageLoadStore:
+        {
+            QMetalTexture *texD = QRHI_RES(QMetalTexture, b->u.simage.tex);
+            id<MTLTexture> t = texD->d->viewForLevel(b->u.simage.level);
+            if (b->stage.testFlag(QRhiShaderResourceBinding::VertexStage))
+                res[0].textures.feed(b->binding, t);
+            if (b->stage.testFlag(QRhiShaderResourceBinding::FragmentStage))
+                res[1].textures.feed(b->binding, t);
+            if (b->stage.testFlag(QRhiShaderResourceBinding::ComputeStage))
+                res[2].textures.feed(b->binding, t);
+        }
+            break;
+        case QRhiShaderResourceBinding::BufferLoad:
+            Q_FALLTHROUGH();
+        case QRhiShaderResourceBinding::BufferStore:
+            Q_FALLTHROUGH();
+        case QRhiShaderResourceBinding::BufferLoadStore:
+        {
+            QMetalBuffer *bufD = QRHI_RES(QMetalBuffer, b->u.sbuf.buf);
+            id<MTLBuffer> mtlbuf = bufD->d->buf[0];
+            uint offset = b->u.sbuf.offset;
+            if (b->stage.testFlag(QRhiShaderResourceBinding::VertexStage)) {
+                res[0].buffers.feed(b->binding, mtlbuf);
+                res[0].bufferOffsets.feed(b->binding, offset);
+            }
+            if (b->stage.testFlag(QRhiShaderResourceBinding::FragmentStage)) {
+                res[1].buffers.feed(b->binding, mtlbuf);
+                res[1].bufferOffsets.feed(b->binding, offset);
+            }
+            if (b->stage.testFlag(QRhiShaderResourceBinding::ComputeStage)) {
+                res[2].buffers.feed(b->binding, mtlbuf);
+                res[2].bufferOffsets.feed(b->binding, offset);
+            }
         }
             break;
         default:
@@ -645,12 +715,17 @@ void QRhiMetal::enqueueShaderResourceBindings(QMetalShaderResourceBindings *srbD
             const auto &offsetBatch(res[idx].bufferOffsets.batches[i]);
             switch (idx) {
             case 0:
-                [cbD->d->currentPassEncoder setVertexBuffers: bufferBatch.resources.constData()
+                [cbD->d->currentRenderPassEncoder setVertexBuffers: bufferBatch.resources.constData()
                   offsets: offsetBatch.resources.constData()
                   withRange: NSMakeRange(bufferBatch.startBinding, bufferBatch.resources.count())];
                 break;
             case 1:
-                [cbD->d->currentPassEncoder setFragmentBuffers: bufferBatch.resources.constData()
+                [cbD->d->currentRenderPassEncoder setFragmentBuffers: bufferBatch.resources.constData()
+                  offsets: offsetBatch.resources.constData()
+                  withRange: NSMakeRange(bufferBatch.startBinding, bufferBatch.resources.count())];
+                break;
+            case 2:
+                [cbD->d->currentComputePassEncoder setBuffers: bufferBatch.resources.constData()
                   offsets: offsetBatch.resources.constData()
                   withRange: NSMakeRange(bufferBatch.startBinding, bufferBatch.resources.count())];
                 break;
@@ -670,11 +745,15 @@ void QRhiMetal::enqueueShaderResourceBindings(QMetalShaderResourceBindings *srbD
             const auto &batch(res[idx].textures.batches[i]);
             switch (idx) {
             case 0:
-                [cbD->d->currentPassEncoder setVertexTextures: batch.resources.constData()
+                [cbD->d->currentRenderPassEncoder setVertexTextures: batch.resources.constData()
                   withRange: NSMakeRange(batch.startBinding, batch.resources.count())];
                 break;
             case 1:
-                [cbD->d->currentPassEncoder setFragmentTextures: batch.resources.constData()
+                [cbD->d->currentRenderPassEncoder setFragmentTextures: batch.resources.constData()
+                  withRange: NSMakeRange(batch.startBinding, batch.resources.count())];
+                break;
+            case 2:
+                [cbD->d->currentComputePassEncoder setTextures: batch.resources.constData()
                   withRange: NSMakeRange(batch.startBinding, batch.resources.count())];
                 break;
             default:
@@ -686,11 +765,15 @@ void QRhiMetal::enqueueShaderResourceBindings(QMetalShaderResourceBindings *srbD
             const auto &batch(res[idx].samplers.batches[i]);
             switch (idx) {
             case 0:
-                [cbD->d->currentPassEncoder setVertexSamplerStates: batch.resources.constData()
+                [cbD->d->currentRenderPassEncoder setVertexSamplerStates: batch.resources.constData()
                   withRange: NSMakeRange(batch.startBinding, batch.resources.count())];
                 break;
             case 1:
-                [cbD->d->currentPassEncoder setFragmentSamplerStates: batch.resources.constData()
+                [cbD->d->currentRenderPassEncoder setFragmentSamplerStates: batch.resources.constData()
+                  withRange: NSMakeRange(batch.startBinding, batch.resources.count())];
+                break;
+            case 2:
+                [cbD->d->currentComputePassEncoder setSamplerStates: batch.resources.constData()
                   withRange: NSMakeRange(batch.startBinding, batch.resources.count())];
                 break;
             default:
@@ -703,19 +786,19 @@ void QRhiMetal::enqueueShaderResourceBindings(QMetalShaderResourceBindings *srbD
 
 void QRhiMetal::setGraphicsPipeline(QRhiCommandBuffer *cb, QRhiGraphicsPipeline *ps)
 {
-    Q_ASSERT(inPass);
-
     QMetalCommandBuffer *cbD = QRHI_RES(QMetalCommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QMetalCommandBuffer::RenderPass);
     QMetalGraphicsPipeline *psD = QRHI_RES(QMetalGraphicsPipeline, ps);
 
-    if (cbD->currentPipeline != ps || cbD->currentPipelineGeneration != psD->generation) {
+    if (cbD->currentGraphicsPipeline != ps || cbD->currentPipelineGeneration != psD->generation) {
-        cbD->currentPipeline = ps;
+        cbD->currentGraphicsPipeline = ps;
+        cbD->currentComputePipeline = nullptr;
         cbD->currentPipelineGeneration = psD->generation;
 
-        [cbD->d->currentPassEncoder setRenderPipelineState: psD->d->ps];
+        [cbD->d->currentRenderPassEncoder setRenderPipelineState: psD->d->ps];
-        [cbD->d->currentPassEncoder setDepthStencilState: psD->d->ds];
+        [cbD->d->currentRenderPassEncoder setDepthStencilState: psD->d->ds];
-        [cbD->d->currentPassEncoder setCullMode: psD->d->cullMode];
+        [cbD->d->currentRenderPassEncoder setCullMode: psD->d->cullMode];
-        [cbD->d->currentPassEncoder setFrontFacingWinding: psD->d->winding];
+        [cbD->d->currentRenderPassEncoder setFrontFacingWinding: psD->d->winding];
     }
 
     psD->lastActiveFrameSlot = currentFrameSlot;
@@ -725,12 +808,17 @@ void QRhiMetal::setShaderResources(QRhiCommandBuffer *cb, QRhiShaderResourceBind
                                    int dynamicOffsetCount,
                                    const QRhiCommandBuffer::DynamicOffset *dynamicOffsets)
 {
-    Q_ASSERT(inPass);
-
     QMetalCommandBuffer *cbD = QRHI_RES(QMetalCommandBuffer, cb);
-    Q_ASSERT(cbD->currentPipeline);
+    Q_ASSERT(cbD->recordingPass != QMetalCommandBuffer::NoPass);
-    if (!srb)
+    QMetalGraphicsPipeline *gfxPsD = QRHI_RES(QMetalGraphicsPipeline, cbD->currentGraphicsPipeline);
-        srb = QRHI_RES(QMetalGraphicsPipeline, cbD->currentPipeline)->m_shaderResourceBindings;
+    QMetalComputePipeline *compPsD = QRHI_RES(QMetalComputePipeline, cbD->currentComputePipeline);
+
+    if (!srb) {
+        if (gfxPsD)
+            srb = gfxPsD->m_shaderResourceBindings;
+        else
+            srb = compPsD->m_shaderResourceBindings;
+    }
 
     QMetalShaderResourceBindings *srbD = QRHI_RES(QMetalShaderResourceBindings, srb);
     bool hasSlottedResourceInSrb = false;
@@ -747,7 +835,7 @@ void QRhiMetal::setShaderResources(QRhiCommandBuffer *cb, QRhiShaderResourceBind
             QMetalBuffer *bufD = QRHI_RES(QMetalBuffer, b->u.ubuf.buf);
             Q_ASSERT(bufD->m_usage.testFlag(QRhiBuffer::UniformBuffer));
             executeBufferHostWritesForCurrentFrame(bufD);
-            if (bufD->m_type != QRhiBuffer::Immutable)
+            if (bufD->d->slotted)
                 hasSlottedResourceInSrb = true;
             if (b->u.ubuf.hasDynamicOffset)
                 hasDynamicOffsetInSrb = true;
@@ -778,6 +866,38 @@ void QRhiMetal::setShaderResources(QRhiCommandBuffer *cb, QRhiShaderResourceBind
             samplerD->lastActiveFrameSlot = currentFrameSlot;
         }
             break;
+        case QRhiShaderResourceBinding::ImageLoad:
+            Q_FALLTHROUGH();
+        case QRhiShaderResourceBinding::ImageStore:
+            Q_FALLTHROUGH();
+        case QRhiShaderResourceBinding::ImageLoadStore:
+        {
+            QMetalTexture *texD = QRHI_RES(QMetalTexture, b->u.simage.tex);
+            if (texD->generation != bd.simage.generation || texD->m_id != bd.simage.id) {
+                resNeedsRebind = true;
+                bd.simage.id = texD->m_id;
+                bd.simage.generation = texD->generation;
+            }
+            texD->lastActiveFrameSlot = currentFrameSlot;
+        }
+            break;
+        case QRhiShaderResourceBinding::BufferLoad:
+            Q_FALLTHROUGH();
+        case QRhiShaderResourceBinding::BufferStore:
+            Q_FALLTHROUGH();
+        case QRhiShaderResourceBinding::BufferLoadStore:
+        {
+            QMetalBuffer *bufD = QRHI_RES(QMetalBuffer, b->u.sbuf.buf);
+            Q_ASSERT(bufD->m_usage.testFlag(QRhiBuffer::StorageBuffer));
+            executeBufferHostWritesForCurrentFrame(bufD);
+            if (bufD->generation != bd.sbuf.generation || bufD->m_id != bd.sbuf.id) {
+                resNeedsRebind = true;
+                bd.sbuf.id = bufD->m_id;
+                bd.sbuf.generation = bufD->generation;
+            }
+            bufD->lastActiveFrameSlot = currentFrameSlot;
+        }
+            break;
         default:
             Q_UNREACHABLE();
             break;
@@ -789,15 +909,22 @@ void QRhiMetal::setShaderResources(QRhiCommandBuffer *cb, QRhiShaderResourceBind
     if (hasSlottedResourceInSrb && cbD->currentResSlot != resSlot)
         resNeedsRebind = true;
 
-    const bool srbChange = cbD->currentSrb != srb || cbD->currentSrbGeneration != srbD->generation;
+    const bool srbChanged = gfxPsD ? (cbD->currentGraphicsSrb != srb) : (cbD->currentComputeSrb != srb);
+    const bool srbRebuilt = cbD->currentSrbGeneration != srbD->generation;
 
     // dynamic uniform buffer offsets always trigger a rebind
-    if (hasDynamicOffsetInSrb || resNeedsRebind || srbChange) {
+    if (hasDynamicOffsetInSrb || resNeedsRebind || srbChanged || srbRebuilt) {
-        cbD->currentSrb = srb;
+        if (gfxPsD) {
+            cbD->currentGraphicsSrb = srb;
+            cbD->currentComputeSrb = nullptr;
+        } else {
+            cbD->currentGraphicsSrb = nullptr;
+            cbD->currentComputeSrb = srb;
+        }
         cbD->currentSrbGeneration = srbD->generation;
         cbD->currentResSlot = resSlot;
 
-        const bool offsetOnlyChange = hasDynamicOffsetInSrb && !resNeedsRebind && !srbChange;
+        const bool offsetOnlyChange = hasDynamicOffsetInSrb && !resNeedsRebind && !srbChanged && !srbRebuilt;
         enqueueShaderResourceBindings(srbD, cbD, dynamicOffsetCount, dynamicOffsets, offsetOnlyChange);
     }
 }
@@ -806,9 +933,8 @@ void QRhiMetal::setVertexInput(QRhiCommandBuffer *cb,
                                int startBinding, int bindingCount, const QRhiCommandBuffer::VertexInput *bindings,
                                QRhiBuffer *indexBuf, quint32 indexOffset, QRhiCommandBuffer::IndexFormat indexFormat)
 {
-    Q_ASSERT(inPass);
     QMetalCommandBuffer *cbD = QRHI_RES(QMetalCommandBuffer, cb);
-    Q_ASSERT(cbD->currentPipeline);
+    Q_ASSERT(cbD->recordingPass == QMetalCommandBuffer::RenderPass);
 
     QRhiBatchedBindings<id<MTLBuffer> > buffers;
     QRhiBatchedBindings<NSUInteger> offsets;
@@ -816,7 +942,7 @@ void QRhiMetal::setVertexInput(QRhiCommandBuffer *cb,
         QMetalBuffer *bufD = QRHI_RES(QMetalBuffer, bindings[i].first);
         executeBufferHostWritesForCurrentFrame(bufD);
         bufD->lastActiveFrameSlot = currentFrameSlot;
-        id<MTLBuffer> mtlbuf = bufD->d->buf[bufD->m_type == QRhiBuffer::Immutable ? 0 : currentFrameSlot];
+        id<MTLBuffer> mtlbuf = bufD->d->buf[bufD->d->slotted ? currentFrameSlot : 0];
         buffers.feed(startBinding + i, mtlbuf);
         offsets.feed(startBinding + i, bindings[i].second);
     }
@@ -824,12 +950,12 @@ void QRhiMetal::setVertexInput(QRhiCommandBuffer *cb,
     offsets.finish();
 
     // same binding space for vertex and constant buffers - work it around
-    QRhiShaderResourceBindings *srb = cbD->currentSrb;
+    QRhiShaderResourceBindings *srb = cbD->currentGraphicsSrb;
     // There's nothing guaranteeing setShaderResources() was called before
     // setVertexInput()... but whatever srb will get bound will have to be
     // layout-compatible anyways so maxBinding is the same.
     if (!srb)
-        srb = cbD->currentPipeline->shaderResourceBindings();
+        srb = cbD->currentGraphicsPipeline->shaderResourceBindings();
     const int firstVertexBinding = QRHI_RES(QMetalShaderResourceBindings, srb)->maxBinding + 1;
 
     if (firstVertexBinding != cbD->d->currentFirstVertexBinding
@@ -843,7 +969,7 @@ void QRhiMetal::setVertexInput(QRhiCommandBuffer *cb,
         for (int i = 0, ie = buffers.batches.count(); i != ie; ++i) {
             const auto &bufferBatch(buffers.batches[i]);
             const auto &offsetBatch(offsets.batches[i]);
-            [cbD->d->currentPassEncoder setVertexBuffers:
+            [cbD->d->currentRenderPassEncoder setVertexBuffers:
                 bufferBatch.resources.constData()
               offsets: offsetBatch.resources.constData()
               withRange: NSMakeRange(firstVertexBinding + bufferBatch.startBinding, bufferBatch.resources.count())];
@@ -864,9 +990,8 @@ void QRhiMetal::setVertexInput(QRhiCommandBuffer *cb,
 
 void QRhiMetal::setViewport(QRhiCommandBuffer *cb, const QRhiViewport &viewport)
 {
-    Q_ASSERT(inPass);
     QMetalCommandBuffer *cbD = QRHI_RES(QMetalCommandBuffer, cb);
-    Q_ASSERT(cbD->currentPipeline && cbD->currentTarget);
+    Q_ASSERT(cbD->recordingPass == QMetalCommandBuffer::RenderPass);
     const QSize outputSize = cbD->currentTarget->pixelSize();
 
     // x,y is top-left in MTLViewportRect but bottom-left in QRhiViewport
@@ -882,24 +1007,23 @@ void QRhiMetal::setViewport(QRhiCommandBuffer *cb, const QRhiViewport &viewport)
     vp.znear = viewport.minDepth();
     vp.zfar = viewport.maxDepth();
 
-    [cbD->d->currentPassEncoder setViewport: vp];
+    [cbD->d->currentRenderPassEncoder setViewport: vp];
 
-    if (!QRHI_RES(QMetalGraphicsPipeline, cbD->currentPipeline)->m_flags.testFlag(QRhiGraphicsPipeline::UsesScissor)) {
+    if (!QRHI_RES(QMetalGraphicsPipeline, cbD->currentGraphicsPipeline)->m_flags.testFlag(QRhiGraphicsPipeline::UsesScissor)) {
         MTLScissorRect s;
         s.x = x;
         s.y = y;
         s.width = w;
         s.height = h;
-        [cbD->d->currentPassEncoder setScissorRect: s];
+        [cbD->d->currentRenderPassEncoder setScissorRect: s];
     }
 }
 
 void QRhiMetal::setScissor(QRhiCommandBuffer *cb, const QRhiScissor &scissor)
 {
-    Q_ASSERT(inPass);
     QMetalCommandBuffer *cbD = QRHI_RES(QMetalCommandBuffer, cb);
-    Q_ASSERT(cbD->currentPipeline && cbD->currentTarget);
+    Q_ASSERT(cbD->recordingPass == QMetalCommandBuffer::RenderPass);
-    Q_ASSERT(QRHI_RES(QMetalGraphicsPipeline, cbD->currentPipeline)->m_flags.testFlag(QRhiGraphicsPipeline::UsesScissor));
+    Q_ASSERT(QRHI_RES(QMetalGraphicsPipeline, cbD->currentGraphicsPipeline)->m_flags.testFlag(QRhiGraphicsPipeline::UsesScissor));
     const QSize outputSize = cbD->currentTarget->pixelSize();
 
     // x,y is top-left in MTLScissorRect but bottom-left in QRhiScissor
@@ -913,38 +1037,42 @@ void QRhiMetal::setScissor(QRhiCommandBuffer *cb, const QRhiScissor &scissor)
     s.width = w;
     s.height = h;
 
-    [cbD->d->currentPassEncoder setScissorRect: s];
+    [cbD->d->currentRenderPassEncoder setScissorRect: s];
 }
 
 void QRhiMetal::setBlendConstants(QRhiCommandBuffer *cb, const QColor &c)
 {
-    Q_ASSERT(inPass);
     QMetalCommandBuffer *cbD = QRHI_RES(QMetalCommandBuffer, cb);
-    [cbD->d->currentPassEncoder setBlendColorRed: c.redF() green: c.greenF() blue: c.blueF() alpha: c.alphaF()];
+    Q_ASSERT(cbD->recordingPass == QMetalCommandBuffer::RenderPass);
+
+    [cbD->d->currentRenderPassEncoder setBlendColorRed: c.redF() green: c.greenF() blue: c.blueF() alpha: c.alphaF()];
 }
 
 void QRhiMetal::setStencilRef(QRhiCommandBuffer *cb, quint32 refValue)
 {
-    Q_ASSERT(inPass);
     QMetalCommandBuffer *cbD = QRHI_RES(QMetalCommandBuffer, cb);
-    [cbD->d->currentPassEncoder setStencilReferenceValue: refValue];
+    Q_ASSERT(cbD->recordingPass == QMetalCommandBuffer::RenderPass);
+
+    [cbD->d->currentRenderPassEncoder setStencilReferenceValue: refValue];
 }
 
 void QRhiMetal::draw(QRhiCommandBuffer *cb, quint32 vertexCount,
                      quint32 instanceCount, quint32 firstVertex, quint32 firstInstance)
 {
-    Q_ASSERT(inPass);
     QMetalCommandBuffer *cbD = QRHI_RES(QMetalCommandBuffer, cb);
-    [cbD->d->currentPassEncoder drawPrimitives:
+    Q_ASSERT(cbD->recordingPass == QMetalCommandBuffer::RenderPass);
-        QRHI_RES(QMetalGraphicsPipeline, cbD->currentPipeline)->d->primitiveType
+
+    [cbD->d->currentRenderPassEncoder drawPrimitives:
+        QRHI_RES(QMetalGraphicsPipeline, cbD->currentGraphicsPipeline)->d->primitiveType
       vertexStart: firstVertex vertexCount: vertexCount instanceCount: instanceCount baseInstance: firstInstance];
 }
 
 void QRhiMetal::drawIndexed(QRhiCommandBuffer *cb, quint32 indexCount,
                             quint32 instanceCount, quint32 firstIndex, qint32 vertexOffset, quint32 firstInstance)
 {
-    Q_ASSERT(inPass);
     QMetalCommandBuffer *cbD = QRHI_RES(QMetalCommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QMetalCommandBuffer::RenderPass);
+
     if (!cbD->currentIndexBuffer)
         return;
 
@@ -952,9 +1080,9 @@ void QRhiMetal::drawIndexed(QRhiCommandBuffer *cb, quint32 indexCount,
     Q_ASSERT(indexOffset == aligned(indexOffset, 4));
 
     QMetalBuffer *ibufD = QRHI_RES(QMetalBuffer, cbD->currentIndexBuffer);
-    id<MTLBuffer> mtlbuf = ibufD->d->buf[ibufD->m_type == QRhiBuffer::Immutable ? 0 : currentFrameSlot];
+    id<MTLBuffer> mtlbuf = ibufD->d->buf[ibufD->d->slotted ? currentFrameSlot : 0];
 
-    [cbD->d->currentPassEncoder drawIndexedPrimitives: QRHI_RES(QMetalGraphicsPipeline, cbD->currentPipeline)->d->primitiveType
+    [cbD->d->currentRenderPassEncoder drawIndexedPrimitives: QRHI_RES(QMetalGraphicsPipeline, cbD->currentGraphicsPipeline)->d->primitiveType
       indexCount: indexCount
       indexType: cbD->currentIndexFormat == QRhiCommandBuffer::IndexUInt16 ? MTLIndexTypeUInt16 : MTLIndexTypeUInt32
       indexBuffer: mtlbuf
@@ -971,8 +1099,8 @@ void QRhiMetal::debugMarkBegin(QRhiCommandBuffer *cb, const QByteArray &name)
 
     NSString *str = [NSString stringWithUTF8String: name.constData()];
     QMetalCommandBuffer *cbD = QRHI_RES(QMetalCommandBuffer, cb);
-    if (inPass) {
+    if (cbD->recordingPass != QMetalCommandBuffer::NoPass) {
-        [cbD->d->currentPassEncoder pushDebugGroup: str];
+        [cbD->d->currentRenderPassEncoder pushDebugGroup: str];
     } else {
         if (@available(macOS 10.13, iOS 11.0, *))
             [cbD->d->cb pushDebugGroup: str];
@@ -985,8 +1113,8 @@ void QRhiMetal::debugMarkEnd(QRhiCommandBuffer *cb)
         return;
 
     QMetalCommandBuffer *cbD = QRHI_RES(QMetalCommandBuffer, cb);
-    if (inPass) {
+    if (cbD->recordingPass != QMetalCommandBuffer::NoPass) {
-        [cbD->d->currentPassEncoder popDebugGroup];
+        [cbD->d->currentRenderPassEncoder popDebugGroup];
     } else {
         if (@available(macOS 10.13, iOS 11.0, *))
             [cbD->d->cb popDebugGroup];
@@ -998,10 +1126,9 @@ void QRhiMetal::debugMarkMsg(QRhiCommandBuffer *cb, const QByteArray &msg)
     if (!debugMarkers)
         return;
 
-    if (inPass) {
+    QMetalCommandBuffer *cbD = QRHI_RES(QMetalCommandBuffer, cb);
-        QMetalCommandBuffer *cbD = QRHI_RES(QMetalCommandBuffer, cb);
+    if (cbD->recordingPass != QMetalCommandBuffer::NoPass)
-        [cbD->d->currentPassEncoder insertDebugSignpost: [NSString stringWithUTF8String: msg.constData()]];
+        [cbD->d->currentRenderPassEncoder insertDebugSignpost: [NSString stringWithUTF8String: msg.constData()]];
-    }
 }
 
 const QRhiNativeHandles *QRhiMetal::nativeHandles(QRhiCommandBuffer *cb)
@@ -1023,8 +1150,6 @@ void QRhiMetal::endExternal(QRhiCommandBuffer *cb)
 QRhi::FrameOpResult QRhiMetal::beginFrame(QRhiSwapChain *swapChain, QRhi::BeginFrameFlags flags)
 {
     Q_UNUSED(flags);
-    Q_ASSERT(!inFrame);
-    inFrame = true;
 
     QMetalSwapChain *swapChainD = QRHI_RES(QMetalSwapChain, swapChain);
 
@@ -1077,9 +1202,6 @@ QRhi::FrameOpResult QRhiMetal::beginFrame(QRhiSwapChain *swapChain, QRhi::BeginF
 
 QRhi::FrameOpResult QRhiMetal::endFrame(QRhiSwapChain *swapChain, QRhi::EndFrameFlags flags)
 {
-    Q_ASSERT(inFrame);
-    inFrame = false;
-
     QMetalSwapChain *swapChainD = QRHI_RES(QMetalSwapChain, swapChain);
     Q_ASSERT(currentSwapChain == swapChainD);
 
@@ -1110,9 +1232,6 @@ QRhi::FrameOpResult QRhiMetal::endFrame(QRhiSwapChain *swapChain, QRhi::EndFrame
 
 QRhi::FrameOpResult QRhiMetal::beginOffscreenFrame(QRhiCommandBuffer **cb)
 {
-    Q_ASSERT(!inFrame);
-    inFrame = true;
-
     currentFrameSlot = (currentFrameSlot + 1) % QMTL_FRAMES_IN_FLIGHT;
     if (swapchains.count() > 1) {
         for (QMetalSwapChain *sc : qAsConst(swapchains)) {
@@ -1140,8 +1259,6 @@ QRhi::FrameOpResult QRhiMetal::endOffscreenFrame()
 {
     Q_ASSERT(d->ofr.active);
     d->ofr.active = false;
-    Q_ASSERT(inFrame);
-    inFrame = false;
 
     [d->ofr.cbWrapper.d->cb commit];
 
@@ -1155,17 +1272,17 @@ QRhi::FrameOpResult QRhiMetal::endOffscreenFrame()
 
 QRhi::FrameOpResult QRhiMetal::finish()
 {
-    Q_ASSERT(!inPass);
-
     id<MTLCommandBuffer> cb = nil;
     QMetalSwapChain *swapChainD = nullptr;
     if (inFrame) {
         if (d->ofr.active) {
             Q_ASSERT(!currentSwapChain);
+            Q_ASSERT(d->ofr.cbWrapper.recordingPass == QMetalCommandBuffer::NoPass);
             cb = d->ofr.cbWrapper.d->cb;
         } else {
             Q_ASSERT(currentSwapChain);
             swapChainD = currentSwapChain;
+            Q_ASSERT(swapChainD->cbWrapper.recordingPass == QMetalCommandBuffer::NoPass);
             cb = swapChainD->cbWrapper.d->cb;
         }
     }
@@ -1373,11 +1490,13 @@ void QRhiMetal::enqueueResourceUpdates(QRhiCommandBuffer *cb, QRhiResourceUpdate
             bufD->d->pendingUpdates[i].append(u);
     }
 
+    // Due to the Metal API the handling of static and dynamic buffers is
+    // basically the same. So go through the same pendingUpdates machinery.
     for (const QRhiResourceUpdateBatchPrivate::StaticBufferUpload &u : ud->staticBufferUploads) {
         QMetalBuffer *bufD = QRHI_RES(QMetalBuffer, u.buf);
         Q_ASSERT(bufD->m_type != QRhiBuffer::Dynamic);
         Q_ASSERT(u.offset + u.data.size() <= bufD->m_size);
-        for (int i = 0, ie = bufD->m_type == QRhiBuffer::Immutable ? 1 : QMTL_FRAMES_IN_FLIGHT; i != ie; ++i)
+        for (int i = 0, ie = bufD->d->slotted ? QMTL_FRAMES_IN_FLIGHT : 1; i != ie; ++i)
             bufD->d->pendingUpdates[i].append({ u.buf, u.offset, u.data.size(), u.data.constData() });
     }
 
@@ -1516,9 +1635,10 @@ void QRhiMetal::enqueueResourceUpdates(QRhiCommandBuffer *cb, QRhiResourceUpdate
     ud->free();
 }
 
+// this handles all types of buffers, not just Dynamic
 void QRhiMetal::executeBufferHostWritesForCurrentFrame(QMetalBuffer *bufD)
 {
-    const int idx = bufD->m_type == QRhiBuffer::Immutable ? 0 : currentFrameSlot;
+    const int idx = bufD->d->slotted ? currentFrameSlot : 0;
     QVector<QRhiResourceUpdateBatchPrivate::DynamicBufferUpdate> &updates(bufD->d->pendingUpdates[idx]);
     if (updates.isEmpty())
         return;
@@ -1542,7 +1662,7 @@ void QRhiMetal::executeBufferHostWritesForCurrentFrame(QMetalBuffer *bufD)
 
 void QRhiMetal::resourceUpdate(QRhiCommandBuffer *cb, QRhiResourceUpdateBatch *resourceUpdates)
 {
-    Q_ASSERT(inFrame && !inPass);
+    Q_ASSERT(QRHI_RES(QMetalCommandBuffer, cb)->recordingPass == QMetalCommandBuffer::NoPass);
 
     enqueueResourceUpdates(cb, resourceUpdates);
 }
@@ -1553,13 +1673,12 @@ void QRhiMetal::beginPass(QRhiCommandBuffer *cb,
                           const QRhiDepthStencilClearValue &depthStencilClearValue,
                           QRhiResourceUpdateBatch *resourceUpdates)
 {
-    Q_ASSERT(inFrame && !inPass);
+    QMetalCommandBuffer *cbD = QRHI_RES(QMetalCommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QMetalCommandBuffer::NoPass);
 
     if (resourceUpdates)
         enqueueResourceUpdates(cb, resourceUpdates);
 
-    QMetalCommandBuffer *cbD = QRHI_RES(QMetalCommandBuffer, cb);
-
     QMetalRenderTargetData *rtD = nullptr;
     switch (rt->resourceType()) {
     case QRhiResource::RenderTarget:
@@ -1639,28 +1758,80 @@ void QRhiMetal::beginPass(QRhiCommandBuffer *cb,
             cbD->d->currentPassRpDesc.depthAttachment.storeAction = MTLStoreActionStore;
     }
 
-    cbD->d->currentPassEncoder = [cbD->d->cb renderCommandEncoderWithDescriptor: cbD->d->currentPassRpDesc];
+    cbD->d->currentRenderPassEncoder = [cbD->d->cb renderCommandEncoderWithDescriptor: cbD->d->currentPassRpDesc];
 
     cbD->resetPerPassState();
 
+    cbD->recordingPass = QMetalCommandBuffer::RenderPass;
     cbD->currentTarget = rt;
-    inPass = true;
 }
 
 void QRhiMetal::endPass(QRhiCommandBuffer *cb, QRhiResourceUpdateBatch *resourceUpdates)
 {
-    Q_ASSERT(inPass);
-    inPass = false;
-
     QMetalCommandBuffer *cbD = QRHI_RES(QMetalCommandBuffer, cb);
-    [cbD->d->currentPassEncoder endEncoding];
+    Q_ASSERT(cbD->recordingPass == QMetalCommandBuffer::RenderPass);
 
+    [cbD->d->currentRenderPassEncoder endEncoding];
+
+    cbD->recordingPass = QMetalCommandBuffer::NoPass;
     cbD->currentTarget = nullptr;
 
     if (resourceUpdates)
         enqueueResourceUpdates(cb, resourceUpdates);
 }
 
+void QRhiMetal::beginComputePass(QRhiCommandBuffer *cb, QRhiResourceUpdateBatch *resourceUpdates)
+{
+    QMetalCommandBuffer *cbD = QRHI_RES(QMetalCommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QMetalCommandBuffer::NoPass);
+
+    if (resourceUpdates)
+        enqueueResourceUpdates(cb, resourceUpdates);
+
+    cbD->d->currentComputePassEncoder = [cbD->d->cb computeCommandEncoder];
+    cbD->resetPerPassState();
+    cbD->recordingPass = QMetalCommandBuffer::ComputePass;
+}
+
+void QRhiMetal::endComputePass(QRhiCommandBuffer *cb, QRhiResourceUpdateBatch *resourceUpdates)
+{
+    QMetalCommandBuffer *cbD = QRHI_RES(QMetalCommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QMetalCommandBuffer::ComputePass);
+
+    [cbD->d->currentComputePassEncoder endEncoding];
+    cbD->recordingPass = QMetalCommandBuffer::NoPass;
+
+    if (resourceUpdates)
+        enqueueResourceUpdates(cb, resourceUpdates);
+}
+
+void QRhiMetal::setComputePipeline(QRhiCommandBuffer *cb, QRhiComputePipeline *ps)
+{
+    QMetalCommandBuffer *cbD = QRHI_RES(QMetalCommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QMetalCommandBuffer::ComputePass);
+    QMetalComputePipeline *psD = QRHI_RES(QMetalComputePipeline, ps);
+
+    if (cbD->currentComputePipeline != ps || cbD->currentPipelineGeneration != psD->generation) {
+        cbD->currentGraphicsPipeline = nullptr;
+        cbD->currentComputePipeline = ps;
+        cbD->currentPipelineGeneration = psD->generation;
+
+        [cbD->d->currentComputePassEncoder setComputePipelineState: psD->d->ps];
+    }
+
+    psD->lastActiveFrameSlot = currentFrameSlot;
+}
+
+void QRhiMetal::dispatch(QRhiCommandBuffer *cb, int x, int y, int z)
+{
+    QMetalCommandBuffer *cbD = QRHI_RES(QMetalCommandBuffer, cb);
+    Q_ASSERT(cbD->recordingPass == QMetalCommandBuffer::ComputePass);
+    QMetalComputePipeline *psD = QRHI_RES(QMetalComputePipeline, cbD->currentComputePipeline);
+
+    [cbD->d->currentComputePassEncoder dispatchThreadgroups: MTLSizeMake(x, y, z)
+      threadsPerThreadgroup: psD->d->localSize];
+}
+
 static void qrhimtl_releaseBuffer(const QRhiMetalData::DeferredReleaseEntry &e)
 {
     for (int i = 0; i < QMTL_FRAMES_IN_FLIGHT; ++i)
@@ -1677,6 +1848,8 @@ static void qrhimtl_releaseTexture(const QRhiMetalData::DeferredReleaseEntry &e)
     [e.texture.texture release];
     for (int i = 0; i < QMTL_FRAMES_IN_FLIGHT; ++i)
         [e.texture.stagingBuffers[i] release];
+    for (int i = 0; i < QRhi::MAX_LEVELS; ++i)
+        [e.texture.views[i] release];
 }
 
 static void qrhimtl_releaseSampler(const QRhiMetalData::DeferredReleaseEntry &e)
@@ -1782,6 +1955,11 @@ bool QMetalBuffer::build()
     if (d->buf[0])
         release();
 
+    if (m_usage.testFlag(QRhiBuffer::StorageBuffer) && m_type == Dynamic) {
+        qWarning("StorageBuffer cannot be combined with Dynamic");
+        return false;
+    }
+
     const int nonZeroSize = m_size <= 0 ? 256 : m_size;
     const int roundedSize = m_usage.testFlag(QRhiBuffer::UniformBuffer) ? aligned(nonZeroSize, 256) : nonZeroSize;
 
@@ -1794,15 +1972,17 @@ bool QMetalBuffer::build()
     }
 #endif
 
+    // Immutable and Static only has buf[0] and pendingUpdates[0] in use.
+    // Dynamic uses all.
+    d->slotted = m_type == Dynamic;
+
     QRHI_RES_RHI(QRhiMetal);
     for (int i = 0; i < QMTL_FRAMES_IN_FLIGHT; ++i) {
-        // Immutable only has buf[0] and pendingUpdates[0] in use.
+        if (i == 0 || d->slotted) {
-        // Static and Dynamic use all.
-        if (i == 0 || m_type != Immutable) {
             d->buf[i] = [rhiD->d->dev newBufferWithLength: roundedSize options: opts];
             d->pendingUpdates[i].reserve(16);
             if (!m_objectName.isEmpty()) {
-                if (m_type == Immutable) {
+                if (!d->slotted) {
                     d->buf[i].label = [NSString stringWithUTF8String: m_objectName.constData()];
                 } else {
                     const QByteArray name = m_objectName + '/' + QByteArray::number(i);
@@ -1813,7 +1993,7 @@ bool QMetalBuffer::build()
     }
 
     QRHI_PROF;
-    QRHI_PROF_F(newBuffer(this, roundedSize, m_type == Immutable ? 1 : QMTL_FRAMES_IN_FLIGHT, 0));
+    QRHI_PROF_F(newBuffer(this, roundedSize, d->slotted ? QMTL_FRAMES_IN_FLIGHT : 1, 0));
 
     lastActiveFrameSlot = -1;
     generation += 1;
@@ -1919,10 +2099,13 @@ QRhiTexture::Format QMetalRenderBuffer::backingFormat() const
 QMetalTexture::QMetalTexture(QRhiImplementation *rhi, Format format, const QSize &pixelSize,
                              int sampleCount, Flags flags)
     : QRhiTexture(rhi, format, pixelSize, sampleCount, flags),
-      d(new QMetalTextureData)
+      d(new QMetalTextureData(this))
 {
     for (int i = 0; i < QMTL_FRAMES_IN_FLIGHT; ++i)
         d->stagingBuf[i] = nil;
+
+    for (int i = 0; i < QRhi::MAX_LEVELS; ++i)
+        d->perLevelViews[i] = nil;
 }
 
 QMetalTexture::~QMetalTexture()
@@ -1949,6 +2132,11 @@ void QMetalTexture::release()
         d->stagingBuf[i] = nil;
     }
 
+    for (int i = 0; i < QRhi::MAX_LEVELS; ++i) {
+        e.texture.views[i] = d->perLevelViews[i];
+        d->perLevelViews[i] = nil;
+    }
+
     QRHI_RES_RHI(QRhiMetal);
     rhiD->d->releaseQueue.append(e);
     QRHI_PROF;
@@ -2138,6 +2326,8 @@ bool QMetalTexture::build()
     desc.usage = MTLTextureUsageShaderRead;
     if (m_flags.testFlag(RenderTarget))
         desc.usage |= MTLTextureUsageRenderTarget;
+    if (m_flags.testFlag(UsedWithLoadStore))
+        desc.usage |= MTLTextureUsageShaderWrite;
 
     QRHI_RES_RHI(QRhiMetal);
     d->tex = [rhiD->d->dev newTextureWithDescriptor: desc];
@@ -2187,6 +2377,21 @@ const QRhiNativeHandles *QMetalTexture::nativeHandles()
     return &nativeHandlesStruct;
 }
 
+id<MTLTexture> QMetalTextureData::viewForLevel(int level)
+{
+    Q_ASSERT(level >= 0 && level < int(q->mipLevelCount));
+    if (perLevelViews[level])
+        return perLevelViews[level];
+
+    const MTLTextureType type = [tex textureType];
+    const bool isCube = q->m_flags.testFlag(QRhiTexture::CubeMap);
+    id<MTLTexture> view = [tex newTextureViewWithPixelFormat: format textureType: type
+            levels: NSMakeRange(level, 1) slices: NSMakeRange(0, isCube ? 6 : 1)];
+
+    perLevelViews[level] = view;
+    return view;
+}
+
 QMetalSampler::QMetalSampler(QRhiImplementation *rhi, Filter magFilter, Filter minFilter, Filter mipmapMode,
                              AddressMode u, AddressMode v)
     : QRhiSampler(rhi, magFilter, minFilter, mipmapMode, u, v),
@@ -2538,6 +2743,28 @@ bool QMetalShaderResourceBindings::build()
             bd.stex.samplerGeneration = samplerD->generation;
         }
             break;
+        case QRhiShaderResourceBinding::ImageLoad:
+            Q_FALLTHROUGH();
+        case QRhiShaderResourceBinding::ImageStore:
+            Q_FALLTHROUGH();
+        case QRhiShaderResourceBinding::ImageLoadStore:
+        {
+            QMetalTexture *texD = QRHI_RES(QMetalTexture, b->u.simage.tex);
+            bd.simage.id = texD->m_id;
+            bd.simage.generation = texD->generation;
+        }
+            break;
+        case QRhiShaderResourceBinding::BufferLoad:
+            Q_FALLTHROUGH();
+        case QRhiShaderResourceBinding::BufferStore:
+            Q_FALLTHROUGH();
+        case QRhiShaderResourceBinding::BufferLoadStore:
+        {
+            QMetalBuffer *bufD = QRHI_RES(QMetalBuffer, b->u.sbuf.buf);
+            bd.sbuf.id = bufD->m_id;
+            bd.sbuf.generation = bufD->generation;
+        }
+            break;
         default:
             Q_UNREACHABLE();
             break;
@@ -2874,21 +3101,12 @@ bool QMetalGraphicsPipeline::build()
 
     rpDesc.vertexDescriptor = inputLayout;
 
-    if (@available(macOS 10.13, iOS 11.0, *)) {
+    // mutability cannot be determined (slotted buffers could be set as
-        // Everything is immutable because we can guarantee that "neither the
+    // MTLMutabilityImmutable, but then we potentially need a different
-        // CPU nor the GPU will modify a buffer's contents between the time the
+    // descriptor for each buffer combination as this depends on the actual
-        // buffer is set in a function's argument table and the time its
+    // buffers not just the resource binding layout) so leave it at the default
-        // associated command buffer completes execution" (as that's the point
-        // of our Vulkan-style buffer juggling in the first place).
-        const int vertexBufferCount = firstVertexBinding + bindings.count(); // cbuf + vbuf
-        const int fragmentBufferCount = firstVertexBinding; // cbuf
-        for (int i = 0; i < vertexBufferCount; ++i)
-            rpDesc.vertexBuffers[i].mutability = MTLMutabilityImmutable;
-        for (int i = 0; i < fragmentBufferCount; ++i)
-            rpDesc.fragmentBuffers[i].mutability = MTLMutabilityImmutable;
-    }
 
-    for (const QRhiGraphicsShaderStage &shaderStage : qAsConst(m_shaderStages)) {
+    for (const QRhiShaderStage &shaderStage : qAsConst(m_shaderStages)) {
         QString error;
         QByteArray entryPoint;
         id<MTLLibrary> lib = rhiD->d->createMetalLib(shaderStage.shader(), shaderStage.shaderVariant(), &error, &entryPoint);
@@ -2903,12 +3121,12 @@ bool QMetalGraphicsPipeline::build()
             return false;
         }
         switch (shaderStage.type()) {
-        case QRhiGraphicsShaderStage::Vertex:
+        case QRhiShaderStage::Vertex:
             rpDesc.vertexFunction = func;
             d->vsLib = lib;
             d->vsFunc = func;
             break;
-        case QRhiGraphicsShaderStage::Fragment:
+        case QRhiShaderStage::Fragment:
             rpDesc.fragmentFunction = func;
             d->fsLib = lib;
             d->fsFunc = func;
@@ -3000,6 +3218,83 @@ bool QMetalGraphicsPipeline::build()
     return true;
 }
 
+QMetalComputePipeline::QMetalComputePipeline(QRhiImplementation *rhi)
+    : QRhiComputePipeline(rhi),
+      d(new QMetalComputePipelineData)
+{
+}
+
+QMetalComputePipeline::~QMetalComputePipeline()
+{
+    release();
+    delete d;
+}
+
+void QMetalComputePipeline::release()
+{
+    QRHI_RES_RHI(QRhiMetal);
+
+    if (d->csFunc) {
+        [d->csFunc release];
+        d->csFunc = nil;
+    }
+    if (d->csLib) {
+        [d->csLib release];
+        d->csLib = nil;
+    }
+
+    if (!d->ps)
+        return;
+
+    if (d->ps) {
+        [d->ps release];
+        d->ps = nil;
+    }
+
+    rhiD->unregisterResource(this);
+}
+
+bool QMetalComputePipeline::build()
+{
+    if (d->ps)
+        release();
+
+    QRHI_RES_RHI(QRhiMetal);
+
+    const QShader shader = m_shaderStage.shader();
+    QString error;
+    QByteArray entryPoint;
+    id<MTLLibrary> lib = rhiD->d->createMetalLib(shader, m_shaderStage.shaderVariant(),
+                                                 &error, &entryPoint);
+    if (!lib) {
+        qWarning("MSL shader compilation failed: %s", qPrintable(error));
+        return false;
+    }
+    id<MTLFunction> func = rhiD->d->createMSLShaderFunction(lib, entryPoint);
+    if (!func) {
+        qWarning("MSL function for entry point %s not found", entryPoint.constData());
+        [lib release];
+        return false;
+    }
+    d->csLib = lib;
+    d->csFunc = func;
+    std::array<uint, 3> localSize = shader.description().computeShaderLocalSize();
+    d->localSize = MTLSizeMake(localSize[0], localSize[1], localSize[2]);
+
+    NSError *err = nil;
+    d->ps = [rhiD->d->dev newComputePipelineStateWithFunction: d->csFunc error: &err];
+    if (!d->ps) {
+        const QString msg = QString::fromNSString(err.localizedDescription);
+        qWarning("Failed to create render pipeline state: %s", qPrintable(msg));
+        return false;
+    }
+
+    lastActiveFrameSlot = -1;
+    generation += 1;
+    rhiD->registerResource(this);
+    return true;
+}
+
 QMetalCommandBuffer::QMetalCommandBuffer(QRhiImplementation *rhi)
     : QRhiCommandBuffer(rhi),
       d(new QMetalCommandBufferData)
@@ -3021,28 +3316,32 @@ void QMetalCommandBuffer::release()
 const QRhiNativeHandles *QMetalCommandBuffer::nativeHandles()
 {
     nativeHandlesStruct.commandBuffer = d->cb;
-    nativeHandlesStruct.encoder = d->currentPassEncoder;
+    nativeHandlesStruct.encoder = d->currentRenderPassEncoder;
     return &nativeHandlesStruct;
 }
 
 void QMetalCommandBuffer::resetState()
 {
-    d->currentPassEncoder = nil;
+    d->currentRenderPassEncoder = nil;
+    d->currentComputePassEncoder = nil;
     d->currentPassRpDesc = nil;
     resetPerPassState();
 }
 
 void QMetalCommandBuffer::resetPerPassState()
 {
+    recordingPass = NoPass;
     currentTarget = nullptr;
     resetPerPassCachedState();
 }
 
 void QMetalCommandBuffer::resetPerPassCachedState()
 {
-    currentPipeline = nullptr;
+    currentGraphicsPipeline = nullptr;
+    currentComputePipeline = nullptr;
     currentPipelineGeneration = 0;
-    currentSrb = nullptr;
+    currentGraphicsSrb = nullptr;
+    currentComputeSrb = nullptr;
     currentSrbGeneration = 0;
     currentResSlot = -1;
     currentIndexBuffer = nullptr;

src/gui/rhi/qrhimetal_p_p.h

@@ -113,6 +113,7 @@ struct QMetalTexture : public QRhiTexture
     int lastActiveFrameSlot = -1;
     friend class QRhiMetal;
     friend struct QMetalShaderResourceBindings;
+    friend struct QMetalTextureData;
 };
 
 struct QMetalSamplerData;
@@ -200,10 +201,20 @@ struct QMetalShaderResourceBindings : public QRhiShaderResourceBindings
         quint64 samplerId;
         uint samplerGeneration;
     };
+    struct BoundStorageImageData {
+        quint64 id;
+        uint generation;
+    };
+    struct BoundStorageBufferData {
+        quint64 id;
+        uint generation;
+    };
     struct BoundResourceData {
         union {
             BoundUniformBufferData ubuf;
             BoundSampledTextureData stex;
+            BoundStorageImageData simage;
+            BoundStorageBufferData sbuf;
         };
     };
     QVector<BoundResourceData> boundResourceData;
@@ -227,6 +238,21 @@ struct QMetalGraphicsPipeline : public QRhiGraphicsPipeline
     friend class QRhiMetal;
 };
 
+struct QMetalComputePipelineData;
+
+struct QMetalComputePipeline : public QRhiComputePipeline
+{
+    QMetalComputePipeline(QRhiImplementation *rhi);
+    ~QMetalComputePipeline();
+    void release() override;
+    bool build() override;
+
+    QMetalComputePipelineData *d;
+    uint generation = 0;
+    int lastActiveFrameSlot = -1;
+    friend class QRhiMetal;
+};
+
 struct QMetalCommandBufferData;
 struct QMetalSwapChain;
 
@@ -239,10 +265,19 @@ struct QMetalCommandBuffer : public QRhiCommandBuffer
     QMetalCommandBufferData *d = nullptr;
     QRhiMetalCommandBufferNativeHandles nativeHandlesStruct;
 
+    enum PassType {
+        NoPass,
+        RenderPass,
+        ComputePass
+    };
+
+    PassType recordingPass;
     QRhiRenderTarget *currentTarget;
-    QRhiGraphicsPipeline *currentPipeline;
+    QRhiGraphicsPipeline *currentGraphicsPipeline;
+    QRhiComputePipeline *currentComputePipeline;
     uint currentPipelineGeneration;
-    QRhiShaderResourceBindings *currentSrb;
+    QRhiShaderResourceBindings *currentGraphicsSrb;
+    QRhiShaderResourceBindings *currentComputeSrb;
     uint currentSrbGeneration;
     int currentResSlot;
     QRhiBuffer *currentIndexBuffer;
@@ -296,6 +331,7 @@ public:
     void destroy() override;
 
     QRhiGraphicsPipeline *createGraphicsPipeline() override;
+    QRhiComputePipeline *createComputePipeline() override;
     QRhiShaderResourceBindings *createShaderResourceBindings() override;
     QRhiBuffer *createBuffer(QRhiBuffer::Type type,
                              QRhiBuffer::UsageFlags usage,
@@ -360,6 +396,11 @@ public:
     void debugMarkEnd(QRhiCommandBuffer *cb) override;
     void debugMarkMsg(QRhiCommandBuffer *cb, const QByteArray &msg) override;
 
+    void beginComputePass(QRhiCommandBuffer *cb, QRhiResourceUpdateBatch *resourceUpdates) override;
+    void endComputePass(QRhiCommandBuffer *cb, QRhiResourceUpdateBatch *resourceUpdates) override;
+    void setComputePipeline(QRhiCommandBuffer *cb, QRhiComputePipeline *ps) override;
+    void dispatch(QRhiCommandBuffer *cb, int x, int y, int z) override;
+
     const QRhiNativeHandles *nativeHandles(QRhiCommandBuffer *cb) override;
     void beginExternal(QRhiCommandBuffer *cb) override;
     void endExternal(QRhiCommandBuffer *cb) override;
@@ -393,8 +434,6 @@ public:
 
     bool importedDevice = false;
     bool importedCmdQueue = false;
-    bool inFrame = false;
-    bool inPass = false;
     QMetalSwapChain *currentSwapChain = nullptr;
     QSet<QMetalSwapChain *> swapchains;
     QRhiMetalNativeHandles nativeHandlesStruct;

src/gui/rhi/qrhinull.cpp

@@ -201,6 +201,11 @@ QRhiGraphicsPipeline *QRhiNull::createGraphicsPipeline()
     return new QNullGraphicsPipeline(this);
 }
 
+QRhiComputePipeline *QRhiNull::createComputePipeline()
+{
+    return new QNullComputePipeline(this);
+}
+
 QRhiShaderResourceBindings *QRhiNull::createShaderResourceBindings()
 {
     return new QNullShaderResourceBindings(this);
@@ -297,6 +302,20 @@ void QRhiNull::debugMarkMsg(QRhiCommandBuffer *cb, const QByteArray &msg)
     Q_UNUSED(msg);
 }
 
+void QRhiNull::setComputePipeline(QRhiCommandBuffer *cb, QRhiComputePipeline *ps)
+{
+    Q_UNUSED(cb);
+    Q_UNUSED(ps);
+}
+
+void QRhiNull::dispatch(QRhiCommandBuffer *cb, int x, int y, int z)
+{
+    Q_UNUSED(cb);
+    Q_UNUSED(x);
+    Q_UNUSED(y);
+    Q_UNUSED(z);
+}
+
 const QRhiNativeHandles *QRhiNull::nativeHandles(QRhiCommandBuffer *cb)
 {
     Q_UNUSED(cb);
@@ -395,6 +414,18 @@ void QRhiNull::endPass(QRhiCommandBuffer *cb, QRhiResourceUpdateBatch *resourceU
         resourceUpdate(cb, resourceUpdates);
 }
 
+void QRhiNull::beginComputePass(QRhiCommandBuffer *cb, QRhiResourceUpdateBatch *resourceUpdates)
+{
+    if (resourceUpdates)
+        resourceUpdate(cb, resourceUpdates);
+}
+
+void QRhiNull::endComputePass(QRhiCommandBuffer *cb, QRhiResourceUpdateBatch *resourceUpdates)
+{
+    if (resourceUpdates)
+        resourceUpdate(cb, resourceUpdates);
+}
+
 QNullBuffer::QNullBuffer(QRhiImplementation *rhi, Type type, UsageFlags usage, int size)
     : QRhiBuffer(rhi, type, usage, size)
 {
@@ -647,6 +678,25 @@ bool QNullGraphicsPipeline::build()
     return true;
 }
 
+QNullComputePipeline::QNullComputePipeline(QRhiImplementation *rhi)
+    : QRhiComputePipeline(rhi)
+{
+}
+
+QNullComputePipeline::~QNullComputePipeline()
+{
+    release();
+}
+
+void QNullComputePipeline::release()
+{
+}
+
+bool QNullComputePipeline::build()
+{
+    return true;
+}
+
 QNullCommandBuffer::QNullCommandBuffer(QRhiImplementation *rhi)
     : QRhiCommandBuffer(rhi)
 {

src/gui/rhi/qrhinull_p_p.h

@@ -154,6 +154,14 @@ struct QNullGraphicsPipeline : public QRhiGraphicsPipeline
     bool build() override;
 };
 
+struct QNullComputePipeline : public QRhiComputePipeline
+{
+    QNullComputePipeline(QRhiImplementation *rhi);
+    ~QNullComputePipeline();
+    void release() override;
+    bool build() override;
+};
+
 struct QNullCommandBuffer : public QRhiCommandBuffer
 {
     QNullCommandBuffer(QRhiImplementation *rhi);
@@ -189,6 +197,7 @@ public:
     void destroy() override;
 
     QRhiGraphicsPipeline *createGraphicsPipeline() override;
+    QRhiComputePipeline *createComputePipeline() override;
     QRhiShaderResourceBindings *createShaderResourceBindings() override;
     QRhiBuffer *createBuffer(QRhiBuffer::Type type,
                              QRhiBuffer::UsageFlags usage,
@@ -253,6 +262,11 @@ public:
     void debugMarkEnd(QRhiCommandBuffer *cb) override;
     void debugMarkMsg(QRhiCommandBuffer *cb, const QByteArray &msg) override;
 
+    void beginComputePass(QRhiCommandBuffer *cb, QRhiResourceUpdateBatch *resourceUpdates) override;
+    void endComputePass(QRhiCommandBuffer *cb, QRhiResourceUpdateBatch *resourceUpdates) override;
+    void setComputePipeline(QRhiCommandBuffer *cb, QRhiComputePipeline *ps) override;
+    void dispatch(QRhiCommandBuffer *cb, int x, int y, int z) override;
+
     const QRhiNativeHandles *nativeHandles(QRhiCommandBuffer *cb) override;
     void beginExternal(QRhiCommandBuffer *cb) override;
     void endExternal(QRhiCommandBuffer *cb) override;

src/gui/rhi/qrhivulkan_p_p.h

@@ -61,6 +61,8 @@ static const int QVK_FRAMES_IN_FLIGHT = 2;
 static const int QVK_DESC_SETS_PER_POOL = 128;
 static const int QVK_UNIFORM_BUFFERS_PER_POOL = 256;
 static const int QVK_COMBINED_IMAGE_SAMPLERS_PER_POOL = 256;
+static const int QVK_STORAGE_BUFFERS_PER_POOL = 128;
+static const int QVK_STORAGE_IMAGES_PER_POOL = 128;
 
 static const int QVK_MAX_ACTIVE_TIMESTAMP_PAIRS = 16;
 
@@ -123,12 +125,14 @@ struct QVkTexture : public QRhiTexture
 
     bool prepareBuild(QSize *adjustedSize = nullptr);
     bool finishBuild();
+    VkImageView imageViewForLevel(int level);
 
     VkImage image = VK_NULL_HANDLE;
     VkImageView imageView = VK_NULL_HANDLE;
     QVkAlloc imageAlloc = nullptr;
     VkBuffer stagingBuffers[QVK_FRAMES_IN_FLIGHT];
     QVkAlloc stagingAllocations[QVK_FRAMES_IN_FLIGHT];
+    VkImageView perLevelImageViews[QRhi::MAX_LEVELS];
     bool owns = true;
     QRhiVulkanTextureNativeHandles nativeHandlesStruct;
     struct UsageState {
@@ -246,10 +250,20 @@ struct QVkShaderResourceBindings : public QRhiShaderResourceBindings
         quint64 samplerId;
         uint samplerGeneration;
     };
+    struct BoundStorageImageData {
+        quint64 id;
+        uint generation;
+    };
+    struct BoundStorageBufferData {
+        quint64 id;
+        uint generation;
+    };
     struct BoundResourceData {
         union {
             BoundUniformBufferData ubuf;
             BoundSampledTextureData stex;
+            BoundStorageImageData simage;
+            BoundStorageBufferData sbuf;
         };
     };
     QVector<BoundResourceData> boundResourceData[QVK_FRAMES_IN_FLIGHT];
@@ -273,6 +287,20 @@ struct QVkGraphicsPipeline : public QRhiGraphicsPipeline
     friend class QRhiVulkan;
 };
 
+struct QVkComputePipeline : public QRhiComputePipeline
+{
+    QVkComputePipeline(QRhiImplementation *rhi);
+    ~QVkComputePipeline();
+    void release() override;
+    bool build() override;
+
+    VkPipelineLayout layout = VK_NULL_HANDLE;
+    VkPipeline pipeline = VK_NULL_HANDLE;
+    int lastActiveFrameSlot = -1;
+    uint generation = 0;
+    friend class QRhiVulkan;
+};
+
 struct QVkCommandBuffer : public QRhiCommandBuffer
 {
     QVkCommandBuffer(QRhiImplementation *rhi);
@@ -287,16 +315,25 @@ struct QVkCommandBuffer : public QRhiCommandBuffer
         return &nativeHandlesStruct;
     }
 
+    enum PassType {
+        NoPass,
+        RenderPass,
+        ComputePass
+    };
+
     void resetState() {
         resetCommands();
+        recordingPass = NoPass;
         currentTarget = nullptr;
         resetCachedState();
     }
 
     void resetCachedState() {
-        currentPipeline = nullptr;
+        currentGraphicsPipeline = nullptr;
+        currentComputePipeline = nullptr;
         currentPipelineGeneration = 0;
-        currentSrb = nullptr;
+        currentGraphicsSrb = nullptr;
+        currentComputeSrb = nullptr;
         currentSrbGeneration = 0;
         currentDescSetSlot = -1;
         currentIndexBuffer = VK_NULL_HANDLE;
@@ -306,10 +343,13 @@ struct QVkCommandBuffer : public QRhiCommandBuffer
         memset(currentVertexOffsets, 0, sizeof(currentVertexOffsets));
     }
 
+    PassType recordingPass;
     QRhiRenderTarget *currentTarget;
-    QRhiGraphicsPipeline *currentPipeline;
+    QRhiGraphicsPipeline *currentGraphicsPipeline;
+    QRhiComputePipeline *currentComputePipeline;
     uint currentPipelineGeneration;
-    QRhiShaderResourceBindings *currentSrb;
+    QRhiShaderResourceBindings *currentGraphicsSrb;
+    QRhiShaderResourceBindings *currentComputeSrb;
     uint currentSrbGeneration;
     int currentDescSetSlot;
     VkBuffer currentIndexBuffer;
@@ -343,7 +383,8 @@ struct QVkCommandBuffer : public QRhiCommandBuffer
             DebugMarkerBegin,
             DebugMarkerEnd,
             DebugMarkerInsert,
-            TransitionPassResources
+            TransitionPassResources,
+            Dispatch
         };
         Cmd cmd;
 
@@ -456,6 +497,9 @@ struct QVkCommandBuffer : public QRhiCommandBuffer
             struct {
                 int trackerIndex;
             } transitionResources;
+            struct {
+                int x, y, z;
+            } dispatch;
         } args;
     };
     QVector<Command> commands;
@@ -532,7 +576,12 @@ struct QVkSwapChain : public QRhiSwapChain
         VkFramebuffer fb = VK_NULL_HANDLE;
         VkImage msaaImage = VK_NULL_HANDLE;
         VkImageView msaaImageView = VK_NULL_HANDLE;
-        bool transferSource = false;
+        enum LastUse {
+            ScImageUseNone,
+            ScImageUseRender,
+            ScImageUseTransferSource
+        };
+        LastUse lastUse = ScImageUseNone;
     } imageRes[MAX_BUFFER_COUNT];
 
     struct FrameResources {
@@ -565,6 +614,7 @@ public:
     void destroy() override;
 
     QRhiGraphicsPipeline *createGraphicsPipeline() override;
+    QRhiComputePipeline *createComputePipeline() override;
     QRhiShaderResourceBindings *createShaderResourceBindings() override;
     QRhiBuffer *createBuffer(QRhiBuffer::Type type,
                              QRhiBuffer::UsageFlags usage,
@@ -629,6 +679,11 @@ public:
     void debugMarkEnd(QRhiCommandBuffer *cb) override;
     void debugMarkMsg(QRhiCommandBuffer *cb, const QByteArray &msg) override;
 
+    void beginComputePass(QRhiCommandBuffer *cb, QRhiResourceUpdateBatch *resourceUpdates) override;
+    void endComputePass(QRhiCommandBuffer *cb, QRhiResourceUpdateBatch *resourceUpdates) override;
+    void setComputePipeline(QRhiCommandBuffer *cb, QRhiComputePipeline *ps) override;
+    void dispatch(QRhiCommandBuffer *cb, int x, int y, int z) override;
+
     const QRhiNativeHandles *nativeHandles(QRhiCommandBuffer *cb) override;
     void beginExternal(QRhiCommandBuffer *cb) override;
     void endExternal(QRhiCommandBuffer *cb) override;
@@ -722,6 +777,7 @@ public:
     VkCommandPool cmdPool = VK_NULL_HANDLE;
     int gfxQueueFamilyIdx = -1;
     VkQueue gfxQueue = VK_NULL_HANDLE;
+    bool hasCompute = false;
     quint32 timestampValidBits = 0;
     bool importedAllocator = false;
     QVkAllocator allocator = nullptr;
@@ -765,8 +821,6 @@ public:
     VkFormat optimalDsFormat = VK_FORMAT_UNDEFINED;
     QMatrix4x4 clipCorrectMatrix;
 
-    bool inFrame = false;
-    bool inPass = false;
     QVkSwapChain *currentSwapChain = nullptr;
     QSet<QVkSwapChain *> swapchains;
     QRhiVulkanNativeHandles nativeHandlesStruct;
@@ -830,6 +884,7 @@ public:
                 QVkAlloc allocation;
                 VkBuffer stagingBuffers[QVK_FRAMES_IN_FLIGHT];
                 QVkAlloc stagingAllocations[QVK_FRAMES_IN_FLIGHT];
+                VkImageView extraImageViews[QRhi::MAX_LEVELS];
             } texture;
             struct {
                 VkSampler sampler;

src/gui/rhi/qshaderdescription.cpp

@@ -522,6 +522,21 @@ QVector<QShaderDescription::InOutVariable> QShaderDescription::storageImages() c
     return d->storageImages;
 }
 
+/*!
+    Returns the local size of a compute shader.
+
+    For example, for a compute shader with the following declaration the
+    function returns { 256, 16, 1}.
+
+    \badcode
+        layout(local_size_x = 256, local_size_y = 16, local_size_z = 1) in;
+    \endcode
+ */
+std::array<uint, 3> QShaderDescription::computeShaderLocalSize() const
+{
+    return d->localSize;
+}
+
 static struct TypeTab {
     QString k;
     QShaderDescription::VariableType v;
@@ -799,6 +814,7 @@ static const QString pushConstantBlocksKey = QLatin1String("pushConstantBlocks")
 static const QString storageBlocksKey = QLatin1String("storageBlocks");
 static const QString combinedImageSamplersKey = QLatin1String("combinedImageSamplers");
 static const QString storageImagesKey = QLatin1String("storageImages");
+static const QString localSizeKey = QLatin1String("localSize");
 
 static void addDeco(QJsonObject *obj, const QShaderDescription::InOutVariable &v)
 {
@@ -941,6 +957,11 @@ QJsonDocument QShaderDescriptionPrivate::makeDoc()
     if (!jstorageImages.isEmpty())
         root[storageImagesKey] = jstorageImages;
 
+    QJsonArray jlocalSize;
+    for (int i = 0; i < 3; ++i)
+        jlocalSize.append(QJsonValue(int(localSize[i])));
+    root[localSizeKey] = jlocalSize;
+
     return QJsonDocument(root);
 }
 
@@ -1082,6 +1103,14 @@ void QShaderDescriptionPrivate::loadDoc(const QJsonDocument &doc)
         for (int i = 0; i < images.count(); ++i)
             storageImages.append(inOutVar(images[i].toObject()));
     }
+
+    if (root.contains(localSizeKey)) {
+        QJsonArray localSizeArr = root[localSizeKey].toArray();
+        if (localSizeArr.count() == 3) {
+            for (int i = 0; i < 3; ++i)
+                localSize[i] = localSizeArr[i].toInt();
+        }
+    }
 }
 
 QT_END_NAMESPACE

src/gui/rhi/qshaderdescription_p.h

@@ -51,6 +51,7 @@
 #include <QtGui/qtguiglobal.h>
 #include <QtCore/QString>
 #include <QtCore/QVector>
+#include <array>
 
 QT_BEGIN_NAMESPACE
 
@@ -254,6 +255,8 @@ public:
     QVector<InOutVariable> combinedImageSamplers() const;
     QVector<InOutVariable> storageImages() const;
 
+    std::array<uint, 3> computeShaderLocalSize() const;
+
 private:
     QShaderDescriptionPrivate *d;
     friend struct QShaderDescriptionPrivate;

src/gui/rhi/qshaderdescription_p_p.h

@@ -60,6 +60,7 @@ struct Q_GUI_EXPORT QShaderDescriptionPrivate
     QShaderDescriptionPrivate()
         : ref(1)
     {
+        localSize[0] = localSize[1] = localSize[2] = 0;
     }
 
     QShaderDescriptionPrivate(const QShaderDescriptionPrivate *other)
@@ -70,7 +71,8 @@ struct Q_GUI_EXPORT QShaderDescriptionPrivate
           pushConstantBlocks(other->pushConstantBlocks),
           storageBlocks(other->storageBlocks),
           combinedImageSamplers(other->combinedImageSamplers),
-          storageImages(other->storageImages)
+          storageImages(other->storageImages),
+          localSize(other->localSize)
     {
     }
 
@@ -88,6 +90,7 @@ struct Q_GUI_EXPORT QShaderDescriptionPrivate
     QVector<QShaderDescription::StorageBlock> storageBlocks;
     QVector<QShaderDescription::InOutVariable> combinedImageSamplers;
     QVector<QShaderDescription::InOutVariable> storageImages;
+    std::array<uint, 3> localSize;
 };
 
 QT_END_NAMESPACE

tests/manual/rhi/compressedtexture_bc1/compressedtexture_bc1.cpp

@@ -113,8 +113,8 @@ void Window::customInit()
         qFatal("Failed to load shader pack (fragment)");
 
     d.ps->setShaderStages({
-        { QRhiGraphicsShaderStage::Vertex, vs },
+        { QRhiShaderStage::Vertex, vs },
-        { QRhiGraphicsShaderStage::Fragment, fs }
+        { QRhiShaderStage::Fragment, fs }
     });
 
     QRhiVertexInputLayout inputLayout;

tests/manual/rhi/compressedtexture_bc1_subupload/compressedtexture_bc1_subupload.cpp

@@ -117,8 +117,8 @@ void Window::customInit()
         qFatal("Failed to load shader pack (fragment)");
 
     d.ps->setShaderStages({
-        { QRhiGraphicsShaderStage::Vertex, vs },
+        { QRhiShaderStage::Vertex, vs },
-        { QRhiGraphicsShaderStage::Fragment, fs }
+        { QRhiShaderStage::Fragment, fs }
     });
 
     QRhiVertexInputLayout inputLayout;

tests/manual/rhi/computebuffer/buffer.comp

@@ -0,0 +1,41 @@
+#version 440
+
+layout (local_size_x = 256) in;
+
+struct Data
+{
+    vec2 pos;
+    float dir;
+};
+
+layout(std140, binding = 0) buffer StorageBuffer
+{
+    Data d[];
+} buf;
+
+layout(std140, binding = 1) uniform UniformBuffer
+{
+    float step;
+    uint count;
+} ubuf;
+
+void main()
+{
+    uint index = gl_GlobalInvocationID.x;
+    if (index < ubuf.count) {
+        vec2 p = buf.d[index].pos;
+        float dir = buf.d[index].dir;
+
+        p.x += dir * ubuf.step * 0.01;
+        if (p.x > 1.0) {
+            p.x = 1.0;
+            buf.d[index].dir *= -1.0;
+        }
+        if (p.x < -1.0) {
+            p.x = -1.0;
+            buf.d[index].dir *= -1.0;
+        }
+
+        buf.d[index].pos = p;
+    }
+}

tests/manual/rhi/computebuffer/buildshaders.bat

@@ -0,0 +1,3 @@
+qsb --glsl "310 es,430" --hlsl 50 --msl 12 buffer.comp -o buffer.comp.qsb
+qsb --glsl "310 es,430" --hlsl 50 --msl 12 main.vert -o main.vert.qsb
+qsb --glsl "310 es,430" --hlsl 50 --msl 12 main.frag -o main.frag.qsb

tests/manual/rhi/computebuffer/computebuffer.cpp

@@ -0,0 +1,202 @@
+/****************************************************************************
+**
+** Copyright (C) 2019 The Qt Company Ltd.
+** Contact: https://www.qt.io/licensing/
+**
+** This file is part of the examples of the Qt Toolkit.
+**
+** $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 "../shared/examplefw.h"
+#include <QRandomGenerator>
+
+// Compute shader example. Writes to a storage buffer from a compute shader,
+// then uses the same buffer as vertex buffer in the vertex stage. This would
+// be typical when implementing particles for example. Here we just simply move
+// the positions back and forth along the X axis.
+
+// Note that the example relies on gl_PointSize which is not supported
+// everywhere. So in some cases the points will be of size 1.
+
+struct {
+    QVector<QRhiResource *> releasePool;
+    QRhiBuffer *sbuf = nullptr;
+    QRhiBuffer *computeUniBuf = nullptr;
+    QRhiShaderResourceBindings *computeBindings = nullptr;
+    QRhiComputePipeline *computePipeline = nullptr;
+    QRhiShaderResourceBindings *graphicsBindings = nullptr;
+    QRhiGraphicsPipeline *graphicsPipeline = nullptr;
+    QRhiResourceUpdateBatch *initialUpdates = nullptr;
+    float step = 0.2f;
+} d;
+
+// these struct must match the std140 packing rules
+struct Data {
+    float pos[2];
+    float dir;
+    quint32 pad[1];
+};
+struct ComputeUBuf {
+    float step;
+    quint32 count;
+};
+
+const int DATA_COUNT = 256 * 128;
+
+const int COMPUTE_UBUF_SIZE = 8;
+
+void Window::customInit()
+{
+    if (!m_r->isFeatureSupported(QRhi::Compute))
+        qFatal("Compute is not supported");
+
+    // compute pass
+
+    d.sbuf = m_r->newBuffer(QRhiBuffer::Immutable,
+                            QRhiBuffer::StorageBuffer | QRhiBuffer::VertexBuffer,
+                            sizeof(Data) * DATA_COUNT);
+    d.sbuf->build();
+    d.releasePool << d.sbuf;
+
+    d.computeUniBuf = m_r->newBuffer(QRhiBuffer::Dynamic, QRhiBuffer::UniformBuffer, COMPUTE_UBUF_SIZE);
+    d.computeUniBuf->build();
+    d.releasePool << d.computeUniBuf;
+
+    d.initialUpdates = m_r->nextResourceUpdateBatch();
+
+    QByteArray data;
+    data.resize(sizeof(Data) * DATA_COUNT);
+    Data *p = reinterpret_cast<Data *>(data.data());
+    QRandomGenerator *rgen = QRandomGenerator::global();
+    for (int i = 0; i < DATA_COUNT; ++i) {
+        p->pos[0] = rgen->bounded(1000) / 500.0f - 1.0f;
+        p->pos[1] = rgen->bounded(1000) / 500.0f - 1.0f;
+        p->dir = rgen->bounded(2) ? 1 : -1;
+        ++p;
+    }
+    d.initialUpdates->uploadStaticBuffer(d.sbuf, data.constData());
+
+    ComputeUBuf ud;
+    ud.step = d.step;
+    ud.count = DATA_COUNT;
+    d.initialUpdates->updateDynamicBuffer(d.computeUniBuf, 0, COMPUTE_UBUF_SIZE, &ud);
+
+    d.computeBindings = m_r->newShaderResourceBindings();
+    d.computeBindings->setBindings({
+                                       QRhiShaderResourceBinding::bufferLoadStore(0, QRhiShaderResourceBinding::ComputeStage, d.sbuf),
+                                       QRhiShaderResourceBinding::uniformBuffer(1, QRhiShaderResourceBinding::ComputeStage, d.computeUniBuf)
+                                   });
+    d.computeBindings->build();
+    d.releasePool << d.computeBindings;
+
+    d.computePipeline = m_r->newComputePipeline();
+    d.computePipeline->setShaderResourceBindings(d.computeBindings);
+    d.computePipeline->setShaderStage({ QRhiShaderStage::Compute, getShader(QLatin1String(":/buffer.comp.qsb")) });
+    d.computePipeline->build();
+    d.releasePool << d.computePipeline;
+
+    // graphics pass
+
+    d.graphicsBindings = m_r->newShaderResourceBindings();
+    d.graphicsBindings->build();
+    d.releasePool << d.graphicsBindings;
+
+    d.graphicsPipeline = m_r->newGraphicsPipeline();
+    d.graphicsPipeline->setTopology(QRhiGraphicsPipeline::Points);
+    d.graphicsPipeline->setShaderStages({
+        { QRhiShaderStage::Vertex, getShader(QLatin1String(":/main.vert.qsb")) },
+        { QRhiShaderStage::Fragment, getShader(QLatin1String(":/main.frag.qsb")) }
+    });
+    QRhiVertexInputLayout inputLayout;
+    inputLayout.setBindings({
+        { 2 * sizeof(float) }
+    });
+    inputLayout.setAttributes({
+        { 0, 0, QRhiVertexInputAttribute::Float2, 0 },
+    });
+    d.graphicsPipeline->setVertexInputLayout(inputLayout);
+    d.graphicsPipeline->setShaderResourceBindings(d.graphicsBindings);
+    d.graphicsPipeline->setRenderPassDescriptor(m_rp);
+    d.graphicsPipeline->build();
+    d.releasePool << d.graphicsPipeline;
+}
+
+void Window::customRelease()
+{
+    qDeleteAll(d.releasePool);
+    d.releasePool.clear();
+}
+
+void Window::customRender()
+{
+    const QSize outputSizeInPixels = m_sc->currentPixelSize();
+    QRhiCommandBuffer *cb = m_sc->currentFrameCommandBuffer();
+    QRhiResourceUpdateBatch *u = m_r->nextResourceUpdateBatch();
+    if (d.initialUpdates) {
+        u->merge(d.initialUpdates);
+        d.initialUpdates->release();
+        d.initialUpdates = nullptr;
+    }
+
+#if 0
+    u->updateDynamicBuffer(d.computeUniBuf, 0, sizeof(float), &d.step);
+    d.step += 0.01f;
+#endif
+
+    // compute pass
+    cb->beginComputePass(u);
+    cb->setComputePipeline(d.computePipeline);
+    cb->setShaderResources();
+    cb->dispatch(DATA_COUNT / 256, 1, 1);
+    cb->endComputePass();
+
+    // graphics pass
+    cb->beginPass(m_sc->currentFrameRenderTarget(), QColor::fromRgbF(0.4f, 0.7f, 0.0f, 1.0f), { 1.0f, 0 });
+    cb->setGraphicsPipeline(d.graphicsPipeline);
+    cb->setViewport({ 0, 0, float(outputSizeInPixels.width()), float(outputSizeInPixels.height()) });
+    QRhiCommandBuffer::VertexInput vbufBinding(d.sbuf, 0);
+    cb->setVertexInput(0, 1, &vbufBinding);
+    cb->draw(DATA_COUNT);
+    cb->endPass();
+}

tests/manual/rhi/computebuffer/computebuffer.pro

@@ -0,0 +1,8 @@
+TEMPLATE = app
+
+QT += gui-private
+
+SOURCES = \
+    computebuffer.cpp
+
+RESOURCES = computebuffer.qrc

tests/manual/rhi/computebuffer/computebuffer.qrc

@@ -0,0 +1,7 @@
+<!DOCTYPE RCC><RCC version="1.0">
+  <qresource>
+    <file>buffer.comp.qsb</file>
+    <file>main.vert.qsb</file>
+    <file>main.frag.qsb</file>
+</qresource>
+</RCC>

tests/manual/rhi/computebuffer/main.frag

@@ -0,0 +1,8 @@
+#version 440
+
+layout(location = 0) out vec4 fragColor;
+
+void main()
+{
+    fragColor = vec4(1.0);
+}

tests/manual/rhi/computebuffer/main.vert

@@ -0,0 +1,11 @@
+#version 440
+
+layout(location = 0) in vec4 position;
+
+out gl_PerVertex { vec4 gl_Position; float gl_PointSize; };
+
+void main()
+{
+    gl_PointSize = 4.0; // required with Vulkan when drawing points
+    gl_Position = position;
+}

tests/manual/rhi/computeimage/buildshaders.bat

@@ -0,0 +1 @@
+qsb --glsl "310 es,430" --hlsl 50 --msl 12 image.comp -o image.comp.qsb

tests/manual/rhi/computeimage/computeimage.cpp

@@ -0,0 +1,228 @@
+/****************************************************************************
+**
+** Copyright (C) 2019 The Qt Company Ltd.
+** Contact: https://www.qt.io/licensing/
+**
+** This file is part of the examples of the Qt Toolkit.
+**
+** $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 "../shared/examplefw.h"
+
+// Compute shader example with image load/store. The texture sampled in the
+// fragment shader is generated by the compute shader.
+
+struct {
+    QVector<QRhiResource *> releasePool;
+
+    QRhiTexture *texIn = nullptr;
+    QRhiTexture *texOut = nullptr;
+    QRhiBuffer *computeUBuf = nullptr;
+    QRhiShaderResourceBindings *computeBindings = nullptr;
+    QRhiComputePipeline *computePipeline = nullptr;
+
+    QRhiBuffer *vbuf = nullptr;
+    QRhiBuffer *ibuf = nullptr;
+    QRhiBuffer *ubuf = nullptr;
+    QRhiSampler *sampler = nullptr;
+    QRhiShaderResourceBindings *srb = nullptr;
+    QRhiGraphicsPipeline *ps = nullptr;
+
+    QRhiResourceUpdateBatch *initialUpdates = nullptr;
+    QSize imageSize;
+    QMatrix4x4 winProj;
+    float factor = 1.0f;
+} d;
+
+static float quadVertexData[] =
+{ // Y up, CCW
+  -0.5f,   0.5f,   0.0f, 0.0f,
+  -0.5f,  -0.5f,   0.0f, 1.0f,
+  0.5f,   -0.5f,   1.0f, 1.0f,
+  0.5f,   0.5f,    1.0f, 0.0f
+};
+
+static quint16 quadIndexData[] =
+{
+    0, 1, 2, 0, 2, 3
+};
+
+void Window::customInit()
+{
+    if (!m_r->isFeatureSupported(QRhi::Compute))
+        qFatal("Compute is not supported");
+
+    d.initialUpdates = m_r->nextResourceUpdateBatch();
+
+    // compute pass
+
+    const QImage image = QImage(QLatin1String(":/qt256.png")).convertToFormat(QImage::Format_RGBA8888);
+    d.imageSize = image.size();
+    d.texIn = m_r->newTexture(QRhiTexture::RGBA8, d.imageSize, 1, QRhiTexture::UsedWithLoadStore);
+    d.texIn->build();
+    d.releasePool << d.texIn;
+
+    d.texOut = m_r->newTexture(QRhiTexture::RGBA8, d.imageSize, 1, QRhiTexture::UsedWithLoadStore);
+    d.texOut->build();
+    d.releasePool << d.texOut;
+
+    d.initialUpdates->uploadTexture(d.texIn, image);
+
+    d.computeUBuf = m_r->newBuffer(QRhiBuffer::Dynamic, QRhiBuffer::UniformBuffer, 4);
+    d.computeUBuf->build();
+    d.releasePool << d.computeUBuf;
+
+    d.computeBindings = m_r->newShaderResourceBindings();
+    d.computeBindings->setBindings({
+                                       QRhiShaderResourceBinding::uniformBuffer(0, QRhiShaderResourceBinding::ComputeStage, d.computeUBuf),
+                                       QRhiShaderResourceBinding::imageLoad(1, QRhiShaderResourceBinding::ComputeStage, d.texIn, 0),
+                                       QRhiShaderResourceBinding::imageStore(2, QRhiShaderResourceBinding::ComputeStage, d.texOut, 0)
+                                   });
+    d.computeBindings->build();
+    d.releasePool << d.computeBindings;
+
+    d.computePipeline = m_r->newComputePipeline();
+    d.computePipeline->setShaderResourceBindings(d.computeBindings);
+    d.computePipeline->setShaderStage({ QRhiShaderStage::Compute, getShader(QLatin1String(":/image.comp.qsb")) });
+    d.computePipeline->build();
+    d.releasePool << d.computePipeline;
+
+    // graphics pass
+
+    d.vbuf = m_r->newBuffer(QRhiBuffer::Immutable, QRhiBuffer::VertexBuffer, sizeof(quadVertexData));
+    d.vbuf->build();
+    d.releasePool << d.vbuf;
+
+    d.initialUpdates->uploadStaticBuffer(d.vbuf, quadVertexData);
+
+    d.ibuf = m_r->newBuffer(QRhiBuffer::Immutable, QRhiBuffer::IndexBuffer, sizeof(quadIndexData));
+    d.ibuf->build();
+    d.releasePool << d.ibuf;
+
+    d.initialUpdates->uploadStaticBuffer(d.ibuf, quadIndexData);
+
+    d.ubuf = m_r->newBuffer(QRhiBuffer::Dynamic, QRhiBuffer::UniformBuffer, 68);
+    d.ubuf->build();
+    d.releasePool << d.ubuf;
+
+    qint32 flip = m_r->isYUpInFramebuffer() ? 1 : 0;
+    d.initialUpdates->updateDynamicBuffer(d.ubuf, 64, 4, &flip);
+
+    d.sampler = m_r->newSampler(QRhiSampler::Linear, QRhiSampler::Linear, QRhiSampler::None,
+                                QRhiSampler::ClampToEdge, QRhiSampler::ClampToEdge);
+    d.releasePool << d.sampler;
+    d.sampler->build();
+
+    d.srb = m_r->newShaderResourceBindings();
+    d.releasePool << d.srb;
+    d.srb->setBindings({
+        QRhiShaderResourceBinding::uniformBuffer(0, QRhiShaderResourceBinding::VertexStage | QRhiShaderResourceBinding::FragmentStage, d.ubuf),
+        QRhiShaderResourceBinding::sampledTexture(1, QRhiShaderResourceBinding::FragmentStage, d.texOut, d.sampler)
+    });
+    d.srb->build();
+
+    d.ps = m_r->newGraphicsPipeline();
+    d.releasePool << d.ps;
+    d.ps->setShaderStages({
+        { QRhiShaderStage::Vertex, getShader(QLatin1String(":/texture.vert.qsb")) },
+        { QRhiShaderStage::Fragment, getShader(QLatin1String(":/texture.frag.qsb")) }
+    });
+    QRhiVertexInputLayout inputLayout;
+    inputLayout.setBindings({
+        { 4 * sizeof(float) }
+    });
+    inputLayout.setAttributes({
+        { 0, 0, QRhiVertexInputAttribute::Float2, 0 },
+        { 0, 1, QRhiVertexInputAttribute::Float2, 2 * sizeof(float) }
+    });
+    d.ps->setVertexInputLayout(inputLayout);
+    d.ps->setShaderResourceBindings(d.srb);
+    d.ps->setRenderPassDescriptor(m_rp);
+    d.ps->build();
+}
+
+void Window::customRelease()
+{
+    qDeleteAll(d.releasePool);
+    d.releasePool.clear();
+}
+
+void Window::customRender()
+{
+    const QSize outputSizeInPixels = m_sc->currentPixelSize();
+    QRhiCommandBuffer *cb = m_sc->currentFrameCommandBuffer();
+    QRhiResourceUpdateBatch *u = m_r->nextResourceUpdateBatch();
+    if (d.initialUpdates) {
+        u->merge(d.initialUpdates);
+        d.initialUpdates->release();
+        d.initialUpdates = nullptr;
+    }
+
+    if (d.winProj != m_proj) {
+        d.winProj = m_proj;
+        QMatrix4x4 mvp = m_proj;
+        mvp.scale(2.5f);
+        u->updateDynamicBuffer(d.ubuf, 0, 64, mvp.constData());
+    }
+
+    u->updateDynamicBuffer(d.computeUBuf, 0, 4, &d.factor);
+    d.factor += 0.1f;
+    if (d.factor >= 50.0f)
+        d.factor = 1.0f;
+
+    cb->beginComputePass(u);
+    cb->setComputePipeline(d.computePipeline);
+    cb->setShaderResources();
+    cb->dispatch(d.imageSize.width() / 16, d.imageSize.height() / 16, 1);
+    cb->endComputePass();
+
+    cb->beginPass(m_sc->currentFrameRenderTarget(), QColor::fromRgbF(0.4f, 0.7f, 0.0f, 1.0f), { 1.0f, 0 });
+    cb->setGraphicsPipeline(d.ps);
+    cb->setViewport({ 0, 0, float(outputSizeInPixels.width()), float(outputSizeInPixels.height()) });
+    cb->setShaderResources();
+    QRhiCommandBuffer::VertexInput vbufBinding(d.vbuf, 0);
+    cb->setVertexInput(0, 1, &vbufBinding, d.ibuf, 0, QRhiCommandBuffer::IndexUInt16);
+    cb->drawIndexed(6);
+    cb->endPass();
+}

tests/manual/rhi/computeimage/computeimage.pro

@@ -0,0 +1,8 @@
+TEMPLATE = app
+
+QT += gui-private
+
+SOURCES = \
+    computeimage.cpp
+
+RESOURCES = computeimage.qrc

tests/manual/rhi/computeimage/computeimage.qrc

@@ -0,0 +1,8 @@
+<!DOCTYPE RCC><RCC version="1.0">
+  <qresource>
+    <file>image.comp.qsb</file>
+    <file alias="texture.vert.qsb">../shared/texture.vert.qsb</file>
+    <file alias="texture.frag.qsb">../shared/texture.frag.qsb</file>
+    <file alias="qt256.png">../shared/qt256.png</file>
+</qresource>
+</RCC>

tests/manual/rhi/computeimage/image.comp

@@ -0,0 +1,20 @@
+#version 440
+
+layout (local_size_x = 16, local_size_y = 16) in;
+
+layout(std140, binding = 0) uniform UniformBuffer
+{
+    float factor;
+} ubuf;
+
+layout (binding = 1, rgba8) uniform readonly image2D texIn;
+layout (binding = 2, rgba8) uniform writeonly image2D texOut;
+
+void main()
+{
+    ivec2 pos = ivec2(gl_GlobalInvocationID.xy);
+    ivec2 d = ivec2(1, 1);
+    vec4 diff = imageLoad(texIn, pos + d) - imageLoad(texIn, pos - d);
+    float c = (diff.x + diff.y + diff.z) / ubuf.factor + 0.5f;
+    imageStore(texOut, pos, vec4(c, c, c, 1.0));
+}

tests/manual/rhi/cubemap/cubemap.cpp

@@ -121,8 +121,8 @@ void Window::customInit()
     QShader fs = getShader(QLatin1String(":/cubemap.frag.qsb"));
     Q_ASSERT(fs.isValid());
     d.ps->setShaderStages({
-        { QRhiGraphicsShaderStage::Vertex, vs },
+        { QRhiShaderStage::Vertex, vs },
-        { QRhiGraphicsShaderStage::Fragment, fs }
+        { QRhiShaderStage::Fragment, fs }
     });
 
     QRhiVertexInputLayout inputLayout;

tests/manual/rhi/floattexture/floattexture.cpp

@@ -266,8 +266,8 @@ void Window::customInit()
     d.ps = m_r->newGraphicsPipeline();
     d.releasePool << d.ps;
     d.ps->setShaderStages({
-        { QRhiGraphicsShaderStage::Vertex, getShader(QLatin1String(":/texture.vert.qsb")) },
+        { QRhiShaderStage::Vertex, getShader(QLatin1String(":/texture.vert.qsb")) },
-        { QRhiGraphicsShaderStage::Fragment, getShader(QLatin1String(":/texture.frag.qsb")) }
+        { QRhiShaderStage::Fragment, getShader(QLatin1String(":/texture.frag.qsb")) }
     });
     QRhiVertexInputLayout inputLayout;
     inputLayout.setBindings({

tests/manual/rhi/hellominimalcrossgfxtriangle/hellominimalcrossgfxtriangle.cpp

@@ -339,8 +339,8 @@ void Window::init()
         qFatal("Failed to load shader pack (fragment)");
 
     m_ps->setShaderStages({
-        { QRhiGraphicsShaderStage::Vertex, vs },
+        { QRhiShaderStage::Vertex, vs },
-        { QRhiGraphicsShaderStage::Fragment, fs }
+        { QRhiShaderStage::Fragment, fs }
     });
 
     QRhiVertexInputLayout inputLayout;

tests/manual/rhi/mrt/mrt.cpp

@@ -157,8 +157,8 @@ void Window::customInit()
     d.ps = m_r->newGraphicsPipeline();
     d.releasePool << d.ps;
     d.ps->setShaderStages({
-        { QRhiGraphicsShaderStage::Vertex, getShader(QLatin1String(":/texture.vert.qsb")) },
+        { QRhiShaderStage::Vertex, getShader(QLatin1String(":/texture.vert.qsb")) },
-        { QRhiGraphicsShaderStage::Fragment, getShader(QLatin1String(":/texture.frag.qsb")) }
+        { QRhiShaderStage::Fragment, getShader(QLatin1String(":/texture.frag.qsb")) }
     });
     QRhiVertexInputLayout inputLayout;
     inputLayout.setBindings({
@@ -197,8 +197,8 @@ void Window::customInit()
     d.triPs = m_r->newGraphicsPipeline();
     d.releasePool << d.triPs;
     d.triPs->setShaderStages({
-        { QRhiGraphicsShaderStage::Vertex, getShader(QLatin1String(":/mrt.vert.qsb")) },
+        { QRhiShaderStage::Vertex, getShader(QLatin1String(":/mrt.vert.qsb")) },
-        { QRhiGraphicsShaderStage::Fragment, getShader(QLatin1String(":/mrt.frag.qsb")) }
+        { QRhiShaderStage::Fragment, getShader(QLatin1String(":/mrt.frag.qsb")) }
     });
     QVector<QRhiGraphicsPipeline::TargetBlend> blends;
     for (int i = 0; i < ATTCOUNT; ++i) {

tests/manual/rhi/msaarenderbuffer/msaarenderbuffer.cpp

@@ -146,8 +146,8 @@ void Window::customInit()
     d.releasePool << d.triPs;
     d.triPs->setSampleCount(4); // must match the render target
     d.triPs->setShaderStages({
-        { QRhiGraphicsShaderStage::Vertex, getShader(QLatin1String(":/color.vert.qsb")) },
+        { QRhiShaderStage::Vertex, getShader(QLatin1String(":/color.vert.qsb")) },
-        { QRhiGraphicsShaderStage::Fragment, getShader(QLatin1String(":/color.frag.qsb")) }
+        { QRhiShaderStage::Fragment, getShader(QLatin1String(":/color.frag.qsb")) }
     });
     QRhiVertexInputLayout inputLayout;
     inputLayout.setBindings({
@@ -178,8 +178,8 @@ void Window::customInit()
     d.ps = m_r->newGraphicsPipeline();
     d.releasePool << d.ps;
     d.ps->setShaderStages({
-        { QRhiGraphicsShaderStage::Vertex, getShader(QLatin1String(":/texture.vert.qsb")) },
+        { QRhiShaderStage::Vertex, getShader(QLatin1String(":/texture.vert.qsb")) },
-        { QRhiGraphicsShaderStage::Fragment, getShader(QLatin1String(":/texture.frag.qsb")) }
+        { QRhiShaderStage::Fragment, getShader(QLatin1String(":/texture.frag.qsb")) }
     });
     inputLayout.setBindings({
         { 4 * sizeof(float) }

tests/manual/rhi/msaatexture/msaatexture.cpp

@@ -164,8 +164,8 @@ void Window::customInit()
     d.psLeft = m_r->newGraphicsPipeline();
     d.releasePool << d.psLeft;
     d.psLeft->setShaderStages({
-        { QRhiGraphicsShaderStage::Vertex, getShader(QLatin1String(":/texture.vert.qsb")) },
+        { QRhiShaderStage::Vertex, getShader(QLatin1String(":/texture.vert.qsb")) },
-        { QRhiGraphicsShaderStage::Fragment, getShader(QLatin1String(":/texture.frag.qsb")) }
+        { QRhiShaderStage::Fragment, getShader(QLatin1String(":/texture.frag.qsb")) }
     });
     QRhiVertexInputLayout inputLayout;
     inputLayout.setBindings({ { 4 * sizeof(float) } });
@@ -181,11 +181,11 @@ void Window::customInit()
     d.psRight = m_r->newGraphicsPipeline();
     d.releasePool << d.psRight;
     d.psRight->setShaderStages({
-        { QRhiGraphicsShaderStage::Vertex, getShader(QLatin1String(":/texture.vert.qsb")) },
+        { QRhiShaderStage::Vertex, getShader(QLatin1String(":/texture.vert.qsb")) },
 #ifndef NO_MSAA
-        { QRhiGraphicsShaderStage::Fragment, getShader(QLatin1String(":/texture_ms4.frag.qsb")) }
+        { QRhiShaderStage::Fragment, getShader(QLatin1String(":/texture_ms4.frag.qsb")) }
 #else
-        { QRhiGraphicsShaderStage::Fragment, getShader(QLatin1String(":/texture.frag.qsb")) }
+        { QRhiShaderStage::Fragment, getShader(QLatin1String(":/texture.frag.qsb")) }
 #endif
     });
     d.psRight->setVertexInputLayout(d.psLeft->vertexInputLayout());
@@ -219,8 +219,8 @@ void Window::customInit()
     d.releasePool << d.triPs;
     d.triPs->setSampleCount(1);
     d.triPs->setShaderStages({
-        { QRhiGraphicsShaderStage::Vertex, getShader(QLatin1String(":/color.vert.qsb")) },
+        { QRhiShaderStage::Vertex, getShader(QLatin1String(":/color.vert.qsb")) },
-        { QRhiGraphicsShaderStage::Fragment, getShader(QLatin1String(":/color.frag.qsb")) }
+        { QRhiShaderStage::Fragment, getShader(QLatin1String(":/color.frag.qsb")) }
     });
     inputLayout.setBindings({
         { 5 * sizeof(float) }

tests/manual/rhi/multiwindow/multiwindow.cpp

@@ -234,8 +234,8 @@ void ensureSharedResources(QRhiRenderPassDescriptor *rp)
             qFatal("Failed to load shader pack (fragment)");
 
         d.ps->setShaderStages({
-                                  { QRhiGraphicsShaderStage::Vertex, vs },
+                                  { QRhiShaderStage::Vertex, vs },
-                                  { QRhiGraphicsShaderStage::Fragment, fs }
+                                  { QRhiShaderStage::Fragment, fs }
                               });
 
         QRhiVertexInputLayout inputLayout;

tests/manual/rhi/multiwindow_threaded/multiwindow_threaded.cpp

@@ -484,8 +484,8 @@ void Renderer::init()
     m_ps->setFrontFace(QRhiGraphicsPipeline::CCW);
 
     m_ps->setShaderStages({
-        { QRhiGraphicsShaderStage::Vertex, getShader(QLatin1String(":/texture.vert.qsb")) },
+        { QRhiShaderStage::Vertex, getShader(QLatin1String(":/texture.vert.qsb")) },
-        { QRhiGraphicsShaderStage::Fragment, getShader(QLatin1String(":/texture.frag.qsb")) }
+        { QRhiShaderStage::Fragment, getShader(QLatin1String(":/texture.frag.qsb")) }
     });
 
     QRhiVertexInputLayout inputLayout;

tests/manual/rhi/offscreen/offscreen.cpp

@@ -264,8 +264,8 @@ int main(int argc, char **argv)
         qFatal("Failed to load shader pack (fragment)");
 
     ps->setShaderStages({
-        { QRhiGraphicsShaderStage::Vertex, vs },
+        { QRhiShaderStage::Vertex, vs },
-        { QRhiGraphicsShaderStage::Fragment, fs }
+        { QRhiShaderStage::Fragment, fs }
     });
 
     QRhiVertexInputLayout inputLayout;

tests/manual/rhi/rhi.pro

@@ -14,7 +14,9 @@ SUBDIRS += \
     offscreen \
     floattexture \
     mrt \
-    shadowmap
+    shadowmap \
+    computebuffer \
+    computeimage
 
 qtConfig(widgets) {
     SUBDIRS += \

tests/manual/rhi/shadowmap/shadowmap.cpp

@@ -128,8 +128,8 @@ void Window::customInit()
     d.ps = m_r->newGraphicsPipeline();
     d.releasePool << d.ps;
     d.ps->setShaderStages({
-        { QRhiGraphicsShaderStage::Vertex, getShader(QLatin1String(":/main.vert.qsb")) },
+        { QRhiShaderStage::Vertex, getShader(QLatin1String(":/main.vert.qsb")) },
-        { QRhiGraphicsShaderStage::Fragment, getShader(QLatin1String(":/main.frag.qsb")) }
+        { QRhiShaderStage::Fragment, getShader(QLatin1String(":/main.frag.qsb")) }
     });
     d.ps->setDepthTest(true);
     d.ps->setDepthWrite(true);
@@ -168,8 +168,8 @@ void Window::customInit()
     d.shadowPs = m_r->newGraphicsPipeline();
     d.releasePool << d.shadowPs;
     d.shadowPs->setShaderStages({
-        { QRhiGraphicsShaderStage::Vertex, getShader(QLatin1String(":/shadowmap.vert.qsb")) },
+        { QRhiShaderStage::Vertex, getShader(QLatin1String(":/shadowmap.vert.qsb")) },
-        { QRhiGraphicsShaderStage::Fragment, getShader(QLatin1String(":/shadowmap.frag.qsb")) }
+        { QRhiShaderStage::Fragment, getShader(QLatin1String(":/shadowmap.frag.qsb")) }
     });
     d.shadowPs->setDepthTest(true);
     d.shadowPs->setDepthWrite(true);

tests/manual/rhi/texuploads/texuploads.cpp

@@ -124,8 +124,8 @@ void Window::customInit()
         qFatal("Failed to load shader pack (fragment)");
 
     d.ps->setShaderStages({
-        { QRhiGraphicsShaderStage::Vertex, vs },
+        { QRhiShaderStage::Vertex, vs },
-        { QRhiGraphicsShaderStage::Fragment, fs }
+        { QRhiShaderStage::Fragment, fs }
     });
 
     QRhiVertexInputLayout inputLayout;

tests/manual/rhi/triquadcube/texturedcuberenderer.cpp

@@ -121,8 +121,8 @@ void TexturedCubeRenderer::initResources(QRhiRenderPassDescriptor *rp)
     QShader fs = getShader(QLatin1String(":/texture.frag.qsb"));
     Q_ASSERT(fs.isValid());
     m_ps->setShaderStages({
-        { QRhiGraphicsShaderStage::Vertex, vs },
+        { QRhiShaderStage::Vertex, vs },
-        { QRhiGraphicsShaderStage::Fragment, fs }
+        { QRhiShaderStage::Fragment, fs }
     });
 
     QRhiVertexInputLayout inputLayout;

tests/manual/rhi/triquadcube/triangleoncuberenderer.cpp

@@ -138,8 +138,8 @@ void TriangleOnCubeRenderer::initResources(QRhiRenderPassDescriptor *rp)
     QShader fs = getShader(QLatin1String(":/texture.frag.qsb"));
     Q_ASSERT(fs.isValid());
     m_ps->setShaderStages({
-        { QRhiGraphicsShaderStage::Vertex, vs },
+        { QRhiShaderStage::Vertex, vs },
-        { QRhiGraphicsShaderStage::Fragment, fs }
+        { QRhiShaderStage::Fragment, fs }
     });
 
     QRhiVertexInputLayout inputLayout;

tests/manual/rhi/triquadcube/trianglerenderer.cpp

@@ -112,8 +112,8 @@ void TriangleRenderer::initResources(QRhiRenderPassDescriptor *rp)
     QShader fs = getShader(QLatin1String(":/color.frag.qsb"));
     Q_ASSERT(fs.isValid());
     m_ps->setShaderStages({
-        { QRhiGraphicsShaderStage::Vertex, vs },
+        { QRhiShaderStage::Vertex, vs },
-        { QRhiGraphicsShaderStage::Fragment, fs }
+        { QRhiShaderStage::Fragment, fs }
     });
 
     QRhiVertexInputLayout inputLayout;

tests/manual/rhi/triquadcube/triquadcube.cpp

@@ -171,6 +171,7 @@ void Window::customInit()
     qDebug("isFeatureSupported(NPOTTextureRepeat): %d", m_r->isFeatureSupported(QRhi::NPOTTextureRepeat));
     qDebug("isFeatureSupported(RedOrAlpha8IsRed): %d", m_r->isFeatureSupported(QRhi::RedOrAlpha8IsRed));
     qDebug("isFeatureSupported(ElementIndexUint): %d", m_r->isFeatureSupported(QRhi::ElementIndexUint));
+    qDebug("isFeatureSupported(Compute): %d", m_r->isFeatureSupported(QRhi::Compute));
     qDebug("Min 2D texture width/height: %d", m_r->resourceLimit(QRhi::TextureSizeMin));
     qDebug("Max 2D texture width/height: %d", m_r->resourceLimit(QRhi::TextureSizeMax));
     qDebug("Max color attachment count: %d", m_r->resourceLimit(QRhi::MaxColorAttachments));