Update timer Query call back

Moved them out of DriverEnums because only the vk backend needs
them.

NEW_RELEASE_NOTES.md

@@ -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

RELEASE_NOTES.md

@@ -7,11 +7,9 @@ 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.62.0
+## v1.61.2
 
 - 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
 

android/filament-android/src/main/cpp/MaterialInstance.cpp

@@ -60,14 +60,6 @@ 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,
@@ -272,13 +264,6 @@ 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(

android/filament-android/src/main/java/com/google/android/filament/MaterialInstance.java

@@ -402,16 +402,6 @@ 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.
      *
@@ -931,9 +921,6 @@ 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);
 

build/common/upload-release-assets/package-lock.json

@@ -204,9 +204,10 @@
       "license": "Apache-2.0"
     },
     "node_modules/brace-expansion": {
-      "version": "1.1.11",
-      "resolved": "https://registry.npmjs.org/brace-expansion/-/brace-expansion-1.1.11.tgz",
-      "integrity": "sha512-iCuPHDFgrHX7H2vEI/5xpz07zSHB00TpugqhmYtVmMO6518mCuRMoOYFldEBl0g187ufozdaHgWKcYFb61qGiA==",
+      "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",
       "dependencies": {
         "balanced-match": "^1.0.0",
         "concat-map": "0.0.1"
@@ -409,9 +410,9 @@
       "integrity": "sha512-Nik3Sc0ncrMK4UUdXQmAnRtzmNQTAAXmXIopizwZ1W1t8QmfJj+zL4OA2I7XPTPW5z5TDqv4hRo/JzouDJnX3A=="
     },
     "brace-expansion": {
-      "version": "1.1.11",
-      "resolved": "https://registry.npmjs.org/brace-expansion/-/brace-expansion-1.1.11.tgz",
-      "integrity": "sha512-iCuPHDFgrHX7H2vEI/5xpz07zSHB00TpugqhmYtVmMO6518mCuRMoOYFldEBl0g187ufozdaHgWKcYFb61qGiA==",
+      "version": "1.1.12",
+      "resolved": "https://registry.npmjs.org/brace-expansion/-/brace-expansion-1.1.12.tgz",
+      "integrity": "sha512-9T9UjW3r0UW5c1Q7GTwllptXwhvYmEzFhzMfZ9H7FQWt+uZePjZPjBP/W1ZEyZ1twGWom5/56TF4lPcqjnDHcg==",
       "requires": {
         "balanced-match": "^1.0.0",
         "concat-map": "0.0.1"

docs_src/markdeep/Materials.md.html

@@ -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,37 +1091,30 @@ Value
     `bool` or `number`, depending on the `type` of the constant. The type must be one of the types
     described in table [materialConstantsTypes].
 
-    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             |
+         Type          |            Description                   |      Default
+:----------------------|:-----------------------------------------|:------------------
+int                    | A signed, 32 bit GLSL int                |         0
+float                  | A single-precision GLSL float            |         0.0
+bool                   | A GLSL bool                              |         false
 [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. 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.
+    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.
 
 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ JSON
 material {
     constants : [
         {
            name : overrideAlpha,
-           type : bool,
-           mutable : true,
+           type : bool
         },
         {
            name : customAlpha,
@@ -2560,7 +2553,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": [
     {
@@ -2576,7 +2569,7 @@ $ matc --reflect parameters filament/src/materials/skybox.mat
     }
   ]
 }
-```
+``` 
 
 ### --variant-filter
 

filament/backend/CMakeLists.txt

@@ -257,6 +257,7 @@ 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
@@ -282,6 +283,8 @@ if (FILAMENT_SUPPORTS_WEBGPU)
             src/webgpu/WebGPUSwapChain.h
             src/webgpu/WebGPUTexture.cpp
             src/webgpu/WebGPUTexture.h
+            src/webgpu/WebGPUTimerQuery.cpp
+            src/webgpu/WebGPUTimerQuery.h
             src/webgpu/WebGPUVertexBuffer.cpp
             src/webgpu/WebGPUVertexBuffer.h
             src/webgpu/WebGPUVertexBufferInfo.cpp

filament/backend/include/backend/DriverEnums.h

@@ -495,18 +495,12 @@ struct DescriptorSetLayoutBinding {
     DescriptorFlags flags = DescriptorFlags::NONE;
     uint16_t count = 0;
 
-    //  TODO: uncomment when needed.  Note that this class is used as hash key.  We need to ensure
-    //  no uninitialized padding bytes.
-    //    uint8_t externalSamplerDataIndex = EXTERNAL_SAMPLER_DATA_INDEX_UNUSED;
-
     friend bool operator==(DescriptorSetLayoutBinding const& lhs,
             DescriptorSetLayoutBinding const& rhs) noexcept {
         return lhs.type == rhs.type &&
                lhs.flags == rhs.flags &&
                lhs.count == rhs.count &&
                lhs.stageFlags == rhs.stageFlags;
-//               lhs.stageFlags == rhs.stageFlags &&
-//               lhs.externalSamplerDataIndex == rhs.externalSamplerDataIndex;
     }
 };
 
@@ -1254,26 +1248,6 @@ enum class SamplerCompareFunc : uint8_t {
     N           //!< Never. The depth / stencil test always fails.
 };
 
-//! this API is copied from (and only applies to) the Vulkan spec.
-//! These specify YUV to RGB conversions.
-enum class SamplerYcbcrModelConversion : uint8_t {
-    RGB_IDENTITY = 0,
-    YCBCR_IDENTITY = 1,
-    YCBCR_709 = 2,
-    YCBCR_601 = 3,
-    YCBCR_2020 = 4,
-};
-
-enum class SamplerYcbcrRange : uint8_t {
-    ITU_FULL = 0,
-    ITU_NARROW = 1,
-};
-
-enum class ChromaLocation : uint8_t {
-    COSITED_EVEN = 0,
-    MIDPOINT = 1,
-};
-
 //! Sampler parameters
 struct SamplerParams {             // NOLINT
     SamplerMagFilter filterMag      : 1;    //!< magnification filter (NEAREST)
@@ -1342,94 +1316,9 @@ static_assert(sizeof(SamplerParams) == 4);
 static_assert(sizeof(SamplerParams) <= sizeof(uint64_t),
         "SamplerParams must be no more than 64 bits");
 
-//! Sampler parameters
-struct SamplerYcbcrConversion {// NOLINT
-    SamplerYcbcrModelConversion ycbcrModel : 4;
-    TextureSwizzle r : 4;
-    TextureSwizzle g : 4;
-    TextureSwizzle b : 4;
-    TextureSwizzle a : 4;
-    SamplerYcbcrRange ycbcrRange : 1;
-    ChromaLocation xChromaOffset : 1;
-    ChromaLocation yChromaOffset : 1;
-    SamplerMagFilter chromaFilter : 1;
-    uint8_t padding;
-
-    struct Hasher {
-        size_t operator()(const SamplerYcbcrConversion p) const noexcept {
-            // we don't use std::hash<> here, so we don't have to include <functional>
-            return *reinterpret_cast<uint32_t const*>(reinterpret_cast<char const*>(&p));
-        }
-    };
-
-    struct EqualTo {
-        bool operator()(SamplerYcbcrConversion lhs, SamplerYcbcrConversion rhs) const noexcept {
-            assert_invariant(lhs.padding == 0);
-            auto* pLhs = reinterpret_cast<uint32_t const*>(reinterpret_cast<char const*>(&lhs));
-            auto* pRhs = reinterpret_cast<uint32_t const*>(reinterpret_cast<char const*>(&rhs));
-            return *pLhs == *pRhs;
-        }
-    };
-
-    struct LessThan {
-        bool operator()(SamplerYcbcrConversion lhs, SamplerYcbcrConversion rhs) const noexcept {
-            assert_invariant(lhs.padding == 0);
-            auto* pLhs = reinterpret_cast<uint32_t const*>(reinterpret_cast<char const*>(&lhs));
-            auto* pRhs = reinterpret_cast<uint32_t const*>(reinterpret_cast<char const*>(&rhs));
-            return *pLhs < *pRhs;
-        }
-    };
-
-private:
-    friend bool operator == (SamplerYcbcrConversion lhs, SamplerYcbcrConversion rhs)
-            noexcept {
-        return SamplerYcbcrConversion::EqualTo{}(lhs, rhs);
-    }
-    friend bool operator != (SamplerYcbcrConversion lhs, SamplerYcbcrConversion rhs)
-            noexcept {
-        return  !SamplerYcbcrConversion::EqualTo{}(lhs, rhs);
-    }
-    friend bool operator < (SamplerYcbcrConversion lhs, SamplerYcbcrConversion rhs)
-            noexcept {
-        return SamplerYcbcrConversion::LessThan{}(lhs, rhs);
-    }
-};
-
-static_assert(sizeof(SamplerYcbcrConversion) == 4);
-
-static_assert(sizeof(SamplerYcbcrConversion) <= sizeof(uint64_t),
-        "SamplerYcbcrConversion must be no more than 64 bits");
-
-struct ExternalSamplerDatum {
-    ExternalSamplerDatum(SamplerYcbcrConversion ycbcr, SamplerParams spm, uint32_t extFmt)
-        : YcbcrConversion(ycbcr),
-          samplerParams(spm),
-          externalFormat(extFmt) {}
-    bool operator==(ExternalSamplerDatum const& rhs) const {
-        return (YcbcrConversion == rhs.YcbcrConversion && samplerParams == rhs.samplerParams &&
-                externalFormat == rhs.externalFormat);
-    }
-    struct EqualTo {
-        bool operator()(const ExternalSamplerDatum& lhs,
-                const ExternalSamplerDatum& rhs) const noexcept {
-            return (lhs.YcbcrConversion == rhs.YcbcrConversion &&
-                lhs.samplerParams == rhs.samplerParams &&
-                lhs.externalFormat == rhs.externalFormat);
-        }
-    };
-    SamplerYcbcrConversion YcbcrConversion;
-    SamplerParams samplerParams;
-    uint32_t externalFormat;
-};
-// No implicit padding allowed due to it being a hash key.
-static_assert(sizeof(ExternalSamplerDatum) == 12);
-
 struct DescriptorSetLayout {
     std::variant<utils::StaticString, utils::CString, std::monostate> label;
     utils::FixedCapacityVector<DescriptorSetLayoutBinding> bindings;
-
-//  TODO: uncomment when needed
-//    utils::FixedCapacityVector<ExternalSamplerDatum> externalSamplerData;
 };
 
 //! blending equation function

filament/backend/include/backend/Platform.h

@@ -149,6 +149,13 @@ 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;

filament/backend/include/backend/Program.h

@@ -20,7 +20,6 @@
 #include <utils/CString.h>
 #include <utils/FixedCapacityVector.h>
 #include <utils/Invocable.h>
-#include <utils/bitset.h>
 
 #include <backend/DriverEnums.h>
 
@@ -67,7 +66,6 @@ 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>;
 
@@ -104,8 +102,7 @@ public:
     Program& descriptorBindings(backend::descriptor_set_t set,
             DescriptorBindingsInfo descriptorBindings) noexcept;
 
-    Program& specializationConstants(SpecializationConstantsInfo specConstants,
-            uint32_t firstMutableId, MutableSpecConstantsInfo mutableSpecConstants) noexcept;
+    Program& specializationConstants(SpecializationConstantsInfo specConstants) noexcept;
 
     struct PushConstant {
         utils::CString name;

filament/backend/src/Program.cpp

@@ -80,18 +80,8 @@ Program& Program::attributes(AttributesInfo attributes) noexcept {
     return *this;
 }
 
-Program& Program::specializationConstants(SpecializationConstantsInfo specConstants,
-        uint32_t firstMutableId, MutableSpecConstantsInfo mutableSpecConstants) noexcept {
-    // String the two lists together.
+Program& Program::specializationConstants(SpecializationConstantsInfo specConstants) noexcept {
     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;
 }
 

filament/backend/src/metal/PlatformMetal.mm

@@ -53,6 +53,9 @@ 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);
 }
 

filament/backend/src/opengl/ShaderCompilerService.cpp

@@ -122,12 +122,18 @@ 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() {
@@ -259,6 +265,7 @@ 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;
     }
 
@@ -337,12 +344,15 @@ GLuint ShaderCompilerService::getProgram(program_token_t& token) {
 
     assert_invariant(token);// This function should be called when the token is still alive.
 
-    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();
+    // 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();
+            }
         }
     }
 

filament/backend/src/vulkan/VulkanYcbcrConversionCache.cpp

@@ -29,6 +29,8 @@ using namespace bluevk;
 
 namespace filament::backend {
 
+using namespace fvkutils;
+
 VulkanYcbcrConversionCache::VulkanYcbcrConversionCache(VkDevice device)
     : mDevice(device) {}
 

filament/backend/src/vulkan/VulkanYcbcrConversionCache.h

@@ -17,6 +17,8 @@
 #ifndef TNT_FILAMENT_BACKEND_VULKANYCBCRCONVERSIONCACHE_H
 #define TNT_FILAMENT_BACKEND_VULKANYCBCRCONVERSIONCACHE_H
 
+#include "utils/Definitions.h"
+
 #include <backend/DriverEnums.h>
 
 #include <utils/Hash.h>
@@ -30,7 +32,7 @@ namespace filament::backend {
 class VulkanYcbcrConversionCache {
 public:
     struct Params {
-        SamplerYcbcrConversion conversion = {}; // 4
+        fvkutils::SamplerYcbcrConversion conversion = {}; // 4
         VkFormat format;                        // 4
         uint64_t externalFormat = 0;            // 8
     };
@@ -45,16 +47,15 @@ private:
 
     struct ConversionEqualTo {
         bool operator()(Params lhs, Params rhs) const noexcept {
-            SamplerYcbcrConversion::EqualTo equal;
+            fvkutils::SamplerYcbcrConversion::EqualTo equal;
             return equal(lhs.conversion, rhs.conversion) &&
-                   lhs.externalFormat == rhs.externalFormat &&
-                   lhs.format == rhs.format;
+                   lhs.externalFormat == rhs.externalFormat && lhs.format == rhs.format;
         }
     };
     using ConversionHashFn = utils::hash::MurmurHashFn<Params>;
     tsl::robin_map<Params, VkSamplerYcbcrConversion, ConversionHashFn, ConversionEqualTo> mCache;
 };
 
-}// namespace filament::backend
+} // namespace filament::backend
 
 #endif// TNT_FILAMENT_BACKEND_VULKANYCBCRCONVERSIONCACHE_H

filament/backend/src/vulkan/utils/Conversion.h

@@ -17,6 +17,8 @@
 #ifndef TNT_FILAMENT_BACKEND_VULKAN_UTILS_CONVERSION_H
 #define TNT_FILAMENT_BACKEND_VULKAN_UTILS_CONVERSION_H
 
+#include "Definitions.h"
+
 #include <backend/DriverEnums.h>
 
 #include <private/backend/BackendUtils.h>  // for getFormatSize()

filament/backend/src/vulkan/utils/Definitions.h

@@ -17,6 +17,8 @@
 #ifndef TNT_FILAMENT_BACKEND_VULKAN_UTILS_DEFINITIONS_H
 #define TNT_FILAMENT_BACKEND_VULKAN_UTILS_DEFINITIONS_H
 
+#include <backend/DriverEnums.h>
+
 #include <utils/bitset.h>
 #include <utils/FixedCapacityVector.h>
 
@@ -379,6 +381,81 @@ static constexpr uint8_t getFragmentStageShift() noexcept {
 // We have at most 4 descriptor sets. This is to indicate which ones are active.
 using DescriptorSetMask = utils::bitset8;
 
+//! this API is copied from (and only applies to) the Vulkan spec.
+//! These specify YUV to RGB conversions.
+enum class SamplerYcbcrModelConversion : uint8_t {
+    RGB_IDENTITY = 0,
+    YCBCR_IDENTITY = 1,
+    YCBCR_709 = 2,
+    YCBCR_601 = 3,
+    YCBCR_2020 = 4,
+};
+
+enum class SamplerYcbcrRange : uint8_t {
+    ITU_FULL = 0,
+    ITU_NARROW = 1,
+};
+
+enum class ChromaLocation : uint8_t {
+    COSITED_EVEN = 0,
+    MIDPOINT = 1,
+};
+
+//! Sampler parameters
+struct SamplerYcbcrConversion { // NOLINT
+    SamplerYcbcrModelConversion ycbcrModel : 4;
+    TextureSwizzle r : 4;
+    TextureSwizzle g : 4;
+    TextureSwizzle b : 4;
+    TextureSwizzle a : 4;
+    SamplerYcbcrRange ycbcrRange : 1;
+    ChromaLocation xChromaOffset : 1;
+    ChromaLocation yChromaOffset : 1;
+    SamplerMagFilter chromaFilter : 1;
+    uint8_t padding;
+
+    struct Hasher {
+        size_t operator()(const SamplerYcbcrConversion p) const noexcept {
+            // we don't use std::hash<> here, so we don't have to include <functional>
+            return *reinterpret_cast<uint32_t const*>(reinterpret_cast<char const*>(&p));
+        }
+    };
+
+    struct EqualTo {
+        bool operator()(SamplerYcbcrConversion lhs, SamplerYcbcrConversion rhs) const noexcept {
+            assert_invariant(lhs.padding == 0);
+            auto* pLhs = reinterpret_cast<uint32_t const*>(reinterpret_cast<char const*>(&lhs));
+            auto* pRhs = reinterpret_cast<uint32_t const*>(reinterpret_cast<char const*>(&rhs));
+            return *pLhs == *pRhs;
+        }
+    };
+
+    struct LessThan {
+        bool operator()(SamplerYcbcrConversion lhs, SamplerYcbcrConversion rhs) const noexcept {
+            assert_invariant(lhs.padding == 0);
+            auto* pLhs = reinterpret_cast<uint32_t const*>(reinterpret_cast<char const*>(&lhs));
+            auto* pRhs = reinterpret_cast<uint32_t const*>(reinterpret_cast<char const*>(&rhs));
+            return *pLhs < *pRhs;
+        }
+    };
+
+private:
+    friend bool operator==(SamplerYcbcrConversion lhs, SamplerYcbcrConversion rhs) noexcept {
+        return SamplerYcbcrConversion::EqualTo{}(lhs, rhs);
+    }
+    friend bool operator!=(SamplerYcbcrConversion lhs, SamplerYcbcrConversion rhs) noexcept {
+        return !SamplerYcbcrConversion::EqualTo{}(lhs, rhs);
+    }
+    friend bool operator<(SamplerYcbcrConversion lhs, SamplerYcbcrConversion rhs) noexcept {
+        return SamplerYcbcrConversion::LessThan{}(lhs, rhs);
+    }
+};
+
+static_assert(sizeof(SamplerYcbcrConversion) == 4);
+
+static_assert(sizeof(SamplerYcbcrConversion) <= sizeof(uint64_t),
+        "SamplerYcbcrConversion must be no more than 64 bits");
+
 } // namespace filament::backend::fvkutils
 
 #endif // TNT_FILAMENT_BACKEND_VULKAN_UTILS_DEFINITIONS_H

filament/backend/src/webgpu/SpdMipmapGenerator/Readme.md

@@ -0,0 +1,5 @@
+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. 

filament/backend/src/webgpu/SpdMipmapGenerator/SpdMipmapGenerator.cpp

@@ -0,0 +1,956 @@
+#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

filament/backend/src/webgpu/SpdMipmapGenerator/SpdMipmapGenerator.h

@@ -0,0 +1,99 @@
+#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

filament/backend/src/webgpu/SpdMipmapGenerator/license.txt

@@ -0,0 +1,21 @@
+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.

filament/backend/src/webgpu/WebGPUDriver.cpp

@@ -25,6 +25,7 @@
 #include "WebGPURenderTarget.h"
 #include "WebGPUSwapChain.h"
 #include "WebGPUTexture.h"
+#include "WebGPUTimerQuery.h"
 #include "WebGPUVertexBuffer.h"
 #include "WebGPUVertexBufferInfo.h"
 #include <backend/platforms/WebGPUPlatform.h>
@@ -53,6 +54,8 @@
 #include <sstream>
 #include <utility>
 
+using namespace std::chrono_literals;
+
 namespace filament::backend {
 
 Driver* WebGPUDriver::create(WebGPUPlatform& platform, const Platform::DriverConfig& driverConfig) noexcept {
@@ -64,13 +67,14 @@ Driver* WebGPUDriver::create(WebGPUPlatform& platform, const Platform::DriverCon
 
 WebGPUDriver::WebGPUDriver(WebGPUPlatform& platform,
         const Platform::DriverConfig& driverConfig) noexcept
-    : mPlatform(platform),
-      mHandleAllocator("Handles", driverConfig.handleArenaSize,
-              driverConfig.disableHandleUseAfterFreeCheck, driverConfig.disableHeapHandleTags) {
-    mAdapter = mPlatform.requestAdapter(nullptr);
-    mDevice = mPlatform.requestDevice(mAdapter);
+    : mPlatform{ platform },
+      mAdapter{ mPlatform.requestAdapter(nullptr) },
+      mDevice{ mPlatform.requestDevice(mAdapter) },
+      mQueue{ mDevice.GetQueue() },
+      mMipMapGenerator{ mDevice },
+      mHandleAllocator{ "Handles", driverConfig.handleArenaSize,
+          driverConfig.disableHandleUseAfterFreeCheck, driverConfig.disableHeapHandleTags } {
     mDevice.GetLimits(&mDeviceLimits);
-    mQueue = mDevice.GetQueue();
 }
 
 WebGPUDriver::~WebGPUDriver() noexcept = default;
@@ -98,8 +102,9 @@ template class ConcreteDispatcher<WebGPUDriver>;
 void WebGPUDriver::terminate() {
 }
 
-void WebGPUDriver::tick(int) {
+void WebGPUDriver::tick(int /*dummy*/) {
     mDevice.Tick();
+    mAdapter.GetInstance().ProcessEvents();
 }
 
 void WebGPUDriver::beginFrame(int64_t monotonic_clock_ns,
@@ -213,9 +218,6 @@ void WebGPUDriver::destroyStream(Handle<HwStream> sh) {
     //TODO
 }
 
-void WebGPUDriver::destroyTimerQuery(Handle<HwTimerQuery> tqh) {
-}
-
 void WebGPUDriver::destroyDescriptorSetLayout(
         Handle<HwDescriptorSetLayout> descriptorSetLayoutHandle) {
     if (descriptorSetLayoutHandle) {
@@ -265,7 +267,31 @@ Handle<HwFence> WebGPUDriver::createFenceS() noexcept {
 }
 
 Handle<HwTimerQuery> WebGPUDriver::createTimerQueryS() noexcept {
-    return Handle<HwTimerQuery>((Handle<HwTimerQuery>::HandleId) mNextFakeHandle++);
+    return allocAndConstructHandle<WebGPUTimerQuery, HwTimerQuery>();
+}
+
+void WebGPUDriver::createTimerQueryR(Handle<HwTimerQuery> timerQueryHandle, int /*dummy*/) {
+    // nothing to do, timer query was constructed in createTimerQueryS
+}
+
+void WebGPUDriver::destroyTimerQuery(Handle<HwTimerQuery> timerQueryHandle) {
+    if (timerQueryHandle) {
+        destructHandle<WebGPUTimerQuery>(timerQueryHandle);
+    }
+}
+
+TimerQueryResult WebGPUDriver::getTimerQueryValue(Handle<HwTimerQuery> timerQueryHandle, uint64_t* elapsedTime) {
+    auto* timerQuery = handleCast<WebGPUTimerQuery>(timerQueryHandle);
+    return timerQuery->getQueryResult(elapsedTime) ? TimerQueryResult::AVAILABLE
+                                               : TimerQueryResult::NOT_READY;
+}
+
+void WebGPUDriver::beginTimerQuery(Handle<HwTimerQuery> timerQueryHandle) {
+    mTimerQuery = handleCast<WebGPUTimerQuery>(timerQueryHandle);
+}
+
+void WebGPUDriver::endTimerQuery(Handle<HwTimerQuery> timerQueryHandle) {
+    mTimerQuery = handleCast<WebGPUTimerQuery>(timerQueryHandle);
 }
 
 Handle<HwIndexBuffer> WebGPUDriver::createIndexBufferS() noexcept {
@@ -459,8 +485,6 @@ void WebGPUDriver::createFenceR(Handle<HwFence> fh, int) {
     //todo
 }
 
-void WebGPUDriver::createTimerQueryR(Handle<HwTimerQuery> tqh, int) {}
-
 void WebGPUDriver::createDescriptorSetLayoutR(
         Handle<HwDescriptorSetLayout> descriptorSetLayoutHandle,
         backend::DescriptorSetLayout&& info) {
@@ -529,8 +553,8 @@ bool WebGPUDriver::isTextureSwizzleSupported() {
 }
 
 bool WebGPUDriver::isTextureFormatMipmappable(const TextureFormat format) {
-    //todo
-    return true;
+    return WebGPUTexture::supportsMultipleMipLevelsViaStorageBinding(
+            WebGPUTexture::fToWGPUTextureFormat(format));
 }
 
 bool WebGPUDriver::isRenderTargetFormatSupported(const TextureFormat format) {
@@ -605,7 +629,7 @@ uint8_t WebGPUDriver::getMaxDrawBuffers() {
 }
 
 size_t WebGPUDriver::getMaxUniformBufferSize() {
-    return 16384u;
+    return mDeviceLimits.maxUniformBufferBindingSize;
 }
 
 size_t WebGPUDriver::getMaxTextureSize(const SamplerType target) {
@@ -724,10 +748,6 @@ void WebGPUDriver::setupExternalImage(void* image) {
     //todo
 }
 
-TimerQueryResult WebGPUDriver::getTimerQueryValue(Handle<HwTimerQuery> tqh, uint64_t* elapsedTime) {
-    return TimerQueryResult::ERROR;
-}
-
 void WebGPUDriver::setupExternalImage2(Platform::ExternalImageHandleRef image) {
     //todo
 }
@@ -738,59 +758,46 @@ 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);
 
-    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 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());
 
-    uint32_t mipLevelCount = wgpuTexture.GetMipLevelCount();
-    if (mipLevelCount <= 1) {
+    const uint32_t totalMipLevels = wgpuTexture.GetMipLevelCount();
+    if (totalMipLevels <= 1) {
         return;
     }
 
-    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;
+    // We will record all passes into a single command encoder.
+    wgpu::CommandEncoderDescriptor encoderDesc = {};
+    encoderDesc.label = "Mipmap Command Encoder";
+    wgpu::CommandEncoder encoder = mDevice.CreateCommandEncoder(&encoderDesc);
 
-    for (uint32_t mipLevel = 0; mipLevel < mipLevelCount - 1; ++mipLevel) {
-        wgpu::TexelCopyTextureInfo sourceCopyInfo{
-            .texture = wgpuTexture,
-            .mipLevel = mipLevel,
-            .aspect = texture->getAspect(),
-        };
+    spd::SPDPassConfig spdConfig = { .filter = spd::SPDFilter::Average,
+        .targetTexture = wgpuTexture,
+        .numMips = totalMipLevels,
+        .halfPrecision = false,
+        .sourceMipLevel = 0 };
 
-        wgpu::TexelCopyTextureInfo destinationCopyInfo{
-            .texture = wgpuTexture,
-            .mipLevel = mipLevel + 1,
-            .aspect = texture->getAspect(),
-        };
+    mMipMapGenerator.Generate(encoder, wgpuTexture, spdConfig);
 
-        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;
-    }
+    // 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);
 }
 
 void WebGPUDriver::compilePrograms(CompilerPriorityQueue priority,
@@ -853,7 +860,7 @@ void WebGPUDriver::beginRenderPass(Handle<HwRenderTarget> renderTargetHandle,
                 const uint32_t depthArrayLayer = depthInfo.layer;
                 customDepthView =
                         depthTexture->getOrMakeTextureView(depthMipLevel, depthArrayLayer);
-                customDepthFormat = depthTexture->getWebGPUFormat();
+                customDepthFormat = depthTexture->getViewFormat();
             }
         }
 
@@ -870,7 +877,7 @@ void WebGPUDriver::beginRenderPass(Handle<HwRenderTarget> renderTargetHandle,
                 const uint32_t stencilArrayLayer = stencilInfo.layer;
                 customStencilView =
                         stencilTexture->getOrMakeTextureView(stencilMipLevel, stencilArrayLayer);
-                customStencilFormat = stencilTexture->getWebGPUFormat();
+                customStencilFormat = stencilTexture->getViewFormat();
             }
         }
     }
@@ -924,12 +931,6 @@ 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);
 }
 
@@ -940,7 +941,34 @@ void WebGPUDriver::commit(Handle<HwSwapChain> sch) {
     mCommandBuffer = mCommandEncoder.Finish(&commandBufferDescriptor);
     assert_invariant(mCommandBuffer);
     mCommandEncoder = nullptr;
+    if (mTimerQuery) {
+        mTimerQuery->beginTimeElapsedQuery();
+    }
     mQueue.Submit(1, &mCommandBuffer);
+    static bool firstRender = true;
+    // For the first frame rendered, we need to make sure the work is done before presenting or we
+    // get a purple flash
+    if (firstRender) {
+        auto f = mQueue.OnSubmittedWorkDone(wgpu::CallbackMode::WaitAnyOnly,
+                [=](wgpu::QueueWorkDoneStatus) {});
+        const wgpu::Instance instance = mAdapter.GetInstance();
+        auto wStatus = instance.WaitAny(f,
+                std::chrono::duration_cast<std::chrono::nanoseconds>(1s).count());
+        if (wStatus != wgpu::WaitStatus::Success) {
+            FWGPU_LOGW << "Waiting for first frame work to finish resulted in an error"
+                       << static_cast<uint32_t>(wStatus);
+        }
+        firstRender = false;
+    } else {
+        mQueue.OnSubmittedWorkDone(wgpu::CallbackMode::AllowSpontaneous,
+            [=](wgpu::QueueWorkDoneStatus status) {
+                if (status == wgpu::QueueWorkDoneStatus::Success) {
+                    if (mTimerQuery) {
+                        mTimerQuery->endTimeElapsedQuery();
+                    }
+                }
+            });
+    }
     mCommandBuffer = nullptr;
     mTextureView = nullptr;
     assert_invariant(mSwapChain);
@@ -1151,12 +1179,6 @@ void WebGPUDriver::scissor(
     //todo
 }
 
-void WebGPUDriver::beginTimerQuery(Handle<HwTimerQuery> tqh) {
-}
-
-void WebGPUDriver::endTimerQuery(Handle<HwTimerQuery> tqh) {
-}
-
 void WebGPUDriver::resetState(int) {
     //todo
 }

filament/backend/src/webgpu/WebGPUDriver.h

@@ -29,6 +29,8 @@
 
 #include <utils/compiler.h>
 
+#include "SpdMipmapGenerator/SpdMipmapGenerator.h"
+#include <tsl/robin_map.h>
 #include <webgpu/webgpu_cpp.h>
 
 #include <cstdint>
@@ -41,6 +43,7 @@
 namespace filament::backend {
 
 class WebGPUSwapChain;
+class WebGPUTimerQuery;
 
 /**
  * WebGPU backend (driver) implementation
@@ -70,12 +73,13 @@ 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;
+    WebGPUTimerQuery* mTimerQuery = nullptr;
 
     tsl::robin_map<uint32_t, wgpu::RenderPipeline> mPipelineMap;
 
@@ -114,6 +118,11 @@ private:
         return mHandleAllocator.allocate<D>();
     }
 
+    template<typename D, typename B, typename... ARGS>
+    Handle<B> allocAndConstructHandle(ARGS&&... args) {
+        return mHandleAllocator.allocateAndConstruct<D>(std::forward<ARGS>(args)...);
+    }
+
     template<typename D, typename B, typename... ARGS>
     D* constructHandle(Handle<B>& handle, ARGS&&... args) noexcept {
         return mHandleAllocator.construct<D>(handle, std::forward<ARGS>(args)...);

filament/backend/src/webgpu/WebGPUStrings.h

@@ -21,50 +21,27 @@
 
 #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>
-#endif
+#include <string>
 #include <string_view>
 
 /**
  * Reusable set of convenience functions for strings -- generally string views, literals, &
- * streaming -- used in the WebGPU backend
+ * strings -- 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>
-utils::io::ostream& streamInsertWebGPUPrintable(utils::io::ostream& out,
-        const WebGPUPrintable printable) {
-    std::stringstream printableStream;
-    printableStream << printable;
-    out << printableStream.str();
-    return out;
+[[nodiscard]] inline std::string webGPUPrintableToString(const WebGPUPrintable printable) {
+    std::stringstream out;
+    out << printable;
+    return out.str();
 }
 #endif
 
@@ -115,46 +92,17 @@ utils::io::ostream& streamInsertWebGPUPrintable(utils::io::ostream& 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 {
+[[nodiscard]] inline std::string adapterOptionsToString(
+        wgpu::RequestAdapterOptions const& options) {
+    std::stringstream out;
     out << "power preference " << powerPreferenceToString(options.powerPreference)
         << " force fallback adapter " << bool(options.forceFallbackAdapter) << " backend type "
         << backendTypeToString(options.backendType);
-    return out;
+    return out.str();
 }
 
-/**
- * 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 {
+[[nodiscard]] std::string adapterInfoToString(wgpu::AdapterInfo const& info) {
+    std::stringstream out;
     out << "vendor (" << info.vendorID << ") '" << info.vendor
         << "' device (" << info.deviceID << ") '" << info.device
         << "' adapter " << adapterTypeToString(info.adapterType)
@@ -162,7 +110,7 @@ STREAM_TYPE& streamInsertRequestAdapterInfo(STREAM_TYPE& out,
         << " architecture '" << info.architecture
         << "' subgroupMinSize " << info.subgroupMinSize
         << " subgroupMaxSize " << info.subgroupMaxSize;
-    return out;
+    return out.str();
 }
 
 [[nodiscard]] constexpr std::string_view deviceLostReasonToString(
@@ -175,7 +123,7 @@ STREAM_TYPE& streamInsertRequestAdapterInfo(STREAM_TYPE& out,
     }
 }
 
-[[nodiscard]] constexpr std::string_view filamentShaderStageToString(ShaderStage stage) {
+[[nodiscard]] constexpr std::string_view filamentShaderStageToString(const ShaderStage stage) {
     switch (stage) {
         case ShaderStage::VERTEX:   return "vertex";
         case ShaderStage::FRAGMENT: return "fragment";

filament/backend/src/webgpu/WebGPUSwapChain.cpp

@@ -24,43 +24,25 @@
 #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: " << capabilities.usages;
+    FWGPU_LOGI << "  surface usages: " << webGPUPrintableToString(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 << "    " << format;
+                    FWGPU_LOGI << "    " << webGPUPrintableToString(format);
                 });
     }
     FWGPU_LOGI << "  surface present modes (" << capabilities.presentModeCount << "):";
@@ -68,7 +50,7 @@ void printSurfaceCapabilitiesDetails(wgpu::SurfaceCapabilities const& capabiliti
         std::for_each(capabilities.presentModes,
                 capabilities.presentModes + capabilities.presentModeCount,
                 [](wgpu::PresentMode const presentMode) {
-                    FWGPU_LOGI << "    " << presentMode;
+                    FWGPU_LOGI << "    " << webGPUPrintableToString(presentMode);
                 });
     }
     FWGPU_LOGI << "  surface alpha modes (" << capabilities.alphaModeCount << "):";
@@ -76,7 +58,7 @@ void printSurfaceCapabilitiesDetails(wgpu::SurfaceCapabilities const& capabiliti
         std::for_each(capabilities.alphaModes,
                 capabilities.alphaModes + capabilities.alphaModeCount,
                 [](wgpu::CompositeAlphaMode const alphaMode) {
-                    FWGPU_LOGI << "    " << alphaMode;
+                    FWGPU_LOGI << "    " << webGPUPrintableToString(alphaMode);
                 });
     }
 }
@@ -84,20 +66,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: " << config.format;
-    FWGPU_LOGI << "  surface usage: " << config.usage;
+    FWGPU_LOGI << "  surface format: " << webGPUPrintableToString(config.format);
+    FWGPU_LOGI << "  surface usage: " << webGPUPrintableToString(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 << "    " << viewFormat;
+                    FWGPU_LOGI << "    " << webGPUPrintableToString(viewFormat);
                 });
     }
-    FWGPU_LOGI << "  surface alpha mode: " << config.alphaMode;
+    FWGPU_LOGI << "  surface alpha mode: " << webGPUPrintableToString(config.alphaMode);
     FWGPU_LOGI << "  surface width: " << config.width;
     FWGPU_LOGI << "  surface height: " << config.height;
-    FWGPU_LOGI << "  surface present mode: " << config.presentMode;
-    FWGPU_LOGI << "WebGPU selected depth format: " << depthFormat;
+    FWGPU_LOGI << "  surface present mode: " << webGPUPrintableToString(config.presentMode);
+    FWGPU_LOGI << "WebGPU selected depth format: " << webGPUPrintableToString(depthFormat);
 }
 #endif// FWGPU_ENABLED(FWGPU_PRINT_SYSTEM)
 
@@ -257,8 +239,6 @@ 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),

filament/backend/src/webgpu/WebGPUTexture.cpp

@@ -25,9 +25,12 @@
 #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)
+#include <webgpu/webgpu_cpp_print.h>
+#include <sstream>
+#endif
 
 #include <cstdint>
 
@@ -59,21 +62,108 @@ namespace {
     }
 }
 
+/**
+ * @param format texture format to potentially be used for storage binding
+ * @return true if the format is compatible with wgpu::TextureUsage::StorageBinding
+ */
+[[nodiscard]] constexpr bool isFormatStorageCompatible(const wgpu::TextureFormat format) {
+    switch (format) {
+        // List of formats that support storage binding
+        case wgpu::TextureFormat::R32Float:
+        case wgpu::TextureFormat::R32Sint:
+        case wgpu::TextureFormat::R32Uint:
+        case wgpu::TextureFormat::RG32Float:
+        case wgpu::TextureFormat::RG32Sint:
+        case wgpu::TextureFormat::RG32Uint:
+        case wgpu::TextureFormat::RGBA16Float:
+        case wgpu::TextureFormat::RGBA16Sint:
+        case wgpu::TextureFormat::RGBA16Uint:
+        case wgpu::TextureFormat::RGBA32Float:
+        case wgpu::TextureFormat::RGBA32Sint:
+        case wgpu::TextureFormat::RGBA32Uint:
+        case wgpu::TextureFormat::RGBA8Unorm:
+        case wgpu::TextureFormat::RGBA8Snorm:
+        case wgpu::TextureFormat::RGBA8Uint:
+        case wgpu::TextureFormat::RGBA8Sint:
+            return true;
+        default:
+            // All other formats, including packed floats (RG11B10Ufloat),
+            // depth/stencil, and sRGB formats do not support storage.
+            return false;
+    }
+}
+
+/**
+ * @param viewFormat potential view format to a texture
+ * @return an underlying texture format that is storage binding compatible for which this view
+ *         texture format can be used. If no such format exists, Undefined is returned.
+ *         If the view texture is already storage binding compatible then the view texture itself
+ *         is returned.
+ */
+[[nodiscard]] constexpr wgpu::TextureFormat storageBindingCompatibleFormatForViewFormat(
+        const wgpu::TextureFormat viewFormat) {
+    switch (viewFormat) {
+        case wgpu::TextureFormat::RGBA8UnormSrgb:      return wgpu::TextureFormat::RGBA8Unorm;
+        case wgpu::TextureFormat::BGRA8UnormSrgb:      return wgpu::TextureFormat::BGRA8Unorm;
+        case wgpu::TextureFormat::BC1RGBAUnormSrgb:    return wgpu::TextureFormat::BC1RGBAUnorm;
+        case wgpu::TextureFormat::BC2RGBAUnormSrgb:    return wgpu::TextureFormat::BC2RGBAUnorm;
+        case  wgpu::TextureFormat::BC3RGBAUnormSrgb:   return wgpu::TextureFormat::BC3RGBAUnorm;
+        case wgpu::TextureFormat::BC7RGBAUnormSrgb:    return wgpu::TextureFormat::BC7RGBAUnorm;
+        case wgpu::TextureFormat::ETC2RGB8UnormSrgb:   return wgpu::TextureFormat::ETC2RGB8Unorm;
+        case wgpu::TextureFormat::ETC2RGB8A1UnormSrgb: return wgpu::TextureFormat::ETC2RGB8A1Unorm;
+        case wgpu::TextureFormat::ETC2RGBA8UnormSrgb:  return wgpu::TextureFormat::ETC2RGBA8Unorm;
+        case wgpu::TextureFormat::ASTC4x4UnormSrgb:    return wgpu::TextureFormat::ASTC4x4Unorm;
+        case wgpu::TextureFormat::ASTC5x4UnormSrgb:    return wgpu::TextureFormat::ASTC5x4Unorm;
+        case wgpu::TextureFormat::ASTC5x5UnormSrgb:    return wgpu::TextureFormat::ASTC5x5Unorm;
+        case wgpu::TextureFormat::ASTC6x5UnormSrgb:    return wgpu::TextureFormat::ASTC6x5Unorm;
+        case wgpu::TextureFormat::ASTC6x6UnormSrgb:    return wgpu::TextureFormat::ASTC6x6Unorm;
+        case wgpu::TextureFormat::ASTC8x5UnormSrgb:    return wgpu::TextureFormat::ASTC8x5Unorm;
+        case wgpu::TextureFormat::ASTC8x6UnormSrgb:    return wgpu::TextureFormat::ASTC8x6Unorm;
+        case wgpu::TextureFormat::ASTC8x8UnormSrgb:    return wgpu::TextureFormat::ASTC8x8Unorm;
+        case wgpu::TextureFormat::ASTC10x5UnormSrgb:   return wgpu::TextureFormat::ASTC10x5Unorm;
+        case wgpu::TextureFormat::ASTC10x6UnormSrgb:   return wgpu::TextureFormat::ASTC10x6Unorm;
+        case wgpu::TextureFormat::ASTC10x8UnormSrgb:   return wgpu::TextureFormat::ASTC10x8Unorm;
+        case wgpu::TextureFormat::ASTC10x10UnormSrgb:  return wgpu::TextureFormat::ASTC10x10Unorm;
+        case wgpu::TextureFormat::ASTC12x10UnormSrgb:  return wgpu::TextureFormat::ASTC12x10Unorm;
+        case wgpu::TextureFormat::ASTC12x12UnormSrgb:  return wgpu::TextureFormat::ASTC12x12Unorm;
+        default:
+            if (isFormatStorageCompatible(viewFormat)) {
+                return viewFormat; // view format is itself already storage binding compatible
+            } else {
+                return wgpu::TextureFormat::Undefined; // no valid format in this case
+            }
+    }
+}
+
+/**
+ * @param fUsage Filament's requested texture usage
+ * @param samples How many samples to use for MSAA
+ * @param needsComputeStorageSupport if we need to use this texture as storage binding in something
+ *                                   like a compute shader
+ * @return The appropriate texture usage flags for the underlying texture
+ */
 [[nodiscard]] wgpu::TextureUsage fToWGPUTextureUsage(TextureUsage const& fUsage,
-        const uint8_t samples) {
+        const uint8_t samples, const bool needsComputeStorageSupport) {
     wgpu::TextureUsage retUsage = wgpu::TextureUsage::None;
 
+    // if needsComputeStorageSupport we need to read and write to the texture in a shader and thus
+    // require CopySrc, CopyDst, TextureBinding, & StorageBinding
+
     // Basing this mapping off of VulkanTexture.cpp's getUsage func and suggestions from Gemini
     // TODO Validate assumptions, revisit if issues.
-    if (any(TextureUsage::BLIT_SRC & fUsage)) {
+    if (needsComputeStorageSupport || any(TextureUsage::BLIT_SRC & fUsage)) {
         retUsage |= wgpu::TextureUsage::CopySrc;
     }
-    if (any((TextureUsage::BLIT_DST | TextureUsage::UPLOADABLE) & fUsage)) {
+    if (needsComputeStorageSupport ||
+            any((TextureUsage::BLIT_DST | TextureUsage::UPLOADABLE) & fUsage)) {
         retUsage |= wgpu::TextureUsage::CopyDst;
     }
-    if (any(TextureUsage::SAMPLEABLE & fUsage)) {
+    if (needsComputeStorageSupport || any(TextureUsage::SAMPLEABLE & fUsage)) {
         retUsage |= wgpu::TextureUsage::TextureBinding;
     }
+    if (needsComputeStorageSupport) {
+        retUsage |= wgpu::TextureUsage::StorageBinding;
+    }
     // WGPU Render attachment covers either color or stencil situation dependant
     // NOTE: Depth attachment isn't used this way in Vulkan but logically maps to WGPU docs. If
     // issues, investigate here.
@@ -100,7 +190,6 @@ namespace {
         retUsage |= wgpu::TextureUsage::TransientAttachment;
     }
     // NOTE: Unused wgpu flags:
-    //  StorageBinding
     //  StorageAttachment
 
     // NOTE: Unused Filament flags:
@@ -205,11 +294,19 @@ WebGPUTexture::WebGPUTexture(const SamplerType samplerType, const uint8_t levels
         const uint32_t height, const uint32_t depth, const TextureUsage usage,
         wgpu::Device const& device) noexcept
     : HwTexture{ samplerType, levels, samples, width, height, depth, format, usage },
-      mWebGPUFormat{ fToWGPUTextureFormat(format) },
+      mViewFormat{ fToWGPUTextureFormat(format) },
+      mSupportsMultipleMipLevels{ levels > 1 &&
+                                  supportsMultipleMipLevelsViaStorageBinding(mViewFormat) },
+      mWebGPUFormat{ mSupportsMultipleMipLevels
+                             ? storageBindingCompatibleFormatForViewFormat(mViewFormat)
+                             : mViewFormat },
       mAspect{ fToWGPUTextureViewAspect(usage, format) },
-      mWebGPUUsage{ fToWGPUTextureUsage(usage, samples) },
+      mWebGPUUsage{ fToWGPUTextureUsage(usage, samples, mSupportsMultipleMipLevels) },
+      mViewUsage{ fToWGPUTextureUsage(usage, samples, false) },
       mBlockWidth{ filament::backend::getBlockWidth(format) },
-      mBlockHeight{ filament::backend::getBlockHeight(format) } {
+      mBlockHeight{ filament::backend::getBlockHeight(format) },
+      mDefaultMipLevel{ 0 },
+      mDefaultBaseArrayLayer{ 0 } {
     assert_invariant(
             samples == 1 ||
             samples == 4 &&
@@ -225,10 +322,8 @@ WebGPUTexture::WebGPUTexture(const SamplerType samplerType, const uint8_t levels
         .format = mWebGPUFormat,
         .mipLevelCount = levels,
         .sampleCount = samples,
-        // TODO Is this fine? Could do all-the-things, a naive mapping or get something from
-        // Filament
-        .viewFormatCount = 0,
-        .viewFormats = nullptr,
+        .viewFormatCount = 1,
+        .viewFormats = &mViewFormat,
     };
     mArrayLayerCount = toArrayLayerCount(samplerType, textureDescriptor.size.depthOrArrayLayers);
     assert_invariant(textureDescriptor.format != wgpu::TextureFormat::Undefined &&
@@ -240,23 +335,40 @@ WebGPUTexture::WebGPUTexture(const SamplerType samplerType, const uint8_t levels
             mArrayLayerCount, samplerType);
     FILAMENT_CHECK_POSTCONDITION(mDefaultTextureView)
             << "Failed to create default texture view for " << textureDescriptor.label;
+#if FWGPU_ENABLED(FWGPU_PRINT_SYSTEM)
+    if (mViewFormat != mWebGPUFormat) {
+        std::stringstream viewFormatStream;
+        viewFormatStream << mViewFormat;
+        std::stringstream textureFormatStream;
+        textureFormatStream << mWebGPUFormat;
+        FWGPU_LOGD << "Texture '" << textureDescriptor.label << "' has view format "
+                   << viewFormatStream.str() << " and texture format " << textureFormatStream.str();
+    }
+#endif
 }
 
 WebGPUTexture::WebGPUTexture(WebGPUTexture const* src, const uint8_t baseLevel,
         const uint8_t levelCount) noexcept
     : HwTexture{ src->target, levelCount, src->samples, src->width, src->height, src->depth,
           src->format, src->usage },
-      mTexture{ src->mTexture },
+      mViewFormat{ src->mViewFormat },
+      mSupportsMultipleMipLevels{ src->mSupportsMultipleMipLevels },
       mWebGPUFormat{ src->mWebGPUFormat },
       mAspect{ src->mAspect },
       mWebGPUUsage{ src->mWebGPUUsage },
-      mArrayLayerCount{ src->mArrayLayerCount },
+      mViewUsage{ src->mViewUsage },
       mBlockWidth{ src->mBlockWidth },
       mBlockHeight{ src->mBlockHeight },
-      mDefaultTextureView{ makeTextureView(baseLevel, levelCount, 0, src->mArrayLayerCount,
-              src->target) },
+      mArrayLayerCount{ src->mArrayLayerCount },
+      mTexture{ src->mTexture },
       mDefaultMipLevel{ baseLevel },
-      mDefaultBaseArrayLayer{ src->mArrayLayerCount } {}
+      mDefaultBaseArrayLayer{ src->mArrayLayerCount },
+      mDefaultTextureView{ makeTextureView(mDefaultMipLevel, levelCount, 0, mDefaultBaseArrayLayer,
+              src->target) } {}
+
+bool WebGPUTexture::supportsMultipleMipLevelsViaStorageBinding(const wgpu::TextureFormat format) {
+    return storageBindingCompatibleFormatForViewFormat(format) != wgpu::TextureFormat::Undefined;
+}
 
 wgpu::TextureView WebGPUTexture::getOrMakeTextureView(const uint8_t mipLevel,
         const uint32_t arrayLayer) {
@@ -356,51 +468,43 @@ 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)"
-                       << utils::io::endl;
+                          "wgpu::TextureFormat::Undefined (no direct mapping in wgpu)";
             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)"
-                       << utils::io::endl;
+                          "wgpu::TextureFormat::Undefined (no direct mapping in wgpu)";
             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)"
-                       << utils::io::endl;
+                          "wgpu::TextureFormat::Undefined (no direct mapping in wgpu)";
             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)"
-                       << utils::io::endl;
+                          "without alpha)";
             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)"
-                       << utils::io::endl;
+                          "without alpha)";
             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)"
-                    << utils::io::endl;
+                       "wgpu::TextureFormat::RGBA8Snorm (no direct mapping in wgpu without alpha)";
             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)"
-                       << utils::io::endl;
+                          "wgpu::TextureFormat::RGBA8Uint (no direct mapping in wgpu without alpha)";
             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)"
-                       << utils::io::endl;
+                          "wgpu::TextureFormat::RGBA8Sint (no direct mapping in wgpu without alpha)";
             return wgpu::TextureFormat::RGBA8Sint;
         case TextureFormat::R11F_G11F_B10F:          return wgpu::TextureFormat::RG11B10Ufloat;
         case TextureFormat::UNUSED:                  return wgpu::TextureFormat::Undefined;
@@ -408,38 +512,32 @@ 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)"
-                    << utils::io::endl;
+                       "wgpu::TextureFormat::RGBA16Float (no direct mapping in wgpu without alpha)";
             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)"
-                    << utils::io::endl;
+                       "wgpu::TextureFormat::RGBA16Uint (no direct mapping in wgpu without alpha)";
             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)"
-                    << utils::io::endl;
+                       "wgpu::TextureFormat::RGBA16Sint (no direct mapping in wgpu without alpha)";
             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)"
-                    << utils::io::endl;
+                       "wgpu::TextureFormat::RGBA32Float (no direct mapping in wgpu without alpha)";
             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)"
-                    << utils::io::endl;
+                       "wgpu::TextureFormat::RGBA32Uint (no direct mapping in wgpu without alpha)";
             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)"
-                    << utils::io::endl;
+                       "wgpu::TextureFormat::RGBA32Sint (no direct mapping in wgpu without alpha)";
             return wgpu::TextureFormat::RGBA32Sint;
         case TextureFormat::DXT1_RGB:                return wgpu::TextureFormat::BC1RGBAUnorm;
         case TextureFormat::DXT1_RGBA:               return wgpu::TextureFormat::BC1RGBAUnorm;
@@ -455,16 +553,23 @@ wgpu::TextureFormat WebGPUTexture::fToWGPUTextureFormat(TextureFormat const& fFo
 wgpu::TextureView WebGPUTexture::makeTextureView(const uint8_t& baseLevel,
         const uint8_t& levelCount, const uint32_t& baseArrayLayer, const uint32_t& arrayLayerCount,
         const SamplerType samplerType) const noexcept {
+#if FWGPU_ENABLED(FWGPU_DEBUG_VALIDATION)
+    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.";
+    }
+#endif
     const wgpu::TextureViewDescriptor textureViewDescriptor{
         .label = getUserTextureViewLabel(target),
-        .format = mWebGPUFormat,
+        .format = mViewFormat,
         .dimension = toWebGPUTextureViewDimension(samplerType),
         .baseMipLevel = baseLevel,
         .mipLevelCount = levelCount,
         .baseArrayLayer = baseArrayLayer,
         .arrayLayerCount = arrayLayerCount,
         .aspect = mAspect,
-        .usage = mWebGPUUsage };
+        .usage = mViewUsage };
     wgpu::TextureView textureView = mTexture.CreateView(&textureViewDescriptor);
     FILAMENT_CHECK_POSTCONDITION(textureView)
             << "Failed to create texture view " << textureViewDescriptor.label;

filament/backend/src/webgpu/WebGPUTexture.h

@@ -47,29 +47,46 @@ public:
 
     [[nodiscard]] wgpu::TextureView getOrMakeTextureView(uint8_t mipLevel, uint32_t arrayLayer);
 
-    [[nodiscard]] wgpu::TextureFormat getWebGPUFormat() const { return mWebGPUFormat; }
+    [[nodiscard]] wgpu::TextureFormat getViewFormat() const { return mViewFormat; }
 
     [[nodiscard]] uint32_t getArrayLayerCount() const { return mArrayLayerCount; }
 
+    [[nodiscard]] bool supportsMultipleMipLevels() const { return mSupportsMultipleMipLevels; }
+
     [[nodiscard]] static wgpu::TextureFormat fToWGPUTextureFormat(
             filament::backend::TextureFormat const& fFormat);
 
+    /**
+     * @param format a required texture format (can be a view to a different underlying texture
+     * format)
+     * @return true if multiple mip levels can be generated via a compute shader + storage binding
+     * on the texture, false otherwise.
+     */
+    [[nodiscard]] static bool supportsMultipleMipLevelsViaStorageBinding(
+            wgpu::TextureFormat format);
+
 private:
-    // CreateTextureR has info for a texture and sampler. Texture Views are needed for binding,
-    // along with a sampler Current plan: Inherit the sampler and Texture to always exist (It is a
-    // ref counted pointer) when making views. View is optional
-    wgpu::Texture mTexture = nullptr;
+    // the texture view's format, which may differ from the underlying texture's format itself,
+    // an example of when this is needed is when certain underlying shaders don't support
+    // certain formats (e.g. compute shaders), but we need a view to that format elsewhere
+    // (e.g. another shader or render target)
+    wgpu::TextureFormat mViewFormat = wgpu::TextureFormat::Undefined;
+    // this is true IFF levels > 1 AND we can actually generate mipmaps for the texture
+    bool mSupportsMultipleMipLevels = false;
     // format is inherited from HwTexture. This naming is to distinguish it from Filament's format
+    // this is the underlying texture's format, not necessarily the "view" of it... see mViewFormat
     wgpu::TextureFormat mWebGPUFormat = wgpu::TextureFormat::Undefined;
     wgpu::TextureAspect mAspect = wgpu::TextureAspect::Undefined;
     // usage is inherited from HwTexture. This naming is to distinguish it from Filament's usage
     wgpu::TextureUsage mWebGPUUsage = wgpu::TextureUsage::None;
-    uint32_t mArrayLayerCount = 1;
+    wgpu::TextureUsage mViewUsage = wgpu::TextureUsage::None;
     size_t mBlockWidth = 0;
     size_t mBlockHeight = 0;
-    wgpu::TextureView mDefaultTextureView = nullptr;
+    uint32_t mArrayLayerCount = 1;
+    wgpu::Texture mTexture = nullptr;
     uint32_t mDefaultMipLevel = 0;
     uint32_t mDefaultBaseArrayLayer = 0;
+    wgpu::TextureView mDefaultTextureView = nullptr;
 
     [[nodiscard]] wgpu::TextureView makeTextureView(const uint8_t& baseLevel,
             const uint8_t& levelCount, const uint32_t& baseArrayLayer,

filament/backend/src/webgpu/WebGPUTimerQuery.cpp

@@ -0,0 +1,60 @@
+/*
+ * 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 "WebGPUTimerQuery.h"
+
+#include <chrono>
+#include <cstdint>
+#include <memory>
+
+namespace filament::backend {
+
+void WebGPUTimerQuery::beginTimeElapsedQuery() {
+    mStatus->elapsedNanoseconds = 0;
+    // Capture the timer query status via a weak_ptr because the WGPUTimerQuery could be destroyed
+    // before the block executes.
+    std::weak_ptr<WebGPUTimerQuery::Status> statusPtr = mStatus;
+
+    if (auto s = statusPtr.lock()) {
+        s->elapsedNanoseconds = std::chrono::steady_clock::now().time_since_epoch().count();
+    }
+}
+
+void WebGPUTimerQuery::endTimeElapsedQuery() {
+    // Capture the timer query status via a weak_ptr because the WGPUTimerQuery could be destroyed
+    // before the block executes.
+    if (mStatus->elapsedNanoseconds != 0) {
+        std::weak_ptr<WebGPUTimerQuery::Status> statusPtr = mStatus;
+        if (auto s = statusPtr.lock()) {
+            s->previousElapsed = s->elapsedNanoseconds =
+                    std::chrono::steady_clock::now().time_since_epoch().count() -
+                    s->elapsedNanoseconds;
+        }
+    }
+}
+
+bool WebGPUTimerQuery::getQueryResult(uint64_t* outElapsedTime) {
+    if (mStatus->previousElapsed == 0) {
+        return false;
+    }
+    if (outElapsedTime) {
+        *outElapsedTime = mStatus->previousElapsed;
+        mStatus->previousElapsed = 0;
+    }
+    return true;
+}
+
+}// namespace filament::backend

filament/backend/src/webgpu/WebGPUTimerQuery.h

@@ -0,0 +1,50 @@
+/*
+ * 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_FILAMENT_BACKEND_WEBGPUTIMERQUERY_H
+#define TNT_FILAMENT_BACKEND_WEBGPUTIMERQUERY_H
+
+#include "DriverBase.h"
+
+#include <webgpu/webgpu_cpp.h>
+
+#include <atomic>
+#include <cstdint>
+#include <memory>
+
+namespace filament::backend {
+
+class WebGPUTimerQuery : public HwTimerQuery {
+public:
+    WebGPUTimerQuery()
+        : mStatus(std::make_shared<Status>()) {}
+
+    void beginTimeElapsedQuery();
+    void endTimeElapsedQuery();
+    bool getQueryResult(uint64_t* outElapsedTimeNanoseconds);
+
+private:
+    struct Status {
+        std::atomic<uint64_t> elapsedNanoseconds{ 0 };
+        std::atomic<uint64_t> previousElapsed{ 0 };
+    };
+
+    std::shared_ptr<Status> mStatus;
+};
+
+} // namespace filament::backend
+
+#endif //TNT_FILAMENT_BACKEND_WEBGPUTIMERQUERY_H

filament/backend/src/webgpu/WebGPUVertexBufferInfo.cpp

@@ -22,7 +22,6 @@
 #include <backend/DriverEnums.h>
 
 #include <utils/Panic.h>
-#include <utils/ostream.h>
 
 #include <webgpu/webgpu_cpp.h>
 
@@ -57,26 +56,22 @@ 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)"
-                        << utils::io::endl;
+                           "wgpu::VertexFormat::Snorm8x4 (no direct mapping in wgpu for x3 byte)";
                 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)"
-                        << utils::io::endl;
+                           "wgpu::VertexFormat::Unorm8x4 (no direct mapping in wgpu for x3 byte)";
                 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)"
-                           << utils::io::endl;
+                              "half/short)";
                 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)"
-                           << utils::io::endl;
+                              "half/short)";
                 return VertexFormat::Unorm16x4; // NOT MINSPEC
             // Four Component Types
             case ElementType::BYTE4:   return VertexFormat::Snorm8x4;
@@ -97,26 +92,22 @@ 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)"
-                       << utils::io::endl;
+                          "wgpu::VertexFormat::Float16 (no direct mapping in wgpu for 8 bit float)";
             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)"
-                       << utils::io::endl;
+                          "wgpu::VertexFormat::Float16 (no direct mapping in wgpu for 8 bit float)";
             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?)"
-                       << utils::io::endl;
+                          "wgpu::VertexFormat::Float16 (potential loss of precision?)";
             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?)"
-                       << utils::io::endl;
+                          "wgpu::VertexFormat::Float16 (potential loss of precision?)";
             return VertexFormat::Float16;
         case ElementType::HALF:    return VertexFormat::Float16;
         case ElementType::INT:     return VertexFormat::Sint32;
@@ -127,58 +118,49 @@ 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)"
-                    << utils::io::endl;
+                       "wgpu::VertexFormat::Float16x2 (no direct mapping in wgpu for 8 bit float)";
             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)"
-                    << utils::io::endl;
+                       "wgpu::VertexFormat::Float16x2 (no direct mapping in wgpu for 8 bit float)";
             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?)"
-                       << utils::io::endl;
+                          "wgpu::VertexFormat::Float16x2 (potential loss of precision?)";
             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?)"
-                       << utils::io::endl;
+                          "wgpu::VertexFormat::Float16x2 (potential loss of precision?)";
             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)"
-                       << utils::io::endl;
+                          "wgpu::VertexFormat::Sint8x4 (no direct mapping in wgpu for x3 byte)";
             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)"
-                       << utils::io::endl;
+                          "wgpu::VertexFormat::Uint8x4 (no direct mapping in wgpu for x3 byte)";
             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)"
-                    << utils::io::endl;
+                       "wgpu::VertexFormat::Sint16x4 (no direct mapping in wgpu for x3 half/short)";
             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)"
-                    << utils::io::endl;
+                       "wgpu::VertexFormat::Uint16x4 (no direct mapping in wgpu for x3 half/short)";
             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)"
-                    << utils::io::endl;
+                       "wgpu::VertexFormat::Float16x4 (no direct mapping in wgpu for x3 half/short)";
             return VertexFormat::Float16x4; // NOT MINSPEC
         case ElementType::FLOAT3:  return VertexFormat::Float32x3;
         // Four Component Types
@@ -186,27 +168,23 @@ 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)"
-                    << utils::io::endl;
+                       "wgpu::VertexFormat::Float16x4 (no direct mapping in wgpu for 8 bit float)";
             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)"
-                    << utils::io::endl;
+                       "wgpu::VertexFormat::Float16x4 (no direct mapping in wgpu for 8 bit float)";
             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?)"
-                       << utils::io::endl;
+                          "wgpu::VertexFormat::Float16x4 (potential loss of precision?)";
             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?)"
-                       << utils::io::endl;
+                          "wgpu::VertexFormat::Float16x4 (potential loss of precision?)";
             return VertexFormat::Float16x4;
         case ElementType::HALF4:   return VertexFormat::Float16x4;
         case ElementType::FLOAT4:  return VertexFormat::Float32x4;

filament/backend/src/webgpu/platform/WebGPUPlatform.cpp

@@ -24,7 +24,6 @@
 #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>
@@ -39,6 +38,7 @@
 #include <utility>
 #include <vector>
 #if FWGPU_ENABLED(FWGPU_PRINT_SYSTEM)
+#include <string>
 #include <string_view>
 #include <variant>
 #endif
@@ -53,6 +53,9 @@ 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*/
@@ -117,28 +120,6 @@ 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) {
@@ -259,7 +240,7 @@ void printInstanceDetails(wgpu::Instance const& instance) {
                     supportedWGSLLanguageFeatures.features +
                             supportedWGSLLanguageFeatures.featureCount,
                     [](wgpu::WGSLLanguageFeatureName const featureName) {
-                        FWGPU_LOGI << "  " << featureName;
+                        FWGPU_LOGI << "  " << webGPUPrintableToString(featureName);
                     });
         }
     }
@@ -299,23 +280,22 @@ 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) {
-    FWGPU_LOGI << "  " << name.data() << ": ";
+    static constexpr std::string_view indent = "  ";
     bool undefined = true;
     if (std::holds_alternative<uint32_t>(value)) {
         if (std::get<uint32_t>(value) != WGPU_LIMIT_U32_UNDEFINED) {
             undefined = false;
-            FWGPU_LOGI << std::get<uint32_t>(value);
+            FWGPU_LOGI << indent << name.data() << ": " << 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 << std::get<uint64_t>(value);
+            FWGPU_LOGI << indent << name.data() << ": " << std::get<uint64_t>(value);
         }
     }
     if (undefined) {
-        FWGPU_LOGI << "UNDEFINED";
+        FWGPU_LOGI << indent << name.data() << ": UNDEFINED";
     }
-    FWGPU_LOGI << "";
 }
 #endif// FWGPU_ENABLED(FWGPU_PRINT_SYSTEM)
 
@@ -393,16 +373,16 @@ struct AdapterDetails final {
     }
 };
 
-template<class STREAM_TYPE>
-STREAM_TYPE& operator<<(STREAM_TYPE& out, AdapterDetails const& details) noexcept {
-    out << details.info
+[[nodiscard]] std::string toString(AdapterDetails const& details) {
+    std::stringstream out;
+    out << adapterInfoToString(details.info)
         << " power preference " << powerPreferenceToString(details.powerPreference);
-    return out;
+    return out.str();
 }
 
 #if FWGPU_ENABLED(FWGPU_PRINT_SYSTEM)
 void printAdapterDetails(AdapterDetails const& details) {
-    FWGPU_LOGI << "Selected WebGPU adapter info: " << details;
+    FWGPU_LOGI << "Selected WebGPU adapter info: " << toString(details);
     wgpu::SupportedFeatures supportedFeatures{};
     details.adapter.GetFeatures(&supportedFeatures);
     FWGPU_LOGI << "WebGPU adapter supported features (" << supportedFeatures.featureCount
@@ -411,7 +391,7 @@ void printAdapterDetails(AdapterDetails const& details) {
         std::for_each(supportedFeatures.features,
                 supportedFeatures.features + supportedFeatures.featureCount,
                 [](wgpu::FeatureName const featureName) {
-                    FWGPU_LOGI << "  " << featureName;
+                    FWGPU_LOGI << "  " << webGPUPrintableToString(featureName);
                 });
     }
     wgpu::Limits supportedLimits{};
@@ -444,12 +424,12 @@ struct AdapterDetailsHash final {
                     return details.adapter.HasFeature(featureName);
                 })) {
 #if FWGPU_ENABLED(FWGPU_PRINT_SYSTEM)
-        FWGPU_LOGI << "WebGPU adapter " << details
+        FWGPU_LOGI << "WebGPU adapter " << toString(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 << requiredFeature;
+                FWGPU_LOGI << webGPUPrintableToString(requiredFeature);
             }
         }
         FWGPU_LOGI << "";
@@ -458,10 +438,10 @@ struct AdapterDetailsHash final {
     }
     wgpu::Limits supportedLimits {};
     FILAMENT_CHECK_POSTCONDITION(details.adapter.GetLimits(&supportedLimits))
-            << "Failed to get limits for WebGPU adapter: " << details;
+            << "Failed to get limits for WebGPU adapter: " << toString(details);
     if (!satisfiesLimits(supportedLimits)) {
 #if FWGPU_ENABLED(FWGPU_PRINT_SYSTEM)
-        FWGPU_LOGI << " (for WebGPU adapter " << details << ")";
+        FWGPU_LOGI << " (for WebGPU adapter " << toString(details) << ")";
 #endif
         return false;
     }
@@ -487,14 +467,15 @@ struct AdapterDetailsHash final {
                             status != wgpu::RequestAdapterStatus::CallbackCancelled)
                             << "Failed to request a WebGPU adapter due to the request callback "
                                "being cancelled? Options: "
-                            << options << " " << message.data;
+                            << adapterOptionsToString(options) << " " << message.data;
                     FILAMENT_CHECK_POSTCONDITION(status != wgpu::RequestAdapterStatus::Error)
                             << "Failed to request a WebGPU adapter due to an error. Options: "
-                            << options << " Error: " << message.data;
+                            << adapterOptionsToString(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: " << options << ")";
+                                << "Failed to get info for adapter (options: "
+                                << adapterOptionsToString(options) << ")";
                         const std::lock_guard<std::mutex> lock(adaptersMutex);
                         compatibleAdapters.emplace(std::move(details.info), details.powerPreference,
                                 std::move(details.adapter));
@@ -509,9 +490,11 @@ 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 " << options;
+                << "Timed out requesting a WebGPU adapter with options "
+                << adapterOptionsToString(options);
         FILAMENT_CHECK_POSTCONDITION(status != wgpu::WaitStatus::Error)
-                << "Failed to request a WebGPU adapter with options " << options
+                << "Failed to request a WebGPU adapter with options "
+                << adapterOptionsToString(options)
                 << " due to an error (as request was made synchronous)";
         assert_invariant(status == wgpu::WaitStatus::Success);
     }
@@ -519,7 +502,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: " << details;
+        FWGPU_LOGI << "  WebGPU adapter: " << toString(details);
     }
 #endif
     return compatibleAdapters;
@@ -586,7 +569,7 @@ void printDeviceDetails(wgpu::Device const& device) {
         std::for_each(supportedFeatures.features,
                 supportedFeatures.features + supportedFeatures.featureCount,
                 [](wgpu::FeatureName const featureName) {
-                    FWGPU_LOGI << "  " << featureName;
+                    FWGPU_LOGI << "  " << webGPUPrintableToString(featureName);
                 });
     }
     wgpu::Limits supportedLimits{};
@@ -628,7 +611,21 @@ wgpu::Device WebGPUPlatform::requestDevice(wgpu::Adapter const& adapter) {
     deviceDescriptor.defaultQueue.label = "default_queue";
     deviceDescriptor.requiredFeatureCount = enabledFeatures.size();
     deviceDescriptor.requiredFeatures = enabledFeatures.data();
-    deviceDescriptor.requiredLimits = &REQUIRED_LIMITS;
+
+    // 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.SetDeviceLostCallback(wgpu::CallbackMode::AllowSpontaneous,
             [](wgpu::Device const&, wgpu::DeviceLostReason const& reason,
                     wgpu::StringView message) {
@@ -684,7 +681,6 @@ 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))

filament/backend/test/BackendTest.cpp

@@ -40,18 +40,24 @@ using namespace filament::math;
 
 using namespace image;
 #endif
+#include <iostream>
 
 namespace test {
 
 Backend BackendTest::sBackend = Backend::NOOP;
 OperatingSystem BackendTest::sOperatingSystem = OperatingSystem::OTHER;
 bool BackendTest::sIsMobilePlatform = false;
+int BackendTest::sArgc = 0;
+char** BackendTest::sArgv = nullptr;
 std::vector<std::string> BackendTest::sFailedImages;
 
-void BackendTest::init(Backend backend, OperatingSystem operatingSystem, bool isMobilePlatform) {
+void BackendTest::init(Backend backend, OperatingSystem operatingSystem, bool isMobilePlatform,
+        int argc, char** argv) {
     sBackend = backend;
     sOperatingSystem = operatingSystem;
     sIsMobilePlatform = isMobilePlatform;
+    sArgc = argc;
+    sArgv = argv;
 }
 
 BackendTest::BackendTest() : commandBufferQueue(CONFIG_MIN_COMMAND_BUFFERS_SIZE,
@@ -160,6 +166,11 @@ void BackendTest::markImageAsFailure(std::string failedImageName) {
     sFailedImages.emplace_back(std::move(failedImageName));
 }
 
+std::filesystem::path BackendTest::binaryDirectory() {
+    assert(sArgc >= 1);
+    return std::filesystem::path(sArgv[0]).remove_filename().string();
+}
+
 void BackendTest::recordFailedImages() {
     if (!sFailedImages.empty()) {
         std::string failedImages;
@@ -187,8 +198,9 @@ public:
     }
 };
 
-void initTests(Backend backend, OperatingSystem operatingSystem, bool isMobile, int& argc, char* argv[]) {
-    BackendTest::init(backend, operatingSystem, isMobile);
+void initTests(Backend backend, OperatingSystem operatingSystem, bool isMobile, int& argc,
+        char* argv[]) {
+    BackendTest::init(backend, operatingSystem, isMobile, argc, argv);
     ::testing::InitGoogleTest(&argc, argv);
     ::testing::AddGlobalTestEnvironment(new Environment);
 }

filament/backend/test/BackendTest.h

@@ -19,6 +19,8 @@
 
 #include <gtest/gtest.h>
 
+#include <filesystem>
+
 #include <backend/Platform.h>
 
 #include "private/backend/CommandBufferQueue.h"
@@ -31,16 +33,20 @@ namespace test {
 
 class BackendTest : public ::testing::Test {
 public:
-
-    static void init(Backend backend, OperatingSystem operatingSystem, bool isMobilePlatform);
+    static void init(Backend backend, OperatingSystem operatingSystem, bool isMobilePlatform,
+            int argc, char** argv);
 
     static Backend sBackend;
     static OperatingSystem sOperatingSystem;
     static bool sIsMobilePlatform;
+    static int sArgc;
+    static char** sArgv;
 
     // Takes the name of the image that wasn't correct, without the .png suffix
     static void markImageAsFailure(std::string failedImageName);
 
+    static std::filesystem::path binaryDirectory();
+
 protected:
 
     BackendTest();

filament/backend/test/ImageExpectations.cpp

@@ -59,28 +59,28 @@ uint32_t ScreenshotParams::expectedHash() const {
     return mExpectedPixelHash;
 }
 
-std::string ScreenshotParams::actualDirectoryPath() {
-    return "images/actual_images";
+std::filesystem::path ScreenshotParams::actualDirectoryPath() {
+    return BackendTest::binaryDirectory().append("images/actual_images");
 }
 
 std::string ScreenshotParams::actualFileName() const {
     return absl::StrFormat("%s_actual.png", mFileName);
 }
 
-std::string ScreenshotParams::actualFilePath() const {
-    return absl::StrFormat("%s/%s", actualDirectoryPath(), actualFileName());
+std::filesystem::path ScreenshotParams::actualFilePath() const {
+    return actualDirectoryPath().append(actualFileName());
 }
 
-std::string ScreenshotParams::expectedDirectoryPath() {
-    return "images/expected_images";
+std::filesystem::path ScreenshotParams::expectedDirectoryPath() {
+    return BackendTest::binaryDirectory().append("images/expected_images");
 }
 
 std::string ScreenshotParams::expectedFileName() const {
     return absl::StrFormat("%s.png", mFileName);
 }
 
-std::string ScreenshotParams::expectedFilePath() const {
-    return absl::StrFormat("%s/%s", expectedDirectoryPath(), expectedFileName());
+std::filesystem::path ScreenshotParams::expectedFilePath() const {
+    return expectedDirectoryPath().append(expectedFileName());
 }
 
 const std::string ScreenshotParams::filePrefix() const {

filament/backend/test/ImageExpectations.h

@@ -17,6 +17,7 @@
 #ifndef TNT_IMAGE_EXPECTATIONS_H
 #define TNT_IMAGE_EXPECTATIONS_H
 
+#include <filesystem>
 #include <vector>
 
 #include "gtest/gtest.h"
@@ -51,12 +52,12 @@ public:
     bool isSrgb() const;
     uint32_t expectedHash() const;
 
-    static std::string actualDirectoryPath();
+    static std::filesystem::path actualDirectoryPath();
     std::string actualFileName() const;
-    std::string actualFilePath() const;
-    static std::string expectedDirectoryPath();
+    std::filesystem::path actualFilePath() const;
+    static std::filesystem::path expectedDirectoryPath();
     std::string expectedFileName() const;
-    std::string expectedFilePath() const;
+    std::filesystem::path expectedFilePath() const;
     const std::string filePrefix() const;
 
 private:

filament/backend/test/test_ReadPixels.cpp

@@ -139,6 +139,24 @@ TEST_F(ReadPixelsTest, ReadPixels) {
             return bufferDimension;
         }
 
+        void exportScreenshot(void* pixelData) const {
+#ifndef FILAMENT_IOS
+            const size_t width = readRect.width, height = readRect.height;
+            LinearImage image(width, height, 4);
+            if (format == PixelDataFormat::RGBA && type == PixelDataType::UBYTE) {
+                image = toLinearWithAlpha<uint8_t>(width, height, width * 4, (uint8_t*)pixelData);
+            }
+            if (format == PixelDataFormat::RGBA && type == PixelDataType::FLOAT) {
+                memcpy(image.getPixelRef(), pixelData, width * height * sizeof(math::float4));
+            }
+            std::string png = std::string(testName) + ".png";
+            std::filesystem::path path = ScreenshotParams::actualDirectoryPath();
+            path.append(png);
+            std::ofstream outputStream(path.c_str(), std::ios::binary | std::ios::trunc);
+            ImageEncoder::encode(outputStream, ImageEncoder::Format::PNG, image, "", png);
+#endif
+        }
+
         // The format and type for the readPixels call.
         PixelDataFormat format = PixelDataFormat::RGBA;
         PixelDataType type = PixelDataType::UBYTE;
@@ -311,6 +329,8 @@ TEST_F(ReadPixelsTest, ReadPixels) {
                     const auto* test = (const TestCase*)user;
                     assert_invariant(test);
 
+                    test->exportScreenshot(buffer);
+
                     // Hash the contents of the buffer and check that they match.
                     uint32_t hash = utils::hash::murmur3((const uint32_t*)buffer, size / 4, 0);
 

filament/include/filament/Engine.h

@@ -317,6 +317,15 @@ 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.

filament/include/filament/MaterialInstance.h

@@ -86,9 +86,6 @@ 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
@@ -240,13 +237,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>
@@ -264,53 +261,6 @@ 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.
      *

filament/src/MaterialInstance.cpp

@@ -232,46 +232,6 @@ 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();
 }

filament/src/MaterialParser.cpp

@@ -229,11 +229,6 @@ 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);
@@ -756,38 +751,6 @@ 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) {

filament/src/MaterialParser.h

@@ -49,7 +49,6 @@ class BufferInterfaceBlock;
 class SamplerInterfaceBlock;
 struct SubpassInfo;
 struct MaterialConstant;
-struct MaterialMutableConstant;
 struct MaterialPushConstant;
 
 class MaterialParser {
@@ -80,8 +79,6 @@ 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;
 
@@ -240,13 +237,6 @@ 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);

filament/src/PostProcessManager.cpp

@@ -187,10 +187,7 @@ FMaterial* PostProcessManager::PostProcessMaterial::getMaterial(FEngine& engine,
     if (UTILS_UNLIKELY(mSize)) {
         loadMaterial(engine);
     }
-    mMaterial->prepareProgram(Variant{ Variant::type_t(variant) },
-            // TODO(exv): allow post process materials to use mutable spec constants?
-            Program::MutableSpecConstantsInfo(0)
-            );
+    mMaterial->prepareProgram(Variant{ Variant::type_t(variant) });
     return mMaterial;
 }
 
@@ -426,10 +423,7 @@ UTILS_NOINLINE
 PipelineState PostProcessManager::getPipelineState(
         FMaterial const* const ma, PostProcessVariant variant) const noexcept {
     return {
-            .program = ma->getProgram(Variant{ Variant::type_t(variant) },
-                    // TODO(exv): allow post process materials to use mutable spec constants?
-                    Program::MutableSpecConstantsInfo(0)
-                    ),
+            .program = ma->getProgram(Variant{ Variant::type_t(variant) }),
             .vertexBufferInfo = mFullScreenQuadVbih,
             .pipelineLayout = {
                     .setLayout = {

filament/src/RenderPass.cpp

@@ -246,8 +246,7 @@ 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,
-                    first->info.mi->getMutableSpecConstants());
+            ma->prepareProgram(first->info.materialVariant);
         }
     }
 }
@@ -1099,8 +1098,7 @@ void RenderPass::Executor::execute(FEngine const& engine, DriverApi& driver,
                 }
 
                 assert_invariant(ma);
-                pipeline.program = ma->getProgram(info.materialVariant,
-                        info.mi->getMutableSpecConstants());
+                pipeline.program = ma->getProgram(info.materialVariant);
 
                 if (UTILS_UNLIKELY(memcmp(&pipeline, &currentPipeline, sizeof(PipelineState)) != 0)) {
                     currentPipeline = pipeline;

filament/src/details/Engine.cpp

@@ -140,6 +140,7 @@ 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);
 
@@ -759,6 +760,7 @@ 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);
 

filament/src/details/Material.cpp

@@ -342,24 +342,8 @@ 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
@@ -465,17 +449,13 @@ 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)) {
-                    for (int spec = 0; spec < mutableConstantVariations; spec++) {
-                        prepareProgram(variant, Program::MutableSpecConstantsInfo(spec), priority);
-                    }
+                    prepareProgram(variant, priority);
                 }
             }
         }
@@ -516,11 +496,6 @@ 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) ||
@@ -561,15 +536,14 @@ 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, mutableSpecConstants, priorityQueue);
+            getSurfaceProgramSlow(variant, priorityQueue);
             break;
         case MaterialDomain::POST_PROCESS:
-            getPostProcessProgramSlow(variant, mutableSpecConstants, priorityQueue);
+            getPostProcessProgramSlow(variant, priorityQueue);
             break;
         case MaterialDomain::COMPUTE:
             // TODO: implement MaterialDomain::COMPUTE
@@ -578,7 +552,6 @@ 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
@@ -589,25 +562,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, mutableSpecConstants) };
+    Program pb{ getProgramWithVariants(variant, vertexVariant, fragmentVariant) };
     pb.priorityQueue(priorityQueue);
     pb.multiview(
             mEngine.getConfig().stereoscopicType == StereoscopicType::MULTIVIEW &&
             Variant::isStereoVariant(variant));
-    createAndCacheProgram(std::move(pb), variant, mutableSpecConstants);
+    createAndCacheProgram(std::move(pb), variant);
 }
 
 void FMaterial::getPostProcessProgramSlow(Variant const variant,
-        const backend::Program::MutableSpecConstantsInfo mutableSpecConstants,
         CompilerPriorityQueue const priorityQueue) const noexcept {
-    Program pb{ getProgramWithVariants(variant, variant, variant, mutableSpecConstants) };
+    Program pb{ getProgramWithVariants(variant, variant, variant) };
     pb.priorityQueue(priorityQueue);
-    createAndCacheProgram(std::move(pb), variant, mutableSpecConstants);
+    createAndCacheProgram(std::move(pb), variant);
 }
 
-Program FMaterial::getProgramWithVariants(Variant variant,
-        Variant vertexVariant, Variant fragmentVariant,
-        const backend::Program::MutableSpecConstantsInfo mutableSpecConstants) const {
+Program FMaterial::getProgramWithVariants(
+        Variant variant,
+        Variant vertexVariant,
+        Variant fragmentVariant) const {
     FEngine const& engine = mEngine;
     const ShaderModel sm = engine.getShaderModel();
     const bool isNoop = engine.getBackend() == Backend::NOOP;
@@ -668,34 +641,29 @@ Program FMaterial::getProgramWithVariants(Variant variant,
             mProgramDescriptorBindings[+DescriptorSetBindingPoints::PER_RENDERABLE]);
     program.descriptorBindings(+DescriptorSetBindingPoints::PER_MATERIAL,
             mProgramDescriptorBindings[+DescriptorSetBindingPoints::PER_MATERIAL]);
-    program.specializationConstants(mSpecializationConstants,
-            CONFIG_MAX_RESERVED_SPEC_CONSTANTS + mMaterialConstants.size(), mutableSpecConstants);
+    program.specializationConstants(mSpecializationConstants);
 
     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),
-                    hash::combine(variant.key, mutableSpecConstants.getValue())));
+    program.cacheId(hash::combine(size_t(mCacheId), variant.key));
 
     return program;
 }
 
-void FMaterial::createAndCacheProgram(Program&& p, Variant const variant,
-        Program::MutableSpecConstantsInfo const mutableSpecConstants) const noexcept {
+void FMaterial::createAndCacheProgram(Program&& p, Variant const variant) 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[cacheIndex];
+            auto const program = pDefaultMaterial->mCachedPrograms[variant.key];
             if (program) {
-                mCachedPrograms[cacheIndex] = program;
+                mCachedPrograms[variant.key] = program;
                 return;
             }
         }
@@ -704,17 +672,17 @@ void FMaterial::createAndCacheProgram(Program&& p, Variant const variant,
     auto const program = driverApi.createProgram(std::move(p));
     driverApi.setDebugTag(program.getId(), mName);
     assert_invariant(program);
-    mCachedPrograms[cacheIndex] = program;
+    mCachedPrograms[variant.key] = 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[cacheIndex]) {
+        if (pDefaultMaterial && !pDefaultMaterial->mCachedPrograms[variant.key]) {
             // set the tag to the default material name
             driverApi.setDebugTag(program.getId(), mName);
-            pDefaultMaterial->mCachedPrograms[cacheIndex] = program;
+            pDefaultMaterial->mCachedPrograms[variant.key] = program;
         }
     }
 }
@@ -762,14 +730,6 @@ 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
@@ -828,9 +788,7 @@ void FMaterial::onQueryCallback(void* userdata, VariantList* pActiveVariants) {
 
 #endif // FILAMENT_ENABLE_MATDBG
 
-[[nodiscard]] Handle<HwProgram> FMaterial::getProgramWithMATDBG(Variant const variant,
-        Program::MutableSpecConstantsInfo const mutableSpecConstants) const noexcept {
-    size_t cacheIndex = getCachedProgramIndex(variant, mutableSpecConstants);
+[[nodiscard]] Handle<HwProgram> FMaterial::getProgramWithMATDBG(Variant const variant) const noexcept {
 #if FILAMENT_ENABLE_MATDBG
     assert_invariant((size_t)variant.key < VARIANT_COUNT);
     std::unique_lock lock(mActiveProgramsLock);
@@ -845,13 +803,13 @@ void FMaterial::onQueryCallback(void* userdata, VariantList* pActiveVariants) {
     lock.unlock();
     if (isSharedVariant(variant)) {
         FMaterial const* const pDefaultMaterial = mEngine.getDefaultMaterial();
-        if (pDefaultMaterial && pDefaultMaterial->mCachedPrograms[cacheIndex]) {
-            return pDefaultMaterial->getProgram(variant, mutableSpecConstants);
+        if (pDefaultMaterial && pDefaultMaterial->mCachedPrograms[variant.key]) {
+            return pDefaultMaterial->getProgram(variant);
         }
     }
 #endif
-    assert_invariant(mCachedPrograms[cacheIndex]);
-    return mCachedPrograms[cacheIndex];
+    assert_invariant(mCachedPrograms[variant.key]);
+    return mCachedPrograms[variant.key];
 }
 
 void FMaterial::destroyPrograms(FEngine& engine,
@@ -864,7 +822,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 = cachedPrograms.size(); k < n; ++k) {
+                for (size_t k = 0, n = VARIANT_COUNT; 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.
@@ -882,7 +840,7 @@ void FMaterial::destroyPrograms(FEngine& engine,
                 UTILS_UNUSED_IN_RELEASE
                 auto const UTILS_NULLABLE pDefaultMaterial = engine.getDefaultMaterial();
 
-                for (size_t k = 0, n = cachedPrograms.size(); k < n; ++k) {
+                for (size_t k = 0, n = VARIANT_COUNT; 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.
@@ -894,7 +852,7 @@ void FMaterial::destroyPrograms(FEngine& engine,
                                 // During Engine::shutdown, auto-cleanup destroys the
                                 // default material first
                                 assert_invariant(!pDefaultMaterial ||
-                                        pDefaultMaterial->mCachedPrograms[k & (VARIANT_COUNT - 1)]);
+                                        pDefaultMaterial->mCachedPrograms[k]);
                                 // we don't own this variant, skip, but clear the entry.
                                 cachedPrograms[k].clear();
                                 continue;
@@ -908,7 +866,7 @@ void FMaterial::destroyPrograms(FEngine& engine,
             break;
         }
         case MaterialDomain::POST_PROCESS: {
-            for (size_t k = 0, n = cachedPrograms.size(); k < n; ++k) {
+            for (size_t k = 0, n = POST_PROCESS_VARIANT_COUNT; 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.
@@ -920,13 +878,8 @@ void FMaterial::destroyPrograms(FEngine& engine,
             break;
         }
         case MaterialDomain::COMPUTE: {
-            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]));
-                }
-            }
+            // Compute programs don't have variants
+            driverApi.destroyProgram(std::move(cachedPrograms[0]));
             break;
         }
     }
@@ -1161,19 +1114,6 @@ 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)];
@@ -1229,8 +1169,7 @@ void FMaterial::precacheDepthVariants(FEngine& engine) {
             }
             assert_invariant(Variant::isValidDepthVariant(variant));
             if (hasVariant(variant)) {
-                // The default material shouldn't ever have mutable spec constants.
-                prepareProgram(variant, Program::MutableSpecConstantsInfo(0));
+                prepareProgram(variant);
             }
         }
         return;
@@ -1243,16 +1182,9 @@ 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));
-            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];
-            }
+            mCachedPrograms[variant.key] = pDefaultMaterial->mCachedPrograms[variant.key];
         }
     }
 }

filament/src/details/Material.h

@@ -122,8 +122,9 @@ public:
 
     FEngine& getEngine() const noexcept  { return mEngine; }
 
-    bool isCached(Variant const variant,
-            const backend::Program::MutableSpecConstantsInfo mutableSpecConstants) const noexcept;
+    bool isCached(Variant const variant) const noexcept {
+        return bool(mCachedPrograms[variant.key]);
+    }
 
     void invalidate(Variant::type_t variantMask = 0, Variant::type_t variantValue = 0) noexcept;
 
@@ -131,31 +132,26 @@ 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, mutableSpecConstants))) {
-            prepareProgramSlow(variant, mutableSpecConstants, priorityQueue);
+        if (UTILS_UNLIKELY(!isCached(variant))) {
+            prepareProgramSlow(variant, priorityQueue);
         }
     }
 
-    // getProgram returns the backend program for the material's given variant and set of mutable
-    // spec constants.
+    // getProgram returns the backend program for the material's given variant.
     // Must be called after prepareProgram().
     [[nodiscard]]
-    backend::Handle<backend::HwProgram> getProgram(Variant const variant,
-            backend::Program::MutableSpecConstantsInfo const mutableSpecConstants) const noexcept {
+    backend::Handle<backend::HwProgram> getProgram(Variant const variant) const noexcept {
 #if FILAMENT_ENABLE_MATDBG
-        return getProgramWithMATDBG(variant, mutableSpecConstants);
+        return getProgramWithMATDBG(variant);
 #endif
-        size_t index = getCachedProgramIndex(variant, mutableSpecConstants);
-        assert_invariant(mCachedPrograms[index]);
-        return mCachedPrograms[index];
+        assert_invariant(mCachedPrograms[variant.key]);
+        return mCachedPrograms[variant.key];
     }
 
     [[nodiscard]]
-    backend::Handle<backend::HwProgram> getProgramWithMATDBG(Variant variant,
-            backend::Program::MutableSpecConstantsInfo const mutableSpecConstants) const noexcept;
+    backend::Handle<backend::HwProgram> getProgramWithMATDBG(Variant variant) const noexcept;
 
     bool isVariantLit() const noexcept { return mIsVariantLit; }
 
@@ -212,12 +208,6 @@ 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,
@@ -271,48 +261,34 @@ 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,
-            backend::Program::MutableSpecConstantsInfo mutableSpecConstants) const;
+            Variant vertexVariant, Variant fragmentVariant) 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,
-            backend::Program::MutableSpecConstantsInfo mutableSpecConstants) const noexcept;
+    void createAndCacheProgram(backend::Program&& p, Variant variant) 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 utils::FixedCapacityVector<backend::Handle<backend::HwProgram>> mCachedPrograms;
+    mutable std::array<backend::Handle<backend::HwProgram>, VARIANT_COUNT> mCachedPrograms;
     DescriptorSetLayout mPerViewDescriptorSetLayout;
     DescriptorSetLayout mPerViewDescriptorSetLayoutVsm;
     DescriptorSetLayout mDescriptorSetLayout;
@@ -348,7 +324,6 @@ 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;
@@ -373,13 +348,6 @@ 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>

filament/src/details/MaterialInstance.cpp

@@ -63,7 +63,6 @@ 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),
@@ -129,7 +128,6 @@ 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),

filament/src/details/MaterialInstance.h

@@ -31,7 +31,6 @@
 
 #include <backend/DriverEnums.h>
 #include <backend/Handle.h>
-#include <backend/Program.h>
 
 #include <utils/BitmaskEnum.h>
 #include <utils/bitset.h>
@@ -67,7 +66,7 @@ public:
     void terminate(FEngine& engine);
 
     void commitStreamUniformAssociations(FEngine::DriverApi& driver);
-
+    
     void commit(FEngine& engine) const;
 
     void commit(FEngine::DriverApi& driver) const;
@@ -243,10 +242,6 @@ public:
 
     using MaterialInstance::setParameter;
 
-    const backend::Program::MutableSpecConstantsInfo getMutableSpecConstants() const noexcept {
-        return mConstants;
-    }
-
 private:
     friend class FMaterial;
     friend class MaterialInstance;
@@ -269,12 +264,6 @@ 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;
 
@@ -288,7 +277,6 @@ private:
     mutable DescriptorSet mDescriptorSet;
     UniformBuffer mUniforms;
     bool mHasStreamUniformAssociations = false;
-    backend::Program::MutableSpecConstantsInfo mConstants;
 
     backend::PolygonOffset mPolygonOffset{};
     backend::StencilState mStencilState{};

filament/src/ds/DescriptorSet.cpp

@@ -66,6 +66,7 @@ DescriptorSet& DescriptorSet::operator=(DescriptorSet&& rhs) noexcept {
         mDirty = rhs.mDirty;
         mValid = rhs.mValid;
         mSetAfterCommitWarning = rhs.mSetAfterCommitWarning;
+        mSetUndefinedParameterWarning = rhs.mSetUndefinedParameterWarning;
     }
     return *this;
 }
@@ -106,11 +107,12 @@ void DescriptorSet::commitSlow(DescriptorSetLayout const& layout,
     });
 
     auto const unsetValidDescriptors = layout.getValidDescriptors() & ~mValid;
-    if (UTILS_VERY_UNLIKELY(!unsetValidDescriptors.empty())) {
+    if (UTILS_VERY_UNLIKELY(!unsetValidDescriptors.empty() && !mSetUndefinedParameterWarning)) {
         unsetValidDescriptors.forEachSetBit([&](auto i) {
             LOG(WARNING) << (layout.isSampler(i) ? "Sampler" : "Buffer") << " descriptor " << i
                          << " of " << mName.c_str() << " is not set. Please report this issue.";
         });
+        mSetUndefinedParameterWarning = true;
     }
 }
 

filament/src/ds/DescriptorSet.h

@@ -110,6 +110,7 @@ private:
     mutable utils::bitset64 mValid;                         //  8
     backend::DescriptorSetHandle mDescriptorSetHandle;      //  4
     mutable bool mSetAfterCommitWarning = false;            //  1
+    mutable bool mSetUndefinedParameterWarning = false;     //  1
     utils::StaticString mName;                              // 16
 };
 

libs/filabridge/include/filament/MaterialChunkType.h

@@ -54,7 +54,6 @@ 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"),

libs/filabridge/include/filament/MaterialEnums.h

@@ -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 = 60;
+static constexpr size_t MATERIAL_VERSION = 61;
 
 /**
  * Supported shading models

libs/filabridge/include/private/filament/ConstantInfo.h

@@ -30,19 +30,7 @@ struct MaterialConstant {
     ConstantType type;
 
     MaterialConstant() = default;
-    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) {}
+    MaterialConstant(const char* name, ConstantType type) : name(name), type(type)  {}
 };
 
 }

libs/filamat/include/filamat/MaterialBuilder.h

@@ -307,15 +307,14 @@ public:
     MaterialBuilder& parameter(const char* name, size_t size, UniformType type,
             ParameterPrecision precision = ParameterPrecision::DEFAULT);
 
-    //! Add a (mutable) constant parameter to this material.
+    //! Add a 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, bool isMutable,
-            T defaultValue = 0);
+    MaterialBuilder& constant(const char *name, ConstantType type, T defaultValue = 0);
 
     /**
      * Add a sampler parameter to this material.
@@ -872,7 +871,6 @@ private:
     PropertyList mProperties;
     ParameterList mParameters;
     ConstantList mConstants;
-    ConstantList mMutableConstants;
     PushConstantList mPushConstants;
     SubpassList mSubpasses;
     VariableList mVariables;

libs/filamat/src/MaterialBuilder.cpp

@@ -317,19 +317,12 @@ MaterialBuilder& MaterialBuilder::parameter(const char* name, SamplerType sample
 }
 
 template<typename T, typename>
-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())
+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())
             << "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,
@@ -361,20 +354,15 @@ MaterialBuilder& MaterialBuilder::constant(const char* name, ConstantType const
         assert_invariant(false);
     }
 
-    if (isMutable) {
-        FILAMENT_CHECK_POSTCONDITION(type == ConstantType::BOOL)
-                << "Mutable spec constants must be bool type.";
-    }
-
-    constants.push_back(constant);
+    mConstants.push_back(constant);
     return *this;
 }
 template MaterialBuilder& MaterialBuilder::constant<int32_t>(
-        const char* name, ConstantType type, bool isMutable, int32_t defaultValue);
+        const char* name, ConstantType type, int32_t defaultValue);
 template MaterialBuilder& MaterialBuilder::constant<float>(
-        const char* name, ConstantType type, bool isMutable, float defaultValue);
+        const char* name, ConstantType type, float defaultValue);
 template MaterialBuilder& MaterialBuilder::constant<bool>(
-        const char* name, ConstantType type, bool isMutable, bool defaultValue);
+        const char* name, ConstantType type, bool defaultValue);
 
 MaterialBuilder& MaterialBuilder::buffer(BufferInterfaceBlock bib) {
     FILAMENT_CHECK_POSTCONDITION(mBuffers.size() < MAX_BUFFERS_COUNT) << "Too many buffers";
@@ -919,7 +907,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, mMutableConstants, mPushConstants,
+    ShaderGenerator sg(mProperties, mVariables, mOutputs, mDefines, mConstants, mPushConstants,
             mMaterialFragmentCode.getResolved(), mMaterialFragmentCode.getLineOffset(),
             mMaterialVertexCode.getResolved(), mMaterialVertexCode.getLineOffset(),
             mMaterialDomain);
@@ -1504,7 +1492,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, mMutableConstants, mPushConstants,
+    ShaderGenerator const sg(properties, mVariables, mOutputs, mDefines, mConstants, mPushConstants,
             mMaterialFragmentCode.getResolved(), mMaterialFragmentCode.getLineOffset(),
             mMaterialVertexCode.getResolved(), mMaterialVertexCode.getLineOffset(),
             mMaterialDomain);
@@ -1630,14 +1618,9 @@ 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.type); });
+            [](Constant const& c) { return MaterialConstant(c.name.c_str(), 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) {

libs/filamat/src/eiff/MaterialInterfaceBlockChunk.cpp

@@ -117,20 +117,6 @@ 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),

libs/filamat/src/eiff/MaterialInterfaceBlockChunk.h

@@ -18,7 +18,6 @@
 #define TNT_FILAMAT_MAT_INTEFFACE_BLOCK_CHUNK_H
 
 #include "Chunk.h"
-#include "private/filament/ConstantInfo.h"
 
 #include <backend/Program.h>
 
@@ -92,20 +91,6 @@ 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,

libs/filamat/src/shaders/ShaderGenerator.cpp

@@ -283,17 +283,10 @@ void ShaderGenerator::appendShader(io::sstream& ss,
     }
 }
 
-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) {
+void ShaderGenerator::generateUserSpecConstants(
+        const CodeGenerator& cg, io::sstream& fs, MaterialBuilder::ConstantList const& constants) {
     // Constants 0 to CONFIG_MAX_RESERVED_SPEC_CONSTANTS - 1 are reserved by Filament.
-    size_t index = CONFIG_MAX_RESERVED_SPEC_CONSTANTS + offset;
+    size_t index = CONFIG_MAX_RESERVED_SPEC_CONSTANTS;
     for (const auto& constant : constants) {
         std::string const fullName = std::string("materialConstants_") + constant.name.c_str();
         switch (constant.type) {
@@ -321,7 +314,6 @@ 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,
@@ -344,7 +336,6 @@ ShaderGenerator::ShaderGenerator(
     mMaterialDomain = materialDomain;
     mDefines = defines;
     mConstants = constants;
-    mMutableConstants = mutableConstants;
     mPushConstants = pushConstants;
 
     if (mMaterialFragmentCode.empty()) {
@@ -401,7 +392,7 @@ std::string ShaderGenerator::createSurfaceVertexProgram(ShaderModel const shader
 
     cg.generateCommonProlog(vs, ShaderStage::VERTEX, material, variant);
 
-    generateAllUserSpecConstants(cg, vs, mConstants, mMutableConstants);
+    generateUserSpecConstants(cg, vs, mConstants);
 
     // 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
@@ -534,7 +525,7 @@ std::string ShaderGenerator::createSurfaceFragmentProgram(ShaderModel const shad
     io::sstream fs;
     cg.generateCommonProlog(fs, ShaderStage::FRAGMENT, material, variant);
 
-    generateAllUserSpecConstants(cg, fs, mConstants, mMutableConstants);
+    generateUserSpecConstants(cg, fs, mConstants);
 
     generateSurfaceMaterialVariantDefines(
             fs, ShaderStage::FRAGMENT, featureLevel, material, variant);
@@ -690,7 +681,7 @@ std::string ShaderGenerator::createSurfaceComputeProgram(ShaderModel const shade
 
     cg.generateCommonProlog(s, ShaderStage::COMPUTE, material, {});
 
-    generateAllUserSpecConstants(cg, s, mConstants, mMutableConstants);
+    generateUserSpecConstants(cg, s, mConstants);
 
     CodeGenerator::generateSurfaceTypes(s, ShaderStage::COMPUTE);
 
@@ -730,7 +721,7 @@ std::string ShaderGenerator::createPostProcessVertexProgram(ShaderModel const sm
     io::sstream vs;
     cg.generateCommonProlog(vs, ShaderStage::VERTEX, material, {});
 
-    generateAllUserSpecConstants(cg, vs, mConstants, mMutableConstants);
+    generateUserSpecConstants(cg, vs, mConstants);
 
     CodeGenerator::generateDefine(vs, "LOCATION_POSITION", uint32_t(POSITION));
 
@@ -774,7 +765,7 @@ std::string ShaderGenerator::createPostProcessFragmentProgram(ShaderModel const
     io::sstream fs;
     cg.generateCommonProlog(fs, ShaderStage::FRAGMENT, material, {});
 
-    generateAllUserSpecConstants(cg, fs, mConstants, mMutableConstants);
+    generateUserSpecConstants(cg, fs, mConstants);
 
     generatePostProcessMaterialVariantDefines(fs, PostProcessVariant(variant));
 

libs/filamat/src/shaders/ShaderGenerator.h

@@ -51,7 +51,6 @@ 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,
@@ -105,14 +104,9 @@ 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,
-            size_t offset);
+            MaterialBuilder::ConstantList const& constants);
 
     std::string createPostProcessVertexProgram(filament::backend::ShaderModel sm,
             MaterialBuilder::TargetApi targetApi, MaterialBuilder::TargetLanguage targetLanguage,
@@ -141,7 +135,6 @@ 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;

libs/filamat/tests/test_filamat.cpp

@@ -888,10 +888,10 @@ TEST_F(MaterialCompiler, ConstantParameter) {
         }
     )");
   filamat::MaterialBuilder builder;
-  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.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.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, /*isMutable=*/false, 1.0f);
-        builder.constant("myFloatConstant", ConstantType::FLOAT, /*isMutable=*/false, 1.0f);
+        builder.constant("myFloatConstant", ConstantType::FLOAT, 1.0f);
+        builder.constant("myFloatConstant", ConstantType::FLOAT, 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, /*isMutable=*/false, 10);
+        builder.constant("myFloatConstant", ConstantType::FLOAT, 10);
     }, utils::PostconditionPanic);
 #endif
 }

libs/gltfio/include/gltfio/MaterialProvider.h

@@ -21,6 +21,7 @@
 #include <filament/Material.h>
 #include <filament/MaterialInstance.h>
 
+#include <utils/FixedCapacityVector.h>
 #include <utils/compiler.h>
 
 #include <array>
@@ -196,6 +197,7 @@ 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.
@@ -203,7 +205,8 @@ void processShaderString(std::string* shader, const UvMap& uvmap,
  * @see createUbershaderProvider
  */
 UTILS_PUBLIC
-MaterialProvider* createJitShaderProvider(Engine* engine, bool optimizeShaders = false);
+MaterialProvider* createJitShaderProvider(Engine* engine, bool optimizeShaders = false,
+        utils::FixedCapacityVector<char const*> const& variantFilters = {});
 
 /**
  * Creates a material provider that loads a small set of pre-built materials.

libs/gltfio/src/JitShaderProvider.cpp

@@ -16,6 +16,7 @@
 
 #include <gltfio/MaterialProvider.h>
 
+#include <filament/MaterialEnums.h>
 #include <filamat/MaterialBuilder.h>
 
 #include <utils/Hash.h>
@@ -23,6 +24,7 @@
 #include <tsl/robin_map.h>
 
 #include <string>
+#include <unordered_map>
 
 using namespace filamat;
 using namespace filament;
@@ -33,7 +35,8 @@ namespace {
 
 class JitShaderProvider : public MaterialProvider {
 public:
-    explicit JitShaderProvider(Engine* engine, bool optimizeShaders);
+    explicit JitShaderProvider(Engine* engine, bool optimizeShaders,
+            utils::FixedCapacityVector<char const*> const& variantFilters);
     ~JitShaderProvider() override;
 
     MaterialInstance* createMaterialInstance(MaterialKey* config, UvMap* uvmap,
@@ -49,15 +52,37 @@ 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) : mEngine(engine),
-        mOptimizeShaders(optimizeShaders) {
+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);
+    }
     MaterialBuilder::init();
 }
 
@@ -344,24 +369,25 @@ std::string shaderFromKey(const MaterialKey& config) {
 }
 
 Material* createMaterial(Engine* engine, const MaterialKey& config, const UvMap& uvmap,
-        const char* name, bool optimizeShaders) {
+        const char* name, bool optimizeShaders, filament::UserVariantFilterMask variantFilter) {
     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);
+            .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);
 
     if (!optimizeShaders) {
         builder.optimization(MaterialBuilder::Optimization::NONE);
@@ -589,13 +615,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);
+        Material* mat =
+                createMaterial(mEngine, *config, *uvmap, label, optimizeShaders, mVariantFilter);
         mCache.emplace(std::make_pair(*config, mat));
         mMaterials.push_back(mat);
         return mat;
@@ -612,8 +638,9 @@ MaterialInstance* JitShaderProvider::createMaterialInstance(MaterialKey* config,
 
 namespace filament::gltfio {
 
-MaterialProvider* createJitShaderProvider(filament::Engine* engine, bool optimizeShaders) {
-    return new JitShaderProvider(engine, optimizeShaders);
+MaterialProvider* createJitShaderProvider(Engine* engine, bool optimizeShaders,
+        utils::FixedCapacityVector<char const*> const& variantFilters) {
+    return new JitShaderProvider(engine, optimizeShaders, variantFilters);
 }
 
 } // namespace filament::gltfio

samples/CMakeLists.txt

@@ -19,7 +19,6 @@ set(RESOURCE_BINS)
 set(MATERIAL_SRCS
         materials/aiDefaultMat.mat
         materials/bakedColor.mat
-        materials/bakedColorWithMutableConstants.mat
         materials/bakedTexture.mat
         materials/aoPreview.mat
         materials/arrayTexture.mat
@@ -223,6 +222,8 @@ 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
@@ -279,7 +280,6 @@ 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)

samples/common/arguments.cpp

@@ -14,6 +14,8 @@
  * limitations under the License.
  */
 
+#include "arguments.h"
+
 #include <filament/Engine.h>
 
 #include <utils/Path.h>

samples/common/arguments.h

@@ -22,8 +22,10 @@
 #include <string>
 
 namespace samples {
-    filament::Engine::Backend parseArgumentsForBackend(int argc, char* argv[]);
-    filament::Engine::Backend parseArgumentsForBackend(const std::string& backend);
-    std::string getBackendAPIArgumentsUsage();
-}
+
+filament::Engine::Backend parseArgumentsForBackend(int argc, char* argv[]);
+filament::Engine::Backend parseArgumentsForBackend(const std::string& backend);
+std::string getBackendAPIArgumentsUsage();
+
+} // namespace samples
 #endif //TNT_SAMPLES_ARGUMENTS_H

samples/common/configuration.cpp

@@ -0,0 +1,65 @@
+/*
+ * 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

samples/common/configuration.h

@@ -0,0 +1,31 @@
+/*
+ * 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

samples/gltf_instances.cpp

@@ -15,6 +15,7 @@
  */
 
 #include "common/arguments.h"
+#include "common/configuration.h"
 
 #include <filamentapp/Config.h>
 #include <filamentapp/FilamentApp.h>
@@ -244,8 +245,11 @@ 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) :
-                createUbershaderProvider(engine, UBERARCHIVE_DEFAULT_DATA, UBERARCHIVE_DEFAULT_SIZE);
+        app.materials = (app.materialSource == JITSHADER)
+                                ? createJitShaderProvider(engine, false /* optimize */,
+                                          samples::getJitMaterialVariantFilter(app.config.backend))
+                                : createUbershaderProvider(engine, UBERARCHIVE_DEFAULT_DATA,
+                                          UBERARCHIVE_DEFAULT_SIZE);
 
         app.loader = AssetLoader::create({engine, app.materials, app.names });
         if (filename.isEmpty()) {

samples/gltf_viewer.cpp

@@ -15,6 +15,7 @@
  */
 
 #include "common/arguments.h"
+#include "common/configuration.h"
 
 #include <filamentapp/Config.h>
 #include <filamentapp/FilamentApp.h>
@@ -778,11 +779,13 @@ int main(int argc, char** argv) {
             }
         }
 
-        app.materials = (app.materialSource == JITSHADER) ?
-                createJitShaderProvider(engine, OPTIMIZE_MATERIALS) :
-                createUbershaderProvider(engine, UBERARCHIVE_DEFAULT_DATA, UBERARCHIVE_DEFAULT_SIZE);
+        app.materials = (app.materialSource == JITSHADER)
+                                ? createJitShaderProvider(engine, OPTIMIZE_MATERIALS,
+                                          samples::getJitMaterialVariantFilter(app.config.backend))
+                                : 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(

samples/materials/bakedColorWithMutableConstants.mat

@@ -1,44 +0,0 @@
-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;
-        }
-    }
-}

samples/mutable_constants.cpp

@@ -1,214 +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 <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;
-}

samples/simple_gltf_test.cpp

@@ -13,7 +13,9 @@
  * 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>
@@ -72,6 +74,7 @@ struct App {
     gltfio::TextureProvider* stbDecoder = nullptr;
     gltfio::TextureProvider* ktxDecoder = nullptr;
     FilamentInstance* instance;
+    bool enableMSAA = false;
 };
 
 static const char* DEFAULT_IBL = "assets/ibl/lightroom_14b";
@@ -87,7 +90,9 @@ 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");
+                      "       Override the built-in IBL\n\n"
+                      "   --msaa, -m\n"
+                      "       setMultiSampleAntiAliasingOptions\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);
@@ -101,9 +106,11 @@ static void printUsage(char* name) {
 }
 
 static int handleCommandLineArguments(int argc, char* argv[], App* app) {
-    static constexpr const char* OPTSTR = "ha:i:un:m:";
+    static constexpr const char* OPTSTR = "hma:i:";
     static const struct option OPTIONS[] = { { "help", no_argument, nullptr, 'h' },
-        { "api", required_argument, nullptr, 'a' }, { "ibl", required_argument, nullptr, 'i' },
+        { "api", required_argument, nullptr, 'a' },
+        { "ibl", required_argument, nullptr, 'i' },
+        { "msaa", no_argument, nullptr, 'm' },
         { nullptr, 0, nullptr, 0 } };
     int opt;
     int option_index = 0;
@@ -122,6 +129,9 @@ static int handleCommandLineArguments(int argc, char* argv[], App* app) {
             case 'i':
                 app->config.iblDirectory = arg;
                 break;
+            case 'm':
+                app->enableMSAA = true;
+                break;
         }
     }
     return optind;
@@ -205,13 +215,19 @@ int main(int argc, char** argv) {
     };
 
     auto setup = [&](Engine* engine, View* view, Scene* scene) {
-        view->setMultiSampleAntiAliasingOptions({ .enabled = true});
+        if (app.enableMSAA) {
+            // TODO Investigate why
+            // Enabling MSAA fixes the "transmission" sample but breaks the others
+            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)
+                                ? createJitShaderProvider(engine, false /* optimize */,
+                                          samples::getJitMaterialVariantFilter(app.config.backend))
                                 : createUbershaderProvider(engine, UBERARCHIVE_DEFAULT_DATA,
                                           UBERARCHIVE_DEFAULT_SIZE);
 

tools/matc/src/matc/ParametersProcessor.cpp

@@ -154,7 +154,7 @@ static bool processParameter(MaterialBuilder& builder, const JsonishObject& json
         std::cerr << "parameters: name value must be STRING." << std::endl;
         return false;
     }
-
+    
     const JsonishValue* transformNameValue = jsonObject.getValue("transformName");
     if (transformNameValue && transformNameValue->getType() != JsonishValue::STRING) {
         std::cerr << "parameters: transformName value must be STRING." << std::endl;
@@ -328,19 +328,6 @@ static bool processConstant(MaterialBuilder& builder, const JsonishObject& jsonO
         return false;
     }
 
-    const JsonishValue* mutableValue = jsonObject.getValue("mutable");
-    bool isMutable;
-    if (mutableValue) {
-        const JsonishBool* mutableBool = mutableValue->toJsonBool();
-        if (!mutableBool) {
-            std::cerr << "constants: mutable value must be BOOL." << std::endl;
-            return false;
-        }
-        isMutable = mutableBool->getBool();
-    } else {
-        isMutable = false;
-    }
-
     auto typeString = typeValue->toJsonString()->getString();
     auto nameString = nameValue->toJsonString()->getString();
     const JsonishValue* defaultValue = jsonObject.getValue("default");
@@ -359,7 +346,7 @@ static bool processConstant(MaterialBuilder& builder, const JsonishObject& jsonO
                     // FIXME: Jsonish doesn't distinguish between integers and floats.
                     intDefault = (int32_t)defaultValue->toJsonNumber()->getFloat();
                 }
-                builder.constant(nameString.c_str(), type, isMutable, intDefault);
+                builder.constant(nameString.c_str(), type, intDefault);
                 break;
             }
             case ConstantType::FLOAT: {
@@ -372,7 +359,7 @@ static bool processConstant(MaterialBuilder& builder, const JsonishObject& jsonO
                     }
                     floatDefault = defaultValue->toJsonNumber()->getFloat();
                 }
-                builder.constant(nameString.c_str(), type, isMutable, floatDefault);
+                builder.constant(nameString.c_str(), type, floatDefault);
                 break;
             }
             case ConstantType::BOOL:
@@ -385,7 +372,7 @@ static bool processConstant(MaterialBuilder& builder, const JsonishObject& jsonO
                     }
                     boolDefault = defaultValue->toJsonBool()->getBool();
                 }
-                builder.constant(nameString.c_str(), type, isMutable, boolDefault);
+                builder.constant(nameString.c_str(), type, boolDefault);
                 break;
         }
     } else {