vk: update vulkan memory allocator lib

2025-06-12 15:23:38 -07:00
102 changed files with 22512 additions and 22016 deletions
@@ -6,4 +6,4 @@ We are chaning the way Vulkan buffers are handled. We need to switch over to a m
 **If you are cherry-picking a commit into an rc/ branch**: add the release note under the
 appropriate header in [RELEASE_NOTES.md](./RELEASE_NOTES.md).

-## Release notes for next branch cut
+## Release notes for next branch cut
@@ -7,9 +7,11 @@ A new header is inserted each time a *tag* is created.
 Instead, if you are authoring a PR for the main branch, add your release note to
 [NEW_RELEASE_NOTES.md](./NEW_RELEASE_NOTES.md).

-## v1.61.2
+## v1.62.0

 - samples: samples now have a CLI to select backend api
+- materials: sampler now export their type in the material binary [⚠️ **New Material Version**]
+- materials: new mutable specialization constants feature. See the [materials documentation](https://google.github.io/filament/Materials.html) for details. [⚠️ **New Material Version**]

 ## v1.61.1

@@ -60,6 +60,14 @@ static void setParameter(JNIEnv* env, jlong nativeMaterialInstance, jstring name
    env->ReleaseStringUTFChars(name_, name);
 }

+template<typename T>
+static void setConstant(JNIEnv* env, jlong nativeMaterialInstance, jstring name_, T v) {
+    MaterialInstance* instance = (MaterialInstance*) nativeMaterialInstance;
+    const char *name = env->GetStringUTFChars(name_, 0);
+    instance->setConstant(name, v);
+    env->ReleaseStringUTFChars(name_, name);
+}
+
 extern "C"
 JNIEXPORT void JNICALL
 Java_com_google_android_filament_MaterialInstance_nSetParameterBool(JNIEnv *env, jclass,
@@ -264,6 +272,13 @@ Java_com_google_android_filament_MaterialInstance_nSetParameterTexture(
    env->ReleaseStringUTFChars(name_, name);
 }

+extern "C"
+JNIEXPORT void JNICALL
+Java_com_google_android_filament_MaterialInstance_nSetConstantBool(JNIEnv *env, jclass,
+        jlong nativeMaterialInstance, jstring name_, jboolean x) {
+    setConstant(env, nativeMaterialInstance, name_, bool(x));
+}
+
 extern "C"
 JNIEXPORT void JNICALL
 Java_com_google_android_filament_MaterialInstance_nSetScissor(
@@ -402,6 +402,16 @@ public class MaterialInstance {
        nSetParameterFloat4(getNativeObject(), name, color[0], color[1], color[2], color[3]);
    }

+    /**
+     * Sets the value of a bool constant.
+     *
+     * @param name the name of the material constant
+     * @param x    the value of the material constant
+     */
+    public void setConstant(@NonNull String name, boolean x) {
+        nSetConstantBool(getNativeObject(), name, x);
+    }
+
    /**
     * Set-up a custom scissor rectangle; by default it is disabled.
     *
@@ -921,6 +931,9 @@ public class MaterialInstance {
            @NonNull String name, int element, @NonNull @Size(min = 1) float[] v,
            @IntRange(from = 0) int offset, @IntRange(from = 1) int count);

+    private static native void nSetConstantBool(long nativeMaterialInstance,
+            @NonNull String name, boolean x);
+
    private static native void nSetParameterTexture(long nativeMaterialInstance,
            @NonNull String name, long nativeTexture, long sampler);

@@ -204,10 +204,9 @@
      "license": "Apache-2.0"
    },
    "node_modules/brace-expansion": {
-      "version": "1.1.12",
-      "resolved": "https://registry.npmjs.org/brace-expansion/-/brace-expansion-1.1.12.tgz",
-      "integrity": "sha512-9T9UjW3r0UW5c1Q7GTwllptXwhvYmEzFhzMfZ9H7FQWt+uZePjZPjBP/W1ZEyZ1twGWom5/56TF4lPcqjnDHcg==",
-      "license": "MIT",
+      "version": "1.1.11",
+      "resolved": "https://registry.npmjs.org/brace-expansion/-/brace-expansion-1.1.11.tgz",
+      "integrity": "sha512-iCuPHDFgrHX7H2vEI/5xpz07zSHB00TpugqhmYtVmMO6518mCuRMoOYFldEBl0g187ufozdaHgWKcYFb61qGiA==",
      "dependencies": {
        "balanced-match": "^1.0.0",
        "concat-map": "0.0.1"
@@ -410,9 +409,9 @@
      "integrity": "sha512-Nik3Sc0ncrMK4UUdXQmAnRtzmNQTAAXmXIopizwZ1W1t8QmfJj+zL4OA2I7XPTPW5z5TDqv4hRo/JzouDJnX3A=="
    },
    "brace-expansion": {
-      "version": "1.1.12",
-      "resolved": "https://registry.npmjs.org/brace-expansion/-/brace-expansion-1.1.12.tgz",
-      "integrity": "sha512-9T9UjW3r0UW5c1Q7GTwllptXwhvYmEzFhzMfZ9H7FQWt+uZePjZPjBP/W1ZEyZ1twGWom5/56TF4lPcqjnDHcg==",
+      "version": "1.1.11",
+      "resolved": "https://registry.npmjs.org/brace-expansion/-/brace-expansion-1.1.11.tgz",
+      "integrity": "sha512-iCuPHDFgrHX7H2vEI/5xpz07zSHB00TpugqhmYtVmMO6518mCuRMoOYFldEBl0g187ufozdaHgWKcYFb61qGiA==",
      "requires": {
        "balanced-match": "^1.0.0",
        "concat-map": "0.0.1"
@@ -211,7 +211,7 @@ when using textures.
 This property can dramatically change the appearance of a surface. Non-metallic surfaces have
 chromatic diffuse reflection and achromatic specular reflection (reflected light does not change
 color). Metallic surfaces do not have any diffuse reflection and chromatic specular reflection
-(reflected light takes on the color of the surfaced as defined by `baseColor`). 
+(reflected light takes on the color of the surfaced as defined by `baseColor`).

 The effect of `metallic` is shown in figure [metallicProperty] (click on the image to see a
 larger version).
@@ -247,7 +247,7 @@ The effect of `roughness` on metallic surfaces is shown in figure [roughnessCond
 When refraction through an object is enabled (using a `refractonType` of `thin` or `solid`), the
 `roughness` property will also affect the refractions, as shown in figure
 [roughnessRefractionProperty] (click on the image to see a larger version).
-                             
+
 ![Figure [roughnessRefractionProperty]: Refractive sphere with `roughness` varying from 0.0
 (left) to 1.0 (right)](images/materials/refraction_roughness.png)

@@ -306,7 +306,7 @@ The sheen color controls the color appearance and strength of an optional sheen
 base layer described by the properties above. The sheen layer always sits below the clear coat layer
 if such a layer is present.

-The sheen layer can be used to represent cloth and fabric materials. Please refer to 
+The sheen layer can be used to represent cloth and fabric materials. Please refer to
 section [Cloth model] for more information about cloth and fabric materials.

 The effect of `sheenColor` is shown in figure [materialSheenColor]
@@ -519,14 +519,14 @@ light to bend further away from the initial path.

 Table [commonMatIOR] describes acceptable refractive indices for various types of materials.

-Material                   |        IOR       
+Material                   |        IOR
 --------------------------:|:-----------------
 Air                        | 1.0
-Water                      | 1.33             
-Common liquids             | 1.33 to 1.5      
-Common gemstones           | 1.58 to 2.33     
-Plastics, glass            | 1.5 to 1.58      
-Other dielectric materials | 1.33 to 1.58     
+Water                      | 1.33
+Common liquids             | 1.33 to 1.5
+Common gemstones           | 1.58 to 2.33
+Plastics, glass            | 1.5 to 1.58
+Other dielectric materials | 1.33 to 1.58
 [Table [commonMatIOR]: Index of refraction of common materials]

 The appearance of a refractive material will greatly depend on the `refractionType` and
@@ -1091,30 +1091,37 @@ Value
    `bool` or `number`, depending on the `type` of the constant. The type must be one of the types
    described in table [materialConstantsTypes].

-         Type          |            Description                   |      Default
-:----------------------|:-----------------------------------------|:------------------
-int                    | A signed, 32 bit GLSL int                |         0
-float                  | A single-precision GLSL float            |         0.0
-bool                   | A GLSL bool                              |         false
+    Constants may also be specified as mutable by setting the `mutable` property to `true`. Only
+    `bool` spec constants may be specified as mutable.
+
+| Type  | Description                   | Default | May be mutable? |
+|:------|:------------------------------|:--------|-----------------|
+| int   | A signed, 32 bit GLSL int     |         | no              |
+| float | A single-precision GLSL float | 0.0     | no              |
+| bool  | A GLSL bool                   | false   | yes             |
 [Table [materialConstantsTypes]: Material constants types]

 Description
 :   Lists the constant parameters accepted by your material. These constants can be set, or
    "specialized", at runtime when loading a material package. Multiple materials can be loaded from
-    the same material package with differing constant parameter specializations. Once a material is
-    loaded from a material package, its constant parameters cannot be changed. Compared to regular
-    parameters, constant parameters allow the compiler to generate more efficient code. Access
-    constant parameters from the shader by prefixing the name with `materialConstant_`. For example,
-    a constant parameter named `myConstant` is accessed in the shader as
-    `materialConstant_myConstant`. If a constant parameter is not set at runtime, the default is
-    used.
+    the same material package with differing constant parameter specializations. If a constant
+    parameter is specialized as mutable, it may be changed at any time on a per-instance basis via
+    `MaterialInstance::setConstant(name, value)`. Otherwise, once a material is loaded from a
+    material package, its constant parameters cannot be changed.
+
+    Compared to regular parameters, constant parameters allow the compiler to generate more
+    efficient code. Access constant parameters from the shader by prefixing the name with
+    `materialConstant_`. For example, a constant parameter named `myConstant` is accessed in the
+    shader as `materialConstant_myConstant`. If a constant parameter is not set at runtime, the
+    default is used.

 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ JSON
 material {
    constants : [
        {
           name : overrideAlpha,
-           type : bool
+           type : bool,
+           mutable : true,
        },
        {
           name : customAlpha,
@@ -2553,7 +2560,7 @@ standard skybox material. It produces a list of 2 parameters, named `showSun` an
 respectively a boolean and a cubemap texture.

 ```text
-$ matc --reflect parameters filament/src/materials/skybox.mat 
+$ matc --reflect parameters filament/src/materials/skybox.mat
 {
  "parameters": [
    {
@@ -2569,7 +2576,7 @@ $ matc --reflect parameters filament/src/materials/skybox.mat
    }
  ]
 }
-``` 
+```

 ### --variant-filter

@@ -257,7 +257,6 @@ if (FILAMENT_SUPPORTS_WEBGPU)
    list(APPEND SRCS
            include/backend/platforms/WebGPUPlatform.h
            src/webgpu/platform/WebGPUPlatform.cpp
-            src/webgpu/SpdMipmapGenerator/SpdMipmapGenerator.cpp
            src/webgpu/WebGPUBufferBase.cpp
            src/webgpu/WebGPUBufferBase.h
            src/webgpu/WebGPUBufferObject.cpp
@@ -474,12 +474,7 @@ constexpr std::string_view to_string(DescriptorType type) noexcept {

 enum class DescriptorFlags : uint8_t {
    NONE = 0x00,
-
-    // Indicate a UNIFORM_BUFFER will have dynamic offsets.
-    DYNAMIC_OFFSET = 0x01,
-
-    // To indicate a texture/sampler type should be unfiltered.
-    UNFILTERABLE = 0x02,
+    DYNAMIC_OFFSET = 0x01
 };

 using descriptor_set_t = uint8_t;
@@ -149,13 +149,6 @@ public:
         *      - PlatformEGLAndroid
         */
        bool assertNativeWindowIsValid = false;
-
-        /**
-         * The action to take if a Drawable cannot be acquired. If true, the
-         * frame is aborted instead of panic. This is only supported for:
-         *      - PlatformMetal
-         */
-        bool metalDisablePanicOnDrawableFailure = false;
    };

    Platform() noexcept;
@@ -20,6 +20,7 @@
 #include <utils/CString.h>
 #include <utils/FixedCapacityVector.h>
 #include <utils/Invocable.h>
+#include <utils/bitset.h>

 #include <backend/DriverEnums.h>

@@ -66,6 +67,7 @@ public:
    using DescriptorBindingsInfo = utils::FixedCapacityVector<Descriptor>;
    using DescriptorSetInfo = std::array<DescriptorBindingsInfo, MAX_DESCRIPTOR_SET_COUNT>;
    using SpecializationConstantsInfo = utils::FixedCapacityVector<SpecializationConstant>;
+    using MutableSpecConstantsInfo = utils::bitset8;
    using ShaderBlob = utils::FixedCapacityVector<uint8_t>;
    using ShaderSource = std::array<ShaderBlob, SHADER_TYPE_COUNT>;

@@ -102,7 +104,8 @@ public:
    Program& descriptorBindings(backend::descriptor_set_t set,
            DescriptorBindingsInfo descriptorBindings) noexcept;

-    Program& specializationConstants(SpecializationConstantsInfo specConstants) noexcept;
+    Program& specializationConstants(SpecializationConstantsInfo specConstants,
+            uint32_t firstMutableId, MutableSpecConstantsInfo mutableSpecConstants) noexcept;

    struct PushConstant {
        utils::CString name;
@@ -80,8 +80,18 @@ Program& Program::attributes(AttributesInfo attributes) noexcept {
    return *this;
 }

-Program& Program::specializationConstants(SpecializationConstantsInfo specConstants) noexcept {
+Program& Program::specializationConstants(SpecializationConstantsInfo specConstants,
+        uint32_t firstMutableId, MutableSpecConstantsInfo mutableSpecConstants) noexcept {
+    // String the two lists together.
    mSpecializationConstants = std::move(specConstants);
+    uint32_t firstMutableIndex = specConstants.size();
+    mSpecializationConstants.reserve(specConstants.size() + mutableSpecConstants.size());
+    for (uint32_t i = 0; i < mutableSpecConstants.size(); i++) {
+        mSpecializationConstants[i + firstMutableIndex] = SpecializationConstant {
+                .id = i + firstMutableId,
+                .value = mutableSpecConstants[i],
+        };
+    }
    return *this;
 }

@@ -53,9 +53,6 @@ PlatformMetal::~PlatformMetal() noexcept {
 }

 Driver* PlatformMetal::createDriver(void* /*sharedContext*/, const Platform::DriverConfig& driverConfig) noexcept {
-    pImpl->mDrawableFailureBehavior = driverConfig.metalDisablePanicOnDrawableFailure
-            ? DrawableFailureBehavior::ABORT_FRAME
-            : DrawableFailureBehavior::PANIC;
    return MetalDriverFactory::create(this, driverConfig);
 }

@@ -122,18 +122,12 @@ struct ShaderCompilerService::OpenGLProgramToken : ProgramToken {
    }

    std::optional<CallbackManager::Handle> handle{};
-
-    // Only valid when the blob functions are provided by users. The validity of this variable
-    // doesn't guarantee that the program was created from the cache blob.
    BlobCacheKey key;

    // Used for the `THREAD_POOL` mode.
    mutable Mutex lock;
    mutable Condition cond;
    bool signaled = false;
-
-    // Indicate this program was created from the cache blob.
-    bool retrievedFromBlobCache = false;
 };

 ShaderCompilerService::OpenGLProgramToken::~OpenGLProgramToken() {
@@ -265,7 +259,6 @@ ShaderCompilerService::program_token_t ShaderCompilerService::createProgram(
    // Try retrieving the cached program blob if available.
    token->gl.program = mBlobCache.retrieve(&token->key, mDriver.mPlatform, program);
    if (token->gl.program) {
-        token->retrievedFromBlobCache = true;
        return token;
    }

@@ -344,15 +337,12 @@ GLuint ShaderCompilerService::getProgram(program_token_t& token) {

    assert_invariant(token);// This function should be called when the token is still alive.

-    // Finalize any pending shader compilation tasks only when the token was created without cache.
-    if (!token->retrievedFromBlobCache) {
-        if (token->compiler.mMode == Mode::THREAD_POOL) {
-            auto const job = token->compiler.mCompilerThreadPool.dequeue(token);
-            if (!job) {
-                // It's likely that the job was already completed. But it may be still being
-                // executed at this moment. Just try waiting for it to avoid a race.
-                token->wait();
-            }
+    if (token->compiler.mMode == Mode::THREAD_POOL) {
+        auto const job = token->compiler.mCompilerThreadPool.dequeue(token);
+        if (!job) {
+            // It's likely that the job was already completed. But it may be still being
+            // executed at this moment. Just try waiting for it to avoid a race.
+            token->wait();
        }
    }

@@ -27,8 +27,7 @@ namespace filament::backend {

 VulkanBufferProxy::VulkanBufferProxy(VmaAllocator allocator, VulkanStagePool& stagePool,
        VulkanBufferCache& bufferCache, VulkanBufferUsage usage, uint32_t numBytes)
-    : mAllocator(allocator),
-      mStagePool(stagePool),
+    : mStagePool(stagePool),
      mBufferCache(bufferCache),
      mBuffer(mBufferCache.acquire(usage, numBytes)),
      mUpdatedOffset(0),
@@ -39,10 +38,8 @@ void VulkanBufferProxy::loadFromCpu(VulkanCommandBuffer& commands, const void* c
    // Note: this should be stored within the command buffer before going out of
    // scope, so that the command buffer can manage its lifecycle.
    fvkmemory::resource_ptr<VulkanStage::Segment> stage = mStagePool.acquireStage(numBytes);
-    assert_invariant(stage->memory());
    commands.acquire(stage);
-    memcpy(stage->mapping(), cpuData, numBytes);
-    vmaFlushAllocation(mAllocator, stage->memory(), stage->offset(), numBytes);
+    stage->copy(0, cpuData, numBytes);

    // If there was a previous update, then we need to make sure the following write is properly
    // synced with the previous read.
@@ -40,7 +40,6 @@ public:
    VulkanBufferUsage getUsage() const noexcept;

 private:
-    VmaAllocator mAllocator;
    VulkanStagePool& mStagePool;
    VulkanBufferCache& mBufferCache;

@@ -95,7 +95,7 @@
 #endif

 #ifndef NDEBUG
-#define FVK_DEBUG_FLAGS (FVK_DEBUG_PERFORMANCE | FVK_DEBUG_FORWARDED_FLAG)
+#define FVK_DEBUG_FLAGS (FVK_DEBUG_PERFORMANCE | FVK_DEBUG_FORWARDED_FLAG | FVK_DEBUG_DEBUG_UTILS | FVK_DEBUG_GROUP_MARKERS | FVK_DEBUG_VALIDATION)
 #else
 #define FVK_DEBUG_FLAGS 0
 #endif
@@ -38,15 +38,23 @@ constexpr uint32_t STAGE_SIZE = 1048576;

 fvkmemory::resource_ptr<VulkanStage::Segment> VulkanStage::acquireSegment(
        fvkmemory::ResourceManager* resManager, uint32_t numBytes) {
-    auto segment = fvkmemory::resource_ptr<Segment>::construct(
-        resManager, this, numBytes, mCurrentOffset, [this](uint32_t offset) {
-            mSegments.erase(offset);
-    });
-    mSegments.insert({mCurrentOffset, segment.get()});
+    auto segment = fvkmemory::resource_ptr<Segment>::construct(resManager, mAllocator, this,
+            mCurrentOffset, [this](uint32_t offset) { mSegments.erase(offset); });
+    mSegments.insert({ mCurrentOffset, segment.get() });
+    // constexpr uint32_t BLOCK_SIZE = 16;
+    // uint32_t const additional = BLOCK_SIZE - (numBytes % BLOCK_SIZE);
+    // numBytes += (additional != BLOCK_SIZE ? additional : 0);
    mCurrentOffset += numBytes;
    return segment;
 }

+VkResult VulkanStage::Segment::copy(size_t dstOffset, void const* src, size_t writeSize) {
+    //    uint8_t* mapped = ((uint8_t*) mapping()) + dstOffset;
+    //    memcpy(mapped, src, writeSize);
+    //    return vmaFlushAllocation(mAllocator, memory(), offset() + dstOffset, writeSize);
+    return vmaCopyMemoryToAllocation(mAllocator, src, memory(), offset() + dstOffset, writeSize);
+}
+
 VulkanStagePool::VulkanStagePool(VmaAllocator allocator, fvkmemory::ResourceManager* resManager,
        VulkanCommands* commands, const VkPhysicalDeviceLimits* deviceLimits)
    : mAllocator(allocator),
@@ -125,7 +133,7 @@ VulkanStage* VulkanStagePool::allocateNewStage(uint32_t capacity) {
 #endif
    }

-    return new VulkanStage(memory, buffer, capacity, pMapping);
+    return new VulkanStage(mAllocator, memory, buffer, capacity, pMapping);
 }

 void VulkanStagePool::destroyStage(VulkanStage const*&& stage) {
@@ -35,8 +35,10 @@ class VulkanCommands;
 // into smaller buffers as needed.
 class VulkanStage {
 public:
-    VulkanStage(VmaAllocation memory, VkBuffer buffer, uint32_t capacity, void* mapping)
-        : mMemory(memory),
+    VulkanStage(VmaAllocator allocator, VmaAllocation memory, VkBuffer buffer, uint32_t capacity,
+            void* mapping)
+        : mAllocator(allocator),
+          mMemory(memory),
          mBuffer(buffer),
          mCapacity(capacity),
          mMapping(mapping) {}
@@ -51,10 +53,10 @@ public:
    public:
        using OnRecycle = std::function<void(uint32_t offset)>;

-        Segment(VulkanStage* parentStage, uint32_t capacity, uint32_t offset,
+        Segment(VmaAllocator allocator, VulkanStage* parentStage, uint32_t offset,
                OnRecycle&& onRecycleFn)
-            : mParentStage(parentStage),
-              mCapacity(capacity),
+            : mAllocator(allocator),
+              mParentStage(parentStage),
              mOffset(offset),
              mOnRecycleFn(onRecycleFn) {}

@@ -70,12 +72,12 @@ public:
        Segment& operator=(const Segment& other) = delete;
        Segment& operator=(Segment&& other) = delete;

-        inline VulkanStage* parentStage() const { return mParentStage; }
-        inline VkBuffer buffer() const { return parentStage()->buffer(); }
-        inline VmaAllocation memory() const { return parentStage()->memory(); }
-        inline uint32_t capacity() const { return mCapacity; }
+        inline VkBuffer buffer() const { return mParentStage->buffer(); }
+        inline VmaAllocation memory() const { return mParentStage->memory(); }
        inline uint32_t offset() const { return mOffset; }

+        VkResult copy(size_t dstOffset, void const* src, size_t writeSize);
+
        inline void* mapping() const {
            return reinterpret_cast<void*>(
                    reinterpret_cast<char*>(mParentStage->mapping()) + offset());
@@ -85,9 +87,9 @@ public:
        // Ensure parent class can access the terminate method.
        friend class VulkanStage;

+        VmaAllocator const mAllocator;
        VulkanStage* const mParentStage;
-        const uint32_t mCapacity;
-        const uint32_t mOffset;
+        uint32_t const mOffset;
        OnRecycle mOnRecycleFn;
    };

@@ -109,6 +111,7 @@ public:
            uint32_t numBytes);

 private:
+    VmaAllocator const mAllocator;
    const VmaAllocation mMemory;
    const VkBuffer mBuffer;
    const uint32_t mCapacity;
@@ -219,8 +219,8 @@ VkImageUsageFlags getUsage(VulkanContext const& context, uint8_t samples,
    return usage;
 }

-void adjustedMemcpy(void* mapped, PixelBufferDescriptor const& p, size_t width, size_t height,
-        size_t depth) {
+void adjustedMemcpy(PixelBufferDescriptor const& p, size_t width, size_t height,
+        size_t depth, std::function<void(size_t, uint8_t*, size_t)> cpy) {
    uint8_t* buf = (uint8_t*) p.buffer;
    size_t const pixelSize = PixelBufferDescriptor::computeDataSize(p.format, p.type, 1, 1, 1);
    size_t const pbdStride = p.stride ? p.stride : width;
@@ -243,13 +243,13 @@ void adjustedMemcpy(void* mapped, PixelBufferDescriptor const& p, size_t width,
            for (size_t y = p.top; y < pbdHeight; y++) {
                uint8_t* buf = (uint8_t*) p.buffer +
                               ((p.left * pixelSize) + (y * pbdRowSize) + (z * pbdLayerSize));
-                uint8_t* curMapped = (uint8_t*) mapped + ((y - p.top) * rowSize + z * layerSize);
-                memcpy(curMapped, buf, writeSize);
+                uint32_t const offset = (y - p.top) * rowSize + z * layerSize;
+                cpy(offset, buf, writeSize);
            }
        }
    } else {
        size_t const writeSize = pixelSize * (width * height * depth);
-        memcpy(mapped, buf, writeSize);
+        cpy(0, buf, writeSize);
    }
 }

@@ -525,9 +525,12 @@ void VulkanTexture::updateImage(const PixelBufferDescriptor& data, uint32_t widt
    fvkmemory::resource_ptr<VulkanStage::Segment> stageSegment =
            mState->mStagePool.acquireStage(writeSize);
    assert_invariant(stageSegment->memory());
-    adjustedMemcpy(stageSegment->mapping(), *hostData, width, height, depth);
-    vmaFlushAllocation(mState->mAllocator, stageSegment->memory(), stageSegment->offset(),
-            writeSize);
+    adjustedMemcpy(*hostData, width, height, depth,
+            [&stageSegment](size_t dstOffset, uint8_t* src, size_t numBytes) {
+                stageSegment->copy(dstOffset, src, numBytes);
+            });
+//    vmaFlushAllocation(mState->mAllocator, stageSegment->memory(), stageSegment->offset(),
+//            writeSize);

    VulkanCommandBuffer& commands = mState->mCommands->get();
    VkCommandBuffer const cmdbuf = commands.buffer();
@@ -597,7 +600,10 @@ void VulkanTexture::updateImageWithBlit(const PixelBufferDescriptor& data, uint3
    VulkanStageImage const* stage
            = mState->mStagePool.acquireImage(data.format, data.type, width, height);
    vmaMapMemory(mState->mAllocator, stage->memory, &mapped);
-    adjustedMemcpy(mapped, data, width, height, depth);
+    adjustedMemcpy(data, width, height, depth,
+            [&mapped](size_t, uint8_t* src, size_t numBytes) {
+                memcpy(mapped, src, numBytes);
+            });
    vmaUnmapMemory(mState->mAllocator, stage->memory);
    vmaFlushAllocation(mState->mAllocator, stage->memory, 0, writeSize);

@@ -1,5 +0,0 @@
-WebGPU does not include a trivial way to create mipmaps, and thus we must provide a render pass or compute pass to generate them.
-
-This directory is the result of porting https://github.com/JolifantoBambla/webgpu-spd to C++, which itself is a WebGPU port of https://github.com/GPUOpen-LibrariesAndSDKs/FidelityFX-SDK/blob/main/docs/samples/single-pass-downsampler.md. 
-
-The first version of our port is primarily from a Gemini "Convert this to C++" request, which worked as an MVP with very few changes. Future commits will likely alter this version significantly, but it is being merged in this early state to facilitate rapid iteration. 
@@ -1,956 +0,0 @@
-#include "SpdMipmapGenerator.h"
-#include <sstream>
-#include <stdexcept>
-// C++ port of https://github.com/JolifantoBambla/webgpu-spd for early experiments
-namespace spd {
-
-// Helper to convert enums to strings for map keys or shader code
-const char* to_string(wgpu::TextureFormat format) {
-    switch (format) {
-        case wgpu::TextureFormat::RGBA8Unorm:
-            return "rgba8unorm";
-        case wgpu::TextureFormat::BGRA8Unorm:
-            return "bgra8unorm";
-        case wgpu::TextureFormat::R32Float:
-            return "r32float";
-        case wgpu::TextureFormat::RG32Float:
-            return "rg32float";
-        case wgpu::TextureFormat::RGBA32Float:
-            return "rgba32float";
-        case wgpu::TextureFormat::R16Float:
-            return "r16float";
-        case wgpu::TextureFormat::RG16Float:
-            return "rg16float";
-        case wgpu::TextureFormat::RGBA16Float:
-            return "rgba16float";
-        // Add other formats as needed
-        default:
-            return "rgba8unorm";
-    }
-}
-
-
-std::string MipmapGenerator::GetFilterCode(SPDFilter filter) {
-    switch (filter) {
-        case SPDFilter::Min:
-            return R"(
-    fn spd_reduce_4(v0: vec4<SPDScalar>, v1: vec4<SPDScalar>, v2: vec4<SPDScalar>, v3: vec4<SPDScalar>) -> vec4<SPDScalar> {
-        return min(min(v0, v1), min(v2, v3));
-    }
-    )";
-        case SPDFilter::Max:
-            return R"(
-    fn spd_reduce_4(v0: vec4<SPDScalar>, v1: vec4<SPDScalar>, v2: vec4<SPDScalar>, v3: vec4<SPDScalar>) -> vec4<SPDScalar> {
-        return max(max(v0, v1), max(v2, v3));
-    }
-    )";
-        case SPDFilter::MinMax:
-            return R"(
-    fn spd_reduce_4(v0: vec4<SPDScalar>, v1: vec4<SPDScalar>, v2: vec4<SPDScalar>, v3: vec4<SPDScalar>) -> vec4<SPDScalar> {
-        let max4 = max(max(v0.xy, v1.xy), max(v2.xy, v3.xy));
-        return vec4<SPDScalar>(min(min(v0.x, v1.x), min(v2.x, v3.x)), max(max4.x, max4.y), 0.0, 0.0);
-    }
-    )";
-        case SPDFilter::Average:
-        default:
-            return R"(
-    fn spd_reduce_4(v0: vec4<SPDScalar>, v1: vec4<SPDScalar>, v2: vec4<SPDScalar>, v3: vec4<SPDScalar>) -> vec4<SPDScalar> {
-        return (v0 + v1 + v2 + v3) * 0.25;
-    }
-    )";
-    }
-}
-
-
-MipmapGenerator::MipmapGenerator(const wgpu::Device& device)
-    : m_device(device) {
-    wgpu::BindGroupLayoutEntry bglEntry{};
-    bglEntry.binding = 0;
-    bglEntry.visibility = wgpu::ShaderStage::Compute;
-    bglEntry.buffer.type = wgpu::BufferBindingType::Uniform;
-    bglEntry.buffer.minBindingSize = 16;
-
-    wgpu::BindGroupLayoutDescriptor bglDesc{};
-    bglDesc.entryCount = 1;
-    bglDesc.entries = &bglEntry;
-    m_internalResourcesBindGroupLayout =
-            m_device.CreateBindGroupLayout(&bglDesc); // Basic BGL for <= 6 mips
-
-    wgpu::Limits deviceLimits;
-    m_device.GetLimits(&deviceLimits);
-    // Max number of *output* mips we can generate in a single pass.
-    // This is device.maxStorageTexturesPerShaderStage - 1 (for the source texture binding),
-    // further capped at 11 output mips (which means 12 total texture bindings: 1 source + 11
-    // outputs). WebGPU minspec for maxStorageTexturesPerShaderStage is 4, so
-    // deviceLimits.maxStorageTexturesPerShaderStage >= 4. Thus, m_maxMipsPerPass will be at least
-    // std::min(4u - 1, 11u) = 3.
-    m_maxMipsPerPass =
-            std::min(static_cast<uint32_t>(deviceLimits.maxStorageTexturesPerShaderStage) - 1, 11u);
-    m_maxArrayLayers = deviceLimits.maxStorageTexturesPerShaderStage;
-
-
-    // If we can generate more than 6 mips in a pass, we need an advanced BGL
-    // that includes atomic counter and mid-mip buffer for levels > 5.
-    if (m_maxMipsPerPass > 6) {
-        std::vector<wgpu::BindGroupLayoutEntry> entriesAdvanced(3);
-        // Entry 0: Uniform buffer (DownsamplePassMeta)
-        entriesAdvanced[0].binding = 0;
-        entriesAdvanced[0].visibility = wgpu::ShaderStage::Compute;
-        entriesAdvanced[0].buffer.type = wgpu::BufferBindingType::Uniform;
-        entriesAdvanced[0].buffer.minBindingSize = 16; // sizeof(DownsamplePassMeta) effectively
-
-        // Entry 1: Storage buffer (spd_global_counter)
-        entriesAdvanced[1].binding = 1;
-        entriesAdvanced[1].visibility = wgpu::ShaderStage::Compute;
-        entriesAdvanced[1].buffer.type = wgpu::BufferBindingType::Storage;
-        entriesAdvanced[1].buffer.minBindingSize = 4; // sizeof(atomic<u32>)
-
-        // Entry 2: Storage buffer (mip_dst_6_buffer)
-        entriesAdvanced[2].binding = 2;
-        entriesAdvanced[2].visibility = wgpu::ShaderStage::Compute;
-        entriesAdvanced[2].buffer.type = wgpu::BufferBindingType::Storage;
-        // Size for array<array<array<vec4<f32>, 64>, 64>> which is 16 (bytes for vec4f32) * 64 * 64
-        // * arrayLayers. The shader defines it as [slice][uv.y][uv.x], so 64x64 per slice.
-        // minBindingSize here is for one slice as a baseline. Actual buffer will be larger for
-        // multiple array layers.
-        entriesAdvanced[2].buffer.minBindingSize = 16 * 64 * 64;
-
-        wgpu::BindGroupLayoutDescriptor bglDescAdvanced{};
-        bglDescAdvanced.label = "SPD Internal BGL (Advanced)";
-        bglDescAdvanced.entryCount = static_cast<uint32_t>(entriesAdvanced.size());
-        bglDescAdvanced.entries = entriesAdvanced.data();
-        m_internalResourcesBindGroupLayout_Advanced =
-                m_device.CreateBindGroupLayout(&bglDescAdvanced);
-    }
-}
-
-void MipmapGenerator::PreparePipelines(wgpu::TextureFormat format, SPDFilter filter,
-        bool halfPrecision) {
-    SPDScalarType scalarType = SanitizeScalarType(format, halfPrecision);
-    // Prepare for a reasonable number of mips, up to m_maxMipsPerPass
-    // The loop should go from 1 up to and including m_maxMipsPerPass.
-    // If m_maxMipsPerPass is 0 (e.g. maxStorageTexturesPerShaderStage <=1 ), this loop won't run.
-    for (uint32_t i = 1; i <= m_maxMipsPerPass; ++i) {
-        GetOrCreatePipeline({ format, scalarType, filter, i });
-    }
-}
-
-SPDScalarType MipmapGenerator::SanitizeScalarType(wgpu::TextureFormat format, bool halfPrecision) {
-    std::string formatStr = to_string(format);
-    std::transform(formatStr.begin(), formatStr.end(), formatStr.begin(), ::tolower);
-
-    SPDScalarType texelType = SPDScalarType::F32; // Default to F32
-    if (formatStr.find("sint") != std::string::npos) {
-        texelType = SPDScalarType::I32;
-    } else if (formatStr.find("uint") != std::string::npos) {
-        texelType = SPDScalarType::U32;
-    }
-
-    if (halfPrecision && texelType == SPDScalarType::F32) {
-        if (m_device.HasFeature(wgpu::FeatureName::ShaderF16)) {
-            return SPDScalarType::F16;
-        } else {
-            // Optional: Log a warning here that halfPrecision was requested for a float format
-            // but the device does not support ShaderF16, so F32 will be used.
-        }
-    } else if (halfPrecision && texelType != SPDScalarType::F32) {
-        // Optional: Log a warning here that halfPrecision was requested for a non-float format.
-    }
-    return texelType;
-}
-
-SPDPipeline& MipmapGenerator::GetOrCreatePipeline(const PipelineCacheKey& key) {
-    uint32_t numMipsForPipeline = std::min(key.numMips, m_maxMipsPerPass);
-    if (m_pipelines.find(key) == m_pipelines.end()) {
-        // Create the pipeline
-        SPDPipeline spdPipeline;
-
-        std::vector<wgpu::BindGroupLayoutEntry> mipsBglEntries;
-        // Bindings are 0 for source, 1 to numMipsForPipeline for destinations
-        for (uint32_t i = 0; i <= numMipsForPipeline; ++i) {
-            wgpu::BindGroupLayoutEntry entry{};
-            entry.binding = i;
-            entry.visibility = wgpu::ShaderStage::Compute;
-            if (i == 0) {
-                entry.texture.sampleType = (key.scalarType == SPDScalarType::I32)
-                                                   ? wgpu::TextureSampleType::Sint
-                                           : (key.scalarType == SPDScalarType::U32)
-                                                   ? wgpu::TextureSampleType::Uint
-                                                   : wgpu::TextureSampleType::UnfilterableFloat;
-                entry.texture.viewDimension = wgpu::TextureViewDimension::e2DArray;
-            } else {
-                entry.storageTexture.access = wgpu::StorageTextureAccess::WriteOnly;
-                entry.storageTexture.format = key.format;
-                entry.storageTexture.viewDimension = wgpu::TextureViewDimension::e2DArray;
-            }
-            mipsBglEntries.push_back(entry);
-        }
-
-        wgpu::BindGroupLayoutDescriptor mipsBglDesc{};
-        mipsBglDesc.entryCount = mipsBglEntries.size();
-        mipsBglDesc.entries = mipsBglEntries.data();
-        spdPipeline.mipsBindGroupLayout = m_device.CreateBindGroupLayout(&mipsBglDesc);
-
-        std::string shaderCode = MakeShaderCode(key.format, GetFilterCode(key.filter),
-                numMipsForPipeline, key.scalarType);
-        wgpu::ShaderModuleWGSLDescriptor wgslDesc{};
-        wgslDesc.code = shaderCode.c_str();
-
-        wgpu::ShaderModuleDescriptor shaderModuleDesc{};
-        shaderModuleDesc.nextInChain = &wgslDesc;
-        wgpu::ShaderModule shaderModule = m_device.CreateShaderModule(&shaderModuleDesc);
-
-        wgpu::BindGroupLayout chosenInternalBGL =
-                (numMipsForPipeline > 6 && m_internalResourcesBindGroupLayout_Advanced)
-                        ? m_internalResourcesBindGroupLayout_Advanced
-                        : m_internalResourcesBindGroupLayout;
-
-        wgpu::BindGroupLayout bgls[] = { spdPipeline.mipsBindGroupLayout, chosenInternalBGL };
-        wgpu::PipelineLayoutDescriptor layoutDesc{};
-        layoutDesc.bindGroupLayoutCount = 2;
-        layoutDesc.bindGroupLayouts = bgls;
-
-        wgpu::ComputePipelineDescriptor pipelineDesc{};
-        pipelineDesc.layout = m_device.CreatePipelineLayout(&layoutDesc);
-        pipelineDesc.compute.module = shaderModule;
-        pipelineDesc.compute.entryPoint = "downsample";
-
-        spdPipeline.pipeline = m_device.CreateComputePipeline(&pipelineDesc);
-        m_pipelines[key] = std::move(spdPipeline);
-    }
-    return m_pipelines[key];
-}
-
-
-void MipmapGenerator::Generate(wgpu::CommandEncoder& commandEncoder, wgpu::Texture srcTexture,
-        const SPDPassConfig& config) {
-    if (config.numMips <= 1 || config.numMips <= config.sourceMipLevel + 1) {
-        return;
-    }
-
-    if (m_maxMipsPerPass == 0) {
-        return;
-    }
-
-    PreparePipelines(srcTexture.GetFormat(), config.filter, config.halfPrecision);
-
-    uint32_t totalSrcArrayLayers = srcTexture.GetDepthOrArrayLayers();
-
-    // Loop over array layer chunks
-    for (uint32_t currentBaseArrayLayer = 0; currentBaseArrayLayer < totalSrcArrayLayers;
-            currentBaseArrayLayer += m_maxArrayLayers) {
-
-        uint32_t arrayLayersThisPass =
-                std::min(m_maxArrayLayers, totalSrcArrayLayers - currentBaseArrayLayer);
-
-        uint32_t mipsToGenerateCount = config.numMips - (config.sourceMipLevel + 1);
-        uint32_t currentSourceMipLevel = config.sourceMipLevel;
-
-        while (mipsToGenerateCount > 0) {
-            uint32_t mipsThisPass = std::min(mipsToGenerateCount, m_maxMipsPerPass);
-
-            SPDPassConfig passSpecificConfig = {};
-            passSpecificConfig.filter = config.filter;
-            passSpecificConfig.halfPrecision = config.halfPrecision;
-            passSpecificConfig.targetTexture =
-                    config.targetTexture ? config.targetTexture : srcTexture;
-            passSpecificConfig.numMips = mipsThisPass;
-            passSpecificConfig.sourceMipLevel = currentSourceMipLevel;
-
-            // Pass baseArrayLayer and numArrayLayers directly to generatePass
-            generatePass(commandEncoder, srcTexture, passSpecificConfig, currentBaseArrayLayer,
-                    arrayLayersThisPass);
-
-            mipsToGenerateCount -= mipsThisPass;
-            currentSourceMipLevel += mipsThisPass;
-        }
-    }
-}
-
-void MipmapGenerator::generatePass(wgpu::CommandEncoder& commandEncoder, wgpu::Texture srcTexture,
-        const SPDPassConfig& passConfig, uint32_t baseArrayLayer, uint32_t numArrayLayers) {
-    uint32_t width = srcTexture.GetWidth() >> passConfig.sourceMipLevel;
-    uint32_t height = srcTexture.GetHeight() >> passConfig.sourceMipLevel;
-
-    uint32_t arrayLayerCountForThisPass = numArrayLayers;
-
-    wgpu::Texture target = passConfig.targetTexture ? passConfig.targetTexture : srcTexture;
-    uint32_t numMipsThisPass = passConfig.numMips;
-
-    if (numMipsThisPass == 0) return;
-
-    SPDScalarType scalarType = SanitizeScalarType(srcTexture.GetFormat(), passConfig.halfPrecision);
-    SPDPipeline& spdPipeline = GetOrCreatePipeline(
-            { target.GetFormat(), scalarType, passConfig.filter, numMipsThisPass });
-
-    // --- Create Bind Group 0 (Mips) ---
-    std::vector<wgpu::BindGroupEntry> mipEntries;
-
-    wgpu::TextureViewDescriptor srcViewDesc{};
-    srcViewDesc.dimension = wgpu::TextureViewDimension::e2DArray;
-    srcViewDesc.baseMipLevel = passConfig.sourceMipLevel;
-    srcViewDesc.mipLevelCount = 1;
-    srcViewDesc.baseArrayLayer = baseArrayLayer;
-    srcViewDesc.arrayLayerCount = arrayLayerCountForThisPass;
-
-    wgpu::BindGroupEntry srcEntry{};
-    srcEntry.binding = 0;
-    srcEntry.textureView = srcTexture.CreateView(&srcViewDesc);
-    mipEntries.push_back(srcEntry);
-
-    for (uint32_t i = 0; i < numMipsThisPass; ++i) {
-        wgpu::TextureViewDescriptor dstViewDesc{};
-        dstViewDesc.dimension = wgpu::TextureViewDimension::e2DArray;
-        dstViewDesc.baseMipLevel = passConfig.sourceMipLevel + i + 1;
-        dstViewDesc.mipLevelCount = 1;
-        dstViewDesc.baseArrayLayer = baseArrayLayer;
-        dstViewDesc.arrayLayerCount = arrayLayerCountForThisPass;
-
-        wgpu::BindGroupEntry dstEntry{};
-        dstEntry.binding = i + 1;
-        dstEntry.textureView = target.CreateView(&dstViewDesc);
-        mipEntries.push_back(dstEntry);
-    }
-
-    wgpu::BindGroupDescriptor mipBindGroupDesc{};
-    mipBindGroupDesc.layout = spdPipeline.mipsBindGroupLayout;
-    mipBindGroupDesc.entryCount = mipEntries.size();
-    mipBindGroupDesc.entries = mipEntries.data();
-    wgpu::BindGroup mipsBindGroup = m_device.CreateBindGroup(&mipBindGroupDesc);
-
-    // --- Create Bind Group 1 (Internal Resources) ---
-    uint32_t widthDiv64 = (width + 63) / 64;
-    uint32_t heightDiv64 = (height + 63) / 64;
-    uint32_t numWorkGroups = widthDiv64 * heightDiv64;
-
-    struct DownsamplePassMeta {
-        uint32_t work_group_offset[2] = { 0, 0 };
-        uint32_t num_work_groups;
-        uint32_t mips;
-        uint32_t padding[12];
-    } meta;
-    meta.num_work_groups = numWorkGroups;
-    meta.mips = numMipsThisPass;
-
-    wgpu::BufferDescriptor metaBufferDesc{};
-    metaBufferDesc.size = sizeof(DownsamplePassMeta);
-    metaBufferDesc.usage = wgpu::BufferUsage::Uniform | wgpu::BufferUsage::CopyDst;
-    wgpu::Buffer metaBuffer = m_device.CreateBuffer(&metaBufferDesc);
-    m_device.GetQueue().WriteBuffer(metaBuffer, 0, &meta, sizeof(meta));
-
-    wgpu::BindGroup internalBindGroup;
-    wgpu::BindGroupLayout chosenInternalBGL =
-            (numMipsThisPass > 6 && m_internalResourcesBindGroupLayout_Advanced)
-                    ? m_internalResourcesBindGroupLayout_Advanced
-                    : m_internalResourcesBindGroupLayout;
-
-    if (numMipsThisPass > 6 && m_internalResourcesBindGroupLayout_Advanced) {
-        wgpu::BufferDescriptor atomicCounterBufferDesc{};
-        atomicCounterBufferDesc.size = sizeof(uint32_t) * arrayLayerCountForThisPass;
-        atomicCounterBufferDesc.usage = wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopyDst;
-        wgpu::Buffer atomicCounterBuffer = m_device.CreateBuffer(&atomicCounterBufferDesc);
-
-        commandEncoder.ClearBuffer(atomicCounterBuffer, 0, WGPU_WHOLE_SIZE);
-
-        wgpu::BufferDescriptor midMipBufferDesc{};
-        midMipBufferDesc.size = 64 * 64 * 16 * arrayLayerCountForThisPass;
-        midMipBufferDesc.usage = wgpu::BufferUsage::Storage;
-        wgpu::Buffer midMipBuffer = m_device.CreateBuffer(&midMipBufferDesc);
-
-        std::vector<wgpu::BindGroupEntry> internalEntries(3);
-        internalEntries[0].binding = 0;
-        internalEntries[0].buffer = metaBuffer;
-        internalEntries[0].size = sizeof(DownsamplePassMeta);
-
-        internalEntries[1].binding = 1;
-        internalEntries[1].buffer = atomicCounterBuffer;
-        internalEntries[1].size = atomicCounterBuffer.GetSize();
-
-        internalEntries[2].binding = 2;
-        internalEntries[2].buffer = midMipBuffer;
-        internalEntries[2].size = midMipBuffer.GetSize();
-
-        wgpu::BindGroupDescriptor internalBindGroupDesc{};
-        internalBindGroupDesc.layout = chosenInternalBGL;
-        internalBindGroupDesc.entryCount = static_cast<uint32_t>(internalEntries.size());
-        internalBindGroupDesc.entries = internalEntries.data();
-        internalBindGroup = m_device.CreateBindGroup(&internalBindGroupDesc);
-    } else {
-        wgpu::BindGroupEntry metaEntry{};
-        metaEntry.binding = 0;
-        metaEntry.buffer = metaBuffer;
-        metaEntry.size = sizeof(DownsamplePassMeta);
-
-        wgpu::BindGroupDescriptor internalBindGroupDesc{};
-        internalBindGroupDesc.layout = chosenInternalBGL;
-        internalBindGroupDesc.entryCount = 1;
-        internalBindGroupDesc.entries = &metaEntry;
-        internalBindGroup = m_device.CreateBindGroup(&internalBindGroupDesc);
-    }
-
-    // --- Dispatch ---
-    wgpu::ComputePassEncoder pass = commandEncoder.BeginComputePass();
-    pass.SetPipeline(spdPipeline.pipeline);
-    pass.SetBindGroup(0, mipsBindGroup);
-    pass.SetBindGroup(1, internalBindGroup);
-    pass.DispatchWorkgroups(widthDiv64, heightDiv64, arrayLayerCountForThisPass);
-    pass.End();
-}
-// Main shader generation logic
-
-// Helper function to check if a string is in a vector of strings
-bool includes(const std::vector<std::string>& vec, const std::string& str) {
-    for (const auto& s: vec) {
-        if (s == str) {
-            return true;
-        }
-    }
-    return false;
-}
-
-std::string MakeShaderCode(wgpu::TextureFormat outputFormat, const std::string& filterOp,
-        unsigned int numMips,       // Assuming a default value for numMips
-        SPDScalarType scalarType) { // Default scalarType
-    std::stringstream ss;
-
-    std::string texelType;
-    if (scalarType == SPDScalarType::I32) {
-        texelType = "i32";
-    } else if (scalarType == SPDScalarType::U32) {
-        texelType = "u32";
-    } else {
-        texelType = "f32";
-    }
-
-    bool useF16 = (scalarType == SPDScalarType::F16);
-
-    std::string filterCode = filterOp;
-    if (filterOp == SPD_FILTER_AVERAGE && !includes({ "f32", "f16" }, texelType)) {
-        // Replace "* 0.25" with "/ 4"
-        size_t pos = filterCode.find("* 0.25");
-        if (pos != std::string::npos) {
-            filterCode.replace(pos, std::string("* 0.25").length(), "/ 4");
-        }
-    }
-
-    // Generate mipsBindings
-    std::string mipsBindings;
-    for (unsigned int i = 0; i < numMips; ++i) {
-        mipsBindings += "@group(0) @binding(" + std::to_string(i + 1) + ") var dst_mip_" +
-                        std::to_string(i + 1) + ": texture_storage_2d_array<" +
-                        to_string(outputFormat) + ", write>;\n";
-    }
-
-    // Generate mipsAccessorBody
-    std::string mipsAccessorBody;
-    for (unsigned int i = 0; i < numMips; ++i) {
-        if (i == 5 && numMips > 6) {
-            mipsAccessorBody += " else if mip == 6 {\n";
-            mipsAccessorBody += "                textureStore(dst_mip_6, uv, slice, " +
-                                (useF16 ? "vec4<" + texelType + ">(value)" : "value") + ");\n";
-            mipsAccessorBody += "                mip_dst_6_buffer[slice][uv.y][uv.x] = value;\n";
-            mipsAccessorBody += "            }";
-        } else {
-            if (i != 0) {
-                mipsAccessorBody += " else ";
-            }
-            mipsAccessorBody += "if mip == " + std::to_string(i + 1) + " {\n";
-            mipsAccessorBody += "                textureStore(dst_mip_" + std::to_string(i + 1) +
-                                ", uv, slice, " +
-                                (useF16 ? "vec4<" + texelType + ">(value)" : "value") + ");\n";
-            mipsAccessorBody += "            }";
-        }
-    }
-
-    std::string mipsAccessor =
-            "fn store_dst_mip(value: vec4<SPDScalar>, uv: vec2<u32>, slice: u32, mip: u32) {\n" +
-            mipsAccessorBody + "\n}";
-    std::string midMipAccessor = "return mip_dst_6_buffer[slice][uv.y][uv.x];";
-
-    // Start building the final shader code string
-    ss << R"(
-// This file is part of the FidelityFX SDK.
-//
-// Copyright (C) 2023 Advanced Micro Devices, Inc.
-//
-// Permission is hereby granted, free of charge, to any person obtaining a copy
-// of this software and associated documentation files(the “Software”), to deal
-// in the Software without restriction, including without limitation the rights
-// to use, copy, modify, merge, publish, distribute, sublicense, and /or sell
-// copies of the Software, and to permit persons to whom the Software is
-// furnished to do so, subject to the following conditions :
-//
-// The above copyright notice and this permission notice shall be included in
-// all copies or substantial portions of the Software.
-//
-// THE SOFTWARE IS PROVIDED “AS IS”, WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
-// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
-// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.IN NO EVENT SHALL THE
-// AUTHORS COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
-// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
-// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
-// THE SOFTWARE.
-
-
-// Definitions --------------------------------------------------------------------------------------------------------
-
-)";
-    if (useF16) {
-        ss << "enable f16;\n";
-    }
-    ss << "alias SPDScalar = " << texelType
-       << ";\n\n"; // Using texelType here, assuming SPDScalar maps to it.
-
-    ss << R"(
-// Helpers ------------------------------------------------------------------------------------------------------------
-
-/**
- * A helper function performing a remap 64x1 to 8x8 remapping which is necessary for 2D wave reductions.
- * * The 64-wide lane indices to 8x8 remapping is performed as follows:
- * 00 01 08 09 10 11 18 19
- * 02 03 0a 0b 12 13 1a 1b
- * 04 05 0c 0d 14 15 1c 1d
- * 06 07 0e 0f 16 17 1e 1f
- * 20 21 28 29 30 31 38 39
- * 22 23 2a 2b 32 33 3a 3b
- * 24 25 2c 2d 34 35 3c 3d
- * 26 27 2e 2f 36 37 3e 3f
- * * @param a: The input 1D coordinate to remap.
- *
- * @returns The remapped 2D coordinates.
- */
-fn remap_for_wave_reduction(a: u32) -> vec2<u32> {
-    return vec2<u32>(
-        insertBits(extractBits(a, 2u, 3u), a, 0u, 1u),
-        insertBits(extractBits(a, 3u, 3u), extractBits(a, 1u, 2u), 0u, 2u)
-    );
-}
-
-fn map_to_xy(local_invocation_index: u32) -> vec2<u32> {
-    let sub_xy: vec2<u32> = remap_for_wave_reduction(local_invocation_index % 64);
-    return vec2<u32>(
-        sub_xy.x + 8 * ((local_invocation_index >> 6) % 2),
-        sub_xy.y + 8 * ((local_invocation_index >> 7))
-    );
-}
-
-/*
- * Compute a linear value from a SRGB value.
- * * @param value: The value to convert to linear from SRGB.
- * * @returns A value in SRGB space.
- */
-/*
-fn srgb_to_linear(value: SPDScalar) -> SPDScalar {
-    let j = vec3<SPDScalar>(0.0031308 * 12.92, 12.92, 1.0 / 2.4);
-    let k = vec2<SPDScalar>(1.055, -0.055);
-    return clamp(j.x, value * j.y, pow(value, j.z) * k.x + k.y);
-}
-*/
-
-// Resources & Accessors -----------------------------------------------------------------------------------------------
-struct DownsamplePassMeta {
-    work_group_offset: vec2<u32>,
-    num_work_groups: u32,
-    mips: u32,
-}
-
-// In the original version dst_mip_i is an image2Darray [SPD_MAX_MIP_LEVELS+1], i.e., 12+1, but WGSL doesn't support arrays of textures yet
-// Also these are read_write because for mips 7-13, the workgroup reads from mip level 6 - since most formats don't support read_write access in WGSL yet, we use a single read_write buffer in such cases instead
-@group(0) @binding(0) var src_mip_0: texture_2d_array<)"
-       << texelType << R"(>;
-)" << mipsBindings
-       << R"(
-
-@group(1) @binding(0) var<uniform> downsample_pass_meta : DownsamplePassMeta;
-@group(1) @binding(1) var<storage, read_write> spd_global_counter: array<atomic<u32>>;
-@group(1) @binding(2) var<storage, read_write> mip_dst_6_buffer: array<array<array<vec4<f32>, 64>, 64>>;
-
-fn get_mips() -> u32 {
-    return downsample_pass_meta.mips;
-}
-
-fn get_num_work_groups() -> u32 {
-    return downsample_pass_meta.num_work_groups;
-}
-
-fn get_work_group_offset() -> vec2<u32> {
-    return downsample_pass_meta.work_group_offset;
-}
-
-fn load_src_image(uv: vec2<u32>, slice: u32) -> vec4<SPDScalar> {
-    return vec4<SPDScalar>(textureLoad(src_mip_0, uv, slice, 0));
-}
-
-fn load_mid_mip_image(uv: vec2<u32>, slice: u32) -> vec4<SPDScalar> {
-    )";
-    if (numMips > 6) {
-        ss << midMipAccessor;
-    } else {
-        ss << "return vec4<SPDScalar>();";
-    }
-    ss << R"(
-}
-
-)" << mipsAccessor
-       << R"(
-
-// Workgroup -----------------------------------------------------------------------------------------------------------
-
-var<workgroup> spd_intermediate: array<array<vec4<SPDScalar>, 16>, 16>;
-var<workgroup> spd_counter: atomic<u32>;
-
-fn spd_increase_atomic_counter(slice: u32) {
-    atomicStore(&spd_counter, atomicAdd(&spd_global_counter[slice], 1));
-}
-
-fn spd_get_atomic_counter() -> u32 {
-    return atomicLoad(&spd_counter);
-}
-
-fn spd_reset_atomic_counter(slice: u32) {
-    atomicStore(&spd_global_counter[slice], 0);
-}
-
-// Cotnrol flow --------------------------------------------------------------------------------------------------------
-
-fn spd_barrier() {
-    // in glsl this does: groupMemoryBarrier(); barrier();
-    workgroupBarrier();
-}
-
-// Only last active workgroup should proceed
-fn spd_exit_workgroup(num_work_groups: u32, local_invocation_index: u32, slice: u32) -> bool {
-    // global atomic counter
-    if (local_invocation_index == 0) {
-        spd_increase_atomic_counter(slice);
-    }
-    spd_barrier();
-    return spd_get_atomic_counter() != (num_work_groups - 1);
-}
-
-// Pixel access --------------------------------------------------------------------------------------------------------
-
-)" << filterCode
-       << R"(
-
-fn spd_store(pix: vec2<u32>, out_value: vec4<SPDScalar>, mip: u32, slice: u32) {
-    store_dst_mip(out_value, pix, slice, mip + 1);
-}
-
-fn spd_load_intermediate(x: u32, y: u32) -> vec4<SPDScalar> {
-    return spd_intermediate[x][y];
-}
-
-fn spd_store_intermediate(x: u32, y: u32, value: vec4<SPDScalar>) {
-    spd_intermediate[x][y] = value;
-}
-
-fn spd_reduce_intermediate(i0: vec2<u32>, i1: vec2<u32>, i2: vec2<u32>, i3: vec2<u32>) -> vec4<SPDScalar> {
-    let v0 = spd_load_intermediate(i0.x, i0.y);
-    let v1 = spd_load_intermediate(i1.x, i1.y);
-    let v2 = spd_load_intermediate(i2.x, i2.y);
-    let v3 = spd_load_intermediate(i3.x, i3.y);
-    return spd_reduce_4(v0, v1, v2, v3);
-}
-
-fn spd_reduce_load_4(base: vec2<u32>, slice: u32) -> vec4<SPDScalar> {
-    let v0 = load_src_image(base + vec2<u32>(0, 0), slice);
-    let v1 = load_src_image(base + vec2<u32>(0, 1), slice);
-    let v2 = load_src_image(base + vec2<u32>(1, 0), slice);
-    let v3 = load_src_image(base + vec2<u32>(1, 1), slice);
-    return spd_reduce_4(v0, v1, v2, v3);
-}
-
-fn spd_reduce_load_mid_mip_4(base: vec2<u32>, slice: u32) -> vec4<SPDScalar> {
-    let v0 = load_mid_mip_image(base + vec2<u32>(0, 0), slice);
-    let v1 = load_mid_mip_image(base + vec2<u32>(0, 1), slice);
-    let v2 = load_mid_mip_image(base + vec2<u32>(1, 0), slice);
-    let v3 = load_mid_mip_image(base + vec2<u32>(1, 1), slice);
-    return spd_reduce_4(v0, v1, v2, v3);
-}
-
-// Main logic ---------------------------------------------------------------------------------------------------------
-
-fn spd_downsample_mips_0_1(x: u32, y: u32, workgroup_id: vec2<u32>, local_invocation_index: u32, mip: u32, slice: u32) {
-    var v: array<vec4<SPDScalar>, 4>;
-
-    let workgroup64 = workgroup_id.xy * 64;
-    let workgroup32 = workgroup_id.xy * 32;
-    let workgroup16 = workgroup_id.xy * 16;
-
-    var tex = workgroup64 + vec2<u32>(x * 2, y * 2);
-    var pix = workgroup32 + vec2<u32>(x, y);
-    v[0] = spd_reduce_load_4(tex, slice);
-    spd_store(pix, v[0], 0, slice);
-
-    tex = workgroup64 + vec2<u32>(x * 2 + 32, y * 2);
-    pix = workgroup32 + vec2<u32>(x + 16, y);
-    v[1] = spd_reduce_load_4(tex, slice);
-    spd_store(pix, v[1], 0, slice);
-
-    tex = workgroup64 + vec2<u32>(x * 2, y * 2 + 32);
-    pix = workgroup32 + vec2<u32>(x, y + 16);
-    v[2] = spd_reduce_load_4(tex, slice);
-    spd_store(pix, v[2], 0, slice);
-
-    tex = workgroup64 + vec2<u32>(x * 2 + 32, y * 2 + 32);
-    pix = workgroup32 + vec2<u32>(x + 16, y + 16);
-    v[3] = spd_reduce_load_4(tex, slice);
-    spd_store(pix, v[3], 0, slice);
-
-    if mip <= 1 {
-        return;
-    }
-
-    for (var i = 0u; i < 4u; i++) {
-        spd_store_intermediate(x, y, v[i]);
-        spd_barrier();
-        if local_invocation_index < 64 {
-            v[i] = spd_reduce_intermediate(
-                vec2<u32>(x * 2 + 0, y * 2 + 0),
-                vec2<u32>(x * 2 + 1, y * 2 + 0),
-                vec2<u32>(x * 2 + 0, y * 2 + 1),
-                vec2<u32>(x * 2 + 1, y * 2 + 1)
-            );
-            spd_store(workgroup16 + vec2<u32>(x + (i % 2) * 8, y + (i / 2) * 8), v[i], 1, slice);
-        }
-        spd_barrier();
-    }
-
-    if local_invocation_index < 64 {
-        spd_store_intermediate(x + 0, y + 0, v[0]);
-        spd_store_intermediate(x + 8, y + 0, v[1]);
-        spd_store_intermediate(x + 0, y + 8, v[2]);
-        spd_store_intermediate(x + 8, y + 8, v[3]);
-    }
-}
-
-fn spd_downsample_mip_2(x: u32, y: u32, workgroup_id: vec2<u32>, local_invocation_index: u32, mip: u32, slice: u32) {
-    if local_invocation_index < 64u {
-        let v = spd_reduce_intermediate(
-            vec2<u32>(x * 2 + 0, y * 2 + 0),
-            vec2<u32>(x * 2 + 1, y * 2 + 0),
-            vec2<u32>(x * 2 + 0, y * 2 + 1),
-            vec2<u32>(x * 2 + 1, y * 2 + 1)
-        );
-        spd_store(workgroup_id.xy * 8 + vec2<u32>(x, y), v, mip, slice);
-        // store to LDS, try to reduce bank conflicts
-        // x 0 x 0 x 0 x 0 x 0 x 0 x 0 x 0
-        // 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
-        // 0 x 0 x 0 x 0 x 0 x 0 x 0 x 0 x
-        // 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
-        // x 0 x 0 x 0 x 0 x 0 x 0 x 0 x 0
-        // ...
-        // x 0 x 0 x 0 x 0 x 0 x 0 x 0 x 0
-        spd_store_intermediate(x * 2 + y % 2, y * 2, v);
-    }
-}
-
-fn spd_downsample_mip_3(x: u32, y: u32, workgroup_id: vec2<u32>, local_invocation_index: u32, mip: u32, slice: u32) {
-    if local_invocation_index < 16u {
-        // x 0 x 0
-        // 0 0 0 0
-        // 0 x 0 x
-        // 0 0 0 0
-        let v = spd_reduce_intermediate(
-            vec2<u32>(x * 4 + 0 + 0, y * 4 + 0),
-            vec2<u32>(x * 4 + 2 + 0, y * 4 + 0),
-            vec2<u32>(x * 4 + 0 + 1, y * 4 + 2),
-            vec2<u32>(x * 4 + 2 + 1, y * 4 + 2)
-        );
-        spd_store(workgroup_id.xy * 4 + vec2<u32>(x, y), v, mip, slice);
-        // store to LDS
-        // x 0 0 0 x 0 0 0 x 0 0 0 x 0 0 0
-        // 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
-        // 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
-        // 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
-        // 0 x 0 0 0 x 0 0 0 x 0 0 0 x 0 0
-        // ...
-        // 0 0 x 0 0 0 x 0 0 0 x 0 0 0 x 0
-        // ...
-        // 0 0 0 x 0 0 0 x 0 0 0 x 0 0 0 x
-        // ...
-        spd_store_intermediate(x * 4 + y, y * 4, v);
-    }
-}
-
-fn spd_downsample_mip_4(x: u32, y: u32, workgroup_id: vec2<u32>, local_invocation_index: u32, mip: u32, slice: u32) {
-    if local_invocation_index < 4u {
-        // x 0 0 0 x 0 0 0
-        // ...
-        // 0 x 0 0 0 x 0 0
-        let v = spd_reduce_intermediate(
-            vec2<u32>(x * 8 + 0 + 0 + y * 2, y * 8 + 0),
-            vec2<u32>(x * 8 + 4 + 0 + y * 2, y * 8 + 0),
-            vec2<u32>(x * 8 + 0 + 1 + y * 2, y * 8 + 4),
-            vec2<u32>(x * 8 + 4 + 1 + y * 2, y * 8 + 4)
-        );
-        spd_store(workgroup_id.xy * 2 + vec2<u32>(x, y), v, mip, slice);
-        // store to LDS
-        // x x x x 0 ...
-        // 0 ...
-        spd_store_intermediate(x + y * 2, 0, v);
-    }
-}
-
-fn spd_downsample_mip_5(workgroup_id: vec2<u32>, local_invocation_index: u32, mip: u32, slice: u32) {
-    if local_invocation_index < 1u {
-        // x x x x 0 ...
-        // 0 ...
-        let v = spd_reduce_intermediate(vec2<u32>(0, 0), vec2<u32>(1, 0), vec2<u32>(2, 0), vec2<u32>(3, 0));
-        spd_store(workgroup_id.xy, v, mip, slice);
-    }
-}
-
-fn spd_downsample_next_four(x: u32, y: u32, workgroup_id: vec2<u32>, local_invocation_index: u32, base_mip: u32, mips: u32, slice: u32) {
-    if mips <= base_mip {
-        return;
-    }
-    spd_barrier();
-    spd_downsample_mip_2(x, y, workgroup_id, local_invocation_index, base_mip, slice);
-
-    if mips <= base_mip + 1 {
-        return;
-    }
-    spd_barrier();
-    spd_downsample_mip_3(x, y, workgroup_id, local_invocation_index, base_mip + 1, slice);
-
-    if mips <= base_mip + 2 {
-        return;
-    }
-    spd_barrier();
-    spd_downsample_mip_4(x, y, workgroup_id, local_invocation_index, base_mip + 2, slice);
-
-    if mips <= base_mip + 3 {
-        return;
-    }
-    spd_barrier();
-    spd_downsample_mip_5(workgroup_id, local_invocation_index, base_mip + 3, slice);
-}
-
-fn spd_downsample_last_four(x: u32, y: u32, workgroup_id: vec2<u32>, local_invocation_index: u32, base_mip: u32, mips: u32, slice: u32, exit: bool) {
-    if mips <= base_mip {
-        return;
-    }
-    spd_barrier();
-    if !exit {
-        spd_downsample_mip_2(x, y, workgroup_id, local_invocation_index, base_mip, slice);
-    }
-
-    if mips <= base_mip + 1 {
-        return;
-    }
-    spd_barrier();
-    if !exit {
-        spd_downsample_mip_3(x, y, workgroup_id, local_invocation_index, base_mip + 1, slice);
-    }
-
-    if mips <= base_mip + 2 {
-        return;
-    }
-    spd_barrier();
-    if !exit {
-        spd_downsample_mip_4(x, y, workgroup_id, local_invocation_index, base_mip + 2, slice);
-    }
-
-    if mips <= base_mip + 3 {
-        return;
-    }
-    spd_barrier();
-    if !exit {
-        spd_downsample_mip_5(workgroup_id, local_invocation_index, base_mip + 3, slice);
-    }
-}
-
-fn spd_downsample_mips_6_7(x: u32, y: u32, mips: u32, slice: u32) {
-    var tex = vec2<u32>(x * 4 + 0, y * 4 + 0);
-    var pix = vec2<u32>(x * 2 + 0, y * 2 + 0);
-    let v0 = spd_reduce_load_mid_mip_4(tex, slice);
-    spd_store(pix, v0, 6, slice);
-
-    tex = vec2<u32>(x * 4 + 2, y * 4 + 0);
-    pix = vec2<u32>(x * 2 + 1, y * 2 + 0);
-    let v1 = spd_reduce_load_mid_mip_4(tex, slice);
-    spd_store(pix, v1, 6, slice);
-
-    tex = vec2<u32>(x * 4 + 0, y * 4 + 2);
-    pix = vec2<u32>(x * 2 + 0, y * 2 + 1);
-    let v2 = spd_reduce_load_mid_mip_4(tex, slice);
-    spd_store(pix, v2, 6, slice);
-
-    tex = vec2<u32>(x * 4 + 2, y * 4 + 2);
-    pix = vec2<u32>(x * 2 + 1, y * 2 + 1);
-    let v3 = spd_reduce_load_mid_mip_4(tex, slice);
-    spd_store(pix, v3, 6, slice);
-
-    if mips <= 7 {
-        return;
-    }
-    // no barrier needed, working on values only from the same thread
-
-    let v = spd_reduce_4(v0, v1, v2, v3);
-    spd_store(vec2<u32>(x, y), v, 7, slice);
-    spd_store_intermediate(x, y, v);
-}
-
-fn spd_downsample_last_6(x: u32, y: u32, local_invocation_index: u32, mips: u32, num_work_groups: u32, slice: u32) {
-    if mips <= 6 {
-        return;
-    }
-
-    // increase the global atomic counter for the given slice and check if it's the last remaining thread group:
-    // terminate if not, continue if yes.
-    let exit = spd_exit_workgroup(num_work_groups, local_invocation_index, slice);
-
-    // can't exit directly because subsequent barrier calls break uniform control flow...
-    if !exit {
-        // reset the global atomic counter back to 0 for the next spd dispatch
-        spd_reset_atomic_counter(slice);
-
-        // After mip 5 there is only a single workgroup left that downsamples the remaining up to 64x64 texels.
-        // compute MIP level 6 and 7
-        spd_downsample_mips_6_7(x, y, mips, slice);
-    }
-
-    // compute MIP level 8, 9, 10, 11
-    spd_downsample_last_four(x, y, vec2<u32>(0, 0), local_invocation_index, 8, mips, slice, exit);
-}
-
-/// Downsamples a 64x64 tile based on the work group id.
-/// If after downsampling it's the last active thread group, computes the remaining MIP levels.
-///
-/// @param [in] workGroupID        index of the work group / thread group
-/// @param [in] localInvocationIndex   index of the thread within the thread group in 1D
-/// @param [in] mips             the number of total MIP levels to compute for the input texture
-/// @param [in] numWorkGroups        the total number of dispatched work groups / thread groups for this slice
-/// @param [in] slice             the slice of the input texture
-fn spd_downsample(workgroup_id: vec2<u32>, local_invocation_index: u32, mips: u32, num_work_groups: u32, slice: u32) {
-    let xy = map_to_xy(local_invocation_index);
-    spd_downsample_mips_0_1(xy.x, xy.y, workgroup_id, local_invocation_index, mips, slice);
-    spd_downsample_next_four(xy.x, xy.y, workgroup_id, local_invocation_index, 2, mips, slice);
-)";
-    if (numMips > 6) {
-        ss << "    spd_downsample_last_6(xy.x, xy.y, local_invocation_index, mips, "
-              "num_work_groups, slice);\n";
-    }
-    ss << R"(}
-
-// Entry points -------------------------------------------------------------------------------------------------------
-
-@compute
-@workgroup_size(256, 1, 1)
-fn downsample(@builtin(local_invocation_index) local_invocation_index: u32, @builtin(workgroup_id) workgroup_id: vec3<u32>) {
-    spd_downsample(
-        workgroup_id.xy + get_work_group_offset(),
-        local_invocation_index,
-        get_mips(),
-        get_num_work_groups(),
-        workgroup_id.z
-    );
-}
-)";
-    return ss.str();
-}
-
-
-} // namespace spd
@@ -1,99 +0,0 @@
-#pragma once
-
-#include <optional>
-#include <string>
-#include <tsl/robin_map.h>
-#include <vector>
-#include <webgpu/webgpu_cpp.h>
-// C++ port of https://github.com/JolifantoBambla/webgpu-spd for early experiments
-namespace spd {
-
-// Enum for selecting the downsampling filter.
-enum class SPDFilter { Average, Min, Max, MinMax };
-
-// Enum for shader scalar types.
-enum class SPDScalarType { F32, F16, I32, U32 };
-
-// Configuration for a single mipmap generation pass.
-struct SPDPassConfig {
-    SPDFilter filter = SPDFilter::Average;
-    wgpu::Texture targetTexture = nullptr;
-    uint32_t numMips = 0; // For the public API, this is total desired mips. Internally, mips for
-                          // the current pass.
-    bool halfPrecision = false;
-    uint32_t sourceMipLevel = 0;
-};
-
-// Holds a pipeline and its corresponding bind group layout.
-struct SPDPipeline {
-    wgpu::BindGroupLayout mipsBindGroupLayout = nullptr;
-    wgpu::ComputePipeline pipeline = nullptr;
-};
-
-// Manages pipeline creation, caching, and execution for mipmap generation.
-class MipmapGenerator {
-public:
-    MipmapGenerator(const wgpu::Device& device);
-
-    // Pre-creates pipelines for specified formats and filters.
-    void PreparePipelines(wgpu::TextureFormat format, SPDFilter filter, bool halfPrecision = false);
-
-    // Generates a compute pass for creating mipmaps.
-    void Generate(wgpu::CommandEncoder& commandEncoder, wgpu::Texture srcTexture,
-            const SPDPassConfig& config);
-
-private:
-    void generatePass(wgpu::CommandEncoder& commandEncoder, wgpu::Texture srcTexture,
-            const SPDPassConfig& passConfig, uint32_t baseArrayLayer,
-            uint32_t numArrayLayers); // Internal pass config might differ slightly if needed
-
-    wgpu::Device m_device;
-    wgpu::BindGroupLayout m_internalResourcesBindGroupLayout;
-    wgpu::BindGroupLayout m_internalResourcesBindGroupLayout_Advanced = nullptr;
-
-    // Maximum number of mips that can be generated in a single pass.
-    uint32_t m_maxMipsPerPass;
-    uint32_t m_maxArrayLayers;
-
-    // Key for the pipeline cache.
-    struct PipelineCacheKey {
-        wgpu::TextureFormat format;
-        SPDScalarType scalarType;
-        SPDFilter filter;
-        uint32_t numMips;
-
-        bool operator==(const PipelineCacheKey& other) const {
-            return format == other.format && scalarType == other.scalarType &&
-                   filter == other.filter && numMips == other.numMips;
-        }
-    };
-
-    // Hash function for PipelineCacheKey.
-    struct PipelineCacheKeyHash {
-        std::size_t operator()(const PipelineCacheKey& key) const {
-            std::size_t h1 = std::hash<wgpu::TextureFormat>()(key.format);
-            std::size_t h2 = std::hash<SPDScalarType>()(key.scalarType);
-            std::size_t h3 = std::hash<SPDFilter>()(key.filter);
-            std::size_t h4 = std::hash<uint32_t>()(key.numMips);
-            // A simple way to combine hashes.
-            return h1 ^ (h2 << 1) ^ (h3 << 2) ^ (h4 << 3);
-        }
-    };
-
-    // Cached pipelines.
-    tsl::robin_map<PipelineCacheKey, SPDPipeline, PipelineCacheKeyHash> m_pipelines;
-
-    // Helper methods
-    SPDPipeline& GetOrCreatePipeline(const PipelineCacheKey& key);
-    SPDScalarType SanitizeScalarType(wgpu::TextureFormat format, bool halfPrecision);
-    std::string GetFilterCode(SPDFilter filter);
-};
-// Assuming SPD_FILTER_AVERAGE is a string constant
-const std::string SPD_FILTER_AVERAGE = "value * 0.25"; // Original filter operation
-
-// Generates the WGSL shader code dynamically.
-std::string MakeShaderCode(wgpu::TextureFormat outputFormat,
-        const std::string& filterOp = SPD_FILTER_AVERAGE, uint32_t numMips = 0,
-        SPDScalarType scalarType = SPDScalarType::F32);
-
-} // namespace spd
@@ -1,21 +0,0 @@
-MIT License
-
-Copyright (c) 2024 Lukas Herzberger
-
-Permission is hereby granted, free of charge, to any person obtaining a copy
-of this software and associated documentation files (the "Software"), to deal
-in the Software without restriction, including without limitation the rights
-to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
-copies of the Software, and to permit persons to whom the Software is
-furnished to do so, subject to the following conditions:
-
-The above copyright notice and this permission notice shall be included in all
-copies or substantial portions of the Software.
-
-THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
-IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
-FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
-AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
-LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
-OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
-SOFTWARE.
@@ -15,7 +15,6 @@
 */

 #include "WebGPUDescriptorSetLayout.h"
-#include "WebGPUConstants.h"

 #include <backend/DriverEnums.h>

@@ -118,10 +117,10 @@ WebGPUDescriptorSetLayout::WebGPUDescriptorSetLayout(DescriptorSetLayout const&
                samplerEntryInfo.binding = samplerEntry.binding;
                samplerEntry.visibility = wEntry.visibility;
                wEntry.texture.multisampled = isMultiSampledTypeDescriptor(fEntry.type);
+                // TODO: Set once we have the filtering values
                if (isDepthDescriptor(fEntry.type)) {
                    samplerEntry.sampler.type = wgpu::SamplerBindingType::Comparison;
-                } else if (isIntDescriptor(fEntry.type) ||
-                        any(fEntry.flags & DescriptorFlags::UNFILTERABLE)) {
+                } else if (isIntDescriptor(fEntry.type)) {
                    samplerEntry.sampler.type = wgpu::SamplerBindingType::NonFiltering;
                } else {
                    samplerEntry.sampler.type = wgpu::SamplerBindingType::Filtering;
@@ -152,11 +151,8 @@ WebGPUDescriptorSetLayout::WebGPUDescriptorSetLayout(DescriptorSetLayout const&
        if (isDepthDescriptor(fEntry.type)) {
            wEntry.texture.sampleType = wgpu::TextureSampleType::Depth;
        } else if (isFloatDescriptor(fEntry.type)) {
-            if (any(fEntry.flags & DescriptorFlags::UNFILTERABLE)) {
-                wEntry.texture.sampleType = wgpu::TextureSampleType::UnfilterableFloat;
-            } else {
-                wEntry.texture.sampleType = wgpu::TextureSampleType::Float;
-            }
+            // TODO: Set once we have the filtering values
+            wEntry.texture.sampleType = wgpu::TextureSampleType::Float;
        } else if (isIntDescriptor(fEntry.type)) {
            wEntry.texture.sampleType = wgpu::TextureSampleType::Sint;
        } else if (isUnsignedIntDescriptor(fEntry.type)) {
@@ -32,7 +32,6 @@
 #include "CommandStreamDispatcher.h"
 #include "DriverBase.h"
 #include "private/backend/Dispatcher.h"
-#include "webgpu/WebGPUConstants.h"
 #include <backend/DriverEnums.h>
 #include <backend/Handle.h>
 #include <backend/TargetBufferInfo.h>
@@ -65,14 +64,13 @@ Driver* WebGPUDriver::create(WebGPUPlatform& platform, const Platform::DriverCon

 WebGPUDriver::WebGPUDriver(WebGPUPlatform& platform,
        const Platform::DriverConfig& driverConfig) noexcept
-    : mPlatform{ platform },
-      mAdapter{ mPlatform.requestAdapter(nullptr) },
-      mDevice{ mPlatform.requestDevice(mAdapter) },
-      mQueue{ mDevice.GetQueue() },
-      mMipMapGenerator{ mDevice },
-      mHandleAllocator{ "Handles", driverConfig.handleArenaSize,
-          driverConfig.disableHandleUseAfterFreeCheck, driverConfig.disableHeapHandleTags } {
+    : mPlatform(platform),
+      mHandleAllocator("Handles", driverConfig.handleArenaSize,
+              driverConfig.disableHandleUseAfterFreeCheck, driverConfig.disableHeapHandleTags) {
+    mAdapter = mPlatform.requestAdapter(nullptr);
+    mDevice = mPlatform.requestDevice(mAdapter);
    mDevice.GetLimits(&mDeviceLimits);
+    mQueue = mDevice.GetQueue();
 }

 WebGPUDriver::~WebGPUDriver() noexcept = default;
@@ -740,46 +738,67 @@ void WebGPUDriver::setExternalStream(Handle<HwTexture> textureHandle,
 }

 void WebGPUDriver::generateMipmaps(Handle<HwTexture> textureHandle) {
+    if (!mCommandEncoder) {
+        mMipQueue.push_back(textureHandle);
+        return;
+    }
    auto texture = handleCast<WebGPUTexture>(textureHandle);
    assert_invariant(texture);
    wgpu::Texture wgpuTexture = texture->getTexture();
    assert_invariant(wgpuTexture);

-    const auto usage = wgpuTexture.GetUsage();
-    FILAMENT_CHECK_PRECONDITION(usage & wgpu::TextureUsage::TextureBinding)
-            << "Texture for mipmap generation must have TextureBinding usage. "
-               "supportsMultipleMipLevels for texture format "
-            << static_cast<uint32_t>(wgpuTexture.GetFormat()) << " is "
-            << WebGPUTexture::supportsMultipleMipLevelsViaStorageBinding(wgpuTexture.GetFormat());
-    FILAMENT_CHECK_PRECONDITION(usage & wgpu::TextureUsage::StorageBinding)
-            << "Texture for mipmap generation must have StorageBinding usage. "
-               "supportsMultipleMipLevels for texture format "
-            << static_cast<uint32_t>(wgpuTexture.GetFormat()) << " is "
-            << WebGPUTexture::supportsMultipleMipLevelsViaStorageBinding(wgpuTexture.GetFormat());
+    FILAMENT_CHECK_PRECONDITION(wgpuTexture.GetUsage() & wgpu::TextureUsage::CopySrc)
+            << "Texture intended for mipmap generation (as source) must have CopySrc usage.";
+    FILAMENT_CHECK_PRECONDITION(wgpuTexture.GetUsage() & wgpu::TextureUsage::CopyDst)
+            << "Texture intended for mipmap generation (as destination) must have CopyDst usage.";

-    const uint32_t totalMipLevels = wgpuTexture.GetMipLevelCount();
-    if (totalMipLevels <= 1) {
+    uint32_t mipLevelCount = wgpuTexture.GetMipLevelCount();
+    if (mipLevelCount <= 1) {
        return;
    }
+    FILAMENT_CHECK_PRECONDITION(texture->supportsMultipleMipLevels())
+            << "Calling generateMipmaps(...) on a texture that doesn't support generating them. "
+               "sampler type "
+            << to_string(texture->target) << " levels " << static_cast<uint32_t>(texture->levels)
+            << " wgpu format (enum value) "
+            << static_cast<uint32_t>(texture->getTexture().GetFormat()) << " width "
+            << texture->width << " height " << texture->height << " depth " << texture->depth
+            << " usage flags (as a number) " << static_cast<uint32_t>(texture->usage);

-    // We will record all passes into a single command encoder.
-    wgpu::CommandEncoderDescriptor encoderDesc = {};
-    encoderDesc.label = "Mipmap Command Encoder";
-    wgpu::CommandEncoder encoder = mDevice.CreateCommandEncoder(&encoderDesc);
+    uint32_t width = wgpuTexture.GetWidth();
+    uint32_t height = wgpuTexture.GetHeight();
+    // For 3D textures, depth is > 1. For 2D/Cube/Array, effectively 1 for mip-level copies.
+    uint32_t depth =
+            (texture->target == SamplerType::SAMPLER_3D) ? wgpuTexture.GetDepthOrArrayLayers() : 1;

-    spd::SPDPassConfig spdConfig = { .filter = spd::SPDFilter::Average,
-        .targetTexture = wgpuTexture,
-        .numMips = totalMipLevels,
-        .halfPrecision = false,
-        .sourceMipLevel = 0 };
+    for (uint32_t mipLevel = 0; mipLevel < mipLevelCount - 1; ++mipLevel) {
+        wgpu::TexelCopyTextureInfo sourceCopyInfo{
+            .texture = wgpuTexture,
+            .mipLevel = mipLevel,
+            .aspect = texture->getAspect(),
+        };

-    mMipMapGenerator.Generate(encoder, wgpuTexture, spdConfig);
+        wgpu::TexelCopyTextureInfo destinationCopyInfo{
+            .texture = wgpuTexture,
+            .mipLevel = mipLevel + 1,
+            .aspect = texture->getAspect(),
+        };

-    // Finish the encoder and submit all the passes at once.
-    wgpu::CommandBufferDescriptor cmdBufferDesc = {};
-    cmdBufferDesc.label = "Mipmap Command Buffer";
-    wgpu::CommandBuffer commandBuffer = encoder.Finish(&cmdBufferDesc);
-    mQueue.Submit(1, &commandBuffer);
+        uint32_t dstWidth = std::max(1u, width >> 1);
+        uint32_t dstHeight = std::max(1u, height >> 1);
+        uint32_t dstDepth = std::max(1u, depth >> 1);
+
+        wgpu::Extent3D copySize{ .width = dstWidth,
+            .height = dstHeight,
+            .depthOrArrayLayers = (texture->target == SamplerType::SAMPLER_3D)
+                                          ? dstDepth
+                                          : texture->getArrayLayerCount() };
+        mCommandEncoder.CopyTextureToTexture(&sourceCopyInfo, &destinationCopyInfo, &copySize);
+
+        width = dstWidth;
+        height = dstHeight;
+        depth = dstDepth;
+    }
 }

 void WebGPUDriver::compilePrograms(CompilerPriorityQueue priority,
@@ -913,6 +932,12 @@ void WebGPUDriver::makeCurrent(Handle<HwSwapChain> drawSch, Handle<HwSwapChain>
        .label = "command_encoder"
    };
    mCommandEncoder = mDevice.CreateCommandEncoder(&commandEncoderDescriptor);
+    if (!mMipQueue.empty()) {
+        for (auto& handle: mMipQueue) {
+            generateMipmaps(handle);
+        }
+        mMipQueue.clear();
+    }
    assert_invariant(mCommandEncoder);
 }

@@ -1167,7 +1192,6 @@ void WebGPUDriver::updateDescriptorSetBuffer(Handle<HwDescriptorSet> descriptorS
 void WebGPUDriver::updateDescriptorSetTexture(Handle<HwDescriptorSet> descriptorSetHandle,
        const backend::descriptor_binding_t binding, Handle<HwTexture> textureHandle,
        const SamplerParams params) {
-
    auto bindGroup = handleCast<WebGPUDescriptorSet>(descriptorSetHandle);
    auto texture = handleCast<WebGPUTexture>(textureHandle);

@@ -29,7 +29,6 @@

 #include <utils/compiler.h>

-#include "SpdMipmapGenerator/SpdMipmapGenerator.h"
 #include <tsl/robin_map.h>
 #include <webgpu/webgpu_cpp.h>

@@ -72,12 +71,12 @@ private:
    WebGPUSwapChain* mSwapChain = nullptr;
    uint64_t mNextFakeHandle = 1;
    wgpu::CommandEncoder mCommandEncoder = nullptr;
+    std::vector<Handle<HwTexture>> mMipQueue;
    wgpu::TextureView mTextureView = nullptr;
    wgpu::RenderPassEncoder mRenderPassEncoder = nullptr;
    wgpu::CommandBuffer mCommandBuffer = nullptr;
    WebGPURenderTarget* mDefaultRenderTarget = nullptr;
    WebGPURenderTarget* mCurrentRenderTarget = nullptr;
-    spd::MipmapGenerator mMipMapGenerator;

    tsl::robin_map<uint32_t, wgpu::RenderPipeline> mPipelineMap;

@@ -91,20 +91,16 @@ void WebGPURenderTarget::setUpRenderPassAttachments(wgpu::RenderPassDescriptor&
            .clearValue = { params.clearColor.r, params.clearColor.g, params.clearColor.b,
                params.clearColor.a } });

-        bool const depthReadonly =
-                (params.readOnlyDepthStencil & RenderPassParams::READONLY_DEPTH) > 0;
-
        if (defaultDepthStencilTextureView) {
            mDepthStencilAttachmentDescriptor = {
                .view = defaultDepthStencilTextureView,
-                .depthLoadOp = depthReadonly ? wgpu::LoadOp::Undefined
-                                             : WebGPURenderTarget::getLoadOperation(params,
-                                                       TargetBufferFlags::DEPTH),
-                .depthStoreOp = depthReadonly ? wgpu::StoreOp::Undefined
-                                              : WebGPURenderTarget::getStoreOperation(params,
-                                                        TargetBufferFlags::DEPTH),
+                .depthLoadOp =
+                        WebGPURenderTarget::getLoadOperation(params, TargetBufferFlags::DEPTH),
+                .depthStoreOp =
+                        WebGPURenderTarget::getStoreOperation(params, TargetBufferFlags::DEPTH),
                .depthClearValue = static_cast<float>(params.clearDepth),
-                .depthReadOnly = depthReadonly,
+                .depthReadOnly =
+                        (params.readOnlyDepthStencil & RenderPassParams::READONLY_DEPTH) > 0,
                .stencilLoadOp = wgpu::LoadOp::Undefined,
                .stencilStoreOp = wgpu::StoreOp::Undefined,
                .stencilClearValue = params.clearStencil,
@@ -160,19 +156,14 @@ void WebGPURenderTarget::setUpRenderPassAttachments(wgpu::RenderPassDescriptor&
                    customDepthTextureView ? customDepthTextureView : customStencilTextureView;

            if (hasDepth) {
-                bool const depthReadonly =
-                        (params.readOnlyDepthStencil & RenderPassParams::READONLY_DEPTH) > 0;
-
-
                mDepthStencilAttachmentDescriptor.depthLoadOp =
-                        depthReadonly ? wgpu::LoadOp::Undefined :
                        WebGPURenderTarget::getLoadOperation(params, TargetBufferFlags::DEPTH);
                mDepthStencilAttachmentDescriptor.depthStoreOp =
-                        depthReadonly ? wgpu::StoreOp::Undefined :
                        WebGPURenderTarget::getStoreOperation(params, TargetBufferFlags::DEPTH);
                mDepthStencilAttachmentDescriptor.depthClearValue =
                        static_cast<float>(params.clearDepth);
-                mDepthStencilAttachmentDescriptor.depthReadOnly = depthReadonly;
+                mDepthStencilAttachmentDescriptor.depthReadOnly =
+                        (params.readOnlyDepthStencil & RenderPassParams::READONLY_DEPTH) > 0;
            } else {
                mDepthStencilAttachmentDescriptor.depthLoadOp = wgpu::LoadOp::Undefined;
                mDepthStencilAttachmentDescriptor.depthStoreOp = wgpu::StoreOp::Undefined;
@@ -21,27 +21,50 @@

 #include <backend/DriverEnums.h>

+#include <utils/Panic.h>
+#include <utils/ostream.h>
+
+#if FWGPU_ENABLED(FWGPU_PRINT_SYSTEM)
 #include <dawn/webgpu_cpp_print.h>
+#endif
 #include <webgpu/webgpu_cpp.h>

 #include <cstdint>
+#if FWGPU_ENABLED(FWGPU_PRINT_SYSTEM)
 #include <sstream>
-#include <string>
+#endif
 #include <string_view>

 /**
 * Reusable set of convenience functions for strings -- generally string views, literals, &
- * strings -- used in the WebGPU backend
+ * streaming -- used in the WebGPU backend
 */

 namespace filament::backend {

 #if FWGPU_ENABLED(FWGPU_PRINT_SYSTEM)
+/**
+ * Convenience template to print WGPU types (which don't support Filament's utils::io::ostream)
+ * @tparam WebGPUPrintable e.g. wgpu::FeatureName type
+ * @param out stream to print to
+ * @param printable the instance to print
+ * @return the stream printed to
+ *
+ * Example usage (because including the template operator override itself does not compile):
+ * #if FWGPU_ENABLED(FWGPU_PRINT_SYSTEM)
+ * utils::io::ostream& operator<<(utils::io::ostream& out,
+ *       const wgpu::FeatureName featureName) noexcept {
+ *   return streamInsertWebGPUPrintable(out, featureName);
+ * }
+ * #endif
+ */
 template<typename WebGPUPrintable>
-[[nodiscard]] inline std::string webGPUPrintableToString(const WebGPUPrintable printable) {
-    std::stringstream out;
-    out << printable;
-    return out.str();
+utils::io::ostream& streamInsertWebGPUPrintable(utils::io::ostream& out,
+        const WebGPUPrintable printable) {
+    std::stringstream printableStream;
+    printableStream << printable;
+    out << printableStream.str();
+    return out;
 }
 #endif

@@ -92,17 +115,46 @@ template<typename WebGPUPrintable>
    }
 }

-[[nodiscard]] inline std::string adapterOptionsToString(
-        wgpu::RequestAdapterOptions const& options) {
-    std::stringstream out;
+/**
+ * Convenience template to print wgpu::RequestAdapterOptions
+ * @tparam STREAM_TYPE useful namely for streaming to either an utils::io::ostream OR
+ * utils::details::PanicStream, but also for any kind of stream
+ * @param out stream to print to
+ * @param options instance to be printed
+ * @return stream printed to
+ *
+ * Example usage (because including the template operator override itself does not compile):
+ * template<typename STREAM_TYPE>
+ * STREAM_TYPE& operator<<(STREAM_TYPE& out, wgpu::RequestAdapterOptions const& options) noexcept {
+ *   return streamInsertRequestAdapterOptions(out, options);
+ * }
+ */
+template<typename STREAM_TYPE>
+STREAM_TYPE& streamInsertRequestAdapterOptions(STREAM_TYPE& out,
+        wgpu::RequestAdapterOptions const& options) noexcept {
    out << "power preference " << powerPreferenceToString(options.powerPreference)
        << " force fallback adapter " << bool(options.forceFallbackAdapter) << " backend type "
        << backendTypeToString(options.backendType);
-    return out.str();
+    return out;
 }

-[[nodiscard]] std::string adapterInfoToString(wgpu::AdapterInfo const& info) {
-    std::stringstream out;
+/**
+ * Convenience template to print wgpu::AdapterInfo
+ * @tparam STREAM_TYPE useful namely for streaming to either an utils::io::ostream OR
+ * utils::details::PanicStream, but also for any kind of stream
+ * @param out stream to print to
+ * @param options instance to be printed
+ * @return stream printed to
+ *
+ * Example usage (because including the template operator override itself does not compile):
+ * template<typename STREAM_TYPE>
+ * STREAM_TYPE& operator<<(STREAM_TYPE& out, wgpu::AdapterInfo const& info) noexcept {
+ *   return streamInsertRequestAdapterInfo(out, info);
+ * }
+ */
+template<class STREAM_TYPE>
+STREAM_TYPE& streamInsertRequestAdapterInfo(STREAM_TYPE& out,
+        wgpu::AdapterInfo const& info) noexcept {
    out << "vendor (" << info.vendorID << ") '" << info.vendor
        << "' device (" << info.deviceID << ") '" << info.device
        << "' adapter " << adapterTypeToString(info.adapterType)
@@ -110,7 +162,7 @@ template<typename WebGPUPrintable>
        << " architecture '" << info.architecture
        << "' subgroupMinSize " << info.subgroupMinSize
        << " subgroupMaxSize " << info.subgroupMaxSize;
-    return out.str();
+    return out;
 }

 [[nodiscard]] constexpr std::string_view deviceLostReasonToString(
@@ -123,7 +175,7 @@ template<typename WebGPUPrintable>
    }
 }

-[[nodiscard]] constexpr std::string_view filamentShaderStageToString(const ShaderStage stage) {
+[[nodiscard]] constexpr std::string_view filamentShaderStageToString(ShaderStage stage) {
    switch (stage) {
        case ShaderStage::VERTEX:   return "vertex";
        case ShaderStage::FRAGMENT: return "fragment";
@@ -24,25 +24,43 @@
 #include "backend/DriverEnums.h"

 #include <utils/Panic.h>
+#include <utils/ostream.h>

 #include <webgpu/webgpu_cpp.h>

 #include <algorithm>
 #include <cstdint>

-namespace filament::backend {
-
 namespace {

 #if FWGPU_ENABLED(FWGPU_PRINT_SYSTEM)
+utils::io::ostream& operator<<(utils::io::ostream& out, const wgpu::TextureFormat format) noexcept {
+    return filament::backend::streamInsertWebGPUPrintable(out, format);
+}
+
+utils::io::ostream& operator<<(utils::io::ostream& out,
+        const wgpu::TextureUsage textureUsage) noexcept {
+    return filament::backend::streamInsertWebGPUPrintable(out, textureUsage);
+}
+
+utils::io::ostream& operator<<(utils::io::ostream& out,
+        const wgpu::PresentMode presentMode) noexcept {
+    return filament::backend::streamInsertWebGPUPrintable(out, presentMode);
+}
+
+utils::io::ostream& operator<<(utils::io::ostream& out,
+        const wgpu::CompositeAlphaMode alphaMode) noexcept {
+    return filament::backend::streamInsertWebGPUPrintable(out, alphaMode);
+}
+
 void printSurfaceCapabilitiesDetails(wgpu::SurfaceCapabilities const& capabilities) {
    FWGPU_LOGI << "WebGPU surface capabilities:";
-    FWGPU_LOGI << "  surface usages: " << webGPUPrintableToString(capabilities.usages);
+    FWGPU_LOGI << "  surface usages: " << capabilities.usages;
    FWGPU_LOGI << "  surface formats (" << capabilities.formatCount << "):";
    if (capabilities.formatCount > 0 && capabilities.formats != nullptr) {
        std::for_each(capabilities.formats, capabilities.formats + capabilities.formatCount,
                [](wgpu::TextureFormat const format) {
-                    FWGPU_LOGI << "    " << webGPUPrintableToString(format);
+                    FWGPU_LOGI << "    " << format;
                });
    }
    FWGPU_LOGI << "  surface present modes (" << capabilities.presentModeCount << "):";
@@ -50,7 +68,7 @@ void printSurfaceCapabilitiesDetails(wgpu::SurfaceCapabilities const& capabiliti
        std::for_each(capabilities.presentModes,
                capabilities.presentModes + capabilities.presentModeCount,
                [](wgpu::PresentMode const presentMode) {
-                    FWGPU_LOGI << "    " << webGPUPrintableToString(presentMode);
+                    FWGPU_LOGI << "    " << presentMode;
                });
    }
    FWGPU_LOGI << "  surface alpha modes (" << capabilities.alphaModeCount << "):";
@@ -58,7 +76,7 @@ void printSurfaceCapabilitiesDetails(wgpu::SurfaceCapabilities const& capabiliti
        std::for_each(capabilities.alphaModes,
                capabilities.alphaModes + capabilities.alphaModeCount,
                [](wgpu::CompositeAlphaMode const alphaMode) {
-                    FWGPU_LOGI << "    " << webGPUPrintableToString(alphaMode);
+                    FWGPU_LOGI << "    " << alphaMode;
                });
    }
 }
@@ -66,20 +84,20 @@ void printSurfaceCapabilitiesDetails(wgpu::SurfaceCapabilities const& capabiliti
 void printSurfaceConfiguration(wgpu::SurfaceConfiguration const& config,
        wgpu::TextureFormat depthFormat) {
    FWGPU_LOGI << "WebGPU surface configuration:";
-    FWGPU_LOGI << "  surface format: " << webGPUPrintableToString(config.format);
-    FWGPU_LOGI << "  surface usage: " << webGPUPrintableToString(config.usage);
+    FWGPU_LOGI << "  surface format: " << config.format;
+    FWGPU_LOGI << "  surface usage: " << config.usage;
    FWGPU_LOGI << "  surface view formats (" << config.viewFormatCount << "):";
    if (config.viewFormatCount > 0 && config.viewFormats != nullptr) {
        std::for_each(config.viewFormats, config.viewFormats + config.viewFormatCount,
                [](wgpu::TextureFormat const viewFormat) {
-                    FWGPU_LOGI << "    " << webGPUPrintableToString(viewFormat);
+                    FWGPU_LOGI << "    " << viewFormat;
                });
    }
-    FWGPU_LOGI << "  surface alpha mode: " << webGPUPrintableToString(config.alphaMode);
+    FWGPU_LOGI << "  surface alpha mode: " << config.alphaMode;
    FWGPU_LOGI << "  surface width: " << config.width;
    FWGPU_LOGI << "  surface height: " << config.height;
-    FWGPU_LOGI << "  surface present mode: " << webGPUPrintableToString(config.presentMode);
-    FWGPU_LOGI << "WebGPU selected depth format: " << webGPUPrintableToString(depthFormat);
+    FWGPU_LOGI << "  surface present mode: " << config.presentMode;
+    FWGPU_LOGI << "WebGPU selected depth format: " << depthFormat;
 }
 #endif// FWGPU_ENABLED(FWGPU_PRINT_SYSTEM)

@@ -239,6 +257,8 @@ void initConfig(wgpu::SurfaceConfiguration& config, wgpu::Device const& device,

 }// namespace

+namespace filament::backend {
+
 WebGPUSwapChain::WebGPUSwapChain(wgpu::Surface&& surface, wgpu::Extent2D const& surfaceSize,
        wgpu::Adapter const& adapter, wgpu::Device const& device, uint64_t flags)
    : mDevice(device),
@@ -25,6 +25,7 @@
 #include <utils/BitmaskEnum.h>
 #include <utils/Panic.h>
 #include <utils/debug.h>
+#include <utils/ostream.h>

 #include <webgpu/webgpu_cpp.h>
 #if FWGPU_ENABLED(FWGPU_PRINT_SYSTEM)
@@ -342,7 +343,8 @@ WebGPUTexture::WebGPUTexture(const SamplerType samplerType, const uint8_t levels
        std::stringstream textureFormatStream;
        textureFormatStream << mWebGPUFormat;
        FWGPU_LOGD << "Texture '" << textureDescriptor.label << "' has view format "
-                   << viewFormatStream.str() << " and texture format " << textureFormatStream.str();
+                   << viewFormatStream.str() << " and texture format " << textureFormatStream.str()
+                   << utils::io::endl;
    }
 #endif
 }
@@ -468,43 +470,51 @@ wgpu::TextureFormat WebGPUTexture::fToWGPUTextureFormat(TextureFormat const& fFo
            // No direct mapping in wgpu. Could potentially map to RGBA8Unorm
            // and discard the alpha and lower precision.
            FWGPU_LOGW << "Requested Filament texture format RGB565 but getting "
-                          "wgpu::TextureFormat::Undefined (no direct mapping in wgpu)";
+                          "wgpu::TextureFormat::Undefined (no direct mapping in wgpu)"
+                       << utils::io::endl;
            return wgpu::TextureFormat::Undefined;
        case TextureFormat::RGB9_E5: return wgpu::TextureFormat::RGB9E5Ufloat;
        case TextureFormat::RGB5_A1:
            // No direct mapping in wgpu. Could potentially map to RGBA8Unorm
            // and handle the packing/unpacking in shaders.
            FWGPU_LOGW << "Requested Filament texture format RGB5_A1 but getting "
-                          "wgpu::TextureFormat::Undefined (no direct mapping in wgpu)";
+                          "wgpu::TextureFormat::Undefined (no direct mapping in wgpu)"
+                       << utils::io::endl;
            return wgpu::TextureFormat::Undefined;
        case TextureFormat::RGBA4:
            // No direct mapping in wgpu. Could potentially map to RGBA8Unorm
            // and handle the packing/unpacking in shaders.
            FWGPU_LOGW << "Requested Filament texture format RGBA4 but getting "
-                          "wgpu::TextureFormat::Undefined (no direct mapping in wgpu)";
+                          "wgpu::TextureFormat::Undefined (no direct mapping in wgpu)"
+                       << utils::io::endl;
            return wgpu::TextureFormat::Undefined;
        case TextureFormat::RGB8:
            FWGPU_LOGW << "Requested Filament texture format RGB8 but getting "
                          "wgpu::TextureFormat::RGBA8Unorm (no direct sRGB equivalent in wgpu "
-                          "without alpha)";
+                          "without alpha)"
+                       << utils::io::endl;
            return wgpu::TextureFormat::RGBA8Unorm;
        case TextureFormat::SRGB8:
            FWGPU_LOGW << "Requested Filament texture format SRGB8 but getting "
                          "wgpu::TextureFormat::RGBA8UnormSrgb (no direct sRGB equivalent in wgpu "
-                          "without alpha)";
+                          "without alpha)"
+                       << utils::io::endl;
            return wgpu::TextureFormat::RGBA8UnormSrgb;
        case TextureFormat::RGB8_SNORM:
            FWGPU_LOGW
                    << "Requested Filament texture format RGB8_SNORM but getting "
-                       "wgpu::TextureFormat::RGBA8Snorm (no direct mapping in wgpu without alpha)";
+                       "wgpu::TextureFormat::RGBA8Snorm (no direct mapping in wgpu without alpha)"
+                    << utils::io::endl;
            return wgpu::TextureFormat::RGBA8Snorm;
        case TextureFormat::RGB8UI:
            FWGPU_LOGW << "Requested Filament texture format RGB8UI but getting "
-                          "wgpu::TextureFormat::RGBA8Uint (no direct mapping in wgpu without alpha)";
+                          "wgpu::TextureFormat::RGBA8Uint (no direct mapping in wgpu without alpha)"
+                       << utils::io::endl;
            return wgpu::TextureFormat::RGBA8Uint;
        case TextureFormat::RGB8I:
            FWGPU_LOGW << "Requested Filament texture format RGB8I but getting "
-                          "wgpu::TextureFormat::RGBA8Sint (no direct mapping in wgpu without alpha)";
+                          "wgpu::TextureFormat::RGBA8Sint (no direct mapping in wgpu without alpha)"
+                       << utils::io::endl;
            return wgpu::TextureFormat::RGBA8Sint;
        case TextureFormat::R11F_G11F_B10F:          return wgpu::TextureFormat::RG11B10Ufloat;
        case TextureFormat::UNUSED:                  return wgpu::TextureFormat::Undefined;
@@ -512,32 +522,38 @@ wgpu::TextureFormat WebGPUTexture::fToWGPUTextureFormat(TextureFormat const& fFo
        case TextureFormat::RGB16F:
            FWGPU_LOGW
                    << "Requested Filament texture format RGB16F but getting "
-                       "wgpu::TextureFormat::RGBA16Float (no direct mapping in wgpu without alpha)";
+                       "wgpu::TextureFormat::RGBA16Float (no direct mapping in wgpu without alpha)"
+                    << utils::io::endl;
            return wgpu::TextureFormat::RGBA16Float;
        case TextureFormat::RGB16UI:
            FWGPU_LOGW
                    << "Requested Filament texture format RGB16UI but getting "
-                       "wgpu::TextureFormat::RGBA16Uint (no direct mapping in wgpu without alpha)";
+                       "wgpu::TextureFormat::RGBA16Uint (no direct mapping in wgpu without alpha)"
+                    << utils::io::endl;
            return wgpu::TextureFormat::RGBA16Uint;
        case TextureFormat::RGB16I:
            FWGPU_LOGW
                    << "Requested Filament texture format RGB16I but getting "
-                       "wgpu::TextureFormat::RGBA16Sint (no direct mapping in wgpu without alpha)";
+                       "wgpu::TextureFormat::RGBA16Sint (no direct mapping in wgpu without alpha)"
+                    << utils::io::endl;
            return wgpu::TextureFormat::RGBA16Sint;
        case TextureFormat::RGB32F:
            FWGPU_LOGW
                    << "Requested Filament texture format RGB32F but getting "
-                       "wgpu::TextureFormat::RGBA32Float (no direct mapping in wgpu without alpha)";
+                       "wgpu::TextureFormat::RGBA32Float (no direct mapping in wgpu without alpha)"
+                    << utils::io::endl;
            return wgpu::TextureFormat::RGBA32Float;
        case TextureFormat::RGB32UI:
            FWGPU_LOGW
                    << "Requested Filament texture format RGB32UI but getting "
-                       "wgpu::TextureFormat::RGBA32Uint (no direct mapping in wgpu without alpha)";
+                       "wgpu::TextureFormat::RGBA32Uint (no direct mapping in wgpu without alpha)"
+                    << utils::io::endl;
            return wgpu::TextureFormat::RGBA32Uint;
        case TextureFormat::RGB32I:
            FWGPU_LOGW
                    << "Requested Filament texture format RGB32I but getting "
-                       "wgpu::TextureFormat::RGBA32Sint (no direct mapping in wgpu without alpha)";
+                       "wgpu::TextureFormat::RGBA32Sint (no direct mapping in wgpu without alpha)"
+                    << utils::io::endl;
            return wgpu::TextureFormat::RGBA32Sint;
        case TextureFormat::DXT1_RGB:                return wgpu::TextureFormat::BC1RGBAUnorm;
        case TextureFormat::DXT1_RGBA:               return wgpu::TextureFormat::BC1RGBAUnorm;
@@ -557,7 +573,7 @@ wgpu::TextureView WebGPUTexture::makeTextureView(const uint8_t& baseLevel,
    if (baseLevel > 0 && !mSupportsMultipleMipLevels) {
        FWGPU_LOGW << "Trying to make a texture view into a level ("
                   << static_cast<uint32_t>(baseLevel)
-                   << ") for which we cannot generate mip levels.";
+                   << ") for which we cannot generate mip levels." << utils::io::endl;
    }
 #endif
    const wgpu::TextureViewDescriptor textureViewDescriptor{
@@ -22,6 +22,7 @@
 #include <backend/DriverEnums.h>

 #include <utils/Panic.h>
+#include <utils/ostream.h>

 #include <webgpu/webgpu_cpp.h>

@@ -56,22 +57,26 @@ namespace {
            case ElementType::BYTE3:
                FWGPU_LOGW
                        << "Requested Filament vertex format BYTE3 (normalized) but getting "
-                           "wgpu::VertexFormat::Snorm8x4 (no direct mapping in wgpu for x3 byte)";
+                           "wgpu::VertexFormat::Snorm8x4 (no direct mapping in wgpu for x3 byte)"
+                        << utils::io::endl;
                return VertexFormat::Snorm8x4; // NOT MINSPEC
            case ElementType::UBYTE3:
                FWGPU_LOGW
                        << "Requested Filament vertex format UBYTE3 (normalized) but getting "
-                           "wgpu::VertexFormat::Unorm8x4 (no direct mapping in wgpu for x3 byte)";
+                           "wgpu::VertexFormat::Unorm8x4 (no direct mapping in wgpu for x3 byte)"
+                        << utils::io::endl;
                return VertexFormat::Unorm8x4; // NOT MINSPEC
            case ElementType::SHORT3:
                FWGPU_LOGW << "Requested Filament vertex format SHORT3 (normalized) but getting "
                              "wgpu::VertexFormat::Snorm16x4 (no direct mapping in wgpu for x3 "
-                              "half/short)";
+                              "half/short)"
+                           << utils::io::endl;
                return VertexFormat::Snorm16x4; // NOT MINSPEC
            case ElementType::USHORT3:
                FWGPU_LOGW << "Requested Filament vertex format USHORT3 (normalized) but getting "
                              "wgpu::VertexFormat::Unorm16x4 (no direct mapping in wgpu for x3 "
-                              "half/short)";
+                              "half/short)"
+                           << utils::io::endl;
                return VertexFormat::Unorm16x4; // NOT MINSPEC
            // Four Component Types
            case ElementType::BYTE4:   return VertexFormat::Snorm8x4;
@@ -92,22 +97,26 @@ namespace {
        case ElementType::BYTE:
            if (integer) return VertexFormat::Sint8;
            FWGPU_LOGW << "Requested Filament vertex format BYTE (float) but getting "
-                          "wgpu::VertexFormat::Float16 (no direct mapping in wgpu for 8 bit float)";
+                          "wgpu::VertexFormat::Float16 (no direct mapping in wgpu for 8 bit float)"
+                       << utils::io::endl;
            return wgpu::VertexFormat::Float16;
        case ElementType::UBYTE:
            if (integer) return VertexFormat::Uint8;
            FWGPU_LOGW << "Requested Filament vertex format UBYTE (float) but getting "
-                          "wgpu::VertexFormat::Float16 (no direct mapping in wgpu for 8 bit float)";
+                          "wgpu::VertexFormat::Float16 (no direct mapping in wgpu for 8 bit float)"
+                       << utils::io::endl;
            return VertexFormat::Float16;
        case ElementType::SHORT:
            if (integer) return VertexFormat::Sint16;
            FWGPU_LOGW << "Requested Filament vertex format SHORT (float) and getting "
-                          "wgpu::VertexFormat::Float16 (potential loss of precision?)";
+                          "wgpu::VertexFormat::Float16 (potential loss of precision?)"
+                       << utils::io::endl;
            return VertexFormat ::Float16;
        case ElementType::USHORT:
            if (integer) return VertexFormat::Uint16;
            FWGPU_LOGW << "Requested Filament vertex format USHORT (float) and getting "
-                          "wgpu::VertexFormat::Float16 (potential loss of precision?)";
+                          "wgpu::VertexFormat::Float16 (potential loss of precision?)"
+                       << utils::io::endl;
            return VertexFormat::Float16;
        case ElementType::HALF:    return VertexFormat::Float16;
        case ElementType::INT:     return VertexFormat::Sint32;
@@ -118,49 +127,58 @@ namespace {
            if (integer) return VertexFormat::Sint8x2;
            FWGPU_LOGW
                    << "Requested Filament vertex format BYTE2 (float) but getting "
-                       "wgpu::VertexFormat::Float16x2 (no direct mapping in wgpu for 8 bit float)";
+                       "wgpu::VertexFormat::Float16x2 (no direct mapping in wgpu for 8 bit float)"
+                    << utils::io::endl;
            return VertexFormat::Float16x2;
        case ElementType::UBYTE2:
            if (integer) return VertexFormat::Uint8x2;
            FWGPU_LOGW
                    << "Requested Filament vertex format UBYTE2 (float) but getting "
-                       "wgpu::VertexFormat::Float16x2 (no direct mapping in wgpu for 8 bit float)";
+                       "wgpu::VertexFormat::Float16x2 (no direct mapping in wgpu for 8 bit float)"
+                    << utils::io::endl;
            return VertexFormat::Float16x2;
        case ElementType::SHORT2:
            if (integer) return VertexFormat::Sint16x2;
            FWGPU_LOGW << "Requested Filament vertex format SHORT2 (float) but getting "
-                          "wgpu::VertexFormat::Float16x2 (potential loss of precision?)";
+                          "wgpu::VertexFormat::Float16x2 (potential loss of precision?)"
+                       << utils::io::endl;
            return VertexFormat::Float16x2;
        case ElementType::USHORT2:
            if (integer) return VertexFormat::Uint16x2;
            FWGPU_LOGW << "Requested Filament vertex format USHORT2 (float) but getting "
-                          "wgpu::VertexFormat::Float16x2 (potential loss of precision?)";
+                          "wgpu::VertexFormat::Float16x2 (potential loss of precision?)"
+                       << utils::io::endl;
            return VertexFormat::Float16x2;
        case ElementType::HALF2:   return VertexFormat::Float16x2;
        case ElementType::FLOAT2:  return VertexFormat::Float32x2;
        // Three Component Types
        case ElementType::BYTE3:
            FWGPU_LOGW << "Requested Filament vertex format BYTE3 but getting "
-                          "wgpu::VertexFormat::Sint8x4 (no direct mapping in wgpu for x3 byte)";
+                          "wgpu::VertexFormat::Sint8x4 (no direct mapping in wgpu for x3 byte)"
+                       << utils::io::endl;
            return VertexFormat::Sint8x4; // NOT MINSPEC
        case ElementType::UBYTE3:
            FWGPU_LOGW << "Requested Filament vertex format UBYTE3 but getting "
-                          "wgpu::VertexFormat::Uint8x4 (no direct mapping in wgpu for x3 byte)";
+                          "wgpu::VertexFormat::Uint8x4 (no direct mapping in wgpu for x3 byte)"
+                       << utils::io::endl;
            return VertexFormat::Uint8x4; // NOT MINSPEC
        case ElementType::SHORT3:
            FWGPU_LOGW
                    << "Requested Filament vertex format SHORT3 but getting "
-                       "wgpu::VertexFormat::Sint16x4 (no direct mapping in wgpu for x3 half/short)";
+                       "wgpu::VertexFormat::Sint16x4 (no direct mapping in wgpu for x3 half/short)"
+                    << utils::io::endl;
            return VertexFormat::Sint16x4; // NOT MINSPEC
        case ElementType::USHORT3:
            FWGPU_LOGW
                    << "Requested Filament vertex format USHORT3 but getting "
-                       "wgpu::VertexFormat::Uint16x4 (no direct mapping in wgpu for x3 half/short)";
+                       "wgpu::VertexFormat::Uint16x4 (no direct mapping in wgpu for x3 half/short)"
+                    << utils::io::endl;
            return VertexFormat::Uint16x4; // NOT MINSPEC
        case ElementType::HALF3:
            FWGPU_LOGW
                    << "Requested Filament vertex format HALF3 but getting "
-                       "wgpu::VertexFormat::Float16x4 (no direct mapping in wgpu for x3 half/short)";
+                       "wgpu::VertexFormat::Float16x4 (no direct mapping in wgpu for x3 half/short)"
+                    << utils::io::endl;
            return VertexFormat::Float16x4; // NOT MINSPEC
        case ElementType::FLOAT3:  return VertexFormat::Float32x3;
        // Four Component Types
@@ -168,23 +186,27 @@ namespace {
            if (integer) return VertexFormat::Sint8x4;
            FWGPU_LOGW
                    << "Requested Filament vertex format BYTE4 (float) but getting "
-                       "wgpu::VertexFormat::Float16x4 (no direct mapping in wgpu for 8 bit float)";
+                       "wgpu::VertexFormat::Float16x4 (no direct mapping in wgpu for 8 bit float)"
+                    << utils::io::endl;
            return VertexFormat::Float16x4;
        case ElementType::UBYTE4:
            if (integer) return VertexFormat::Uint8x4;
            FWGPU_LOGW
                    << "Requested Filament vertex format UBYTE4 (float) but getting "
-                       "wgpu::VertexFormat::Float16x4 (no direct mapping in wgpu for 8 bit float)";
+                       "wgpu::VertexFormat::Float16x4 (no direct mapping in wgpu for 8 bit float)"
+                    << utils::io::endl;
            return VertexFormat::Float16x4;
        case ElementType::SHORT4:
            if (integer) return VertexFormat::Sint16x4;
            FWGPU_LOGW << "Requested Filament vertex format SHORT4 (float) but getting "
-                          "wgpu::VertexFormat::Float16x4 (potential loss of precision?)";
+                          "wgpu::VertexFormat::Float16x4 (potential loss of precision?)"
+                       << utils::io::endl;
            return VertexFormat::Float16x4;
        case ElementType::USHORT4:
            if (integer) return VertexFormat::Uint16x4;
            FWGPU_LOGW << "Requested Filament vertex format USHORT4 (float) but getting "
-                          "wgpu::VertexFormat::Float16x4 (potential loss of precision?)";
+                          "wgpu::VertexFormat::Float16x4 (potential loss of precision?)"
+                       << utils::io::endl;
            return VertexFormat::Float16x4;
        case ElementType::HALF4:   return VertexFormat::Float16x4;
        case ElementType::FLOAT4:  return VertexFormat::Float32x4;
@@ -24,6 +24,7 @@
 #include <backend/Platform.h>
 #include <utils/Hash.h>
 #include <utils/Panic.h>
+#include <utils/ostream.h>

 #include <dawn/webgpu_cpp_print.h>
 #include <webgpu/webgpu_cpp.h>
@@ -38,7 +39,6 @@
 #include <utility>
 #include <vector>
 #if FWGPU_ENABLED(FWGPU_PRINT_SYSTEM)
-#include <string>
 #include <string_view>
 #include <variant>
 #endif
@@ -53,9 +53,6 @@ namespace filament::backend {

 namespace {

-// The SPD Algorithm can make use of up to 12 storage texture attachments
-constexpr uint32_t MAX_MIPMAP_STORAGE_TEXTURES_PER_STAGE = 12u;
-
 constexpr std::array REQUIRED_FEATURES = { wgpu::FeatureName::TransientAttachments,
    /*To make filtering assumptions like we want while waiting for Filament to provide that info,
       float 32 needs to be filterable*/
@@ -120,6 +117,28 @@ constexpr void forEachLimitToValidate(std::function<bool(LimitToValidate const)>
    }
 }

+#if FWGPU_ENABLED(FWGPU_PRINT_SYSTEM)
+utils::io::ostream& operator<<(utils::io::ostream& out,
+        const wgpu::WGSLLanguageFeatureName languageFeatureName) noexcept {
+    return streamInsertWebGPUPrintable(out, languageFeatureName);
+}
+
+utils::io::ostream& operator<<(utils::io::ostream& out,
+        const wgpu::FeatureName featureName) noexcept {
+    return streamInsertWebGPUPrintable(out, featureName);
+}
+#endif
+
+template<typename STREAM_TYPE>
+STREAM_TYPE& operator<<(STREAM_TYPE& out, wgpu::RequestAdapterOptions const& options) noexcept {
+    return streamInsertRequestAdapterOptions(out, options);
+}
+
+template<class STREAM_TYPE>
+STREAM_TYPE& operator<<(STREAM_TYPE& out, wgpu::AdapterInfo const& info) noexcept {
+    return streamInsertRequestAdapterInfo(out, info);
+}
+
 constexpr bool isDefined(uint32_t limit) { return limit < wgpu::kLimitU32Undefined; }

 [[maybe_unused]] constexpr bool isDefined(uint64_t limit) {
@@ -240,7 +259,7 @@ void printInstanceDetails(wgpu::Instance const& instance) {
                    supportedWGSLLanguageFeatures.features +
                            supportedWGSLLanguageFeatures.featureCount,
                    [](wgpu::WGSLLanguageFeatureName const featureName) {
-                        FWGPU_LOGI << "  " << webGPUPrintableToString(featureName);
+                        FWGPU_LOGI << "  " << featureName;
                    });
        }
    }
@@ -280,22 +299,23 @@ void printInstanceDetails(wgpu::Instance const& instance) {

 #if FWGPU_ENABLED(FWGPU_PRINT_SYSTEM)
 void printLimit(std::string_view name, const std::variant<uint32_t, uint64_t> value) {
-    static constexpr std::string_view indent = "  ";
+    FWGPU_LOGI << "  " << name.data() << ": ";
    bool undefined = true;
    if (std::holds_alternative<uint32_t>(value)) {
        if (std::get<uint32_t>(value) != WGPU_LIMIT_U32_UNDEFINED) {
            undefined = false;
-            FWGPU_LOGI << indent << name.data() << ": " << std::get<uint32_t>(value);
+            FWGPU_LOGI << std::get<uint32_t>(value);
        }
    } else if (std::holds_alternative<uint64_t>(value)) {
        if (std::get<uint64_t>(value) != WGPU_LIMIT_U64_UNDEFINED) {
            undefined = false;
-            FWGPU_LOGI << indent << name.data() << ": " << std::get<uint64_t>(value);
+            FWGPU_LOGI << std::get<uint64_t>(value);
        }
    }
    if (undefined) {
-        FWGPU_LOGI << indent << name.data() << ": UNDEFINED";
+        FWGPU_LOGI << "UNDEFINED";
    }
+    FWGPU_LOGI << "";
 }
 #endif// FWGPU_ENABLED(FWGPU_PRINT_SYSTEM)

@@ -373,16 +393,16 @@ struct AdapterDetails final {
    }
 };

-[[nodiscard]] std::string toString(AdapterDetails const& details) {
-    std::stringstream out;
-    out << adapterInfoToString(details.info)
+template<class STREAM_TYPE>
+STREAM_TYPE& operator<<(STREAM_TYPE& out, AdapterDetails const& details) noexcept {
+    out << details.info
        << " power preference " << powerPreferenceToString(details.powerPreference);
-    return out.str();
+    return out;
 }

 #if FWGPU_ENABLED(FWGPU_PRINT_SYSTEM)
 void printAdapterDetails(AdapterDetails const& details) {
-    FWGPU_LOGI << "Selected WebGPU adapter info: " << toString(details);
+    FWGPU_LOGI << "Selected WebGPU adapter info: " << details;
    wgpu::SupportedFeatures supportedFeatures{};
    details.adapter.GetFeatures(&supportedFeatures);
    FWGPU_LOGI << "WebGPU adapter supported features (" << supportedFeatures.featureCount
@@ -391,7 +411,7 @@ void printAdapterDetails(AdapterDetails const& details) {
        std::for_each(supportedFeatures.features,
                supportedFeatures.features + supportedFeatures.featureCount,
                [](wgpu::FeatureName const featureName) {
-                    FWGPU_LOGI << "  " << webGPUPrintableToString(featureName);
+                    FWGPU_LOGI << "  " << featureName;
                });
    }
    wgpu::Limits supportedLimits{};
@@ -424,12 +444,12 @@ struct AdapterDetailsHash final {
                    return details.adapter.HasFeature(featureName);
                })) {
 #if FWGPU_ENABLED(FWGPU_PRINT_SYSTEM)
-        FWGPU_LOGI << "WebGPU adapter " << toString(details)
+        FWGPU_LOGI << "WebGPU adapter " << details
                   << " does not have the minimum required features.";
        FWGPU_LOGI << "Missing required feature(s): ";
        for (wgpu::FeatureName const& requiredFeature: REQUIRED_FEATURES) {
            if (!details.adapter.HasFeature(requiredFeature)) {
-                FWGPU_LOGI << webGPUPrintableToString(requiredFeature);
+                FWGPU_LOGI << requiredFeature;
            }
        }
        FWGPU_LOGI << "";
@@ -438,10 +458,10 @@ struct AdapterDetailsHash final {
    }
    wgpu::Limits supportedLimits {};
    FILAMENT_CHECK_POSTCONDITION(details.adapter.GetLimits(&supportedLimits))
-            << "Failed to get limits for WebGPU adapter: " << toString(details);
+            << "Failed to get limits for WebGPU adapter: " << details;
    if (!satisfiesLimits(supportedLimits)) {
 #if FWGPU_ENABLED(FWGPU_PRINT_SYSTEM)
-        FWGPU_LOGI << " (for WebGPU adapter " << toString(details) << ")";
+        FWGPU_LOGI << " (for WebGPU adapter " << details << ")";
 #endif
        return false;
    }
@@ -467,15 +487,14 @@ struct AdapterDetailsHash final {
                            status != wgpu::RequestAdapterStatus::CallbackCancelled)
                            << "Failed to request a WebGPU adapter due to the request callback "
                               "being cancelled? Options: "
-                            << adapterOptionsToString(options) << " " << message.data;
+                            << options << " " << message.data;
                    FILAMENT_CHECK_POSTCONDITION(status != wgpu::RequestAdapterStatus::Error)
                            << "Failed to request a WebGPU adapter due to an error. Options: "
-                            << adapterOptionsToString(options) << " Error: " << message.data;
+                            << options << " Error: " << message.data;
                    if (status == wgpu::RequestAdapterStatus::Success) {
                        AdapterDetails details = AdapterDetails(readyAdapter);
                        FILAMENT_CHECK_POSTCONDITION(readyAdapter.GetInfo(&details.info))
-                                << "Failed to get info for adapter (options: "
-                                << adapterOptionsToString(options) << ")";
+                                << "Failed to get info for adapter (options: " << options << ")";
                        const std::lock_guard<std::mutex> lock(adaptersMutex);
                        compatibleAdapters.emplace(std::move(details.info), details.powerPreference,
                                std::move(details.adapter));
@@ -490,11 +509,9 @@ struct AdapterDetailsHash final {
        wgpu::Future& future = futures[i];
        wgpu::WaitStatus status = instance.WaitAny(future, REQUEST_ADAPTER_TIMEOUT_NANOSECONDS);
        FILAMENT_CHECK_POSTCONDITION(status != wgpu::WaitStatus::TimedOut)
-                << "Timed out requesting a WebGPU adapter with options "
-                << adapterOptionsToString(options);
+                << "Timed out requesting a WebGPU adapter with options " << options;
        FILAMENT_CHECK_POSTCONDITION(status != wgpu::WaitStatus::Error)
-                << "Failed to request a WebGPU adapter with options "
-                << adapterOptionsToString(options)
+                << "Failed to request a WebGPU adapter with options " << options
                << " due to an error (as request was made synchronous)";
        assert_invariant(status == wgpu::WaitStatus::Success);
    }
@@ -502,7 +519,7 @@ struct AdapterDetailsHash final {
 #if FWGPU_ENABLED(FWGPU_PRINT_SYSTEM)
    FWGPU_LOGI << compatibleAdapters.size() << " WebGPU adapter(s) found:";
    for (auto& details: compatibleAdapters) {
-        FWGPU_LOGI << "  WebGPU adapter: " << toString(details);
+        FWGPU_LOGI << "  WebGPU adapter: " << details;
    }
 #endif
    return compatibleAdapters;
@@ -569,7 +586,7 @@ void printDeviceDetails(wgpu::Device const& device) {
        std::for_each(supportedFeatures.features,
                supportedFeatures.features + supportedFeatures.featureCount,
                [](wgpu::FeatureName const featureName) {
-                    FWGPU_LOGI << "  " << webGPUPrintableToString(featureName);
+                    FWGPU_LOGI << "  " << featureName;
                });
    }
    wgpu::Limits supportedLimits{};
@@ -611,21 +628,7 @@ wgpu::Device WebGPUPlatform::requestDevice(wgpu::Adapter const& adapter) {
    deviceDescriptor.defaultQueue.label = "default_queue";
    deviceDescriptor.requiredFeatureCount = enabledFeatures.size();
    deviceDescriptor.requiredFeatures = enabledFeatures.data();
-
-    // It is helpful to increase maxStorageTexturesPerShaderStage if the hardware supports it, but
-    // not required. This allows our mipmap generation to perform as well as possible for the given
-    // hardware. If we have more cases of "Nice to increase" limits like this we can generalize it
-    // in the future.
-    wgpu::Limits supportedLimits{};
-    FILAMENT_CHECK_POSTCONDITION(
-            adapter.GetLimits(&supportedLimits).status == wgpu::Status::Success)
-            << "Failed to get limits for WebGPU adapter";
-    auto limitsToRequest = REQUIRED_LIMITS;
-    limitsToRequest.maxStorageTexturesPerShaderStage =
-            std::min(MAX_MIPMAP_STORAGE_TEXTURES_PER_STAGE,
-                    supportedLimits.maxStorageTexturesPerShaderStage);
-    deviceDescriptor.requiredLimits = &limitsToRequest;
-
+    deviceDescriptor.requiredLimits = &REQUIRED_LIMITS;
    deviceDescriptor.SetDeviceLostCallback(wgpu::CallbackMode::AllowSpontaneous,
            [](wgpu::Device const&, wgpu::DeviceLostReason const& reason,
                    wgpu::StringView message) {
@@ -681,6 +684,7 @@ wgpu::Device WebGPUPlatform::requestDevice(wgpu::Adapter const& adapter) {
                            "adapter should support them: %s\n",
                missingFeatures, featureNamesStream.str().data());
    }
+    wgpu::Limits supportedLimits {};
    FILAMENT_CHECK_POSTCONDITION(device.GetLimits(&supportedLimits))
            << "Failed to get limits for the device?";
    FILAMENT_CHECK_POSTCONDITION(satisfiesLimits(supportedLimits))
@@ -317,15 +317,6 @@ public:
         */
        size_t metalUploadBufferSizeBytes = 512 * 1024;

-        /**
-         * The action to take if a Drawable cannot be acquired.
-         *
-         * Each frame rendered requires a CAMetalDrawable texture, which is
-         * presented on-screen at the completion of each frame. These are
-         * limited and provided round-robin style by the system.
-         */
-        bool metalDisablePanicOnDrawableFailure = false;
-
        /**
         * Set to `true` to forcibly disable parallel shader compilation in the backend.
         * Currently only honored by the GL and Metal backends.
@@ -86,6 +86,9 @@ public:
            std::is_same_v<math::mat3f, T>
    >;

+    template<typename T>
+    using is_supported_constant_parameter_t = std::enable_if_t<std::is_same_v<bool, T>>;
+
    /**
     * Creates a new MaterialInstance using another MaterialInstance as a template for initialization.
     * The new MaterialInstance is an instance of the same Material of the template instance and
@@ -237,13 +240,13 @@ public:

    /**
     * Gets the value of a parameter by name.
-     * 
+     *
     * Note: Only supports non-texture parameters such as numeric and math types.
-     * 
+     *
     * @param name          Name of the parameter as defined by Material. Cannot be nullptr.
     * @param nameLength    Length in `char` of the name parameter.
     * @throws utils::PreConditionPanic if name doesn't exist or no-op if exceptions are disabled.
-     * 
+     *
     * @see Material::hasParameter
     */
    template<typename T>
@@ -261,6 +264,53 @@ public:
        return getParameter<T>(name, strlen(name));
    }

+    /**
+     * Set a mutable constant parameter by name.
+     *
+     * @param name          Name of the parameter as defined by Material. Cannot be nullptr.
+     * @param nameLength    Length in `char` of the name parameter.
+     * @param value         Value of the parameter to set.
+     * @throws utils::PreConditionPanic if name doesn't exist or no-op if exceptions are disabled.
+     */
+    template<typename T, typename = is_supported_constant_parameter_t<T>>
+    void setConstant(const char* UTILS_NONNULL name, size_t nameLength, T const& value);
+
+    /** inline helper to provide the name as a null-terminated string literal */
+    template<typename T, typename = is_supported_constant_parameter_t<T>>
+    void setConstant(StringLiteral const name, T const& value) {
+        setConstant<T>(name.data, name.size, value);
+    }
+
+    /** inline helper to provide the name as a null-terminated C string */
+    template<typename T, typename = is_supported_constant_parameter_t<T>>
+    void setConstant(const char* UTILS_NONNULL name, T const& value) {
+        setConstant<T>(name, strlen(name), value);
+    }
+
+    /**
+     * Gets the value of a mutable constant parameter by name.
+     *
+     * @param name          Name of the parameter as defined by Material. Cannot be nullptr.
+     * @param nameLength    Length in `char` of the name parameter.
+     * @throws utils::PreConditionPanic if name doesn't exist or no-op if exceptions are disabled.
+     *
+     * @see Material::hasConstant
+     */
+    template<typename T, typename = is_supported_constant_parameter_t<T>>
+    T getConstant(const char* UTILS_NONNULL name, size_t nameLength) const;
+
+    /** inline helper to provide the name as a null-terminated C string */
+    template<typename T, typename = is_supported_constant_parameter_t<T>>
+    T getConstant(StringLiteral const name) const {
+        return getConstant<T>(name.data, name.size);
+    }
+
+    /** inline helper to provide the name as a null-terminated C string */
+    template<typename T, typename = is_supported_constant_parameter_t<T>>
+    T getConstant(const char* UTILS_NONNULL name) const {
+        return getConstant<T>(name, strlen(name));
+    }
+
    /**
     * Set-up a custom scissor rectangle; by default it is disabled.
     *
@@ -232,6 +232,46 @@ template UTILS_PUBLIC mat3f MaterialInstance::getParameter<mat3f>      (const ch

 // ------------------------------------------------------------------------------------------------

+template<>
+inline void FMaterialInstance::setConstantImpl<bool>(std::string_view const name, bool const& value) {
+    std::optional<uint32_t> id = mMaterial->getMutableConstantId(name);
+    FILAMENT_CHECK_PRECONDITION(id.has_value()) << "No mutable constant with name " << name;
+    mConstants.set(*id, value);
+}
+
+template<typename T, typename>
+void MaterialInstance::setConstant(const char* name, size_t nameLength, T const& value) {
+    downcast(this)->setConstantImpl({ name, nameLength }, value);
+}
+
+// Explicit template instantiation of our supported types.
+//
+// Mutable spec constants will probably only ever allow bools, but it's nice to keep the API
+// forwards-compatible.
+template UTILS_PUBLIC void MaterialInstance::setConstant<bool>(const char* name, size_t nameLength, bool const& v);
+
+// ------------------------------------------------------------------------------------------------
+
+template<>
+inline bool FMaterialInstance::getConstantImpl<bool>(std::string_view const name) const {
+    std::optional<uint32_t> id = mMaterial->getMutableConstantId(name);
+    FILAMENT_CHECK_PRECONDITION(id.has_value()) << "No mutable constant with name " << name;
+    return mConstants[*id];
+}
+
+template<typename T, typename>
+T MaterialInstance::getConstant(const char* name, size_t nameLength) const {
+    return downcast(this)->getConstantImpl<T>({ name, nameLength });
+}
+
+// Explicit template instantiation of our supported types.
+//
+// Mutable spec constants will probably only ever allow bools, but it's nice to keep the API
+// forwards-compatible.
+template UTILS_PUBLIC bool MaterialInstance::getConstant<bool>(const char* name, size_t nameLength) const;
+
+// ------------------------------------------------------------------------------------------------
+
 Material const* MaterialInstance::getMaterial() const noexcept {
    return downcast(this)->getMaterial();
 }
@@ -229,6 +229,11 @@ bool MaterialParser::getConstants(FixedCapacityVector<MaterialConstant>* contain
    return get<ChunkMaterialConstants>(container);
 }

+bool MaterialParser::getMutableConstants(
+        FixedCapacityVector<MaterialMutableConstant>* container) const noexcept {
+    return get<ChunkMaterialMutableConstants>(container);
+}
+
 bool MaterialParser::getPushConstants(CString* structVarName,
        FixedCapacityVector<MaterialPushConstant>* value) const noexcept {
    auto [start, end] = mImpl.mChunkContainer.getChunkRange(MaterialPushConstants);
@@ -481,7 +486,6 @@ bool ChunkSamplerInterfaceBlock::unflatten(Unflattener& unflattener,
        uint8_t fieldType = 0;
        uint8_t fieldFormat = 0;
        uint8_t fieldPrecision = 0;
-        bool fieldUnfilterable = false;
        bool fieldMultisample = false;

        if (!unflattener.read(&fieldName)) {
@@ -504,10 +508,6 @@ bool ChunkSamplerInterfaceBlock::unflatten(Unflattener& unflattener,
            return false;
        }

-        if (!unflattener.read(&fieldUnfilterable)) {
-            return false;
-        }
-
        if (!unflattener.read(&fieldMultisample)) {
            return false;
        }
@@ -517,7 +517,6 @@ bool ChunkSamplerInterfaceBlock::unflatten(Unflattener& unflattener,
                SamplerInterfaceBlock::Type(fieldType),
                SamplerInterfaceBlock::Format(fieldFormat),
                SamplerInterfaceBlock::Precision(fieldPrecision),
-                fieldUnfilterable,
                fieldMultisample);
    }

@@ -757,6 +756,38 @@ bool ChunkMaterialConstants::unflatten(Unflattener& unflattener,
    return true;
 }

+bool ChunkMaterialMutableConstants::unflatten(Unflattener& unflattener,
+        FixedCapacityVector<MaterialMutableConstant>* materialConstants) {
+    assert_invariant(materialConstants);
+
+    // Read number of constants.
+    uint64_t numConstants = 0;
+    if (!unflattener.read(&numConstants)) {
+        return false;
+    }
+
+    materialConstants->reserve(numConstants);
+    materialConstants->resize(numConstants);
+
+    for (uint64_t i = 0; i < numConstants; i++) {
+        CString constantName;
+        bool defaultValue;
+
+        if (!unflattener.read(&constantName)) {
+            return false;
+        }
+
+        if (!unflattener.read(&defaultValue)) {
+            return false;
+        }
+
+        (*materialConstants)[i].name = std::move(constantName);
+        (*materialConstants)[i].defaultValue = defaultValue;
+    }
+
+    return true;
+}
+
 bool ChunkMaterialPushConstants::unflatten(Unflattener& unflattener,
        CString* structVarName,
        FixedCapacityVector<MaterialPushConstant>* materialPushConstants) {
@@ -49,6 +49,7 @@ class BufferInterfaceBlock;
 class SamplerInterfaceBlock;
 struct SubpassInfo;
 struct MaterialConstant;
+struct MaterialMutableConstant;
 struct MaterialPushConstant;

 class MaterialParser {
@@ -79,6 +80,8 @@ public:
    bool getShaderModels(uint32_t* value) const noexcept;
    bool getMaterialProperties(uint64_t* value) const noexcept;
    bool getConstants(utils::FixedCapacityVector<MaterialConstant>* value) const noexcept;
+    bool getMutableConstants(utils::FixedCapacityVector<MaterialMutableConstant>* value)
+        const noexcept;
    bool getPushConstants(utils::CString* structVarName,
            utils::FixedCapacityVector<MaterialPushConstant>* value) const noexcept;

@@ -237,6 +240,13 @@ struct ChunkMaterialConstants {
    static filamat::ChunkType const tag = filamat::MaterialConstants;
 };

+struct ChunkMaterialMutableConstants {
+    static bool unflatten(filaflat::Unflattener& unflattener,
+            utils::FixedCapacityVector<MaterialMutableConstant>* materialConstants);
+    using Container = utils::FixedCapacityVector<MaterialMutableConstant>;
+    static filamat::ChunkType const tag = filamat::MaterialMutableConstants;
+};
+
 struct ChunkMaterialPushConstants {
    static bool unflatten(filaflat::Unflattener& unflattener, utils::CString* structVarName,
            utils::FixedCapacityVector<MaterialPushConstant>* materialPushConstants);
@@ -187,7 +187,10 @@ FMaterial* PostProcessManager::PostProcessMaterial::getMaterial(FEngine& engine,
    if (UTILS_UNLIKELY(mSize)) {
        loadMaterial(engine);
    }
-    mMaterial->prepareProgram(Variant{ Variant::type_t(variant) });
+    mMaterial->prepareProgram(Variant{ Variant::type_t(variant) },
+            // TODO(exv): allow post process materials to use mutable spec constants?
+            Program::MutableSpecConstantsInfo(0)
+            );
    return mMaterial;
 }

@@ -423,7 +426,10 @@ UTILS_NOINLINE
 PipelineState PostProcessManager::getPipelineState(
        FMaterial const* const ma, PostProcessVariant variant) const noexcept {
    return {
-            .program = ma->getProgram(Variant{ Variant::type_t(variant) }),
+            .program = ma->getProgram(Variant{ Variant::type_t(variant) },
+                    // TODO(exv): allow post process materials to use mutable spec constants?
+                    Program::MutableSpecConstantsInfo(0)
+                    ),
            .vertexBufferInfo = mFullScreenQuadVbih,
            .pipelineLayout = {
                    .setLayout = {
@@ -246,7 +246,8 @@ void RenderPass::appendCommands(FEngine const& engine,
    for (Command const* first = curr, *last = curr + commandCount ; first != last ; ++first) {
        if (UTILS_LIKELY((first->key & CUSTOM_MASK) == uint64_t(CustomCommand::PASS))) {
            auto ma = first->info.mi->getMaterial();
-            ma->prepareProgram(first->info.materialVariant);
+            ma->prepareProgram(first->info.materialVariant,
+                    first->info.mi->getMutableSpecConstants());
        }
    }
 }
@@ -1098,7 +1099,8 @@ void RenderPass::Executor::execute(FEngine const& engine, DriverApi& driver,
                }

                assert_invariant(ma);
-                pipeline.program = ma->getProgram(info.materialVariant);
+                pipeline.program = ma->getProgram(info.materialVariant,
+                        info.mi->getMutableSpecConstants());

                if (UTILS_UNLIKELY(memcmp(&pipeline, &currentPipeline, sizeof(PipelineState)) != 0)) {
                    currentPipeline = pipeline;
@@ -140,7 +140,6 @@ Engine* FEngine::create(Builder const& builder) {
                .forceGLES2Context = instance->getConfig().forceGLES2Context,
                .stereoscopicType = instance->getConfig().stereoscopicType,
                .assertNativeWindowIsValid = instance->features.backend.opengl.assert_native_window_is_valid,
-                .metalDisablePanicOnDrawableFailure = instance->getConfig().metalDisablePanicOnDrawableFailure,
        };
        instance->mDriver = platform->createDriver(sharedContext, driverConfig);

@@ -760,7 +759,6 @@ int FEngine::loop() {
            .forceGLES2Context = mConfig.forceGLES2Context,
            .stereoscopicType =  mConfig.stereoscopicType,
            .assertNativeWindowIsValid = features.backend.opengl.assert_native_window_is_valid,
-            .metalDisablePanicOnDrawableFailure = mConfig.metalDisablePanicOnDrawableFailure,
    };
    mDriver = mPlatform->createDriver(mSharedGLContext, driverConfig);

@@ -342,8 +342,24 @@ FMaterial::FMaterial(FEngine& engine, const Builder& builder,

    processBlendingMode(parser);
    processSpecializationConstants(engine, builder, parser);
+    processMutableSpecializationConstants(engine, builder, parser);
    processPushConstants(engine, parser);
    processDescriptorSets(engine, parser);
+
+    switch (mMaterialDomain) {
+        case filament::MaterialDomain::SURFACE:
+            mVariantBits = VARIANT_BITS;
+            break;
+        case filament::MaterialDomain::POST_PROCESS:
+            mVariantBits = POST_PROCESS_VARIANT_BITS;
+            break;
+        case filament::MaterialDomain::COMPUTE:
+            mVariantBits = 0;
+            break;
+    }
+    mCachedPrograms = FixedCapacityVector<backend::Handle<backend::HwProgram>>(
+            1 << (mVariantBits + mMaterialMutableConstants.size()));
+
    precacheDepthVariants(engine);

 #if FILAMENT_ENABLE_MATDBG
@@ -449,13 +465,17 @@ void FMaterial::compile(CompilerPriorityQueue const priority,
    UserVariantFilterMask const variantFilter =
            ~variantSpec & UserVariantFilterMask(UserVariantFilterBit::ALL);

+    int mutableConstantVariations = 1 << mMaterialMutableConstants.size();
+
    if (UTILS_LIKELY(mEngine.getDriverApi().isParallelShaderCompileSupported())) {
        auto const& variants = isVariantLit() ?
                VariantUtils::getLitVariants() : VariantUtils::getUnlitVariants();
        for (auto const variant: variants) {
            if (!variantFilter || variant == Variant::filterUserVariant(variant, variantFilter)) {
                if (hasVariant(variant)) {
-                    prepareProgram(variant, priority);
+                    for (int spec = 0; spec < mutableConstantVariations; spec++) {
+                        prepareProgram(variant, Program::MutableSpecConstantsInfo(spec), priority);
+                    }
                }
            }
        }
@@ -496,6 +516,11 @@ FMaterialInstance* FMaterial::getDefaultInstance() noexcept {
    return mDefaultMaterialInstance;
 }

+bool FMaterial::isCached(Variant const variant,
+        const backend::Program::MutableSpecConstantsInfo mutableSpecConstants) const noexcept {
+    return bool(mCachedPrograms[getCachedProgramIndex(variant, mutableSpecConstants)]);
+}
+
 bool FMaterial::hasParameter(const char* name) const noexcept {
    return mUniformInterfaceBlock.hasField(name) ||
           mSamplerInterfaceBlock.hasSampler(name) ||
@@ -536,14 +561,15 @@ bool FMaterial::hasVariant(Variant const variant) const noexcept {
 }

 void FMaterial::prepareProgramSlow(Variant const variant,
+        const backend::Program::MutableSpecConstantsInfo mutableSpecConstants,
        backend::CompilerPriorityQueue const priorityQueue) const noexcept {
    assert_invariant(mEngine.hasFeatureLevel(mFeatureLevel));
    switch (getMaterialDomain()) {
        case MaterialDomain::SURFACE:
-            getSurfaceProgramSlow(variant, priorityQueue);
+            getSurfaceProgramSlow(variant, mutableSpecConstants, priorityQueue);
            break;
        case MaterialDomain::POST_PROCESS:
-            getPostProcessProgramSlow(variant, priorityQueue);
+            getPostProcessProgramSlow(variant, mutableSpecConstants, priorityQueue);
            break;
        case MaterialDomain::COMPUTE:
            // TODO: implement MaterialDomain::COMPUTE
@@ -552,6 +578,7 @@ void FMaterial::prepareProgramSlow(Variant const variant,
 }

 void FMaterial::getSurfaceProgramSlow(Variant const variant,
+        const backend::Program::MutableSpecConstantsInfo mutableSpecConstants,
        CompilerPriorityQueue const priorityQueue) const noexcept {
    // filterVariant() has already been applied in generateCommands(), shouldn't be needed here
    // if we're unlit, we don't have any bits that correspond to lit materials
@@ -562,25 +589,25 @@ void FMaterial::getSurfaceProgramSlow(Variant const variant,
    Variant const vertexVariant   = Variant::filterVariantVertex(variant);
    Variant const fragmentVariant = Variant::filterVariantFragment(variant);

-    Program pb{ getProgramWithVariants(variant, vertexVariant, fragmentVariant) };
+    Program pb{ getProgramWithVariants(variant, vertexVariant, fragmentVariant, mutableSpecConstants) };
    pb.priorityQueue(priorityQueue);
    pb.multiview(
            mEngine.getConfig().stereoscopicType == StereoscopicType::MULTIVIEW &&
            Variant::isStereoVariant(variant));
-    createAndCacheProgram(std::move(pb), variant);
+    createAndCacheProgram(std::move(pb), variant, mutableSpecConstants);
 }

 void FMaterial::getPostProcessProgramSlow(Variant const variant,
+        const backend::Program::MutableSpecConstantsInfo mutableSpecConstants,
        CompilerPriorityQueue const priorityQueue) const noexcept {
-    Program pb{ getProgramWithVariants(variant, variant, variant) };
+    Program pb{ getProgramWithVariants(variant, variant, variant, mutableSpecConstants) };
    pb.priorityQueue(priorityQueue);
-    createAndCacheProgram(std::move(pb), variant);
+    createAndCacheProgram(std::move(pb), variant, mutableSpecConstants);
 }

-Program FMaterial::getProgramWithVariants(
-        Variant variant,
-        Variant vertexVariant,
-        Variant fragmentVariant) const {
+Program FMaterial::getProgramWithVariants(Variant variant,
+        Variant vertexVariant, Variant fragmentVariant,
+        const backend::Program::MutableSpecConstantsInfo mutableSpecConstants) const {
    FEngine const& engine = mEngine;
    const ShaderModel sm = engine.getShaderModel();
    const bool isNoop = engine.getBackend() == Backend::NOOP;
@@ -641,29 +668,34 @@ Program FMaterial::getProgramWithVariants(
            mProgramDescriptorBindings[+DescriptorSetBindingPoints::PER_RENDERABLE]);
    program.descriptorBindings(+DescriptorSetBindingPoints::PER_MATERIAL,
            mProgramDescriptorBindings[+DescriptorSetBindingPoints::PER_MATERIAL]);
-    program.specializationConstants(mSpecializationConstants);
+    program.specializationConstants(mSpecializationConstants,
+            CONFIG_MAX_RESERVED_SPEC_CONSTANTS + mMaterialConstants.size(), mutableSpecConstants);

    program.pushConstants(ShaderStage::VERTEX, mPushConstants[uint8_t(ShaderStage::VERTEX)]);
    program.pushConstants(ShaderStage::FRAGMENT, mPushConstants[uint8_t(ShaderStage::FRAGMENT)]);

-    program.cacheId(hash::combine(size_t(mCacheId), variant.key));
+    program.cacheId(
+            hash::combine(size_t(mCacheId),
+                    hash::combine(variant.key, mutableSpecConstants.getValue())));

    return program;
 }

-void FMaterial::createAndCacheProgram(Program&& p, Variant const variant) const noexcept {
+void FMaterial::createAndCacheProgram(Program&& p, Variant const variant,
+        Program::MutableSpecConstantsInfo const mutableSpecConstants) const noexcept {
    FEngine const& engine = mEngine;
    DriverApi& driverApi = mEngine.getDriverApi();

    bool const isShared = isSharedVariant(variant);
+    size_t const cacheIndex = getCachedProgramIndex(variant, mutableSpecConstants);

    // Check if the default material has this program cached
    if (isShared) {
        FMaterial const* const pDefaultMaterial = engine.getDefaultMaterial();
        if (pDefaultMaterial) {
-            auto const program = pDefaultMaterial->mCachedPrograms[variant.key];
+            auto const program = pDefaultMaterial->mCachedPrograms[cacheIndex];
            if (program) {
-                mCachedPrograms[variant.key] = program;
+                mCachedPrograms[cacheIndex] = program;
                return;
            }
        }
@@ -672,17 +704,17 @@ void FMaterial::createAndCacheProgram(Program&& p, Variant const variant) const
    auto const program = driverApi.createProgram(std::move(p));
    driverApi.setDebugTag(program.getId(), mName);
    assert_invariant(program);
-    mCachedPrograms[variant.key] = program;
+    mCachedPrograms[cacheIndex] = program;

    // If the default material doesn't already have this program cached, and all caching conditions
    // are met (Surface Domain and no custom depth shader), cache it now.
    // New Materials will inherit these program automatically.
    if (isShared) {
        FMaterial const* const pDefaultMaterial = engine.getDefaultMaterial();
-        if (pDefaultMaterial && !pDefaultMaterial->mCachedPrograms[variant.key]) {
+        if (pDefaultMaterial && !pDefaultMaterial->mCachedPrograms[cacheIndex]) {
            // set the tag to the default material name
            driverApi.setDebugTag(program.getId(), mName);
-            pDefaultMaterial->mCachedPrograms[variant.key] = program;
+            pDefaultMaterial->mCachedPrograms[cacheIndex] = program;
        }
    }
 }
@@ -730,6 +762,14 @@ size_t FMaterial::getParameters(ParameterInfo* parameters, size_t count) const n
    return count;
 }

+std::optional<uint32_t> FMaterial::getMutableConstantId(std::string_view name) const noexcept {
+    auto const pos = mMutableSpecializationConstantsNameToIndex.find(name);
+    if (pos != mMutableSpecializationConstantsNameToIndex.end()) {
+        return pos->second;
+    }
+    return std::nullopt;
+}
+
 #if FILAMENT_ENABLE_MATDBG

 // Swaps in an edited version of the original package that was used to create the material. The
@@ -788,7 +828,9 @@ void FMaterial::onQueryCallback(void* userdata, VariantList* pActiveVariants) {

 #endif // FILAMENT_ENABLE_MATDBG

-[[nodiscard]] Handle<HwProgram> FMaterial::getProgramWithMATDBG(Variant const variant) const noexcept {
+[[nodiscard]] Handle<HwProgram> FMaterial::getProgramWithMATDBG(Variant const variant,
+        Program::MutableSpecConstantsInfo const mutableSpecConstants) const noexcept {
+    size_t cacheIndex = getCachedProgramIndex(variant, mutableSpecConstants);
 #if FILAMENT_ENABLE_MATDBG
    assert_invariant((size_t)variant.key < VARIANT_COUNT);
    std::unique_lock lock(mActiveProgramsLock);
@@ -803,13 +845,13 @@ void FMaterial::onQueryCallback(void* userdata, VariantList* pActiveVariants) {
    lock.unlock();
    if (isSharedVariant(variant)) {
        FMaterial const* const pDefaultMaterial = mEngine.getDefaultMaterial();
-        if (pDefaultMaterial && pDefaultMaterial->mCachedPrograms[variant.key]) {
-            return pDefaultMaterial->getProgram(variant);
+        if (pDefaultMaterial && pDefaultMaterial->mCachedPrograms[cacheIndex]) {
+            return pDefaultMaterial->getProgram(variant, mutableSpecConstants);
        }
    }
 #endif
-    assert_invariant(mCachedPrograms[variant.key]);
-    return mCachedPrograms[variant.key];
+    assert_invariant(mCachedPrograms[cacheIndex]);
+    return mCachedPrograms[cacheIndex];
 }

 void FMaterial::destroyPrograms(FEngine& engine,
@@ -822,7 +864,7 @@ void FMaterial::destroyPrograms(FEngine& engine,
        case MaterialDomain::SURFACE: {
            if (mIsDefaultMaterial || mHasCustomDepthShader) {
                // default material, or we have custom depth shaders, we destroy all variants
-                for (size_t k = 0, n = VARIANT_COUNT; k < n; ++k) {
+                for (size_t k = 0, n = cachedPrograms.size(); k < n; ++k) {
                    if ((k & variantMask) == variantValue) {
                        // Only destroy if the handle is valid. Not strictly needed, but we have a lot
                        // of variants, and this generates traffic in the command queue.
@@ -840,7 +882,7 @@ void FMaterial::destroyPrograms(FEngine& engine,
                UTILS_UNUSED_IN_RELEASE
                auto const UTILS_NULLABLE pDefaultMaterial = engine.getDefaultMaterial();

-                for (size_t k = 0, n = VARIANT_COUNT; k < n; ++k) {
+                for (size_t k = 0, n = cachedPrograms.size(); k < n; ++k) {
                    if ((k & variantMask) == variantValue) {
                        // Only destroy if the handle is valid. Not strictly needed, but we have a lot
                        // of variant, and this generates traffic in the command queue.
@@ -852,7 +894,7 @@ void FMaterial::destroyPrograms(FEngine& engine,
                                // During Engine::shutdown, auto-cleanup destroys the
                                // default material first
                                assert_invariant(!pDefaultMaterial ||
-                                        pDefaultMaterial->mCachedPrograms[k]);
+                                        pDefaultMaterial->mCachedPrograms[k & (VARIANT_COUNT - 1)]);
                                // we don't own this variant, skip, but clear the entry.
                                cachedPrograms[k].clear();
                                continue;
@@ -866,7 +908,7 @@ void FMaterial::destroyPrograms(FEngine& engine,
            break;
        }
        case MaterialDomain::POST_PROCESS: {
-            for (size_t k = 0, n = POST_PROCESS_VARIANT_COUNT; k < n; ++k) {
+            for (size_t k = 0, n = cachedPrograms.size(); k < n; ++k) {
                if ((k & variantMask) == variantValue) {
                    // Only destroy if the handle is valid. Not strictly needed, but we have a lot
                    // of variant, and this generates traffic in the command queue.
@@ -878,8 +920,13 @@ void FMaterial::destroyPrograms(FEngine& engine,
            break;
        }
        case MaterialDomain::COMPUTE: {
-            // Compute programs don't have variants
-            driverApi.destroyProgram(std::move(cachedPrograms[0]));
+            for (size_t k = 0, n = cachedPrograms.size(); k < n; ++k) {
+                // Only destroy if the handle is valid. Not strictly needed, but we have a lot
+                // of variant, and this generates traffic in the command queue.
+                if (cachedPrograms[k]) {
+                    driverApi.destroyProgram(std::move(cachedPrograms[k]));
+                }
+            }
            break;
        }
    }
@@ -1114,6 +1161,19 @@ void FMaterial::processSpecializationConstants(FEngine& engine, Builder const& b
    }
 }

+void FMaterial::processMutableSpecializationConstants(FEngine& engine, Builder const& builder,
+        MaterialParser const* const parser) {
+    // Older materials won't have a mutable constants chunk, but that's okay.
+    parser->getMutableConstants(&mMaterialMutableConstants);
+    for (size_t i = 0, c = mMaterialMutableConstants.size(); i < c; i++) {
+        auto& item = mMaterialMutableConstants[i];
+        // the key can be a string_view because mMaterialMutableConstants owns the CString
+        std::string_view const key{ item.name.data(), item.name.size() };
+        mMutableSpecializationConstantsNameToIndex[key] = i;
+        mDefaultMutableSpecializationConstants.set(i, item.defaultValue);
+    }
+}
+
 void FMaterial::processPushConstants(FEngine&, MaterialParser const* parser) {
    FixedCapacityVector<Program::PushConstant>& vertexConstants =
            mPushConstants[uint8_t(ShaderStage::VERTEX)];
@@ -1169,7 +1229,8 @@ void FMaterial::precacheDepthVariants(FEngine& engine) {
            }
            assert_invariant(Variant::isValidDepthVariant(variant));
            if (hasVariant(variant)) {
-                prepareProgram(variant);
+                // The default material shouldn't ever have mutable spec constants.
+                prepareProgram(variant, Program::MutableSpecConstantsInfo(0));
            }
        }
        return;
@@ -1182,9 +1243,16 @@ void FMaterial::precacheDepthVariants(FEngine& engine) {
        FMaterial const* const pDefaultMaterial = engine.getDefaultMaterial();
        assert_invariant(pDefaultMaterial);
        auto const allDepthVariants = VariantUtils::getDepthVariants();
+        int mutableConstantVariations = 1 << mMaterialMutableConstants.size();
        for (auto const variant: allDepthVariants) {
            assert_invariant(Variant::isValidDepthVariant(variant));
-            mCachedPrograms[variant.key] = pDefaultMaterial->mCachedPrograms[variant.key];
+            size_t defaultCacheIndex = pDefaultMaterial->getCachedProgramIndex(variant,
+                    Program::MutableSpecConstantsInfo(0));
+            for (int spec = 0; spec < mutableConstantVariations; spec++) {
+                size_t cacheIndex = getCachedProgramIndex(variant,
+                        Program::MutableSpecConstantsInfo(spec));
+                mCachedPrograms[cacheIndex] = pDefaultMaterial->mCachedPrograms[defaultCacheIndex];
+            }
        }
    }
 }
@@ -122,9 +122,8 @@ public:

    FEngine& getEngine() const noexcept  { return mEngine; }

-    bool isCached(Variant const variant) const noexcept {
-        return bool(mCachedPrograms[variant.key]);
-    }
+    bool isCached(Variant const variant,
+            const backend::Program::MutableSpecConstantsInfo mutableSpecConstants) const noexcept;

    void invalidate(Variant::type_t variantMask = 0, Variant::type_t variantValue = 0) noexcept;

@@ -132,26 +131,31 @@ public:
    // Must be called outside of backend render pass.
    // Must be called before getProgram() below.
    void prepareProgram(Variant const variant,
+            backend::Program::MutableSpecConstantsInfo const mutableSpecConstants,
            backend::CompilerPriorityQueue const priorityQueue = CompilerPriorityQueue::HIGH) const noexcept {
        // prepareProgram() is called for each RenderPrimitive in the scene, so it must be efficient.
-        if (UTILS_UNLIKELY(!isCached(variant))) {
-            prepareProgramSlow(variant, priorityQueue);
+        if (UTILS_UNLIKELY(!isCached(variant, mutableSpecConstants))) {
+            prepareProgramSlow(variant, mutableSpecConstants, priorityQueue);
        }
    }

-    // getProgram returns the backend program for the material's given variant.
+    // getProgram returns the backend program for the material's given variant and set of mutable
+    // spec constants.
    // Must be called after prepareProgram().
    [[nodiscard]]
-    backend::Handle<backend::HwProgram> getProgram(Variant const variant) const noexcept {
+    backend::Handle<backend::HwProgram> getProgram(Variant const variant,
+            backend::Program::MutableSpecConstantsInfo const mutableSpecConstants) const noexcept {
 #if FILAMENT_ENABLE_MATDBG
-        return getProgramWithMATDBG(variant);
+        return getProgramWithMATDBG(variant, mutableSpecConstants);
 #endif
-        assert_invariant(mCachedPrograms[variant.key]);
-        return mCachedPrograms[variant.key];
+        size_t index = getCachedProgramIndex(variant, mutableSpecConstants);
+        assert_invariant(mCachedPrograms[index]);
+        return mCachedPrograms[index];
    }

    [[nodiscard]]
-    backend::Handle<backend::HwProgram> getProgramWithMATDBG(Variant variant) const noexcept;
+    backend::Handle<backend::HwProgram> getProgramWithMATDBG(Variant variant,
+            backend::Program::MutableSpecConstantsInfo const mutableSpecConstants) const noexcept;

    bool isVariantLit() const noexcept { return mIsVariantLit; }

@@ -208,6 +212,12 @@ public:
    }
    size_t getParameters(ParameterInfo* parameters, size_t count) const noexcept;

+    uint32_t getMutableConstantCount() const noexcept { return mMaterialMutableConstants.size(); }
+    std::optional<uint32_t> getMutableConstantId(std::string_view name) const noexcept;
+    backend::Program::MutableSpecConstantsInfo getDefaultMutableConstants() const noexcept {
+        return mDefaultMutableSpecializationConstants;
+    }
+
    uint32_t generateMaterialInstanceId() const noexcept { return mMaterialInstanceId++; }

    void destroyPrograms(FEngine& engine,
@@ -261,34 +271,48 @@ public:
 private:
    bool hasVariant(Variant variant) const noexcept;
    void prepareProgramSlow(Variant variant,
+            backend::Program::MutableSpecConstantsInfo mutableSpecConstants,
            CompilerPriorityQueue priorityQueue) const noexcept;
    void getSurfaceProgramSlow(Variant variant,
+            backend::Program::MutableSpecConstantsInfo mutableSpecConstants,
            CompilerPriorityQueue priorityQueue) const noexcept;
    void getPostProcessProgramSlow(Variant variant,
+            backend::Program::MutableSpecConstantsInfo mutableSpecConstants,
            CompilerPriorityQueue priorityQueue) const noexcept;
    backend::Program getProgramWithVariants(Variant variant,
-            Variant vertexVariant, Variant fragmentVariant) const;
+            Variant vertexVariant, Variant fragmentVariant,
+            backend::Program::MutableSpecConstantsInfo mutableSpecConstants) const;

    void processBlendingMode(MaterialParser const* parser);

    void processSpecializationConstants(FEngine& engine, Builder const& builder,
            MaterialParser const* parser);

+    void processMutableSpecializationConstants(FEngine& engine, Builder const& builder,
+            MaterialParser const* parser);
+
    void processPushConstants(FEngine& engine, MaterialParser const* parser);

    void precacheDepthVariants(FEngine& engine);

    void processDescriptorSets(FEngine& engine, MaterialParser const* parser);

-    void createAndCacheProgram(backend::Program&& p, Variant variant) const noexcept;
+    void createAndCacheProgram(backend::Program&& p, Variant variant,
+            backend::Program::MutableSpecConstantsInfo mutableSpecConstants) const noexcept;

    inline bool isSharedVariant(Variant const variant) const {
        return (mMaterialDomain == MaterialDomain::SURFACE) && !mIsDefaultMaterial &&
               !mHasCustomDepthShader && Variant::isValidDepthVariant(variant);
    }

+    inline size_t getCachedProgramIndex(const Variant variant,
+            const backend::Program::MutableSpecConstantsInfo mutableSpecConstants) const noexcept {
+        return static_cast<size_t>(variant.key) |
+                (static_cast<size_t>(mutableSpecConstants.getValue()) << mVariantBits);
+    }
+
    // try to order by frequency of use
-    mutable std::array<backend::Handle<backend::HwProgram>, VARIANT_COUNT> mCachedPrograms;
+    mutable utils::FixedCapacityVector<backend::Handle<backend::HwProgram>> mCachedPrograms;
    DescriptorSetLayout mPerViewDescriptorSetLayout;
    DescriptorSetLayout mPerViewDescriptorSetLayoutVsm;
    DescriptorSetLayout mDescriptorSetLayout;
@@ -324,6 +348,7 @@ private:
    bool mHasCustomDepthShader = false;
    bool mIsDefaultMaterial = false;
    bool mSpecularAntiAliasing = false;
+    uint8_t mVariantBits;

    // reserve some space to construct the default material instance
    mutable FMaterialInstance* mDefaultMaterialInstance = nullptr;
@@ -348,6 +373,13 @@ private:
    // current specialization constants for the HwProgram
    utils::FixedCapacityVector<backend::Program::SpecializationConstant> mSpecializationConstants;

+    // Mutable constants defined by this Material.
+    utils::FixedCapacityVector<MaterialMutableConstant> mMaterialMutableConstants;
+    // A map from the Constant name to the mMutableMaterialConstant index
+    std::unordered_map<std::string_view, uint32_t> mMutableSpecializationConstantsNameToIndex;
+    // Default values for the mutable constants.
+    backend::Program::MutableSpecConstantsInfo mDefaultMutableSpecializationConstants;
+
    // current push constants for the HwProgram
    std::array<utils::FixedCapacityVector<backend::Program::PushConstant>,
            backend::Program::SHADER_TYPE_COUNT>
@@ -63,6 +63,7 @@ FMaterialInstance::FMaterialInstance(FEngine& engine, FMaterial const* material,
                                     const char* name) noexcept
        : mMaterial(material),
          mDescriptorSet("MaterialInstance", material->getDescriptorSetLayout()),
+          mConstants(material->getDefaultMutableConstants()),
          mCulling(CullingMode::BACK),
          mShadowCulling(CullingMode::BACK),
          mDepthFunc(RasterState::DepthFunc::LE),
@@ -128,6 +129,7 @@ FMaterialInstance::FMaterialInstance(FEngine& engine,
          mTextureParameters(other->mTextureParameters),
          mDescriptorSet(other->mDescriptorSet.duplicate(
                "MaterialInstance", mMaterial->getDescriptorSetLayout())),
+          mConstants(other->mConstants),
          mPolygonOffset(other->mPolygonOffset),
          mStencilState(other->mStencilState),
          mMaskThreshold(other->mMaskThreshold),
@@ -31,6 +31,7 @@

 #include <backend/DriverEnums.h>
 #include <backend/Handle.h>
+#include <backend/Program.h>

 #include <utils/BitmaskEnum.h>
 #include <utils/bitset.h>
@@ -66,7 +67,7 @@ public:
    void terminate(FEngine& engine);

    void commitStreamUniformAssociations(FEngine::DriverApi& driver);
-    
+
    void commit(FEngine& engine) const;

    void commit(FEngine::DriverApi& driver) const;
@@ -242,6 +243,10 @@ public:

    using MaterialInstance::setParameter;

+    const backend::Program::MutableSpecConstantsInfo getMutableSpecConstants() const noexcept {
+        return mConstants;
+    }
+
 private:
    friend class FMaterial;
    friend class MaterialInstance;
@@ -264,6 +269,12 @@ private:
    template<typename T>
    T getParameterImpl(std::string_view name) const;

+    template<typename T>
+    void setConstantImpl(std::string_view name, T const& value);
+
+    template<typename T>
+    T getConstantImpl(std::string_view name) const;
+
    // keep these grouped, they're accessed together in the render-loop
    FMaterial const* mMaterial = nullptr;

@@ -277,6 +288,7 @@ private:
    mutable DescriptorSet mDescriptorSet;
    UniformBuffer mUniforms;
    bool mHasStreamUniformAssociations = false;
+    backend::Program::MutableSpecConstantsInfo mConstants;

    backend::PolygonOffset mPolygonOffset{};
    backend::StencilState mStencilState{};
@@ -8,8 +8,7 @@ material {
        {
            type : sampler2d,
            name : depth,
-            precision: high,
-            unfilterable: true
+            precision: high
        },
        {
            type : float,
@@ -4,8 +4,7 @@ material {
        {
            type : sampler2d,
            name : depth,
-            precision: medium,
-            unfilterable: true
+            precision: medium
        },
        {
            type : float4,
@@ -9,8 +9,7 @@ material {
        {
            type : sampler2d,
            name : depth,
-            precision: high,
-            unfilterable: true
+            precision: high
        },
        {
            type : float4,
@@ -78,3 +77,4 @@ fragment {
        postProcess.color = color;
    }
 }
+
@@ -9,8 +9,7 @@ material {
        {
            type : sampler2d,
            name : depth,
-            precision: medium,
-            unfilterable: true
+            precision: medium
        },
        {
            type : float2,
@@ -9,8 +9,7 @@ material {
        {
            type : sampler2d,
            name : depth,
-            precision: medium,
-            unfilterable: true
+            precision: medium
        },
        {
            type : float2,
@@ -5,8 +5,7 @@ material {
            type : sampler2d,
            name : depth,
            precision: high,
-            multisample : true,
-            unfilterable: true
+            multisample : true
        }
    ],
    outputs : [
@@ -4,8 +4,7 @@ material {
        {
            type : sampler2d,
            name : depth,
-            precision: high,
-            unfilterable: true
+            precision: high
        },
        {
            type : mat4,
@@ -4,8 +4,7 @@ material {
        {
            type : sampler2d,
            name : depth,
-            precision: high,
-            unfilterable: true
+            precision: high
        },
        {
            type : mat4,
@@ -4,8 +4,7 @@ material {
        {
            type : sampler2d,
            name : depth,
-            precision: high,
-            unfilterable: true
+            precision: high
        }
    ],
    variables : [
@@ -4,8 +4,7 @@ material {
        {
            type : sampler2d,
            name : depth,
-            precision: high,
-            unfilterable: true
+            precision: high
        },
        {
            type : mat4,
@@ -4,8 +4,7 @@ material {
        {
            type : sampler2d,
            name : depth,
-            precision: high,
-            unfilterable: true
+            precision: high
        },
        {
            type : mat4,
@@ -54,6 +54,7 @@ enum UTILS_PUBLIC ChunkType : uint64_t {
    MaterialDescriptorSetLayoutInfo = charTo64bitNum("MAT_DSLI"),
    MaterialProperties = charTo64bitNum("MAT_PROP"),
    MaterialConstants = charTo64bitNum("MAT_CONS"),
+    MaterialMutableConstants = charTo64bitNum("MAT_MCNS"),
    MaterialPushConstants = charTo64bitNum("MAT_PCON"),

    MaterialName = charTo64bitNum("MAT_NAME"),
@@ -28,7 +28,7 @@
 namespace filament {

 // update this when a new version of filament wouldn't work with older materials
-static constexpr size_t MATERIAL_VERSION = 61;
+static constexpr size_t MATERIAL_VERSION = 60;

 /**
 * Supported shading models
@@ -30,7 +30,19 @@ struct MaterialConstant {
    ConstantType type;

    MaterialConstant() = default;
-    MaterialConstant(const char* name, ConstantType type) : name(name), type(type)  {}
+    MaterialConstant(utils::CString name, ConstantType type) : name(std::move(name)), type(type)  {}
+};
+
+struct MaterialMutableConstant {
+    utils::CString name;
+    // The Filament engine stores mutable spec constants as a bitfield, meaning that must override
+    // all of the default values built into in the shader souces. Therefore, unlike the immutable
+    // constants above, we have to store the default values in the material metadata.
+    bool defaultValue;
+
+    MaterialMutableConstant() = default;
+    MaterialMutableConstant(utils::CString name, bool defaultValue)
+            : name(std::move(name)), defaultValue(defaultValue) {}
 };

 }
@@ -61,7 +61,6 @@ public:
        Type type;                  // type of this sampler
        Format format;              // format of this sampler
        Precision precision;        // precision of this sampler
-        bool unfilterable;          // whether the sampling should be unfiltered.
        bool multisample;           // multisample capable
        ShaderStageFlags stages;    // stages the sampler can be accessed from
    };
@@ -84,7 +83,6 @@ public:
            Type type;                      // type of this sampler
            Format format;                  // format of this sampler
            Precision precision;            // precision of this sampler
-            bool unfilterable = false;      // whether the sampling should be unfiltered.
            bool multisample = false;       // multisample capable
            ShaderStageFlags stages =
                    ShaderStageFlags::ALL_SHADER_STAGE_FLAGS; // shader stages using this sampler
@@ -97,8 +95,7 @@ public:

        // Add a sampler
        Builder& add(std::string_view samplerName, Binding binding, Type type, Format format,
-                Precision precision = Precision::MEDIUM, bool unfilterable = false,
-                bool multisample = false,
+                Precision precision = Precision::MEDIUM, bool multisample = false,
                ShaderStageFlags stages = ShaderStageFlags::ALL_SHADER_STAGE_FLAGS) noexcept;

        // Add multiple samplers
@@ -50,12 +50,12 @@ SamplerInterfaceBlock::Builder::stageFlags(backend::ShaderStageFlags stageFlags)
 }

 SamplerInterfaceBlock::Builder& SamplerInterfaceBlock::Builder::add(std::string_view samplerName,
-        Binding binding, Type type, Format format, Precision precision, bool unfilterable,
-        bool multisample, ShaderStageFlags stages) noexcept {
+        Binding binding, Type type, Format format, Precision precision, bool multisample,
+        ShaderStageFlags stages) noexcept {
    mEntries.push_back({
-        { samplerName.data(), samplerName.size() }, // name
-        {},                                         // uniform name
-        binding, type, format, precision, unfilterable, multisample, stages });
+            { samplerName.data(), samplerName.size() }, // name
+            { }, // uniform name
+            binding, type, format, precision, multisample, stages });
    return *this;
 }

@@ -65,9 +65,8 @@ SamplerInterfaceBlock SamplerInterfaceBlock::Builder::build() {

 SamplerInterfaceBlock::Builder& SamplerInterfaceBlock::Builder::add(
        std::initializer_list<ListEntry> list) noexcept {
-    for (auto& e: list) {
-        add(e.name, e.binding, e.type, e.format, e.precision, e.unfilterable, e.multisample,
-                e.stages);
+    for (auto& e : list) {
+        add(e.name, e.binding, e.type, e.format, e.precision, e.multisample, e.stages);
    }
    return *this;
 }
@@ -307,14 +307,15 @@ public:
    MaterialBuilder& parameter(const char* name, size_t size, UniformType type,
            ParameterPrecision precision = ParameterPrecision::DEFAULT);

-    //! Add a constant parameter to this material.
+    //! Add a (mutable) constant parameter to this material.
    template<typename T>
    using is_supported_constant_parameter_t = typename std::enable_if<
            std::is_same<int32_t, T>::value ||
            std::is_same<float, T>::value ||
            std::is_same<bool, T>::value>::type;
    template<typename T, typename = is_supported_constant_parameter_t<T>>
-    MaterialBuilder& constant(const char *name, ConstantType type, T defaultValue = 0);
+    MaterialBuilder& constant(const char* name, ConstantType type, bool isMutable,
+            T defaultValue = 0);

    /**
     * Add a sampler parameter to this material.
@@ -323,8 +324,8 @@ public:
     */
    MaterialBuilder& parameter(const char* name, SamplerType samplerType,
            SamplerFormat format = SamplerFormat::FLOAT,
-            ParameterPrecision precision = ParameterPrecision::DEFAULT, bool unfilterable = false,
-            bool multisample = false, const char* transformName = "",
+            ParameterPrecision precision = ParameterPrecision::DEFAULT, bool multisample = false,
+            const char* transformName = "",
            ShaderStageFlags stages = ShaderStageFlags::ALL_SHADER_STAGE_FLAGS);

    MaterialBuilder& buffer(filament::BufferInterfaceBlock bib);
@@ -658,17 +659,9 @@ public:

        // Sampler
        Parameter(const char* paramName, SamplerType t, SamplerFormat f, ParameterPrecision p,
-                bool unfilterable, bool ms, const char* tn, ShaderStageFlags s)
-            : name(paramName),
-              size(1),
-              precision(p),
-              samplerType(t),
-              format(f),
-              parameterType(SAMPLER),
-              unfilterable(unfilterable),
-              multisample(ms),
-              transformName(tn),
-              stages(s) {}
+                bool ms, const char* tn, ShaderStageFlags s)
+            : name(paramName), size(1), precision(p), samplerType(t), format(f),
+              parameterType(SAMPLER), multisample(ms), transformName(tn), stages(s) { }

        // Uniform
        Parameter(const char* paramName, UniformType t, size_t typeSize, ParameterPrecision p)
@@ -684,9 +677,8 @@ public:
        ParameterPrecision precision;
        SamplerType samplerType;
        SubpassType subpassType;
-        SamplerFormat format = SamplerFormat::INT; // 0 of the enum
-        bool unfilterable = false;
-        bool multisample = false;
+        SamplerFormat format;
+        bool multisample;
        utils::CString transformName;
        ShaderStageFlags stages;
        enum {
@@ -880,6 +872,7 @@ private:
    PropertyList mProperties;
    ParameterList mParameters;
    ConstantList mConstants;
+    ConstantList mMutableConstants;
    PushConstantList mPushConstants;
    SubpassList mSubpasses;
    VariableList mVariables;
@@ -300,29 +300,36 @@ MaterialBuilder& MaterialBuilder::parameter(const char* name, UniformType const
    return parameter(name, 1, type, precision);
 }

+
 MaterialBuilder& MaterialBuilder::parameter(const char* name, SamplerType samplerType,
-        SamplerFormat format, ParameterPrecision precision, bool unfilterable, bool multisample,
+        SamplerFormat format, ParameterPrecision precision, bool multisample,
        const char* transformName, ShaderStageFlags stages) {
-    FILAMENT_CHECK_PRECONDITION(
-            !multisample || (format != SamplerFormat::SHADOW &&
-                                    (samplerType == SamplerType::SAMPLER_2D ||
-                                            samplerType == SamplerType::SAMPLER_2D_ARRAY)))
+    FILAMENT_CHECK_PRECONDITION(!multisample ||
+            (format != SamplerFormat::SHADOW &&
+                    (samplerType == SamplerType::SAMPLER_2D ||
+                            samplerType == SamplerType::SAMPLER_2D_ARRAY)))
            << "multisample samplers only possible with SAMPLER_2D or SAMPLER_2D_ARRAY,"
               " as long as type is not SHADOW";

    FILAMENT_CHECK_POSTCONDITION(mParameterCount < MAX_PARAMETERS_COUNT) << "Too many parameters";
-    mParameters[mParameterCount++] = { name, samplerType, format, precision, unfilterable,
-        multisample, transformName, stages };
+    mParameters[mParameterCount++] = { name, samplerType, format, precision, multisample, transformName, stages };
    return *this;
 }

 template<typename T, typename>
-MaterialBuilder& MaterialBuilder::constant(const char* name, ConstantType const type, T defaultValue) {
-    auto result = std::find_if(mConstants.begin(), mConstants.end(), [name](const Constant& c) {
-        return c.name == CString(name);
-    });
-    FILAMENT_CHECK_POSTCONDITION(result == mConstants.end())
+MaterialBuilder& MaterialBuilder::constant(const char* name, ConstantType const type,
+        bool isMutable, T defaultValue) {
+    ConstantList& constants = isMutable ? mMutableConstants : mConstants;
+    const ConstantList& otherConstants = isMutable ? mConstants : mMutableConstants;
+    auto result = std::find_if(constants.begin(), constants.end(),
+            [name](const Constant& c) { return c.name == CString(name); });
+    FILAMENT_CHECK_POSTCONDITION(result == constants.end())
            << "There is already a constant parameter present with the name " << name << ".";
+    auto otherResult = std::find_if(otherConstants.begin(), otherConstants.end(),
+            [name](const Constant& c) { return c.name == CString(name); });
+    FILAMENT_CHECK_POSTCONDITION(otherResult == otherConstants.end())
+            << "There is already a constant parameter present with the name " << name << ".";
+
    Constant constant {
            .name = CString(name),
            .type = type,
@@ -354,15 +361,20 @@ MaterialBuilder& MaterialBuilder::constant(const char* name, ConstantType const
        assert_invariant(false);
    }

-    mConstants.push_back(constant);
+    if (isMutable) {
+        FILAMENT_CHECK_POSTCONDITION(type == ConstantType::BOOL)
+                << "Mutable spec constants must be bool type.";
+    }
+
+    constants.push_back(constant);
    return *this;
 }
 template MaterialBuilder& MaterialBuilder::constant<int32_t>(
-        const char* name, ConstantType type, int32_t defaultValue);
+        const char* name, ConstantType type, bool isMutable, int32_t defaultValue);
 template MaterialBuilder& MaterialBuilder::constant<float>(
-        const char* name, ConstantType type, float defaultValue);
+        const char* name, ConstantType type, bool isMutable, float defaultValue);
 template MaterialBuilder& MaterialBuilder::constant<bool>(
-        const char* name, ConstantType type, bool defaultValue);
+        const char* name, ConstantType type, bool isMutable, bool defaultValue);

 MaterialBuilder& MaterialBuilder::buffer(BufferInterfaceBlock bib) {
    FILAMENT_CHECK_POSTCONDITION(mBuffers.size() < MAX_BUFFERS_COUNT) << "Too many buffers";
@@ -639,12 +651,10 @@ void MaterialBuilder::prepareToBuild(MaterialInfo& info) noexcept {
        assert_invariant(!param.isSubpass());
        if (param.isSampler()) {
            sbb.add({ param.name.data(), param.name.size() }, binding, param.samplerType,
-                    param.format, param.precision, param.unfilterable, param.multisample,
-                    param.stages);
+                    param.format, param.precision, param.multisample, param.stages);
            if (!param.transformName.empty()) {
-                ibb.add({ { { param.transformName.data(), param.transformName.size() },
-                    uint8_t(binding), 0, UniformType::MAT3, Precision::DEFAULT,
-                    FeatureLevel::FEATURE_LEVEL_0 } });
+                ibb.add({{{ param.transformName.data(), param.transformName.size() }, uint8_t(binding),
+                          0, UniformType::MAT3, Precision::DEFAULT, FeatureLevel::FEATURE_LEVEL_0 }});
            }
            binding++;
        } else if (param.isUniform()) {
@@ -909,7 +919,7 @@ bool MaterialBuilder::generateShaders(JobSystem& jobSystem, const std::vector<Va
    BlobDictionary spirvDictionary;
    // End: must be protected by lock

-    ShaderGenerator sg(mProperties, mVariables, mOutputs, mDefines, mConstants, mPushConstants,
+    ShaderGenerator sg(mProperties, mVariables, mOutputs, mDefines, mConstants, mMutableConstants, mPushConstants,
            mMaterialFragmentCode.getResolved(), mMaterialFragmentCode.getLineOffset(),
            mMaterialVertexCode.getResolved(), mMaterialVertexCode.getLineOffset(),
            mMaterialDomain);
@@ -1494,7 +1504,7 @@ bool MaterialBuilder::needsStandardDepthProgram() const noexcept {
 std::string MaterialBuilder::peek(backend::ShaderStage const stage,
        const CodeGenParams& params, const PropertyList& properties) noexcept {

-    ShaderGenerator const sg(properties, mVariables, mOutputs, mDefines, mConstants, mPushConstants,
+    ShaderGenerator const sg(properties, mVariables, mOutputs, mDefines, mConstants, mMutableConstants, mPushConstants,
            mMaterialFragmentCode.getResolved(), mMaterialFragmentCode.getLineOffset(),
            mMaterialVertexCode.getResolved(), mMaterialVertexCode.getLineOffset(),
            mMaterialDomain);
@@ -1620,9 +1630,14 @@ void MaterialBuilder::writeCommonChunks(ChunkContainer& container, MaterialInfo&
    // User constant parameters
    FixedCapacityVector<MaterialConstant> constantsEntry(mConstants.size());
    std::transform(mConstants.begin(), mConstants.end(), constantsEntry.begin(),
-            [](Constant const& c) { return MaterialConstant(c.name.c_str(), c.type); });
+            [](Constant const& c) { return MaterialConstant(c.name, c.type); });
    container.push<MaterialConstantParametersChunk>(std::move(constantsEntry));

+    FixedCapacityVector<MaterialMutableConstant> mutableConstantsEntry(mMutableConstants.size());
+    std::transform(mMutableConstants.begin(), mMutableConstants.end(), mutableConstantsEntry.begin(),
+            [](Constant const& c) { return MaterialMutableConstant(c.name, c.defaultValue.b); });
+    container.push<MaterialMutableConstantParametersChunk>(std::move(mutableConstantsEntry));
+
    FixedCapacityVector<MaterialPushConstant> pushConstantsEntry(mPushConstants.size());
    std::transform(mPushConstants.begin(), mPushConstants.end(), pushConstantsEntry.begin(),
            [](PushConstant const& c) {
@@ -77,7 +77,6 @@ void MaterialSamplerInterfaceBlockChunk::flatten(Flattener& f) {
        f.writeUint8(static_cast<uint8_t>(sInfo.type));
        f.writeUint8(static_cast<uint8_t>(sInfo.format));
        f.writeUint8(static_cast<uint8_t>(sInfo.precision));
-        f.writeBool(sInfo.unfilterable);
        f.writeBool(sInfo.multisample);
    }
 }
@@ -118,6 +117,20 @@ void MaterialConstantParametersChunk::flatten(Flattener& f) {

 // ------------------------------------------------------------------------------------------------

+MaterialMutableConstantParametersChunk::MaterialMutableConstantParametersChunk(
+        FixedCapacityVector<MaterialMutableConstant> constants)
+    : Chunk(MaterialMutableConstants), mConstants(std::move(constants)) {}
+
+void MaterialMutableConstantParametersChunk::flatten(Flattener& f) {
+    f.writeUint64(mConstants.size());
+    for (const auto& constant : mConstants) {
+        f.writeString(constant.name.c_str());
+        f.writeBool(constant.defaultValue);
+    }
+}
+
+// ------------------------------------------------------------------------------------------------
+
 MaterialPushConstantParametersChunk::MaterialPushConstantParametersChunk(
        CString const& structVarName, FixedCapacityVector<MaterialPushConstant> constants)
    : Chunk(MaterialPushConstants),
@@ -231,11 +244,7 @@ void MaterialDescriptorSetLayoutChunk::flatten(Flattener& f) {
        f.writeUint8(uint8_t(descriptor_sets::getDescriptorType(entry.type, entry.format)));
        f.writeUint8(uint8_t(entry.stages));
        f.writeUint8(entry.binding);
-        if (entry.unfilterable) {
-            f.writeUint8(uint8_t(DescriptorFlags::UNFILTERABLE));
-        } else {
-            f.writeUint8(uint8_t(DescriptorFlags::NONE));
-        }
+        f.writeUint8(uint8_t(DescriptorFlags::NONE));
        f.writeUint16(0);
    }
 }
@@ -18,6 +18,7 @@
 #define TNT_FILAMAT_MAT_INTEFFACE_BLOCK_CHUNK_H

 #include "Chunk.h"
+#include "private/filament/ConstantInfo.h"

 #include <backend/Program.h>

@@ -91,6 +92,20 @@ private:

 // ------------------------------------------------------------------------------------------------

+class MaterialMutableConstantParametersChunk final : public Chunk {
+public:
+    explicit MaterialMutableConstantParametersChunk(
+            FixedCapacityVector<filament::MaterialMutableConstant> constants);
+    ~MaterialMutableConstantParametersChunk() override = default;
+
+private:
+    void flatten(Flattener&) override;
+
+    FixedCapacityVector<filament::MaterialMutableConstant> mConstants;
+};
+
+// ------------------------------------------------------------------------------------------------
+
 class MaterialPushConstantParametersChunk final : public Chunk {
 public:
    explicit MaterialPushConstantParametersChunk(CString const& structVarName,
@@ -283,10 +283,17 @@ void ShaderGenerator::appendShader(io::sstream& ss,
    }
 }

-void ShaderGenerator::generateUserSpecConstants(
-        const CodeGenerator& cg, io::sstream& fs, MaterialBuilder::ConstantList const& constants) {
+void ShaderGenerator::generateAllUserSpecConstants(const CodeGenerator& cg, io::sstream& fs,
+        MaterialBuilder::ConstantList const& constants,
+        MaterialBuilder::ConstantList const& mutableConstants) {
+    generateUserSpecConstants(cg, fs, constants, 0);
+    generateUserSpecConstants(cg, fs, mutableConstants, constants.size());
+}
+
+void ShaderGenerator::generateUserSpecConstants(const CodeGenerator& cg, io::sstream& fs,
+        MaterialBuilder::ConstantList const& constants, size_t offset) {
    // Constants 0 to CONFIG_MAX_RESERVED_SPEC_CONSTANTS - 1 are reserved by Filament.
-    size_t index = CONFIG_MAX_RESERVED_SPEC_CONSTANTS;
+    size_t index = CONFIG_MAX_RESERVED_SPEC_CONSTANTS + offset;
    for (const auto& constant : constants) {
        std::string const fullName = std::string("materialConstants_") + constant.name.c_str();
        switch (constant.type) {
@@ -314,6 +321,7 @@ ShaderGenerator::ShaderGenerator(
        MaterialBuilder::OutputList const& outputs,
        MaterialBuilder::PreprocessorDefineList const& defines,
        MaterialBuilder::ConstantList const& constants,
+        MaterialBuilder::ConstantList const& mutableConstants,
        MaterialBuilder::PushConstantList const& pushConstants,
        CString const& materialCode, size_t const lineOffset,
        CString const& materialVertexCode, size_t const vertexLineOffset,
@@ -336,6 +344,7 @@ ShaderGenerator::ShaderGenerator(
    mMaterialDomain = materialDomain;
    mDefines = defines;
    mConstants = constants;
+    mMutableConstants = mutableConstants;
    mPushConstants = pushConstants;

    if (mMaterialFragmentCode.empty()) {
@@ -392,7 +401,7 @@ std::string ShaderGenerator::createSurfaceVertexProgram(ShaderModel const shader

    cg.generateCommonProlog(vs, ShaderStage::VERTEX, material, variant);

-    generateUserSpecConstants(cg, vs, mConstants);
+    generateAllUserSpecConstants(cg, vs, mConstants, mMutableConstants);

    // note: even if the user vertex shader is empty, we can't use the "optimized" version if
    // we're in masked mode because fragment shader needs the color varyings
@@ -525,7 +534,7 @@ std::string ShaderGenerator::createSurfaceFragmentProgram(ShaderModel const shad
    io::sstream fs;
    cg.generateCommonProlog(fs, ShaderStage::FRAGMENT, material, variant);

-    generateUserSpecConstants(cg, fs, mConstants);
+    generateAllUserSpecConstants(cg, fs, mConstants, mMutableConstants);

    generateSurfaceMaterialVariantDefines(
            fs, ShaderStage::FRAGMENT, featureLevel, material, variant);
@@ -681,7 +690,7 @@ std::string ShaderGenerator::createSurfaceComputeProgram(ShaderModel const shade

    cg.generateCommonProlog(s, ShaderStage::COMPUTE, material, {});

-    generateUserSpecConstants(cg, s, mConstants);
+    generateAllUserSpecConstants(cg, s, mConstants, mMutableConstants);

    CodeGenerator::generateSurfaceTypes(s, ShaderStage::COMPUTE);

@@ -721,7 +730,7 @@ std::string ShaderGenerator::createPostProcessVertexProgram(ShaderModel const sm
    io::sstream vs;
    cg.generateCommonProlog(vs, ShaderStage::VERTEX, material, {});

-    generateUserSpecConstants(cg, vs, mConstants);
+    generateAllUserSpecConstants(cg, vs, mConstants, mMutableConstants);

    CodeGenerator::generateDefine(vs, "LOCATION_POSITION", uint32_t(POSITION));

@@ -765,7 +774,7 @@ std::string ShaderGenerator::createPostProcessFragmentProgram(ShaderModel const
    io::sstream fs;
    cg.generateCommonProlog(fs, ShaderStage::FRAGMENT, material, {});

-    generateUserSpecConstants(cg, fs, mConstants);
+    generateAllUserSpecConstants(cg, fs, mConstants, mMutableConstants);

    generatePostProcessMaterialVariantDefines(fs, PostProcessVariant(variant));

@@ -51,6 +51,7 @@ public:
            MaterialBuilder::OutputList const& outputs,
            MaterialBuilder::PreprocessorDefineList const& defines,
            MaterialBuilder::ConstantList const& constants,
+            MaterialBuilder::ConstantList const& mutableConstants,
            MaterialBuilder::PushConstantList const& pushConstants,
            utils::CString const& materialCode,
            size_t lineOffset,
@@ -104,9 +105,14 @@ private:
    static void generatePostProcessMaterialVariantDefines(utils::io::sstream& out,
            filament::PostProcessVariant variant) noexcept;

+    static void generateAllUserSpecConstants(const CodeGenerator& cg, utils::io::sstream& fs,
+            MaterialBuilder::ConstantList const& constants,
+            MaterialBuilder::ConstantList const& mutableConstants);
+
    static void generateUserSpecConstants(
            const CodeGenerator& cg, utils::io::sstream& fs,
-            MaterialBuilder::ConstantList const& constants);
+            MaterialBuilder::ConstantList const& constants,
+            size_t offset);

    std::string createPostProcessVertexProgram(filament::backend::ShaderModel sm,
            MaterialBuilder::TargetApi targetApi, MaterialBuilder::TargetLanguage targetLanguage,
@@ -135,6 +141,7 @@ private:
    MaterialBuilder::MaterialDomain mMaterialDomain;
    MaterialBuilder::PreprocessorDefineList mDefines;
    MaterialBuilder::ConstantList mConstants;
+    MaterialBuilder::ConstantList mMutableConstants;
    MaterialBuilder::PushConstantList mPushConstants;
    utils::CString mMaterialFragmentCode;   // fragment or compute code
    utils::CString mMaterialVertexCode;
@@ -888,10 +888,10 @@ TEST_F(MaterialCompiler, ConstantParameter) {
        }
    )");
  filamat::MaterialBuilder builder;
-  builder.constant("myFloatConstant", ConstantType::FLOAT, 1.0f);
-  builder.constant("myIntConstant", ConstantType::INT, 123);
-  builder.constant("myBoolConstant", ConstantType::BOOL, true);
-  builder.constant<bool>("myOtherBoolConstant", ConstantType::BOOL);
+  builder.constant("myFloatConstant", ConstantType::FLOAT, /*isMutable=*/false, 1.0f);
+  builder.constant("myIntConstant", ConstantType::INT, /*isMutable=*/false, 123);
+  builder.constant("myBoolConstant", ConstantType::BOOL, /*isMutable=*/false, true);
+  builder.constant<bool>("myOtherBoolConstant", ConstantType::BOOL, /*isMutable=*/false);

  builder.shading(filament::Shading::LIT);
  builder.material(shaderCode.c_str());
@@ -904,8 +904,8 @@ TEST_F(MaterialCompiler, ConstantParameterSameName) {
 #ifdef __EXCEPTIONS
    EXPECT_THROW({
        filamat::MaterialBuilder builder;
-        builder.constant("myFloatConstant", ConstantType::FLOAT, 1.0f);
-        builder.constant("myFloatConstant", ConstantType::FLOAT, 1.0f);
+        builder.constant("myFloatConstant", ConstantType::FLOAT, /*isMutable=*/false, 1.0f);
+        builder.constant("myFloatConstant", ConstantType::FLOAT, /*isMutable=*/false, 1.0f);
    }, utils::PostconditionPanic);
 #endif
 }
@@ -914,7 +914,7 @@ TEST_F(MaterialCompiler, ConstantParameterWrongType) {
 #ifdef __EXCEPTIONS
    EXPECT_THROW({
        filamat::MaterialBuilder builder;
-        builder.constant("myFloatConstant", ConstantType::FLOAT, 10);
+        builder.constant("myFloatConstant", ConstantType::FLOAT, /*isMutable=*/false, 10);
    }, utils::PostconditionPanic);
 #endif
 }
@@ -21,7 +21,6 @@
 #include <filament/Material.h>
 #include <filament/MaterialInstance.h>

-#include <utils/FixedCapacityVector.h>
 #include <utils/compiler.h>

 #include <array>
@@ -197,7 +196,6 @@ void processShaderString(std::string* shader, const UvMap& uvmap,
 * Creates a material provider that builds materials on the fly, composing GLSL at run time.
 *
 * @param optimizeShaders Optimizes shaders, but at significant cost to construction time.
- * @param variantFilters  Filter out variants that are not required.
 * @return New material provider that can build materials at run time.
 *
 * Requires \c libfilamat to be linked in. Not available in \c libgltfio_core.
@@ -205,8 +203,7 @@ void processShaderString(std::string* shader, const UvMap& uvmap,
 * @see createUbershaderProvider
 */
 UTILS_PUBLIC
-MaterialProvider* createJitShaderProvider(Engine* engine, bool optimizeShaders = false,
-        utils::FixedCapacityVector<char const*> const& variantFilters = {});
+MaterialProvider* createJitShaderProvider(Engine* engine, bool optimizeShaders = false);

 /**
 * Creates a material provider that loads a small set of pre-built materials.
@@ -16,7 +16,6 @@

 #include <gltfio/MaterialProvider.h>

-#include <filament/MaterialEnums.h>
 #include <filamat/MaterialBuilder.h>

 #include <utils/Hash.h>
@@ -24,7 +23,6 @@
 #include <tsl/robin_map.h>

 #include <string>
-#include <unordered_map>

 using namespace filamat;
 using namespace filament;
@@ -35,8 +33,7 @@ namespace {

 class JitShaderProvider : public MaterialProvider {
 public:
-    explicit JitShaderProvider(Engine* engine, bool optimizeShaders,
-            utils::FixedCapacityVector<char const*> const& variantFilters);
+    explicit JitShaderProvider(Engine* engine, bool optimizeShaders);
    ~JitShaderProvider() override;

    MaterialInstance* createMaterialInstance(MaterialKey* config, UvMap* uvmap,
@@ -52,37 +49,15 @@ public:
        return false;
    }

-private:
    using HashFn = hash::MurmurHashFn<MaterialKey>;
    tsl::robin_map<MaterialKey, Material*, HashFn> mCache;
    std::vector<Material*> mMaterials;
    Engine* const mEngine;
    const bool mOptimizeShaders;
-    filament::UserVariantFilterMask mVariantFilter;
 };

-JitShaderProvider::JitShaderProvider(Engine* engine, bool optimizeShaders,
-        utils::FixedCapacityVector<char const*> const& variantFilters)
-    : mEngine(engine),
-      mOptimizeShaders(optimizeShaders) {
-
-    // Note that this is the same as the list in tools/matc/src/ParametersProcessor.cpp
-    static const std::unordered_map<std::string, filament::UserVariantFilterBit> strToEnum  = [] {
-        std::unordered_map<std::string, filament::UserVariantFilterBit> strToEnum;
-        strToEnum["directionalLighting"]    = filament::UserVariantFilterBit::DIRECTIONAL_LIGHTING;
-        strToEnum["dynamicLighting"]        = filament::UserVariantFilterBit::DYNAMIC_LIGHTING;
-        strToEnum["shadowReceiver"]         = filament::UserVariantFilterBit::SHADOW_RECEIVER;
-        strToEnum["skinning"]               = filament::UserVariantFilterBit::SKINNING;
-        strToEnum["vsm"]                    = filament::UserVariantFilterBit::VSM;
-        strToEnum["fog"]                    = filament::UserVariantFilterBit::FOG;
-        strToEnum["ssr"]                    = filament::UserVariantFilterBit::SSR;
-        strToEnum["stereo"]                 = filament::UserVariantFilterBit::STE;
-        return strToEnum;
-    }();
-
-    for (auto& filterStr : variantFilters) {
-        mVariantFilter |= (uint32_t)strToEnum.at(filterStr);
-    }
+JitShaderProvider::JitShaderProvider(Engine* engine, bool optimizeShaders) : mEngine(engine),
+        mOptimizeShaders(optimizeShaders) {
    MaterialBuilder::init();
 }

@@ -369,25 +344,24 @@ std::string shaderFromKey(const MaterialKey& config) {
 }

 Material* createMaterial(Engine* engine, const MaterialKey& config, const UvMap& uvmap,
-        const char* name, bool optimizeShaders, filament::UserVariantFilterMask variantFilter) {
+        const char* name, bool optimizeShaders) {
    std::string shader = shaderFromKey(config);
    processShaderString(&shader, uvmap, config);
    MaterialBuilder builder;
    builder.name(name)
-            .flipUV(false)
-            .specularAmbientOcclusion(MaterialBuilder::SpecularAmbientOcclusion::SIMPLE)
-            .specularAntiAliasing(true)
-            .clearCoatIorChange(false)
-            .material(shader.c_str())
-            .doubleSided(config.doubleSided)
-            .transparencyMode(config.doubleSided
-                                      ? MaterialBuilder::TransparencyMode::TWO_PASSES_TWO_SIDES
-                                      : MaterialBuilder::TransparencyMode::DEFAULT)
-            .reflectionMode(MaterialBuilder::ReflectionMode::SCREEN_SPACE)
-            .targetApi(filamat::targetApiFromBackend(engine->getBackend()))
-            .stereoscopicType(engine->getConfig().stereoscopicType)
-            .stereoscopicEyeCount(engine->getConfig().stereoscopicEyeCount)
-            .variantFilter(variantFilter);
+           .flipUV(false)
+           .specularAmbientOcclusion(MaterialBuilder::SpecularAmbientOcclusion::SIMPLE)
+           .specularAntiAliasing(true)
+           .clearCoatIorChange(false)
+           .material(shader.c_str())
+           .doubleSided(config.doubleSided)
+           .transparencyMode(config.doubleSided ?
+           MaterialBuilder::TransparencyMode::TWO_PASSES_TWO_SIDES :
+           MaterialBuilder::TransparencyMode::DEFAULT)
+           .reflectionMode(MaterialBuilder::ReflectionMode::SCREEN_SPACE)
+           .targetApi(filamat::targetApiFromBackend(engine->getBackend()))
+           .stereoscopicType(engine->getConfig().stereoscopicType)
+           .stereoscopicEyeCount(engine->getConfig().stereoscopicEyeCount);

    if (!optimizeShaders) {
        builder.optimization(MaterialBuilder::Optimization::NONE);
@@ -615,13 +589,13 @@ Material* JitShaderProvider::getMaterial(MaterialKey* config, UvMap* uvmap, cons
    constrainMaterial(config, uvmap);
    auto iter = mCache.find(*config);
    if (iter == mCache.end()) {
+
        bool optimizeShaders = mOptimizeShaders;
 #ifndef NDEBUG
        optimizeShaders = false;
 #endif

-        Material* mat =
-                createMaterial(mEngine, *config, *uvmap, label, optimizeShaders, mVariantFilter);
+        Material* mat = createMaterial(mEngine, *config, *uvmap, label, optimizeShaders);
        mCache.emplace(std::make_pair(*config, mat));
        mMaterials.push_back(mat);
        return mat;
@@ -638,9 +612,8 @@ MaterialInstance* JitShaderProvider::createMaterialInstance(MaterialKey* config,

 namespace filament::gltfio {

-MaterialProvider* createJitShaderProvider(Engine* engine, bool optimizeShaders,
-        utils::FixedCapacityVector<char const*> const& variantFilters) {
-    return new JitShaderProvider(engine, optimizeShaders, variantFilters);
+MaterialProvider* createJitShaderProvider(filament::Engine* engine, bool optimizeShaders) {
+    return new JitShaderProvider(engine, optimizeShaders);
 }

 } // namespace filament::gltfio
@@ -383,10 +383,8 @@ ViewerGui::ViewerGui(filament::Engine* engine, filament::Scene* scene, filament:
    mSettings.view.vsmShadowOptions.anisotropy = 0;
    mSettings.view.dithering = Dithering::TEMPORAL;
    mSettings.view.antiAliasing = AntiAliasing::FXAA;
-//    mSettings.view.msaa = { .enabled = true, .sampleCount = 4 };
-    mSettings.view.msaa = { .enabled = false, .sampleCount = 4 };
-//    mSettings.view.ssao.enabled = true;
-    mSettings.view.ssao.enabled = false;
+    mSettings.view.msaa = { .enabled = true, .sampleCount = 4 };
+    mSettings.view.ssao.enabled = true;
    mSettings.view.bloom.enabled = true;

    DebugRegistry& debug = mEngine->getDebugRegistry();
@@ -19,6 +19,7 @@ set(RESOURCE_BINS)
 set(MATERIAL_SRCS
        materials/aiDefaultMat.mat
        materials/bakedColor.mat
+        materials/bakedColorWithMutableConstants.mat
        materials/bakedTexture.mat
        materials/aoPreview.mat
        materials/arrayTexture.mat
@@ -222,8 +223,6 @@ target_sources(gltf-demo-resources PRIVATE ${RESGEN_SOURCE})

 add_library(samples-common STATIC
    common/arguments.cpp
-    common/configuration.h
-    common/configuration.cpp
 )

 target_include_directories(samples-common PRIVATE
@@ -280,6 +279,7 @@ if (NOT ANDROID)
    add_demo(lightbulb)
    add_demo(material_sandbox)
    add_demo(multiple_windows)
+    add_demo(mutable_constants)
    # Tint does not support setting gl_PointSize, disable the relevant sample if using WebGPU
    # https://github.com/gpuweb/gpuweb/issues/1190
    if (NOT FILAMENT_SUPPORTS_WEBGPU)
@@ -14,8 +14,6 @@
 * limitations under the License.
 */

-#include "arguments.h"
-
 #include <filament/Engine.h>

 #include <utils/Path.h>
@@ -22,10 +22,8 @@
 #include <string>

 namespace samples {
-
-filament::Engine::Backend parseArgumentsForBackend(int argc, char* argv[]);
-filament::Engine::Backend parseArgumentsForBackend(const std::string& backend);
-std::string getBackendAPIArgumentsUsage();
-
-} // namespace samples
+    filament::Engine::Backend parseArgumentsForBackend(int argc, char* argv[]);
+    filament::Engine::Backend parseArgumentsForBackend(const std::string& backend);
+    std::string getBackendAPIArgumentsUsage();
+}
 #endif //TNT_SAMPLES_ARGUMENTS_H
@@ -1,65 +0,0 @@
-/*
- * Copyright (C) 2025 The Android Open Source Project
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- *      http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-#include "configuration.h"
-
-#include <filament/Engine.h>
-
-#include <utils/FixedCapacityVector.h>
-
-#include <unordered_map>
-
-namespace samples {
-
-namespace {
-
-using FilterVector = utils::FixedCapacityVector<char const*>;
-std::unordered_map<filament::Engine::Backend, FilterVector> gFilters;
-
-// TODO: This is not necessary once we can support these variants in webgpu.
-char const* WEBGPU_VARIANT_FILTERS[] = {
-    "skinning",
-    "stereo",
-};
-
-} // namespace anonymous
-
-FilterVector const& getJitMaterialVariantFilter(filament::Engine::Backend backend) {
-    if (auto itr = gFilters.find(backend); itr != gFilters.end()) {
-        return itr->second;
-    }
-    auto build = [&](char const** filters, size_t filterCount) -> FilterVector {
-        utils::FixedCapacityVector<char const*> filterStr;
-        filterStr.clear();
-        filterStr.reserve(filterCount);
-        for (size_t i = 0; i < filterCount; ++i) {
-            filterStr.push_back(filters[i]);
-        }
-        return filterStr;
-    };
-
-    switch (backend) {
-        case filament::Engine::Backend::WEBGPU:
-            gFilters[backend] = build(WEBGPU_VARIANT_FILTERS,
-                    sizeof(WEBGPU_VARIANT_FILTERS) / sizeof(char const*));
-            break;
-        default:
-            gFilters[backend] = build(nullptr, 0);
-    }
-    return gFilters[backend];
-}
-
-} // namespace samples
@@ -1,31 +0,0 @@
-/*
- * Copyright (C) 2025 The Android Open Source Project
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- *      http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-#ifndef TNT_SAMPLES_CONFIGURATION_H
-#define TNT_SAMPLES_CONFIGURATION_H
-
-#include <filament/Engine.h>
-
-#include <utils/FixedCapacityVector.h>
-
-namespace samples {
-
-utils::FixedCapacityVector<char const*> const& getJitMaterialVariantFilter(
-        filament::Engine::Backend backend);
-
-}
-
-#endif // TNT_SAMPLES_CONFIGURATION_H
@@ -15,7 +15,6 @@
 */

 #include "common/arguments.h"
-#include "common/configuration.h"

 #include <filamentapp/Config.h>
 #include <filamentapp/FilamentApp.h>
@@ -245,11 +244,8 @@ int main(int argc, char** argv) {
        app.names = new NameComponentManager(EntityManager::get());
        app.viewer = new ViewerGui(engine, scene, view);

-        app.materials = (app.materialSource == JITSHADER)
-                                ? createJitShaderProvider(engine, false /* optimize */,
-                                          samples::getJitMaterialVariantFilter(app.config.backend))
-                                : createUbershaderProvider(engine, UBERARCHIVE_DEFAULT_DATA,
-                                          UBERARCHIVE_DEFAULT_SIZE);
+        app.materials = (app.materialSource == JITSHADER) ? createJitShaderProvider(engine) :
+                createUbershaderProvider(engine, UBERARCHIVE_DEFAULT_DATA, UBERARCHIVE_DEFAULT_SIZE);

        app.loader = AssetLoader::create({engine, app.materials, app.names });
        if (filename.isEmpty()) {
@@ -15,7 +15,6 @@
 */

 #include "common/arguments.h"
-#include "common/configuration.h"

 #include <filamentapp/Config.h>
 #include <filamentapp/FilamentApp.h>
@@ -779,13 +778,11 @@ int main(int argc, char** argv) {
            }
        }

-        app.materials = (app.materialSource == JITSHADER)
-                                ? createJitShaderProvider(engine, OPTIMIZE_MATERIALS,
-                                          samples::getJitMaterialVariantFilter(app.config.backend))
-                                : createUbershaderProvider(engine, UBERARCHIVE_DEFAULT_DATA,
-                                          UBERARCHIVE_DEFAULT_SIZE);
+        app.materials = (app.materialSource == JITSHADER) ?
+                createJitShaderProvider(engine, OPTIMIZE_MATERIALS) :
+                createUbershaderProvider(engine, UBERARCHIVE_DEFAULT_DATA, UBERARCHIVE_DEFAULT_SIZE);

-        app.assetLoader = AssetLoader::create({ engine, app.materials, app.names });
+        app.assetLoader = AssetLoader::create({engine, app.materials, app.names });
        app.mainCamera = &view->getCamera();
        if (filename.isEmpty()) {
            app.asset = app.assetLoader->createAsset(
@@ -0,0 +1,44 @@
+material {
+    name : BakedColorWithMutableConstants,
+    requires : [
+        color
+    ],
+    constants : [
+        {
+            name: swizzle,
+            type: bool,
+            mutable: true,
+            default: true,
+        },
+        {
+            name: darken,
+            type: bool,
+            mutable: true,
+        },
+        {
+            // In generated shader sources, we first specify immutable
+            // specialization constants, then mutable specialization constants.
+            // By including this dummy constant, we offset the ID of the first
+            // mutable constant by one. This is done here purely for testing
+            // purposes.
+            name: dummy,
+            type: bool,
+        }
+    ],
+    shadingModel : unlit,
+    culling : none,
+}
+
+fragment {
+    void material(inout MaterialInputs material) {
+        prepareMaterial(material);
+        if (materialConstants_swizzle) {
+            material.baseColor = getColor().gbra;
+        } else {
+            material.baseColor = getColor();
+        }
+        if (materialConstants_darken) {
+            material.baseColor.rgb *= 0.5;
+        }
+    }
+}
@@ -0,0 +1,214 @@
+/*
+ * Copyright (C) 2025 The Android Open Source Project
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *      http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include <filament/Camera.h>
+#include <filament/Engine.h>
+#include <filament/IndexBuffer.h>
+#include <filament/Material.h>
+#include <filament/MaterialInstance.h>
+#include <filament/RenderableManager.h>
+#include <filament/Scene.h>
+#include <filament/Skybox.h>
+#include <filament/TransformManager.h>
+#include <filament/VertexBuffer.h>
+#include <filament/View.h>
+
+#include <imgui.h>
+
+#include <utils/EntityManager.h>
+
+#include <filamentapp/Config.h>
+#include <filamentapp/FilamentApp.h>
+
+#include <getopt/getopt.h>
+
+#include <cmath>
+#include <iostream>
+
+
+#include "generated/resources/resources.h"
+
+using namespace filament;
+using utils::Entity;
+using utils::EntityManager;
+
+struct App {
+    Config config;
+    VertexBuffer* vb;
+    IndexBuffer* ib;
+    Material* mat;
+    Camera* cam;
+    Entity camera;
+    Skybox* skybox;
+    Entity renderable;
+};
+
+struct Vertex {
+    filament::math::float2 position;
+    uint32_t color;
+};
+
+static const Vertex TRIANGLE_VERTICES[3] = {
+    {{1, 0}, 0xffff0000u},
+    {{cos(M_PI * 2 / 3), sin(M_PI * 2 / 3)}, 0xff00ff00u},
+    {{cos(M_PI * 4 / 3), sin(M_PI * 4 / 3)}, 0xff0000ffu},
+};
+
+static constexpr uint16_t TRIANGLE_INDICES[3] = { 0, 1, 2 };
+
+static void printUsage(char* name) {
+    std::string exec_name(utils::Path(name).getName());
+    std::string usage(
+            "MUTABLE_CONSTANTS renders a spinning colored triangle\n"
+            "Usage:\n"
+            "    MUTABLE_CONSTANTS [options]\n"
+            "Options:\n"
+            "   --help, -h\n"
+            "       Prints this message\n\n"
+            "   --api, -a\n"
+            "       Specify the backend API: opengl, vulkan, metal, or webgpu\n"
+            "       (note: webgpu is a no-op for now, printing backend\n"
+            "        component info if FILAMENT_BACKEND_DEBUG_FLAG, \n"
+            "        set at build time, includes the \n"
+            "        FWGPU_PRINT_SYSTEM bit flag 0x2)\n"
+    );
+    const std::string from("MUTABLE_CONSTANTS");
+    for (size_t pos = usage.find(from); pos != std::string::npos; pos = usage.find(from, pos)) {
+        usage.replace(pos, from.length(), exec_name);
+    }
+    std::cout << usage;
+}
+
+static int handleCommandLineArguments(int argc, char* argv[], App* app) {
+    static constexpr const char* OPTSTR = "ha:";
+    static const struct option OPTIONS[] = {
+            { "help", no_argument,       nullptr, 'h' },
+            { "api",  required_argument, nullptr, 'a' },
+            { nullptr, 0,                nullptr, 0 }
+    };
+    int opt;
+    int option_index = 0;
+    while ((opt = getopt_long(argc, argv, OPTSTR, OPTIONS, &option_index)) >= 0) {
+        std::string arg(optarg ? optarg : "");
+        switch (opt) {
+            default:
+            case 'h':
+                printUsage(argv[0]);
+                exit(0);
+            case 'a':
+                if (arg == "opengl") {
+                    app->config.backend = Engine::Backend::OPENGL;
+                } else if (arg == "vulkan") {
+                    app->config.backend = Engine::Backend::VULKAN;
+                } else if (arg == "metal") {
+                    app->config.backend = Engine::Backend::METAL;
+                } else if (arg == "webgpu") {
+                    app->config.backend = Engine::Backend::WEBGPU;
+                } else {
+                    std::cerr << "Unrecognized backend. Must be "
+                                 "'opengl'|'vulkan'|'metal'|'webgpu'.\n";
+                    exit(1);
+                }
+                break;
+        }
+    }
+    return optind;
+}
+
+int main(int argc, char** argv) {
+    App app{};
+    app.config.title = "mutable_constants";
+    handleCommandLineArguments(argc, argv, &app);
+
+    auto setup = [&app](Engine* engine, View* view, Scene* scene) {
+        app.skybox = Skybox::Builder().color({0.1, 0.125, 0.25, 1.0}).build(*engine);
+        scene->setSkybox(app.skybox);
+        view->setPostProcessingEnabled(false);
+        static_assert(sizeof(Vertex) == 12, "Strange vertex size.");
+        app.vb = VertexBuffer::Builder()
+                .vertexCount(3)
+                .bufferCount(1)
+                .attribute(VertexAttribute::POSITION, 0, VertexBuffer::AttributeType::FLOAT2, 0, 12)
+                .attribute(VertexAttribute::COLOR, 0, VertexBuffer::AttributeType::UBYTE4, 8, 12)
+                .normalized(VertexAttribute::COLOR)
+                .build(*engine);
+        app.vb->setBufferAt(*engine, 0,
+                VertexBuffer::BufferDescriptor(TRIANGLE_VERTICES, 36, nullptr));
+        app.ib = IndexBuffer::Builder()
+                .indexCount(3)
+                .bufferType(IndexBuffer::IndexType::USHORT)
+                .build(*engine);
+        app.ib->setBuffer(*engine,
+                IndexBuffer::BufferDescriptor(TRIANGLE_INDICES, 6, nullptr));
+        app.mat = Material::Builder().package(RESOURCES_BAKEDCOLORWITHMUTABLECONSTANTS_DATA,
+                RESOURCES_BAKEDCOLORWITHMUTABLECONSTANTS_SIZE)
+                .build(*engine);
+        app.renderable = EntityManager::get().create();
+        RenderableManager::Builder(1)
+                .boundingBox({{ -1, -1, -1 }, { 1, 1, 1 }})
+                .material(0, app.mat->getDefaultInstance())
+                .geometry(0, RenderableManager::PrimitiveType::TRIANGLES, app.vb, app.ib, 0, 3)
+                .culling(false)
+                .receiveShadows(false)
+                .castShadows(false)
+                .build(*engine, app.renderable);
+        scene->addEntity(app.renderable);
+        app.camera = utils::EntityManager::get().create();
+        app.cam = engine->createCamera(app.camera);
+        view->setCamera(app.cam);
+    };
+
+    auto cleanup = [&app](Engine* engine, View*, Scene*) {
+        engine->destroy(app.skybox);
+        engine->destroy(app.renderable);
+        engine->destroy(app.mat);
+        engine->destroy(app.vb);
+        engine->destroy(app.ib);
+        engine->destroyCameraComponent(app.camera);
+        utils::EntityManager::get().destroy(app.camera);
+    };
+
+    auto gui = [&app](Engine* engine, View* view) {
+        MaterialInstance* mi = app.mat->getDefaultInstance();
+        bool swizzle = mi->getConstant<bool>("swizzle");
+        bool darken = mi->getConstant<bool>("darken");
+        ImGui::Checkbox("Swizzle", &swizzle);
+        ImGui::Checkbox("Darken", &darken);
+        if (swizzle != mi->getConstant<bool>("swizzle")) {
+            mi->setConstant("swizzle", swizzle);
+        }
+        if (darken != mi->getConstant<bool>("darken")) {
+            mi->setConstant("darken", darken);
+        }
+    };
+
+    FilamentApp::get().animate([&app](Engine* engine, View* view, double now) {
+        constexpr float ZOOM = 1.5f;
+        const uint32_t w = view->getViewport().width;
+        const uint32_t h = view->getViewport().height;
+        const float aspect = (float) w / h;
+        app.cam->setProjection(Camera::Projection::ORTHO,
+            -aspect * ZOOM, aspect * ZOOM,
+            -ZOOM, ZOOM, 0, 1);
+        auto& tcm = engine->getTransformManager();
+        tcm.setTransform(tcm.getInstance(app.renderable),
+                filament::math::mat4f::rotation(now, filament::math::float3{ 0, 0, 1 }));
+    });
+
+    FilamentApp::get().run(app.config, setup, cleanup, gui);
+
+    return 0;
+}
@@ -13,9 +13,7 @@
 * See the License for the specific language governing permissions and
 * limitations under the License.
 */
-
 #include "common/arguments.h"
-#include "common/configuration.h"

 #include <filamentapp/Config.h>
 #include <filamentapp/FilamentApp.h>
@@ -74,7 +72,6 @@ struct App {
    gltfio::TextureProvider* stbDecoder = nullptr;
    gltfio::TextureProvider* ktxDecoder = nullptr;
    FilamentInstance* instance;
-    bool enableMSAA = false;
 };

 static const char* DEFAULT_IBL = "assets/ibl/lightroom_14b";
@@ -90,9 +87,7 @@ static void printUsage(char* name) {
                      "       Prints this message\n\n"
                      "API_USAGE"
                      "   --ibl=<path to cmgen IBL>, -i <path>\n"
-                      "       Override the built-in IBL\n\n"
-                      "   --msaa, -m\n"
-                      "       setMultiSampleAntiAliasingOptions\n\n");
+                      "       Override the built-in IBL\n\n");
    const std::string from("SHOWCASE");
    for (size_t pos = usage.find(from); pos != std::string::npos; pos = usage.find(from, pos)) {
        usage.replace(pos, from.length(), exec_name);
@@ -106,11 +101,9 @@ static void printUsage(char* name) {
 }

 static int handleCommandLineArguments(int argc, char* argv[], App* app) {
-    static constexpr const char* OPTSTR = "hma:i:";
+    static constexpr const char* OPTSTR = "ha:i:un:m:";
    static const struct option OPTIONS[] = { { "help", no_argument, nullptr, 'h' },
-        { "api", required_argument, nullptr, 'a' },
-        { "ibl", required_argument, nullptr, 'i' },
-        { "msaa", no_argument, nullptr, 'm' },
+        { "api", required_argument, nullptr, 'a' }, { "ibl", required_argument, nullptr, 'i' },
        { nullptr, 0, nullptr, 0 } };
    int opt;
    int option_index = 0;
@@ -129,9 +122,6 @@ static int handleCommandLineArguments(int argc, char* argv[], App* app) {
            case 'i':
                app->config.iblDirectory = arg;
                break;
-            case 'm':
-                app->enableMSAA = true;
-                break;
        }
    }
    return optind;
@@ -215,19 +205,13 @@ int main(int argc, char** argv) {
    };

    auto setup = [&](Engine* engine, View* view, Scene* scene) {
-        if (app.enableMSAA) {
-            // TODO Investigate why
-            // Enabling MSAA fixes the "transmission" sample but breaks the others
-            view->setMultiSampleAntiAliasingOptions({ .enabled = true});
-        }
+        view->setMultiSampleAntiAliasingOptions({ .enabled = true});
        app.engine = engine;
        app.names = new NameComponentManager(EntityManager::get());
        app.viewer = new ViewerGui(engine, scene, view);

-
        app.materials = (app.materialSource == JITSHADER)
-                                ? createJitShaderProvider(engine, false /* optimize */,
-                                          samples::getJitMaterialVariantFilter(app.config.backend))
+                                ? createJitShaderProvider(engine)
                                : createUbershaderProvider(engine, UBERARCHIVE_DEFAULT_DATA,
                                          UBERARCHIVE_DEFAULT_SIZE);

@@ -1,3 +1,62 @@
+# 3.3.0 (2025-05-12)
+
+Additions to the library API:
+
+- Added function `vmaImportVulkanFunctionsFromVolk`, useful for loading pointers to Vulkan functions with [volk library](https://github.com/zeux/volk).
+
+Other changes:
+
+- Added macro `VMA_DEBUG_DONT_EXCEED_HEAP_SIZE_WITH_ALLOCATION_SIZE` with default value 1.
+- Changed macro `VMA_DEBUG_DONT_EXCEED_MAX_MEMORY_ALLOCATION_COUNT` default value from 0 to 1.
+- Added documentation chapter "Frequenty asked questions".
+- Other fixes and improvements, including compatibility with various platforms and compilers.
+
+# 3.2.1 (2025-02-05)
+
+Changes:
+
+- Fixed an assert in `vmaCreateAllocator` function incorrectly failing when Vulkan version 1.4 is used (#457).
+- Fix for importing function `vkGetPhysicalDeviceMemoryProperties2` / `vkGetPhysicalDeviceMemoryProperties2KHR` when `VMA_DYNAMIC_VULKAN_FUNCTIONS` macro is enabled (#410).
+- Other minor fixes and improvements...
+
+# 3.2.0 (2024-12-30)
+
+Additions to the library API:
+
+- Added support for Vulkan 1.4.
+- Added support for VK_KHR_external_memory_win32 extension - `VMA_ALLOCATOR_CREATE_KHR_EXTERNAL_MEMORY_WIN32_BIT` flag, `vmaGetMemoryWin32Handle` function, and a whole new documentation chapter about it (#442).
+
+Other changes:
+
+- Fixed thread safety issue (#451).
+- Many other bug fixes and improvements in the library code, documentation, sample app, Cmake script, mostly to improve compatibility with various compilers and GPUs.
+
+# 3.1.0 (2024-05-27)
+
+This release gathers fixes and improvements made during many months of continuous development on the main branch, mostly based on issues and pull requests on GitHub.
+
+Additions to the library API:
+
+- Added convenience functions `vmaCopyMemoryToAllocation`, `vmaCopyAllocationToMemory`.
+- Added functions `vmaCreateAliasingBuffer2`, `vmaCreateAliasingImage2` that offer creating a buffer/image in an existing allocation with additional `allocationLocalOffset`.
+- Added function `vmaGetAllocationInfo2`, structure `VmaAllocationInfo2` that return additional information about an allocation, useful for interop with other APIs (#383, #340).
+- Added callback `VmaDefragmentationInfo::pfnBreakCallback` that allows breaking long execution of `vmaBeginDefragmentation`.
+  Also added `PFN_vmaCheckDefragmentationBreakFunction`, `VmaDefragmentationInfo::pBreakCallbackUserData`.
+- Added support for VK_KHR_maintenance4 extension - `VMA_ALLOCATOR_CREATE_KHR_MAINTENANCE4_BIT` flag (#397).
+- Added support for VK_KHR_maintenance5 extension - `VMA_ALLOCATOR_CREATE_KHR_MAINTENANCE5_BIT` flag (#411).
+
+Other changes:
+
+- Changes in debug and configuration macros:
+  - Split macros into separate `VMA_DEBUG_LOG` and `VMA_DEBUG_LOG_FORMAT` (#297).
+  - Added macros `VMA_ASSERT_LEAK`, `VMA_LEAK_LOG_FORMAT` separate from normal `VMA_ASSERT`, `VMA_DEBUG_LOG_FORMAT` (#379, #385).
+  - Added macro `VMA_EXTENDS_VK_STRUCT` (#347).
+- Countless bug fixes and improvements in the code and documentation, mostly to improve compatibility with various compilers and GPUs, including:
+  - Fixed missing `#include` that resulted in compilation error about `snprintf` not declared on some compilers (#312).
+  - Fixed main memory type selection algorithm for GPUs that have no `HOST_CACHED` memory type, like Raspberry Pi (#362).
+- Major changes in Cmake script.
+- Fixes in GpuMemDumpVis.py script.
+
 # 3.0.1 (2022-05-26)

 - Fixes in defragmentation algorithm.
@@ -1,55 +1,71 @@
-cmake_minimum_required(VERSION 3.9)
+#
+# Copyright (c) 2017-2025 Advanced Micro Devices, Inc. All rights reserved.
+#
+# Permission is hereby granted, free of charge, to any person obtaining a copy
+# of this software and associated documentation files (the "Software"), to deal
+# in the Software without restriction, including without limitation the rights
+# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+# copies of the Software, and to permit persons to whom the Software is
+# furnished to do so, subject to the following conditions:
+#
+# The above copyright notice and this permission notice shall be included in
+# all copies or substantial portions of the Software.
+#
+# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+# THE SOFTWARE.
+#

-project(VulkanMemoryAllocator)
+cmake_minimum_required(VERSION 3.15...3.26)

-find_package(Vulkan REQUIRED)
-include_directories(${Vulkan_INCLUDE_DIR})
+project(VMA VERSION 3.3.0 LANGUAGES CXX)

-# VulkanMemoryAllocator contains an sample application which is not built by default
-option(VMA_BUILD_SAMPLE "Build VulkanMemoryAllocator sample application" OFF)
-option(VMA_BUILD_SAMPLE_SHADERS "Build VulkanMemoryAllocator sample application's shaders" OFF)
+add_library(VulkanMemoryAllocator INTERFACE)
+add_library(GPUOpen::VulkanMemoryAllocator ALIAS VulkanMemoryAllocator)

-message(STATUS "VMA_BUILD_SAMPLE = ${VMA_BUILD_SAMPLE}")
-message(STATUS "VMA_BUILD_SAMPLE_SHADERS = ${VMA_BUILD_SAMPLE_SHADERS}")
+target_include_directories(VulkanMemoryAllocator INTERFACE $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/include>)

-option(VMA_STATIC_VULKAN_FUNCTIONS "Link statically with Vulkan API" ON)
-option(VMA_DYNAMIC_VULKAN_FUNCTIONS "Fetch pointers to Vulkan functions internally (no static linking)" OFF)
-option(VMA_DEBUG_ALWAYS_DEDICATED_MEMORY "Every allocation will have its own memory block" OFF)
-option(VMA_DEBUG_INITIALIZE_ALLOCATIONS "Automatically fill new allocations and destroyed allocations with some bit pattern" OFF)
-option(VMA_DEBUG_GLOBAL_MUTEX "Enable single mutex protecting all entry calls to the library" OFF)
-option(VMA_DEBUG_DONT_EXCEED_MAX_MEMORY_ALLOCATION_COUNT "Never exceed VkPhysicalDeviceLimits::maxMemoryAllocationCount and return error" OFF)
+if (CMAKE_VERSION VERSION_LESS "3.21")
+    # https://cmake.org/cmake/help/latest/variable/PROJECT_IS_TOP_LEVEL.html
+    string(COMPARE EQUAL ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_SOURCE_DIR} PROJECT_IS_TOP_LEVEL)
+endif()

-message(STATUS "VMA_STATIC_VULKAN_FUNCTIONS = ${VMA_STATIC_VULKAN_FUNCTIONS}")
-message(STATUS "VMA_DYNAMIC_VULKAN_FUNCTIONS = ${VMA_DYNAMIC_VULKAN_FUNCTIONS}")
-message(STATUS "VMA_DEBUG_ALWAYS_DEDICATED_MEMORY = ${VMA_DEBUG_ALWAYS_DEDICATED_MEMORY}")
-message(STATUS "VMA_DEBUG_INITIALIZE_ALLOCATIONS = ${VMA_DEBUG_INITIALIZE_ALLOCATIONS}")
-message(STATUS "VMA_DEBUG_GLOBAL_MUTEX = ${VMA_DEBUG_GLOBAL_MUTEX}")
-message(STATUS "VMA_DEBUG_DONT_EXCEED_MAX_MEMORY_ALLOCATION_COUNT = ${VMA_DEBUG_DONT_EXCEED_MAX_MEMORY_ALLOCATION_COUNT}")
+option(VMA_ENABLE_INSTALL "Install VulkanMemoryAllocator" ${PROJECT_IS_TOP_LEVEL})
+if (VMA_ENABLE_INSTALL)
+    include(GNUInstallDirs)
+    include(CMakePackageConfigHelpers)

-if(VMA_BUILD_SAMPLE)
-    set(VMA_BUILD_SAMPLE_SHADERS ON)
-endif(VMA_BUILD_SAMPLE)
+    install(DIRECTORY "${CMAKE_CURRENT_SOURCE_DIR}/include/" DESTINATION ${CMAKE_INSTALL_INCLUDEDIR})
+    install(TARGETS VulkanMemoryAllocator EXPORT VulkanMemoryAllocatorConfig INCLUDES DESTINATION ${CMAKE_INSTALL_INCLUDEDIR})
+    install(EXPORT VulkanMemoryAllocatorConfig NAMESPACE "GPUOpen::" DESTINATION "share/cmake/VulkanMemoryAllocator")

-find_package(Doxygen)
-option(BUILD_DOCUMENTATION "Create and install the HTML based API documentation (requires Doxygen)" OFF)
+    write_basic_package_version_file(VulkanMemoryAllocatorConfigVersion.cmake COMPATIBILITY SameMajorVersion ARCH_INDEPENDENT)
+    install(FILES "${CMAKE_CURRENT_BINARY_DIR}/VulkanMemoryAllocatorConfigVersion.cmake" DESTINATION "share/cmake/VulkanMemoryAllocator")

-if(BUILD_DOCUMENTATION)
-    if(DOXYGEN_FOUND)
+    option(VMA_BUILD_DOCUMENTATION "Create and install the HTML based API documentation")
+    if(VMA_BUILD_DOCUMENTATION)
+        find_package(Doxygen REQUIRED)
        # set input and output files
        set(DOXYGEN_IN ${CMAKE_CURRENT_SOURCE_DIR}/Doxyfile)
        set(DOXYGEN_OUT ${CMAKE_CURRENT_BINARY_DIR}/Doxyfile)
-
        # request to configure the file
        configure_file(${DOXYGEN_IN} ${DOXYGEN_OUT} @ONLY)
        # note the option ALL which allows to build the docs together with the application
-        add_custom_target( doc_doxygen ALL
+        add_custom_target(doc_doxygen ALL
            COMMAND ${DOXYGEN_EXECUTABLE} ${DOXYGEN_OUT}
            WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}
            COMMENT "Generating API documentation with Doxygen"
-            VERBATIM )
-    else()
-        message("Doxygen need to be installed to generate the doxygen documentation")
+            VERBATIM
+        )
+        install(DIRECTORY docs/ DESTINATION "${CMAKE_INSTALL_DATADIR}/doc/VulkanMemoryAllocator" PATTERN ".nojekyll" EXCLUDE)
+    endif()
+
+    option(VMA_BUILD_SAMPLES "Build samples")
+    if (VMA_BUILD_SAMPLES)
+        add_subdirectory(src)
    endif()
 endif()
-
-add_subdirectory(src)
@@ -1,4 +1,4 @@
-# Doxyfile 1.9.2
+# Doxyfile 1.10.0

 # This file describes the settings to be used by the documentation system
 # doxygen (www.doxygen.org) for a project.
@@ -12,6 +12,16 @@
 # For lists, items can also be appended using:
 # TAG += value [value, ...]
 # Values that contain spaces should be placed between quotes (\" \").
+#
+# Note:
+#
+# Use doxygen to compare the used configuration file with the template
+# configuration file:
+# doxygen -x [configFile]
+# Use doxygen to compare the used configuration file with the template
+# configuration file without replacing the environment variables or CMake type
+# replacement variables:
+# doxygen -x_noenv [configFile]

 #---------------------------------------------------------------------------
 # Project related configuration options
@@ -53,23 +63,41 @@ PROJECT_BRIEF          =

 PROJECT_LOGO           =

+# With the PROJECT_ICON tag one can specify an icon that is included in the tabs
+# when the HTML document is shown. Doxygen will copy the logo to the output
+# directory.
+
+PROJECT_ICON           =
+
 # The OUTPUT_DIRECTORY tag is used to specify the (relative or absolute) path
 # into which the generated documentation will be written. If a relative path is
 # entered, it will be relative to the location where doxygen was started. If
 # left blank the current directory will be used.

-OUTPUT_DIRECTORY       = "@CMAKE_SOURCE_DIR@/docs"
+OUTPUT_DIRECTORY       = @CMAKE_SOURCE_DIR@/docs

-# If the CREATE_SUBDIRS tag is set to YES then doxygen will create 4096 sub-
-# directories (in 2 levels) under the output directory of each output format and
-# will distribute the generated files over these directories. Enabling this
+# If the CREATE_SUBDIRS tag is set to YES then doxygen will create up to 4096
+# sub-directories (in 2 levels) under the output directory of each output format
+# and will distribute the generated files over these directories. Enabling this
 # option can be useful when feeding doxygen a huge amount of source files, where
 # putting all generated files in the same directory would otherwise causes
-# performance problems for the file system.
+# performance problems for the file system. Adapt CREATE_SUBDIRS_LEVEL to
+# control the number of sub-directories.
 # The default value is: NO.

 CREATE_SUBDIRS         = NO

+# Controls the number of sub-directories that will be created when
+# CREATE_SUBDIRS tag is set to YES. Level 0 represents 16 directories, and every
+# level increment doubles the number of directories, resulting in 4096
+# directories at level 8 which is the default and also the maximum value. The
+# sub-directories are organized in 2 levels, the first level always has a fixed
+# number of 16 directories.
+# Minimum value: 0, maximum value: 8, default value: 8.
+# This tag requires that the tag CREATE_SUBDIRS is set to YES.
+
+CREATE_SUBDIRS_LEVEL   = 8
+
 # If the ALLOW_UNICODE_NAMES tag is set to YES, doxygen will allow non-ASCII
 # characters to appear in the names of generated files. If set to NO, non-ASCII
 # characters will be escaped, for example _xE3_x81_x84 will be used for Unicode
@@ -81,14 +109,14 @@ ALLOW_UNICODE_NAMES    = NO
 # The OUTPUT_LANGUAGE tag is used to specify the language in which all
 # documentation generated by doxygen is written. Doxygen will use this
 # information to generate all constant output in the proper language.
-# Possible values are: Afrikaans, Arabic, Armenian, Brazilian, Catalan, Chinese,
-# Chinese-Traditional, Croatian, Czech, Danish, Dutch, English (United States),
-# Esperanto, Farsi (Persian), Finnish, French, German, Greek, Hungarian,
-# Indonesian, Italian, Japanese, Japanese-en (Japanese with English messages),
-# Korean, Korean-en (Korean with English messages), Latvian, Lithuanian,
-# Macedonian, Norwegian, Persian (Farsi), Polish, Portuguese, Romanian, Russian,
-# Serbian, Serbian-Cyrillic, Slovak, Slovene, Spanish, Swedish, Turkish,
-# Ukrainian and Vietnamese.
+# Possible values are: Afrikaans, Arabic, Armenian, Brazilian, Bulgarian,
+# Catalan, Chinese, Chinese-Traditional, Croatian, Czech, Danish, Dutch, English
+# (United States), Esperanto, Farsi (Persian), Finnish, French, German, Greek,
+# Hindi, Hungarian, Indonesian, Italian, Japanese, Japanese-en (Japanese with
+# English messages), Korean, Korean-en (Korean with English messages), Latvian,
+# Lithuanian, Macedonian, Norwegian, Persian (Farsi), Polish, Portuguese,
+# Romanian, Russian, Serbian, Serbian-Cyrillic, Slovak, Slovene, Spanish,
+# Swedish, Turkish, Ukrainian and Vietnamese.
 # The default value is: English.

 OUTPUT_LANGUAGE        = English
@@ -341,6 +369,17 @@ MARKDOWN_SUPPORT       = YES

 TOC_INCLUDE_HEADINGS   = 0

+# The MARKDOWN_ID_STYLE tag can be used to specify the algorithm used to
+# generate identifiers for the Markdown headings. Note: Every identifier is
+# unique.
+# Possible values are: DOXYGEN use a fixed 'autotoc_md' string followed by a
+# sequence number starting at 0 and GITHUB use the lower case version of title
+# with any whitespace replaced by '-' and punctuation characters removed.
+# The default value is: DOXYGEN.
+# This tag requires that the tag MARKDOWN_SUPPORT is set to YES.
+
+MARKDOWN_ID_STYLE      = DOXYGEN
+
 # When enabled doxygen tries to link words that correspond to documented
 # classes, or namespaces to their corresponding documentation. Such a link can
 # be prevented in individual cases by putting a % sign in front of the word or
@@ -452,7 +491,7 @@ TYPEDEF_HIDES_STRUCT   = NO

 LOOKUP_CACHE_SIZE      = 0

-# The NUM_PROC_THREADS specifies the number threads doxygen is allowed to use
+# The NUM_PROC_THREADS specifies the number of threads doxygen is allowed to use
 # during processing. When set to 0 doxygen will based this on the number of
 # cores available in the system. You can set it explicitly to a value larger
 # than 0 to get more control over the balance between CPU load and processing
@@ -465,6 +504,14 @@ LOOKUP_CACHE_SIZE      = 0

 NUM_PROC_THREADS       = 1

+# If the TIMESTAMP tag is set different from NO then each generated page will
+# contain the date or date and time when the page was generated. Setting this to
+# NO can help when comparing the output of multiple runs.
+# Possible values are: YES, NO, DATETIME and DATE.
+# The default value is: NO.
+
+TIMESTAMP              = NO
+
 #---------------------------------------------------------------------------
 # Build related configuration options
 #---------------------------------------------------------------------------
@@ -546,7 +593,8 @@ HIDE_UNDOC_MEMBERS     = NO
 # If the HIDE_UNDOC_CLASSES tag is set to YES, doxygen will hide all
 # undocumented classes that are normally visible in the class hierarchy. If set
 # to NO, these classes will be included in the various overviews. This option
-# has no effect if EXTRACT_ALL is enabled.
+# will also hide undocumented C++ concepts if enabled. This option has no effect
+# if EXTRACT_ALL is enabled.
 # The default value is: NO.

 HIDE_UNDOC_CLASSES     = NO
@@ -577,14 +625,15 @@ INTERNAL_DOCS          = NO
 # filesystem is case sensitive (i.e. it supports files in the same directory
 # whose names only differ in casing), the option must be set to YES to properly
 # deal with such files in case they appear in the input. For filesystems that
-# are not case sensitive the option should be be set to NO to properly deal with
+# are not case sensitive the option should be set to NO to properly deal with
 # output files written for symbols that only differ in casing, such as for two
 # classes, one named CLASS and the other named Class, and to also support
 # references to files without having to specify the exact matching casing. On
 # Windows (including Cygwin) and MacOS, users should typically set this option
 # to NO, whereas on Linux or other Unix flavors it should typically be set to
 # YES.
-# The default value is: system dependent.
+# Possible values are: SYSTEM, NO and YES.
+# The default value is: SYSTEM.

 CASE_SENSE_NAMES       = NO

@@ -836,11 +885,26 @@ WARN_IF_INCOMPLETE_DOC = YES

 WARN_NO_PARAMDOC       = NO

+# If WARN_IF_UNDOC_ENUM_VAL option is set to YES, doxygen will warn about
+# undocumented enumeration values. If set to NO, doxygen will accept
+# undocumented enumeration values. If EXTRACT_ALL is set to YES then this flag
+# will automatically be disabled.
+# The default value is: NO.
+
+WARN_IF_UNDOC_ENUM_VAL = NO
+
 # If the WARN_AS_ERROR tag is set to YES then doxygen will immediately stop when
 # a warning is encountered. If the WARN_AS_ERROR tag is set to FAIL_ON_WARNINGS
 # then doxygen will continue running as if WARN_AS_ERROR tag is set to NO, but
 # at the end of the doxygen process doxygen will return with a non-zero status.
-# Possible values are: NO, YES and FAIL_ON_WARNINGS.
+# If the WARN_AS_ERROR tag is set to FAIL_ON_WARNINGS_PRINT then doxygen behaves
+# like FAIL_ON_WARNINGS but in case no WARN_LOGFILE is defined doxygen will not
+# write the warning messages in between other messages but write them at the end
+# of a run, in case a WARN_LOGFILE is defined the warning messages will be
+# besides being in the defined file also be shown at the end of a run, unless
+# the WARN_LOGFILE is defined as - i.e. standard output (stdout) in that case
+# the behavior will remain as with the setting FAIL_ON_WARNINGS.
+# Possible values are: NO, YES, FAIL_ON_WARNINGS and FAIL_ON_WARNINGS_PRINT.
 # The default value is: NO.

 WARN_AS_ERROR          = NO
@@ -851,13 +915,27 @@ WARN_AS_ERROR          = NO
 # and the warning text. Optionally the format may contain $version, which will
 # be replaced by the version of the file (if it could be obtained via
 # FILE_VERSION_FILTER)
+# See also: WARN_LINE_FORMAT
 # The default value is: $file:$line: $text.

 WARN_FORMAT            = "$file:$line: $text"

+# In the $text part of the WARN_FORMAT command it is possible that a reference
+# to a more specific place is given. To make it easier to jump to this place
+# (outside of doxygen) the user can define a custom "cut" / "paste" string.
+# Example:
+# WARN_LINE_FORMAT = "'vi $file +$line'"
+# See also: WARN_FORMAT
+# The default value is: at line $line of file $file.
+
+WARN_LINE_FORMAT       = "at line $line of file $file"
+
 # The WARN_LOGFILE tag can be used to specify a file to which warning and error
 # messages should be written. If left blank the output is written to standard
-# error (stderr).
+# error (stderr). In case the file specified cannot be opened for writing the
+# warning and error messages are written to standard error. When as file - is
+# specified the warning and error messages are written to standard output
+# (stdout).

 WARN_LOGFILE           =

@@ -878,10 +956,21 @@ INPUT                  = "@CMAKE_SOURCE_DIR@/include/vk_mem_alloc.h"
 # libiconv (or the iconv built into libc) for the transcoding. See the libiconv
 # documentation (see:
 # https://www.gnu.org/software/libiconv/) for the list of possible encodings.
+# See also: INPUT_FILE_ENCODING
 # The default value is: UTF-8.

 INPUT_ENCODING         = UTF-8

+# This tag can be used to specify the character encoding of the source files
+# that doxygen parses The INPUT_FILE_ENCODING tag can be used to specify
+# character encoding on a per file pattern basis. Doxygen will compare the file
+# name with each pattern and apply the encoding instead of the default
+# INPUT_ENCODING) if there is a match. The character encodings are a list of the
+# form: pattern=encoding (like *.php=ISO-8859-1). See cfg_input_encoding
+# "INPUT_ENCODING" for further information on supported encodings.
+
+INPUT_FILE_ENCODING    =
+
 # If the value of the INPUT tag contains directories, you can use the
 # FILE_PATTERNS tag to specify one or more wildcard patterns (like *.cpp and
 # *.h) to filter out the source-files in the directories.
@@ -893,12 +982,12 @@ INPUT_ENCODING         = UTF-8
 # Note the list of default checked file patterns might differ from the list of
 # default file extension mappings.
 #
-# If left blank the following patterns are tested:*.c, *.cc, *.cxx, *.cpp,
-# *.c++, *.java, *.ii, *.ixx, *.ipp, *.i++, *.inl, *.idl, *.ddl, *.odl, *.h,
-# *.hh, *.hxx, *.hpp, *.h++, *.l, *.cs, *.d, *.php, *.php4, *.php5, *.phtml,
-# *.inc, *.m, *.markdown, *.md, *.mm, *.dox (to be provided as doxygen C
-# comment), *.py, *.pyw, *.f90, *.f95, *.f03, *.f08, *.f18, *.f, *.for, *.vhd,
-# *.vhdl, *.ucf, *.qsf and *.ice.
+# If left blank the following patterns are tested:*.c, *.cc, *.cxx, *.cxxm,
+# *.cpp, *.cppm, *.ccm, *.c++, *.c++m, *.java, *.ii, *.ixx, *.ipp, *.i++, *.inl,
+# *.idl, *.ddl, *.odl, *.h, *.hh, *.hxx, *.hpp, *.h++, *.ixx, *.l, *.cs, *.d,
+# *.php, *.php4, *.php5, *.phtml, *.inc, *.m, *.markdown, *.md, *.mm, *.dox (to
+# be provided as doxygen C comment), *.py, *.pyw, *.f90, *.f95, *.f03, *.f08,
+# *.f18, *.f, *.for, *.vhd, *.vhdl, *.ucf, *.qsf and *.ice.

 FILE_PATTERNS          = *.c \
                         *.cc \
@@ -980,10 +1069,7 @@ EXCLUDE_PATTERNS       =
 # (namespaces, classes, functions, etc.) that should be excluded from the
 # output. The symbol name can be a fully qualified name, a word, or if the
 # wildcard * is used, a substring. Examples: ANamespace, AClass,
-# AClass::ANamespace, ANamespace::*Test
-#
-# Note that the wildcards are matched against the file with absolute path, so to
-# exclude all test directories use the pattern */test/*
+# ANamespace::AClass, ANamespace::*Test

 EXCLUDE_SYMBOLS        =

@@ -1028,6 +1114,11 @@ IMAGE_PATH             =
 # code is scanned, but not when the output code is generated. If lines are added
 # or removed, the anchors will not be placed correctly.
 #
+# Note that doxygen will use the data processed and written to standard output
+# for further processing, therefore nothing else, like debug statements or used
+# commands (so in case of a Windows batch file always use @echo OFF), should be
+# written to standard output.
+#
 # Note that for custom extensions or not directly supported extensions you also
 # need to set EXTENSION_MAPPING for the extension otherwise the files are not
 # properly processed by doxygen.
@@ -1069,6 +1160,15 @@ FILTER_SOURCE_PATTERNS =

 USE_MDFILE_AS_MAINPAGE =

+# The Fortran standard specifies that for fixed formatted Fortran code all
+# characters from position 72 are to be considered as comment. A common
+# extension is to allow longer lines before the automatic comment starts. The
+# setting FORTRAN_COMMENT_AFTER will also make it possible that longer lines can
+# be processed before the automatic comment starts.
+# Minimum value: 7, maximum value: 10000, default value: 72.
+
+FORTRAN_COMMENT_AFTER  = 72
+
 #---------------------------------------------------------------------------
 # Configuration options related to source browsing
 #---------------------------------------------------------------------------
@@ -1083,7 +1183,8 @@ USE_MDFILE_AS_MAINPAGE =
 SOURCE_BROWSER         = NO

 # Setting the INLINE_SOURCES tag to YES will include the body of functions,
-# classes and enums directly into the documentation.
+# multi-line macros, enums or list initialized variables directly into the
+# documentation.
 # The default value is: NO.

 INLINE_SOURCES         = NO
@@ -1206,10 +1307,11 @@ CLANG_DATABASE_PATH    =

 ALPHABETICAL_INDEX     = YES

-# In case all classes in a project start with a common prefix, all classes will
-# be put under the same header in the alphabetical index. The IGNORE_PREFIX tag
-# can be used to specify a prefix (or a list of prefixes) that should be ignored
-# while generating the index headers.
+# The IGNORE_PREFIX tag can be used to specify a prefix (or a list of prefixes)
+# that should be ignored while generating the index headers. The IGNORE_PREFIX
+# tag works for classes, function and member names. The entity will be placed in
+# the alphabetical list under the first letter of the entity name that remains
+# after removing the prefix.
 # This tag requires that the tag ALPHABETICAL_INDEX is set to YES.

 IGNORE_PREFIX          =
@@ -1288,7 +1390,12 @@ HTML_STYLESHEET        =
 # Doxygen will copy the style sheet files to the output directory.
 # Note: The order of the extra style sheet files is of importance (e.g. the last
 # style sheet in the list overrules the setting of the previous ones in the
-# list). For an example see the documentation.
+# list).
+# Note: Since the styling of scrollbars can currently not be overruled in
+# Webkit/Chromium, the styling will be left out of the default doxygen.css if
+# one or more extra stylesheets have been specified. So if scrollbar
+# customization is desired it has to be added explicitly. For an example see the
+# documentation.
 # This tag requires that the tag GENERATE_HTML is set to YES.

 HTML_EXTRA_STYLESHEET  =
@@ -1303,6 +1410,19 @@ HTML_EXTRA_STYLESHEET  =

 HTML_EXTRA_FILES       =

+# The HTML_COLORSTYLE tag can be used to specify if the generated HTML output
+# should be rendered with a dark or light theme.
+# Possible values are: LIGHT always generate light mode output, DARK always
+# generate dark mode output, AUTO_LIGHT automatically set the mode according to
+# the user preference, use light mode if no preference is set (the default),
+# AUTO_DARK automatically set the mode according to the user preference, use
+# dark mode if no preference is set and TOGGLE allow to user to switch between
+# light and dark mode via a button.
+# The default value is: AUTO_LIGHT.
+# This tag requires that the tag GENERATE_HTML is set to YES.
+
+HTML_COLORSTYLE        = AUTO_LIGHT
+
 # The HTML_COLORSTYLE_HUE tag controls the color of the HTML output. Doxygen
 # will adjust the colors in the style sheet and background images according to
 # this color. Hue is specified as an angle on a color-wheel, see
@@ -1333,15 +1453,6 @@ HTML_COLORSTYLE_SAT    = 100

 HTML_COLORSTYLE_GAMMA  = 80

-# If the HTML_TIMESTAMP tag is set to YES then the footer of each generated HTML
-# page will contain the date and time when the page was generated. Setting this
-# to YES can help to show when doxygen was last run and thus if the
-# documentation is up to date.
-# The default value is: NO.
-# This tag requires that the tag GENERATE_HTML is set to YES.
-
-HTML_TIMESTAMP         = NO
-
 # If the HTML_DYNAMIC_MENUS tag is set to YES then the generated HTML
 # documentation will contain a main index with vertical navigation menus that
 # are dynamically created via JavaScript. If disabled, the navigation index will
@@ -1361,6 +1472,33 @@ HTML_DYNAMIC_MENUS     = YES

 HTML_DYNAMIC_SECTIONS  = NO

+# If the HTML_CODE_FOLDING tag is set to YES then classes and functions can be
+# dynamically folded and expanded in the generated HTML source code.
+# The default value is: YES.
+# This tag requires that the tag GENERATE_HTML is set to YES.
+
+HTML_CODE_FOLDING      = YES
+
+# If the HTML_COPY_CLIPBOARD tag is set to YES then doxygen will show an icon in
+# the top right corner of code and text fragments that allows the user to copy
+# its content to the clipboard. Note this only works if supported by the browser
+# and the web page is served via a secure context (see:
+# https://www.w3.org/TR/secure-contexts/), i.e. using the https: or file:
+# protocol.
+# The default value is: YES.
+# This tag requires that the tag GENERATE_HTML is set to YES.
+
+HTML_COPY_CLIPBOARD    = YES
+
+# Doxygen stores a couple of settings persistently in the browser (via e.g.
+# cookies). By default these settings apply to all HTML pages generated by
+# doxygen across all projects. The HTML_PROJECT_COOKIE tag can be used to store
+# the settings under a project specific key, such that the user preferences will
+# be stored separately.
+# This tag requires that the tag GENERATE_HTML is set to YES.
+
+HTML_PROJECT_COOKIE    =
+
 # With HTML_INDEX_NUM_ENTRIES one can control the preferred number of entries
 # shown in the various tree structured indices initially; the user can expand
 # and collapse entries dynamically later on. Doxygen will expand the tree to
@@ -1397,6 +1535,13 @@ GENERATE_DOCSET        = NO

 DOCSET_FEEDNAME        = "Doxygen generated docs"

+# This tag determines the URL of the docset feed. A documentation feed provides
+# an umbrella under which multiple documentation sets from a single provider
+# (such as a company or product suite) can be grouped.
+# This tag requires that the tag GENERATE_DOCSET is set to YES.
+
+DOCSET_FEEDURL         =
+
 # This tag specifies a string that should uniquely identify the documentation
 # set bundle. This should be a reverse domain-name style string, e.g.
 # com.mycompany.MyDocSet. Doxygen will append .docset to the name.
@@ -1484,6 +1629,16 @@ BINARY_TOC             = NO

 TOC_EXPAND             = NO

+# The SITEMAP_URL tag is used to specify the full URL of the place where the
+# generated documentation will be placed on the server by the user during the
+# deployment of the documentation. The generated sitemap is called sitemap.xml
+# and placed on the directory specified by HTML_OUTPUT. In case no SITEMAP_URL
+# is specified no sitemap is generated. For information about the sitemap
+# protocol see https://www.sitemaps.org
+# This tag requires that the tag GENERATE_HTML is set to YES.
+
+SITEMAP_URL            =
+
 # If the GENERATE_QHP tag is set to YES and both QHP_NAMESPACE and
 # QHP_VIRTUAL_FOLDER are set, an additional index file will be generated that
 # can be used as input for Qt's qhelpgenerator to generate a Qt Compressed Help
@@ -1601,7 +1756,7 @@ GENERATE_TREEVIEW      = NO
 # area (value NO) or if it should extend to the full height of the window (value
 # YES). Setting this to YES gives a layout similar to
 # https://docs.readthedocs.io with more room for contents, but less room for the
-# project logo, title, and description. If either GENERATOR_TREEVIEW or
+# project logo, title, and description. If either GENERATE_TREEVIEW or
 # DISABLE_INDEX is set to NO, this option has no effect.
 # The default value is: NO.
 # This tag requires that the tag GENERATE_HTML is set to YES.
@@ -1632,6 +1787,13 @@ TREEVIEW_WIDTH         = 250

 EXT_LINKS_IN_WINDOW    = NO

+# If the OBFUSCATE_EMAILS tag is set to YES, doxygen will obfuscate email
+# addresses.
+# The default value is: YES.
+# This tag requires that the tag GENERATE_HTML is set to YES.
+
+OBFUSCATE_EMAILS       = YES
+
 # If the HTML_FORMULA_FORMAT option is set to svg, doxygen will use the pdf2svg
 # tool (see https://github.com/dawbarton/pdf2svg) or inkscape (see
 # https://inkscape.org) to generate formulas as SVG images instead of PNGs for
@@ -1652,17 +1814,6 @@ HTML_FORMULA_FORMAT    = png

 FORMULA_FONTSIZE       = 10

-# Use the FORMULA_TRANSPARENT tag to determine whether or not the images
-# generated for formulas are transparent PNGs. Transparent PNGs are not
-# supported properly for IE 6.0, but are supported on all modern browsers.
-#
-# Note that when changing this option you need to delete any form_*.png files in
-# the HTML output directory before the changes have effect.
-# The default value is: YES.
-# This tag requires that the tag GENERATE_HTML is set to YES.
-
-FORMULA_TRANSPARENT    = YES
-
 # The FORMULA_MACROFILE can contain LaTeX \newcommand and \renewcommand commands
 # to create new LaTeX commands to be used in formulas as building blocks. See
 # the section "Including formulas" for details.
@@ -1976,9 +2127,16 @@ PDF_HYPERLINKS         = YES

 USE_PDFLATEX           = YES

-# If the LATEX_BATCHMODE tag is set to YES, doxygen will add the \batchmode
-# command to the generated LaTeX files. This will instruct LaTeX to keep running
-# if errors occur, instead of asking the user for help.
+# The LATEX_BATCHMODE tag signals the behavior of LaTeX in case of an error.
+# Possible values are: NO same as ERROR_STOP, YES same as BATCH, BATCH In batch
+# mode nothing is printed on the terminal, errors are scrolled as if <return> is
+# hit at every error; missing files that TeX tries to input or request from
+# keyboard input (\read on a not open input stream) cause the job to abort,
+# NON_STOP In nonstop mode the diagnostic message will appear on the terminal,
+# but there is no possibility of user interaction just like in batch mode,
+# SCROLL In scroll mode, TeX will stop only for missing files to input or if
+# keyboard input is necessary and ERROR_STOP In errorstop mode, TeX will stop at
+# each error, asking for user intervention.
 # The default value is: NO.
 # This tag requires that the tag GENERATE_LATEX is set to YES.

@@ -1999,14 +2157,6 @@ LATEX_HIDE_INDICES     = NO

 LATEX_BIB_STYLE        = plain

-# If the LATEX_TIMESTAMP tag is set to YES then the footer of each generated
-# page will contain the date and time when the page was generated. Setting this
-# to NO can help when comparing the output of multiple runs.
-# The default value is: NO.
-# This tag requires that the tag GENERATE_LATEX is set to YES.
-
-LATEX_TIMESTAMP        = NO
-
 # The LATEX_EMOJI_DIRECTORY tag is used to specify the (relative or absolute)
 # path from which the emoji images will be read. If a relative path is entered,
 # it will be relative to the LATEX_OUTPUT directory. If left blank the
@@ -2172,13 +2322,39 @@ DOCBOOK_OUTPUT         = docbook
 #---------------------------------------------------------------------------

 # If the GENERATE_AUTOGEN_DEF tag is set to YES, doxygen will generate an
-# AutoGen Definitions (see http://autogen.sourceforge.net/) file that captures
+# AutoGen Definitions (see https://autogen.sourceforge.net/) file that captures
 # the structure of the code including all documentation. Note that this feature
 # is still experimental and incomplete at the moment.
 # The default value is: NO.

 GENERATE_AUTOGEN_DEF   = NO

+#---------------------------------------------------------------------------
+# Configuration options related to Sqlite3 output
+#---------------------------------------------------------------------------
+
+# If the GENERATE_SQLITE3 tag is set to YES doxygen will generate a Sqlite3
+# database with symbols found by doxygen stored in tables.
+# The default value is: NO.
+
+GENERATE_SQLITE3       = NO
+
+# The SQLITE3_OUTPUT tag is used to specify where the Sqlite3 database will be
+# put. If a relative path is entered the value of OUTPUT_DIRECTORY will be put
+# in front of it.
+# The default directory is: sqlite3.
+# This tag requires that the tag GENERATE_SQLITE3 is set to YES.
+
+SQLITE3_OUTPUT         = sqlite3
+
+# The SQLITE3_RECREATE_DB tag is set to YES, the existing doxygen_sqlite3.db
+# database file will be recreated with each doxygen run. If set to NO, doxygen
+# will warn if a database file is already found and not modify it.
+# The default value is: YES.
+# This tag requires that the tag GENERATE_SQLITE3 is set to YES.
+
+SQLITE3_RECREATE_DB    = YES
+
 #---------------------------------------------------------------------------
 # Configuration options related to the Perl module output
 #---------------------------------------------------------------------------
@@ -2253,7 +2429,8 @@ SEARCH_INCLUDES        = YES

 # The INCLUDE_PATH tag can be used to specify one or more directories that
 # contain include files that are not input files but should be processed by the
-# preprocessor.
+# preprocessor. Note that the INCLUDE_PATH is not recursive, so the setting of
+# RECURSIVE has no effect here.
 # This tag requires that the tag SEARCH_INCLUDES is set to YES.

 INCLUDE_PATH           =
@@ -2282,8 +2459,18 @@ PREDEFINED             = VMA_CALL_PRE= \
                         VMA_NOT_NULL_NON_DISPATCHABLE= \
                         VMA_NULLABLE_NON_DISPATCHABLE= \
                         VMA_VULKAN_VERSION=1003000 \
+                         VMA_DEDICATED_ALLOCATION=1 \
+                         VMA_BIND_MEMORY2=1 \
+                         VMA_MEMORY_BUDGET=1 \
+                         VMA_BUFFER_DEVICE_ADDRESS=1 \
+                         VMA_MEMORY_PRIORITY=1 \
+                         VMA_KHR_MAINTENANCE4=1 \
+                         VMA_KHR_MAINTENANCE5=1 \
+                         VMA_EXTERNAL_MEMORY_WIN32=1 \
                         VMA_EXTERNAL_MEMORY=1 \
-                         VMA_MEMORY_PRIORITY=1
+                         VMA_EXTENDS_VK_STRUCT= \
+                         VMA_STATS_STRING_ENABLED=1 \
+                         VOLK_HEADER_VERSION=304

 # If the MACRO_EXPANSION and EXPAND_ONLY_PREDEF tags are set to YES then this
 # tag can be used to specify a list of macro names that should be expanded. The
@@ -2329,15 +2516,15 @@ TAGFILES               =

 GENERATE_TAGFILE       =

-# If the ALLEXTERNALS tag is set to YES, all external class will be listed in
-# the class index. If set to NO, only the inherited external classes will be
-# listed.
+# If the ALLEXTERNALS tag is set to YES, all external classes and namespaces
+# will be listed in the class and namespace index. If set to NO, only the
+# inherited external classes will be listed.
 # The default value is: NO.

 ALLEXTERNALS           = NO

 # If the EXTERNAL_GROUPS tag is set to YES, all external groups will be listed
-# in the modules index. If set to NO, only the current project's groups will be
+# in the topic index. If set to NO, only the current project's groups will be
 # listed.
 # The default value is: YES.

@@ -2351,25 +2538,9 @@ EXTERNAL_GROUPS        = YES
 EXTERNAL_PAGES         = YES

 #---------------------------------------------------------------------------
-# Configuration options related to the dot tool
+# Configuration options related to diagram generator tools
 #---------------------------------------------------------------------------

-# If the CLASS_DIAGRAMS tag is set to YES, doxygen will generate a class diagram
-# (in HTML and LaTeX) for classes with base or super classes. Setting the tag to
-# NO turns the diagrams off. Note that this option also works with HAVE_DOT
-# disabled, but it is recommended to install and use dot, since it yields more
-# powerful graphs.
-# The default value is: YES.
-
-CLASS_DIAGRAMS         = YES
-
-# You can include diagrams made with dia in doxygen documentation. Doxygen will
-# then run dia to produce the diagram and insert it in the documentation. The
-# DIA_PATH tag allows you to specify the directory where the dia binary resides.
-# If left empty dia is assumed to be found in the default search path.
-
-DIA_PATH               =
-
 # If set to YES the inheritance and collaboration graphs will hide inheritance
 # and usage relations if the target is undocumented or is not a class.
 # The default value is: YES.
@@ -2378,7 +2549,7 @@ HIDE_UNDOC_RELATIONS   = YES

 # If you set the HAVE_DOT tag to YES then doxygen will assume the dot tool is
 # available from the path. This tool is part of Graphviz (see:
-# http://www.graphviz.org/), a graph visualization toolkit from AT&T and Lucent
+# https://www.graphviz.org/), a graph visualization toolkit from AT&T and Lucent
 # Bell Labs. The other options in this section have no effect if this option is
 # set to NO
 # The default value is: NO.
@@ -2395,49 +2566,77 @@ HAVE_DOT               = NO

 DOT_NUM_THREADS        = 0

-# When you want a differently looking font in the dot files that doxygen
-# generates you can specify the font name using DOT_FONTNAME. You need to make
-# sure dot is able to find the font, which can be done by putting it in a
-# standard location or by setting the DOTFONTPATH environment variable or by
-# setting DOT_FONTPATH to the directory containing the font.
-# The default value is: Helvetica.
+# DOT_COMMON_ATTR is common attributes for nodes, edges and labels of
+# subgraphs. When you want a differently looking font in the dot files that
+# doxygen generates you can specify fontname, fontcolor and fontsize attributes.
+# For details please see <a href=https://graphviz.org/doc/info/attrs.html>Node,
+# Edge and Graph Attributes specification</a> You need to make sure dot is able
+# to find the font, which can be done by putting it in a standard location or by
+# setting the DOTFONTPATH environment variable or by setting DOT_FONTPATH to the
+# directory containing the font. Default graphviz fontsize is 14.
+# The default value is: fontname=Helvetica,fontsize=10.
 # This tag requires that the tag HAVE_DOT is set to YES.

-DOT_FONTNAME           = Helvetica
+DOT_COMMON_ATTR        = "fontname=Helvetica,fontsize=10"

-# The DOT_FONTSIZE tag can be used to set the size (in points) of the font of
-# dot graphs.
-# Minimum value: 4, maximum value: 24, default value: 10.
+# DOT_EDGE_ATTR is concatenated with DOT_COMMON_ATTR. For elegant style you can
+# add 'arrowhead=open, arrowtail=open, arrowsize=0.5'. <a
+# href=https://graphviz.org/doc/info/arrows.html>Complete documentation about
+# arrows shapes.</a>
+# The default value is: labelfontname=Helvetica,labelfontsize=10.
 # This tag requires that the tag HAVE_DOT is set to YES.

-DOT_FONTSIZE           = 10
+DOT_EDGE_ATTR          = "labelfontname=Helvetica,labelfontsize=10"

-# By default doxygen will tell dot to use the default font as specified with
-# DOT_FONTNAME. If you specify a different font using DOT_FONTNAME you can set
-# the path where dot can find it using this tag.
+# DOT_NODE_ATTR is concatenated with DOT_COMMON_ATTR. For view without boxes
+# around nodes set 'shape=plain' or 'shape=plaintext' <a
+# href=https://www.graphviz.org/doc/info/shapes.html>Shapes specification</a>
+# The default value is: shape=box,height=0.2,width=0.4.
+# This tag requires that the tag HAVE_DOT is set to YES.
+
+DOT_NODE_ATTR          = "shape=box,height=0.2,width=0.4"
+
+# You can set the path where dot can find font specified with fontname in
+# DOT_COMMON_ATTR and others dot attributes.
 # This tag requires that the tag HAVE_DOT is set to YES.

 DOT_FONTPATH           =

-# If the CLASS_GRAPH tag is set to YES then doxygen will generate a graph for
-# each documented class showing the direct and indirect inheritance relations.
-# Setting this tag to YES will force the CLASS_DIAGRAMS tag to NO.
+# If the CLASS_GRAPH tag is set to YES or GRAPH or BUILTIN then doxygen will
+# generate a graph for each documented class showing the direct and indirect
+# inheritance relations. In case the CLASS_GRAPH tag is set to YES or GRAPH and
+# HAVE_DOT is enabled as well, then dot will be used to draw the graph. In case
+# the CLASS_GRAPH tag is set to YES and HAVE_DOT is disabled or if the
+# CLASS_GRAPH tag is set to BUILTIN, then the built-in generator will be used.
+# If the CLASS_GRAPH tag is set to TEXT the direct and indirect inheritance
+# relations will be shown as texts / links. Explicit enabling an inheritance
+# graph or choosing a different representation for an inheritance graph of a
+# specific class, can be accomplished by means of the command \inheritancegraph.
+# Disabling an inheritance graph can be accomplished by means of the command
+# \hideinheritancegraph.
+# Possible values are: NO, YES, TEXT, GRAPH and BUILTIN.
 # The default value is: YES.
-# This tag requires that the tag HAVE_DOT is set to YES.

 CLASS_GRAPH            = YES

 # If the COLLABORATION_GRAPH tag is set to YES then doxygen will generate a
 # graph for each documented class showing the direct and indirect implementation
 # dependencies (inheritance, containment, and class references variables) of the
-# class with other documented classes.
+# class with other documented classes. Explicit enabling a collaboration graph,
+# when COLLABORATION_GRAPH is set to NO, can be accomplished by means of the
+# command \collaborationgraph. Disabling a collaboration graph can be
+# accomplished by means of the command \hidecollaborationgraph.
 # The default value is: YES.
 # This tag requires that the tag HAVE_DOT is set to YES.

 COLLABORATION_GRAPH    = YES

 # If the GROUP_GRAPHS tag is set to YES then doxygen will generate a graph for
-# groups, showing the direct groups dependencies.
+# groups, showing the direct groups dependencies. Explicit enabling a group
+# dependency graph, when GROUP_GRAPHS is set to NO, can be accomplished by means
+# of the command \groupgraph. Disabling a directory graph can be accomplished by
+# means of the command \hidegroupgraph. See also the chapter Grouping in the
+# manual.
 # The default value is: YES.
 # This tag requires that the tag HAVE_DOT is set to YES.

@@ -2479,8 +2678,8 @@ DOT_UML_DETAILS        = NO

 # The DOT_WRAP_THRESHOLD tag can be used to set the maximum number of characters
 # to display on a single line. If the actual line length exceeds this threshold
-# significantly it will wrapped across multiple lines. Some heuristics are apply
-# to avoid ugly line breaks.
+# significantly it will be wrapped across multiple lines. Some heuristics are
+# applied to avoid ugly line breaks.
 # Minimum value: 0, maximum value: 1000, default value: 17.
 # This tag requires that the tag HAVE_DOT is set to YES.

@@ -2497,7 +2696,9 @@ TEMPLATE_RELATIONS     = NO
 # If the INCLUDE_GRAPH, ENABLE_PREPROCESSING and SEARCH_INCLUDES tags are set to
 # YES then doxygen will generate a graph for each documented file showing the
 # direct and indirect include dependencies of the file with other documented
-# files.
+# files. Explicit enabling an include graph, when INCLUDE_GRAPH is is set to NO,
+# can be accomplished by means of the command \includegraph. Disabling an
+# include graph can be accomplished by means of the command \hideincludegraph.
 # The default value is: YES.
 # This tag requires that the tag HAVE_DOT is set to YES.

@@ -2506,7 +2707,10 @@ INCLUDE_GRAPH          = YES
 # If the INCLUDED_BY_GRAPH, ENABLE_PREPROCESSING and SEARCH_INCLUDES tags are
 # set to YES then doxygen will generate a graph for each documented file showing
 # the direct and indirect include dependencies of the file with other documented
-# files.
+# files. Explicit enabling an included by graph, when INCLUDED_BY_GRAPH is set
+# to NO, can be accomplished by means of the command \includedbygraph. Disabling
+# an included by graph can be accomplished by means of the command
+# \hideincludedbygraph.
 # The default value is: YES.
 # This tag requires that the tag HAVE_DOT is set to YES.

@@ -2546,16 +2750,26 @@ GRAPHICAL_HIERARCHY    = YES
 # If the DIRECTORY_GRAPH tag is set to YES then doxygen will show the
 # dependencies a directory has on other directories in a graphical way. The
 # dependency relations are determined by the #include relations between the
-# files in the directories.
+# files in the directories. Explicit enabling a directory graph, when
+# DIRECTORY_GRAPH is set to NO, can be accomplished by means of the command
+# \directorygraph. Disabling a directory graph can be accomplished by means of
+# the command \hidedirectorygraph.
 # The default value is: YES.
 # This tag requires that the tag HAVE_DOT is set to YES.

 DIRECTORY_GRAPH        = YES

+# The DIR_GRAPH_MAX_DEPTH tag can be used to limit the maximum number of levels
+# of child directories generated in directory dependency graphs by dot.
+# Minimum value: 1, maximum value: 25, default value: 1.
+# This tag requires that the tag DIRECTORY_GRAPH is set to YES.
+
+DIR_GRAPH_MAX_DEPTH    = 1
+
 # The DOT_IMAGE_FORMAT tag can be used to set the image format of the images
 # generated by dot. For an explanation of the image formats see the section
 # output formats in the documentation of the dot tool (Graphviz (see:
-# http://www.graphviz.org/)).
+# https://www.graphviz.org/)).
 # Note: If you choose svg you need to set HTML_FILE_EXTENSION to xhtml in order
 # to make the SVG files visible in IE 9+ (other browsers do not have this
 # requirement).
@@ -2592,11 +2806,12 @@ DOT_PATH               =

 DOTFILE_DIRS           =

-# The MSCFILE_DIRS tag can be used to specify one or more directories that
-# contain msc files that are included in the documentation (see the \mscfile
-# command).
+# You can include diagrams made with dia in doxygen documentation. Doxygen will
+# then run dia to produce the diagram and insert it in the documentation. The
+# DIA_PATH tag allows you to specify the directory where the dia binary resides.
+# If left empty dia is assumed to be found in the default search path.

-MSCFILE_DIRS           =
+DIA_PATH               =

 # The DIAFILE_DIRS tag can be used to specify one or more directories that
 # contain dia files that are included in the documentation (see the \diafile
@@ -2605,10 +2820,10 @@ MSCFILE_DIRS           =
 DIAFILE_DIRS           =

 # When using plantuml, the PLANTUML_JAR_PATH tag should be used to specify the
-# path where java can find the plantuml.jar file. If left blank, it is assumed
-# PlantUML is not used or called during a preprocessing step. Doxygen will
-# generate a warning when it encounters a \startuml command in this case and
-# will not generate output for the diagram.
+# path where java can find the plantuml.jar file or to the filename of jar file
+# to be used. If left blank, it is assumed PlantUML is not used or called during
+# a preprocessing step. Doxygen will generate a warning when it encounters a
+# \startuml command in this case and will not generate output for the diagram.

 PLANTUML_JAR_PATH      =

@@ -2646,18 +2861,6 @@ DOT_GRAPH_MAX_NODES    = 50

 MAX_DOT_GRAPH_DEPTH    = 0

-# Set the DOT_TRANSPARENT tag to YES to generate images with a transparent
-# background. This is disabled by default, because dot on Windows does not seem
-# to support this out of the box.
-#
-# Warning: Depending on the platform used, enabling this option may lead to
-# badly anti-aliased labels on the edges of a graph (i.e. they become hard to
-# read).
-# The default value is: NO.
-# This tag requires that the tag HAVE_DOT is set to YES.
-
-DOT_TRANSPARENT        = NO
-
 # Set the DOT_MULTI_TARGETS tag to YES to allow dot to generate multiple output
 # files in one run (i.e. multiple -o and -T options on the command line). This
 # makes dot run faster, but since only newer versions of dot (>1.8.10) support
@@ -2670,6 +2873,8 @@ DOT_MULTI_TARGETS      = NO
 # If the GENERATE_LEGEND tag is set to YES doxygen will generate a legend page
 # explaining the meaning of the various boxes and arrows in the dot generated
 # graphs.
+# Note: This tag requires that UML_LOOK isn't set, i.e. the doxygen internal
+# graphical representation for inheritance and collaboration diagrams is used.
 # The default value is: YES.
 # This tag requires that the tag HAVE_DOT is set to YES.

@@ -2683,3 +2888,19 @@ GENERATE_LEGEND        = YES
 # The default value is: YES.

 DOT_CLEANUP            = YES
+
+# You can define message sequence charts within doxygen comments using the \msc
+# command. If the MSCGEN_TOOL tag is left empty (the default), then doxygen will
+# use a built-in version of mscgen tool to produce the charts. Alternatively,
+# the MSCGEN_TOOL tag can also specify the name an external tool. For instance,
+# specifying prog as the value, doxygen will call the tool as prog -T
+# <outfile_format> -o <outputfile> <inputfile>. The external tool should support
+# output file formats "png", "eps", "svg", and "ismap".
+
+MSCGEN_TOOL            =
+
+# The MSCFILE_DIRS tag can be used to specify one or more directories that
+# contain msc files that are included in the documentation (see the \mscfile
+# command).
+
+MSCFILE_DIRS           =
@@ -1,4 +1,4 @@
-Copyright (c) 2017-2022 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2017-2025 Advanced Micro Devices, Inc. All rights reserved.

 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal
@@ -10,11 +10,6 @@ Easy to integrate Vulkan memory allocation library.

 **Product page:** [Vulkan Memory Allocator on GPUOpen](https://gpuopen.com/gaming-product/vulkan-memory-allocator/)

-**Build status:**
-
- Windows: [![Build status](https://ci.appveyor.com/api/projects/status/4vlcrb0emkaio2pn/branch/master?svg=true)](https://ci.appveyor.com/project/adam-sawicki-amd/vulkanmemoryallocator/branch/master)  
- Linux: [![Build Status](https://app.travis-ci.com/GPUOpen-LibrariesAndSDKs/VulkanMemoryAllocator.svg?branch=master)](https://app.travis-ci.com/GPUOpen-LibrariesAndSDKs/VulkanMemoryAllocator)
-
 [![Average time to resolve an issue](http://isitmaintained.com/badge/resolution/GPUOpen-LibrariesAndSDKs/VulkanMemoryAllocator.svg)](http://isitmaintained.com/project/GPUOpen-LibrariesAndSDKs/VulkanMemoryAllocator "Average time to resolve an issue")

 # Problem
@@ -48,13 +43,17 @@ Additional features:
 - Support for sparse binding and sparse residency: Convenience functions that allocate or free multiple memory pages at once.
 - Custom memory pools: Create a pool with desired parameters (e.g. fixed or limited maximum size) and allocate memory out of it.
 - Linear allocator: Create a pool with linear algorithm and use it for much faster allocations and deallocations in free-at-once, stack, double stack, or ring buffer fashion.
- Support for Vulkan 1.0, 1.1, 1.2, 1.3.
+- Support for Vulkan 1.0...1.4.
 - Support for extensions (and equivalent functionality included in new Vulkan versions):
   - VK_KHR_dedicated_allocation: Just enable it and it will be used automatically by the library.
-   - VK_KHR_buffer_device_address: Flag `VK_MEMORY_ALLOCATE_DEVICE_ADDRESS_BIT_KHR` is automatically added to memory allocations where needed.
+   - VK_KHR_bind_memory2.
+   - VK_KHR_maintenance4.
+   - VK_KHR_maintenance5, including `VkBufferUsageFlags2CreateInfoKHR`.
   - VK_EXT_memory_budget: Used internally if available to query for current usage and budget. If not available, it falls back to an estimation based on memory heap sizes.
+   - VK_KHR_buffer_device_address: Flag `VK_MEMORY_ALLOCATE_DEVICE_ADDRESS_BIT_KHR` is automatically added to memory allocations where needed.
   - VK_EXT_memory_priority: Set `priority` of allocations or custom pools and it will be set automatically using this extension.
-   - VK_AMD_device_coherent_memory
+   - VK_AMD_device_coherent_memory.
+   - VK_KHR_external_memory_win32.
 - Defragmentation of GPU and CPU memory: Let the library move data around to free some memory blocks and make your allocations better compacted.
 - Statistics: Obtain brief or detailed statistics about the amount of memory used, unused, number of allocated blocks, number of allocations etc. - globally, per memory heap, and per memory type.
 - Debug annotations: Associate custom `void* pUserData` and debug `char* pName` with each allocation.
@@ -99,42 +98,51 @@ With this one function call:

 # How to build

-On Windows it is recommended to use [CMake UI](https://cmake.org/runningcmake/). Alternatively you can generate a Visual Studio project map using CMake in command line: `cmake -B./build/ -DCMAKE_BUILD_TYPE=Debug -G "Visual Studio 16 2019" -A x64 ./`
+On Windows it is recommended to use [CMake GUI](https://cmake.org/runningcmake/).
+
+Alternatively you can generate/open a Visual Studio from the command line:
+
+```sh
+# By default CMake picks the newest version of Visual Studio it can use
+cmake -S .  -B build -D VMA_BUILD_SAMPLES=ON
+cmake --open build
+```

 On Linux:

-```
-mkdir build
-cd build
-cmake ..
-make
+```sh
+cmake -S . -B build
+# Since VMA has no source files, you can skip to installation immediately
+cmake --install build --prefix build/install
 ```

-The following targets are available
+## How to use

-| Target | Description | CMake option | Default setting |
-| ------------- | ------------- | ------------- | ------------- |
-| VmaSample | VMA sample application | `VMA_BUILD_SAMPLE` | `OFF` |
-| VmaBuildSampleShaders | Shaders for VmaSample | `VMA_BUILD_SAMPLE_SHADERS` | `OFF` |
+After calling either `find_package` or `add_subdirectory` simply link the library.
+This automatically handles configuring the include directory. Example:

-Please note that while VulkanMemoryAllocator library is supported on other platforms besides Windows, VmaSample is not.
+```cmake
+find_package(VulkanMemoryAllocator CONFIG REQUIRED)
+target_link_libraries(YourGameEngine PRIVATE GPUOpen::VulkanMemoryAllocator)
+```

-These CMake options are available
+For more info on using CMake visit the official [CMake documentation](https://cmake.org/cmake/help/latest/index.html).

-| CMake option | Description | Default setting |
-| ------------- | ------------- | ------------- |
-| `VMA_RECORDING_ENABLED` | Enable VMA memory recording for debugging | `OFF` |
-| `VMA_USE_STL_CONTAINERS` | Use C++ STL containers instead of VMA's containers | `OFF` |
-| `VMA_STATIC_VULKAN_FUNCTIONS` | Link statically with Vulkan API | `OFF` |
-| `VMA_DYNAMIC_VULKAN_FUNCTIONS` | Fetch pointers to Vulkan functions internally (no static linking) | `ON` |
-| `VMA_DEBUG_ALWAYS_DEDICATED_MEMORY` | Every allocation will have its own memory block | `OFF` |
-| `VMA_DEBUG_INITIALIZE_ALLOCATIONS` | Automatically fill new allocations and destroyed allocations with some bit pattern | `OFF` |
-| `VMA_DEBUG_GLOBAL_MUTEX` | Enable single mutex protecting all entry calls to the library | `OFF` |
-| `VMA_DEBUG_DONT_EXCEED_MAX_MEMORY_ALLOCATION_COUNT` | Never exceed [VkPhysicalDeviceLimits::maxMemoryAllocationCount](https://www.khronos.org/registry/vulkan/specs/1.1-extensions/html/vkspec.html#limits-maxMemoryAllocationCount) and return error | `OFF` |
+## Building using vcpkg
+
+You can download and install VulkanMemoryAllocator using the [vcpkg](https://github.com/Microsoft/vcpkg) dependency manager:
+
+    git clone https://github.com/Microsoft/vcpkg.git
+    cd vcpkg
+    ./bootstrap-vcpkg.sh
+    ./vcpkg integrate install
+    ./vcpkg install vulkan-memory-allocator
+
+The VulkanMemoryAllocator port in vcpkg is kept up to date by Microsoft team members and community contributors. If the version is out of date, please [create an issue or pull request](https://github.com/Microsoft/vcpkg) on the vcpkg repository.

 # Binaries

-The release comes with precompiled binary executable for "VulkanSample" application which contains test suite. It is compiled using Visual Studio 2019, so it requires appropriate libraries to work, including "MSVCP140.dll", "VCRUNTIME140.dll", "VCRUNTIME140_1.dll". If the launch fails with error message telling about those files missing, please download and install [Microsoft Visual C++ Redistributable for Visual Studio 2015, 2017 and 2019](https://support.microsoft.com/en-us/help/2977003/the-latest-supported-visual-c-downloads), "x64" version.
+The release comes with precompiled binary executable for "VulkanSample" application which contains test suite. It is compiled using Visual Studio 2022, so it requires appropriate libraries to work, including "MSVCP140.dll", "VCRUNTIME140.dll", "VCRUNTIME140_1.dll". If the launch fails with error message telling about those files missing, please download and install [Microsoft Visual C++ Redistributable](https://support.microsoft.com/en-us/help/2977003/the-latest-supported-visual-c-downloads), "X64" version.

 # Read more

@@ -142,24 +150,31 @@ See **[Documentation](https://gpuopen-librariesandsdks.github.io/VulkanMemoryAll

 # Software using this library

+- **[Blender](https://www.blender.org)**
+- **[Qt Project](https://github.com/qt)**
+- **[Baldur's Gate III](https://www.mobygames.com/game/150689/baldurs-gate-iii/credits/windows/?autoplatform=true)**
+- **[Cyberpunk 2077](https://www.mobygames.com/game/128136/cyberpunk-2077/credits/windows/?autoplatform=true)**
 - **[X-Plane](https://x-plane.com/)**
 - **[Detroit: Become Human](https://gpuopen.com/learn/porting-detroit-3/)**
 - **[Vulkan Samples](https://github.com/LunarG/VulkanSamples)** - official Khronos Vulkan samples. License: Apache-style.
+- **[GFXReconstruct](https://github.com/LunarG/gfxreconstruct)** - a tools for the capture and replay of graphics API calls. License: MIT.
 - **[Anvil](https://github.com/GPUOpen-LibrariesAndSDKs/Anvil)** - cross-platform framework for Vulkan. License: MIT.
 - **[Filament](https://github.com/google/filament)** - physically based rendering engine for Android, Windows, Linux and macOS, from Google. Apache License 2.0.
 - **[Atypical Games - proprietary game engine](https://developer.samsung.com/galaxy-gamedev/gamedev-blog/infinitejet.html)**
 - **[Flax Engine](https://flaxengine.com/)**
 - **[Godot Engine](https://github.com/godotengine/godot/)** - multi-platform 2D and 3D game engine. License: MIT.
 - **[Lightweight Java Game Library (LWJGL)](https://www.lwjgl.org/)** - includes binding of the library for Java. License: BSD.
+- **[LightweightVK](https://github.com/corporateshark/lightweightvk)** - lightweight C++ bindless Vulkan 1.3 wrapper. License: MIT.
 - **[PowerVR SDK](https://github.com/powervr-graphics/Native_SDK)** - C++ cross-platform 3D graphics SDK, from Imagination. License: MIT.
 - **[Skia](https://github.com/google/skia)** - complete 2D graphic library for drawing Text, Geometries, and Images, from Google.
 - **[The Forge](https://github.com/ConfettiFX/The-Forge)** - cross-platform rendering framework. Apache License 2.0.
- **[VK9](https://github.com/disks86/VK9)** - Direct3D 9 compatibility layer using Vulkan. Zlib lincese.
+- **[VK9](https://github.com/disks86/VK9)** - Direct3D 9 compatibility layer using Vulkan. Zlib license.
 - **[vkDOOM3](https://github.com/DustinHLand/vkDOOM3)** - Vulkan port of GPL DOOM 3 BFG Edition. License: GNU GPL.
 - **[vkQuake2](https://github.com/kondrak/vkQuake2)** - vanilla Quake 2 with Vulkan support. License: GNU GPL.
 - **[Vulkan Best Practice for Mobile Developers](https://github.com/ARM-software/vulkan_best_practice_for_mobile_developers)** from ARM. License: MIT.
 - **[RPCS3](https://github.com/RPCS3/rpcs3)** - PlayStation 3 emulator/debugger. License: GNU GPLv2.
 - **[PPSSPP](https://github.com/hrydgard/ppsspp)** - Playstation Portable emulator/debugger. License: GNU GPLv2+.
+- **[Wicked Engine](https://github.com/turanszkij/WickedEngine)** - 3D engine with modern graphics 

 [Many other projects on GitHub](https://github.com/search?q=AMD_VULKAN_MEMORY_ALLOCATOR_H&type=Code) and some game development studios that use Vulkan in their games.

@@ -167,7 +182,8 @@ See **[Documentation](https://gpuopen-librariesandsdks.github.io/VulkanMemoryAll

 - **[D3D12 Memory Allocator](https://github.com/GPUOpen-LibrariesAndSDKs/D3D12MemoryAllocator)** - equivalent library for Direct3D 12. License: MIT.
 - **[Awesome Vulkan](https://github.com/vinjn/awesome-vulkan)** - a curated list of awesome Vulkan libraries, debuggers and resources.
- **[VulkanMemoryAllocator-Hpp](https://github.com/malte-v/VulkanMemoryAllocator-Hpp)** - C++ binding for this library. License: CC0-1.0.
+- **[vcpkg](https://github.com/Microsoft/vcpkg)** dependency manager from Microsoft also offers a port of this library.
+- **[VulkanMemoryAllocator-Hpp](https://github.com/YaaZ/VulkanMemoryAllocator-Hpp)** - C++ binding for this library. License: CC0-1.0.
 - **[PyVMA](https://github.com/realitix/pyvma)** - Python wrapper for this library. Author: Jean-Sébastien B. (@realitix). License: Apache 2.0.
 - **[vk-mem](https://github.com/gwihlidal/vk-mem-rs)** - Rust binding for this library. Author: Graham Wihlidal. License: Apache 2.0 or MIT.
 - **[Haskell bindings](https://hackage.haskell.org/package/VulkanMemoryAllocator)**, **[github](https://github.com/expipiplus1/vulkan/tree/master/VulkanMemoryAllocator)** - Haskell bindings for this library. Author: Ellie Hermaszewska (@expipiplus1). License BSD-3-Clause.
@@ -1,109 +1,101 @@
-set(VMA_LIBRARY_SOURCE_FILES
+#
+# Copyright (c) 2017-2025 Advanced Micro Devices, Inc. All rights reserved.
+#
+# Permission is hereby granted, free of charge, to any person obtaining a copy
+# of this software and associated documentation files (the "Software"), to deal
+# in the Software without restriction, including without limitation the rights
+# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+# copies of the Software, and to permit persons to whom the Software is
+# furnished to do so, subject to the following conditions:
+#
+# The above copyright notice and this permission notice shall be included in
+# all copies or substantial portions of the Software.
+#
+# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+# THE SOFTWARE.
+#
+
+option(VMA_STATIC_VULKAN_FUNCTIONS "Link statically with Vulkan API" ON)
+option(VMA_DYNAMIC_VULKAN_FUNCTIONS "Fetch pointers to Vulkan functions internally (no static linking)" OFF)
+set(VMA_VOLK_HEADER_PATH "" CACHE STRING "Path to volk.h file from the volk library (optional)")
+option(VMA_DEBUG_ALWAYS_DEDICATED_MEMORY "Every allocation will have its own memory block" OFF)
+option(VMA_DEBUG_INITIALIZE_ALLOCATIONS "Automatically fill new allocations and destroyed allocations with some bit pattern" OFF)
+option(VMA_DEBUG_GLOBAL_MUTEX "Enable single mutex protecting all entry calls to the library" OFF)
+option(VMA_DEBUG_DONT_EXCEED_MAX_MEMORY_ALLOCATION_COUNT "Never exceed VkPhysicalDeviceLimits::maxMemoryAllocationCount and return error" OFF)
+
+message(STATUS "VMA_STATIC_VULKAN_FUNCTIONS = ${VMA_STATIC_VULKAN_FUNCTIONS}")
+message(STATUS "VMA_DYNAMIC_VULKAN_FUNCTIONS = ${VMA_DYNAMIC_VULKAN_FUNCTIONS}")
+message(STATUS "VMA_DEBUG_ALWAYS_DEDICATED_MEMORY = ${VMA_DEBUG_ALWAYS_DEDICATED_MEMORY}")
+message(STATUS "VMA_DEBUG_INITIALIZE_ALLOCATIONS = ${VMA_DEBUG_INITIALIZE_ALLOCATIONS}")
+message(STATUS "VMA_DEBUG_GLOBAL_MUTEX = ${VMA_DEBUG_GLOBAL_MUTEX}")
+message(STATUS "VMA_DEBUG_DONT_EXCEED_MAX_MEMORY_ALLOCATION_COUNT = ${VMA_DEBUG_DONT_EXCEED_MAX_MEMORY_ALLOCATION_COUNT}")
+
+set(CMAKE_CXX_STANDARD 14)
+set(CMAKE_CXX_STANDARD_REQUIRED ON)
+set(CMAKE_CXX_EXTENSIONS OFF)
+
+find_package(Vulkan REQUIRED)
+
+add_executable(VmaSample)
+target_sources(VmaSample PRIVATE
+    Common.cpp
+    Common.h
+    SparseBindingTest.cpp
+    SparseBindingTest.h
+    Tests.cpp
+    Tests.h
    VmaUsage.cpp
    VmaUsage.h
-    "${PROJECT_SOURCE_DIR}/include/vk_mem_alloc.h"
+    VolkUsage.cpp
+    VulkanSample.cpp
+    ../include/vk_mem_alloc.h
 )

-# Visual Studio only debugger symbols
-if(WIN32 AND ${CMAKE_GENERATOR} MATCHES "Visual Studio.*")
-    set(VMA_LIBRARY_SOURCE_FILES ${VMA_LIBRARY_SOURCE_FILES} vk_mem_alloc.natvis)
+# Only link to Vulkan library if volk loader is not used, but always add Vulkan headers directory.
+if("${VMA_VOLK_HEADER_PATH}" STREQUAL "")
+    target_link_libraries(VmaSample PUBLIC Vulkan::Vulkan)
+else()
+    target_link_libraries(VmaSample PUBLIC Vulkan::Headers)
 endif()

-set(CMAKE_DEBUG_POSTFIX d)
-set(CMAKE_RELWITHDEBINFO_POSTFIX rd)
-set(CMAKE_MINSIZEREL_POSTFIX s)
+target_link_libraries(VmaSample PRIVATE GPUOpen::VulkanMemoryAllocator)

-add_library(VulkanMemoryAllocator ${VMA_LIBRARY_SOURCE_FILES})
-
-set_target_properties(
-    VulkanMemoryAllocator PROPERTIES
-
-    CXX_EXTENSIONS OFF
-    # Use C++14
-    CXX_STANDARD 14
-    CXX_STANDARD_REQUIRED ON
-)
-
-target_include_directories(VulkanMemoryAllocator PUBLIC "${PROJECT_SOURCE_DIR}/include")
-
-# Only link to Vulkan if static linking is used
-if(${VMA_STATIC_VULKAN_FUNCTIONS})
-    target_link_libraries(VulkanMemoryAllocator PUBLIC Vulkan::Vulkan)
-endif()
-
-target_compile_definitions(
-    VulkanMemoryAllocator
-
-    PUBLIC
+target_compile_definitions(VmaSample PUBLIC
    VMA_STATIC_VULKAN_FUNCTIONS=$<BOOL:${VMA_STATIC_VULKAN_FUNCTIONS}>
    VMA_DYNAMIC_VULKAN_FUNCTIONS=$<BOOL:${VMA_DYNAMIC_VULKAN_FUNCTIONS}>
    VMA_DEBUG_ALWAYS_DEDICATED_MEMORY=$<BOOL:${VMA_DEBUG_ALWAYS_DEDICATED_MEMORY}>
    VMA_DEBUG_INITIALIZE_ALLOCATIONS=$<BOOL:${VMA_DEBUG_INITIALIZE_ALLOCATIONS}>
    VMA_DEBUG_GLOBAL_MUTEX=$<BOOL:${VMA_DEBUG_GLOBAL_MUTEX}>
    VMA_DEBUG_DONT_EXCEED_MAX_MEMORY_ALLOCATION_COUNT=$<BOOL:${VMA_DEBUG_DONT_EXCEED_MAX_MEMORY_ALLOCATION_COUNT}>
-    VMA_RECORDING_ENABLED=$<BOOL:${VMA_RECORDING_ENABLED}>
 )

-install(TARGETS VulkanMemoryAllocator DESTINATION "lib")
-install(FILES "${PROJECT_SOURCE_DIR}/include/vk_mem_alloc.h" DESTINATION "include")
+# Provides MSVC users nicer debugging support
+target_sources(VmaSample PRIVATE vk_mem_alloc.natvis)

-if(VMA_BUILD_SAMPLE)
-    if(WIN32)
-        set(VMA_SAMPLE_SOURCE_FILES
-            Common.cpp
-            Common.h
-            SparseBindingTest.cpp
-            SparseBindingTest.h
-            Tests.cpp
-            Tests.h
-            VulkanSample.cpp
-        )
+add_subdirectory(Shaders)
+add_dependencies(VmaSample VmaSampleShaders)

-        add_executable(VmaSample ${VMA_SAMPLE_SOURCE_FILES})
-		add_dependencies(VmaSample VulkanMemoryAllocator VmaSampleShaders)
-		
-        # Visual Studio specific settings
-        if(${CMAKE_GENERATOR} MATCHES "Visual Studio.*")
-            # Use Unicode instead of multibyte set
-            add_compile_definitions(UNICODE _UNICODE)
-            
-            # Set VmaSample as startup project
-            set_property(DIRECTORY "${PROJECT_SOURCE_DIR}" PROPERTY VS_STARTUP_PROJECT "VmaSample")
-
-            # Add C++ warnings and security checks
-            set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /permissive- /sdl /W3")
-            
-            # Enable multithreaded compiling
-            target_compile_options(VmaSample PRIVATE "/MP")
-
-            # Set working directory for Visual Studio debugger
-            set_target_properties(
-                VmaSample
-                PROPERTIES VS_DEBUGGER_WORKING_DIRECTORY "${PROJECT_SOURCE_DIR}/bin"
-            )
-        endif()
-
-        set_target_properties(
-            VmaSample PROPERTIES
-
-            CXX_EXTENSIONS OFF
-            # Use C++14
-            CXX_STANDARD 14
-            CXX_STANDARD_REQUIRED ON
-        )
-
-        target_link_libraries(
-            VmaSample
-            PRIVATE
-            
-            VulkanMemoryAllocator
-            Vulkan::Vulkan
-        )
+if(NOT "${VMA_VOLK_HEADER_PATH}" STREQUAL "")
+    if(EXISTS "${VMA_VOLK_HEADER_PATH}")
+        message(STATUS "File volk.h found and used from path: ${VMA_VOLK_HEADER_PATH}")
+        target_compile_definitions(VmaSample PRIVATE VMA_VOLK_HEADER_PATH="${VMA_VOLK_HEADER_PATH}")
    else()
-        message(STATUS "VmaSample application is not supported to Linux")
+        message(FATAL_ERROR "File volk.h not found in path: ${VMA_VOLK_HEADER_PATH}")
    endif()
 endif()

-if(VMA_BUILD_SAMPLE_SHADERS)
-    add_subdirectory(Shaders)
-endif()
+# Use Unicode instead of multibyte set
+add_compile_definitions(UNICODE _UNICODE)
+
+# Add C++ warnings and security checks
+add_compile_options(/permissive- /sdl /W3)
+
+# Set VmaSample as startup project
+set_property(DIRECTORY "${PROJECT_SOURCE_DIR}" PROPERTY VS_STARTUP_PROJECT "VmaSample")
+
+set_target_properties(VmaSample PROPERTIES VS_DEBUGGER_WORKING_DIRECTORY "${PROJECT_SOURCE_DIR}/bin")
@@ -1,5 +1,5 @@
 //
-// Copyright (c) 2017-2022 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2017-2025 Advanced Micro Devices, Inc. All rights reserved.
 //
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to deal
@@ -68,9 +68,9 @@ void PrintMessage(CONSOLE_COLOR color, const char* msg)
 {
    if(color != CONSOLE_COLOR::NORMAL)
        SetConsoleColor(color);
-    
+
    printf("%s\n", msg);
-    
+
    if (color != CONSOLE_COLOR::NORMAL)
        SetConsoleColor(CONSOLE_COLOR::NORMAL);
 }
@@ -79,9 +79,9 @@ void PrintMessage(CONSOLE_COLOR color, const wchar_t* msg)
 {
    if(color != CONSOLE_COLOR::NORMAL)
        SetConsoleColor(color);
-    
+
    wprintf(L"%s\n", msg);
-    
+
    if (color != CONSOLE_COLOR::NORMAL)
        SetConsoleColor(CONSOLE_COLOR::NORMAL);
 }
@@ -1,5 +1,5 @@
 //
-// Copyright (c) 2017-2022 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2017-2025 Advanced Micro Devices, Inc. All rights reserved.
 //
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to deal
@@ -1,3 +1,25 @@
+#
+# Copyright (c) 2017-2025 Advanced Micro Devices, Inc. All rights reserved.
+#
+# Permission is hereby granted, free of charge, to any person obtaining a copy
+# of this software and associated documentation files (the "Software"), to deal
+# in the Software without restriction, including without limitation the rights
+# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+# copies of the Software, and to permit persons to whom the Software is
+# furnished to do so, subject to the following conditions:
+#
+# The above copyright notice and this permission notice shall be included in
+# all copies or substantial portions of the Software.
+#
+# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+# THE SOFTWARE.
+#
+
 # This file will only be executed if VMA_BUILD_SAMPLE_SHADERS is set to ON

 find_program(GLSL_VALIDATOR glslangValidator REQUIRED)
@@ -17,16 +39,24 @@ foreach(SHADER ${SHADERS})
    get_filename_component(FILE_NAME ${SHADER} NAME)
    
    # Put the .spv files into the bin folder
+    set(SPIRV_BIN ${CMAKE_CURRENT_BINARY_DIR}/${FILE_NAME}.spv)
    set(SPIRV ${PROJECT_SOURCE_DIR}/bin/${FILE_NAME}.spv)

+
    add_custom_command(
        OUTPUT ${SPIRV}
        # Use the same file name and append .spv to the compiled shader
        COMMAND ${GLSL_VALIDATOR} -V ${CMAKE_CURRENT_SOURCE_DIR}/${SHADER} -o ${SPIRV}
        DEPENDS ${SHADER}
    )
+    add_custom_command(
+        OUTPUT ${SPIRV_BIN}
+        # Use the same file name and append .spv to the compiled shader
+        COMMAND ${GLSL_VALIDATOR} -V ${CMAKE_CURRENT_SOURCE_DIR}/${SHADER} -o ${SPIRV_BIN}
+        DEPENDS ${SHADER}
+    )

-    list(APPEND SPIRV_FILES ${SPIRV})
+    list(APPEND SPIRV_FILES ${SPIRV} ${SPIRV_BIN})
 endforeach()

 add_custom_target(VmaSampleShaders ALL DEPENDS ${SPIRV_FILES})
@@ -1,5 +1,5 @@
 //
-// Copyright (c) 2017-2022 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2017-2025 Advanced Micro Devices, Inc. All rights reserved.
 //
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to deal
@@ -1,5 +1,5 @@
 //
-// Copyright (c) 2017-2022 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2017-2025 Advanced Micro Devices, Inc. All rights reserved.
 //
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to deal
@@ -1,5 +1,5 @@
 //
-// Copyright (c) 2018-2022 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2018-2025 Advanced Micro Devices, Inc. All rights reserved.
 //
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to deal
@@ -1,5 +1,5 @@
 //
-// Copyright (c) 2017-2022 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2017-2025 Advanced Micro Devices, Inc. All rights reserved.
 //
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to deal
@@ -153,7 +153,7 @@ void BaseImage::UploadContent()
    VmaAllocation srcBufAlloc = nullptr;
    VmaAllocationInfo srcAllocInfo = {};
    TEST( vmaCreateBuffer(g_hAllocator, &srcBufCreateInfo, &srcBufAllocCreateInfo, &srcBuf, &srcBufAlloc, &srcAllocInfo) == VK_SUCCESS );
-    
+
    // Fill texels with: r = x % 255, g = u % 255, b = 13, a = 25
    uint32_t* srcBufPtr = (uint32_t*)srcAllocInfo.pMappedData;
    for(uint32_t y = 0, sizeY = m_CreateInfo.extent.height; y < sizeY; ++y)
@@ -211,7 +211,7 @@ void BaseImage::UploadContent()
        vkCmdCopyBufferToImage(g_hTemporaryCommandBuffer, srcBuf, m_Image,
            VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL, 1, &region);
    }
-    
+
    // Barrier transfer dst to fragment shader read only.
    {
        VkImageMemoryBarrier barrier = { VK_STRUCTURE_TYPE_IMAGE_MEMORY_BARRIER };
@@ -441,7 +441,7 @@ void TraditionalImage::Init(RandomNumberGenerator& rand)
    allocCreateInfo.usage = VMA_MEMORY_USAGE_AUTO;
    // Default BEST_FIT is clearly better.
    //allocCreateInfo.flags |= VMA_ALLOCATION_CREATE_STRATEGY_WORST_FIT_BIT;
-    
+
    ERR_GUARD_VULKAN( vmaCreateImage(g_hAllocator, &m_CreateInfo, &allocCreateInfo,
        &m_Image, &m_Allocation, nullptr) );
 }
@@ -470,15 +470,6 @@ void SparseBindingImage::Init(RandomNumberGenerator& rand)
    VkMemoryRequirements imageMemReq;
    vkGetImageMemoryRequirements(g_hDevice, m_Image, &imageMemReq);

-    // This is just to silence validation layer warning.
-    // But it doesn't help. Looks like a bug in Vulkan validation layers.
-    // See: https://github.com/KhronosGroup/Vulkan-ValidationLayers/issues/364
-    uint32_t sparseMemReqCount = 0;
-    vkGetImageSparseMemoryRequirements(g_hDevice, m_Image, &sparseMemReqCount, nullptr);
-    TEST(sparseMemReqCount <= 8);
-    VkSparseImageMemoryRequirements sparseMemReq[8];
-    vkGetImageSparseMemoryRequirements(g_hDevice, m_Image, &sparseMemReqCount, sparseMemReq);
-
    // According to Vulkan specification, for sparse resources memReq.alignment is also page size.
    const VkDeviceSize pageSize = imageMemReq.alignment;
    const uint32_t pageCount = (uint32_t)ceil_div<VkDeviceSize>(imageMemReq.size, pageSize);
@@ -513,7 +504,7 @@ void SparseBindingImage::Init(RandomNumberGenerator& rand)
    VkBindSparseInfo bindSparseInfo = { VK_STRUCTURE_TYPE_BIND_SPARSE_INFO };
    bindSparseInfo.pImageOpaqueBinds = &imageBindInfo;
    bindSparseInfo.imageOpaqueBindCount = 1;
-    
+
    ERR_GUARD_VULKAN( vkResetFences(g_hDevice, 1, &g_ImmediateFence) );
    ERR_GUARD_VULKAN( vkQueueBindSparse(g_hSparseBindingQueue, 1, &bindSparseInfo, g_ImmediateFence) );
    ERR_GUARD_VULKAN( vkWaitForFences(g_hDevice, 1, &g_ImmediateFence, VK_TRUE, UINT64_MAX) );
@@ -1,5 +1,5 @@
 //
-// Copyright (c) 2017-2022 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2017-2025 Advanced Micro Devices, Inc. All rights reserved.
 //
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to deal
@@ -1,5 +1,5 @@
 //
-// Copyright (c) 2017-2022 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2017-2025 Advanced Micro Devices, Inc. All rights reserved.
 //
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to deal
@@ -1,5 +1,5 @@
 //
-// Copyright (c) 2017-2022 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2017-2025 Advanced Micro Devices, Inc. All rights reserved.
 //
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to deal
@@ -1,5 +1,5 @@
 //
-// Copyright (c) 2017-2022 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2017-2025 Advanced Micro Devices, Inc. All rights reserved.
 //
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to deal
@@ -25,9 +25,15 @@

 #ifdef _WIN32

-#define NOMINMAX
-#define WIN32_LEAN_AND_MEAN
-#include <Windows.h>
+#if !defined(NOMINMAX)
+    #define NOMINMAX
+#endif
+
+#if !defined(WIN32_LEAN_AND_MEAN)
+    #define WIN32_LEAN_AND_MEAN
+#endif
+
+#include <windows.h>
 #if !defined(VK_USE_PLATFORM_WIN32_KHR)
    #define VK_USE_PLATFORM_WIN32_KHR
 #endif // #if !defined(VK_USE_PLATFORM_WIN32_KHR)
@@ -58,6 +64,7 @@ include all public interface declarations. Example:
 //#define VMA_MEMORY_BUDGET 0
 //#define VMA_STATS_STRING_ENABLED 0
 //#define VMA_MAPPING_HYSTERESIS_ENABLED 0
+//#define VMA_KHR_MAINTENANCE5 0

 //#define VMA_VULKAN_VERSION 1003000 // Vulkan 1.3
 //#define VMA_VULKAN_VERSION 1002000 // Vulkan 1.2
@@ -76,6 +83,7 @@ include all public interface declarations. Example:
 #pragma warning(disable: 4100) // unreferenced formal parameter
 #pragma warning(disable: 4189) // local variable is initialized but not referenced
 #pragma warning(disable: 4324) // structure was padded due to alignment specifier
+#pragma warning(disable: 4820) // 'X': 'N' bytes padding added after data member 'X'

 #endif  // #ifdef _MSVC_LANG

@@ -88,7 +96,17 @@ include all public interface declarations. Example:
    #pragma clang diagnostic ignored "-Wnullability-completeness"
 #endif

-#include "../include/vk_mem_alloc.h"
+#ifdef VMA_VOLK_HEADER_PATH
+    #include VMA_VOLK_HEADER_PATH
+#else
+    #include <vulkan/vulkan.h>
+#endif
+
+#ifdef _WIN32
+    #include <vulkan/vulkan_win32.h>
+#endif  // #ifdef _WIN32
+
+#include "vk_mem_alloc.h"

 #ifdef __clang__
    #pragma clang diagnostic pop
@@ -0,0 +1,28 @@
+//
+// Copyright (c) 2017-2025 Advanced Micro Devices, Inc. All rights reserved.
+//
+// Permission is hereby granted, free of charge, to any person obtaining a copy
+// of this software and associated documentation files (the "Software"), to deal
+// in the Software without restriction, including without limitation the rights
+// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the Software is
+// furnished to do so, subject to the following conditions:
+//
+// The above copyright notice and this permission notice shall be included in
+// all copies or substantial portions of the Software.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+// THE SOFTWARE.
+//
+
+#ifdef VMA_VOLK_HEADER_PATH
+
+#define VOLK_IMPLEMENTATION
+#include "VmaUsage.h"
+
+#endif // #ifdef VMA_VOLK_HEADER_PATH
--- a/Show More
+++ b/Show More