18 files changed, 602 insertions, 234 deletions

diff --git a/source/slang/slang-diagnostic-defs.h b/source/slang/slang-diagnostic-defs.h
index 5a841dfc9..d3aaeb948 100644
--- a/source/slang/slang-diagnostic-defs.h
+++ b/source/slang/slang-diagnostic-defs.h
@@ -453,7 +453,7 @@ DIAGNOSTIC(39015, Error, wholeSpaceParameterRequiresZeroBinding, "shader paramet
 DIAGNOSTIC(39013, Error, dontExpectOutParametersForStage, "the '$0' stage does not support `out` or `inout` entry point parameters")
 DIAGNOSTIC(39014, Error, dontExpectInParametersForStage, "the '$0' stage does not support `in` entry point parameters")
 
-DIAGNOSTIC(39016, Error, globalUniformsNotSupported, "'$0' is implicitly a global uniform shader parameter, which is currently unsupported by Slang. If a uniform parameter is intended, use a constant buffer or parameter block. If a global is intended, use the 'static' modifier.")
+DIAGNOSTIC(39016, Warning, globalUniformNotExpected, "'$0' is implicitly a global shader parameter, not a global variable. If a global variable is intended, add the 'static' modifier. If a uniform shader parameter is intended, add the 'uniform' modifier to silence this warning.")
 
 DIAGNOSTIC(39017, Error, tooManyShaderRecordConstantBuffers, "can have at most one 'shader record' attributed constant buffer; found $0.")
 
diff --git a/source/slang/slang-emit-c-like.h b/source/slang/slang-emit-c-like.h
index c37a1514e..b04b075c5 100644
--- a/source/slang/slang-emit-c-like.h
+++ b/source/slang/slang-emit-c-like.h
@@ -278,7 +278,7 @@ public:
     void computeEmitActions(IRModule* module, List<EmitAction>& ioActions);
 
     void executeEmitActions(List<EmitAction> const& actions);
-    void emitModule(IRModule* module) { emitModuleImpl(module); }
+    void emitModule(IRModule* module) { m_irModule = module; emitModuleImpl(module); }
 
     void emitPreprocessorDirectives() { emitPreprocessorDirectivesImpl(); }
     void emitSimpleType(IRType* type);
@@ -352,6 +352,7 @@ public:
     List<IRWitnessTableEntry*> getSortedWitnessTableEntries(IRWitnessTable* witnessTable);
     
     BackEndCompileRequest* m_compileRequest = nullptr;
+    IRModule* m_irModule = nullptr;
 
     // The stage for which we are emitting code.
     //
@@ -370,15 +371,8 @@ public:
     // Where source is written to
     SourceWriter* m_writer;
 
-    // We only want to emit each `import`ed module one time, so
-    // we maintain a set of already-emitted modules.
-    HashSet<ModuleDecl*> m_modulesAlreadyEmitted;
-
-    ModuleDecl* m_program = nullptr;
-
     UInt m_uniqueIDCounter = 1;
     Dictionary<IRInst*, UInt> m_mapIRValueToID;
-    Dictionary<Decl*, UInt> m_mapDeclToID;
 
     HashSet<String> m_irDeclsVisited;
 
diff --git a/source/slang/slang-emit-cpp.cpp b/source/slang/slang-emit-cpp.cpp
index 6b7fc46bf..1ff5af4fb 100644
--- a/source/slang/slang-emit-cpp.cpp
+++ b/source/slang/slang-emit-cpp.cpp
@@ -500,6 +500,16 @@ SlangResult CPPSourceEmitter::calcTypeName(IRType* type, CodeGenTarget target, S
             out << "void*";
             return SLANG_OK;
         }
+        case kIROp_ConstantBufferType:
+        case kIROp_ParameterBlockType:
+        {
+            auto groupType = cast<IRParameterGroupType>(type);
+            auto elementType = groupType->getElementType();
+
+            SLANG_RETURN_ON_FAIL(calcTypeName(elementType, target, out));
+            out << "*";
+            return SLANG_OK;
+        }
         default:
         {
             if (isNominalOp(type->op))
@@ -2357,29 +2367,6 @@ void CPPSourceEmitter::emitOperandImpl(IRInst* inst, EmitOpInfo const&  outerPre
 
     switch (inst->op)
     {
-        case 0: // nothing yet
-        case kIROp_GlobalParam:
-        {            
-            String name = getName(inst);
-
-            if (inst->findDecorationImpl(kIROp_EntryPointParamDecoration))
-            {
-                // It's an entry point parameter
-                // The parameter is held in a struct so always deref
-                m_writer->emit("(*");
-                m_writer->emit(name);
-                m_writer->emit(")");
-            }
-            else
-            {
-                // It's in UniformState
-                m_writer->emit("(");
-                m_writer->emit("uniformState->");
-                m_writer->emit(name);
-                m_writer->emit(")");
-            }
-            break;
-        }
         case kIROp_Param:
         {
             auto varLayout = getVarLayout(inst);
@@ -2443,7 +2430,7 @@ static bool _isFunction(IROp op)
     return op == kIROp_Func;
 }
 
-void CPPSourceEmitter::_emitEntryPointDefinitionStart(IRFunc* func, IRGlobalParam* entryPointGlobalParams, const String& funcName, const UnownedStringSlice& varyingTypeName)
+void CPPSourceEmitter::_emitEntryPointDefinitionStart(IRFunc* func, IRGlobalParam* entryPointParams, IRGlobalParam* globalParams, const String& funcName, const UnownedStringSlice& varyingTypeName)
 {
     auto resultType = func->getResultType();
     
@@ -2456,27 +2443,35 @@ void CPPSourceEmitter::_emitEntryPointDefinitionStart(IRFunc* func, IRGlobalPara
 
     m_writer->emit("(");
     m_writer->emit(varyingTypeName);
-    m_writer->emit("* varyingInput, void* params, void* uniformState)");
+    m_writer->emit("* varyingInput, void* entryPointParams, void* globalParams)");
     emitSemantics(func);
     m_writer->emit("\n{\n");
 
     m_writer->indent();
     // Initialize when constructing so that globals are zeroed
     m_writer->emit("KernelContext context = {};\n");
-    m_writer->emit("context.uniformState = (UniformState*)uniformState;\n");
     
-    if (entryPointGlobalParams)
+    if (entryPointParams)
     {
-        auto varDecl = entryPointGlobalParams;
-        auto rawType = varDecl->getDataType();
+        auto param = entryPointParams;
+        auto paramType = param->getDataType();
 
-        auto varType = rawType;
+        m_writer->emit("context.");
+        m_writer->emit(getName(param));
+        m_writer->emit(" =  (");
+        emitType(paramType);
+        m_writer->emit(")entryPointParams; \n");
+    }
+    if (globalParams)
+    {
+        auto param = globalParams;
+        auto paramType = param->getDataType();
 
         m_writer->emit("context.");
-        m_writer->emit(getName(varDecl));
+        m_writer->emit(getName(param));
         m_writer->emit(" =  (");
-        emitType(varType);
-        m_writer->emit("*)params; \n");
+        emitType(paramType);
+        m_writer->emit(")globalParams; \n");
     }
 }
 
@@ -2677,71 +2672,53 @@ void CPPSourceEmitter::_emitForwardDeclarations(const List<EmitAction>& actions)
     }
 }
 
-void CPPSourceEmitter::_calcGlobalParams(const List<EmitAction>& actions, List<GlobalParamInfo>& outParams, IRGlobalParam** outEntryPointGlobalParams)
+void CPPSourceEmitter::_findShaderParams(
+    IRGlobalParam** outEntryPointParam,
+    IRGlobalParam** outGlobalParam)
 {
-    outParams.clear();
-    *outEntryPointGlobalParams = nullptr;
+    SLANG_ASSERT(outEntryPointParam);
+    SLANG_ASSERT(outGlobalParam);
 
-    IRGlobalParam* entryPointGlobalParams = nullptr;
-    for (auto action : actions)
-    {
-        if (action.level == EmitAction::Level::Definition && action.inst->op == kIROp_GlobalParam)
-        {
-            auto inst = action.inst;
+    IRGlobalParam*& entryPointParam = *outEntryPointParam;
+    IRGlobalParam*& globalParam = *outGlobalParam;
 
-            if (inst->findDecorationImpl(kIROp_EntryPointParamDecoration))
-            {
-                // Should only be one instruction marked this way
-                SLANG_ASSERT(entryPointGlobalParams == nullptr);
-                entryPointGlobalParams = as<IRGlobalParam>(inst);
-                continue;
-            }
-
-            IRVarLayout* varLayout = CLikeSourceEmitter::getVarLayout(action.inst);
-            SLANG_ASSERT(varLayout);
-
-            IRVarOffsetAttr* offsetAttr = varLayout->findOffsetAttr(LayoutResourceKind::Uniform);
-            IRTypeLayout* typeLayout = varLayout->getTypeLayout();
-            IRTypeSizeAttr* sizeAttr = typeLayout->findSizeAttr(LayoutResourceKind::Uniform);
-
-            GlobalParamInfo paramInfo;
-            paramInfo.inst = action.inst;
-            // Index is the byte offset for uniform
-            paramInfo.offset = offsetAttr ? offsetAttr->getOffset() : 0;
-            paramInfo.size = sizeAttr ? sizeAttr->getFiniteSize() : 0;
+    for(auto inst : m_irModule->getGlobalInsts())
+    {
+        auto param = as<IRGlobalParam>(inst);
+        if(!param)
+            continue;
 
-            outParams.add(paramInfo);
+        // Currently, the entry-point parameters
+        // are represented as a single parameter
+        // at the global scope, and the same is
+        // true of the parameters that were
+        // originally declared as globals.
+        //
+        // We need to find capture each of these
+        // parameters, and we need to tell them
+        // apart. Luckily, the logic that
+        // moved the entry-point parameters to
+        // global scope will ahve also marked
+        // the entry-point parameters with
+        // a decoration that we can detect.
+        //
+        if (inst->findDecorationImpl(kIROp_EntryPointParamDecoration))
+        {
+            // Should only be one instruction marked this way
+            SLANG_ASSERT(entryPointParam == nullptr);
+            entryPointParam = param;
+            continue;
         }
-    }
-
-    // We want to sort by layout offset, and insert suitable padding
-    outParams.sort();
-
-    *outEntryPointGlobalParams = entryPointGlobalParams;
-}
-
-void CPPSourceEmitter::_emitUniformStateMembers(const List<EmitAction>& actions, IRGlobalParam** outEntryPointGlobalParams)
-{
-    List<GlobalParamInfo> params;
-    _calcGlobalParams(actions, params, outEntryPointGlobalParams);
-
-    int padIndex = 0;
-    size_t offset = 0;
-    for (const auto& paramInfo : params)
-    {
-        if (offset < paramInfo.offset)
+        else
         {
-            // We want to output some padding
-            StringBuilder builder;
-            builder << "uint8_t _pad" << (padIndex++) << "[" << (paramInfo.offset - offset) << "];\n";
-            m_writer->emit(builder);
+            // There should only be one instruction representing
+            // the global-scope shader parameters.
+            //
+            SLANG_ASSERT(globalParam == nullptr);
+            globalParam = param;
+            continue;
         }
-
-        emitGlobalInst(paramInfo.inst);
-        // Set offset after this 
-        offset = paramInfo.offset + paramInfo.size;
     }
-    m_writer->emit("\n");
 }
 
 void CPPSourceEmitter::emitModuleImpl(IRModule* module)
@@ -2756,26 +2733,15 @@ void CPPSourceEmitter::emitModuleImpl(IRModule* module)
     
     _emitForwardDeclarations(actions);
 
-    IRGlobalParam* entryPointGlobalParams = nullptr;
-    
-    // Output the global parameters in a 'UniformState' structure
-    {
-        m_writer->emit("struct UniformState\n{\n");
-        m_writer->indent();
-
-        _emitUniformStateMembers(actions, &entryPointGlobalParams);
-
-        m_writer->dedent();
-        m_writer->emit("\n};\n\n");
-    }
+    IRGlobalParam* entryPointParams = nullptr;
+    IRGlobalParam* globalParams = nullptr;
+    _findShaderParams(&entryPointParams, &globalParams);
     
     // Output the 'Context' which will be used for execution
     {
         m_writer->emit("struct KernelContext\n{\n");
         m_writer->indent();
 
-        m_writer->emit("UniformState* uniformState;\n");
-
         m_writer->emit("uint3 dispatchThreadID;\n");
 
         //if (m_semanticUsedFlags & SemanticUsedFlag::GroupID)
@@ -2797,9 +2763,14 @@ void CPPSourceEmitter::emitModuleImpl(IRModule* module)
             m_writer->emit("}\n");
         }
 
-        if (entryPointGlobalParams)
+
+        if (globalParams)
+        {
+            emitGlobalInst(globalParams);
+        }
+        if (entryPointParams)
         {
-            emitGlobalInst(entryPointGlobalParams);
+            emitGlobalInst(entryPointParams);
         }
 
         // Output all the thread locals 
@@ -2865,7 +2836,7 @@ void CPPSourceEmitter::emitModuleImpl(IRModule* module)
 
                     String threadFuncName = builder;
 
-                    _emitEntryPointDefinitionStart(func, entryPointGlobalParams, threadFuncName, UnownedStringSlice::fromLiteral("ComputeThreadVaryingInput"));
+                    _emitEntryPointDefinitionStart(func, entryPointParams, globalParams, threadFuncName, UnownedStringSlice::fromLiteral("ComputeThreadVaryingInput"));
 
                     if (m_semanticUsedFlags & SemanticUsedFlag::GroupThreadID)
                     {
@@ -2896,7 +2867,7 @@ void CPPSourceEmitter::emitModuleImpl(IRModule* module)
 
                     String groupFuncName = builder;
 
-                    _emitEntryPointDefinitionStart(func, entryPointGlobalParams, groupFuncName, UnownedStringSlice::fromLiteral("ComputeVaryingInput"));
+                    _emitEntryPointDefinitionStart(func, entryPointParams, globalParams, groupFuncName, UnownedStringSlice::fromLiteral("ComputeVaryingInput"));
 
                     m_writer->emit("const uint3 start = ");
                     _emitInitAxisValues(groupThreadSize, UnownedStringSlice::fromLiteral("varyingInput->startGroupID"), UnownedStringSlice());
@@ -2918,7 +2889,7 @@ void CPPSourceEmitter::emitModuleImpl(IRModule* module)
 
                 // Emit the main version - which takes a dispatch size
                 {
-                    _emitEntryPointDefinitionStart(func, entryPointGlobalParams, funcName, UnownedStringSlice::fromLiteral("ComputeVaryingInput"));
+                    _emitEntryPointDefinitionStart(func, entryPointParams, globalParams, funcName, UnownedStringSlice::fromLiteral("ComputeVaryingInput"));
 
                     m_writer->emit("const uint3 start = ");
                     _emitInitAxisValues(groupThreadSize, UnownedStringSlice::fromLiteral("varyingInput->startGroupID"), UnownedStringSlice());
diff --git a/source/slang/slang-emit-cpp.h b/source/slang/slang-emit-cpp.h
index c64b92b0f..5a18686ec 100644
--- a/source/slang/slang-emit-cpp.h
+++ b/source/slang/slang-emit-cpp.h
@@ -37,18 +37,6 @@ public:
         int colCount;
     };
 
-    struct GlobalParamInfo
-    {
-        typedef GlobalParamInfo ThisType;
-        bool operator<(const ThisType& rhs) const { return offset < rhs.offset; }
-        bool operator==(const ThisType& rhs) const { return offset == rhs.offset; }
-        bool operator!=(const ThisType& rhs) const { return !(*this == rhs); }
-
-        IRInst* inst;
-        UInt offset;
-        UInt size;
-    };
-
     virtual void useType(IRType* type);
     virtual void emitCall(const HLSLIntrinsic* specOp, IRInst* inst, const IRUse* operands, int numOperands, const EmitOpInfo& inOuterPrec);
     virtual void emitTypeDefinition(IRType* type);
@@ -94,8 +82,11 @@ protected:
     void _maybeEmitSpecializedOperationDefinition(const HLSLIntrinsic* specOp);
 
     void _emitForwardDeclarations(const List<EmitAction>& actions);
-    void _calcGlobalParams(const List<EmitAction>& actions, List<GlobalParamInfo>& outParams, IRGlobalParam** outEntryPointGlobalParams);
-    void _emitUniformStateMembers(const List<EmitAction>& actions, IRGlobalParam** outEntryPointGlobalParams);
+
+        /// Find the IR global parameters representing the entry-point and global shader parameters (if any)
+    void _findShaderParams(
+        IRGlobalParam** outEntryPointParam,
+        IRGlobalParam** outGlobalParam);
 
     void _emitAryDefinition(const HLSLIntrinsic* specOp);
 
@@ -127,7 +118,7 @@ protected:
     
     SlangResult _calcCPPTextureTypeName(IRTextureTypeBase* texType, StringBuilder& outName);
 
-    void _emitEntryPointDefinitionStart(IRFunc* func, IRGlobalParam* entryPointGlobalParams, const String& funcName, const UnownedStringSlice& varyingTypeName);
+    void _emitEntryPointDefinitionStart(IRFunc* func, IRGlobalParam* entryPointParams, IRGlobalParam* globalParams, const String& funcName, const UnownedStringSlice& varyingTypeName);
     void _emitEntryPointDefinitionEnd(IRFunc* func);
     void _emitEntryPointGroup(const Int sizeAlongAxis[kThreadGroupAxisCount], const String& funcName);
     void _emitEntryPointGroupRange(const Int sizeAlongAxis[kThreadGroupAxisCount], const String& funcName);
diff --git a/source/slang/slang-emit-cuda.cpp b/source/slang/slang-emit-cuda.cpp
index b79518052..d05c4edfc 100644
--- a/source/slang/slang-emit-cuda.cpp
+++ b/source/slang/slang-emit-cuda.cpp
@@ -711,19 +711,9 @@ void CUDASourceEmitter::emitModuleImpl(IRModule* module)
 
     _emitForwardDeclarations(actions);
 
-    IRGlobalParam* entryPointGlobalParams = nullptr;
-
-    // Output the global parameters in a 'UniformState' structure
-    {
-        m_writer->emit("struct UniformState\n{\n");
-        m_writer->indent();
-
-        // We need these to be prefixed by __device__
-        _emitUniformStateMembers(actions, &entryPointGlobalParams);
-
-        m_writer->dedent();
-        m_writer->emit("\n};\n\n");
-    }
+    IRGlobalParam* entryPointParams = nullptr;
+    IRGlobalParam* globalParams = nullptr;
+    _findShaderParams(&entryPointParams, &globalParams);
 
     // Output group shared variables
 
@@ -742,11 +732,13 @@ void CUDASourceEmitter::emitModuleImpl(IRModule* module)
         m_writer->emit("struct KernelContext\n{\n");
         m_writer->indent();
 
-        m_writer->emit("UniformState* uniformState;\n");
-
-        if (entryPointGlobalParams)
+        if (globalParams)
         {
-            emitGlobalInst(entryPointGlobalParams);
+            emitGlobalInst(globalParams);
+        }
+        if (entryPointParams)
+        {
+            emitGlobalInst(entryPointParams);
         }
 
         // Output all the thread locals 
@@ -813,7 +805,7 @@ void CUDASourceEmitter::emitModuleImpl(IRModule* module)
             #undef CASE
                 }
 
-                if( stage != Stage::Compute )
+                if(globalParams && stage != Stage::Compute )
                 {
                     // Non-compute shaders (currently just OptiX ray tracing kernels)
                     // require parameter data that is shared across multiple kernels
@@ -826,7 +818,20 @@ void CUDASourceEmitter::emitModuleImpl(IRModule* module)
                     // used here (`SLANG_globalParams`) is thus a part of the
                     // binary interface for Slang->OptiX translation.
                     //
-                    m_writer->emit("extern \"C\" { __constant__ UniformState SLANG_globalParams; }\n");
+                    // TODO: We need to make a decision about how indirected
+                    // the parameter passing for global-scope data is going to
+                    // be for CUDA and OptiX (ideally with an answer that is
+                    // consistent across the two). For now we are deciding to
+                    // make this global `__constant__` variable represent the
+                    // global parameter data directly, rather than indirectly.
+                    //
+                    auto globalParamsPtrType = as<IRPointerLikeType>(globalParams->getDataType());
+                    SLANG_ASSERT(globalParamsPtrType);
+                    auto gloablParamsElementType = globalParamsPtrType->getElementType();
+                    //
+                    m_writer->emit("extern \"C\" { __constant__ ");
+                    emitType(gloablParamsElementType, "SLANG_globalParams");
+                    m_writer->emit("; }\n");
                 }
 
                 // As a convenience for anybody reading the generated
@@ -869,7 +874,7 @@ void CUDASourceEmitter::emitModuleImpl(IRModule* module)
                     // `uniform` parameter data, and the second points to the global-scope
                     // parameter data (if any).
                     //
-                    m_writer->emit("(void* entryPointShaderParameters, void* uniformState)");
+                    m_writer->emit("(void* entryPointParams, void* globalParams)");
                 }
                 else
                 {
@@ -891,18 +896,27 @@ void CUDASourceEmitter::emitModuleImpl(IRModule* module)
                 // a compute shader or a ray-tracing shader, so we need to alter how we initialize
                 // the pointer in our `context` based on the stage.
                 //
-                if( stage == Stage::Compute )
-                {
-                    m_writer->emit("context.uniformState = (UniformState*)uniformState;\n");
-                }
-                else
+                if( globalParams )
                 {
-                    m_writer->emit("context.uniformState = &SLANG_globalParams;\n");
+                    if( stage == Stage::Compute )
+                    {
+                        m_writer->emit("context.");
+                        m_writer->emit(getName(globalParams));
+                        m_writer->emit(" = (");
+                        emitType(globalParams->getDataType());
+                        m_writer->emit(")globalParams;\n");
+                    }
+                    else
+                    {
+                        m_writer->emit("context.");
+                        m_writer->emit(getName(globalParams));
+                        m_writer->emit(" = &SLANG_globalParams;\n");
+                    }
                 }
 
-                if (entryPointGlobalParams)
+                if (entryPointParams)
                 {
-                    auto varDecl = entryPointGlobalParams;
+                    auto varDecl = entryPointParams;
                     auto rawType = varDecl->getDataType();
                     auto varType = rawType;
 
@@ -910,7 +924,7 @@ void CUDASourceEmitter::emitModuleImpl(IRModule* module)
                     m_writer->emit(getName(varDecl));
                     m_writer->emit(" =  (");
                     emitType(varType);
-                    m_writer->emit("*)");
+                    m_writer->emit(")");
 
                     // Similar to the case for global parameter data above, the entry-point
                     // uniform parameter data gets passed in differently for compute kernels
@@ -922,7 +936,7 @@ void CUDASourceEmitter::emitModuleImpl(IRModule* module)
                         // in as an explicit parameter on the CUDA kernel, and we simply
                         // cast it to the expected type here.
                         //
-                        m_writer->emit("entryPointShaderParameters");
+                        m_writer->emit("entryPointParams");
                     }
                     else
                     {
diff --git a/source/slang/slang-emit.cpp b/source/slang/slang-emit.cpp
index f824ac50f..7d2260597 100644
--- a/source/slang/slang-emit.cpp
+++ b/source/slang/slang-emit.cpp
@@ -6,6 +6,7 @@
 
 #include "slang-ir-bind-existentials.h"
 #include "slang-ir-byte-address-legalize.h"
+#include "slang-ir-collect-global-uniforms.h"
 #include "slang-ir-dce.h"
 #include "slang-ir-entry-point-uniforms.h"
 #include "slang-ir-glsl-legalize.h"
@@ -224,26 +225,42 @@ Result linkAndOptimizeIR(
 #endif
     validateIRModuleIfEnabled(compileRequest, irModule);
 
-
-
-
-
     // Now that we've linked the IR code, any layout/binding
     // information has been attached to shader parameters
     // and entry points. Now we are safe to make transformations
     // that might move code without worrying about losing
     // the connection between a parameter and its layout.
+
+    // One example of a transformation that needs to wait until
+    // we have layout information is the step where we collect
+    // any global-scope shader parameters with ordinary/uniform
+    // type into an aggregate `struct`, and then (optionally)
+    // wrap that `struct` up in a constant buffer.
+    //
+    // This step allows shaders to declare parameters of ordinary
+    // type as globals in the input file, while ensuring that
+    // downstream passes for graphics APIs like Vulkan and D3D
+    // can assume that all ordinary/uniform data is strictly
+    // passed using constant buffers.
     //
-    // An easy transformation of this kind is to take uniform
+    collectGlobalUniformParameters(irModule, outLinkedIR.globalScopeVarLayout);
+#if 0
+    dumpIRIfEnabled(compileRequest, irModule, "GLOBAL UNIFORMS COLLECTED");
+#endif
+    validateIRModuleIfEnabled(compileRequest, irModule);
+
+    // Another transformation that needed to wait until we
+    // had layout information on parameters is to take uniform
     // parameters of a shader entry point and move them into
     // the global scope instead.
     //
-    moveEntryPointUniformParamsToGlobalScope(irModule, target);
+    moveEntryPointUniformParamsToGlobalScope(irModule);
 #if 0
     dumpIRIfEnabled(compileRequest, irModule, "ENTRY POINT UNIFORMS MOVED");
 #endif
     validateIRModuleIfEnabled(compileRequest, irModule);
 
+
     // Desguar any union types, since these will be illegal on
     // various targets.
     //
diff --git a/source/slang/slang-ir-collect-global-uniforms.cpp b/source/slang/slang-ir-collect-global-uniforms.cpp
new file mode 100644
index 000000000..306bbdfee
--- /dev/null
+++ b/source/slang/slang-ir-collect-global-uniforms.cpp
@@ -0,0 +1,288 @@
+// slang-ir-collect-global-uniforms.cpp
+#include "slang-ir-collect-global-uniforms.h"
+
+#include "slang-ir-insts.h"
+
+namespace Slang
+{
+
+// This file implements a pass that takes input code like:
+//
+//      uniform int gA;
+//      uniform float gB;
+//
+//      void main() { ... gA + gB ... }
+//
+// and transforms it into code like:
+//
+//      struct GlobalParams
+//      {
+//          int gA;
+//          float gB;
+//      }
+//
+//      ConstantBuffer<GlobalParams> globalParams;
+//
+//      void main() { ... globalParams.gA + globalParams.gB ... }
+//
+// The main consequence of this transformation is that we can support
+// global `uniform` shader parameters of "ordinary" data types when
+// compiling for targets that don't directly support that feature
+// (e.g., GLSL/SPIR-V).
+//
+// In addition, on targets that already support a similar transformation
+// (e.g., when compiling to DXBC/DXIL via fxc/dxc), making the `globalParams`
+// constant buffer explicit allows us to control the binding that is
+// assigned to it using our existing logic for automatic layout, rather than
+// being left at the whims of the undocumented defaults of those compilers.
+//
+// A final consequence of this pass is that for targets where *all*
+// shader parameters use "ordinary" data types (because there are no
+// non-first-class types), we end up with a conveniently packaged up
+// single parameter and type that encapsulates all of the shader inputs.
+//
+struct CollectGlobalUniformParametersContext
+{
+    // In orderto perform our transformation, we need access to the module
+    // to be transformed, as well as the layout information representing
+    // the global-scope shader parameters.
+    //
+    IRModule* module;
+    IRVarLayout* globalScopeVarLayout;
+
+    // This is a relatively simple pass, and it is all driven
+    // by a single subroutine.
+    //
+    void processModule()
+    {
+        // We start by looking at the layout that was computed for the global-scope
+        // parameters to determine how the parameters are supposed to be pacakged.
+        //
+        // This step relies on the earlier layout computation logic to have implemented
+        // any target-specific policies around how the global-scope parametesr are
+        // to be passed, and therefore we avoid trying to make any target-specific
+        // decisions in this pass.
+        //
+        auto globalScopeTypeLayout = globalScopeVarLayout->getTypeLayout();
+        auto globalParamsTypeLayout = globalScopeTypeLayout;
+
+        // One example of a difference that might appear in the global-scope layout
+        // depending on the target is that the global-scope parameters might end
+        // up just pacakged as a `struct`, *or* they might be packaged up in a
+        // `ConstantBuffer<...>` or other parameter group that wraps that `struct`.
+        //
+        IRParameterGroupTypeLayout* globalParameterGroupTypeLayout = as<IRParameterGroupTypeLayout>(globalParamsTypeLayout);
+        if( globalParameterGroupTypeLayout )
+        {
+            // In the case where there is a wrapping `ConstantBuffer<...>`, we want to
+            // get at the element type of that constant buffer, because it represents
+            // the packaged-up `struct` that we want to produce.
+            //
+            globalParamsTypeLayout = globalParameterGroupTypeLayout->getElementVarLayout()->getTypeLayout();
+        }
+
+        // As a special case (in order to avoid disruption to any downstream passes),
+        // if the layout for the global-scope parameters doesn't include any "ordinary"
+        // data (represented as `LayoutResourceKind::Uniform`), then we will not do
+        // the "packaging up" step at all.
+        //
+        // This means that the current pass will not change the results for a majority
+        // of targets (notably, all the current graphics APIs) *unless* global shader
+        // parameters are declared that use "ordinary' data.
+        //
+        // TODO: eventually we should remove this special case, and confirm that the resulting
+        // logic works across all shaders (it should). Doing so will be a necessary
+        // step if want to start applying the packaging-up of global-scope parameters on
+        // a per-module basis.
+        //
+        if(!globalParameterGroupTypeLayout && !globalParamsTypeLayout->findSizeAttr(LayoutResourceKind::Uniform))
+            return;
+
+        // We expect that the layout for the global-scope parameters is always
+        // computed for a `struct` type.
+        //
+        auto globalParamsStructTypeLayout = as<IRStructTypeLayout>(globalParamsTypeLayout);
+        SLANG_ASSERT(globalParamsStructTypeLayout);
+
+        // We need to construct a single IR parameter that will represent
+        // the collected global-scope parameters. The `IRBuilder` we construct
+        // for this will also be used when replacing the individual parameters.
+        //
+        SharedIRBuilder sharedBuilder;
+        sharedBuilder.module = this->module;
+        sharedBuilder.session = module->session;
+
+        IRBuilder builderStorage;
+        IRBuilder* builder = &builderStorage;
+
+        builder->sharedBuilder = &sharedBuilder;
+        builder->setInsertInto(module->getModuleInst());
+
+        // The packaged-up global parameters will be turned into fields of
+        // a new global IR `struct` type, which we give a name of `GlobalParams`
+        // so that it is identifiable in the output.
+        //
+        // Note: the equivalent in fxc/dxc is the `$Globals` constant buffer.
+        //
+        auto wrapperStructType = builder->createStructType();
+        builder->addNameHintDecoration(wrapperStructType, UnownedTerminatedStringSlice("GlobalParams"));
+
+        // If the computed layout used a bare `struct` type, then we will use
+        // our `GlobalParams` struct as-is, but if the layout involved an
+        // implicitly defined `ConstantBuffer<...>`, this is where we construct
+        // the type `ConstantBuffer<GlobalParams>`.
+        //
+        IRType* wrapperParamType = wrapperStructType;
+        if( globalParameterGroupTypeLayout )
+        {
+            auto wrapperParamGroupType = builder->getConstantBufferType(wrapperStructType);
+            wrapperParamType = wrapperParamGroupType;
+        }
+
+        // Now that we've determined what the type of the new single global parameter
+        // should be, we can go ahead and emit it into the IR module.
+        //
+        // We will call the implicit parameter for all the globals `globalParams`.
+        //
+        IRGlobalParam* wrapperParam = builder->createGlobalParam(wrapperParamType);
+        builder->addLayoutDecoration(wrapperParam, globalScopeVarLayout);
+        builder->addNameHintDecoration(wrapperParam, UnownedTerminatedStringSlice("globalParams"));
+
+        // With the setup work out of the way, we can iterate over the global
+        // parameters that were present in the layout information (they are
+        // represented as the fields of the global-scope `struct` layout).
+        //
+        for( auto fieldLayoutAttr : globalParamsStructTypeLayout->getFieldLayoutAttrs() )
+        {
+            // We expect the IR layout pass to have encoded field per-field
+            // layout so that the "key" for the field is the corresponding
+            // global shader parameter.
+            //
+            auto globalParam = as<IRGlobalParam>(fieldLayoutAttr->getFieldKey());
+            SLANG_ASSERT(globalParam);
+
+            auto globalParamLayout = fieldLayoutAttr->getLayout();
+
+            // If the given parameter doesn't contribute to uniform/ordinary usage, then
+            // we can safely leave it at the global scope and potentially avoid a lot
+            // of complications that might otherwise arise (that is, we don't need to worry
+            // about downstream passes that might have worked for a simple global parameter,
+            // but that would not work for one nested inside a structure.
+            //
+            // TODO: It would be more consistent and robust to *always* wrap up
+            // these global parameters appropriately, and ensure that all the downstream
+            // passes can handle that case, since they would need to do so in general.
+            //
+            if(!globalParamLayout->getTypeLayout()->findSizeAttr(LayoutResourceKind::Uniform) )
+                continue;
+
+            // Once we have decided to do replacement, we need to
+            // set ourselves up to emit the replacement code.
+            //
+            builder->setInsertBefore(globalParam);
+
+            // This global parameter needs to be turned into a field of the global
+            // parameter structure type, and that field will need a key.
+            //
+            auto fieldKey = builder->createStructKey();
+
+            // The new structure field will need to have whatever decorations
+            // had been put on the global parameter (notably including any name hint)
+            //
+            globalParam->transferDecorationsTo(fieldKey);
+
+            // In order to make sure that the existing IR layout information for
+            // the global scope remains valid, we will swap out the key in the
+            // per-field layout information to reference the key we created
+            // instead of the existing parameter (which we will be removing).
+            //
+            fieldLayoutAttr->setOperand(0, fieldKey);
+
+            // Now we can add a field to the `GlobalParams` type that
+            // will stand in for the parameter: it will have the key we
+            // just generated, and the type of the original parameter.
+            //
+            auto globalParamType = globalParam->getFullType();
+            builder->createStructField(wrapperStructType, fieldKey, globalParamType);
+
+            // Next we need to replace the uses of the parameter will
+            // logic to extract the appropriate field from the aggregated
+            // parameter.
+            //
+            // Unlike trivial cases that can work with `replaceAllUsesWith`,
+            // we are going to need to different code for each use, and that
+            // potentially puts us in the bad case of modifying the use-def
+            // information while also iterating it.
+            //
+            // To worka around the problem, we will make a copy of the list of
+            // uses and work with that instead.
+            //
+            List<IRUse*> uses;
+            for(auto use = globalParam->firstUse; use; use = use->nextUse)
+            {
+                uses.add(use);
+            }
+            for( auto use : uses )
+            {
+                // For each use site for the global parameter, we will
+                // insert new code right before the instruction that uses
+                // the parameter.
+                //
+                // TODO: In some cases we might want to emit a single load of
+                // a global parameter at the start of a function, rather
+                // than individual loads at multiple points in the body
+                // of a function. Ideally we can/should annotate the
+                // `globalParams` parameter with the equivalent of `__restrict__`
+                // so that these loads can be merged/moved without concern
+                // for aliasing.
+                //
+                auto user = use->user;
+                builder->setInsertBefore(user);
+
+                IRInst* value = nullptr;
+                if( globalParameterGroupTypeLayout )
+                {
+                    // If the global parameters are being placed in a
+                    // `ConstantBuffer<GlobalParams>`, then we need to
+                    // emit an instruction to compute a pointer to the
+                    // desired field, and then load from it.
+                    //
+                    auto ptrType = builder->getPtrType(globalParamType);
+                    auto fieldAddr = builder->emitFieldAddress(ptrType, wrapperParam, fieldKey);
+                    value = builder->emitLoad(globalParamType, fieldAddr);
+                }
+                else
+                {
+                    // If the global parameters are being bundled in a
+                    // plain old `struct`, then we simple want to emit
+                    // an instruction to extract the desired field.
+                    //
+                    value = builder->emitFieldExtract(globalParamType, wrapperParam, fieldKey);
+                }
+
+                // Whatever replacement value we computed, we need
+                // to install it as the value to be used at the use site.
+                //
+                use->set(value);
+            }
+
+            // Once we've replaced all uses of the global parameter,
+            // we can remove it from the IR module completely.
+            //
+            globalParam->removeAndDeallocate();
+        }
+    }
+};
+
+void collectGlobalUniformParameters(
+    IRModule*       module,
+    IRVarLayout*    globalScopeVarLayout)
+{
+    CollectGlobalUniformParametersContext context;
+    context.module = module;
+    context.globalScopeVarLayout = globalScopeVarLayout;
+
+    context.processModule();
+}
+
+}
diff --git a/source/slang/slang-ir-collect-global-uniforms.h b/source/slang/slang-ir-collect-global-uniforms.h
new file mode 100644
index 000000000..5a173f4db
--- /dev/null
+++ b/source/slang/slang-ir-collect-global-uniforms.h
@@ -0,0 +1,18 @@
+// slang-ir-collect-global-uniforms.h
+#pragma once
+
+#include "slang-compiler.h"
+
+namespace Slang
+{
+struct IRModule;
+struct IRVarLayout;
+
+    /// Collect global-scope shader parameters that use uniform/ordinary
+    /// storage into a single `struct` (possibly wrapped in a constant buffer).
+    ///
+void collectGlobalUniformParameters(
+    IRModule*       module,
+    IRVarLayout*    globalScopeVarLayout);
+
+}
diff --git a/source/slang/slang-ir-entry-point-uniforms.cpp b/source/slang/slang-ir-entry-point-uniforms.cpp
index 388a7004d..9c3c029a5 100644
--- a/source/slang/slang-ir-entry-point-uniforms.cpp
+++ b/source/slang/slang-ir-entry-point-uniforms.cpp
@@ -98,12 +98,6 @@ struct MoveEntryPointUniformParametersToGlobalScope
     //
     IRModule* module;
 
-    // The target can determine how a variable is moved out into global scope
-    CodeGenTarget codeGenTarget;
-
-    // If true the target needs constant buffer wrapping (for uniforms say)
-    bool targetNeedsConstantBuffer;
-
     // We will process a whole module by visiting all
     // its global functions, looking for entry points.
     //
@@ -168,7 +162,7 @@ struct MoveEntryPointUniformParametersToGlobalScope
         // an explicit IR constant buffer for that wrapper, 
         //
         auto entryPointParamsLayout = entryPointLayout->getParamsLayout();
-        bool needConstantBuffer = targetNeedsConstantBuffer && as<IRParameterGroupTypeLayout>(entryPointParamsLayout->getTypeLayout()); 
+        bool needConstantBuffer = as<IRParameterGroupTypeLayout>(entryPointParamsLayout->getTypeLayout()) != nullptr; 
 
         auto entryPointParamsStructLayout = getScopeStructLayout(entryPointLayout);
 
@@ -237,6 +231,7 @@ struct MoveEntryPointUniformParametersToGlobalScope
                 //
                 builder->setInsertBefore(func);
                 paramStructType = builder->createStructType();
+                builder->addNameHintDecoration(paramStructType, UnownedTerminatedStringSlice("EntryPointParams"));
 
                 if( needConstantBuffer )
                 {
@@ -252,9 +247,6 @@ struct MoveEntryPointUniformParametersToGlobalScope
                     // an instance of `paramStructType`.
                     //
                     globalParam = builder->createGlobalParam(paramStructType);
-
-                    // Mark that this global comes from the entry point
-                    builder->addDecoration(globalParam, kIROp_EntryPointParamDecoration);
                 }
 
                 // No matter what, the global shader parameter should have the layout
@@ -262,6 +254,30 @@ struct MoveEntryPointUniformParametersToGlobalScope
                 // contained parameters will end up in the right place(s).
                 //
                 builder->addLayoutDecoration(globalParam, entryPointParamsLayout);
+
+                // We add a name hint to the global parameter so that it will
+                // emit to more readable code when referenced.
+                //
+                builder->addNameHintDecoration(globalParam, UnownedTerminatedStringSlice("entryPointParams"));
+
+                // We also decorate the parameter for the entry-point parameters
+                // so that we can find it again in downstream passes (like emit
+                // for CPU/CUDA) that might want to treat entry-point parameters
+                // different from other cases.
+                //
+                // TODO: Once we have support for multiple entry points to be emitted
+                // at once, we need a way to associate these per-entry-point parameters
+                // more closely with the original entry point. The two easiest options
+                // are:
+                //
+                // 1. Don't move the new aggregate parameter to the global scope
+                // on those targets, and instead keep it as a parameter of the
+                // entry point.
+                //
+                // 2. Use a decoration on the entry point itself to point at the
+                // global parameter for its per-entry-point parameter data.
+                //
+                builder->addDecoration(globalParam, kIROp_EntryPointParamDecoration);
             }
 
             // Now that we've ensured the global `struct` type and shader paramter
@@ -285,7 +301,7 @@ struct MoveEntryPointUniformParametersToGlobalScope
             // the linker. After all, this pass is traversing the same information
             // anyway, so it could do the work while it is here...
             //
-            auto paramFieldKey = entryPointParamsStructLayout->getFieldLayoutAttrs()[paramIndex]->getFieldKey();
+            auto paramFieldKey = cast<IRStructKey>(entryPointParamsStructLayout->getFieldLayoutAttrs()[paramIndex]->getFieldKey());
 
             auto paramField = builder->createStructField(paramStructType, paramFieldKey, paramType);
             SLANG_UNUSED(paramField);
@@ -435,32 +451,10 @@ struct MoveEntryPointUniformParametersToGlobalScope
 };
 
 void moveEntryPointUniformParamsToGlobalScope(
-    IRModule*   module,
-    CodeGenTarget target)
+    IRModule*   module)
 {
     MoveEntryPointUniformParametersToGlobalScope context;
-    
     context.module = module;
-    context.codeGenTarget = target;
-    context.targetNeedsConstantBuffer = true;
-
-    // Check if this target needs constant buffer wrapping
-    switch (target)
-    {
-        case CodeGenTarget::CPPSource:
-        case CodeGenTarget::CSource:
-        case CodeGenTarget::Executable:
-        case CodeGenTarget::SharedLibrary:
-        case CodeGenTarget::HostCallable:
-        case CodeGenTarget::CUDASource:
-        case CodeGenTarget::PTX:
-        {
-            context.targetNeedsConstantBuffer = false;
-            break;
-        }
-        default: break;
-    }
-
     context.processModule();
 }
 
diff --git a/source/slang/slang-ir-entry-point-uniforms.h b/source/slang/slang-ir-entry-point-uniforms.h
index 0e978b9eb..3b6a0743b 100644
--- a/source/slang/slang-ir-entry-point-uniforms.h
+++ b/source/slang/slang-ir-entry-point-uniforms.h
@@ -9,7 +9,6 @@ struct IRModule;
 
     /// Move any uniform parameters of entry points to the global scope instead.
 void moveEntryPointUniformParamsToGlobalScope(
-    IRModule*   module,
-    CodeGenTarget target);
+    IRModule*   module);
 
 }
diff --git a/source/slang/slang-ir-insts.h b/source/slang/slang-ir-insts.h
index 159021412..93e709c45 100644
--- a/source/slang/slang-ir-insts.h
+++ b/source/slang/slang-ir-insts.h
@@ -793,9 +793,9 @@ struct IRStructFieldLayoutAttr : IRAttr
 {
     IR_LEAF_ISA(StructFieldLayoutAttr)
 
-    IRStructKey* getFieldKey()
+    IRInst* getFieldKey()
     {
-        return cast<IRStructKey>(getOperand(0));
+        return getOperand(0);
     }
 
     IRVarLayout* getLayout()
@@ -836,7 +836,7 @@ struct IRStructTypeLayout : IRTypeLayout
             : Super::Builder(irBuilder)
         {}
 
-        void addField(IRStructKey* key, IRVarLayout* layout)
+        void addField(IRInst* key, IRVarLayout* layout)
         {
             FieldInfo info;
             info.key = key;
@@ -855,7 +855,7 @@ struct IRStructTypeLayout : IRTypeLayout
 
         struct FieldInfo
         {
-            IRStructKey* key;
+            IRInst* key;
             IRVarLayout* layout;
         };
 
@@ -2099,7 +2099,7 @@ struct IRBuilder
     IRPendingLayoutAttr* getPendingLayoutAttr(
         IRLayout* pendingLayout);
     IRStructFieldLayoutAttr* getFieldLayoutAttr(
-        IRStructKey*    key,
+        IRInst*         key,
         IRVarLayout*    layout);
     IRCaseTypeLayoutAttr* getCaseTypeLayoutAttr(
         IRTypeLayout*   layout);
diff --git a/source/slang/slang-ir-link.cpp b/source/slang/slang-ir-link.cpp
index 3935eab07..704917629 100644
--- a/source/slang/slang-ir-link.cpp
+++ b/source/slang/slang-ir-link.cpp
@@ -1403,7 +1403,8 @@ LinkedIR linkIR(
     // responsible for associating layout information to those
     // global symbols via decorations.
     //
-    insertGlobalValueSymbols(sharedContext, targetProgram->getExistingIRModuleForLayout());
+    auto irModuleForLayout = targetProgram->getExistingIRModuleForLayout();
+    insertGlobalValueSymbols(sharedContext, irModuleForLayout);
 
     auto context = state->getContext();
 
@@ -1454,6 +1455,18 @@ LinkedIR linkIR(
     auto entryPointMangledName = program->getEntryPointMangledName(entryPointIndices[0]);
     auto irEntryPoint = specializeIRForEntryPoint(context, entryPointMangledName);
 
+    // Layout information for global shader parameters is also required,
+    // and in particular every global parameter that is part of the layout
+    // should be present in the initial IR module so that steps that
+    // need to operate on all the global parameters can do so.
+    //
+    IRVarLayout* irGlobalScopeVarLayout = nullptr;
+    if( auto irGlobalScopeLayoutDecoration = irModuleForLayout->getModuleInst()->findDecoration<IRLayoutDecoration>() )
+    {
+        auto irOriginalGlobalScopeVarLayout = irGlobalScopeLayoutDecoration->getLayout();
+        irGlobalScopeVarLayout = cast<IRVarLayout>(cloneValue(context, irOriginalGlobalScopeVarLayout));
+    }
+
     // Bindings for global generic parameters are currently represented
     // as stand-alone global-scope instructions in the IR module for
     // `SpecializedComponentType`s. These instructions are required for
@@ -1516,6 +1529,7 @@ LinkedIR linkIR(
     LinkedIR linkedIR;
     linkedIR.module = state->irModule;
     linkedIR.entryPoint = irEntryPoint;
+    linkedIR.globalScopeVarLayout = irGlobalScopeVarLayout;
     return linkedIR;
 }
 
diff --git a/source/slang/slang-ir-link.h b/source/slang/slang-ir-link.h
index dc9ad50a0..ce2af64d8 100644
--- a/source/slang/slang-ir-link.h
+++ b/source/slang/slang-ir-link.h
@@ -5,10 +5,13 @@
 
 namespace Slang
 {
+    struct IRVarLayout;
+
     struct LinkedIR
     {
         RefPtr<IRModule>    module;
         IRFunc*             entryPoint;
+        IRVarLayout*        globalScopeVarLayout;
     };
 
 
diff --git a/source/slang/slang-ir.cpp b/source/slang/slang-ir.cpp
index ef5ecb959..25bf0f9f3 100644
--- a/source/slang/slang-ir.cpp
+++ b/source/slang/slang-ir.cpp
@@ -3665,7 +3665,7 @@ namespace Slang
     }
 
     IRStructFieldLayoutAttr* IRBuilder::getFieldLayoutAttr(
-        IRStructKey*    key,
+        IRInst*         key,
         IRVarLayout*    layout)
     {
         IRInst* operands[] = { key, layout };
@@ -4985,6 +4985,12 @@ namespace Slang
         if(as<IRConstant>(this))
             return false;
 
+        if(as<IRLayout>(this))
+            return false;
+
+        if(as<IRAttr>(this))
+            return false;
+
         switch(op)
         {
         // By default, assume that we might have side effects,
diff --git a/source/slang/slang-lower-to-ir.cpp b/source/slang/slang-lower-to-ir.cpp
index 1ba99dffb..adebefc8b 100644
--- a/source/slang/slang-lower-to-ir.cpp
+++ b/source/slang/slang-lower-to-ir.cpp
@@ -7583,7 +7583,7 @@ struct IRLayoutGenContext : IRGenContext
     {}
 
         /// Cache for custom key instructions used for entry-point parameter layout information.
-    Dictionary<ParamDecl*, IRStructKey*> mapEntryPointParamToKey;
+    Dictionary<ParamDecl*, IRInst*> mapEntryPointParamToKey;
 };
 
     /// Lower an AST-level type layout to an IR-level type layout.
@@ -7654,7 +7654,7 @@ IRTypeLayout* lowerTypeLayout(
         {
             auto fieldDecl = fieldLayout->varDecl;
 
-            IRStructKey* irFieldKey = nullptr;
+            IRInst* irFieldKey = nullptr;
             if(auto paramDecl = as<ParamDecl>(fieldDecl) )
             {
                 // There is a subtle special case here.
@@ -7699,9 +7699,8 @@ IRTypeLayout* lowerTypeLayout(
             }
             else
             {
-                IRInst* irFieldKeyInst = getSimpleVal(context,
+                irFieldKey = getSimpleVal(context,
                     ensureDecl(context, fieldDecl));
-                irFieldKey = as<IRStructKey>(irFieldKeyInst);
             }
             SLANG_ASSERT(irFieldKey);
 
@@ -7762,9 +7761,9 @@ IRTypeLayout* lowerTypeLayout(
 
 IRVarLayout* lowerVarLayout(
     IRLayoutGenContext* context,
-    VarLayout*          varLayout)
+    VarLayout*          varLayout,
+    IRTypeLayout*       irTypeLayout)
 {
-    auto irTypeLayout = lowerTypeLayout(context, varLayout->typeLayout);
     IRVarLayout::Builder irLayoutBuilder(context->irBuilder, irTypeLayout);
 
     for( auto resInfo : varLayout->resourceInfos )
@@ -7806,6 +7805,14 @@ IRVarLayout* lowerVarLayout(
     return irLayoutBuilder.build();
 }
 
+IRVarLayout* lowerVarLayout(
+    IRLayoutGenContext* context,
+    VarLayout*          varLayout)
+{
+    auto irTypeLayout = lowerTypeLayout(context, varLayout->typeLayout);
+    return lowerVarLayout(context, varLayout, irTypeLayout);
+}
+
     /// Handle the lowering of an entry-point result layout to the IR
 IRVarLayout* lowerEntryPointResultLayout(
     IRLayoutGenContext* context,
@@ -7893,10 +7900,12 @@ RefPtr<IRModule> TargetProgram::createIRModuleForLayout(DiagnosticSink* sink)
 
     // Okay, now we need to walk through and decorate everything.
     auto globalStructLayout = getScopeStructLayout(programLayout);
-    for(auto globalVarPair : globalStructLayout->mapVarToLayout)
+
+    IRStructTypeLayout::Builder globalStructTypeLayoutBuilder(builder);
+
+    for(auto varLayout : globalStructLayout->fields)
     {
-        auto varDecl = globalVarPair.Key;
-        auto varLayout = globalVarPair.Value;
+        auto varDecl = varLayout->varDecl;
 
         // Ensure that an `[import(...)]` declaration for the variable
         // has been emitted to this module, so that we will have something
@@ -7909,7 +7918,36 @@ RefPtr<IRModule> TargetProgram::createIRModuleForLayout(DiagnosticSink* sink)
         // Now attach the decoration to the variable.
         //
         builder->addLayoutDecoration(irVar, irLayout);
+
+        // Also add this to our mapping for the global-scope structure type
+        globalStructTypeLayoutBuilder.addField(irVar, irLayout);
     }
+    auto irGlobalStructTypeLayout = _lowerTypeLayoutCommon(context, &globalStructTypeLayoutBuilder, globalStructLayout);
+
+    auto globalScopeVarLayout = programLayout->parametersLayout;
+    auto globalScopeTypeLayout = globalScopeVarLayout->typeLayout;
+    IRTypeLayout* irGlobalScopeTypeLayout = irGlobalStructTypeLayout;
+    if( auto paramGroupTypeLayout = as<ParameterGroupTypeLayout>(globalScopeTypeLayout) )
+    {
+        IRParameterGroupTypeLayout::Builder globalParameterGroupTypeLayoutBuilder(builder);
+
+        auto irElementTypeLayout = irGlobalStructTypeLayout;
+        auto irElementVarLayout = lowerVarLayout(context, paramGroupTypeLayout->elementVarLayout, irElementTypeLayout);
+
+        globalParameterGroupTypeLayoutBuilder.setContainerVarLayout(
+            lowerVarLayout(context, paramGroupTypeLayout->containerVarLayout));
+        globalParameterGroupTypeLayoutBuilder.setElementVarLayout(irElementVarLayout);
+        globalParameterGroupTypeLayoutBuilder.setOffsetElementTypeLayout(
+            lowerTypeLayout(context, paramGroupTypeLayout->offsetElementTypeLayout));
+
+        auto irParamGroupTypeLayout = _lowerTypeLayoutCommon(context, &globalParameterGroupTypeLayoutBuilder, paramGroupTypeLayout);
+
+        irGlobalScopeTypeLayout = irParamGroupTypeLayout;
+    }
+
+    auto irGlobalScopeVarLayout = lowerVarLayout(context, globalScopeVarLayout, irGlobalScopeTypeLayout);
+
+    builder->addLayoutDecoration(irModule->getModuleInst(), irGlobalScopeVarLayout);
 
     for( auto entryPointLayout : programLayout->entryPoints )
     {
diff --git a/source/slang/slang-parameter-binding.cpp b/source/slang/slang-parameter-binding.cpp
index ee3ef1234..6f2a1632e 100644
--- a/source/slang/slang-parameter-binding.cpp
+++ b/source/slang/slang-parameter-binding.cpp
@@ -689,6 +689,7 @@ static RefPtr<VarLayout> _createVarLayout(
     if(auto pendingDataTypeLayout = typeLayout->pendingDataTypeLayout)
     {
         RefPtr<VarLayout> pendingVarLayout = new VarLayout();
+        pendingVarLayout->varDecl = varDeclRef;
         pendingVarLayout->typeLayout = pendingDataTypeLayout;
         varLayout->pendingVarLayout = pendingVarLayout;
     }
@@ -2857,29 +2858,41 @@ static void collectParameters(
     program->acceptVisitor(&visitor, nullptr);
 }
 
-    /// Emit a diagnostic about a uniform parameter at global scope.
+    /// Emit a diagnostic about a uniform/ordinary parameter at global scope.
 void diagnoseGlobalUniform(
     SharedParameterBindingContext*  sharedContext,
     VarDeclBase*                    varDecl)
 {
-    // It is entirely possible for Slang to support uniform parameters at the global scope,
-    // by bundling them into an implicit constant buffer, and indeed the layout algorithm
-    // implemented in this file computes a layout *as if* the Slang compiler does just that.
+    // This subroutine gets invoked if a shader parameter containing
+    // "ordinary" data (sometimes just called "uniform" data) is present
+    // at the global scope.
     //
-    // The missing link is the downstream IR and code generation steps, where we would need
-    // to collect all of the global-scope uniforms into a common `struct` type and then
-    // create a new constant buffer parameter over that type.
+    // Slang can support such parameters by aggregating them into
+    // an implicit constant buffer, but it is also common for programmers
+    // to accidentally declare a global-scope shader parameter when they
+    // meant to declare a global variable instead:
     //
-    // For now it is easier to simply ban this case, since most shader authors have
-    // switched to modern HLSL/GLSL style with `cbuffer` or `uniform` block declarations.
+    //      int gCounter = 0; // this is a shader parameter, not a global
     //
-    // TODO: In the long run it may be best to require *all* global-scope shader parameters
-    // to be marked with a keyword (e.g., `uniform`) so that ordinary global variable syntax can be
-    // used safely.
+    // In order to avoid mistakes, we'd like to warn the user when
+    // they write code like the above, and hint to them that they
+    // should make their intention more explicit with a keyword:
     //
-    getSink(sharedContext)->diagnose(varDecl, Diagnostics::globalUniformsNotSupported, varDecl->getName());
+    //      static int gCounter = 0; // this is now a (static) global
+    //
+    //      uniform int gCounter; // this is now explicitly a shader parameter
+    //
+    // We skip the diagnostic whenever the variable was explicitly `uniform`,
+    // under the assumption that the programmer who added that modifier
+    // knew what they were opting into.
+    //
+    if(varDecl->hasModifier<HLSLUniformModifier>())
+        return;
+
+    getSink(sharedContext)->diagnose(varDecl, Diagnostics::globalUniformNotExpected, varDecl->getName());
 }
 
+
 static int _calcTotalNumUsedRegistersForLayoutResourceKind(ParameterBindingContext* bindingContext, LayoutResourceKind kind)
 {
     int numUsed = 0;
diff --git a/source/slang/slang.vcxproj b/source/slang/slang.vcxproj
index 5850bc2b8..95f98e4ec 100644
--- a/source/slang/slang.vcxproj
+++ b/source/slang/slang.vcxproj
@@ -223,6 +223,7 @@
     <ClInclude Include="slang-ir-bind-existentials.h" />
     <ClInclude Include="slang-ir-byte-address-legalize.h" />
     <ClInclude Include="slang-ir-clone.h" />
+    <ClInclude Include="slang-ir-collect-global-uniforms.h" />
     <ClInclude Include="slang-ir-constexpr.h" />
     <ClInclude Include="slang-ir-dce.h" />
     <ClInclude Include="slang-ir-dominators.h" />
@@ -313,6 +314,7 @@
     <ClCompile Include="slang-ir-bind-existentials.cpp" />
     <ClCompile Include="slang-ir-byte-address-legalize.cpp" />
     <ClCompile Include="slang-ir-clone.cpp" />
+    <ClCompile Include="slang-ir-collect-global-uniforms.cpp" />
     <ClCompile Include="slang-ir-constexpr.cpp" />
     <ClCompile Include="slang-ir-dce.cpp" />
     <ClCompile Include="slang-ir-dominators.cpp" />
diff --git a/source/slang/slang.vcxproj.filters b/source/slang/slang.vcxproj.filters
index fdc1bf45d..561599a8f 100644
--- a/source/slang/slang.vcxproj.filters
+++ b/source/slang/slang.vcxproj.filters
@@ -120,6 +120,9 @@
     <ClInclude Include="slang-ir-clone.h">
       <Filter>Header Files</Filter>
     </ClInclude>
+    <ClInclude Include="slang-ir-collect-global-uniforms.h">
+      <Filter>Header Files</Filter>
+    </ClInclude>
     <ClInclude Include="slang-ir-constexpr.h">
       <Filter>Header Files</Filter>
     </ClInclude>
@@ -386,6 +389,9 @@
     <ClCompile Include="slang-ir-clone.cpp">
       <Filter>Source Files</Filter>
     </ClCompile>
+    <ClCompile Include="slang-ir-collect-global-uniforms.cpp">
+      <Filter>Source Files</Filter>
+    </ClCompile>
     <ClCompile Include="slang-ir-constexpr.cpp">
       <Filter>Source Files</Filter>
     </ClCompile>