Compile append and consume structured buffers to glsl. (#3142)

* Compile append and consume structured buffers to glsl.

* Fix.

* Update CI config.

---------

Co-authored-by: Yong He <yhe@nvidia.com>

40 files changed, 809 insertions, 239 deletions

diff --git a/.github/workflows/windows-selfhosted.yml b/.github/workflows/windows-selfhosted.yml
index 9b7c7b6eb..08f079537 100644
--- a/.github/workflows/windows-selfhosted.yml
+++ b/.github/workflows/windows-selfhosted.yml
@@ -12,8 +12,9 @@ jobs:
   build:
     runs-on: [Windows, self-hosted]
     timeout-minutes: 100
-
+    continue-on-error: true
     strategy:
+      fail-fast: false
       matrix:
         configuration: ['Release']
         platform: ['x64']
@@ -38,11 +39,13 @@ jobs:
         MSBuild.exe slang.sln -v:m -m -property:Configuration=${{matrix.configuration}} -property:Platform=${{matrix.platform}} -property:WindowsTargetPlatformVersion=10.0.19041.0 -maxcpucount:12
     - name: test-spirv-direct
       run: |
-        set PATH=%PATH%;.\external\slang-binaries\spirv-tools\windows-${{matrix.testPlatform}}\bin\
-        ".\bin\windows-${{matrix.testPlatform}}\${{matrix.configuration}}\slang-test.exe" tests/ -use-test-server -emit-spirv-directly -expected-failure-list tests/expected-failure.txt -api vk 2>&1
-      shell: cmd
+        $ErrorActionPreference = "SilentlyContinue"
+        $env:Path += ';.\external\slang-binaries\spirv-tools\windows-${{matrix.testPlatform}}\bin\'
+        .\bin\windows-${{matrix.testPlatform}}\${{matrix.configuration}}\slang-test.exe tests/ -use-test-server -emit-spirv-directly -expected-failure-list tests/expected-failure.txt -api vk
+      
     - name: test
       run: |
-        set PATH=%PATH%;.\external\slang-binaries\spirv-tools\windows-${{matrix.testPlatform}}\bin\
-        ".\bin\windows-${{matrix.testPlatform}}\${{matrix.configuration}}\slang-test.exe" -use-test-server -api vk 2>&1
-      shell: cmd
+        $ErrorActionPreference = "SilentlyContinue"
+        $env:Path += ';.\external\slang-binaries\spirv-tools\windows-${{matrix.testPlatform}}\bin\'
+        .\bin\windows-${{matrix.testPlatform}}\${{matrix.configuration}}\slang-test.exe -use-test-server -api all-cpu
+      
diff --git a/build/visual-studio/slang/slang.vcxproj b/build/visual-studio/slang/slang.vcxproj
index 2e82d97ce..ebbb4f98d 100644
--- a/build/visual-studio/slang/slang.vcxproj
+++ b/build/visual-studio/slang/slang.vcxproj
@@ -408,6 +408,7 @@ IF EXIST ..\..\..\external\slang-glslang\bin\windows-aarch64\release\slang-glsla
     <ClInclude Include="..\..\..\source\slang\slang-ir-liveness.h" />

     <ClInclude Include="..\..\..\source\slang\slang-ir-loop-inversion.h" />

     <ClInclude Include="..\..\..\source\slang\slang-ir-loop-unroll.h" />

+    <ClInclude Include="..\..\..\source\slang\slang-ir-lower-append-consume-structured-buffer.h" />

     <ClInclude Include="..\..\..\source\slang\slang-ir-lower-binding-query.h" />

     <ClInclude Include="..\..\..\source\slang\slang-ir-lower-bit-cast.h" />

     <ClInclude Include="..\..\..\source\slang\slang-ir-lower-buffer-element-type.h" />

@@ -618,6 +619,7 @@ IF EXIST ..\..\..\external\slang-glslang\bin\windows-aarch64\release\slang-glsla
     <ClCompile Include="..\..\..\source\slang\slang-ir-liveness.cpp" />

     <ClCompile Include="..\..\..\source\slang\slang-ir-loop-inversion.cpp" />

     <ClCompile Include="..\..\..\source\slang\slang-ir-loop-unroll.cpp" />

+    <ClCompile Include="..\..\..\source\slang\slang-ir-lower-append-consume-structured-buffer.cpp" />

     <ClCompile Include="..\..\..\source\slang\slang-ir-lower-binding-query.cpp" />

     <ClCompile Include="..\..\..\source\slang\slang-ir-lower-bit-cast.cpp" />

     <ClCompile Include="..\..\..\source\slang\slang-ir-lower-buffer-element-type.cpp" />

diff --git a/build/visual-studio/slang/slang.vcxproj.filters b/build/visual-studio/slang/slang.vcxproj.filters
index 546a03ab0..0c4f8a4d3 100644
--- a/build/visual-studio/slang/slang.vcxproj.filters
+++ b/build/visual-studio/slang/slang.vcxproj.filters
@@ -312,6 +312,9 @@
     <ClInclude Include="..\..\..\source\slang\slang-ir-loop-unroll.h">

       <Filter>Header Files</Filter>

     </ClInclude>

+    <ClInclude Include="..\..\..\source\slang\slang-ir-lower-append-consume-structured-buffer.h">

+      <Filter>Header Files</Filter>

+    </ClInclude>

     <ClInclude Include="..\..\..\source\slang\slang-ir-lower-binding-query.h">

       <Filter>Header Files</Filter>

     </ClInclude>

@@ -938,6 +941,9 @@
     <ClCompile Include="..\..\..\source\slang\slang-ir-loop-unroll.cpp">

       <Filter>Source Files</Filter>

     </ClCompile>

+    <ClCompile Include="..\..\..\source\slang\slang-ir-lower-append-consume-structured-buffer.cpp">

+      <Filter>Source Files</Filter>

+    </ClCompile>

     <ClCompile Include="..\..\..\source\slang\slang-ir-lower-binding-query.cpp">

       <Filter>Source Files</Filter>

     </ClCompile>

diff --git a/source/slang/hlsl.meta.slang b/source/slang/hlsl.meta.slang
index b690a5910..3dcdc6c54 100644
--- a/source/slang/hlsl.meta.slang
+++ b/source/slang/hlsl.meta.slang
@@ -7,15 +7,30 @@ typedef uint UINT;
 [ForceInline] float3 __asFloat3(float3 v) { return v; }
 
 __generic<T>
+__intrinsic_op($(kIROp_StructuredBufferGetDimensions))
+uint2 __structuredBufferGetDimensions(AppendStructuredBuffer<T> buffer);
+
+__generic<T>
+__intrinsic_op($(kIROp_StructuredBufferGetDimensions))
+uint2 __structuredBufferGetDimensions(ConsumeStructuredBuffer<T> buffer);
+
+__generic<T>
 __magic_type(HLSLAppendStructuredBufferType)
 __intrinsic_type($(kIROp_HLSLAppendStructuredBufferType))
 struct AppendStructuredBuffer
 {
+    __intrinsic_op($(kIROp_StructuredBufferAppend))
     void Append(T value);
 
+    [ForceInline]
     void GetDimensions(
         out uint numStructs,
-        out uint stride);
+        out uint stride)
+    {
+        let result = __structuredBufferGetDimensions(this);
+        numStructs = result.x;
+        stride = result.y;
+    }
 };
 
 __magic_type(HLSLByteAddressBufferType)
@@ -257,11 +272,18 @@ __magic_type(HLSLConsumeStructuredBufferType)
 __intrinsic_type($(kIROp_HLSLConsumeStructuredBufferType))
 struct ConsumeStructuredBuffer
 {
+    __intrinsic_op($(kIROp_StructuredBufferConsume))
     T Consume();
 
+    [ForceInline]
     void GetDimensions(
         out uint numStructs,
-        out uint stride);
+        out uint stride)
+    {
+        let result = __structuredBufferGetDimensions(this);
+        numStructs = result.x;
+        stride = result.y;
+    }
 };
 
 __generic<T, let N : int>
diff --git a/source/slang/slang-emit-c-like.cpp b/source/slang/slang-emit-c-like.cpp
index e1f631283..75a15d0c9 100644
--- a/source/slang/slang-emit-c-like.cpp
+++ b/source/slang/slang-emit-c-like.cpp
@@ -455,6 +455,66 @@ void CLikeSourceEmitter::emitRTTIObject(IRRTTIObject* rttiObject)
     // This is only used in targets that support dynamic dispatching.
 }
 
+void CLikeSourceEmitter::defaultEmitInstStmt(IRInst* inst)
+{
+    switch (inst->getOp())
+    {
+    case kIROp_AtomicCounterIncrement:
+        {
+            auto oldValName = getName(inst);
+            m_writer->emit("int ");
+            m_writer->emit(oldValName);
+            m_writer->emit(";\n");
+            m_writer->emit("InterlockedAdd(");
+            emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
+            m_writer->emit(", 1, ");
+            m_writer->emit(oldValName);
+            m_writer->emit(");\n");
+        }
+        break;
+    case kIROp_AtomicCounterDecrement:
+        {
+            auto oldValName = getName(inst);
+            m_writer->emit("int ");
+            m_writer->emit(oldValName);
+            m_writer->emit(";\n");
+            m_writer->emit("InterlockedAdd(");
+            emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
+            m_writer->emit(", -1, ");
+            m_writer->emit(oldValName);
+            m_writer->emit(");\n");
+        }
+        break;
+    case kIROp_StructuredBufferGetDimensions:
+        {
+            auto count = _generateUniqueName(UnownedStringSlice("_elementCount"));
+            auto stride = _generateUniqueName(UnownedStringSlice("_stride"));
+
+            m_writer->emit("uint ");
+            m_writer->emit(count);
+            m_writer->emit(";\n");
+            m_writer->emit("uint ");
+            m_writer->emit(stride);
+            m_writer->emit(";\n");
+            emitOperand(inst->getOperand(0), leftSide(getInfo(EmitOp::General), getInfo(EmitOp::Postfix)));
+            m_writer->emit(".GetDimensions(");
+            m_writer->emit(count);
+            m_writer->emit(", ");
+            m_writer->emit(stride);
+            m_writer->emit(");\n");
+            emitInstResultDecl(inst);
+            m_writer->emit("uint2(");
+            m_writer->emit(count);
+            m_writer->emit(", ");
+            m_writer->emit(stride);
+            m_writer->emit(");\n");
+        }
+        break;
+    default:
+        diagnoseUnhandledInst(inst);
+    }
+}
+
 
 void CLikeSourceEmitter::emitTypeImpl(IRType* type, const StringSliceLoc* nameAndLoc)
 {
@@ -1874,6 +1934,16 @@ void CLikeSourceEmitter::emitInstExpr(IRInst* inst, const EmitOpInfo& inOuterPre
     defaultEmitInstExpr(inst, inOuterPrec);
 }
 
+void CLikeSourceEmitter::emitInstStmt(IRInst* inst)
+{
+    // Try target specific impl first
+    if (tryEmitInstStmtImpl(inst))
+    {
+        return;
+    }
+    defaultEmitInstStmt(inst);
+}
+
 void CLikeSourceEmitter::diagnoseUnhandledInst(IRInst* inst)
 {
     getSink()->diagnose(inst, Diagnostics::unimplemented, "unexpected IR opcode during code emit");
@@ -2193,6 +2263,23 @@ void CLikeSourceEmitter::defaultEmitInstExpr(IRInst* inst, const EmitOpInfo& inO
         }
         break;
 
+    case kIROp_StructuredBufferAppend:
+        {
+            auto outer = getInfo(EmitOp::General);
+            emitOperand(inst->getOperand(0), leftSide(outer, getInfo(EmitOp::Postfix)));
+            m_writer->emit(".Append(");
+            emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
+            m_writer->emit(")");
+        }
+        break;
+    case kIROp_StructuredBufferConsume:
+        {
+            auto outer = getInfo(EmitOp::General);
+            emitOperand(inst->getOperand(0), leftSide(outer, getInfo(EmitOp::Postfix)));
+            m_writer->emit(".Consume()");
+        }
+        break;
+
     case kIROp_Call:
         {
             emitCallExpr((IRCall*)inst, outerPrec);
@@ -2562,7 +2649,10 @@ void CLikeSourceEmitter::_emitInst(IRInst* inst)
 
         // Insts that needs to be emitted as code blocks.
     case kIROp_CudaKernelLaunch:
-        emitInstStmtImpl(inst);
+    case kIROp_AtomicCounterIncrement:
+    case kIROp_AtomicCounterDecrement:
+    case kIROp_StructuredBufferGetDimensions:
+        emitInstStmt(inst);
         break;
 
     case kIROp_LiveRangeStart:
diff --git a/source/slang/slang-emit-c-like.h b/source/slang/slang-emit-c-like.h
index 4f8d23a0d..420132a5d 100644
--- a/source/slang/slang-emit-c-like.h
+++ b/source/slang/slang-emit-c-like.h
@@ -549,7 +549,10 @@ public:
 
     virtual bool tryEmitGlobalParamImpl(IRGlobalParam* varDecl, IRType* varType) { SLANG_UNUSED(varDecl); SLANG_UNUSED(varType); return false; }
     virtual bool tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOuterPrec) { SLANG_UNUSED(inst); SLANG_UNUSED(inOuterPrec); return false; }
-    virtual void emitInstStmtImpl(IRInst* inst) { SLANG_UNUSED(inst); }
+    virtual bool tryEmitInstStmtImpl(IRInst* inst) { SLANG_UNUSED(inst); return false; }
+
+    void defaultEmitInstStmt(IRInst* inst);
+    void emitInstStmt(IRInst* inst);
 
     virtual void emitPostKeywordTypeAttributesImpl(IRInst* inst) { SLANG_UNUSED(inst); }
 
diff --git a/source/slang/slang-emit-glsl.cpp b/source/slang/slang-emit-glsl.cpp
index 0920c236c..e1f74f70d 100644
--- a/source/slang/slang-emit-glsl.cpp
+++ b/source/slang/slang-emit-glsl.cpp
@@ -201,8 +201,11 @@ void GLSLSourceEmitter::_emitGLSLStructuredBuffer(IRGlobalParam* varDecl, IRHLSL
     m_writer->emit("buffer ");
 
     // Generate a dummy name for the block
-    m_writer->emit("_S");
-    m_writer->emit(m_uniqueIDCounter++);
+    StringBuilder blockTypeName;
+    blockTypeName << "StructuredBuffer_";
+    getTypeNameHint(blockTypeName, structuredBufferType->getElementType());
+    blockTypeName << "_t";
+    m_writer->emit(_generateUniqueName(blockTypeName.produceString().getUnownedSlice()));
 
     m_writer->emit(" {\n");
     m_writer->indent();
@@ -2007,6 +2010,37 @@ bool GLSLSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOu
     return false;
 }
 
+bool GLSLSourceEmitter::tryEmitInstStmtImpl(IRInst* inst)
+{
+    switch (inst->getOp())
+    {
+    case kIROp_AtomicCounterIncrement:
+        {
+            auto oldValName = getName(inst);
+            m_writer->emit("int ");
+            m_writer->emit(oldValName);
+            m_writer->emit(" = ");
+            m_writer->emit("atomicAdd(");
+            emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
+            m_writer->emit(", 1);\n");
+            return true;
+        }
+    case kIROp_AtomicCounterDecrement:
+        {
+            auto oldValName = getName(inst);
+            m_writer->emit("int ");
+            m_writer->emit(oldValName);
+            m_writer->emit(" = ");
+            m_writer->emit("atomicAdd(");
+            emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
+            m_writer->emit(", -1);\n");
+            return true;
+        }
+    default:
+        return false;
+    }
+}
+
 void GLSLSourceEmitter::handleRequiredCapabilitiesImpl(IRInst* inst)
 {
     // Does this function declare any requirements on GLSL version or
diff --git a/source/slang/slang-emit-glsl.h b/source/slang/slang-emit-glsl.h
index d0cabfa94..7c1a15315 100644
--- a/source/slang/slang-emit-glsl.h
+++ b/source/slang/slang-emit-glsl.h
@@ -48,6 +48,8 @@ protected:
 
     virtual bool tryEmitGlobalParamImpl(IRGlobalParam* varDecl, IRType* varType) SLANG_OVERRIDE;
     virtual bool tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOuterPrec) SLANG_OVERRIDE;
+    virtual bool tryEmitInstStmtImpl(IRInst* inst) SLANG_OVERRIDE;
+
     virtual void emitGlobalInstImpl(IRInst* inst) override;
     void emitBufferPointerTypeDefinition(IRInst* ptrType);
 
diff --git a/source/slang/slang-emit-torch.cpp b/source/slang/slang-emit-torch.cpp
index bdb650607..ef04f33ba 100644
--- a/source/slang/slang-emit-torch.cpp
+++ b/source/slang/slang-emit-torch.cpp
@@ -65,12 +65,12 @@ void emitTorchScalarTypeName(SourceWriter* m_writer, IRInst* type)
     }
 }
 
-void TorchCppSourceEmitter::emitInstStmtImpl(IRInst* inst)
+bool TorchCppSourceEmitter::tryEmitInstStmtImpl(IRInst* inst)
 {
     switch (inst->getOp())
     {
     default:
-        return;
+        return false;
     case kIROp_CudaKernelLaunch:
         {
             m_writer->emit("AT_CUDA_CHECK(cudaLaunchKernel(");
@@ -101,7 +101,7 @@ void TorchCppSourceEmitter::emitInstStmtImpl(IRInst* inst)
             emitOperand(inst->getOperand(4), getInfo(EmitOp::General));
             m_writer->emit(")));\n");
         
-            break;
+            return true;
         }
     }
 }
diff --git a/source/slang/slang-emit-torch.h b/source/slang/slang-emit-torch.h
index aeb9058a4..9e76e42d1 100644
--- a/source/slang/slang-emit-torch.h
+++ b/source/slang/slang-emit-torch.h
@@ -19,7 +19,7 @@ public:
 
 protected:
     // CPPSourceEmitter overrides 
-    virtual void emitInstStmtImpl(IRInst* inst) override;
+    virtual bool tryEmitInstStmtImpl(IRInst* inst) override;
 
     virtual bool tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOuterPrec) override;
     virtual SlangResult calcTypeName(IRType* type, CodeGenTarget target, StringBuilder& out) override;
diff --git a/source/slang/slang-emit.cpp b/source/slang/slang-emit.cpp
index 6521b05ba..03d62b540 100644
--- a/source/slang/slang-emit.cpp
+++ b/source/slang/slang-emit.cpp
@@ -31,6 +31,7 @@
 #include "slang-ir-legalize-varying-params.h"
 #include "slang-ir-link.h"
 #include "slang-ir-com-interface.h"
+#include "slang-ir-lower-append-consume-structured-buffer.h"
 #include "slang-ir-lower-binding-query.h"
 #include "slang-ir-lower-generics.h"
 #include "slang-ir-lower-tuple-types.h"
@@ -494,6 +495,14 @@ Result linkAndOptimizeIR(
 
     validateIRModuleIfEnabled(codeGenContext, irModule);
 
+    // On non-HLSL targets, there isn't an implementation of `AppendStructuredBuffer`
+    // and `ConsumeStructuredBuffer` types, so we lower them into normal struct types
+    // of `RWStructuredBuffer` typed fields now.
+    if (target != CodeGenTarget::HLSL)
+    {
+        lowerAppendConsumeStructuredBuffers(targetRequest, irModule, sink);
+    }
+
     // We don't need the legalize pass for C/C++ based types
     if(options.shouldLegalizeExistentialAndResourceTypes )
     {
diff --git a/source/slang/slang-ir-byte-address-legalize.cpp b/source/slang/slang-ir-byte-address-legalize.cpp
index 40fe64693..b4de66d77 100644
--- a/source/slang/slang-ir-byte-address-legalize.cpp
+++ b/source/slang/slang-ir-byte-address-legalize.cpp
@@ -741,6 +741,8 @@ struct ByteAddressBufferLegalizationContext
         paramBuilder.setInsertBefore(byteAddressBufferParam);
 
         auto structuredBufferParam = paramBuilder.createGlobalParam(structuredBufferParamType);
+        if (auto nameHint = byteAddressBufferParam->findDecoration<IRNameHintDecoration>())
+            paramBuilder.addNameHintDecoration(structuredBufferParam, nameHint->getName());
 
         // The new parameter needs to be given a layout to match the existing
         // parameter, so that it is given the same `binding` in the generated code.
diff --git a/source/slang/slang-ir-inst-defs.h b/source/slang/slang-ir-inst-defs.h
index 22355bd7e..c1b021181 100644
--- a/source/slang/slang-ir-inst-defs.h
+++ b/source/slang/slang-ir-inst-defs.h
@@ -447,6 +447,14 @@ INST(RWStructuredBufferStore, rwstructuredBufferStore, 3, 0)
 
 INST(RWStructuredBufferGetElementPtr, rwstructuredBufferGetElementPtr, 2, 0)
 
+// Append/Consume-StructuredBuffer operations
+INST(StructuredBufferAppend, StructuredBufferAppend, 1, 0)
+INST(StructuredBufferConsume, StructuredBufferConsume, 1, 0)
+INST(StructuredBufferGetDimensions, StructuredBufferGetDimensions, 1, 0)
+
+INST(AtomicCounterIncrement, AtomicCounterIncrement, 1, 0)
+INST(AtomicCounterDecrement, AtomicCounterDecrement, 1, 0)
+
 INST(MeshOutputRef, meshOutputRef, 2, 0)
 
 // Construct a vector from a scalar
diff --git a/source/slang/slang-ir-insts.h b/source/slang/slang-ir-insts.h
index adfcac7fd..4b0cac182 100644
--- a/source/slang/slang-ir-insts.h
+++ b/source/slang/slang-ir-insts.h
@@ -2144,6 +2144,25 @@ struct IRRWStructuredBufferGetElementPtr : IRInst
     IRInst* getIndex() { return getOperand(1); }
 };
 
+struct IRStructuredBufferAppend : IRInst
+{
+    IR_LEAF_ISA(StructuredBufferAppend);
+    IRInst* getBuffer() { return getOperand(0); }
+    IRInst* getElement() { return getOperand(1); }
+};
+
+struct IRStructuredBufferConsume : IRInst
+{
+    IR_LEAF_ISA(StructuredBufferConsume);
+    IRInst* getBuffer() { return getOperand(0); }
+};
+
+struct IRStructuredBufferGetDimensions : IRInst
+{
+    IR_LEAF_ISA(StructuredBufferGetDimensions);
+    IRInst* getBuffer() { return getOperand(0); }
+};
+
 struct IRLoadReverseGradient : IRInst
 {
     IR_LEAF_ISA(LoadReverseGradient)
diff --git a/source/slang/slang-ir-lower-append-consume-structured-buffer.cpp b/source/slang/slang-ir-lower-append-consume-structured-buffer.cpp
new file mode 100644
index 000000000..fa9f16223
--- /dev/null
+++ b/source/slang/slang-ir-lower-append-consume-structured-buffer.cpp
@@ -0,0 +1,247 @@
+#include "slang-ir-lower-append-consume-structured-buffer.h"
+#include "slang-ir.h"
+#include "slang-ir-insts.h"
+#include "slang-ir-util.h"
+#include "slang-ir-layout.h"
+#include "slang-ir-lower-buffer-element-type.h"
+
+namespace Slang
+{
+    static void lowerStructuredBufferType(TargetRequest* target, IRHLSLStructuredBufferTypeBase* type)
+    {
+        IRBuilder builder(type);
+        builder.setInsertBefore(type);
+
+        auto elementType = type->getElementType();
+
+        // Type.
+        auto structType = builder.createStructType();
+        StringBuilder nameSb;
+        if (type->getOp() == kIROp_HLSLAppendStructuredBufferType)
+            nameSb << "AppendStructuredBuffer_";
+        else
+            nameSb << "ConsumeStructuredBuffer_";
+        getTypeNameHint(nameSb, elementType);
+        nameSb << "_t";
+        builder.addNameHintDecoration(structType, nameSb.produceString().getUnownedSlice());
+
+        auto elementBufferKey = builder.createStructKey();
+        builder.addNameHintDecoration(elementBufferKey, UnownedStringSlice("elements"));
+
+        auto counterBufferKey = builder.createStructKey();
+        builder.addNameHintDecoration(counterBufferKey, UnownedStringSlice("counter"));
+
+        auto elementBufferType = builder.getType(kIROp_HLSLRWStructuredBufferType, elementType);
+        auto counterBufferType = builder.getType(kIROp_HLSLRWStructuredBufferType, builder.getIntType());
+
+        builder.createStructField(structType, elementBufferKey, elementBufferType);
+        builder.createStructField(structType, counterBufferKey, counterBufferType);
+
+        // Type layout.
+        auto layoutRules = getTypeLayoutRuleForBuffer(target, type);
+
+        IRTypeLayout::Builder elementTypeLayoutBuilder(&builder);
+        IRSizeAndAlignment elementSize;
+        getSizeAndAlignment(layoutRules, elementType, &elementSize);
+        elementTypeLayoutBuilder.addResourceUsage(LayoutResourceKind::Uniform, LayoutSize((LayoutSize::RawValue)elementSize.getStride()));
+        auto elementTypeLayout = elementTypeLayoutBuilder.build();
+
+        IRStructuredBufferTypeLayout::Builder elementBufferTypeLayoutBuilder(&builder, elementTypeLayout);
+        elementBufferTypeLayoutBuilder.addResourceUsage(LayoutResourceKind::DescriptorTableSlot, 1);
+        auto elementBufferTypeLayout = elementBufferTypeLayoutBuilder.build();
+
+        IRTypeLayout::Builder counterTypeLayoutBuilder(&builder);
+        counterTypeLayoutBuilder.addResourceUsage(LayoutResourceKind::Uniform, LayoutSize(4));
+        auto counterTypeLayout = counterTypeLayoutBuilder.build();
+
+        IRStructuredBufferTypeLayout::Builder counterBufferTypeLayoutBuilder(&builder, counterTypeLayout);
+        counterBufferTypeLayoutBuilder.addResourceUsage(LayoutResourceKind::DescriptorTableSlot, 1);
+        auto counterBufferTypeLayout = counterBufferTypeLayoutBuilder.build();
+
+        IRVarLayout::Builder elementBufferVarLayoutBuilder(&builder, elementBufferTypeLayout);
+        elementBufferVarLayoutBuilder.findOrAddResourceInfo(LayoutResourceKind::DescriptorTableSlot)->offset = 0;
+        auto elementBufferVarLayout = elementBufferVarLayoutBuilder.build();
+
+        IRVarLayout::Builder counterBufferVarLayoutBuilder(&builder, counterBufferTypeLayout);
+        counterBufferVarLayoutBuilder.findOrAddResourceInfo(LayoutResourceKind::DescriptorTableSlot)->offset = 1;
+        auto counterBufferVarLayout = counterBufferVarLayoutBuilder.build();
+
+        IRStructTypeLayout::Builder layoutBuilder(&builder);
+        layoutBuilder.addField(elementBufferKey, elementBufferVarLayout);
+        layoutBuilder.addField(counterBufferKey, counterBufferVarLayout);
+        auto typeLayout = layoutBuilder.build();
+
+        builder.addLayoutDecoration(structType, typeLayout);
+
+        IRFunc* appendFunc = nullptr;
+        IRFunc* consumeFunc = nullptr;
+        IRFunc* getDimensionsFunc = nullptr;
+
+        if (type->getOp() == kIROp_HLSLAppendStructuredBufferType)
+        {
+            // Append method.
+            appendFunc = builder.createFunc();
+            builder.addNameHintDecoration(appendFunc, UnownedStringSlice("AppendStructuredBuffer_Append"));
+            IRType* paramTypes[] = { structType, elementType };
+            auto funcType = builder.getFuncType(2, paramTypes, builder.getVoidType());
+            appendFunc->setFullType(funcType);
+            builder.setInsertInto(appendFunc);
+            builder.emitBlock();
+            auto bufferParam = builder.emitParam(structType);
+            auto elementParam = builder.emitParam(elementType);
+            auto elementBuffer = builder.emitFieldExtract(elementBufferType, bufferParam, elementBufferKey);
+            auto counterBuffer = builder.emitFieldExtract(counterBufferType, bufferParam, counterBufferKey);
+            IRInst* getCounterPtrArgs[] = { counterBuffer, builder.getIntValue(builder.getIntType(), 0) };
+            auto counterBufferPtr = builder.emitIntrinsicInst(builder.getPtrType(builder.getIntType()), kIROp_RWStructuredBufferGetElementPtr, 2, getCounterPtrArgs);
+            auto oldCounter = builder.emitIntrinsicInst(builder.getIntType(), kIROp_AtomicCounterIncrement, 1, &counterBufferPtr);
+
+            IRInst* getElementPtrArgs[] = { elementBuffer, oldCounter };
+            auto elementBufferPtr = builder.emitIntrinsicInst(builder.getPtrType(elementType), kIROp_RWStructuredBufferGetElementPtr, 2, getElementPtrArgs);
+
+            builder.emitStore(elementBufferPtr, elementParam);
+            builder.emitReturn();
+        }
+        else
+        {
+            // Consume method.
+            consumeFunc = builder.createFunc();
+            builder.addNameHintDecoration(consumeFunc, UnownedStringSlice("ConsumeStructuredBuffer_Consume"));
+            IRType* paramTypes[] = { structType };
+            auto funcType = builder.getFuncType(1, paramTypes, elementType);
+            consumeFunc->setFullType(funcType);
+            builder.setInsertInto(consumeFunc);
+            auto firstBlock = builder.emitBlock();
+            auto bufferParam = builder.emitParam(structType);
+            auto elementBuffer = builder.emitFieldExtract(elementBufferType, bufferParam, elementBufferKey);
+            auto counterBuffer = builder.emitFieldExtract(counterBufferType, bufferParam, counterBufferKey);
+            IRInst* getCounterPtrArgs[] = { counterBuffer, builder.getIntValue(builder.getIntType(), 0) };
+            auto counterBufferPtr = builder.emitIntrinsicInst(builder.getPtrType(builder.getIntType()), kIROp_RWStructuredBufferGetElementPtr, 2, getCounterPtrArgs);
+            auto oldCounter = builder.emitIntrinsicInst(builder.getIntType(), kIROp_AtomicCounterDecrement, 1, &counterBufferPtr);
+            auto index = builder.emitSub(builder.getIntType(), oldCounter, builder.getIntValue(builder.getIntType(), 1));
+
+            // Test if index is greater or equal than 0.
+            auto geq = builder.emitGeq(index, builder.getIntValue(builder.getIntType(), 0));
+            auto trueBlock = builder.emitBlock();
+
+            auto falseBlock = builder.emitBlock();
+            auto mergeBlock = builder.emitBlock();
+
+            builder.setInsertInto(firstBlock);
+            builder.emitIfElse(geq, trueBlock, falseBlock, mergeBlock);
+
+            builder.setInsertInto(trueBlock);
+            IRInst* getElementPtrArgs[] = { elementBuffer, index };
+            auto elementBufferPtr = builder.emitIntrinsicInst(builder.getPtrType(elementType), kIROp_RWStructuredBufferGetElementPtr, 2, getElementPtrArgs);
+            auto val = builder.emitLoad(elementBufferPtr);
+            builder.emitReturn(val);
+
+            builder.setInsertInto(falseBlock);
+            auto defaultVal = builder.emitDefaultConstruct(elementType);
+            builder.emitReturn(defaultVal);
+
+            builder.setInsertInto(mergeBlock);
+            builder.emitUnreachable();
+        }
+
+        // GetDimensions method.
+        {
+            getDimensionsFunc = builder.createFunc();
+            builder.addNameHintDecoration(getDimensionsFunc, UnownedStringSlice("StructuredBuffer_GetDimensions"));
+            IRType* paramTypes[] = { structType };
+            auto uint2Type = builder.getVectorType(builder.getUIntType(), 2);
+            auto funcType = builder.getFuncType(1, paramTypes, uint2Type);
+            getDimensionsFunc->setFullType(funcType);
+            builder.setInsertInto(getDimensionsFunc);
+            builder.emitBlock();
+            auto bufferParam = builder.emitParam(structType);
+            auto counterBuffer = builder.emitFieldExtract(counterBufferType, bufferParam, counterBufferKey);
+            IRInst* getCounterPtrArgs[] = { counterBuffer, builder.getIntValue(builder.getIntType(), 0) };
+            auto counterBufferPtr = builder.emitIntrinsicInst(builder.getPtrType(builder.getIntType()), kIROp_RWStructuredBufferGetElementPtr, 2, getCounterPtrArgs);
+            auto counter = builder.emitLoad(counterBufferPtr);
+            counter = builder.emitCast(builder.getUIntType(), counter);
+            auto stride = builder.getIntValue(builder.getUIntType(), elementSize.getStride());
+            IRInst* vecArgs[] = { counter, stride };
+            builder.emitReturn(builder.emitMakeVector(uint2Type, 2, vecArgs));
+        }
+
+        // Replace all insts with synthesized functions.
+        traverseUsers(type, [&](IRInst* typeUser)
+            {
+                if (typeUser->getFullType() != type)
+                    return;
+                if (auto layoutDecor = typeUser->findDecoration<IRLayoutDecoration>())
+                {
+                    // Replace the original StructuredBufferVarLayout with the new StructTypeVarLayout.
+                    if (auto varLayout = as<IRVarLayout>(layoutDecor->getLayout()))
+                    {
+                        IRBuilder subBuilder(typeUser);
+                        IRVarLayout::Builder newVarLayoutBuilder(&subBuilder, typeLayout);
+                        newVarLayoutBuilder.cloneEverythingButOffsetsFrom(varLayout);
+                        for (auto offsetAttr : varLayout->getOffsetAttrs())
+                        {
+                            auto info = newVarLayoutBuilder.findOrAddResourceInfo(offsetAttr->getResourceKind());
+                            info->offset = offsetAttr->getOffset();
+                            info->space = offsetAttr->getSpace();
+                            info->kind = offsetAttr->getResourceKind();
+                        }
+                        auto newVarLayout = newVarLayoutBuilder.build();
+                        subBuilder.addLayoutDecoration(typeUser, newVarLayout);
+                        varLayout->removeAndDeallocate();
+                    }
+                }
+                traverseUses(typeUser, [&](IRUse* use)
+                    {
+                        auto user = use->getUser();
+                        switch (user->getOp())
+                        {
+                        case kIROp_StructuredBufferAppend:
+                            {
+                                IRBuilder subBuilder(user);
+                                subBuilder.setInsertBefore(user);
+                                IRInst* args[] = { user->getOperand(0), user->getOperand(1) };
+                                auto call = subBuilder.emitCallInst(user->getFullType(), appendFunc, 2, args);
+                                user->replaceUsesWith(call);
+                                user->removeAndDeallocate();
+                                break;
+                            }
+                        case kIROp_StructuredBufferConsume:
+                            {
+                                IRBuilder subBuilder(user);
+                                subBuilder.setInsertBefore(user);
+                                IRInst* args[] = { user->getOperand(0) };
+                                auto call = subBuilder.emitCallInst(user->getFullType(), consumeFunc, 1, args);
+                                user->replaceUsesWith(call);
+                                user->removeAndDeallocate();
+                                break;
+                            }
+                        case kIROp_StructuredBufferGetDimensions:
+                            {
+                                IRBuilder subBuilder(user);
+                                subBuilder.setInsertBefore(user);
+                                IRInst* args[] = { user->getOperand(0) };
+                                auto call = subBuilder.emitCallInst(user->getFullType(), getDimensionsFunc, 1, args);
+                                user->replaceUsesWith(call);
+                                user->removeAndDeallocate();
+                                break;
+                            }
+                        }
+                    });
+            });
+        type->replaceUsesWith(structType);
+    }
+
+    void lowerAppendConsumeStructuredBuffers(TargetRequest* target, IRModule* module, DiagnosticSink* sink)
+    {
+        SLANG_UNUSED(sink);
+        for (auto globalInst : module->getGlobalInsts())
+        {
+            switch (globalInst->getOp())
+            {
+            case kIROp_HLSLAppendStructuredBufferType:
+            case kIROp_HLSLConsumeStructuredBufferType:
+                lowerStructuredBufferType(target, as<IRHLSLStructuredBufferTypeBase>(globalInst));
+                break;
+            }
+        }
+    }
+}
diff --git a/source/slang/slang-ir-lower-append-consume-structured-buffer.h b/source/slang/slang-ir-lower-append-consume-structured-buffer.h
new file mode 100644
index 000000000..81048724d
--- /dev/null
+++ b/source/slang/slang-ir-lower-append-consume-structured-buffer.h
@@ -0,0 +1,17 @@
+// slang-ir-lower-append-consume-structured-buffer.h
+#pragma once
+
+#include "slang-ir.h"
+
+namespace Slang
+{
+    struct IRModule;
+    class DiagnosticSink;
+    class TargetRequest;
+
+    /// For non-hlsl targets, lower append- and consume- structured buffers into `struct` types
+    /// that contains two RWStructuredBuffer typed fields, one to store the elements, and one
+    /// for the atomic buffer.
+    void lowerAppendConsumeStructuredBuffers(TargetRequest* target, IRModule* module, DiagnosticSink* sink);
+
+}
diff --git a/source/slang/slang-type-layout.cpp b/source/slang/slang-type-layout.cpp
index cdf1f3694..978fa6fbb 100644
--- a/source/slang/slang-type-layout.cpp
+++ b/source/slang/slang-type-layout.cpp
@@ -715,6 +715,7 @@ static LayoutResourceKind _getHLSLLayoutResourceKind(ShaderParameterKind kind)
         case ShaderParameterKind::MutableRawBuffer:
         case ShaderParameterKind::MutableBuffer:
         case ShaderParameterKind::MutableTexture:
+        case ShaderParameterKind::AppendConsumeStructuredBuffer:
             return LayoutResourceKind::UnorderedAccess;
 
         case ShaderParameterKind::SamplerState:
@@ -728,6 +729,13 @@ struct GLSLObjectLayoutRulesImpl : ObjectLayoutRulesImpl
 {
     virtual SimpleLayoutInfo GetObjectLayout(ShaderParameterKind kind, const Options& options) override
     {
+        int slotCount = 1;
+
+        // In Vulkan GLSL, pretty much every object is just a descriptor-table slot.
+        // Except for AppendConsumeStructuredBuffer, which takes two slots.
+        if (kind == ShaderParameterKind::AppendConsumeStructuredBuffer)
+            slotCount = 2;
+
         if (options.hlslToVulkanKindFlags)
         {
             // Is this an HLSL kind that might be shifted
@@ -745,14 +753,12 @@ struct GLSLObjectLayoutRulesImpl : ObjectLayoutRulesImpl
                 {
                     // We are going to consume a HLSL layout kind
                     // Later we will do shifting as necessary
-                    return SimpleLayoutInfo(hlslLayoutKind, 1);
+                    return SimpleLayoutInfo(hlslLayoutKind, slotCount);
                 }
             }
         }
 
-        // In Vulkan GLSL, pretty much every object is just a descriptor-table slot.
-        // We can refine this method once we support a case where this isn't true.
-        return SimpleLayoutInfo(LayoutResourceKind::DescriptorTableSlot, 1);
+        return SimpleLayoutInfo(LayoutResourceKind::DescriptorTableSlot, slotCount);
     }
 };
 GLSLObjectLayoutRulesImpl kGLSLObjectLayoutRulesImpl;
@@ -799,6 +805,7 @@ struct HLSLObjectLayoutRulesImpl : ObjectLayoutRulesImpl
         case ShaderParameterKind::MutableRawBuffer:
         case ShaderParameterKind::MutableBuffer:
         case ShaderParameterKind::MutableTexture:
+        case ShaderParameterKind::AppendConsumeStructuredBuffer:
             return SimpleLayoutInfo(LayoutResourceKind::UnorderedAccess, 1);
 
         case ShaderParameterKind::SamplerState:
@@ -974,6 +981,7 @@ struct CPUObjectLayoutRulesImpl : ObjectLayoutRulesImpl
                 
             case ShaderParameterKind::StructuredBuffer:            
             case ShaderParameterKind::MutableStructuredBuffer:
+            case ShaderParameterKind::AppendConsumeStructuredBuffer:
                 // It's a ptr and a size of the amount of elements
                 return SimpleLayoutInfo(LayoutResourceKind::Uniform, sizeof(void*) * 2, SLANG_ALIGN_OF(void*));
 
@@ -1033,6 +1041,7 @@ struct CUDAObjectLayoutRulesImpl : CPUObjectLayoutRulesImpl
 
             case ShaderParameterKind::StructuredBuffer:
             case ShaderParameterKind::MutableStructuredBuffer:
+            case ShaderParameterKind::AppendConsumeStructuredBuffer:
             {
                 // It's a ptr and a count of the amount of elements
                 const size_t size = _roundToAlignment(sizeof(CUDAPtr) + sizeof(CUDACount), sizeof(CUDAPtr));
@@ -3763,8 +3772,8 @@ static TypeLayoutResult _createTypeLayout(
     CASE(HLSLStructuredBufferType,                  StructuredBuffer);
     CASE(HLSLRWStructuredBufferType,                MutableStructuredBuffer);
     CASE(HLSLRasterizerOrderedStructuredBufferType, MutableStructuredBuffer);
-    CASE(HLSLAppendStructuredBufferType,            MutableStructuredBuffer);
-    CASE(HLSLConsumeStructuredBufferType,           MutableStructuredBuffer);
+    CASE(HLSLAppendStructuredBufferType,            AppendConsumeStructuredBuffer);
+    CASE(HLSLConsumeStructuredBufferType,           AppendConsumeStructuredBuffer);
 
 #undef CASE
 
diff --git a/source/slang/slang-type-layout.h b/source/slang/slang-type-layout.h
index af07f3e73..e3dd719d6 100644
--- a/source/slang/slang-type-layout.h
+++ b/source/slang/slang-type-layout.h
@@ -935,6 +935,8 @@ enum class ShaderParameterKind
     MutableImage,
 
     RegisterSpace,
+
+    AppendConsumeStructuredBuffer,
 };
 
 struct SimpleLayoutRulesImpl
diff --git a/tests/bugs/texture2d-ms.hlsl.glsl b/tests/bugs/texture2d-ms.hlsl.glsl
index 40ce5f9de..013d7d7bd 100644
--- a/tests/bugs/texture2d-ms.hlsl.glsl
+++ b/tests/bugs/texture2d-ms.hlsl.glsl
@@ -5,7 +5,7 @@ layout(row_major) buffer;
 layout(binding = 0)
 uniform texture2DMS tex_0;
 
-layout(std430, binding = 1) buffer _S1 {
+layout(std430, binding = 1) buffer StructuredBuffer_float4_t_0 {
     vec4 _data[];
 } outBuffer_0;
 layout(local_size_x = 4, local_size_y = 4, local_size_z = 1) in;
diff --git a/tests/bugs/vk-structured-buffer-binding.hlsl.glsl b/tests/bugs/vk-structured-buffer-binding.hlsl.glsl
index 7298ea594..f108aac00 100644
--- a/tests/bugs/vk-structured-buffer-binding.hlsl.glsl
+++ b/tests/bugs/vk-structured-buffer-binding.hlsl.glsl
@@ -2,25 +2,23 @@
 //TEST_IGNORE_FILE:
 
 #version 450
+layout(row_major) uniform;
+layout(row_major) buffer;
 
-#define gDoneGroups gDoneGroups_0
-#define uv _S3
-#define SV_Target _S2
-
-layout(std430, binding = 3, set = 4)
-buffer _S1
-{
+layout(std430, binding = 3, set = 4) buffer StructuredBuffer_uint_t_0 {
     uint _data[];
-} gDoneGroups;
+} gDoneGroups_0;
 
 layout(location = 0)
-out vec4 SV_Target;
+out vec4 _S1;
 
 layout(location = 0)
-in vec3 uv;
+in vec3 _S2;
 
 void main()
 {
-    SV_Target = vec4(gDoneGroups._data[uint(int(uv.z))]);
+    _S1 = vec4(float(gDoneGroups_0._data[uint(int(_S2.z))]));
     return;
 }
+
+
diff --git a/tests/bugs/vk-structured-buffer-load.hlsl.glsl b/tests/bugs/vk-structured-buffer-load.hlsl.glsl
index 0184e6ed5..35fad779b 100644
--- a/tests/bugs/vk-structured-buffer-load.hlsl.glsl
+++ b/tests/bugs/vk-structured-buffer-load.hlsl.glsl
@@ -2,9 +2,11 @@
 #extension GL_NV_ray_tracing : require
 layout(row_major) uniform;
 layout(row_major) buffer;
-layout(std430, binding = 1) readonly buffer _S1 {
+
+layout(std430, binding = 1) readonly buffer StructuredBuffer_float_t_0 {
     float _data[];
 } gParamBlock_sbuf_0;
+
 float rcp_0(float x_0)
 {
     return 1.0 / x_0;
@@ -15,35 +17,37 @@ struct RayHitInfoPacked_0
     vec4 PackedHitInfoA_0;
 };
 
-rayPayloadInNV RayHitInfoPacked_0 _S2;
+rayPayloadInNV RayHitInfoPacked_0 _S1;
 
 struct BuiltInTriangleIntersectionAttributes_0
 {
     vec2 barycentrics_0;
 };
 
-hitAttributeNV BuiltInTriangleIntersectionAttributes_0 _S3;
+hitAttributeNV BuiltInTriangleIntersectionAttributes_0 _S2;
 
 void main()
 {
     float HitT_0 = ((gl_RayTmaxNV));
-    _S2.PackedHitInfoA_0.x = HitT_0;
+    _S1.PackedHitInfoA_0[0] = HitT_0;
+
+    float offsfloat_0 = gParamBlock_sbuf_0._data[0];
 
-    float offsfloat_0 = ((gParamBlock_sbuf_0)._data[(0)]);
     uint use_rcp_0 = 0U | uint(HitT_0 > 0.0);
+
     if(use_rcp_0 != 0U)
     {
-        _S2.PackedHitInfoA_0.y = rcp_0(offsfloat_0);
+        _S1.PackedHitInfoA_0[1] = rcp_0(offsfloat_0);
     }
     else
     {
         if(use_rcp_0 > 0U&&offsfloat_0 == 0.0)
         {
-            _S2.PackedHitInfoA_0.y = (inversesqrt((offsfloat_0 + 1.0)));
+            _S1.PackedHitInfoA_0[1] = (inversesqrt((offsfloat_0 + 1.0)));
         }
         else
         {
-            _S2.PackedHitInfoA_0.y = (inversesqrt((offsfloat_0)));
+            _S1.PackedHitInfoA_0[1] = (inversesqrt((offsfloat_0)));
         }
     }
     return;
diff --git a/tests/compute/comma-operator.slang.glsl b/tests/compute/comma-operator.slang.glsl
index af7120704..236029b47 100644
--- a/tests/compute/comma-operator.slang.glsl
+++ b/tests/compute/comma-operator.slang.glsl
@@ -3,7 +3,7 @@
 
 //TEST_IGNORE_FILE:
 
-layout(std430, binding = 0) buffer _S1 {
+layout(std430, binding = 0) buffer StructuredBuffer_int_t_0 {
     int _data[];
 } outputBuffer_0;
 
diff --git a/tests/compute/half-texture.slang.glsl b/tests/compute/half-texture.slang.glsl
index 27f63620d..16ff4d6e1 100644
--- a/tests/compute/half-texture.slang.glsl
+++ b/tests/compute/half-texture.slang.glsl
@@ -17,7 +17,7 @@ layout(rgba16f)
 layout(binding = 3)
 uniform image2D halfTexture4_0;
 
-layout(std430, binding = 0) buffer _S1 {
+layout(std430, binding = 0) buffer StructuredBuffer_int_t_0 {
     int _data[];
 } outputBuffer_0;
 
diff --git a/tests/compute/unbounded-array-of-array-syntax.slang.glsl b/tests/compute/unbounded-array-of-array-syntax.slang.glsl
index 73148ae10..2ea90e4a6 100644
--- a/tests/compute/unbounded-array-of-array-syntax.slang.glsl
+++ b/tests/compute/unbounded-array-of-array-syntax.slang.glsl
@@ -1,28 +1,23 @@
 //TEST_IGNORE_FILE:
-
 #version 450
 #extension GL_EXT_nonuniform_qualifier : require
 layout(row_major) uniform;
 layout(row_major) buffer;
-
-layout(std430, binding = 1) buffer _S1 {
+layout(std430, binding = 1) buffer StructuredBuffer_int_t_0 {
     int _data[];
 } g_aoa_0[];
-layout(std430, binding = 0) buffer _S2 {
+layout(std430, binding = 0) buffer StructuredBuffer_int_t_1 {
     int _data[];
 } outputBuffer_0;
-
 layout(local_size_x = 8, local_size_y = 1, local_size_z = 1) in;
 void main()
 {
     int index_0 = int(gl_GlobalInvocationID.x);
     int innerIndex_0 = index_0 & 3;
-
-    int _S3 = nonuniformEXT(index_0 >> 2);
+    int _S1 = nonuniformEXT(index_0 >> 2);
     uint bufferCount_0;
     uint bufferStride_0;
-    (bufferCount_0) = (g_aoa_0[_S3])._data.length(); (bufferStride_0) = 0;
-
+    (bufferCount_0) = (g_aoa_0[_S1])._data.length(); (bufferStride_0) = 0;
     int innerIndex_1;
     if(innerIndex_0 >= int(bufferCount_0))
     {
@@ -32,6 +27,6 @@ void main()
     {
         innerIndex_1 = innerIndex_0;
     }
-    outputBuffer_0._data[uint(index_0)] = g_aoa_0[_S3]._data[uint(innerIndex_1)];
+    outputBuffer_0._data[uint(index_0)] = g_aoa_0[_S1]._data[uint(innerIndex_1)];
     return;
 }
diff --git a/tests/cross-compile/array-of-buffers.slang.glsl b/tests/cross-compile/array-of-buffers.slang.glsl
index a198a5277..fb1a4be2f 100644
--- a/tests/cross-compile/array-of-buffers.slang.glsl
+++ b/tests/cross-compile/array-of-buffers.slang.glsl
@@ -10,8 +10,7 @@ layout(binding = 0)
 layout(std140) uniform _S1
 {
     uint index_0;
-} C_0;
-
+}C_0;
 struct S_0
 {
     vec4 f_0;
@@ -21,23 +20,25 @@ layout(binding = 1)
 layout(std140) uniform _S2
 {
     vec4 f_0;
-} cb_0[3];
-layout(std430, binding = 2) readonly buffer _S3 {
+}cb_0[3];
+layout(std430, binding = 2) readonly buffer StructuredBuffer_S_t_0 {
     S_0 _data[];
 } sb1_0[4];
-layout(std430, binding = 3) buffer _S4 {
+layout(std430, binding = 3) buffer StructuredBuffer_float4_t_0 {
     vec4 _data[];
 } sb2_0[5];
-layout(std430, binding = 4) readonly buffer _S5
+layout(std430, binding = 4) readonly buffer _S3
 {
     uint _data[];
 } bb_0[6];
 layout(location = 0)
-out vec4 _S6;
+out vec4 _S4;
 
 void main()
 {
-    uint _S7 = ((bb_0[C_0.index_0])._data[(int(C_0.index_0 * 4U))/4]);
-    _S6 = cb_0[C_0.index_0].f_0 + sb1_0[C_0.index_0]._data[C_0.index_0].f_0 + sb2_0[C_0.index_0]._data[C_0.index_0] + vec4(float(_S7));
+
+    uint _S5 = ((bb_0[C_0.index_0])._data[(int(C_0.index_0 * 4U))/4]);
+    _S4 = cb_0[C_0.index_0].f_0 + sb1_0[C_0.index_0]._data[C_0.index_0].f_0 + sb2_0[C_0.index_0]._data[C_0.index_0] + vec4(float(_S5));
     return;
 }
+
diff --git a/tests/cross-compile/func-resource-param-array.slang.glsl b/tests/cross-compile/func-resource-param-array.slang.glsl
index 9e396e55f..0c4c44bef 100644
--- a/tests/cross-compile/func-resource-param-array.slang.glsl
+++ b/tests/cross-compile/func-resource-param-array.slang.glsl
@@ -1,37 +1,33 @@
 #version 450
 layout(row_major) uniform;
 layout(row_major) buffer;
-
-layout(std430, binding = 0) buffer _S1 {
+layout(std430, binding = 0) buffer StructuredBuffer_int_t_0 {
     int _data[];
 } a_0;
-
-layout(std430, binding = 1) buffer _S2 {
+layout(std430, binding = 1) buffer StructuredBuffer_int_t_1 {
     int _data[];
 } b_0[3];
-
-layout(std430, binding = 2) buffer _S3 {
+layout(std430, binding = 2) buffer StructuredBuffer_int_t_2 {
     int _data[];
 } c_0[4][3];
-
-int f_0(uint _S4)
+int f_0(uint _S1)
 {
-    return a_0._data[_S4];
+    return a_0._data[_S1];
 }
 
-int f_1(uint _S5, uint _S6)
+int f_1(uint _S2, uint _S3)
 {
-    return b_0[_S5]._data[_S6];
+    return b_0[_S2]._data[_S3];
 }
 
-int g_0(uint _S7, uint _S8)
+int g_0(uint _S4, uint _S5)
 {
-    return b_0[_S7]._data[_S8];
+    return b_0[_S4]._data[_S5];
 }
 
-int g_1(uint _S9, uint _S10, uint _S11)
+int g_1(uint _S6, uint _S7, uint _S8)
 {
-    return c_0[_S9][_S10]._data[_S11];
+    return c_0[_S6][_S7]._data[_S8];
 }
 
 layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
diff --git a/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-assign.slang.1.expected b/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-assign.slang.1.expected
index 4338e70ac..2bb288e11 100644
--- a/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-assign.slang.1.expected
+++ b/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-assign.slang.1.expected
@@ -7,7 +7,7 @@ standard output = {
 #extension GL_NV_shader_invocation_reorder : require
 layout(row_major) uniform;
 layout(row_major) buffer;
-layout(std430, binding = 0) buffer _S1 {
+layout(std430, binding = 0) buffer StructuredBuffer_uint_t_0 {
     uint _data[];
 } outputBuffer_0;
 struct RayDesc_0
@@ -20,23 +20,23 @@ struct RayDesc_0
 
 void main()
 {
-    uvec3 _S2 = ((gl_LaunchIDEXT));
-    ivec2 launchID_0 = ivec2(_S2.xy);
-    uvec3 _S3 = ((gl_LaunchSizeEXT));
+    uvec3 _S1 = ((gl_LaunchIDEXT));
+    ivec2 launchID_0 = ivec2(_S1.xy);
+    uvec3 _S2 = ((gl_LaunchSizeEXT));
     int idx_0 = launchID_0.x;
     RayDesc_0 ray_0;
     ray_0.Origin_0 = vec3(float(idx_0), 0.0, 0.0);
     ray_0.TMin_0 = 0.00999999977648258209;
     ray_0.Direction_0 = vec3(0.0, 1.0, 0.0);
     ray_0.TMax_0 = 10000.0;
-    uint _S4 = uint(idx_0);
+    uint _S3 = uint(idx_0);
     hitObjectNV hitObj_0;
-    hitObjectRecordMissNV(hitObj_0, _S4, ray_0.Origin_0, ray_0.TMin_0, ray_0.Direction_0, ray_0.TMax_0);
-    uint _S5 = uint(idx_0 + 1);
+    hitObjectRecordMissNV(hitObj_0, _S3, ray_0.Origin_0, ray_0.TMin_0, ray_0.Direction_0, ray_0.TMax_0);
+    uint _S4 = uint(idx_0 + 1);
     hitObjectNV hitObj_1;
-    hitObjectRecordMissNV(hitObj_1, _S5, ray_0.Origin_0, ray_0.TMin_0, ray_0.Direction_0, ray_0.TMax_0);
-    bool _S6 = (hitObjectIsMissNV((hitObj_1)));
-    outputBuffer_0._data[_S4] = uint(int(_S6));
+    hitObjectRecordMissNV(hitObj_1, _S4, ray_0.Origin_0, ray_0.TMin_0, ray_0.Direction_0, ray_0.TMax_0);
+    bool _S5 = (hitObjectIsMissNV((hitObj_1)));
+    outputBuffer_0._data[_S3] = uint(int(_S5));
     return;
 }
 
diff --git a/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-make-hit.slang.1.expected b/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-make-hit.slang.1.expected
index 199fdc86b..12922d994 100644
--- a/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-make-hit.slang.1.expected
+++ b/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-make-hit.slang.1.expected
@@ -10,7 +10,7 @@ layout(row_major) buffer;
 layout(binding = 0)
 uniform accelerationStructureEXT scene_0;
 
-layout(std430, binding = 1) buffer _S1 {
+layout(std430, binding = 1) buffer StructuredBuffer_uint_t_0 {
     uint _data[];
 } outputBuffer_0;
 struct SomeValues_0
@@ -33,11 +33,11 @@ struct RayDesc_0
 
 RayDesc_0 HitObject_GetRayDesc_0(hitObjectNV this_0)
 {
-    vec3 _S2 = (hitObjectGetWorldRayOriginNV((this_0)));
-    float _S3 = (hitObjectGetRayTMinNV((this_0)));
-    vec3 _S4 = (hitObjectGetObjectRayDirectionNV((this_0)));
-    float _S5 = (hitObjectGetRayTMaxNV((this_0)));
-    RayDesc_0 ray_0 = { _S2, _S3, _S4, _S5 };
+    vec3 _S1 = (hitObjectGetWorldRayOriginNV((this_0)));
+    float _S2 = (hitObjectGetRayTMinNV((this_0)));
+    vec3 _S3 = (hitObjectGetObjectRayDirectionNV((this_0)));
+    float _S4 = (hitObjectGetRayTMaxNV((this_0)));
+    RayDesc_0 ray_0 = { _S1, _S2, _S3, _S4 };
     return ray_0;
 }
 
@@ -49,9 +49,9 @@ SomeValues_0 HitObject_GetAttributes_0(hitObjectNV this_1)
 
 uint calcValue_0(hitObjectNV hit_0)
 {
-    bool _S6 = (hitObjectIsHitNV((hit_0)));
+    bool _S5 = (hitObjectIsHitNV((hit_0)));
     uint r_0;
-    if(_S6)
+    if(_S5)
     {
         uint instanceIndex_0 = (hitObjectGetInstanceCustomIndexNV((hit_0)));
         uint instanceID_0 = (hitObjectGetInstanceIdNV((hit_0)));
@@ -66,8 +66,8 @@ uint calcValue_0(hitObjectNV hit_0)
     }
     else
     {
-        bool _S7 = (hitObjectIsMissNV((hit_0)));
-        if(_S7)
+        bool _S6 = (hitObjectIsMissNV((hit_0)));
+        if(_S6)
         {
             r_0 = 1U;
         }
@@ -81,30 +81,30 @@ uint calcValue_0(hitObjectNV hit_0)
 
 void main()
 {
-    uvec3 _S8 = ((gl_LaunchIDEXT));
-    ivec2 launchID_0 = ivec2(_S8.xy);
-    uvec3 _S9 = ((gl_LaunchSizeEXT));
+    uvec3 _S7 = ((gl_LaunchIDEXT));
+    ivec2 launchID_0 = ivec2(_S7.xy);
+    uvec3 _S8 = ((gl_LaunchSizeEXT));
     int idx_0 = launchID_0.x;
     RayDesc_0 ray_2;
     ray_2.Origin_0 = vec3(float(idx_0), 0.0, 0.0);
     ray_2.TMin_0 = 0.00999999977648258209;
     ray_2.Direction_0 = vec3(0.0, 1.0, 0.0);
     ray_2.TMax_0 = 10000.0;
-    uint _S10 = uint(idx_0);
-    uint _S11 = uint(idx_0 * 2);
-    uint _S12 = uint(idx_0 * 3);
-    RayDesc_0 _S13 = ray_2;
+    uint _S9 = uint(idx_0);
+    uint _S10 = uint(idx_0 * 2);
+    uint _S11 = uint(idx_0 * 3);
+    RayDesc_0 _S12 = ray_2;
     hitObjectNV hitObj_0;
+    int _S13 = int(_S9);
     int _S14 = int(_S10);
     int _S15 = int(_S11);
-    int _S16 = int(_S12);
-    hitObjectRecordHitWithIndexNV(hitObj_0, scene_0, _S14, _S15, _S16, 0U, 0U, _S13.Origin_0, _S13.TMin_0, _S13.Direction_0, _S13.TMax_0, (0));
+    hitObjectRecordHitWithIndexNV(hitObj_0, scene_0, _S13, _S14, _S15, 0U, 0U, _S12.Origin_0, _S12.TMin_0, _S12.Direction_0, _S12.TMax_0, (0));
     uint r_3 = calcValue_0(hitObj_0);
-    RayDesc_0 _S17 = ray_2;
+    RayDesc_0 _S16 = ray_2;
     hitObjectNV hitObj_1;
-    hitObjectRecordHitNV(hitObj_1, scene_0, _S14, _S16, _S15, 0U, 0U, 4U, _S17.Origin_0, _S17.TMin_0, _S17.Direction_0, _S17.TMax_0, (0));
-    uint _S18 = calcValue_0(hitObj_1);
-    outputBuffer_0._data[_S10] = r_3 + _S18;
+    hitObjectRecordHitNV(hitObj_1, scene_0, _S13, _S15, _S14, 0U, 0U, 4U, _S16.Origin_0, _S16.TMin_0, _S16.Direction_0, _S16.TMax_0, (0));
+    uint _S17 = calcValue_0(hitObj_1);
+    outputBuffer_0._data[_S9] = r_3 + _S17;
     return;
 }
 
diff --git a/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-make-miss.slang.1.expected b/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-make-miss.slang.1.expected
index ed80f38d3..f74c8d5e9 100644
--- a/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-make-miss.slang.1.expected
+++ b/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-make-miss.slang.1.expected
@@ -7,7 +7,7 @@ standard output = {
 #extension GL_NV_shader_invocation_reorder : require
 layout(row_major) uniform;
 layout(row_major) buffer;
-layout(std430, binding = 0) buffer _S1 {
+layout(std430, binding = 0) buffer StructuredBuffer_uint_t_0 {
     uint _data[];
 } outputBuffer_0;
 struct RayDesc_0
@@ -27,11 +27,11 @@ void main()
     ray_0.TMin_0 = 0.00999999977648258209;
     ray_0.Direction_0 = vec3(0.0, 1.0, 0.0);
     ray_0.TMax_0 = 10000.0;
-    uint _S2 = uint(idx_0);
+    uint _S1 = uint(idx_0);
     hitObjectNV hitObj_0;
-    hitObjectRecordMissNV(hitObj_0, _S2, ray_0.Origin_0, ray_0.TMin_0, ray_0.Direction_0, ray_0.TMax_0);
-    bool _S3 = (hitObjectIsMissNV((hitObj_0)));
-    outputBuffer_0._data[_S2] = uint(int(_S3));
+    hitObjectRecordMissNV(hitObj_0, _S1, ray_0.Origin_0, ray_0.TMin_0, ray_0.Direction_0, ray_0.TMax_0);
+    bool _S2 = (hitObjectIsMissNV((hitObj_0)));
+    outputBuffer_0._data[_S1] = uint(int(_S2));
     return;
 }
 
diff --git a/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-make-nop.slang.1.expected b/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-make-nop.slang.1.expected
index c8136645c..2522ee2e5 100644
--- a/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-make-nop.slang.1.expected
+++ b/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-make-nop.slang.1.expected
@@ -7,20 +7,20 @@ standard output = {
 #extension GL_NV_shader_invocation_reorder : require
 layout(row_major) uniform;
 layout(row_major) buffer;
-layout(std430, binding = 0) buffer _S1 {
+layout(std430, binding = 0) buffer StructuredBuffer_uint_t_0 {
     uint _data[];
 } outputBuffer_0;
 void main()
 {
-    uvec3 _S2 = ((gl_LaunchIDEXT));
-    ivec2 launchID_0 = ivec2(_S2.xy);
-    uvec3 _S3 = ((gl_LaunchSizeEXT));
+    uvec3 _S1 = ((gl_LaunchIDEXT));
+    ivec2 launchID_0 = ivec2(_S1.xy);
+    uvec3 _S2 = ((gl_LaunchSizeEXT));
     int idx_0 = launchID_0.x;
     hitObjectNV hitObj_0;
     hitObjectRecordEmptyNV((hitObj_0));
-    uint _S4 = uint(idx_0);
-    bool _S5 = (hitObjectIsEmptyNV((hitObj_0)));
-    outputBuffer_0._data[_S4] = uint(_S5);
+    uint _S3 = uint(idx_0);
+    bool _S4 = (hitObjectIsEmptyNV((hitObj_0)));
+    outputBuffer_0._data[_S3] = uint(_S4);
     return;
 }
 
diff --git a/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-output.slang.1.expected b/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-output.slang.1.expected
index 2340caa21..f1c52d1ad 100644
--- a/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-output.slang.1.expected
+++ b/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-output.slang.1.expected
@@ -10,7 +10,7 @@ layout(row_major) buffer;
 layout(binding = 0)
 uniform accelerationStructureEXT scene_0;
 
-layout(std430, binding = 1) buffer _S1 {
+layout(std430, binding = 1) buffer StructuredBuffer_uint_t_0 {
     uint _data[];
 } outputBuffer_0;
 struct MyAttributes_0
@@ -58,19 +58,19 @@ MyAttributes_0 HitObject_GetAttributes_0(hitObjectNV this_0)
 void accumulate_0(inout uint value_2, hitObjectNV hit_0)
 {
     value_2 = value_2 * 256U;
-    bool _S2 = (hitObjectIsHitNV((hit_0)));
-    if(_S2)
+    bool _S1 = (hitObjectIsHitNV((hit_0)));
+    if(_S1)
     {
-        MyAttributes_0 _S3 = HitObject_GetAttributes_0(hit_0);
-        value_2 = value_2 + (16U + _S3.value_0);
+        MyAttributes_0 _S2 = HitObject_GetAttributes_0(hit_0);
+        value_2 = value_2 + (16U + _S2.value_0);
     }
     return;
 }
 
 void main()
 {
-    uvec3 _S4 = ((gl_LaunchIDEXT));
-    uint idx_1 = _S4.x;
+    uvec3 _S3 = ((gl_LaunchIDEXT));
+    uint idx_1 = _S3.x;
     uint r_0 = 0U;
     RayDesc_0 ray_1 = makeRay_0(idx_1, 0U);
     hitObjectNV hitObj_0;
diff --git a/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-reorder-thread.slang.1.expected b/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-reorder-thread.slang.1.expected
index 00213c62d..28bd5c1d9 100644
--- a/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-reorder-thread.slang.1.expected
+++ b/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-reorder-thread.slang.1.expected
@@ -10,7 +10,7 @@ layout(row_major) buffer;
 layout(binding = 0)
 uniform accelerationStructureEXT scene_0;
 
-layout(std430, binding = 1) buffer _S1 {
+layout(std430, binding = 1) buffer StructuredBuffer_uint_t_0 {
     uint _data[];
 } outputBuffer_0;
 struct SomeValues_0
@@ -39,9 +39,9 @@ SomeValues_0 HitObject_GetAttributes_0(hitObjectNV this_0)
 
 uint calcValue_0(hitObjectNV hit_0)
 {
-    bool _S2 = (hitObjectIsHitNV((hit_0)));
+    bool _S1 = (hitObjectIsHitNV((hit_0)));
     uint r_0;
-    if(_S2)
+    if(_S1)
     {
         uint instanceIndex_0 = (hitObjectGetInstanceCustomIndexNV((hit_0)));
         uint instanceID_0 = (hitObjectGetInstanceIdNV((hit_0)));
@@ -75,45 +75,45 @@ struct RayDesc_0
 
 void main()
 {
-    uvec3 _S3 = ((gl_LaunchIDEXT));
-    ivec2 launchID_0 = ivec2(_S3.xy);
-    uvec3 _S4 = ((gl_LaunchSizeEXT));
+    uvec3 _S2 = ((gl_LaunchIDEXT));
+    ivec2 launchID_0 = ivec2(_S2.xy);
+    uvec3 _S3 = ((gl_LaunchSizeEXT));
     int idx_0 = launchID_0.x;
-    float _S5 = float(idx_0);
-    float _S6 = _S5 * 2.0;
+    float _S4 = float(idx_0);
+    float _S5 = _S4 * 2.0;
     RayDesc_0 ray_0;
-    ray_0.Origin_0 = vec3(_S5, 0.0, 0.0);
+    ray_0.Origin_0 = vec3(_S4, 0.0, 0.0);
     ray_0.TMin_0 = 0.00999999977648258209;
     ray_0.Direction_0 = vec3(0.0, 1.0, 0.0);
     ray_0.TMax_0 = 10000.0;
-    RayDesc_0 _S7 = ray_0;
+    RayDesc_0 _S6 = ray_0;
     hitObjectNV hitObj_0;
     p_1.a_0 = idx_0;
-    p_1.b_0 = _S6;
-    hitObjectTraceRayNV(hitObj_0, scene_0, 20U, 255U, 0U, 4U, 0U, _S7.Origin_0, _S7.TMin_0, _S7.Direction_0, _S7.TMax_0, (1));
+    p_1.b_0 = _S5;
+    hitObjectTraceRayNV(hitObj_0, scene_0, 20U, 255U, 0U, 4U, 0U, _S6.Origin_0, _S6.TMin_0, _S6.Direction_0, _S6.TMax_0, (1));
     uint r_1 = calcValue_0(hitObj_0);
     reorderThreadNV(hitObj_0);
-    float _S8 = _S5 * 4.0;
+    float _S7 = _S4 * 4.0;
     SomeValues_0 otherValues_0;
     otherValues_0.a_0 = idx_0 * -1;
-    otherValues_0.b_0 = _S8;
+    otherValues_0.b_0 = _S7;
     HitObject_Invoke_0(scene_0, hitObj_0, otherValues_0);
-    uint _S9 = calcValue_0(hitObj_0);
-    uint r_2 = r_1 + _S9;
+    uint _S8 = calcValue_0(hitObj_0);
+    uint r_2 = r_1 + _S8;
     reorderThreadNV(hitObj_0, uint(idx_0 & 3), 2U);
-    float _S10 = _S5 * 8.0;
+    float _S9 = _S4 * 8.0;
     otherValues_0.a_0 = idx_0 * -2;
-    otherValues_0.b_0 = _S10;
+    otherValues_0.b_0 = _S9;
     HitObject_Invoke_0(scene_0, hitObj_0, otherValues_0);
-    uint _S11 = calcValue_0(hitObj_0);
-    uint r_3 = r_2 + _S11;
+    uint _S10 = calcValue_0(hitObj_0);
+    uint r_3 = r_2 + _S10;
     reorderThreadNV(uint(idx_0 & 1), 1U);
-    float _S12 = _S5 * 16.0;
+    float _S11 = _S4 * 16.0;
     otherValues_0.a_0 = idx_0 * -4;
-    otherValues_0.b_0 = _S12;
+    otherValues_0.b_0 = _S11;
     HitObject_Invoke_0(scene_0, hitObj_0, otherValues_0);
-    uint _S13 = calcValue_0(hitObj_0);
-    outputBuffer_0._data[uint(idx_0)] = r_3 + _S13;
+    uint _S12 = calcValue_0(hitObj_0);
+    outputBuffer_0._data[uint(idx_0)] = r_3 + _S12;
     return;
 }
 
diff --git a/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-trace-motion-ray.slang.1.expected b/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-trace-motion-ray.slang.1.expected
index 535a7ccb7..872c635f9 100644
--- a/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-trace-motion-ray.slang.1.expected
+++ b/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-trace-motion-ray.slang.1.expected
@@ -11,7 +11,7 @@ layout(row_major) buffer;
 layout(binding = 0)
 uniform accelerationStructureEXT scene_0;
 
-layout(std430, binding = 1) buffer _S1 {
+layout(std430, binding = 1) buffer StructuredBuffer_uint_t_0 {
     uint _data[];
 } outputBuffer_0;
 struct SomeValues_0
@@ -36,9 +36,9 @@ SomeValues_0 HitObject_GetAttributes_0(hitObjectNV this_0)
 
 uint calcValue_0(hitObjectNV hit_0)
 {
-    bool _S2 = (hitObjectIsHitNV((hit_0)));
+    bool _S1 = (hitObjectIsHitNV((hit_0)));
     uint r_0;
-    if(_S2)
+    if(_S1)
     {
         uint instanceIndex_0 = (hitObjectGetInstanceCustomIndexNV((hit_0)));
         uint instanceID_0 = (hitObjectGetInstanceIdNV((hit_0)));
@@ -64,26 +64,26 @@ struct RayDesc_0
 
 void main()
 {
-    uvec3 _S3 = ((gl_LaunchIDEXT));
-    ivec2 launchID_0 = ivec2(_S3.xy);
-    uvec3 _S4 = ((gl_LaunchSizeEXT));
+    uvec3 _S2 = ((gl_LaunchIDEXT));
+    ivec2 launchID_0 = ivec2(_S2.xy);
+    uvec3 _S3 = ((gl_LaunchSizeEXT));
     int idx_0 = launchID_0.x;
     float currentTime_0 = float(idx_0 / 4);
-    float _S5 = float(idx_0);
-    float _S6 = _S5 * 2.0;
+    float _S4 = float(idx_0);
+    float _S5 = _S4 * 2.0;
     RayDesc_0 ray_0;
-    ray_0.Origin_0 = vec3(_S5, 0.0, 0.0);
+    ray_0.Origin_0 = vec3(_S4, 0.0, 0.0);
     ray_0.TMin_0 = 0.00999999977648258209;
     ray_0.Direction_0 = vec3(0.0, 1.0, 0.0);
     ray_0.TMax_0 = 10000.0;
-    RayDesc_0 _S7 = ray_0;
+    RayDesc_0 _S6 = ray_0;
     hitObjectNV hitObj_0;
     p_0.a_0 = idx_0;
-    p_0.b_0 = _S6;
-    hitObjectTraceRayMotionNV(hitObj_0, scene_0, 20U, 255U, 0U, 4U, 0U, _S7.Origin_0, _S7.TMin_0, _S7.Direction_0, _S7.TMax_0, currentTime_0, (0));
-    uint _S8 = uint(idx_0);
-    uint _S9 = calcValue_0(hitObj_0);
-    outputBuffer_0._data[_S8] = _S9;
+    p_0.b_0 = _S5;
+    hitObjectTraceRayMotionNV(hitObj_0, scene_0, 20U, 255U, 0U, 4U, 0U, _S6.Origin_0, _S6.TMin_0, _S6.Direction_0, _S6.TMax_0, currentTime_0, (0));
+    uint _S7 = uint(idx_0);
+    uint _S8 = calcValue_0(hitObj_0);
+    outputBuffer_0._data[_S7] = _S8;
     return;
 }
 
diff --git a/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-trace-ray.slang.1.expected b/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-trace-ray.slang.1.expected
index b1441df72..c3aa6608c 100644
--- a/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-trace-ray.slang.1.expected
+++ b/tests/hlsl-intrinsic/shader-execution-reordering/hit-object-trace-ray.slang.1.expected
@@ -10,7 +10,7 @@ layout(row_major) buffer;
 layout(binding = 0)
 uniform accelerationStructureEXT scene_0;
 
-layout(std430, binding = 1) buffer _S1 {
+layout(std430, binding = 1) buffer StructuredBuffer_uint_t_0 {
     uint _data[];
 } outputBuffer_0;
 struct SomeValues_0
@@ -35,9 +35,9 @@ SomeValues_0 HitObject_GetAttributes_0(hitObjectNV this_0)
 
 uint calcValue_0(hitObjectNV hit_0)
 {
-    bool _S2 = (hitObjectIsHitNV((hit_0)));
+    bool _S1 = (hitObjectIsHitNV((hit_0)));
     uint r_0;
-    if(_S2)
+    if(_S1)
     {
         uint instanceIndex_0 = (hitObjectGetInstanceCustomIndexNV((hit_0)));
         uint instanceID_0 = (hitObjectGetInstanceIdNV((hit_0)));
@@ -63,25 +63,25 @@ struct RayDesc_0
 
 void main()
 {
-    uvec3 _S3 = ((gl_LaunchIDEXT));
-    ivec2 launchID_0 = ivec2(_S3.xy);
-    uvec3 _S4 = ((gl_LaunchSizeEXT));
+    uvec3 _S2 = ((gl_LaunchIDEXT));
+    ivec2 launchID_0 = ivec2(_S2.xy);
+    uvec3 _S3 = ((gl_LaunchSizeEXT));
     int idx_0 = launchID_0.x;
-    float _S5 = float(idx_0);
-    float _S6 = _S5 * 2.0;
+    float _S4 = float(idx_0);
+    float _S5 = _S4 * 2.0;
     RayDesc_0 ray_0;
-    ray_0.Origin_0 = vec3(_S5, 0.0, 0.0);
+    ray_0.Origin_0 = vec3(_S4, 0.0, 0.0);
     ray_0.TMin_0 = 0.00999999977648258209;
     ray_0.Direction_0 = vec3(0.0, 1.0, 0.0);
     ray_0.TMax_0 = 10000.0;
-    RayDesc_0 _S7 = ray_0;
+    RayDesc_0 _S6 = ray_0;
     hitObjectNV hitObj_0;
     p_0.a_0 = idx_0;
-    p_0.b_0 = _S6;
-    hitObjectTraceRayNV(hitObj_0, scene_0, 20U, 255U, 0U, 4U, 0U, _S7.Origin_0, _S7.TMin_0, _S7.Direction_0, _S7.TMax_0, (0));
-    uint _S8 = uint(idx_0);
-    uint _S9 = calcValue_0(hitObj_0);
-    outputBuffer_0._data[_S8] = _S9;
+    p_0.b_0 = _S5;
+    hitObjectTraceRayNV(hitObj_0, scene_0, 20U, 255U, 0U, 4U, 0U, _S6.Origin_0, _S6.TMin_0, _S6.Direction_0, _S6.TMax_0, (0));
+    uint _S7 = uint(idx_0);
+    uint _S8 = calcValue_0(hitObj_0);
+    outputBuffer_0._data[_S7] = _S8;
     return;
 }
 
diff --git a/tests/hlsl/append-structured-buffer.slang b/tests/hlsl/append-structured-buffer.slang
new file mode 100644
index 000000000..8d0352e85
--- /dev/null
+++ b/tests/hlsl/append-structured-buffer.slang
@@ -0,0 +1,43 @@
+
+//TEST:SIMPLE(filecheck=GLSL):-target glsl -profile glsl_450 -stage compute -entry computeMain
+//TEST:SIMPLE(filecheck=SPIRV):-target spirv -profile glsl_450 -stage compute -entry computeMain
+
+//DISABLED_TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -shaderobj -output-using-type -xslang -fvk-use-gl-layout
+
+//TEST_INPUT:ubuffer(data=[0], stride=4):out,name=outputBuffer
+RWStructuredBuffer<float> outputBuffer;
+
+AppendStructuredBuffer<float2> appendBuffer;
+
+// GLSL: layout(std430, binding = 1) buffer StructuredBuffer_float2_t
+// GLSL: vec2 _data[];
+// GLSL: } appendBuffer_elements_0
+
+// GLSL: layout(std430, binding = 2) buffer StructuredBuffer_int_t
+// GLSL: int _data[];
+// GLSL: } appendBuffer_counter
+
+// GLSL: void AppendStructuredBuffer_Append_0(vec2 [[PARAM:[A-Za-z0-9_]+]])
+// GLSL:    int [[COUNTER:[A-Za-z0-9_]+]] = atomicAdd(appendBuffer_counter_0._data[0], 1);
+// GLSL:    appendBuffer_elements_0._data{{\[}}[[COUNTER]]{{\]}} = [[PARAM]];
+
+// GLSL: uvec2 StructuredBuffer_GetDimensions_0()
+// GLSL: {
+// GLSL:     return uvec2(uint(appendBuffer_counter_0._data[0]), 8U);
+// GLSL: }
+
+// SPIRV: OpEntryPoint
+
+//TEST_INPUT:set inBuffer = ubuffer(data=[1.0 2.0 3.0 4.0 5.0 6.0 7.0 8.0], stride=4)
+RWByteAddressBuffer inBuffer;
+
+[numthreads(1, 1, 1)]
+void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID)
+{
+    var g = inBuffer.Load<float4>(4);
+    appendBuffer.Append(g.xy);
+
+    uint numStructs, stride;
+    appendBuffer.GetDimensions(numStructs, stride);
+    outputBuffer[dispatchThreadID.x] = numStructs; // expect 1.0
+}
diff --git a/tests/hlsl/consume-structured-buffer.slang b/tests/hlsl/consume-structured-buffer.slang
new file mode 100644
index 000000000..3027b4184
--- /dev/null
+++ b/tests/hlsl/consume-structured-buffer.slang
@@ -0,0 +1,35 @@
+
+//TEST:SIMPLE(filecheck=GLSL):-target glsl -profile glsl_450 -stage compute -entry computeMain
+//TEST:SIMPLE(filecheck=SPIRV):-target spirv -profile glsl_450 -stage compute -entry computeMain
+
+//DISABLED_TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -shaderobj -output-using-type -xslang -fvk-use-gl-layout
+
+//TEST_INPUT:ubuffer(data=[0], stride=4):out,name=outputBuffer
+RWStructuredBuffer<float> outputBuffer;
+
+ConsumeStructuredBuffer<float2> consumeBuffer;
+
+// GLSL: layout(std430, binding = 1) buffer StructuredBuffer_float2_t
+// GLSL: vec2 _data[];
+// GLSL: } consumeBuffer_elements_0
+
+// GLSL: layout(std430, binding = 2) buffer StructuredBuffer_int_t
+// GLSL: int _data[];
+// GLSL: } consumeBuffer_counter
+
+// GLSL: vec2 ConsumeStructuredBuffer_Consume_0()
+// GLSL:    int [[COUNTER:[A-Za-z0-9_]+]] = atomicAdd(consumeBuffer_counter_0._data[0], -1);
+// GLSL:    int [[COUNTER1:[A-Za-z0-9_]+]] = [[COUNTER]] - 1;
+// GLSL:    if{{\s?}}([[COUNTER1]] >= 0)
+// GLSL:    return consumeBuffer_elements_0._data{{\[}}[[COUNTER1]]{{\]}};
+// GLSL:    else
+// GLSL:    return vec2(0.0);
+
+// SPIRV: OpEntryPoint
+
+[numthreads(1, 1, 1)]
+void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID)
+{
+    var v = consumeBuffer.Consume();
+    outputBuffer[dispatchThreadID.x] = v.x; // expect 1.0
+}
diff --git a/tests/pipeline/ray-tracing/trace-ray-inline.slang.glsl b/tests/pipeline/ray-tracing/trace-ray-inline.slang.glsl
index e1a54a30a..b704eb2b9 100644
--- a/tests/pipeline/ray-tracing/trace-ray-inline.slang.glsl
+++ b/tests/pipeline/ray-tracing/trace-ray-inline.slang.glsl
@@ -28,7 +28,7 @@ layout(std140) uniform _S1
 layout(binding = 0)
 uniform accelerationStructureEXT myAccelerationStructure_0;
 
-layout(std430, binding = 1) buffer _S2 {
+layout(std430, binding = 1) buffer StructuredBuffer_int_t_0 {
     int _data[];
 } resultBuffer_0;
 struct MyProceduralHitAttrs_0
@@ -81,40 +81,37 @@ void main()
     rayQueryEXT query_0;
     MyRayPayload_0 payload_5;
     payload_5.value_1 = -1;
-    rayQueryInitializeEXT((query_0), (myAccelerationStructure_0), (C_0.rayFlags_0 | 512), (C_0.instanceMask_0), (C_0.origin_0), (C_0.tMin_0), (C_0.direction_0), (C_0.tMax_0));
+    rayQueryInitializeEXT((query_0), (myAccelerationStructure_0), (C_0.rayFlags_0 | 512U), (C_0.instanceMask_0), (C_0.origin_0), (C_0.tMin_0), (C_0.direction_0), (C_0.tMax_0));
     MyProceduralHitAttrs_0 committedProceduralAttrs_0;
     for(;;)
     {
-        bool _S3 = rayQueryProceedEXT(query_0);
-        if(!_S3)
+        bool _S2 = rayQueryProceedEXT(query_0);
+        if(!_S2)
         {
             break;
         }
-        uint _S4 = (rayQueryGetIntersectionTypeEXT((query_0), false));
+        uint _S3 = (rayQueryGetIntersectionTypeEXT((query_0), false));
         MyProceduralHitAttrs_0 committedProceduralAttrs_1;
-        switch(_S4)
+        switch(_S3)
         {
         case 1U:
             {
                 MyProceduralHitAttrs_0 candidateProceduralAttrs_0;
                 candidateProceduralAttrs_0.value_0 = 0;
                 float tHit_1 = 0.0;
-                bool _S5 = myProceduralIntersection_0(tHit_1, candidateProceduralAttrs_0);
-                if(_S5)
+                bool _S4 = myProceduralIntersection_0(tHit_1, candidateProceduralAttrs_0);
+                if(_S4)
                 {
-                    bool _S6 = myProceduralAnyHit_0(payload_5);
-                    if(_S6)
+                    bool _S5 = myProceduralAnyHit_0(payload_5);
+                    if(_S5)
                     {
                         rayQueryGenerateIntersectionEXT(query_0, tHit_1);
-                        MyProceduralHitAttrs_0 _S7 = candidateProceduralAttrs_0;
+                        MyProceduralHitAttrs_0 _S6 = candidateProceduralAttrs_0;
                         if(C_0.shouldStopAtFirstHit_0 != 0U)
                         {
                             rayQueryTerminateEXT(query_0);
                         }
-                        else
-                        {
-                        }
-                        committedProceduralAttrs_1 = _S7;
+                        committedProceduralAttrs_1 = _S6;
                     }
                     else
                     {
@@ -129,20 +126,14 @@ void main()
             }
         case 0U:
             {
-                bool _S8 = myTriangleAnyHit_0(payload_5);
-                if(_S8)
+                bool _S7 = myTriangleAnyHit_0(payload_5);
+                if(_S7)
                 {
                     rayQueryConfirmIntersectionEXT(query_0);
                     if(C_0.shouldStopAtFirstHit_0 != 0U)
                     {
                         rayQueryTerminateEXT(query_0);
                     }
-                    else
-                    {
-                    }
-                }
-                else
-                {
                 }
                 committedProceduralAttrs_1 = committedProceduralAttrs_0;
                 break;
@@ -155,8 +146,8 @@ void main()
         }
         committedProceduralAttrs_0 = committedProceduralAttrs_1;
     }
-    uint _S9 = (rayQueryGetIntersectionTypeEXT((query_0), true));
-    switch(_S9)
+    uint _S8 = (rayQueryGetIntersectionTypeEXT((query_0), true));
+    switch(_S8)
     {
     case 1U:
         {
@@ -178,6 +169,8 @@ void main()
             break;
         }
     }
-    ((resultBuffer_0)._data[(index_0)]) = payload_5.value_1;
+    resultBuffer_0._data[index_0] = payload_5.value_1;
     return;
 }
+
+
diff --git a/tests/slang-extension/atomic-float-byte-address-buffer-cross.slang.glsl b/tests/slang-extension/atomic-float-byte-address-buffer-cross.slang.glsl
index 4fb647199..330e76d11 100644
--- a/tests/slang-extension/atomic-float-byte-address-buffer-cross.slang.glsl
+++ b/tests/slang-extension/atomic-float-byte-address-buffer-cross.slang.glsl
@@ -2,34 +2,64 @@
 #extension GL_EXT_shader_atomic_float : require
 layout(row_major) uniform;
 layout(row_major) buffer;
-layout(std430, binding = 1) buffer _S1 {
+
+#line 11 "tests/slang-extension/atomic-float-byte-address-buffer-cross.slang"
+layout(std430, binding = 1) buffer StructuredBuffer_float_t_0 {
     float _data[];
 } anotherBuffer_0;
-layout(std430, binding = 0) buffer _S2 {
+
+#line 11
+layout(std430, binding = 0) buffer StructuredBuffer_float_t_1 {
     float _data[];
-} _S3;
-void RWByteAddressBuffer_InterlockedAddF32_0(uint _S4, float _S5, out float _S6)
+} outputBuffer_0;
+
+#line 1264 "core.meta.slang"
+void RWByteAddressBuffer_InterlockedAddF32_0(uint _S1, float _S2, out float _S3)
 {
-    float _S7 = (atomicAdd((((_S3)._data[(_S4 / 4U)])), (_S5)));
-    _S6 = _S7;
+
+#line 391 "hlsl.meta.slang"
+    float _S4 = (atomicAdd((outputBuffer_0._data[_S1 / 4U]), (_S2)));
+
+#line 391
+    _S3 = _S4;
     return;
 }
 
-void RWByteAddressBuffer_InterlockedAddF32_1(uint _S8, float _S9)
+
+#line 392
+void RWByteAddressBuffer_InterlockedAddF32_1(uint _S5, float _S6)
 {
-    float _S10 = (atomicAdd((((_S3)._data[(_S8 / 4U)])), (_S9)));
+
+#line 406
+    float _S7 = (atomicAdd((outputBuffer_0._data[_S5 / 4U]), (_S6)));
     return;
 }
 
+
+#line 14 "tests/slang-extension/atomic-float-byte-address-buffer-cross.slang"
 layout(local_size_x = 16, local_size_y = 1, local_size_z = 1) in;
 void main()
 {
+
+#line 16
     uint tid_0 = gl_GlobalInvocationID.x;
-    uint _S11 = tid_0 >> 2;
-    int idx_0 = int(tid_0 & 3U ^ _S11);
-    float delta_0 = ((anotherBuffer_0)._data[(uint(idx_0 & 3))]);
+    uint _S8 = tid_0 >> 2;
+
+#line 17
+    int idx_0 = int(tid_0 & 3U ^ _S8);
+
+    float delta_0 = anotherBuffer_0._data[uint(idx_0 & 3)];
+
     float previousValue_0 = 0.0;
+
+#line 21
     RWByteAddressBuffer_InterlockedAddF32_0(uint(idx_0 << 2), 1.0, previousValue_0);
-    RWByteAddressBuffer_InterlockedAddF32_1(uint(int(_S11) << 2), delta_0);
+
+#line 21
+    RWByteAddressBuffer_InterlockedAddF32_1(uint(int(_S8) << 2), delta_0);
+
+#line 27
     return;
 }
+
+
diff --git a/tests/vkray/closesthit.slang.glsl b/tests/vkray/closesthit.slang.glsl
index 6094b3a3d..bb3e81f87 100644
--- a/tests/vkray/closesthit.slang.glsl
+++ b/tests/vkray/closesthit.slang.glsl
@@ -3,17 +3,17 @@
 #extension GL_NV_ray_tracing : require
 
 #define tmp_shaderrecord    _S1
-#define tmp_colors          _S2
-#define tmp_hitattrs        _S3
-#define tmp_payload         _S4
-#define tmp_customidx       _S5
-#define tmp_instanceid      _S6
-#define tmp_add_0           _S7
-#define tmp_primid          _S8
-#define tmp_add_1           _S9
-#define tmp_hitkind         _S10
-#define tmp_hitt            _S11
-#define tmp_tmin            _S12
+#define tmp_colors          StructuredBuffer_float4_t_0
+#define tmp_hitattrs        _S2
+#define tmp_payload         _S3
+#define tmp_customidx       _S4
+#define tmp_instanceid      _S5
+#define tmp_add_0           _S6
+#define tmp_primid          _S7
+#define tmp_add_1           _S8
+#define tmp_hitkind         _S9
+#define tmp_hitt            _S10
+#define tmp_tmin            _S11
 
 struct SLANG_ParameterGroup_ShaderRecord_0
 {
diff --git a/tests/vkray/entry-point-params.slang.glsl b/tests/vkray/entry-point-params.slang.glsl
index 00d2ba630..a505b337e 100644
--- a/tests/vkray/entry-point-params.slang.glsl
+++ b/tests/vkray/entry-point-params.slang.glsl
@@ -15,7 +15,7 @@
 #endif
 
 layout(std430, binding = 0)
-buffer _S1 {
+buffer StructuredBuffer_float_t_0 {
     float _data[];
 } buffer_0;
 
@@ -25,14 +25,14 @@ struct EntryPointParams_0
 };
 
 layout(shaderRecordEXT)
-buffer _S2
+buffer _S1
 {
     float value_0;
-} _S3;
+} _S2;
 
 void main()
 {
-    uvec3 _S4 = gl_LaunchIDEXT;
-    buffer_0._data[_S4.x] = _S3.value_0;
+    uvec3 _S3 = gl_LaunchIDEXT;
+    buffer_0._data[_S3.x] = _S2.value_0;
     return;
 }