[gfx] Support StructuredBuffer<IInterface>. (#1851)

Co-authored-by: T. Foley <tfoleyNV@users.noreply.github.com>

15 files changed, 1113 insertions, 1523 deletions

diff --git a/tools/gfx-util/shader-cursor.cpp b/tools/gfx-util/shader-cursor.cpp
index b188901ec..afb1540d5 100644
--- a/tools/gfx-util/shader-cursor.cpp
+++ b/tools/gfx-util/shader-cursor.cpp
@@ -144,7 +144,18 @@ Result ShaderCursor::getField(const char* name, const char* nameEnd, ShaderCurso
 
 ShaderCursor ShaderCursor::getElement(SlangInt index) const
 {
-    // TODO: need to auto-dereference various buffer types...
+    if (m_containerType != ShaderObjectContainerType::None)
+    {
+        ShaderCursor elementCursor;
+        elementCursor.m_baseObject = m_baseObject;
+        elementCursor.m_typeLayout = m_typeLayout->getElementTypeLayout();
+        elementCursor.m_containerType = m_containerType;
+        elementCursor.m_offset.uniformOffset = index * m_typeLayout->getStride();
+        elementCursor.m_offset.bindingRangeIndex = 0;
+        elementCursor.m_offset.bindingArrayIndex = index;
+        return elementCursor;
+    }
+
     switch( m_typeLayout->getKind() )
     {
     case slang::TypeReflection::Kind::Array:
diff --git a/tools/gfx-util/shader-cursor.h b/tools/gfx-util/shader-cursor.h
index 24008cf24..7512c62ee 100644
--- a/tools/gfx-util/shader-cursor.h
+++ b/tools/gfx-util/shader-cursor.h
@@ -26,6 +26,7 @@ struct ShaderCursor
 {
     IShaderObject* m_baseObject = nullptr;
     slang::TypeLayoutReflection* m_typeLayout = nullptr;
+    ShaderObjectContainerType m_containerType = ShaderObjectContainerType::None;
     ShaderOffset m_offset;
 
     /// Get the type (layout) of the value being pointed at by the cursor
@@ -78,6 +79,7 @@ struct ShaderCursor
     ShaderCursor(IShaderObject* object)
         : m_baseObject(object)
         , m_typeLayout(object->getElementTypeLayout())
+        , m_containerType(object->getContainerType())
     {}
 
     SlangResult setData(void const* data, size_t size) const
@@ -116,9 +118,13 @@ struct ShaderCursor
         /// Produce a cursor to the element or field with the given `index`.
         ///
         /// This is a convenience wrapper around `getElement()`.
-    ShaderCursor operator[](SlangInt index) const
-    {
-        return getElement(index);
-    }
+    ShaderCursor operator[](int64_t index) const { return getElement((SlangInt)index); }
+    ShaderCursor operator[](uint64_t index) const { return getElement((SlangInt)index); }
+    ShaderCursor operator[](int32_t index) const { return getElement((SlangInt)index); }
+    ShaderCursor operator[](uint32_t index) const { return getElement((SlangInt)index); }
+    ShaderCursor operator[](int16_t index) const { return getElement((SlangInt)index); }
+    ShaderCursor operator[](uint16_t index) const { return getElement((SlangInt)index); }
+    ShaderCursor operator[](int8_t index) const { return getElement((SlangInt)index); }
+    ShaderCursor operator[](uint8_t index) const { return getElement((SlangInt)index); }
 };
 }
diff --git a/tools/gfx/cpu/render-cpu.cpp b/tools/gfx/cpu/render-cpu.cpp
index 0bdc06ad6..0c5f119b8 100644
--- a/tools/gfx/cpu/render-cpu.cpp
+++ b/tools/gfx/cpu/render-cpu.cpp
@@ -357,7 +357,7 @@ public:
     void*           m_data = nullptr;
 };
 
-class CPUResourceView : public IResourceView, public ComObject
+class CPUResourceView : public ResourceViewBase
 {
 public:
     enum class Kind
@@ -365,15 +365,6 @@ public:
         Buffer,
         Texture,
     };
-
-    SLANG_COM_OBJECT_IUNKNOWN_ALL
-    IResourceView* getInterface(const Guid& guid)
-    {
-        if (guid == GfxGUID::IID_ISlangUnknown || guid == GfxGUID::IID_IResourceView)
-            return static_cast<IResourceView*>(this);
-        return nullptr;
-    }
-
     Kind getViewKind() const { return m_kind; }
     Desc const& getDesc() const { return m_desc; }
 
@@ -576,6 +567,7 @@ public:
         slang::BindingType bindingType;
         Index count;
         Index baseIndex; // Flat index for sub-ojects
+        Index subObjectIndex;
 
         // TODO: The `uniformOffset` field should be removed,
         // since it cannot be supported by the Slang reflection
@@ -610,10 +602,10 @@ public:
     {
         initBase(renderer, layout);
 
-        Index subObjectCount = 0;
-        Index resourceCount = 0;
+        m_subObjectCount = 0;
+        m_resourceCount = 0;
 
-        m_elementTypeLayout = _unwrapParameterGroups(layout);
+        m_elementTypeLayout = _unwrapParameterGroups(layout, m_containerType);
         m_size = m_elementTypeLayout->getSize();
 
         // Compute the binding ranges that are used to store
@@ -645,18 +637,31 @@ public:
                 descriptorSetIndex, rangeIndexInDescriptorSet);
 
             Index baseIndex = 0;
+            Index subObjectIndex = 0;
             switch (slangBindingType)
             {
             case slang::BindingType::ConstantBuffer:
             case slang::BindingType::ParameterBlock:
             case slang::BindingType::ExistentialValue:
-                baseIndex = subObjectCount;
-                subObjectCount += count;
+                baseIndex = m_subObjectCount;
+                subObjectIndex = baseIndex;
+                m_subObjectCount += count;
+                break;
+            case slang::BindingType::RawBuffer:
+            case slang::BindingType::MutableRawBuffer:
+                if (slangLeafTypeLayout->getType()->getElementType() != nullptr)
+                {
+                    // A structured buffer occupies both a resource slot and
+                    // a sub-object slot.
+                    subObjectIndex = m_subObjectCount;
+                    m_subObjectCount += count;
+                }
+                baseIndex = m_resourceCount;
+                m_resourceCount += count;
                 break;
-
             default:
-                baseIndex = resourceCount;
-                resourceCount += count;
+                baseIndex = m_resourceCount;
+                m_resourceCount += count;
                 break;
             }
 
@@ -665,12 +670,10 @@ public:
             bindingRangeInfo.count = count;
             bindingRangeInfo.baseIndex = baseIndex;
             bindingRangeInfo.uniformOffset = uniformOffset;
+            bindingRangeInfo.subObjectIndex = subObjectIndex;
             m_bindingRanges.add(bindingRangeInfo);
         }
 
-        m_subObjectCount = subObjectCount;
-        m_resourceCount = resourceCount;
-
         SlangInt subObjectRangeCount = m_elementTypeLayout->getSubObjectRangeCount();
         for (SlangInt r = 0; r < subObjectRangeCount; ++r)
         {
@@ -703,6 +706,9 @@ public:
     size_t getSize() { return m_size; }
     Index getResourceCount() const { return m_resourceCount; }
     Index getSubObjectCount() const { return m_subObjectCount; }
+    List<SubObjectRangeInfo>& getSubObjectRanges() { return subObjectRanges; }
+    BindingRangeInfo getBindingRange(Index index) { return m_bindingRanges[index]; }
+    Index getBindingRangeCount() const { return m_bindingRanges.getCount(); }
 };
 
 class CPUEntryPointLayout : public CPUShaderObjectLayout
@@ -763,32 +769,64 @@ public:
     CPUEntryPointLayout* getEntryPoint(Index index) { return m_entryPointLayouts[index]; }
 };
 
-class CPUShaderObject : public ShaderObjectBase
+class CPUShaderObjectData
 {
 public:
-    void* m_data = nullptr;
+    Slang::List<char> m_ordinaryData;
+    // Any "ordinary" / uniform data for this object
+    Slang::RefPtr<CPUBufferResource> m_bufferResource;
+    Slang::RefPtr<CPUBufferView> m_bufferView;
 
-    ~CPUShaderObject()
+    Index getCount() { return m_ordinaryData.getCount(); }
+    void setCount(Index count) { m_ordinaryData.setCount(count); }
+    char* getBuffer() { return m_ordinaryData.getBuffer(); }
+
+    ~CPUShaderObjectData()
     {
-        free(m_data);
+        // m_bufferResource's data is managed by m_ordinaryData so we
+        // set it to null to prevent m_bufferResource from freeing it.
+        if (m_bufferResource)
+            m_bufferResource->m_data = nullptr;
+    }
+
+    /// Returns a StructuredBuffer resource view for GPU access into the buffer content.
+    /// Creates a StructuredBuffer resource if it has not been created.
+    ResourceViewBase* getResourceView(
+        RendererBase* device,
+        slang::TypeLayoutReflection* elementLayout,
+        slang::BindingType bindingType)
+    {
+        SLANG_UNUSED(device);
+        if (!m_bufferResource)
+        {
+            IBufferResource::Desc desc = {};
+            desc.type = IResource::Type::Buffer;
+            desc.elementSize = (int)elementLayout->getSize();
+            m_bufferResource = new CPUBufferResource(desc);
+
+            IResourceView::Desc viewDesc = {};
+            viewDesc.type = IResourceView::Type::UnorderedAccess;
+            viewDesc.format = Format::Unknown;
+            m_bufferView = new CPUBufferView(viewDesc, m_bufferResource);
+            
+        }
+        m_bufferResource->getDesc()->sizeInBytes = m_ordinaryData.getCount();
+        m_bufferResource->m_data = m_ordinaryData.getBuffer();
+        return m_bufferView.Ptr();
     }
+};
+
+class CPUShaderObject
+    : public ShaderObjectBaseImpl<CPUShaderObject, CPUShaderObjectLayout, CPUShaderObjectData>
+{
+    typedef ShaderObjectBaseImpl<CPUShaderObject, CPUShaderObjectLayout, CPUShaderObjectData> Super;
 
-    List<RefPtr<CPUShaderObject>> m_objects;
+public:
     List<RefPtr<CPUResourceView>> m_resources;
 
     virtual SLANG_NO_THROW Result SLANG_MCALL
         init(IDevice* device, CPUShaderObjectLayout* typeLayout);
 
-    CPUShaderObjectLayout* getLayout()
-    {
-        return static_cast<CPUShaderObjectLayout*>(m_layout.Ptr());
-    }
-
-    virtual SLANG_NO_THROW slang::TypeLayoutReflection* SLANG_MCALL getElementTypeLayout() override
-    {
-        return getLayout()->getElementTypeLayout();
-    }
-
     virtual SLANG_NO_THROW UInt SLANG_MCALL getEntryPointCount() override { return 0; }
     virtual SLANG_NO_THROW Result SLANG_MCALL
         getEntryPoint(UInt index, IShaderObject** outEntryPoint) override
@@ -799,143 +837,8 @@ public:
     virtual SLANG_NO_THROW Result SLANG_MCALL
         setData(ShaderOffset const& offset, void const* data, size_t size) override
     {
-        size = Math::Min(size, getLayout()->getSize() - offset.uniformOffset);
-        memcpy((char*)m_data + offset.uniformOffset, data, size);
-        return SLANG_OK;
-    }
-    virtual SLANG_NO_THROW Result SLANG_MCALL getObject(
-        ShaderOffset const& offset,
-        IShaderObject**     outObject) override
-    {
-        auto layout = getLayout();
-
-        auto bindingRangeIndex = offset.bindingRangeIndex;
-        SLANG_ASSERT(bindingRangeIndex >= 0);
-        SLANG_ASSERT(bindingRangeIndex < layout->m_bindingRanges.getCount());
-
-        auto& bindingRange = layout->m_bindingRanges[bindingRangeIndex];
-        auto subObjectIndex = bindingRange.baseIndex + offset.bindingArrayIndex;
-        auto& subObject = m_objects[subObjectIndex];
-
-        returnComPtr(outObject, subObject);
-
-        return SLANG_OK;
-    }
-    virtual SLANG_NO_THROW Result SLANG_MCALL setObject(
-        ShaderOffset const& offset,
-        IShaderObject*      object) override
-    {
-        auto layout = getLayout();
-
-        auto bindingRangeIndex = offset.bindingRangeIndex;
-        SLANG_ASSERT(bindingRangeIndex >= 0);
-        SLANG_ASSERT(bindingRangeIndex < layout->m_bindingRanges.getCount());
-
-        auto& bindingRange = layout->m_bindingRanges[bindingRangeIndex];
-        auto subObjectIndex = bindingRange.baseIndex + offset.bindingArrayIndex;
-
-        CPUShaderObject* subObject = static_cast<CPUShaderObject*>(object);
-        m_objects[subObjectIndex] = subObject;
-
-        switch( bindingRange.bindingType )
-        {
-        default:
-            SLANG_RETURN_ON_FAIL(setData(offset, &subObject->m_data, sizeof(void*)));
-            break;
-
-        // If the range being assigned into represents an interface/existential-type leaf field,
-        // then we need to consider how the `object` being assigned here affects specialization.
-        // We may also need to assign some data from the sub-object into the ordinary data
-        // buffer for the parent object.
-        //
-        case slang::BindingType::ExistentialValue:
-            {
-                auto renderer = getRenderer();
-
-                ComPtr<slang::ISession> slangSession;
-                SLANG_RETURN_ON_FAIL(renderer->getSlangSession(slangSession.writeRef()));
-
-                // A leaf field of interface type is laid out inside of the parent object
-                // as a tuple of `(RTTI, WitnessTable, Payload)`. The layout of these fields
-                // is a contract between the compiler and any runtime system, so we will
-                // need to rely on details of the binary layout.
-
-                // We start by querying the layout/type of the concrete value that the application
-                // is trying to store into the field, and also the layout/type of the leaf
-                // existential-type field itself.
-                //
-                auto concreteTypeLayout = subObject->getElementTypeLayout();
-                auto concreteType = concreteTypeLayout->getType();
-                //
-                auto existentialTypeLayout = layout->getElementTypeLayout()->getBindingRangeLeafTypeLayout(bindingRangeIndex);
-                auto existentialType = existentialTypeLayout->getType();
-
-                // The first field of the tuple (offset zero) is the run-time type information (RTTI)
-                // ID for the concrete type being stored into the field.
-                //
-                // TODO: We need to be able to gather the RTTI type ID from `object` and then
-                // use `setData(offset, &TypeID, sizeof(TypeID))`.
-
-                // The second field of the tuple (offset 8) is the ID of the "witness" for the
-                // conformance of the concrete type to the interface used by this field.
-                //
-                auto witnessTableOffset = offset;
-                witnessTableOffset.uniformOffset += 8;
-                //
-                // Conformances of a type to an interface are computed and then stored by the
-                // Slang runtime, so we can look up the ID for this particular conformance (which
-                // will create it on demand).
-                //
-                // Note: If the type doesn't actually conform to the required interface for
-                // this sub-object range, then this is the point where we will detect that
-                // fact and error out.
-                //
-                uint32_t conformanceID = 0xFFFFFFFF;
-                SLANG_RETURN_ON_FAIL(slangSession->getTypeConformanceWitnessSequentialID(
-                    concreteType, existentialType, &conformanceID));
-                //
-                // Once we have the conformance ID, then we can write it into the object
-                // at the required offset.
-                //
-                SLANG_RETURN_ON_FAIL(setData(witnessTableOffset, &conformanceID, sizeof(conformanceID)));
-
-                // The third field of the tuple (offset 16) is the "payload" that is supposed to
-                // hold the data for a value of the given concrete type.
-                //
-                auto payloadOffset = offset;
-                payloadOffset.uniformOffset += 16;
-
-                // There are two cases we need to consider here for how the payload might be used:
-                //
-                // * If the concrete type of the value being bound is one that can "fit" into the
-                //   available payload space,  then it should be stored in the payload.
-                //
-                // * If the concrete type of the value cannot fit in the payload space, then it
-                //   will need to be stored somewhere else.
-                //
-                if(_doesValueFitInExistentialPayload(concreteTypeLayout, existentialTypeLayout))
-                {
-                    // If the value can fit in the payload area, then we will go ahead and copy
-                    // its bytes into that area.
-                    //
-                    auto valueSize = concreteTypeLayout->getSize();
-                    SLANG_RETURN_ON_FAIL(setData(payloadOffset, subObject->m_data, valueSize));
-                }
-                else
-                {
-                    // If the value cannot fit in the payload area, then we will pass a pointer
-                    // to the sub-object instead.
-                    //
-                    // Note: The Slang compiler does not currently emit code that handles the
-                    // pointer case, but that is the expected implementation for values
-                    // that do not fit into the fixed-size payload.
-                    //
-                    SLANG_RETURN_ON_FAIL(setData(payloadOffset, &subObject->m_data, sizeof(void*)));
-                }
-            }
-            break;
-        }
-
+        size = Math::Min(size, (size_t)m_data.getCount() - offset.uniformOffset);
+        memcpy((char*)m_data.getBuffer() + offset.uniformOffset, data, size);
         return SLANG_OK;
     }
     virtual SLANG_NO_THROW Result SLANG_MCALL
@@ -989,6 +892,31 @@ public:
         return SLANG_OK;
     }
     virtual SLANG_NO_THROW Result SLANG_MCALL
+        setObject(ShaderOffset const& offset, IShaderObject* object) override
+    {
+        SLANG_RETURN_ON_FAIL(Super::setObject(offset, object));
+
+        auto bindingRangeIndex = offset.bindingRangeIndex;
+        auto& bindingRange = getLayout()->m_bindingRanges[bindingRangeIndex];
+
+        CPUShaderObject* subObject = static_cast<CPUShaderObject*>(object);
+
+        switch (bindingRange.bindingType)
+        {
+        default:
+            {
+                void* bufferPtr = subObject->m_data.getBuffer();
+                SLANG_RETURN_ON_FAIL(setData(offset, &bufferPtr, sizeof(void*)));
+            }
+            break;
+        case slang::BindingType::ExistentialValue:
+        case slang::BindingType::RawBuffer:
+        case slang::BindingType::MutableRawBuffer:
+            break;
+        }
+        return SLANG_OK;
+    }
+    virtual SLANG_NO_THROW Result SLANG_MCALL
         setSampler(ShaderOffset const& offset, ISamplerState* sampler) override
     {
         SLANG_UNUSED(sampler);
@@ -1003,52 +931,7 @@ public:
         return SLANG_OK;
     }
 
-    // Appends all types that are used to specialize the element type of this shader object in `args` list.
-    virtual Result collectSpecializationArgs(ExtendedShaderObjectTypeList& args) override
-    {
-        // TODO: the logic here is a copy-paste of `GraphicsCommonShaderObject::collectSpecializationArgs`,
-        // consider moving the implementation to `ShaderObjectBase` and share the logic among different implementations.
-
-        auto& subObjectRanges = getLayout()->subObjectRanges;
-        // The following logic is built on the assumption that all fields that involve existential types (and
-        // therefore require specialization) will results in a sub-object range in the type layout.
-        // This allows us to simply scan the sub-object ranges to find out all specialization arguments.
-        for (Index subObjIndex = 0; subObjIndex < subObjectRanges.getCount(); subObjIndex++)
-        {
-            // Retrieve the corresponding binding range of the sub object.
-            auto bindingRange = getLayout()->m_bindingRanges[subObjectRanges[subObjIndex].bindingRangeIndex];
-            switch (bindingRange.bindingType)
-            {
-            case slang::BindingType::ExistentialValue:
-            {
-                // A binding type of `ExistentialValue` means the sub-object represents a interface-typed field.
-                // In this case the specialization argument for this field is the actual specialized type of the bound
-                // shader object. If the shader object's type is an ordinary type without existential fields, then the
-                // type argument will simply be the ordinary type. But if the sub object's type is itself a specialized
-                // type, we need to make sure to use that type as the specialization argument.
-
-                // TODO: need to implement the case where the field is an array of existential values.
-                ExtendedShaderObjectType specializedSubObjType;
-                SLANG_RETURN_ON_FAIL(m_objects[subObjIndex]->getSpecializedShaderObjectType(&specializedSubObjType));
-                args.add(specializedSubObjType);
-                break;
-            }
-            case slang::BindingType::ParameterBlock:
-            case slang::BindingType::ConstantBuffer:
-                // Currently we only handle the case where the field's type is
-                // `ParameterBlock<SomeStruct>` or `ConstantBuffer<SomeStruct>`, where `SomeStruct` is a struct type
-                // (not directly an interface type). In this case, we just recursively collect the specialization arguments
-                // from the bound sub object.
-                SLANG_RETURN_ON_FAIL(m_objects[subObjIndex]->collectSpecializationArgs(args));
-                // TODO: we need to handle the case where the field is of the form `ParameterBlock<IFoo>`. We should treat
-                // this case the same way as the `ExistentialValue` case here, but currently we lack a mechanism to distinguish
-                // the two scenarios.
-                break;
-            }
-            // TODO: need to handle another case where specialization happens on resources fields e.g. `StructuredBuffer<IFoo>`.
-        }
-        return SLANG_OK;
-    }
+    char* getDataBuffer() { return m_data.getBuffer(); }
 };
 
 class CPUEntryPointShaderObject : public CPUShaderObject
@@ -1173,8 +1056,8 @@ private:
         varyingInput.endGroupID.y = y;
         varyingInput.endGroupID.z = z;
 
-        auto globalParamsData = m_currentRootObject->m_data;
-        auto entryPointParamsData = entryPointObject->m_data;
+        auto globalParamsData = m_currentRootObject->getDataBuffer();
+        auto entryPointParamsData = entryPointObject->getDataBuffer();
         func(&varyingInput, entryPointParamsData, globalParamsData);
     }
 
@@ -1379,10 +1262,7 @@ SlangResult CPUShaderObject::init(IDevice* device, CPUShaderObjectLayout* typeLa
     //
     auto slangLayout = getLayout()->getElementTypeLayout();
     size_t uniformSize = slangLayout->getSize();
-    if (uniformSize)
-    {
-        m_data = malloc(uniformSize);
-    }
+    m_data.setCount(uniformSize);
 
     // If the layout specifies that we have any resources or sub-objects,
     // then we need to size the appropriate arrays to account for them.
@@ -1405,6 +1285,8 @@ SlangResult CPUShaderObject::init(IDevice* device, CPUShaderObjectLayout* typeLa
         //
         if (!subObjectLayout)
             continue;
+        auto _debugname = subObjectLayout->getElementTypeLayout()->getName();
+        
         //
         // Otherwise, we will allocate a sub-object to fill
         // in each entry in this range, based on the layout
diff --git a/tools/gfx/cuda/render-cuda.cpp b/tools/gfx/cuda/render-cuda.cpp
index b919ac6b0..f60be8eda 100644
--- a/tools/gfx/cuda/render-cuda.cpp
+++ b/tools/gfx/cuda/render-cuda.cpp
@@ -70,7 +70,8 @@ struct CUDAErrorInfo
             builder << m_errorString;
         }
 
-        StdWriters::getError().put(builder.getUnownedSlice());
+        getDebugCallback()->handleMessage(DebugMessageType::Error, DebugMessageSource::Driver,
+            builder.getUnownedSlice().begin());
 
         // Slang::signalUnexpectedError(builder.getBuffer());
         return SLANG_FAIL;
@@ -82,7 +83,6 @@ struct CUDAErrorInfo
     const char* m_errorString;
 };
 
-#    if 1
 // If this code path is enabled, CUDA errors will be reported directly to StdWriter::out stream.
 
 static SlangResult _handleCUDAError(CUresult cuResult, const char* file, int line)
@@ -98,27 +98,7 @@ static SlangResult _handleCUDAError(cudaError_t error, const char* file, int lin
     return CUDAErrorInfo(file, line, cudaGetErrorName(error), cudaGetErrorString(error)).handle();
 }
 
-#        define SLANG_CUDA_HANDLE_ERROR(x) _handleCUDAError(_res, __FILE__, __LINE__)
-
-#    else
-// If this code path is enabled, errors are not reported, but can have an assert enabled
-
-static SlangResult _handleCUDAError(CUresult cuResult)
-{
-    SLANG_UNUSED(cuResult);
-    // SLANG_ASSERT(!"Failed CUDA call");
-    return SLANG_FAIL;
-}
-
-static SlangResult _handleCUDAError(cudaError_t error)
-{
-    SLANG_UNUSED(error);
-    // SLANG_ASSERT(!"Failed CUDA call");
-    return SLANG_FAIL;
-}
-
-#        define SLANG_CUDA_HANDLE_ERROR(x) _handleCUDAError(_res)
-#    endif
+#    define SLANG_CUDA_HANDLE_ERROR(x) _handleCUDAError(_res, __FILE__, __LINE__)
 
 #    define SLANG_CUDA_RETURN_ON_FAIL(x)              \
         {                                             \
@@ -251,20 +231,13 @@ public:
     RefPtr<CUDAContext> m_cudaContext;
 };
 
-class CUDAResourceView : public IResourceView, public ComObject
+class CUDAResourceView : public ResourceViewBase
 {
 public:
-    SLANG_COM_OBJECT_IUNKNOWN_ALL
-    IResourceView* getInterface(const Guid& guid)
-    {
-        if (guid == GfxGUID::IID_ISlangUnknown || guid == GfxGUID::IID_IResourceView)
-            return static_cast<IResourceView*>(this);
-        return nullptr;
-    }
-public:
     Desc desc;
     RefPtr<MemoryCUDAResource> memoryResource = nullptr;
     RefPtr<TextureCUDAResource> textureResource = nullptr;
+    void* proxyBuffer = nullptr;
 };
 
 class CUDAShaderObjectLayout : public ShaderObjectLayoutBase
@@ -275,6 +248,7 @@ public:
         slang::BindingType bindingType;
         Index count;
         Index baseIndex; // Flat index for sub-ojects
+        Index subObjectIndex;
 
         // TODO: The `uniformOffset` field should be removed,
         // since it cannot be supported by the Slang reflection
@@ -306,12 +280,9 @@ public:
 
     CUDAShaderObjectLayout(RendererBase* renderer, slang::TypeLayoutReflection* layout)
     {
-        initBase(renderer, layout);
-
-        Index subObjectCount = 0;
-        Index resourceCount = 0;
+        m_elementTypeLayout = _unwrapParameterGroups(layout, m_containerType);
 
-        m_elementTypeLayout = _unwrapParameterGroups(layout);
+        initBase(renderer, m_elementTypeLayout);
 
         // Compute the binding ranges that are used to store
         // the logical contents of the object in memory. These will relate
@@ -342,18 +313,31 @@ public:
                 descriptorSetIndex, rangeIndexInDescriptorSet);
 
             Index baseIndex = 0;
+            Index subObjectIndex = 0;
             switch (slangBindingType)
             {
             case slang::BindingType::ConstantBuffer:
             case slang::BindingType::ParameterBlock:
             case slang::BindingType::ExistentialValue:
-                baseIndex = subObjectCount;
-                subObjectCount += count;
+                baseIndex = m_subObjectCount;
+                subObjectIndex = baseIndex;
+                m_subObjectCount += count;
+                break;
+            case slang::BindingType::RawBuffer:
+            case slang::BindingType::MutableRawBuffer:
+                if (slangLeafTypeLayout->getType()->getElementType() != nullptr)
+                {
+                    // A structured buffer occupies both a resource slot and
+                    // a sub-object slot.
+                    subObjectIndex = m_subObjectCount;
+                    m_subObjectCount += count;
+                }
+                baseIndex = m_resourceCount;
+                m_resourceCount += count;
                 break;
-
             default:
-                baseIndex = resourceCount;
-                resourceCount += count;
+                baseIndex = m_resourceCount;
+                m_resourceCount += count;
                 break;
             }
 
@@ -362,12 +346,10 @@ public:
             bindingRangeInfo.count = count;
             bindingRangeInfo.baseIndex = baseIndex;
             bindingRangeInfo.uniformOffset = uniformOffset;
+            bindingRangeInfo.subObjectIndex = subObjectIndex;
             m_bindingRanges.add(bindingRangeInfo);
         }
 
-        m_subObjectCount = subObjectCount;
-        m_resourceCount = resourceCount;
-
         SlangInt subObjectRangeCount = m_elementTypeLayout->getSubObjectRangeCount();
         for (SlangInt r = 0; r < subObjectRangeCount; ++r)
         {
@@ -399,6 +381,9 @@ public:
 
     Index getResourceCount() const { return m_resourceCount; }
     Index getSubObjectCount() const { return m_subObjectCount; }
+    List<SubObjectRangeInfo>& getSubObjectRanges() { return subObjectRanges; }
+    BindingRangeInfo getBindingRange(Index index) { return m_bindingRanges[index]; }
+    Index getBindingRangeCount() const { return m_bindingRanges.getCount(); }
 };
 
 class CUDAProgramLayout : public CUDAShaderObjectLayout
@@ -439,50 +424,108 @@ public:
     }
 };
 
-class CUDAShaderObject : public ShaderObjectBase
+class CUDAShaderObjectData
 {
 public:
-    RefPtr<MemoryCUDAResource> bufferResource;
-    List<RefPtr<CUDAShaderObject>> objects;
-    List<RefPtr<CUDAResourceView>> resources;
-
-    virtual SLANG_NO_THROW Result SLANG_MCALL
-        init(IDevice* device, CUDAShaderObjectLayout* typeLayout);
-
-    CUDAShaderObjectLayout* getLayout()
+    bool isHostOnly = false;
+    Slang::RefPtr<MemoryCUDAResource> m_bufferResource;
+    Slang::RefPtr<CUDAResourceView> m_bufferView;
+    Slang::List<uint8_t> m_cpuBuffer;
+    void setCount(Index count)
     {
-        return static_cast<CUDAShaderObjectLayout*>(m_layout.Ptr());
-    }
+        if (isHostOnly)
+        {
+            m_cpuBuffer.setCount(count);
+            if (!m_bufferView)
+            {
+                IResourceView::Desc viewDesc = {};
+                viewDesc.type = IResourceView::Type::UnorderedAccess;
+                m_bufferView = new CUDAResourceView();
+                m_bufferView->proxyBuffer = m_cpuBuffer.getBuffer();
+                m_bufferView->desc = viewDesc;
+            }
+            return;
+        }
 
-    virtual SLANG_NO_THROW Result SLANG_MCALL initBuffer(IDevice* device, size_t bufferSize)
-    {
-        BufferResource::Desc bufferDesc;
-        bufferDesc.type = IResource::Type::Buffer;
-        bufferDesc.defaultState = ResourceState::ConstantBuffer;
-        bufferDesc.allowedStates =
-            ResourceStateSet(ResourceState::ConstantBuffer, ResourceState::CopyDestination);
-        bufferDesc.sizeInBytes = bufferSize;
-        bufferDesc.cpuAccessFlags |= IResource::AccessFlag::Write;
-        ComPtr<IBufferResource> constantBuffer;
-        SLANG_RETURN_ON_FAIL(device->createBufferResource(bufferDesc, nullptr, constantBuffer.writeRef()));
-        bufferResource = static_cast<MemoryCUDAResource*>(constantBuffer.get());
-        return SLANG_OK;
+        if (!m_bufferResource)
+        {
+            IBufferResource::Desc desc;
+            desc.type = IResource::Type::Buffer;
+            desc.sizeInBytes = count;
+            m_bufferResource = new MemoryCUDAResource(desc);
+            if (count)
+                cudaMalloc(&m_bufferResource->m_cudaMemory, (size_t)count);
+            IResourceView::Desc viewDesc = {};
+            viewDesc.type = IResourceView::Type::UnorderedAccess;
+            m_bufferView = new CUDAResourceView();
+            m_bufferView->memoryResource = m_bufferResource;
+            m_bufferView->desc = viewDesc;
+        }
+        auto oldSize = m_bufferResource->getDesc()->sizeInBytes;
+        if ((size_t)count != oldSize)
+        {
+            void* newMemory = nullptr;
+            if (count)
+            {
+                cudaMalloc(&newMemory, (size_t)count);
+            }
+            if (oldSize)
+            {
+                cudaMemcpy(
+                    newMemory,
+                    m_bufferResource->m_cudaMemory,
+                    Math::Min((size_t)count, oldSize),
+                    cudaMemcpyDefault);
+            }
+            cudaFree(m_bufferResource->m_cudaMemory);
+            m_bufferResource->m_cudaMemory = newMemory;
+            m_bufferResource->getDesc()->sizeInBytes = count;
+        }
     }
 
-    virtual SLANG_NO_THROW void* SLANG_MCALL getBuffer()
+    Slang::Index getCount()
     {
-        return bufferResource ? bufferResource->m_cudaMemory : nullptr;
+        if (isHostOnly)
+            return m_cpuBuffer.getCount();
+        if (m_bufferResource)
+            return (Slang::Index)(m_bufferResource->getDesc()->sizeInBytes);
+        else
+            return 0;
     }
 
-    virtual SLANG_NO_THROW size_t SLANG_MCALL getBufferSize()
+    void* getBuffer()
     {
-        return bufferResource ? bufferResource->getDesc()->sizeInBytes : 0;
+        if (isHostOnly)
+            return m_cpuBuffer.getBuffer();
+
+        if (m_bufferResource)
+            return m_bufferResource->m_cudaMemory;
+        return nullptr;
     }
 
-    virtual SLANG_NO_THROW slang::TypeLayoutReflection* SLANG_MCALL getElementTypeLayout() override
+    /// Returns a resource view for GPU access into the buffer content.
+    ResourceViewBase* getResourceView(
+        RendererBase* device,
+        slang::TypeLayoutReflection* elementLayout,
+        slang::BindingType bindingType)
     {
-        return getLayout()->getElementTypeLayout();
+        SLANG_UNUSED(device);
+        m_bufferResource->getDesc()->elementSize = (int)elementLayout->getSize();
+        return m_bufferView.Ptr();
     }
+};
+
+class CUDAShaderObject
+    : public ShaderObjectBaseImpl<CUDAShaderObject, CUDAShaderObjectLayout, CUDAShaderObjectData>
+{
+    typedef ShaderObjectBaseImpl<CUDAShaderObject, CUDAShaderObjectLayout, CUDAShaderObjectData>
+        Super;
+
+public:
+    List<RefPtr<CUDAResourceView>> resources;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL
+        init(IDevice* device, CUDAShaderObjectLayout* typeLayout);
 
     virtual SLANG_NO_THROW UInt SLANG_MCALL getEntryPointCount() override { return 0; }
     virtual SLANG_NO_THROW Result SLANG_MCALL
@@ -494,164 +537,17 @@ public:
     virtual SLANG_NO_THROW Result SLANG_MCALL
         setData(ShaderOffset const& offset, void const* data, size_t size) override
     {
-        size = Math::Min(size, bufferResource->getDesc()->sizeInBytes - offset.uniformOffset);
+        size = Math::Min(size, (size_t)m_data.getCount() - offset.uniformOffset);
         SLANG_CUDA_RETURN_ON_FAIL(cudaMemcpy(
-            (uint8_t*)bufferResource->m_cudaMemory + offset.uniformOffset,
-            data,
-            size,
-            cudaMemcpyHostToDevice));
+            (uint8_t*)m_data.getBuffer() + offset.uniformOffset, data, size, cudaMemcpyDefault));
         return SLANG_OK;
     }
     virtual SLANG_NO_THROW Result SLANG_MCALL
-        setDeviceData(size_t offset, void* data, size_t size)
-    {
-        size = Math::Min(size, bufferResource->getDesc()->sizeInBytes - offset);
-        SLANG_CUDA_RETURN_ON_FAIL(cudaMemcpy(
-            (uint8_t*)bufferResource->m_cudaMemory + offset,
-            data,
-            size,
-            cudaMemcpyHostToDevice));
-        return SLANG_OK;
-    }
-    virtual SLANG_NO_THROW Result SLANG_MCALL
-        getObject(ShaderOffset const& offset, IShaderObject** object) override
+        setResource(ShaderOffset const& offset, IResourceView* resourceView) override
     {
-        auto subObjectIndex =
-            getLayout()->m_bindingRanges[offset.bindingRangeIndex].baseIndex + offset.bindingArrayIndex;
-
-        SLANG_ASSERT(subObjectIndex < objects.getCount());
-        if(subObjectIndex >= objects.getCount())
-            return SLANG_E_INVALID_ARG;
-
-        if (subObjectIndex >= objects.getCount())
-        {
-            *object = nullptr;
+        if (!resourceView)
             return SLANG_OK;
-        }
-        returnComPtr(object, objects[subObjectIndex]);
-        return SLANG_OK;
-    }
-    virtual SLANG_NO_THROW Result SLANG_MCALL
-        setObject(ShaderOffset const& offset, IShaderObject* object) override
-    {
-        auto layout = getLayout();
-
-        auto bindingRangeIndex = offset.bindingRangeIndex;
-        SLANG_ASSERT(bindingRangeIndex >= 0);
-        SLANG_ASSERT(bindingRangeIndex < layout->m_bindingRanges.getCount());
-
-        auto& bindingRange = layout->m_bindingRanges[bindingRangeIndex];
-
-        auto subObjectIndex = bindingRange.baseIndex + offset.bindingArrayIndex;
-        auto subObject = dynamic_cast<CUDAShaderObject*>(object);
-
-        // TODO: We should really not need to retain the objects here
-        objects[subObjectIndex] = subObject;
-
-        switch( bindingRange.bindingType )
-        {
-        default:
-            SLANG_RETURN_ON_FAIL(setData(offset, &subObject->bufferResource->m_cudaMemory, sizeof(void*)));
-            break;
-
-        // If the range being assigned into represents an interface/existential-type leaf field,
-        // then we need to consider how the `object` being assigned here affects specialization.
-        // We may also need to assign some data from the sub-object into the ordinary data
-        // buffer for the parent object.
-        //
-        case slang::BindingType::ExistentialValue:
-            {
-                auto renderer = getRenderer();
-
-                ComPtr<slang::ISession> slangSession;
-                SLANG_RETURN_ON_FAIL(renderer->getSlangSession(slangSession.writeRef()));
-
-                // A leaf field of interface type is laid out inside of the parent object
-                // as a tuple of `(RTTI, WitnessTable, Payload)`. The layout of these fields
-                // is a contract between the compiler and any runtime system, so we will
-                // need to rely on details of the binary layout.
-
-                // We start by querying the layout/type of the concrete value that the application
-                // is trying to store into the field, and also the layout/type of the leaf
-                // existential-type field itself.
-                //
-                auto concreteTypeLayout = subObject->getElementTypeLayout();
-                auto concreteType = concreteTypeLayout->getType();
-                //
-                auto existentialTypeLayout = layout->getElementTypeLayout()->getBindingRangeLeafTypeLayout(bindingRangeIndex);
-                auto existentialType = existentialTypeLayout->getType();
-
-                // The first field of the tuple (offset zero) is the run-time type information (RTTI)
-                // ID for the concrete type being stored into the field.
-                //
-                // TODO: We need to be able to gather the RTTI type ID from `object` and then
-                // use `setData(offset, &TypeID, sizeof(TypeID))`.
-
-                // The second field of the tuple (offset 8) is the ID of the "witness" for the
-                // conformance of the concrete type to the interface used by this field.
-                //
-                auto witnessTableOffset = offset;
-                witnessTableOffset.uniformOffset += 8;
-                //
-                // Conformances of a type to an interface are computed and then stored by the
-                // Slang runtime, so we can look up the ID for this particular conformance (which
-                // will create it on demand).
-                //
-                // Note: If the type doesn't actually conform to the required interface for
-                // this sub-object range, then this is the point where we will detect that
-                // fact and error out.
-                //
-                uint32_t conformanceID = 0xFFFFFFFF;
-                SLANG_RETURN_ON_FAIL(slangSession->getTypeConformanceWitnessSequentialID(
-                    concreteType, existentialType, &conformanceID));
-                //
-                // Once we have the conformance ID, then we can write it into the object
-                // at the required offset.
-                //
-                SLANG_RETURN_ON_FAIL(setData(witnessTableOffset, &conformanceID, sizeof(conformanceID)));
-
-                // The third field of the tuple (offset 16) is the "payload" that is supposed to
-                // hold the data for a value of the given concrete type.
-                //
-                auto payloadOffset = offset;
-                payloadOffset.uniformOffset += 16;
-
-                // There are two cases we need to consider here for how the payload might be used:
-                //
-                // * If the concrete type of the value being bound is one that can "fit" into the
-                //   available payload space,  then it should be stored in the payload.
-                //
-                // * If the concrete type of the value cannot fit in the payload space, then it
-                //   will need to be stored somewhere else.
-                //
-                if(_doesValueFitInExistentialPayload(concreteTypeLayout, existentialTypeLayout))
-                {
-                    // If the value can fit in the payload area, then we will go ahead and copy
-                    // its bytes into that area.
-                    //
-                    auto valueSize = concreteTypeLayout->getSize();
-                    SLANG_RETURN_ON_FAIL(setDeviceData(payloadOffset.uniformOffset, subObject->getBuffer(), valueSize));
-                }
-                else
-                {
-                    // If the value cannot fit in the payload area, then we will pass a pointer
-                    // to the sub-object instead.
-                    //
-                    // Note: The Slang compiler does not currently emit code that handles the
-                    // pointer case, but that is the expected implementation for values
-                    // that do not fit into the fixed-size payload.
-                    //
-                    SLANG_RETURN_ON_FAIL(setData(payloadOffset, &subObject->bufferResource->m_cudaMemory, sizeof(void*)));
-                }
-            }
-            break;
-        }
 
-        return SLANG_OK;
-    }
-    virtual SLANG_NO_THROW Result SLANG_MCALL
-        setResource(ShaderOffset const& offset, IResourceView* resourceView) override
-    {
         auto layout = getLayout();
 
         auto bindingRangeIndex = offset.bindingRangeIndex;
@@ -678,7 +574,7 @@ public:
                 setData(offset, &handle, sizeof(uint64_t));
             }
         }
-        else
+        else if (cudaView->memoryResource)
         {
             auto handle = cudaView->memoryResource->getBindlessHandle();
             setData(offset, &handle, sizeof(handle));
@@ -689,7 +585,42 @@ public:
             if (desc.elementSize > 1)
                 size /= desc.elementSize;
             setData(sizeOffset, &size, sizeof(size));
+        }
+        else if (cudaView->proxyBuffer)
+        {
+            auto handle = cudaView->proxyBuffer;
+            setData(offset, &handle, sizeof(handle));
+            auto sizeOffset = offset;
+            sizeOffset.uniformOffset += sizeof(handle);
+            auto& desc = *cudaView->memoryResource->getDesc();
+            size_t size = desc.sizeInBytes;
+            if (desc.elementSize > 1)
+                size /= desc.elementSize;
+            setData(sizeOffset, &size, sizeof(size));
+        }
+        return SLANG_OK;
+    }
+    virtual SLANG_NO_THROW Result SLANG_MCALL
+        setObject(ShaderOffset const& offset, IShaderObject* object) override
+    {
+        SLANG_RETURN_ON_FAIL(Super::setObject(offset, object));
 
+        auto bindingRangeIndex = offset.bindingRangeIndex;
+        auto& bindingRange = getLayout()->m_bindingRanges[bindingRangeIndex];
+
+        CUDAShaderObject* subObject = static_cast<CUDAShaderObject*>(object);
+        switch (bindingRange.bindingType)
+        {
+        default:
+            {
+                void* subObjectDataBuffer = subObject->getBuffer();
+                SLANG_RETURN_ON_FAIL(setData(offset, &subObjectDataBuffer, sizeof(void*)));
+            }
+            break;
+        case slang::BindingType::ExistentialValue:
+        case slang::BindingType::RawBuffer:
+        case slang::BindingType::MutableRawBuffer:
+            break;
         }
         return SLANG_OK;
     }
@@ -707,107 +638,12 @@ public:
         setResource(offset, textureView);
         return SLANG_OK;
     }
-
-    // Appends all types that are used to specialize the element type of this shader object in `args` list.
-    virtual Result collectSpecializationArgs(ExtendedShaderObjectTypeList& args) override
-    {
-        // TODO: the logic here is a copy-paste of `GraphicsCommonShaderObject::collectSpecializationArgs`,
-        // consider moving the implementation to `ShaderObjectBase` and share the logic among different implementations.
-
-        auto& subObjectRanges = getLayout()->subObjectRanges;
-        // The following logic is built on the assumption that all fields that involve existential types (and
-        // therefore require specialization) will results in a sub-object range in the type layout.
-        // This allows us to simply scan the sub-object ranges to find out all specialization arguments.
-        for (Index subObjIndex = 0; subObjIndex < subObjectRanges.getCount(); subObjIndex++)
-        {
-            // Retrieve the corresponding binding range of the sub object.
-            auto bindingRange = getLayout()->m_bindingRanges[subObjectRanges[subObjIndex].bindingRangeIndex];
-            switch (bindingRange.bindingType)
-            {
-            case slang::BindingType::ExistentialValue:
-            {
-                // A binding type of `ExistentialValue` means the sub-object represents a interface-typed field.
-                // In this case the specialization argument for this field is the actual specialized type of the bound
-                // shader object. If the shader object's type is an ordinary type without existential fields, then the
-                // type argument will simply be the ordinary type. But if the sub object's type is itself a specialized
-                // type, we need to make sure to use that type as the specialization argument.
-
-                // TODO: need to implement the case where the field is an array of existential values.
-                ExtendedShaderObjectType specializedSubObjType;
-                SLANG_RETURN_ON_FAIL(objects[subObjIndex]->getSpecializedShaderObjectType(&specializedSubObjType));
-                args.add(specializedSubObjType);
-                break;
-            }
-            case slang::BindingType::ParameterBlock:
-            case slang::BindingType::ConstantBuffer:
-                // Currently we only handle the case where the field's type is
-                // `ParameterBlock<SomeStruct>` or `ConstantBuffer<SomeStruct>`, where `SomeStruct` is a struct type
-                // (not directly an interface type). In this case, we just recursively collect the specialization arguments
-                // from the bound sub object.
-                SLANG_RETURN_ON_FAIL(objects[subObjIndex]->collectSpecializationArgs(args));
-                // TODO: we need to handle the case where the field is of the form `ParameterBlock<IFoo>`. We should treat
-                // this case the same way as the `ExistentialValue` case here, but currently we lack a mechanism to distinguish
-                // the two scenarios.
-                break;
-            }
-            // TODO: need to handle another case where specialization happens on resources fields e.g. `StructuredBuffer<IFoo>`.
-        }
-        return SLANG_OK;
-    }
 };
 
 class CUDAEntryPointShaderObject : public CUDAShaderObject
 {
 public:
-    void* hostBuffer = nullptr;
-    size_t uniformBufferSize = 0;
-    // Override buffer allocation so we store all uniform data on host memory instead of device memory.
-    virtual SLANG_NO_THROW Result SLANG_MCALL initBuffer(IDevice* device, size_t bufferSize) override
-    {
-        SLANG_UNUSED(device);
-        uniformBufferSize = bufferSize;
-        hostBuffer = malloc(bufferSize);
-        return SLANG_OK;
-    }
-
-    virtual SLANG_NO_THROW Result SLANG_MCALL
-        setData(ShaderOffset const& offset, void const* data, size_t size) override
-    {
-        size = Math::Min(size, uniformBufferSize - offset.uniformOffset);
-        memcpy(
-            (uint8_t*)hostBuffer + offset.uniformOffset,
-            data,
-            size);
-        return SLANG_OK;
-    }
-
-    virtual SLANG_NO_THROW Result SLANG_MCALL
-        setDeviceData(size_t offset, void* data, size_t size) override
-    {
-        size = Math::Min(size, uniformBufferSize - offset);
-        SLANG_CUDA_RETURN_ON_FAIL(cudaMemcpy(
-            (uint8_t*)hostBuffer + offset,
-            data,
-            size,
-            cudaMemcpyDeviceToHost));
-        return SLANG_OK;
-    }
-
-
-    virtual SLANG_NO_THROW void* SLANG_MCALL getBuffer() override
-    {
-        return hostBuffer;
-    }
-
-    virtual SLANG_NO_THROW size_t SLANG_MCALL getBufferSize() override
-    {
-        return uniformBufferSize;
-    }
-
-    ~CUDAEntryPointShaderObject()
-    {
-        free(hostBuffer);
-    }
+    CUDAEntryPointShaderObject() { m_data.isHostOnly = true; }
 };
 
 class CUDARootShaderObject : public CUDAShaderObject
@@ -1233,15 +1069,12 @@ public:
                     currentPipeline->shaderProgram->cudaModule,
                     "SLANG_globalParams");
 
-                CUdeviceptr globalParamsCUDAData =
-                    currentRootObject->bufferResource
-                        ? (CUdeviceptr)currentRootObject->bufferResource->getBindlessHandle()
-                        : 0;
+                CUdeviceptr globalParamsCUDAData = (CUdeviceptr)currentRootObject->getBuffer();
                 cudaMemcpyAsync(
                     (void*)globalParamsSymbol,
                     (void*)globalParamsCUDAData,
                     globalParamsSymbolSize,
-                    cudaMemcpyDeviceToDevice,
+                    cudaMemcpyDefault,
                     0);
             }
             //
@@ -1818,7 +1651,7 @@ public:
         SLANG_CUDA_RETURN_ON_FAIL(cudaMallocManaged(&resource->m_cudaMemory, desc.sizeInBytes));
         if (initData)
         {
-            SLANG_CUDA_RETURN_ON_FAIL(cudaMemcpy(resource->m_cudaMemory, initData, desc.sizeInBytes, cudaMemcpyHostToDevice));
+            SLANG_CUDA_RETURN_ON_FAIL(cudaMemcpy(resource->m_cudaMemory, initData, desc.sizeInBytes, cudaMemcpyDefault));
         }
         returnComPtr(outResource, resource);
         return SLANG_OK;
@@ -2080,7 +1913,7 @@ SlangResult CUDAShaderObject::init(IDevice* device, CUDAShaderObjectLayout* type
     size_t uniformSize = slangLayout->getSize();
     if (uniformSize)
     {
-        initBuffer(device, uniformSize);
+        m_data.setCount((Index)uniformSize);
     }
 
     // If the layout specifies that we have any resources or sub-objects,
@@ -2090,10 +1923,7 @@ SlangResult CUDAShaderObject::init(IDevice* device, CUDAShaderObjectLayout* type
     // and not just the number of resource/sub-object ranges.
     //
     resources.setCount(typeLayout->getResourceCount());
-    objects.setCount(typeLayout->getSubObjectCount());
-
-    Index subObjectCount = slangLayout->getSubObjectRangeCount();
-    objects.setCount(subObjectCount);
+    m_objects.setCount(typeLayout->getSubObjectCount());
 
     for (auto subObjectRange : getLayout()->subObjectRanges)
     {
diff --git a/tools/gfx/d3d/d3d-util.cpp b/tools/gfx/d3d/d3d-util.cpp
index 651d61abf..7db1ce585 100644
--- a/tools/gfx/d3d/d3d-util.cpp
+++ b/tools/gfx/d3d/d3d-util.cpp
@@ -139,8 +139,7 @@ D3D12_RESOURCE_STATES D3DUtil::translateResourceState(ResourceState state)
         return D3D12_RESOURCE_STATE_PIXEL_SHADER_RESOURCE |
                D3D12_RESOURCE_STATE_NON_PIXEL_SHADER_RESOURCE;
     case gfx::ResourceState::UnorderedAccess:
-        return D3D12_RESOURCE_STATE_UNORDERED_ACCESS | D3D12_RESOURCE_STATE_PIXEL_SHADER_RESOURCE |
-               D3D12_RESOURCE_STATE_NON_PIXEL_SHADER_RESOURCE;
+        return D3D12_RESOURCE_STATE_UNORDERED_ACCESS;
     case gfx::ResourceState::RenderTarget:
         return D3D12_RESOURCE_STATE_RENDER_TARGET;
     case gfx::ResourceState::DepthRead:
diff --git a/tools/gfx/d3d11/render-d3d11.cpp b/tools/gfx/d3d11/render-d3d11.cpp
index 6f6ba4faf..d4a55543c 100644
--- a/tools/gfx/d3d11/render-d3d11.cpp
+++ b/tools/gfx/d3d11/render-d3d11.cpp
@@ -212,17 +212,9 @@ protected:
     };
 
 
-    class ResourceViewImpl : public IResourceView, public ComObject
+    class ResourceViewImpl : public ResourceViewBase
     {
     public:
-        SLANG_COM_OBJECT_IUNKNOWN_ALL
-        IResourceView* getInterface(const Guid& guid)
-        {
-            if (guid == GfxGUID::IID_ISlangUnknown || guid == GfxGUID::IID_IResourceView)
-                return static_cast<IResourceView*>(this);
-            return nullptr;
-        }
-    public:
         enum class Type
         {
             SRV,
@@ -662,6 +654,10 @@ protected:
                 /// `t` registers.
                 ///
             uint32_t registerOffset;
+
+                /// An index into the sub-object array if this binding range is treated
+                /// as a sub-object.
+            Index subObjectIndex;
         };
 
         // Sometimes we just want to iterate over the ranges that represnet
@@ -753,10 +749,15 @@ protected:
             Index m_subObjectCount = 0;
 
             uint32_t m_totalOrdinaryDataSize = 0;
+            
+                /// The container type of this shader object. When `m_containerType` is
+                /// `StructuredBuffer` or `UnsizedArray`, this shader object represents a collection
+                /// instead of a single object.
+            ShaderObjectContainerType m_containerType = ShaderObjectContainerType::None;
 
             Result setElementTypeLayout(slang::TypeLayoutReflection* typeLayout)
             {
-                typeLayout = _unwrapParameterGroups(typeLayout);
+                typeLayout = _unwrapParameterGroups(typeLayout, m_containerType);
 
                 m_elementTypeLayout = typeLayout;
 
@@ -783,9 +784,31 @@ protected:
                     case slang::BindingType::ParameterBlock:
                     case slang::BindingType::ExistentialValue:
                         bindingRangeInfo.baseIndex = m_subObjectCount;
+                        bindingRangeInfo.subObjectIndex = m_subObjectCount;
                         m_subObjectCount += count;
                         break;
-
+                    case slang::BindingType::RawBuffer:
+                    case slang::BindingType::MutableRawBuffer:
+                        if (slangLeafTypeLayout->getType()->getElementType() != nullptr)
+                        {
+                            // A structured buffer occupies both a resource slot and
+                            // a sub-object slot.
+                            bindingRangeInfo.subObjectIndex = m_subObjectCount;
+                            m_subObjectCount += count;
+                        }
+                        if (slangBindingType == slang::BindingType::RawBuffer)
+                        {
+                            bindingRangeInfo.baseIndex = m_srvCount;
+                            m_srvCount += count;
+                            m_srvRanges.add(r);
+                        }
+                        else
+                        {
+                            bindingRangeInfo.baseIndex = m_uavCount;
+                            m_uavCount += count;
+                            m_uavRanges.add(r);
+                        }
+                        break;
                     case slang::BindingType::Sampler:
                         bindingRangeInfo.baseIndex = m_samplerCount;
                         m_samplerCount += count;
@@ -794,8 +817,6 @@ protected:
 
                     case slang::BindingType::CombinedTextureSampler:
                         break;
-
-                    case slang::BindingType::MutableRawBuffer:
                     case slang::BindingType::MutableTexture:
                     case slang::BindingType::MutableTypedBuffer:
                         bindingRangeInfo.baseIndex = m_uavCount;
@@ -953,8 +974,6 @@ protected:
 
         BindingRangeInfo const& getBindingRange(Index index) { return m_bindingRanges[index]; }
 
-        slang::TypeLayoutReflection* getElementTypeLayout() { return m_elementTypeLayout; }
-
         Index getSRVCount() { return m_srvCount; }
         Index getSamplerCount() { return m_samplerCount; }
         Index getUAVCount() { return m_uavCount; }
@@ -1002,6 +1021,7 @@ protected:
 
             m_totalOrdinaryDataSize = builder->m_totalOrdinaryDataSize;
 
+            m_containerType = builder->m_containerType;
             return SLANG_OK;
         }
 
@@ -1127,7 +1147,11 @@ protected:
         SimpleBindingOffset m_pendingDataOffset;
     };
 
-    class ShaderObjectImpl : public ShaderObjectBase
+    class ShaderObjectImpl
+        : public ShaderObjectBaseImpl<
+              ShaderObjectImpl,
+              ShaderObjectLayoutImpl,
+              SimpleShaderObjectData>
     {
     public:
         static Result create(
@@ -1153,24 +1177,14 @@ protected:
             return SLANG_OK;
         }
 
-        ShaderObjectLayoutImpl* getLayout()
-        {
-            return static_cast<ShaderObjectLayoutImpl*>(m_layout.Ptr());
-        }
-
-        SLANG_NO_THROW slang::TypeLayoutReflection* SLANG_MCALL getElementTypeLayout() SLANG_OVERRIDE
-        {
-            return m_layout->getElementTypeLayout();
-        }
-
         SLANG_NO_THROW Result SLANG_MCALL
             setData(ShaderOffset const& inOffset, void const* data, size_t inSize) SLANG_OVERRIDE
         {
             Index offset = inOffset.uniformOffset;
             Index size = inSize;
 
-            char* dest = m_ordinaryData.getBuffer();
-            Index availableSize = m_ordinaryData.getCount();
+            char* dest = m_data.getBuffer();
+            Index availableSize = m_data.getCount();
 
             // TODO: We really should bounds-check access rather than silently ignoring sets
             // that are too large, but we have several test cases that set more data than
@@ -1191,128 +1205,7 @@ protected:
             return SLANG_OK;
         }
 
-        virtual SLANG_NO_THROW Result SLANG_MCALL
-            setObject(ShaderOffset const& offset, IShaderObject* object)
-            SLANG_OVERRIDE
-        {
-            if (offset.bindingRangeIndex < 0)
-                return SLANG_E_INVALID_ARG;
-            auto layout = getLayout();
-            if (offset.bindingRangeIndex >= layout->getBindingRangeCount())
-                return SLANG_E_INVALID_ARG;
-
-            auto subObject = static_cast<ShaderObjectImpl*>(object);
-
-            auto bindingRangeIndex = offset.bindingRangeIndex;
-            auto& bindingRange = layout->getBindingRange(bindingRangeIndex);
-
-            m_objects[bindingRange.baseIndex + offset.bindingArrayIndex] = subObject;
-
-            // If the range being assigned into represents an interface/existential-type leaf field,
-            // then we need to consider how the `object` being assigned here affects specialization.
-            // We may also need to assign some data from the sub-object into the ordinary data
-            // buffer for the parent object.
-            //
-            if (bindingRange.bindingType == slang::BindingType::ExistentialValue)
-            {
-                // A leaf field of interface type is laid out inside of the parent object
-                // as a tuple of `(RTTI, WitnessTable, Payload)`. The layout of these fields
-                // is a contract between the compiler and any runtime system, so we will
-                // need to rely on details of the binary layout.
-
-                // We start by querying the layout/type of the concrete value that the application
-                // is trying to store into the field, and also the layout/type of the leaf
-                // existential-type field itself.
-                //
-                auto concreteTypeLayout = subObject->getElementTypeLayout();
-                auto concreteType = concreteTypeLayout->getType();
-                //
-                auto existentialTypeLayout = layout->getElementTypeLayout()->getBindingRangeLeafTypeLayout(bindingRangeIndex);
-                auto existentialType = existentialTypeLayout->getType();
-
-                // The first field of the tuple (offset zero) is the run-time type information (RTTI)
-                // ID for the concrete type being stored into the field.
-                //
-                // TODO: We need to be able to gather the RTTI type ID from `object` and then
-                // use `setData(offset, &TypeID, sizeof(TypeID))`.
-
-                // The second field of the tuple (offset 8) is the ID of the "witness" for the
-                // conformance of the concrete type to the interface used by this field.
-                //
-                auto witnessTableOffset = offset;
-                witnessTableOffset.uniformOffset += 8;
-                //
-                // Conformances of a type to an interface are computed and then stored by the
-                // Slang runtime, so we can look up the ID for this particular conformance (which
-                // will create it on demand).
-                //
-                ComPtr<slang::ISession> slangSession;
-                SLANG_RETURN_ON_FAIL(getRenderer()->getSlangSession(slangSession.writeRef()));
-                //
-                // Note: If the type doesn't actually conform to the required interface for
-                // this sub-object range, then this is the point where we will detect that
-                // fact and error out.
-                //
-                uint32_t conformanceID = 0xFFFFFFFF;
-                SLANG_RETURN_ON_FAIL(slangSession->getTypeConformanceWitnessSequentialID(
-                    concreteType, existentialType, &conformanceID));
-                //
-                // Once we have the conformance ID, then we can write it into the object
-                // at the required offset.
-                //
-                SLANG_RETURN_ON_FAIL(setData(witnessTableOffset, &conformanceID, sizeof(conformanceID)));
-
-                // The third field of the tuple (offset 16) is the "payload" that is supposed to
-                // hold the data for a value of the given concrete type.
-                //
-                auto payloadOffset = offset;
-                payloadOffset.uniformOffset += 16;
-
-                // There are two cases we need to consider here for how the payload might be used:
-                //
-                // * If the concrete type of the value being bound is one that can "fit" into the
-                //   available payload space,  then it should be stored in the payload.
-                //
-                // * If the concrete type of the value cannot fit in the payload space, then it
-                //   will need to be stored somewhere else.
-                //
-                if (_doesValueFitInExistentialPayload(concreteTypeLayout, existentialTypeLayout))
-                {
-                    // If the value can fit in the payload area, then we will go ahead and copy
-                    // its bytes into that area.
-                    //
-                    setData(payloadOffset, subObject->m_ordinaryData.getBuffer(), subObject->m_ordinaryData.getCount());
-                }
-                else
-                {
-                    // If the value does *not *fit in the payload area, then there is nothing
-                    // we can do at this point (beyond saving a reference to the sub-object, which
-                    // was handled above).
-                    //
-                    // Once all the sub-objects have been set into the parent object, we can
-                    // compute a specialized layout for it, and that specialized layout can tell
-                    // us where the data for these sub-objects has been laid out.
-                }
-            }
-
-            return SLANG_E_NOT_IMPLEMENTED;
-        }
-
-        virtual SLANG_NO_THROW Result SLANG_MCALL
-            getObject(ShaderOffset const& offset, IShaderObject** outObject)
-            SLANG_OVERRIDE
-        {
-            SLANG_ASSERT(outObject);
-            if (offset.bindingRangeIndex < 0)
-                return SLANG_E_INVALID_ARG;
-            auto layout = getLayout();
-            if (offset.bindingRangeIndex >= layout->getBindingRangeCount())
-                return SLANG_E_INVALID_ARG;
-            auto& bindingRange = layout->getBindingRange(offset.bindingRangeIndex);
-
-            returnComPtr(outObject, m_objects[bindingRange.baseIndex + offset.bindingArrayIndex]);
-            return SLANG_OK;
-        }
+        
 
         SLANG_NO_THROW Result SLANG_MCALL
             setResource(ShaderOffset const& offset, IResourceView* resourceView) SLANG_OVERRIDE
@@ -1359,56 +1252,6 @@ protected:
         }
 
     public:
-        // Appends all types that are used to specialize the element type of this shader object in `args` list.
-        virtual Result collectSpecializationArgs(ExtendedShaderObjectTypeList& args) override
-        {
-            auto& subObjectRanges = getLayout()->getSubObjectRanges();
-            // The following logic is built on the assumption that all fields that involve existential types (and
-            // therefore require specialization) will results in a sub-object range in the type layout.
-            // This allows us to simply scan the sub-object ranges to find out all specialization arguments.
-            Index subObjectRangeCount = subObjectRanges.getCount();
-            for (Index subObjectRangeIndex = 0; subObjectRangeIndex < subObjectRangeCount; subObjectRangeIndex++)
-            {
-                auto const& subObjectRange = subObjectRanges[subObjectRangeIndex];
-                auto const& bindingRange = getLayout()->getBindingRange(subObjectRange.bindingRangeIndex);
-
-                Index count = bindingRange.count;
-                SLANG_ASSERT(count == 1);
-
-                Index subObjectIndexInRange = 0;
-                auto subObject = m_objects[bindingRange.baseIndex + subObjectIndexInRange];
-
-                switch (bindingRange.bindingType)
-                {
-                case slang::BindingType::ExistentialValue:
-                {
-                    // A binding type of `ExistentialValue` means the sub-object represents a interface-typed field.
-                    // In this case the specialization argument for this field is the actual specialized type of the bound
-                    // shader object. If the shader object's type is an ordinary type without existential fields, then the
-                    // type argument will simply be the ordinary type. But if the sub object's type is itself a specialized
-                    // type, we need to make sure to use that type as the specialization argument.
-
-                    ExtendedShaderObjectType specializedSubObjType;
-                    SLANG_RETURN_ON_FAIL(subObject->getSpecializedShaderObjectType(&specializedSubObjType));
-                    args.add(specializedSubObjType);
-                    break;
-                }
-                case slang::BindingType::ParameterBlock:
-                case slang::BindingType::ConstantBuffer:
-                    // Currently we only handle the case where the field's type is
-                    // `ParameterBlock<SomeStruct>` or `ConstantBuffer<SomeStruct>`, where `SomeStruct` is a struct type
-                    // (not directly an interface type). In this case, we just recursively collect the specialization arguments
-                    // from the bound sub object.
-                    SLANG_RETURN_ON_FAIL(subObject->collectSpecializationArgs(args));
-                    // TODO: we need to handle the case where the field is of the form `ParameterBlock<IFoo>`. We should treat
-                    // this case the same way as the `ExistentialValue` case here, but currently we lack a mechanism to distinguish
-                    // the two scenarios.
-                    break;
-                }
-                // TODO: need to handle another case where specialization happens on resources fields e.g. `StructuredBuffer<IFoo>`.
-            }
-            return SLANG_OK;
-        }
 
 
     protected:
@@ -1430,8 +1273,8 @@ protected:
             size_t uniformSize = layout->getElementTypeLayout()->getSize();
             if (uniformSize)
             {
-                m_ordinaryData.setCount(uniformSize);
-                memset(m_ordinaryData.getBuffer(), 0, uniformSize);
+                m_data.setCount(uniformSize);
+                memset(m_data.getBuffer(), 0, uniformSize);
             }
 
             m_srvs.setCount(layout->getSRVCount());
@@ -1467,7 +1310,7 @@ protected:
                     RefPtr<ShaderObjectImpl> subObject;
                     SLANG_RETURN_ON_FAIL(
                         ShaderObjectImpl::create(device, subObjectLayout, subObject.writeRef()));
-                    m_objects[bindingRangeInfo.baseIndex + i] = subObject;
+                    m_objects[bindingRangeInfo.subObjectIndex + i] = subObject;
                 }
             }
 
@@ -1482,8 +1325,8 @@ protected:
         {
             // We start by simply writing in the ordinary data contained directly in this object.
             //
-            auto src = m_ordinaryData.getBuffer();
-            auto srcSize = size_t(m_ordinaryData.getCount());
+            auto src = m_data.getBuffer();
+            auto srcSize = size_t(m_data.getCount());
             SLANG_ASSERT(srcSize <= destSize);
             memcpy(dest, src, srcSize);
 
@@ -1554,7 +1397,7 @@ protected:
 
                 for (Slang::Index i = 0; i < count; ++i)
                 {
-                    auto subObject = m_objects[bindingRangeInfo.baseIndex + i];
+                    auto subObject = m_objects[bindingRangeInfo.subObjectIndex + i];
 
                     RefPtr<ShaderObjectLayoutImpl> subObjectLayout;
                     SLANG_RETURN_ON_FAIL(subObject->_getSpecializedLayout(subObjectLayout.writeRef()));
@@ -1758,7 +1601,7 @@ protected:
                 auto subObjectLayout = subObjectRange.layout;
                 auto const& bindingRange = specializedLayout->getBindingRange(subObjectRange.bindingRangeIndex);
                 Index count = bindingRange.count;
-                Index baseIndex = bindingRange.baseIndex;
+                Index subObjectIndex = bindingRange.subObjectIndex;
 
                 // The starting offset for a sub-object range was computed
                 // from Slang reflection information, so we can apply it here.
@@ -1782,7 +1625,7 @@ protected:
                         BindingOffset objOffset = rangeOffset;
                         for(Index i = 0; i < count; ++i)
                         {
-                            auto subObject = m_objects[ baseIndex + i ];
+                            auto subObject = m_objects[subObjectIndex + i];
 
                             // Unsurprisingly, we bind each object in the range as
                             // a constant buffer.
@@ -1810,7 +1653,7 @@ protected:
 
                         for(Index i = 0; i < count; ++i)
                         {
-                            auto subObject = m_objects[ baseIndex + i ];
+                            auto subObject = m_objects[subObjectIndex + i];
                             subObject->bindAsValue(context, BindingOffset(objOffset), subObjectLayout);
 
                             objOffset += objStride;
@@ -1830,10 +1673,6 @@ protected:
         // and organized as part of each shader object layout,
         // the object itself can store its data in a small number
         // of simple arrays.
-
-            /// Any "ordinary" / uniform data for this object
-        List<char> m_ordinaryData;
-
             /// The shader resource views (SRVs) that are part of the state of this object
         List<RefPtr<ShaderResourceViewImpl>> m_srvs;
 
@@ -1843,9 +1682,6 @@ protected:
             /// The samplers that are part of the state of this object
         List<RefPtr<SamplerStateImpl>> m_samplers;
 
-            /// The sub-objects that are part of the state of this object
-        List<RefPtr<ShaderObjectImpl>> m_objects;
-
             /// A constant buffer used to stored ordinary data for this object
             /// and existential-type sub-objects.
             ///
@@ -1879,7 +1715,9 @@ protected:
             auto renderer = getRenderer();
             RefPtr<ShaderObjectLayoutImpl> layout;
             SLANG_RETURN_ON_FAIL(renderer->getShaderObjectLayout(
-                extendedType.slangType, (ShaderObjectLayoutBase**)layout.writeRef()));
+                extendedType.slangType,
+                m_layout->getContainerType(),
+                (ShaderObjectLayoutBase**)layout.writeRef()));
 
             returnRefPtrMove(outLayout, layout);
             return SLANG_OK;
diff --git a/tools/gfx/d3d12/render-d3d12.cpp b/tools/gfx/d3d12/render-d3d12.cpp
index da169c03b..26b774b1a 100644
--- a/tools/gfx/d3d12/render-d3d12.cpp
+++ b/tools/gfx/d3d12/render-d3d12.cpp
@@ -232,17 +232,9 @@ public:
         }
     };
 
-    class ResourceViewImpl : public IResourceView, public ComObject
+    class ResourceViewImpl : public ResourceViewBase
     {
     public:
-        SLANG_COM_OBJECT_IUNKNOWN_ALL
-        IResourceView* getInterface(const Guid& guid)
-        {
-            if (guid == GfxGUID::IID_ISlangUnknown || guid == GfxGUID::IID_IResourceView)
-                return static_cast<IResourceView*>(this);
-            return nullptr;
-        }
-    public:
         RefPtr<Resource> m_resource;
         D3D12Descriptor m_descriptor;
         RefPtr<D3D12GeneralDescriptorHeap> m_allocator;
@@ -718,7 +710,11 @@ public:
             uint32_t count;
 
                 /// A "flat" index for this range in whatever array provides backing storage for it
-            uint32_t flatIndex;
+            uint32_t baseIndex;
+
+                /// An index into the sub-object array if this binding range is treated
+                /// as a sub-object.
+            uint32_t subObjectIndex;
         };
 
             /// Offset information for a sub-object range
@@ -802,12 +798,17 @@ public:
                 /// The number of root parameter consumed by (transitive) sub-objects
             uint32_t m_childRootParameterCount = 0;
 
-                /// The total size in bytes of the ordinary data for this object and transitive sub-objects
+                /// The total size in bytes of the ordinary data for this object and transitive sub-object.
             uint32_t m_totalOrdinaryDataSize = 0;
 
+                /// The container type of this shader object. When `m_containerType` is
+                /// `StructuredBuffer` or `UnsizedArray`, this shader object represents a collection
+                /// instead of a single object.
+            ShaderObjectContainerType m_containerType = ShaderObjectContainerType::None;
+
             Result setElementTypeLayout(slang::TypeLayoutReflection* typeLayout)
             {
-                typeLayout = _unwrapParameterGroups(typeLayout);
+                typeLayout = _unwrapParameterGroups(typeLayout, m_containerType);
                 m_elementTypeLayout = typeLayout;
 
                 // If the type contains any ordinary data, then we must reserve a buffer
@@ -849,12 +850,24 @@ public:
                     case slang::BindingType::ConstantBuffer:
                     case slang::BindingType::ParameterBlock:
                     case slang::BindingType::ExistentialValue:
-                        bindingRangeInfo.flatIndex = m_subObjectCount;
+                        bindingRangeInfo.baseIndex = m_subObjectCount;
+                        bindingRangeInfo.subObjectIndex = m_subObjectCount;
                         m_subObjectCount += count;
                         break;
-
+                    case slang::BindingType::RawBuffer:
+                    case slang::BindingType::MutableRawBuffer:
+                        if (slangLeafTypeLayout->getType()->getElementType() != nullptr)
+                        {
+                            // A structured buffer occupies both a resource slot and
+                            // a sub-object slot.
+                            bindingRangeInfo.subObjectIndex = m_subObjectCount;
+                            m_subObjectCount += count;
+                        }
+                        bindingRangeInfo.baseIndex = m_ownCounts.resource;
+                        m_ownCounts.resource += count;
+                        break;
                     case slang::BindingType::Sampler:
-                        bindingRangeInfo.flatIndex = m_ownCounts.sampler;
+                        bindingRangeInfo.baseIndex = m_ownCounts.sampler;
                         m_ownCounts.sampler += count;
                         break;
 
@@ -867,7 +880,7 @@ public:
                         break;
 
                     default:
-                        bindingRangeInfo.flatIndex = m_ownCounts.resource;
+                        bindingRangeInfo.baseIndex = m_ownCounts.resource;
                         m_ownCounts.resource += count;
                         break;
                     }
@@ -1092,8 +1105,6 @@ public:
 
         BindingRangeInfo const& getBindingRange(Index index) { return m_bindingRanges[index]; }
 
-        slang::TypeLayoutReflection* getElementTypeLayout() { return m_elementTypeLayout; }
-
         uint32_t getResourceSlotCount() { return m_ownCounts.resource; }
         uint32_t getSamplerSlotCount() { return m_ownCounts.sampler; }
         Index getSubObjectSlotCount() { return m_subObjectCount; }
@@ -1128,6 +1139,8 @@ public:
 
             initBase(renderer, builder->m_elementTypeLayout);
 
+            m_containerType = builder->m_containerType;
+
             m_bindingRanges = _Move(builder->m_bindingRanges);
             m_subObjectRanges = builder->m_subObjectRanges;
 
@@ -2033,7 +2046,11 @@ public:
         RefPtr<RootShaderObjectLayoutImpl> m_rootObjectLayout;
     };
 
-    class ShaderObjectImpl : public ShaderObjectBase
+    class ShaderObjectImpl
+        : public ShaderObjectBaseImpl<
+              ShaderObjectImpl,
+              ShaderObjectLayoutImpl,
+              SimpleShaderObjectData>
     {
     public:
         static Result create(
@@ -2064,25 +2081,14 @@ public:
             return SLANG_OK;
         }
 
-        ShaderObjectLayoutImpl* getLayout()
-        {
-            return static_cast<ShaderObjectLayoutImpl*>(m_layout.Ptr());
-        }
-
-        SLANG_NO_THROW slang::TypeLayoutReflection* SLANG_MCALL getElementTypeLayout()
-            SLANG_OVERRIDE
-        {
-            return m_layout->getElementTypeLayout();
-        }
-
         SLANG_NO_THROW Result SLANG_MCALL
             setData(ShaderOffset const& inOffset, void const* data, size_t inSize) SLANG_OVERRIDE
         {
             Index offset = inOffset.uniformOffset;
             Index size = inSize;
 
-            char* dest = m_ordinaryData.getBuffer();
-            Index availableSize = m_ordinaryData.getCount();
+            char* dest = m_data.getBuffer();
+            Index availableSize = m_data.getCount();
 
             // TODO: We really should bounds-check access rather than silently ignoring sets
             // that are too large, but we have several test cases that set more data than
@@ -2103,133 +2109,6 @@ public:
             return SLANG_OK;
         }
 
-        virtual SLANG_NO_THROW Result SLANG_MCALL
-            setObject(ShaderOffset const& offset, IShaderObject* object) SLANG_OVERRIDE
-        {
-            if (offset.bindingRangeIndex < 0)
-                return SLANG_E_INVALID_ARG;
-            auto layout = getLayout();
-            if (offset.bindingRangeIndex >= layout->getBindingRangeCount())
-                return SLANG_E_INVALID_ARG;
-
-            auto subObject = static_cast<ShaderObjectImpl*>(object);
-
-            auto bindingRangeIndex = offset.bindingRangeIndex;
-            auto& bindingRange = layout->getBindingRange(bindingRangeIndex);
-
-            m_objects[bindingRange.flatIndex + offset.bindingArrayIndex] = subObject;
-
-            // If the range being assigned into represents an interface/existential-type leaf field,
-            // then we need to consider how the `object` being assigned here affects specialization.
-            // We may also need to assign some data from the sub-object into the ordinary data
-            // buffer for the parent object.
-            //
-            if (bindingRange.bindingType == slang::BindingType::ExistentialValue)
-            {
-                // A leaf field of interface type is laid out inside of the parent object
-                // as a tuple of `(RTTI, WitnessTable, Payload)`. The layout of these fields
-                // is a contract between the compiler and any runtime system, so we will
-                // need to rely on details of the binary layout.
-
-                // We start by querying the layout/type of the concrete value that the application
-                // is trying to store into the field, and also the layout/type of the leaf
-                // existential-type field itself.
-                //
-                auto concreteTypeLayout = subObject->getElementTypeLayout();
-                auto concreteType = concreteTypeLayout->getType();
-                //
-                auto existentialTypeLayout =
-                    layout->getElementTypeLayout()->getBindingRangeLeafTypeLayout(
-                        bindingRangeIndex);
-                auto existentialType = existentialTypeLayout->getType();
-
-                // The first field of the tuple (offset zero) is the run-time type information
-                // (RTTI) ID for the concrete type being stored into the field.
-                //
-                // TODO: We need to be able to gather the RTTI type ID from `object` and then
-                // use `setData(offset, &TypeID, sizeof(TypeID))`.
-
-                // The second field of the tuple (offset 8) is the ID of the "witness" for the
-                // conformance of the concrete type to the interface used by this field.
-                //
-                auto witnessTableOffset = offset;
-                witnessTableOffset.uniformOffset += 8;
-                //
-                // Conformances of a type to an interface are computed and then stored by the
-                // Slang runtime, so we can look up the ID for this particular conformance (which
-                // will create it on demand).
-                //
-                ComPtr<slang::ISession> slangSession;
-                SLANG_RETURN_ON_FAIL(getRenderer()->getSlangSession(slangSession.writeRef()));
-                //
-                // Note: If the type doesn't actually conform to the required interface for
-                // this sub-object range, then this is the point where we will detect that
-                // fact and error out.
-                //
-                uint32_t conformanceID = 0xFFFFFFFF;
-                SLANG_RETURN_ON_FAIL(slangSession->getTypeConformanceWitnessSequentialID(
-                    concreteType, existentialType, &conformanceID));
-                //
-                // Once we have the conformance ID, then we can write it into the object
-                // at the required offset.
-                //
-                SLANG_RETURN_ON_FAIL(
-                    setData(witnessTableOffset, &conformanceID, sizeof(conformanceID)));
-
-                // The third field of the tuple (offset 16) is the "payload" that is supposed to
-                // hold the data for a value of the given concrete type.
-                //
-                auto payloadOffset = offset;
-                payloadOffset.uniformOffset += 16;
-
-                // There are two cases we need to consider here for how the payload might be used:
-                //
-                // * If the concrete type of the value being bound is one that can "fit" into the
-                //   available payload space,  then it should be stored in the payload.
-                //
-                // * If the concrete type of the value cannot fit in the payload space, then it
-                //   will need to be stored somewhere else.
-                //
-                if (_doesValueFitInExistentialPayload(concreteTypeLayout, existentialTypeLayout))
-                {
-                    // If the value can fit in the payload area, then we will go ahead and copy
-                    // its bytes into that area.
-                    //
-                    setData(
-                        payloadOffset,
-                        subObject->m_ordinaryData.getBuffer(),
-                        subObject->m_ordinaryData.getCount());
-                }
-                else
-                {
-                    // If the value does *not *fit in the payload area, then there is nothing
-                    // we can do at this point (beyond saving a reference to the sub-object, which
-                    // was handled above).
-                    //
-                    // Once all the sub-objects have been set into the parent object, we can
-                    // compute a specialized layout for it, and that specialized layout can tell
-                    // us where the data for these sub-objects has been laid out.
-                    return SLANG_E_NOT_IMPLEMENTED;
-                }
-            }
-            return SLANG_OK;
-        }
-
-        virtual SLANG_NO_THROW Result SLANG_MCALL
-            getObject(ShaderOffset const& offset, IShaderObject** outObject) SLANG_OVERRIDE
-        {
-            SLANG_ASSERT(outObject);
-            if (offset.bindingRangeIndex < 0)
-                return SLANG_E_INVALID_ARG;
-            auto layout = getLayout();
-            if (offset.bindingRangeIndex >= layout->getBindingRangeCount())
-                return SLANG_E_INVALID_ARG;
-            auto& bindingRange = layout->getBindingRange(offset.bindingRangeIndex);
-
-            returnComPtr(outObject, m_objects[bindingRange.flatIndex + offset.bindingArrayIndex]);
-            return SLANG_OK;
-        }
-
         SLANG_NO_THROW Result SLANG_MCALL
             setResource(ShaderOffset const& offset, IResourceView* resourceView) SLANG_OVERRIDE
         {
@@ -2242,15 +2121,15 @@ public:
             auto resourceViewImpl = static_cast<ResourceViewImpl*>(resourceView);
 
             auto& bindingRange = layout->getBindingRange(offset.bindingRangeIndex);
-            auto descriptorSlotIndex = bindingRange.flatIndex + (int32_t)offset.bindingArrayIndex;
+            auto descriptorSlotIndex = bindingRange.baseIndex + (int32_t)offset.bindingArrayIndex;
             // Hold a reference to the resource to prevent its destruction.
-            m_boundResources[bindingRange.flatIndex + offset.bindingArrayIndex] =
+            m_boundResources[bindingRange.baseIndex + offset.bindingArrayIndex] =
                 resourceViewImpl->m_resource;
             ID3D12Device* d3dDevice = static_cast<D3D12Device*>(getDevice())->m_device;
             d3dDevice->CopyDescriptorsSimple(
                 1,
                 m_descriptorSet.resourceTable.getCpuHandle(
-                    bindingRange.flatIndex + (int32_t)offset.bindingArrayIndex),
+                    bindingRange.baseIndex + (int32_t)offset.bindingArrayIndex),
                 resourceViewImpl->m_descriptor.cpuHandle,
                 D3D12_DESCRIPTOR_HEAP_TYPE_CBV_SRV_UAV);
             return SLANG_OK;
@@ -2270,7 +2149,7 @@ public:
             d3dDevice->CopyDescriptorsSimple(
                 1,
                 m_descriptorSet.samplerTable.getCpuHandle(
-                    bindingRange.flatIndex +
+                    bindingRange.baseIndex +
                     (int32_t)offset.bindingArrayIndex),
                 samplerImpl->m_descriptor.cpuHandle,
                 D3D12_DESCRIPTOR_HEAP_TYPE_SAMPLER);
@@ -2313,66 +2192,6 @@ public:
         }
 
     public:
-        // Appends all types that are used to specialize the element type of this shader object in
-        // `args` list.
-        virtual Result collectSpecializationArgs(ExtendedShaderObjectTypeList& args) override
-        {
-            auto& subObjectRanges = getLayout()->getSubObjectRanges();
-            // The following logic is built on the assumption that all fields that involve
-            // existential types (and therefore require specialization) will results in a sub-object
-            // range in the type layout. This allows us to simply scan the sub-object ranges to find
-            // out all specialization arguments.
-            Index subObjectRangeCount = subObjectRanges.getCount();
-            for (Index subObjectRangeIndex = 0; subObjectRangeIndex < subObjectRangeCount;
-                 subObjectRangeIndex++)
-            {
-                auto const& subObjectRange = subObjectRanges[subObjectRangeIndex];
-                auto const& bindingRange =
-                    getLayout()->getBindingRange(subObjectRange.bindingRangeIndex);
-
-                Index count = bindingRange.count;
-                SLANG_ASSERT(count == 1);
-
-                Index subObjectIndexInRange = 0;
-                auto subObject = m_objects[bindingRange.flatIndex + subObjectIndexInRange];
-
-                switch (bindingRange.bindingType)
-                {
-                case slang::BindingType::ExistentialValue:
-                    {
-                        // A binding type of `ExistentialValue` means the sub-object represents a
-                        // interface-typed field. In this case the specialization argument for this
-                        // field is the actual specialized type of the bound shader object. If the
-                        // shader object's type is an ordinary type without existential fields, then
-                        // the type argument will simply be the ordinary type. But if the sub
-                        // object's type is itself a specialized type, we need to make sure to use
-                        // that type as the specialization argument.
-
-                        ExtendedShaderObjectType specializedSubObjType;
-                        SLANG_RETURN_ON_FAIL(
-                            subObject->getSpecializedShaderObjectType(&specializedSubObjType));
-                        args.add(specializedSubObjType);
-                        break;
-                    }
-                case slang::BindingType::ParameterBlock:
-                case slang::BindingType::ConstantBuffer:
-                    // Currently we only handle the case where the field's type is
-                    // `ParameterBlock<SomeStruct>` or `ConstantBuffer<SomeStruct>`, where
-                    // `SomeStruct` is a struct type (not directly an interface type). In this case,
-                    // we just recursively collect the specialization arguments from the bound sub
-                    // object.
-                    SLANG_RETURN_ON_FAIL(subObject->collectSpecializationArgs(args));
-                    // TODO: we need to handle the case where the field is of the form
-                    // `ParameterBlock<IFoo>`. We should treat this case the same way as the
-                    // `ExistentialValue` case here, but currently we lack a mechanism to
-                    // distinguish the two scenarios.
-                    break;
-                }
-                // TODO: need to handle another case where specialization happens on resources
-                // fields e.g. `StructuredBuffer<IFoo>`.
-            }
-            return SLANG_OK;
-        }
 
     protected:
         Result init(
@@ -2399,8 +2218,8 @@ public:
             size_t uniformSize = layout->getElementTypeLayout()->getSize();
             if (uniformSize)
             {
-                m_ordinaryData.setCount(uniformSize);
-                memset(m_ordinaryData.getBuffer(), 0, uniformSize);
+                m_data.setCount(uniformSize);
+                memset(m_data.getBuffer(), 0, uniformSize);
             }
 
             // Each shader object will own CPU descriptor heap memory
@@ -2458,7 +2277,7 @@ public:
                     RefPtr<ShaderObjectImpl> subObject;
                     SLANG_RETURN_ON_FAIL(
                         ShaderObjectImpl::create(device, subObjectLayout, subObject.writeRef()));
-                    m_objects[bindingRangeInfo.flatIndex + i] = subObject;
+                    m_objects[bindingRangeInfo.subObjectIndex + i] = subObject;
                 }
             }
 
@@ -2474,8 +2293,8 @@ public:
             size_t destSize,
             ShaderObjectLayoutImpl* specializedLayout)
         {
-            auto src = m_ordinaryData.getBuffer();
-            auto srcSize = size_t(m_ordinaryData.getCount());
+            auto src = m_data.getBuffer();
+            auto srcSize = size_t(m_data.getCount());
 
             SLANG_ASSERT(srcSize <= destSize);
 
@@ -2550,7 +2369,7 @@ public:
 
                 for (uint32_t i = 0; i < count; ++i)
                 {
-                    auto subObject = m_objects[bindingRangeInfo.flatIndex + i];
+                    auto subObject = m_objects[bindingRangeInfo.subObjectIndex + i];
 
                     RefPtr<ShaderObjectLayoutImpl> subObjectLayout;
                     SLANG_RETURN_ON_FAIL(
@@ -2863,7 +2682,7 @@ public:
             {
                 auto& subObjectRange = specializedLayout->getSubObjectRange(i);
                 auto& bindingRange = specializedLayout->getBindingRange(subObjectRange.bindingRangeIndex);
-                auto baseIndex = bindingRange.flatIndex;
+                auto subObjectIndex = bindingRange.subObjectIndex;
                 auto subObjectLayout = subObjectRange.layout.Ptr();
 
                 BindingOffset rangeOffset = offset;
@@ -2878,7 +2697,7 @@ public:
                         auto objOffset = rangeOffset;
                         for (uint32_t j = 0; j < bindingRange.count; j++)
                         {
-                            auto& object = m_objects[baseIndex + j];
+                            auto& object = m_objects[subObjectIndex + j];
                             object->bindAsConstantBuffer(context, descriptorSet, objOffset, subObjectLayout);
                             objOffset += rangeStride;
                         }
@@ -2890,7 +2709,7 @@ public:
                         auto objOffset = rangeOffset;
                         for (uint32_t j = 0; j < bindingRange.count; j++)
                         {
-                            auto& object = m_objects[baseIndex + j];
+                            auto& object = m_objects[subObjectIndex + j];
                             object->bindAsParameterBlock(context, objOffset, subObjectLayout);
                             objOffset += rangeStride;
                         }
@@ -2903,7 +2722,7 @@ public:
                         auto objOffset = rangeOffset;
                         for (uint32_t j = 0; j < bindingRange.count; j++)
                         {
-                            auto& object = m_objects[baseIndex + j];
+                            auto& object = m_objects[subObjectIndex + j];
                             object->bindAsValue(context, descriptorSet, objOffset, subObjectLayout);
                             objOffset += rangeStride;
                         }
@@ -2915,10 +2734,6 @@ public:
             return SLANG_OK;
         }
 
-        /// Any "ordinary" / uniform data for this object
-        List<char> m_ordinaryData;
-
-        List<RefPtr<ShaderObjectImpl>> m_objects;
 
             /// A CPU-memory descriptor set holding any descriptors used to represent the resources/samplers in this object's state
         DescriptorSet m_descriptorSet;
@@ -2967,7 +2782,9 @@ public:
             auto renderer = getRenderer();
             RefPtr<ShaderObjectLayoutImpl> layout;
             SLANG_RETURN_ON_FAIL(renderer->getShaderObjectLayout(
-                extendedType.slangType, (ShaderObjectLayoutBase**)layout.writeRef()));
+                extendedType.slangType,
+                m_layout->getContainerType(),
+                (ShaderObjectLayoutBase**)layout.writeRef()));
 
             returnRefPtrMove(outLayout, layout);
             return SLANG_OK;
diff --git a/tools/gfx/debug-layer.cpp b/tools/gfx/debug-layer.cpp
index aa8989623..56ae4fdab 100644
--- a/tools/gfx/debug-layer.cpp
+++ b/tools/gfx/debug-layer.cpp
@@ -373,12 +373,18 @@ Result DebugDevice::createCommandQueue(const ICommandQueue::Desc& desc, ICommand
     return result;
 }
 
-Result DebugDevice::createShaderObject(slang::TypeReflection* type, IShaderObject** outShaderObject)
+Result DebugDevice::createShaderObject(
+    slang::TypeReflection* type,
+    ShaderObjectContainerType containerType,
+    IShaderObject** outShaderObject)
 {
     SLANG_GFX_API_FUNC;
 
     RefPtr<DebugShaderObject> outObject = new DebugShaderObject();
-    auto result = baseObject->createShaderObject(type, outObject->baseObject.writeRef());
+    auto typeName = type->getName();
+    auto result =
+        baseObject->createShaderObject(type, containerType, outObject->baseObject.writeRef());
+    outObject->m_typeName = typeName;
     if (SLANG_FAILED(result))
         return result;
     returnComPtr(outShaderObject, outObject);
@@ -838,6 +844,12 @@ void DebugSwapchain::maybeRebuildImageList()
     }
 }
 
+ShaderObjectContainerType DebugShaderObject::getContainerType()
+{
+    SLANG_GFX_API_FUNC;
+    return baseObject->getContainerType();
+}
+
 slang::TypeLayoutReflection* DebugShaderObject::getElementTypeLayout()
 {
     SLANG_GFX_API_FUNC;
@@ -896,6 +908,7 @@ Result DebugShaderObject::getObject(ShaderOffset const& offset, IShaderObject**
     }
     debugShaderObject = new DebugShaderObject();
     debugShaderObject->baseObject = innerObject;
+    debugShaderObject->m_typeName = innerObject->getElementTypeLayout()->getName();
     m_objects[ShaderOffsetKey{offset}] = debugShaderObject;
     returnComPtr(object, debugShaderObject);
     return resultCode;
diff --git a/tools/gfx/debug-layer.h b/tools/gfx/debug-layer.h
index 12540c31a..a4e201e4f 100644
--- a/tools/gfx/debug-layer.h
+++ b/tools/gfx/debug-layer.h
@@ -75,8 +75,10 @@ public:
         IInputLayout** outLayout) override;
     virtual SLANG_NO_THROW Result SLANG_MCALL
         createCommandQueue(const ICommandQueue::Desc& desc, ICommandQueue** outQueue) override;
-    virtual SLANG_NO_THROW Result SLANG_MCALL
-        createShaderObject(slang::TypeReflection* type, IShaderObject** outObject) override;
+    virtual SLANG_NO_THROW Result SLANG_MCALL createShaderObject(
+        slang::TypeReflection* type,
+        ShaderObjectContainerType container,
+        IShaderObject** outObject) override;
     virtual SLANG_NO_THROW Result SLANG_MCALL
         createProgram(const IShaderProgram::Desc& desc, IShaderProgram** outProgram) override;
     virtual SLANG_NO_THROW Result SLANG_MCALL createGraphicsPipelineState(
@@ -147,6 +149,7 @@ public:
 public:
     IShaderObject* getInterface(const Slang::Guid& guid);
     virtual SLANG_NO_THROW slang::TypeLayoutReflection* SLANG_MCALL getElementTypeLayout() override;
+    virtual SLANG_NO_THROW ShaderObjectContainerType SLANG_MCALL getContainerType() override;
     virtual SLANG_NO_THROW UInt SLANG_MCALL getEntryPointCount() override;
     virtual SLANG_NO_THROW Result SLANG_MCALL
         getEntryPoint(UInt index, IShaderObject** entryPoint) override;
@@ -184,6 +187,7 @@ public:
                     (Slang::HashCode)offset.bindingRangeIndex));
         }
     };
+    Slang::String m_typeName;
     Slang::List<Slang::RefPtr<DebugShaderObject>> m_entryPoints;
     Slang::Dictionary<ShaderOffsetKey, Slang::RefPtr<DebugShaderObject>> m_objects;
     Slang::Dictionary<ShaderOffsetKey, Slang::RefPtr<DebugResourceView>> m_resources;
diff --git a/tools/gfx/open-gl/render-gl.cpp b/tools/gfx/open-gl/render-gl.cpp
index 53e9dd4b1..5b1ea2c3e 100644
--- a/tools/gfx/open-gl/render-gl.cpp
+++ b/tools/gfx/open-gl/render-gl.cpp
@@ -272,17 +272,9 @@ public:
         GLuint m_samplerID;
     };
 
-    class ResourceViewImpl : public IResourceView, public ComObject
+    class ResourceViewImpl : public ResourceViewBase
     {
     public:
-        SLANG_COM_OBJECT_IUNKNOWN_ALL
-        IResourceView* getInterface(const Guid& guid)
-        {
-            if (guid == GfxGUID::IID_ISlangUnknown || guid == GfxGUID::IID_IResourceView)
-                return static_cast<IResourceView*>(this);
-            return nullptr;
-        }
-    public:
         enum class Type
         {
             Texture, Buffer
@@ -620,6 +612,7 @@ public:
             slang::BindingType bindingType;
             Index count;
             Index baseIndex;
+            Index subObjectIndex;
         };
 
         struct SubObjectRangeInfo
@@ -638,6 +631,11 @@ public:
             RendererBase* m_renderer;
             slang::TypeLayoutReflection* m_elementTypeLayout;
 
+            /// The container type of this shader object. When `m_containerType` is
+            /// `StructuredBuffer` or `UnsizedArray`, this shader object represents a collection
+            /// instead of a single object.
+            ShaderObjectContainerType m_containerType = ShaderObjectContainerType::None;
+
             List<BindingRangeInfo> m_bindingRanges;
             List<SubObjectRangeInfo> m_subObjectRanges;
 
@@ -648,7 +646,7 @@ public:
 
             Result setElementTypeLayout(slang::TypeLayoutReflection* typeLayout)
             {
-                typeLayout = _unwrapParameterGroups(typeLayout);
+                typeLayout = _unwrapParameterGroups(typeLayout, m_containerType);
 
                 m_elementTypeLayout = typeLayout;
 
@@ -673,9 +671,21 @@ public:
                     case slang::BindingType::ParameterBlock:
                     case slang::BindingType::ExistentialValue:
                         bindingRangeInfo.baseIndex = m_subObjectCount;
+                        bindingRangeInfo.subObjectIndex = m_subObjectCount;
                         m_subObjectCount += count;
                         break;
-
+                    case slang::BindingType::RawBuffer:
+                    case slang::BindingType::MutableRawBuffer:
+                        if (slangLeafTypeLayout->getType()->getElementType() != nullptr)
+                        {
+                            // A structured buffer occupies both a resource slot and
+                            // a sub-object slot.
+                            bindingRangeInfo.subObjectIndex = m_subObjectCount;
+                            m_subObjectCount += count;
+                        }
+                        bindingRangeInfo.baseIndex = m_storageBufferCount;
+                        m_storageBufferCount += count;
+                        break;
                     case slang::BindingType::Sampler:
                         break;
 
@@ -690,7 +700,6 @@ public:
                         m_imageCount += count;
                         break;
 
-                    case slang::BindingType::MutableRawBuffer:
                     case slang::BindingType::MutableTypedBuffer:
                         bindingRangeInfo.baseIndex = m_storageBufferCount;
                         m_storageBufferCount += count;
@@ -765,8 +774,6 @@ public:
 
         BindingRangeInfo const& getBindingRange(Index index) { return m_bindingRanges[index]; }
 
-        slang::TypeLayoutReflection* getElementTypeLayout() { return m_elementTypeLayout; }
-
         Index getTextureCount() { return m_textureCount; }
         Index getImageCount() { return m_imageCount; }
         Index getStorageBufferCount() { return m_storageBufferCount; }
@@ -795,6 +802,8 @@ public:
             m_storageBufferCount = builder->m_storageBufferCount;
             m_subObjectCount = builder->m_subObjectCount;
             m_subObjectRanges = builder->m_subObjectRanges;
+
+            m_containerType = builder->m_containerType;
             return SLANG_OK;
         }
 
@@ -903,7 +912,11 @@ public:
         List<EntryPointInfo> m_entryPoints;
     };
 
-    class ShaderObjectImpl : public ShaderObjectBase
+    class ShaderObjectImpl
+        : public ShaderObjectBaseImpl<
+              ShaderObjectImpl,
+              ShaderObjectLayoutImpl,
+              SimpleShaderObjectData>
     {
     public:
         static Result create(
@@ -934,19 +947,14 @@ public:
             return static_cast<ShaderObjectLayoutImpl*>(m_layout.Ptr());
         }
 
-        SLANG_NO_THROW slang::TypeLayoutReflection* SLANG_MCALL getElementTypeLayout() SLANG_OVERRIDE
-        {
-            return m_layout->getElementTypeLayout();
-        }
-
         SLANG_NO_THROW Result SLANG_MCALL
             setData(ShaderOffset const& inOffset, void const* data, size_t inSize) SLANG_OVERRIDE
         {
             Index offset = inOffset.uniformOffset;
             Index size = inSize;
 
-            char* dest = m_ordinaryData.getBuffer();
-            Index availableSize = m_ordinaryData.getCount();
+            char* dest = m_data.getBuffer();
+            Index availableSize = m_data.getCount();
 
             // TODO: We really should bounds-check access rather than silently ignoring sets
             // that are too large, but we have several test cases that set more data than
@@ -967,130 +975,6 @@ public:
             return SLANG_OK;
         }
 
-        virtual SLANG_NO_THROW Result SLANG_MCALL
-            setObject(ShaderOffset const& offset, IShaderObject* object)
-            SLANG_OVERRIDE
-        {
-            if (offset.bindingRangeIndex < 0)
-                return SLANG_E_INVALID_ARG;
-            auto layout = getLayout();
-            if (offset.bindingRangeIndex >= layout->getBindingRangeCount())
-                return SLANG_E_INVALID_ARG;
-
-            auto subObject = static_cast<ShaderObjectImpl*>(object);
-
-            auto bindingRangeIndex = offset.bindingRangeIndex;
-            auto& bindingRange = layout->getBindingRange(bindingRangeIndex);
-
-            m_objects[bindingRange.baseIndex + offset.bindingArrayIndex] = subObject;
-
-            // If the range being assigned into represents an interface/existential-type leaf field,
-            // then we need to consider how the `object` being assigned here affects specialization.
-            // We may also need to assign some data from the sub-object into the ordinary data
-            // buffer for the parent object.
-            //
-            if (bindingRange.bindingType == slang::BindingType::ExistentialValue)
-            {
-                // A leaf field of interface type is laid out inside of the parent object
-                // as a tuple of `(RTTI, WitnessTable, Payload)`. The layout of these fields
-                // is a contract between the compiler and any runtime system, so we will
-                // need to rely on details of the binary layout.
-
-                // We start by querying the layout/type of the concrete value that the application
-                // is trying to store into the field, and also the layout/type of the leaf
-                // existential-type field itself.
-                //
-                auto concreteTypeLayout = subObject->getElementTypeLayout();
-                auto concreteType = concreteTypeLayout->getType();
-                //
-                auto existentialTypeLayout = layout->getElementTypeLayout()->getBindingRangeLeafTypeLayout(bindingRangeIndex);
-                auto existentialType = existentialTypeLayout->getType();
-
-                // The first field of the tuple (offset zero) is the run-time type information (RTTI)
-                // ID for the concrete type being stored into the field.
-                //
-                // TODO: We need to be able to gather the RTTI type ID from `object` and then
-                // use `setData(offset, &TypeID, sizeof(TypeID))`.
-
-                // The second field of the tuple (offset 8) is the ID of the "witness" for the
-                // conformance of the concrete type to the interface used by this field.
-                //
-                auto witnessTableOffset = offset;
-                witnessTableOffset.uniformOffset += 8;
-                //
-                // Conformances of a type to an interface are computed and then stored by the
-                // Slang runtime, so we can look up the ID for this particular conformance (which
-                // will create it on demand).
-                //
-                ComPtr<slang::ISession> slangSession;
-                SLANG_RETURN_ON_FAIL(getRenderer()->getSlangSession(slangSession.writeRef()));
-                //
-                // Note: If the type doesn't actually conform to the required interface for
-                // this sub-object range, then this is the point where we will detect that
-                // fact and error out.
-                //
-                uint32_t conformanceID = 0xFFFFFFFF;
-                SLANG_RETURN_ON_FAIL(slangSession->getTypeConformanceWitnessSequentialID(
-                    concreteType, existentialType, &conformanceID));
-                //
-                // Once we have the conformance ID, then we can write it into the object
-                // at the required offset.
-                //
-                SLANG_RETURN_ON_FAIL(setData(witnessTableOffset, &conformanceID, sizeof(conformanceID)));
-
-                // The third field of the tuple (offset 16) is the "payload" that is supposed to
-                // hold the data for a value of the given concrete type.
-                //
-                auto payloadOffset = offset;
-                payloadOffset.uniformOffset += 16;
-
-                // There are two cases we need to consider here for how the payload might be used:
-                //
-                // * If the concrete type of the value being bound is one that can "fit" into the
-                //   available payload space,  then it should be stored in the payload.
-                //
-                // * If the concrete type of the value cannot fit in the payload space, then it
-                //   will need to be stored somewhere else.
-                //
-                if (_doesValueFitInExistentialPayload(concreteTypeLayout, existentialTypeLayout))
-                {
-                    // If the value can fit in the payload area, then we will go ahead and copy
-                    // its bytes into that area.
-                    //
-                    setData(payloadOffset, subObject->m_ordinaryData.getBuffer(), subObject->m_ordinaryData.getCount());
-                }
-                else
-                {
-                    // If the value does *not *fit in the payload area, then there is nothing
-                    // we can do at this point (beyond saving a reference to the sub-object, which
-                    // was handled above).
-                    //
-                    // Once all the sub-objects have been set into the parent object, we can
-                    // compute a specialized layout for it, and that specialized layout can tell
-                    // us where the data for these sub-objects has been laid out.
-                }
-            }
-
-            return SLANG_E_NOT_IMPLEMENTED;
-        }
-
-        virtual SLANG_NO_THROW Result SLANG_MCALL
-            getObject(ShaderOffset const& offset, IShaderObject** outObject)
-            SLANG_OVERRIDE
-        {
-            SLANG_ASSERT(outObject);
-            if (offset.bindingRangeIndex < 0)
-                return SLANG_E_INVALID_ARG;
-            auto layout = getLayout();
-            if (offset.bindingRangeIndex >= layout->getBindingRangeCount())
-                return SLANG_E_INVALID_ARG;
-            auto& bindingRange = layout->getBindingRange(offset.bindingRangeIndex);
-
-            auto object = m_objects[bindingRange.baseIndex + offset.bindingArrayIndex].Ptr();
-            object->addRef();
-            *outObject = object;
-            return SLANG_OK;
-        }
 
         SLANG_NO_THROW Result SLANG_MCALL
             setResource(ShaderOffset const& offset, IResourceView* resourceView) SLANG_OVERRIDE
@@ -1151,57 +1035,6 @@ public:
         }
 
     public:
-        // Appends all types that are used to specialize the element type of this shader object in `args` list.
-        virtual Result collectSpecializationArgs(ExtendedShaderObjectTypeList& args) override
-        {
-            auto& subObjectRanges = getLayout()->getSubObjectRanges();
-            // The following logic is built on the assumption that all fields that involve existential types (and
-            // therefore require specialization) will results in a sub-object range in the type layout.
-            // This allows us to simply scan the sub-object ranges to find out all specialization arguments.
-            Index subObjectRangeCount = subObjectRanges.getCount();
-            for (Index subObjectRangeIndex = 0; subObjectRangeIndex < subObjectRangeCount; subObjectRangeIndex++)
-            {
-                auto const& subObjectRange = subObjectRanges[subObjectRangeIndex];
-                auto const& bindingRange = getLayout()->getBindingRange(subObjectRange.bindingRangeIndex);
-
-                Index count = bindingRange.count;
-                SLANG_ASSERT(count == 1);
-
-                Index subObjectIndexInRange = 0;
-                auto subObject = m_objects[bindingRange.baseIndex + subObjectIndexInRange];
-
-                switch (bindingRange.bindingType)
-                {
-                case slang::BindingType::ExistentialValue:
-                {
-                    // A binding type of `ExistentialValue` means the sub-object represents a interface-typed field.
-                    // In this case the specialization argument for this field is the actual specialized type of the bound
-                    // shader object. If the shader object's type is an ordinary type without existential fields, then the
-                    // type argument will simply be the ordinary type. But if the sub object's type is itself a specialized
-                    // type, we need to make sure to use that type as the specialization argument.
-
-                    ExtendedShaderObjectType specializedSubObjType;
-                    SLANG_RETURN_ON_FAIL(subObject->getSpecializedShaderObjectType(&specializedSubObjType));
-                    args.add(specializedSubObjType);
-                    break;
-                }
-                case slang::BindingType::ParameterBlock:
-                case slang::BindingType::ConstantBuffer:
-                    // Currently we only handle the case where the field's type is
-                    // `ParameterBlock<SomeStruct>` or `ConstantBuffer<SomeStruct>`, where `SomeStruct` is a struct type
-                    // (not directly an interface type). In this case, we just recursively collect the specialization arguments
-                    // from the bound sub object.
-                    SLANG_RETURN_ON_FAIL(subObject->collectSpecializationArgs(args));
-                    // TODO: we need to handle the case where the field is of the form `ParameterBlock<IFoo>`. We should treat
-                    // this case the same way as the `ExistentialValue` case here, but currently we lack a mechanism to distinguish
-                    // the two scenarios.
-                    break;
-                }
-                // TODO: need to handle another case where specialization happens on resources fields e.g. `StructuredBuffer<IFoo>`.
-            }
-            return SLANG_OK;
-        }
-
 
     protected:
         friend class ProgramVars;
@@ -1222,8 +1055,8 @@ public:
             size_t uniformSize = layout->getElementTypeLayout()->getSize();
             if (uniformSize)
             {
-                m_ordinaryData.setCount(uniformSize);
-                memset(m_ordinaryData.getBuffer(), 0, uniformSize);
+                m_data.setCount(uniformSize);
+                memset(m_data.getBuffer(), 0, uniformSize);
             }
 
             m_samplers.setCount(layout->getTextureCount());
@@ -1260,7 +1093,7 @@ public:
                     RefPtr<ShaderObjectImpl> subObject;
                     SLANG_RETURN_ON_FAIL(
                         ShaderObjectImpl::create(device, subObjectLayout, subObject.writeRef()));
-                    m_objects[bindingRangeInfo.baseIndex + i] = subObject;
+                    m_objects[bindingRangeInfo.subObjectIndex + i] = subObject;
                 }
             }
 
@@ -1275,8 +1108,8 @@ public:
             size_t                              destSize,
             ShaderObjectLayoutImpl* specializedLayout)
         {
-            auto src = m_ordinaryData.getBuffer();
-            auto srcSize = size_t(m_ordinaryData.getCount());
+            auto src = m_data.getBuffer();
+            auto srcSize = size_t(m_data.getCount());
 
             SLANG_ASSERT(srcSize <= destSize);
 
@@ -1349,7 +1182,7 @@ public:
 
                 for (Slang::Index i = 0; i < count; ++i)
                 {
-                    auto subObject = m_objects[bindingRangeInfo.baseIndex + i];
+                    auto subObject = m_objects[bindingRangeInfo.subObjectIndex + i];
 
                     RefPtr<ShaderObjectLayoutImpl> subObjectLayout;
                     SLANG_RETURN_ON_FAIL(subObject->_getSpecializedLayout(subObjectLayout.writeRef()));
@@ -1459,15 +1292,31 @@ public:
             for (auto buffer : m_storageBuffers)
                 bindingState->storageBufferBindings.add(buffer ? buffer->m_bufferID : 0);
 
-            for (auto subObject : m_objects)
-                subObject->bindObject(device, bindingState);
+            for (auto const& subObjectRange : layout->getSubObjectRanges())
+            {
+                auto subObjectLayout = subObjectRange.layout;
+                auto const& bindingRange =
+                    layout->getBindingRange(subObjectRange.bindingRangeIndex);
+                
+                switch (bindingRange.bindingType)
+                {
+                case slang::BindingType::ConstantBuffer:
+                case slang::BindingType::ParameterBlock:
+                case slang::BindingType::ExistentialValue:
+                    break;
+                default:
+                    continue;
+                }
+
+                for (Index i = 0; i < bindingRange.count; i++)
+                {
+                    m_objects[i + bindingRange.subObjectIndex]->bindObject(device, bindingState);
+                }
+            }
 
             return SLANG_OK;
         }
 
-        /// Any "ordinary" / uniform data for this object
-        List<char> m_ordinaryData;
-
         List<RefPtr<TextureViewImpl>> m_textures;
 
         List<RefPtr<TextureViewImpl>> m_images;
@@ -1476,8 +1325,6 @@ public:
 
         List<RefPtr<BufferViewImpl>> m_storageBuffers;
 
-        List<RefPtr<ShaderObjectImpl>> m_objects;
-
         /// A constant buffer used to stored ordinary data for this object
         /// and existential-type sub-objects.
         ///
@@ -1511,7 +1358,9 @@ public:
             auto renderer = getRenderer();
             RefPtr<ShaderObjectLayoutImpl> layout;
             SLANG_RETURN_ON_FAIL(renderer->getShaderObjectLayout(
-                extendedType.slangType, (ShaderObjectLayoutBase**)layout.writeRef()));
+                extendedType.slangType,
+                m_layout->getContainerType(),
+                (ShaderObjectLayoutBase**)layout.writeRef()));
 
             returnRefPtrMove(outLayout, layout);
             return SLANG_OK;
diff --git a/tools/gfx/renderer-shared.cpp b/tools/gfx/renderer-shared.cpp
index 72b4c45f5..1b384f8eb 100644
--- a/tools/gfx/renderer-shared.cpp
+++ b/tools/gfx/renderer-shared.cpp
@@ -32,7 +32,7 @@ const Slang::Guid GfxGUID::IID_IResourceCommandEncoder = SLANG_UUID_IResourceCom
 const Slang::Guid GfxGUID::IID_ICommandBuffer = SLANG_UUID_ICommandBuffer;
 const Slang::Guid GfxGUID::IID_ICommandQueue = SLANG_UUID_ICommandQueue;
 
-gfx::StageType translateStage(SlangStage slangStage)
+StageType translateStage(SlangStage slangStage)
 {
     switch (slangStage)
     {
@@ -86,12 +86,12 @@ IResource* TextureResource::getInterface(const Slang::Guid& guid)
 SLANG_NO_THROW IResource::Type SLANG_MCALL TextureResource::getType() { return m_type; }
 SLANG_NO_THROW ITextureResource::Desc* SLANG_MCALL TextureResource::getDesc() { return &m_desc; }
 
-gfx::StageType mapStage(SlangStage stage)
+StageType mapStage(SlangStage stage)
 {
     switch( stage )
     {
     default:
-        return gfx::StageType::Unknown;
+        return StageType::Unknown;
 
     case SLANG_STAGE_AMPLIFICATION:     return gfx::StageType::Amplification;
     case SLANG_STAGE_ANY_HIT:           return gfx::StageType::AnyHit;
@@ -110,14 +110,21 @@ gfx::StageType mapStage(SlangStage stage)
     }
 }
 
-IShaderObject* gfx::ShaderObjectBase::getInterface(const Guid& guid)
+IResourceView* ResourceViewBase::getInterface(const Guid& guid)
+{
+    if (guid == GfxGUID::IID_ISlangUnknown || guid == GfxGUID::IID_IResourceView)
+        return static_cast<IResourceView*>(this);
+    return nullptr;
+}
+
+IShaderObject* ShaderObjectBase::getInterface(const Guid& guid)
 {
     if (guid == GfxGUID::IID_ISlangUnknown || guid == GfxGUID::IID_IShaderObject)
         return static_cast<IShaderObject*>(this);
     return nullptr;
 }
 
-bool gfx::ShaderObjectBase::_doesValueFitInExistentialPayload(
+bool ShaderObjectBase::_doesValueFitInExistentialPayload(
     slang::TypeLayoutReflection*    concreteTypeLayout,
     slang::TypeLayoutReflection*    existentialTypeLayout)
 {
@@ -176,21 +183,21 @@ bool gfx::ShaderObjectBase::_doesValueFitInExistentialPayload(
     return true;
 }
 
-IShaderProgram* gfx::ShaderProgramBase::getInterface(const Guid& guid)
+IShaderProgram* ShaderProgramBase::getInterface(const Guid& guid)
 {
     if (guid == GfxGUID::IID_ISlangUnknown || guid == GfxGUID::IID_IShaderProgram)
         return static_cast<IShaderProgram*>(this);
     return nullptr;
 }
 
-IInputLayout* gfx::InputLayoutBase::getInterface(const Guid& guid)
+IInputLayout* InputLayoutBase::getInterface(const Guid& guid)
 {
     if (guid == GfxGUID::IID_ISlangUnknown || guid == GfxGUID::IID_IInputLayout)
         return static_cast<IInputLayout*>(this);
     return nullptr;
 }
 
-IFramebufferLayout* gfx::FramebufferLayoutBase::getInterface(const Guid& guid)
+IFramebufferLayout* FramebufferLayoutBase::getInterface(const Guid& guid)
 {
     if (guid == GfxGUID::IID_ISlangUnknown || guid == GfxGUID::IID_IFramebufferLayout)
         return static_cast<IFramebufferLayout*>(this);
@@ -260,18 +267,34 @@ SLANG_NO_THROW Result SLANG_MCALL RendererBase::getSlangSession(slang::ISession*
     return SLANG_OK;
 }
 
-SLANG_NO_THROW Result SLANG_MCALL RendererBase::createShaderObject(slang::TypeReflection* type, IShaderObject** outObject)
+SLANG_NO_THROW Result SLANG_MCALL RendererBase::createShaderObject(
+    slang::TypeReflection* type,
+    ShaderObjectContainerType container,
+    IShaderObject** outObject)
 {
     RefPtr<ShaderObjectLayoutBase> shaderObjectLayout;
-    SLANG_RETURN_FALSE_ON_FAIL(getShaderObjectLayout(type, shaderObjectLayout.writeRef()));
+    SLANG_RETURN_FALSE_ON_FAIL(getShaderObjectLayout(type, container, shaderObjectLayout.writeRef()));
     return createShaderObject(shaderObjectLayout, outObject);
 }
 
 Result RendererBase::getShaderObjectLayout(
-    slang::TypeReflection*      type,
-    ShaderObjectLayoutBase**    outLayout)
+    slang::TypeReflection* type,
+    ShaderObjectContainerType container,
+    ShaderObjectLayoutBase** outLayout)
 {
     RefPtr<ShaderObjectLayoutBase> shaderObjectLayout;
+    switch (container)
+    {
+    case ShaderObjectContainerType::StructuredBuffer:
+        type = slangContext.session->getContainerType(type, slang::ContainerType::StructuredBuffer);
+        break;
+    case ShaderObjectContainerType::Array:
+        type = slangContext.session->getContainerType(type, slang::ContainerType::UnsizedArray);
+        break;
+    default:
+        break;
+    }
+
     if( !m_shaderObjectLayoutCache.TryGetValue(type, shaderObjectLayout) )
     {
         auto typeLayout = slangContext.session->getTypeLayout(type);
@@ -373,8 +396,8 @@ Result ShaderObjectBase::_getSpecializedShaderObjectType(ExtendedShaderObjectTyp
     SLANG_RETURN_ON_FAIL(collectSpecializationArgs(specializationArgs));
     if (specializationArgs.getCount() == 0)
     {
-        shaderObjectType.componentID = getLayout()->getComponentID();
-        shaderObjectType.slangType = getLayout()->getElementTypeLayout()->getType();
+        shaderObjectType.componentID = getLayoutBase()->getComponentID();
+        shaderObjectType.slangType = getLayoutBase()->getElementTypeLayout()->getType();
     }
     else
     {
@@ -387,6 +410,94 @@ Result ShaderObjectBase::_getSpecializedShaderObjectType(ExtendedShaderObjectTyp
     return SLANG_OK;
 }
 
+Result ShaderObjectBase::setExistentialHeader(
+    slang::TypeReflection* existentialType,
+    slang::TypeReflection* concreteType,
+    ShaderOffset offset)
+{
+    // The first field of the tuple (offset zero) is the run-time type information
+    // (RTTI) ID for the concrete type being stored into the field.
+    //
+    // TODO: We need to be able to gather the RTTI type ID from `object` and then
+    // use `setData(offset, &TypeID, sizeof(TypeID))`.
+
+    // The second field of the tuple (offset 8) is the ID of the "witness" for the
+    // conformance of the concrete type to the interface used by this field.
+    //
+    auto witnessTableOffset = offset;
+    witnessTableOffset.uniformOffset += 8;
+    //
+    // Conformances of a type to an interface are computed and then stored by the
+    // Slang runtime, so we can look up the ID for this particular conformance (which
+    // will create it on demand).
+    //
+    ComPtr<slang::ISession> slangSession;
+    SLANG_RETURN_ON_FAIL(getRenderer()->getSlangSession(slangSession.writeRef()));
+    //
+    // Note: If the type doesn't actually conform to the required interface for
+    // this sub-object range, then this is the point where we will detect that
+    // fact and error out.
+    //
+    uint32_t conformanceID = 0xFFFFFFFF;
+    SLANG_RETURN_ON_FAIL(slangSession->getTypeConformanceWitnessSequentialID(
+        concreteType, existentialType, &conformanceID));
+    //
+    // Once we have the conformance ID, then we can write it into the object
+    // at the required offset.
+    //
+    SLANG_RETURN_ON_FAIL(setData(witnessTableOffset, &conformanceID, sizeof(conformanceID)));
+
+    return SLANG_OK;
+}
+
+ResourceViewBase* SimpleShaderObjectData::getResourceView(
+    RendererBase* device,
+    slang::TypeLayoutReflection* elementLayout,
+    slang::BindingType bindingType)
+{
+    if (!m_structuredBuffer)
+    {
+        // Create structured buffer resource if it has not been created.
+        IBufferResource::Desc desc = {};
+        desc.allowedStates =
+            ResourceStateSet(ResourceState::ShaderResource, ResourceState::UnorderedAccess);
+        desc.defaultState = ResourceState::ShaderResource;
+        desc.elementSize = (int)elementLayout->getSize();
+        desc.format = Format::Unknown;
+        desc.type = IResource::Type::Buffer;
+        desc.sizeInBytes = (size_t)m_ordinaryData.getCount();
+        ComPtr<IBufferResource> bufferResource;
+        SLANG_RETURN_NULL_ON_FAIL(device->createBufferResource(
+            desc, m_ordinaryData.getBuffer(), bufferResource.writeRef()));
+        m_structuredBuffer = static_cast<BufferResource*>(bufferResource.get());
+
+        // Create read-only (shader-resource) and mutable (unordered access) views.
+        ComPtr<IResourceView> resourceView;
+        IResourceView::Desc viewDesc = {};
+        viewDesc.format = Format::Unknown;
+        viewDesc.type = IResourceView::Type::ShaderResource;
+        SLANG_RETURN_NULL_ON_FAIL(device->createBufferView(
+            bufferResource.get(), viewDesc, resourceView.writeRef()));
+        m_structuredBufferView = static_cast<ResourceViewBase*>(resourceView.get());
+        viewDesc.type = IResourceView::Type::UnorderedAccess;
+        SLANG_RETURN_NULL_ON_FAIL(
+            device->createBufferView(
+            bufferResource.get(), viewDesc, resourceView.writeRef()));
+        m_rwStructuredBufferView = static_cast<ResourceViewBase*>(resourceView.get());
+    }
+
+    switch (bindingType)
+    {
+    case slang::BindingType::RawBuffer:
+        return m_structuredBufferView.Ptr();
+    case slang::BindingType::MutableRawBuffer:
+        return m_rwStructuredBufferView.Ptr();
+    default:
+        SLANG_ASSERT(false && "Invalid binding type.");
+        return nullptr;
+    }
+}
+
 Result RendererBase::maybeSpecializePipeline(
     PipelineStateBase* currentPipeline,
     ShaderObjectBase* rootObject,
@@ -425,11 +536,14 @@ Result RendererBase::maybeSpecializePipeline(
                 specializationArgs.getCount(),
                 specializedComponentType.writeRef(),
                 diagnosticBlob.writeRef());
-            if (compileRs != SLANG_OK)
+            if (diagnosticBlob)
             {
-                printf("%s\n", (char*)diagnosticBlob->getBufferPointer());
-                return SLANG_FAIL;
+                getDebugCallback()->handleMessage(
+                    compileRs == SLANG_OK ? DebugMessageType::Warning : DebugMessageType::Error,
+                    DebugMessageSource::Slang,
+                    (char*)diagnosticBlob->getBufferPointer());
             }
+            SLANG_RETURN_ON_FAIL(compileRs);
 
             // Now create specialized shader program using compiled binaries.
             ComPtr<IShaderProgram> specializedProgram;
@@ -496,4 +610,3 @@ IDebugCallback* _getNullDebugCallback()
 }
 
 } // namespace gfx
-
diff --git a/tools/gfx/renderer-shared.h b/tools/gfx/renderer-shared.h
index 9aedc8c74..ba0327cf4 100644
--- a/tools/gfx/renderer-shared.h
+++ b/tools/gfx/renderer-shared.h
@@ -242,6 +242,15 @@ protected:
     Desc m_desc;
 };
 
+class ResourceViewBase
+    : public IResourceView
+    , public Slang::ComObject
+{
+public:
+    SLANG_COM_OBJECT_IUNKNOWN_ALL
+    IResourceView* getInterface(const Slang::Guid& guid);
+};
+
 class RendererBase;
 
 typedef uint32_t ShaderComponentID;
@@ -262,6 +271,13 @@ struct ExtendedShaderObjectTypeList
         componentIDs.add(component.componentID);
         components.add(slang::SpecializationArg{ slang::SpecializationArg::Kind::Type, component.slangType });
     }
+    void addRange(const ExtendedShaderObjectTypeList& list)
+    {
+        for (Slang::Index i = 0; i < list.getCount(); i++)
+        {
+            add(list[i]);
+        }
+    }
     ExtendedShaderObjectType operator[](Slang::Index index) const
     {
         ExtendedShaderObjectType result;
@@ -274,7 +290,7 @@ struct ExtendedShaderObjectTypeList
         componentIDs.clear();
         components.clear();
     }
-    Slang::Index getCount()
+    Slang::Index getCount() const
     {
         return componentIDs.getCount();
     }
@@ -289,9 +305,19 @@ protected:
     RendererBase* m_renderer;
     slang::TypeLayoutReflection* m_elementTypeLayout = nullptr;
     ShaderComponentID m_componentID = 0;
+
+    /// The container type of this shader object. When `m_containerType` is `StructuredBuffer` or
+    /// `UnsizedArray`, this shader object represents a collection instead of a single object.
+    ShaderObjectContainerType m_containerType = ShaderObjectContainerType::None;
+
 public:
-    static slang::TypeLayoutReflection* _unwrapParameterGroups(slang::TypeLayoutReflection* typeLayout)
+    ShaderObjectContainerType getContainerType() { return m_containerType; }
+
+    static slang::TypeLayoutReflection* _unwrapParameterGroups(
+        slang::TypeLayoutReflection* typeLayout,
+        ShaderObjectContainerType& outContainerType)
     {
+        outContainerType = ShaderObjectContainerType::None;
         for (;;)
         {
             if (!typeLayout->getType())
@@ -299,12 +325,29 @@ public:
                 if (auto elementTypeLayout = typeLayout->getElementTypeLayout())
                     typeLayout = elementTypeLayout;
             }
-
+            switch (typeLayout->getKind())
+            {
+            case slang::TypeReflection::Kind::Array:
+                SLANG_ASSERT(outContainerType == ShaderObjectContainerType::None);
+                outContainerType = ShaderObjectContainerType::Array;
+                typeLayout = typeLayout->getElementTypeLayout();
+                return typeLayout;
+            case slang::TypeReflection::Kind::Resource:
+                {
+                    if (typeLayout->getResourceShape() != SLANG_STRUCTURED_BUFFER)
+                        break;
+                    SLANG_ASSERT(outContainerType == ShaderObjectContainerType::None);
+                    outContainerType = ShaderObjectContainerType::StructuredBuffer;
+                    typeLayout = typeLayout->getElementTypeLayout();
+                }
+                return typeLayout;
+            default:
+                break;
+            }
             switch (typeLayout->getKind())
             {
             default:
                 return typeLayout;
-
             case slang::TypeReflection::Kind::ConstantBuffer:
             case slang::TypeReflection::Kind::ParameterBlock:
                 typeLayout = typeLayout->getElementTypeLayout();
@@ -330,6 +373,29 @@ public:
     void initBase(RendererBase* renderer, slang::TypeLayoutReflection* elementTypeLayout);
 };
 
+class SimpleShaderObjectData
+{
+public:
+    // Any "ordinary" / uniform data for this object
+    Slang::List<char> m_ordinaryData;
+    // The structured buffer resource used when the object represents a structured buffer.
+    Slang::RefPtr<BufferResource> m_structuredBuffer;
+    // The structured buffer resource view used when the object represents a structured buffer.
+    Slang::RefPtr<ResourceViewBase> m_structuredBufferView;
+    Slang::RefPtr<ResourceViewBase> m_rwStructuredBufferView;
+
+    Slang::Index getCount() { return m_ordinaryData.getCount(); }
+    void setCount(Slang::Index count) { m_ordinaryData.setCount(count); }
+    char* getBuffer() { return m_ordinaryData.getBuffer(); }
+
+    /// Returns a StructuredBuffer resource view for GPU access into the buffer content.
+    /// Creates a StructuredBuffer resource if it has not been created.
+    ResourceViewBase* getResourceView(
+        RendererBase* device,
+        slang::TypeLayoutReflection* elementLayout,
+        slang::BindingType bindingType);
+};
+
 class ShaderObjectBase : public IShaderObject, public Slang::ComObject
 {
 protected:
@@ -364,6 +430,8 @@ public:
     // this function will return a specialized type using the bound sub-objects' type as specialization argument.
     virtual Result getSpecializedShaderObjectType(ExtendedShaderObjectType* outType);
 
+    virtual Result collectSpecializationArgs(ExtendedShaderObjectTypeList& args) = 0;
+
     RendererBase* getRenderer() { return m_layout->getDevice(); }
 
     SLANG_NO_THROW UInt SLANG_MCALL getEntryPointCount() SLANG_OVERRIDE { return 0; }
@@ -375,17 +443,346 @@ public:
         return SLANG_OK;
     }
 
-    ShaderObjectLayoutBase* getLayout()
-    {
-        return m_layout;
-    }
+    ShaderObjectLayoutBase* getLayoutBase() { return m_layout; }
 
     SLANG_NO_THROW slang::TypeLayoutReflection* SLANG_MCALL getElementTypeLayout() SLANG_OVERRIDE
     {
         return m_layout->getElementTypeLayout();
     }
 
-    virtual Result collectSpecializationArgs(ExtendedShaderObjectTypeList& args) = 0;
+    virtual SLANG_NO_THROW ShaderObjectContainerType SLANG_MCALL getContainerType() SLANG_OVERRIDE
+    {
+        return m_layout->getContainerType();
+    }
+
+        /// Sets the RTTI ID and RTTI witness table fields of an existential value.
+    Result setExistentialHeader(
+        slang::TypeReflection* existentialType,
+        slang::TypeReflection* concreteType,
+        ShaderOffset offset);
+};
+
+template<typename TShaderObjectImpl, typename TShaderObjectLayoutImpl, typename TShaderObjectData>
+class ShaderObjectBaseImpl : public ShaderObjectBase
+{
+protected:
+    TShaderObjectData m_data;
+    Slang::List<Slang::RefPtr<TShaderObjectImpl>> m_objects;
+
+    // Specialization args for a StructuredBuffer object.
+    ExtendedShaderObjectTypeList m_structuredBufferSpecializationArgs;
+
+public:
+    TShaderObjectLayoutImpl* getLayout()
+    {
+        return static_cast<TShaderObjectLayoutImpl*>(m_layout.Ptr());
+    }
+
+    void* getBuffer() { return m_data.getBuffer(); }
+    size_t getBufferSize() { return (size_t)m_data.getCount(); }
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL
+        getObject(ShaderOffset const& offset, IShaderObject** outObject) SLANG_OVERRIDE
+    {
+        SLANG_ASSERT(outObject);
+        if (offset.bindingRangeIndex < 0)
+            return SLANG_E_INVALID_ARG;
+        auto layout = getLayout();
+        if (offset.bindingRangeIndex >= layout->getBindingRangeCount())
+            return SLANG_E_INVALID_ARG;
+        auto bindingRange = layout->getBindingRange(offset.bindingRangeIndex);
+
+        returnComPtr(outObject, m_objects[bindingRange.subObjectIndex + offset.bindingArrayIndex]);
+        return SLANG_OK;
+    }
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL
+        setObject(ShaderOffset const& offset, IShaderObject* object) SLANG_OVERRIDE
+    {
+        auto layout = getLayout();
+        auto subObject = static_cast<TShaderObjectImpl*>(object);
+        // There are three different cases in `setObject`.
+        // 1. `this` object represents a StructuredBuffer, and `object` is an
+        //    element to be written into the StructuredBuffer.
+        // 2. `object` represents a StructuredBuffer and we are setting it into
+        //    a StructuredBuffer typed field in `this` object.
+        // 3. We are setting `object` as an ordinary sub-object, e.g. an existential
+        //    field, a constant buffer or a parameter block.
+        // We handle each case separately below.
+
+        if (layout->getContainerType() != ShaderObjectContainerType::None)
+        {
+            // Case 1:
+            // We are setting an element into a `StructuredBuffer` object.
+            // We need to hold a reference to the element object, as well as
+            // writing uniform data to the plain buffer.
+            if (offset.bindingArrayIndex >= m_objects.getCount())
+            {
+                m_objects.setCount(offset.bindingArrayIndex + 1);
+                auto stride = layout->getElementTypeLayout()->getStride();
+                m_data.setCount(m_objects.getCount() * stride);
+            }
+            m_objects[offset.bindingArrayIndex] = subObject;
+
+            ExtendedShaderObjectTypeList specializationArgs;
+
+            auto payloadOffset = offset;
+
+            // If the element type of the StructuredBuffer field is an existential type,
+            // we need to make sure to fill in the existential value header (RTTI ID and
+            // witness table IDs).
+            if (layout->getElementTypeLayout()->getKind() == slang::TypeReflection::Kind::Interface)
+            {
+                auto existentialType = layout->getElementTypeLayout()->getType();
+                ExtendedShaderObjectType concreteType;
+                SLANG_RETURN_ON_FAIL(subObject->getSpecializedShaderObjectType(&concreteType));
+                SLANG_RETURN_ON_FAIL(
+                    setExistentialHeader(existentialType, concreteType.slangType, offset));
+                payloadOffset.uniformOffset += 16;
+
+                // If this object is a `StructuredBuffer<ISomeInterface>`, then the
+                // specialization argument should be the specialized type of the sub object
+                // itself.
+                specializationArgs.add(concreteType);
+            }
+            else
+            {
+                // If this object is a `StructuredBuffer<SomeConcreteType>`, then the
+                // specialization
+                // argument should come recursively from the sub object.
+                subObject->collectSpecializationArgs(specializationArgs);
+            }
+            SLANG_RETURN_ON_FAIL(setData(
+                payloadOffset,
+                subObject->m_data.getBuffer(),
+                (size_t)subObject->m_data.getCount()));
+
+            // Compute specialization args for the structured buffer object.
+            // If we haven't filled anything to `m_structuredBufferSpecializationArgs` yet,
+            // use `specializationArgs` directly.
+            if (m_structuredBufferSpecializationArgs.getCount() == 0)
+            {
+                m_structuredBufferSpecializationArgs = Slang::_Move(specializationArgs);
+            }
+            else
+            {
+                // If `m_structuredBufferSpecializationArgs` already contains some arguments, we
+                // need to check if they are the same as `specializationArgs`, and replace
+                // anything that is different with `__Dynamic` because we cannot specialize the
+                // buffer type if the element types are not the same.
+                SLANG_ASSERT(
+                    m_structuredBufferSpecializationArgs.getCount() ==
+                    specializationArgs.getCount());
+                auto device = getRenderer();
+                for (Slang::Index i = 0; i < m_structuredBufferSpecializationArgs.getCount(); i++)
+                {
+                    if (m_structuredBufferSpecializationArgs[i].componentID !=
+                        specializationArgs[i].componentID)
+                    {
+                        auto dynamicType = device->slangContext.session->getDynamicType();
+                        m_structuredBufferSpecializationArgs.componentIDs[i] =
+                            device->shaderCache.getComponentId(dynamicType);
+                        m_structuredBufferSpecializationArgs.components[i] =
+                            slang::SpecializationArg::fromType(dynamicType);
+                    }
+                }
+            }
+            return SLANG_OK;
+        }
+
+        // Case 2 & 3, setting object as an StructuredBuffer, ConstantBuffer, ParameterBlock or
+        // existential value.
+
+        if (offset.bindingRangeIndex < 0)
+            return SLANG_E_INVALID_ARG;
+        if (offset.bindingRangeIndex >= layout->getBindingRangeCount())
+            return SLANG_E_INVALID_ARG;
+
+        auto bindingRangeIndex = offset.bindingRangeIndex;
+        auto bindingRange = layout->getBindingRange(bindingRangeIndex);
+
+        m_objects[bindingRange.subObjectIndex + offset.bindingArrayIndex] = subObject;
+
+        switch (bindingRange.bindingType)
+        {
+        case slang::BindingType::ExistentialValue:
+            {
+                // If the range being assigned into represents an interface/existential-type
+                // leaf field, then we need to consider how the `object` being assigned here
+                // affects specialization. We may also need to assign some data from the
+                // sub-object into the ordinary data buffer for the parent object.
+                //
+                // A leaf field of interface type is laid out inside of the parent object
+                // as a tuple of `(RTTI, WitnessTable, Payload)`. The layout of these fields
+                // is a contract between the compiler and any runtime system, so we will
+                // need to rely on details of the binary layout.
+
+                // We start by querying the layout/type of the concrete value that the
+                // application is trying to store into the field, and also the layout/type of
+                // the leaf existential-type field itself.
+                //
+                auto concreteTypeLayout = subObject->getElementTypeLayout();
+                auto concreteType = concreteTypeLayout->getType();
+                //
+                auto existentialTypeLayout =
+                    layout->getElementTypeLayout()->getBindingRangeLeafTypeLayout(
+                        bindingRangeIndex);
+                auto existentialType = existentialTypeLayout->getType();
+
+                // Fills in the first and second field of the tuple that specify RTTI type ID
+                // and witness table ID.
+                SLANG_RETURN_ON_FAIL(setExistentialHeader(existentialType, concreteType, offset));
+
+                // The third field of the tuple (offset 16) is the "payload" that is supposed to
+                // hold the data for a value of the given concrete type.
+                //
+                auto payloadOffset = offset;
+                payloadOffset.uniformOffset += 16;
+
+                // There are two cases we need to consider here for how the payload might be
+                // used:
+                //
+                // * If the concrete type of the value being bound is one that can "fit" into
+                // the
+                //   available payload space,  then it should be stored in the payload.
+                //
+                // * If the concrete type of the value cannot fit in the payload space, then it
+                //   will need to be stored somewhere else.
+                //
+                if (_doesValueFitInExistentialPayload(concreteTypeLayout, existentialTypeLayout))
+                {
+                    // If the value can fit in the payload area, then we will go ahead and copy
+                    // its bytes into that area.
+                    //
+                    setData(
+                        payloadOffset, subObject->m_data.getBuffer(), subObject->m_data.getCount());
+                }
+                else
+                {
+                    // If the value does *not *fit in the payload area, then there is nothing
+                    // we can do at this point (beyond saving a reference to the sub-object,
+                    // which was handled above).
+                    //
+                    // Once all the sub-objects have been set into the parent object, we can
+                    // compute a specialized layout for it, and that specialized layout can tell
+                    // us where the data for these sub-objects has been laid out.
+                    return SLANG_E_NOT_IMPLEMENTED;
+                }
+            }
+            break;
+        case slang::BindingType::MutableRawBuffer:
+        case slang::BindingType::RawBuffer:
+            {
+                // If we are setting into a `StructuredBuffer` field, make sure we create and set
+                // the StructuredBuffer resource as well.
+                setResource(
+                    offset,
+                    subObject->m_data.getResourceView(
+                        getRenderer(),
+                        subObject->getElementTypeLayout(),
+                        bindingRange.bindingType));
+            }
+            break;
+        }
+        return SLANG_OK;
+    }
+
+    // Appends all types that are used to specialize the element type of this shader object in
+    // `args` list.
+    virtual Result collectSpecializationArgs(ExtendedShaderObjectTypeList& args) override
+    {
+        auto device = getRenderer();
+        auto& subObjectRanges = getLayout()->getSubObjectRanges();
+        // The following logic is built on the assumption that all fields that involve
+        // existential types (and therefore require specialization) will results in a sub-object
+        // range in the type layout. This allows us to simply scan the sub-object ranges to find
+        // out all specialization arguments.
+        Slang::Index subObjectRangeCount = subObjectRanges.getCount();
+
+        for (Slang::Index subObjectRangeIndex = 0; subObjectRangeIndex < subObjectRangeCount;
+             subObjectRangeIndex++)
+        {
+            auto const& subObjectRange = subObjectRanges[subObjectRangeIndex];
+            auto const& bindingRange =
+                getLayout()->getBindingRange(subObjectRange.bindingRangeIndex);
+
+            Slang::Index oldArgsCount = args.getCount();
+
+            Slang::Index count = bindingRange.count;
+
+            for (Slang::Index subObjectIndexInRange = 0; subObjectIndexInRange < count;
+                 subObjectIndexInRange++)
+            {
+                ExtendedShaderObjectTypeList typeArgs;
+
+                auto subObject = m_objects[bindingRange.subObjectIndex + subObjectIndexInRange];
+
+                if (!subObject)
+                    continue;
+
+                switch (bindingRange.bindingType)
+                {
+                case slang::BindingType::ExistentialValue:
+                    {
+                        // A binding type of `ExistentialValue` means the sub-object represents a
+                        // interface-typed field. In this case the specialization argument for this
+                        // field is the actual specialized type of the bound shader object. If the
+                        // shader object's type is an ordinary type without existential fields, then
+                        // the type argument will simply be the ordinary type. But if the sub
+                        // object's type is itself a specialized type, we need to make sure to use
+                        // that type as the specialization argument.
+
+                        ExtendedShaderObjectType specializedSubObjType;
+                        SLANG_RETURN_ON_FAIL(
+                            subObject->getSpecializedShaderObjectType(&specializedSubObjType));
+                        typeArgs.add(specializedSubObjType);
+                        break;
+                    }
+                case slang::BindingType::ParameterBlock:
+                case slang::BindingType::ConstantBuffer:
+                    // Currently we only handle the case where the field's type is
+                    // `ParameterBlock<SomeStruct>` or `ConstantBuffer<SomeStruct>`, where
+                    // `SomeStruct` is a struct type (not directly an interface type). In this case,
+                    // we just recursively collect the specialization arguments from the bound sub
+                    // object.
+                    SLANG_RETURN_ON_FAIL(subObject->collectSpecializationArgs(typeArgs));
+                    // TODO: we need to handle the case where the field is of the form
+                    // `ParameterBlock<IFoo>`. We should treat this case the same way as the
+                    // `ExistentialValue` case here, but currently we lack a mechanism to
+                    // distinguish the two scenarios.
+                    break;
+                case slang::BindingType::RawBuffer:
+                case slang::BindingType::MutableRawBuffer:
+                     typeArgs.addRange(subObject->m_structuredBufferSpecializationArgs);
+                    break;
+                }
+
+                auto addedTypeArgCountForCurrentRange = args.getCount() - oldArgsCount;
+                if (addedTypeArgCountForCurrentRange == 0)
+                {
+                    args.addRange(typeArgs);
+                }
+                else
+                {
+                    // If type arguments for each elements in the array is different, use
+                    // `__Dynamic` type for the differing argument to disable specialization.
+                    SLANG_ASSERT(addedTypeArgCountForCurrentRange == typeArgs.getCount());
+                    for (Slang::Index i = 0; i < addedTypeArgCountForCurrentRange; i++)
+                    {
+                        if (args[i + oldArgsCount].componentID != typeArgs[i].componentID)
+                        {
+                            auto dynamicType = device->slangContext.session->getDynamicType();
+                            args.componentIDs[i + oldArgsCount] =
+                                device->shaderCache.getComponentId(dynamicType);
+                            args.components[i + oldArgsCount] =
+                                slang::SpecializationArg::fromType(dynamicType);
+                        }
+                    }
+                }
+            }
+        }
+        return SLANG_OK;
+    }
 };
 
 class ShaderProgramBase : public IShaderProgram, public Slang::ComObject
@@ -571,10 +968,14 @@ public:
     virtual SLANG_NO_THROW Result SLANG_MCALL getSlangSession(slang::ISession** outSlangSession) SLANG_OVERRIDE;
     IDevice* getInterface(const Slang::Guid& guid);
 
-    virtual SLANG_NO_THROW Result SLANG_MCALL createShaderObject(slang::TypeReflection* type, IShaderObject** outObject) SLANG_OVERRIDE;
+    virtual SLANG_NO_THROW Result SLANG_MCALL createShaderObject(
+        slang::TypeReflection* type,
+        ShaderObjectContainerType containerType,
+        IShaderObject** outObject) SLANG_OVERRIDE;
 
     Result getShaderObjectLayout(
         slang::TypeReflection*      type,
+        ShaderObjectContainerType   container,
         ShaderObjectLayoutBase**    outLayout);
 
 public:
diff --git a/tools/gfx/vulkan/render-vk.cpp b/tools/gfx/vulkan/render-vk.cpp
index f330d95a8..643c41394 100644
--- a/tools/gfx/vulkan/render-vk.cpp
+++ b/tools/gfx/vulkan/render-vk.cpp
@@ -35,6 +35,9 @@
 #ifdef Always
 #undef Always
 #endif
+#ifdef None
+#undef None
+#endif
 
 namespace gfx {
 using namespace Slang;
@@ -224,17 +227,9 @@ public:
         }
     };
 
-    class ResourceViewImpl : public IResourceView, public ComObject
+    class ResourceViewImpl : public ResourceViewBase
     {
     public:
-        SLANG_COM_OBJECT_IUNKNOWN_ALL
-        IResourceView* getInterface(const Guid& guid)
-        {
-            if (guid == GfxGUID::IID_ISlangUnknown || guid == GfxGUID::IID_IResourceView)
-                return static_cast<IResourceView*>(this);
-            return nullptr;
-        }
-    public:
         enum class ViewType
         {
             Texture,
@@ -810,6 +805,10 @@ public:
             Index count;
             Index baseIndex;
 
+                /// An index into the sub-object array if this binding range is treated
+                /// as a sub-object.
+            Index subObjectIndex;
+
                 /// The `binding` offset to apply for this range
             uint32_t bindingOffset;
 
@@ -903,6 +902,11 @@ public:
             VKDevice* m_renderer;
             slang::TypeLayoutReflection* m_elementTypeLayout;
 
+            /// The container type of this shader object. When `m_containerType` is
+            /// `StructuredBuffer` or `UnsizedArray`, this shader object represents a collection
+            /// instead of a single object.
+            ShaderObjectContainerType m_containerType = ShaderObjectContainerType::None;
+
             List<BindingRangeInfo> m_bindingRanges;
             List<SubObjectRangeInfo> m_subObjectRanges;
 
@@ -1240,15 +1244,28 @@ public:
                         typeLayout->getBindingRangeLeafTypeLayout(r);
 
                     Index baseIndex = 0;
+                    Index subObjectIndex = 0;
                     switch (slangBindingType)
                     {
                     case slang::BindingType::ConstantBuffer:
                     case slang::BindingType::ParameterBlock:
                     case slang::BindingType::ExistentialValue:
                         baseIndex = m_subObjectCount;
+                        subObjectIndex = baseIndex;
                         m_subObjectCount += count;
                         break;
-
+                    case slang::BindingType::RawBuffer:
+                    case slang::BindingType::MutableRawBuffer:
+                        if (slangLeafTypeLayout->getType()->getElementType() != nullptr)
+                        {
+                            // A structured buffer occupies both a resource slot and
+                            // a sub-object slot.
+                            subObjectIndex = m_subObjectCount;
+                            m_subObjectCount += count;
+                        }
+                        baseIndex = m_resourceViewCount;
+                        m_resourceViewCount += count;
+                        break;
                     case slang::BindingType::Sampler:
                         baseIndex = m_samplerCount;
                         m_samplerCount += count;
@@ -1281,6 +1298,7 @@ public:
                     bindingRangeInfo.bindingType = slangBindingType;
                     bindingRangeInfo.count = count;
                     bindingRangeInfo.baseIndex = baseIndex;
+                    bindingRangeInfo.subObjectIndex = subObjectIndex;
 
                     // We'd like to extract the information on the GLSL/SPIR-V
                     // `binding` that this range should bind into (or whatever
@@ -1408,7 +1426,7 @@ public:
             Result setElementTypeLayout(
                 slang::TypeLayoutReflection* typeLayout)
             {
-                typeLayout = _unwrapParameterGroups(typeLayout);
+                typeLayout = _unwrapParameterGroups(typeLayout, m_containerType);
                 m_elementTypeLayout = typeLayout;
 
                 m_totalOrdinaryDataSize = (uint32_t) typeLayout->getSize();
@@ -1556,8 +1574,6 @@ public:
 
         BindingRangeInfo const& getBindingRange(Index index) { return m_bindingRanges[index]; }
 
-        slang::TypeLayoutReflection* getElementTypeLayout() { return m_elementTypeLayout; }
-
         Index getResourceViewCount() { return m_resourceViewCount; }
         Index getSamplerCount() { return m_samplerCount; }
         Index getCombinedTextureSamplerCount() { return m_combinedTextureSamplerCount; }
@@ -1593,6 +1609,8 @@ public:
             m_subObjectRanges = builder->m_subObjectRanges;
             m_totalOrdinaryDataSize = builder->m_totalOrdinaryDataSize;
 
+            m_containerType = builder->m_containerType;
+
             // Create VkDescriptorSetLayout for all descriptor sets.
             for (auto& descriptorSetInfo : m_descriptorSetInfos)
             {
@@ -2194,7 +2212,7 @@ public:
         ConstArrayView<VkPushConstantRange> pushConstantRanges;
     };
 
-    class ShaderObjectImpl : public ShaderObjectBase
+    class ShaderObjectImpl : public ShaderObjectBaseImpl<ShaderObjectImpl, ShaderObjectLayoutImpl, SimpleShaderObjectData>
     {
     public:
         static Result create(
@@ -2220,25 +2238,14 @@ public:
             return SLANG_OK;
         }
 
-        ShaderObjectLayoutImpl* getLayout()
-        {
-            return static_cast<ShaderObjectLayoutImpl*>(m_layout.Ptr());
-        }
-
-        SLANG_NO_THROW slang::TypeLayoutReflection* SLANG_MCALL getElementTypeLayout()
-            SLANG_OVERRIDE
-        {
-            return m_layout->getElementTypeLayout();
-        }
-
         SLANG_NO_THROW Result SLANG_MCALL
             setData(ShaderOffset const& inOffset, void const* data, size_t inSize) SLANG_OVERRIDE
         {
             Index offset = inOffset.uniformOffset;
             Index size = inSize;
 
-            char* dest = m_ordinaryData.getBuffer();
-            Index availableSize = m_ordinaryData.getCount();
+            char* dest = m_data.getBuffer();
+            Index availableSize = m_data.getCount();
 
             // TODO: We really should bounds-check access rather than silently ignoring sets
             // that are too large, but we have several test cases that set more data than
@@ -2259,148 +2266,6 @@ public:
             return SLANG_OK;
         }
 
-        virtual SLANG_NO_THROW Result SLANG_MCALL
-            setObject(ShaderOffset const& offset, IShaderObject* object) SLANG_OVERRIDE
-        {
-            if (offset.bindingRangeIndex < 0)
-                return SLANG_E_INVALID_ARG;
-            auto layout = getLayout();
-            if (offset.bindingRangeIndex >= layout->getBindingRangeCount())
-                return SLANG_E_INVALID_ARG;
-
-            auto subObject = static_cast<ShaderObjectImpl*>(object);
-
-            auto bindingRangeIndex = offset.bindingRangeIndex;
-            auto& bindingRange = layout->getBindingRange(bindingRangeIndex);
-
-            m_objects[bindingRange.baseIndex + offset.bindingArrayIndex] = subObject;
-
-            // If the range being assigned into represents an interface/existential-type leaf field,
-            // then we need to consider how the `object` being assigned here affects specialization.
-            // We may also need to assign some data from the sub-object into the ordinary data
-            // buffer for the parent object.
-            //
-            if (bindingRange.bindingType == slang::BindingType::ExistentialValue)
-            {
-                // A leaf field of interface type is laid out inside of the parent object
-                // as a tuple of `(RTTI, WitnessTable, Payload)`. The layout of these fields
-                // is a contract between the compiler and any runtime system, so we will
-                // need to rely on details of the binary layout.
-
-                // We start by querying the layout/type of the concrete value that the application
-                // is trying to store into the field, and also the layout/type of the leaf
-                // existential-type field itself.
-                //
-                auto concreteTypeLayout = subObject->getElementTypeLayout();
-                auto concreteType = concreteTypeLayout->getType();
-                //
-                auto existentialTypeLayout =
-                    layout->getElementTypeLayout()->getBindingRangeLeafTypeLayout(
-                        bindingRangeIndex);
-                auto existentialType = existentialTypeLayout->getType();
-
-                // The first field of the tuple (offset zero) is the run-time type information
-                // (RTTI) ID for the concrete type being stored into the field.
-                //
-                // TODO: We need to be able to gather the RTTI type ID from `object` and then
-                // use `setData(offset, &TypeID, sizeof(TypeID))`.
-
-                // The second field of the tuple (offset 8) is the ID of the "witness" for the
-                // conformance of the concrete type to the interface used by this field.
-                //
-                auto witnessTableOffset = offset;
-                witnessTableOffset.uniformOffset += 8;
-                //
-                // Conformances of a type to an interface are computed and then stored by the
-                // Slang runtime, so we can look up the ID for this particular conformance (which
-                // will create it on demand).
-                //
-                ComPtr<slang::ISession> slangSession;
-                SLANG_RETURN_ON_FAIL(getRenderer()->getSlangSession(slangSession.writeRef()));
-                //
-                // Note: If the type doesn't actually conform to the required interface for
-                // this sub-object range, then this is the point where we will detect that
-                // fact and error out.
-                //
-                uint32_t conformanceID = 0xFFFFFFFF;
-                SLANG_RETURN_ON_FAIL(slangSession->getTypeConformanceWitnessSequentialID(
-                    concreteType, existentialType, &conformanceID));
-                //
-                // Once we have the conformance ID, then we can write it into the object
-                // at the required offset.
-                //
-                SLANG_RETURN_ON_FAIL(
-                    setData(witnessTableOffset, &conformanceID, sizeof(conformanceID)));
-
-                // The third field of the tuple (offset 16) is the "payload" that is supposed to
-                // hold the data for a value of the given concrete type.
-                //
-                auto payloadOffset = offset;
-                payloadOffset.uniformOffset += 16;
-
-                // There are two cases we need to consider here for how the payload might be used:
-                //
-                // * If the concrete type of the value being bound is one that can "fit" into the
-                //   available payload space,  then it should be stored in the payload.
-                //
-                // * If the concrete type of the value cannot fit in the payload space, then it
-                //   will need to be stored somewhere else.
-                //
-                if (_doesValueFitInExistentialPayload(concreteTypeLayout, existentialTypeLayout))
-                {
-                    // If the value can fit in the payload area, then we will go ahead and copy
-                    // its bytes into that area.
-                    //
-                    setData(
-                        payloadOffset,
-                        subObject->m_ordinaryData.getBuffer(),
-                        subObject->m_ordinaryData.getCount());
-                }
-                else
-                {
-                    // If the value does *not *fit in the payload area, then there is nothing
-                    // we can do at this point (beyond saving a reference to the sub-object, which
-                    // was handled above).
-                    //
-                    // Once all the sub-objects have been set into the parent object, we can
-                    // compute a specialized layout for it, and that specialized layout can tell
-                    // us where the data for these sub-objects has been laid out.
-                }
-            }
-
-            return SLANG_OK;
-        }
-
-        virtual SLANG_NO_THROW Result SLANG_MCALL
-            getObject(ShaderOffset const& offset, IShaderObject** outObject) SLANG_OVERRIDE
-        {
-            SLANG_ASSERT(outObject);
-            if (offset.bindingRangeIndex < 0)
-                return SLANG_E_INVALID_ARG;
-            auto layout = getLayout();
-            if (offset.bindingRangeIndex >= layout->getBindingRangeCount())
-                return SLANG_E_INVALID_ARG;
-            auto& bindingRange = layout->getBindingRange(offset.bindingRangeIndex);
-
-            auto object = m_objects[bindingRange.baseIndex + offset.bindingArrayIndex].Ptr();
-            returnComPtr(outObject, object);
-
-            //        auto& subObjectRange =
-            //        m_layout->getSubObjectRange(bindingRange.subObjectRangeIndex); *outObject =
-            //        m_objects[subObjectRange.baseIndex + offset.bindingArrayIndex];
-
-            return SLANG_OK;
-
-#if 0
-        SLANG_ASSERT(bindingRange.descriptorSetIndex >= 0);
-        SLANG_ASSERT(bindingRange.descriptorSetIndex < m_descriptorSets.getCount());
-        auto& descriptorSet = m_descriptorSets[bindingRange.descriptorSetIndex];
-
-        descriptorSet->setConstantBuffer(bindingRange.rangeIndexInDescriptorSet, offset.bindingArrayIndex, buffer);
-        return SLANG_OK;
-#endif
-        }
-
         SLANG_NO_THROW Result SLANG_MCALL
             setResource(ShaderOffset const& offset, IResourceView* resourceView) SLANG_OVERRIDE
         {
@@ -2450,68 +2315,6 @@ public:
             return SLANG_OK;
         }
 
-    public:
-        // Appends all types that are used to specialize the element type of this shader object in
-        // `args` list.
-        virtual Result collectSpecializationArgs(ExtendedShaderObjectTypeList& args) override
-        {
-            auto& subObjectRanges = getLayout()->getSubObjectRanges();
-            // The following logic is built on the assumption that all fields that involve
-            // existential types (and therefore require specialization) will results in a sub-object
-            // range in the type layout. This allows us to simply scan the sub-object ranges to find
-            // out all specialization arguments.
-            Index subObjectRangeCount = subObjectRanges.getCount();
-            for (Index subObjectRangeIndex = 0; subObjectRangeIndex < subObjectRangeCount;
-                 subObjectRangeIndex++)
-            {
-                auto const& subObjectRange = subObjectRanges[subObjectRangeIndex];
-                auto const& bindingRange =
-                    getLayout()->getBindingRange(subObjectRange.bindingRangeIndex);
-
-                Index count = bindingRange.count;
-                SLANG_ASSERT(count == 1);
-
-                Index subObjectIndexInRange = 0;
-                auto subObject = m_objects[bindingRange.baseIndex + subObjectIndexInRange];
-
-                switch (bindingRange.bindingType)
-                {
-                case slang::BindingType::ExistentialValue:
-                    {
-                        // A binding type of `ExistentialValue` means the sub-object represents a
-                        // interface-typed field. In this case the specialization argument for this
-                        // field is the actual specialized type of the bound shader object. If the
-                        // shader object's type is an ordinary type without existential fields, then
-                        // the type argument will simply be the ordinary type. But if the sub
-                        // object's type is itself a specialized type, we need to make sure to use
-                        // that type as the specialization argument.
-
-                        ExtendedShaderObjectType specializedSubObjType;
-                        SLANG_RETURN_ON_FAIL(
-                            subObject->getSpecializedShaderObjectType(&specializedSubObjType));
-                        args.add(specializedSubObjType);
-                        break;
-                    }
-                case slang::BindingType::ParameterBlock:
-                case slang::BindingType::ConstantBuffer:
-                    // Currently we only handle the case where the field's type is
-                    // `ParameterBlock<SomeStruct>` or `ConstantBuffer<SomeStruct>`, where
-                    // `SomeStruct` is a struct type (not directly an interface type). In this case,
-                    // we just recursively collect the specialization arguments from the bound sub
-                    // object.
-                    SLANG_RETURN_ON_FAIL(subObject->collectSpecializationArgs(args));
-                    // TODO: we need to handle the case where the field is of the form
-                    // `ParameterBlock<IFoo>`. We should treat this case the same way as the
-                    // `ExistentialValue` case here, but currently we lack a mechanism to
-                    // distinguish the two scenarios.
-                    break;
-                }
-                // TODO: need to handle another case where specialization happens on resources
-                // fields e.g. `StructuredBuffer<IFoo>`.
-            }
-            return SLANG_OK;
-        }
-
     protected:
         friend class RootShaderObjectLayout;
 
@@ -2533,8 +2336,8 @@ public:
             size_t uniformSize = layout->getElementTypeLayout()->getSize();
             if (uniformSize)
             {
-                m_ordinaryData.setCount(uniformSize);
-                memset(m_ordinaryData.getBuffer(), 0, uniformSize);
+                m_data.setCount(uniformSize);
+                memset(m_data.getBuffer(), 0, uniformSize);
             }
 
 #if 0
@@ -2583,7 +2386,7 @@ public:
                     RefPtr<ShaderObjectImpl> subObject;
                     SLANG_RETURN_ON_FAIL(
                         ShaderObjectImpl::create(device, subObjectLayout, subObject.writeRef()));
-                    m_objects[bindingRangeInfo.baseIndex + i] = subObject;
+                    m_objects[bindingRangeInfo.subObjectIndex + i] = subObject;
                 }
             }
 
@@ -2599,8 +2402,8 @@ public:
             size_t destSize,
             ShaderObjectLayoutImpl* specializedLayout)
         {
-            auto src = m_ordinaryData.getBuffer();
-            auto srcSize = size_t(m_ordinaryData.getCount());
+            auto src = m_data.getBuffer();
+            auto srcSize = size_t(m_data.getCount());
 
             SLANG_ASSERT(srcSize <= destSize);
 
@@ -2675,7 +2478,7 @@ public:
 
                 for (Slang::Index i = 0; i < count; ++i)
                 {
-                    auto subObject = m_objects[bindingRangeInfo.baseIndex + i];
+                    auto subObject = m_objects[bindingRangeInfo.subObjectIndex + i];
 
                     RefPtr<ShaderObjectLayoutImpl> subObjectLayout;
                     SLANG_RETURN_ON_FAIL(
@@ -3083,7 +2886,7 @@ public:
             {
                 auto const& bindingRangeInfo = specializedLayout->getBindingRange(subObjectRange.bindingRangeIndex);
                 auto count = bindingRangeInfo.count;
-                auto baseIndex = bindingRangeInfo.baseIndex;
+                auto subObjectIndex = bindingRangeInfo.subObjectIndex;
 
                 auto subObjectLayout = subObjectRange.layout;
 
@@ -3109,7 +2912,7 @@ public:
                             // the ordinary data buffer (if needed) and any other
                             // bindings it recursively contains.
                             //
-                            ShaderObjectImpl* subObject = m_objects[baseIndex + i];
+                            ShaderObjectImpl* subObject = m_objects[subObjectIndex + i];
                             subObject->bindAsConstantBuffer(encoder, context, objOffset, subObjectLayout);
 
                             // When dealing with arrays of sub-objects, we need to make
@@ -3129,7 +2932,7 @@ public:
                             // from `ConstantBuffer<X>`, except that we call `bindAsParameterBlock`
                             // instead (understandably).
                             //
-                            ShaderObjectImpl* subObject = m_objects[baseIndex + i];
+                            ShaderObjectImpl* subObject = m_objects[subObjectIndex + i];
                             subObject->bindAsParameterBlock(encoder, context, objOffset, subObjectLayout);
 
                             objOffset += rangeStride;
@@ -3162,13 +2965,16 @@ public:
                             // have been handled as part of the buffer for a parent object
                             // already.
                             //
-                            ShaderObjectImpl* subObject = m_objects[baseIndex + i];
+                            ShaderObjectImpl* subObject = m_objects[subObjectIndex + i];
                             subObject->bindAsValue(encoder, context, BindingOffset(objOffset), subObjectLayout);
                             objOffset += objStride;
                         }
                     }
                     break;
-
+                case slang::BindingType::RawBuffer:
+                case slang::BindingType::MutableRawBuffer:
+                    // No action needed for sub-objects bound though a `StructuredBuffer`.
+                    break;
                 default:
                     SLANG_ASSERT(!"unsupported sub-object type");
                     return SLANG_FAIL;
@@ -3297,17 +3103,12 @@ public:
             return SLANG_OK;
         }
 
-        /// Any "ordinary" / uniform data for this object
-        List<char> m_ordinaryData;
-
         List<RefPtr<ResourceViewImpl>> m_resourceViews;
 
         List<RefPtr<SamplerStateImpl>> m_samplers;
 
         List<CombinedTextureSamplerSlot> m_combinedTextureSamplers;
 
-        List<RefPtr<ShaderObjectImpl>> m_objects;
-
         // The version number of the transient resource heap that contains up-to-date
         // constant buffer content for this shader object.
         uint64_t m_upToDateConstantBufferHeapVersion;
@@ -3345,7 +3146,9 @@ public:
             auto device = getDevice();
             RefPtr<ShaderObjectLayoutImpl> layout;
             SLANG_RETURN_ON_FAIL(device->getShaderObjectLayout(
-                extendedType.slangType, (ShaderObjectLayoutBase**)layout.writeRef()));
+                extendedType.slangType,
+                m_layout->getContainerType(),
+                (ShaderObjectLayoutBase**)layout.writeRef()));
 
             returnRefPtrMove(outLayout, layout);
             return SLANG_OK;
@@ -3389,7 +3192,7 @@ public:
             //
             // TODO: Can/should this function be renamed as just `bindAsPushConstantBuffer`?
             //
-            if (m_ordinaryData.getCount())
+            if (m_data.getCount())
             {
                 // The index of the push constant range to bind should be
                 // passed down as part of the `offset`, and we will increment
@@ -3410,9 +3213,9 @@ public:
                 // TODO: This would not be the case if specialization for interface-type
                 // parameters led to the entry point having "pending" ordinary data.
                 //
-                SLANG_ASSERT(pushConstantRange.size == (uint32_t) m_ordinaryData.getCount());
+                SLANG_ASSERT(pushConstantRange.size == (uint32_t)m_data.getCount());
 
-                auto pushConstantData = m_ordinaryData.getBuffer();
+                auto pushConstantData = m_data.getBuffer();
 
                 encoder->m_api->vkCmdPushConstants(
                     encoder->m_commandBuffer->m_commandBuffer,
@@ -6292,9 +6095,10 @@ Result VKDevice::createBufferView(IBufferResource* buffer, IResourceView::Desc c
         return SLANG_FAIL;
 
     case IResourceView::Type::UnorderedAccess:
+    case IResourceView::Type::ShaderResource:
         // Is this a formatted view?
         //
-        if(desc.format == Format::Unknown)
+        if (desc.format == Format::Unknown)
         {
             // Buffer usage that doesn't involve formatting doesn't
             // require a view in Vulkan.
@@ -6310,7 +6114,6 @@ Result VKDevice::createBufferView(IBufferResource* buffer, IResourceView::Desc c
         // it just like we would for a "sampled" buffer:
         //
         // FALLTHROUGH
-    case IResourceView::Type::ShaderResource:
         {
             VkBufferViewCreateInfo info = { VK_STRUCTURE_TYPE_BUFFER_VIEW_CREATE_INFO };
 
diff --git a/tools/render-test/render-test-main.cpp b/tools/render-test/render-test-main.cpp
index c6f775312..5feeeaf21 100644
--- a/tools/render-test/render-test-main.cpp
+++ b/tools/render-test/render-test-main.cpp
@@ -232,13 +232,20 @@ struct AssignValsFromLayoutContext
     SlangResult assignTexture(ShaderCursor const& dstCursor, ShaderInputLayout::TextureVal* srcVal)
     {
         ComPtr<ITextureResource> texture;
-        SLANG_RETURN_ON_FAIL(ShaderRendererUtil::generateTextureResource(
-            srcVal->textureDesc, ResourceState::ShaderResource, device, texture));
+        ResourceState defaultState = ResourceState::ShaderResource;
+        IResourceView::Type viewType = IResourceView::Type::ShaderResource;
+
+        if (srcVal->textureDesc.isRWTexture)
+        {
+            defaultState = ResourceState::UnorderedAccess;
+            viewType = IResourceView::Type::UnorderedAccess;
+        }
 
-        // TODO: support UAV textures...
+        SLANG_RETURN_ON_FAIL(ShaderRendererUtil::generateTextureResource(
+            srcVal->textureDesc, defaultState, device, texture));
 
         IResourceView::Desc viewDesc;
-        viewDesc.type = IResourceView::Type::ShaderResource;
+        viewDesc.type = viewType;
         viewDesc.format = texture->getDesc()->format;
         auto textureView = device->createTextureView(
             texture,
@@ -1078,7 +1085,14 @@ static SlangResult _innerMain(Slang::StdWriters* stdWriters, SlangSession* sessi
 
         desc.requiredFeatures = requiredFeatureList.getBuffer();
         desc.requiredFeatureCount = (int)requiredFeatureList.getCount();
-
+        for (int i = 0; i < options.slangArgCount; i++)
+        {
+            if (UnownedStringSlice(options.slangArgs[i]) == "-matrix-layout-column-major")
+            {
+                desc.slang.defaultMatrixLayoutMode = SLANG_MATRIX_LAYOUT_COLUMN_MAJOR;
+            }
+        }
+        
         desc.nvapiExtnSlot = int(nvapiExtnSlot);
         desc.slang.slangGlobalSession = session;
 
diff --git a/tools/render-test/shader-input-layout.cpp b/tools/render-test/shader-input-layout.cpp
index 43dfee804..bece936dd 100644
--- a/tools/render-test/shader-input-layout.cpp
+++ b/tools/render-test/shader-input-layout.cpp
@@ -416,7 +416,17 @@ namespace renderer_test
 
         String parseTypeName(TokenReader& parser)
         {
-            return parser.ReadWord();
+            String typeName = parser.ReadWord();
+            if (parser.AdvanceIf("<"))
+            {
+                StringBuilder sb;
+                sb << typeName << "<";
+                sb << parseTypeName(parser);
+                sb << ">";
+                parser.Read(">");
+                return sb.ProduceString();
+            }
+            return typeName;
         }
 
         RefPtr<ShaderInputLayout::Val> parseValExpr(TokenReader& parser)