Add support for global uniform shader parameters (#1433)

* Adding support for global uniform shader parameters

This change adds support for Slang programmers to declare shader parameters of "ordinary" types at global scope:

```hlsl
uniform float gScaleFactor;

void main() { ... *= gScaleFactor; ... }
```

The generated HLSL/GLSL/DXIL/SPIR-V output will be something along the lines of:

```hlsl
struct GlobalParams
{
    float gScaleFactor;
}

cbuffer globalParams
{
    GlobalParams globalParams;
}

void main() { ... *= globalParams.gScaleFactor; ... }
```

The binding information used for the implicit `globalParams` constant buffer will be determined by the existing implicit parameter binding logic (which already had support for this kind of transformation).

The reason this change is being pursued right now is because it is one step toward removing the implicit `KernelContext` type that is used to wrap the generated code for our CPU and CUDA C++ targets. Handling global-scope parameters of ordinary type requires an IR pass that synthesizes the `GlobalParams` structure type above, and that step ends up removing the need for the similar `UniformState` structure that was being used in the CPU/CUDA emit logic.

A more detailed guide to the changes included follows:

* The diagnostic for a global-scope variable that is implicitly a shader parameter was kept, but changed to a warning. Users can opt out of the warning by decorating their parameter as a `uniform` (since that keyword is already being used to mark entry-point parameters that should be treated as uniform shader parameters).

* To simplify the task of finding the global shader parameters, the `CLikeSourceEmitter` type has been given an `m_irModule` member. The previous emit logic for `UniformState` was having to do a roundabout solution involving the `EmitAction`s to deal with not having direct access to the module.

* Removed a few dead declarations in the emit logic (related to a much earlier point where emit was based on the AST instead of the IR).

* Made the computation of type names in C++ emit take into account `ConstantBuffer<T>` and `ParameterBlock<T>`. As far as I can tell, these were being handled with some special-case hacks in the emit logic instead of being supported more fundamentally. It might actually be good to pass these through as `ConstantBuffer<T>` and `ParameterBlock<T>` in the C++ output, and allow the prelude to customize their translation (defaulting to defining them as `T*`).

* Removed the special-case C++ emit logic for references to global shader parameters. There are now at most two global shader parameters to deal with, and the default emit logic (referring to them by name) does the Right Thing.

* Changed the handling of entry points for C++ (both CPU and CUDA) so that it handles the bundled-up shader paameters for the global and entry-point scopes the same way. The main complication here is OptiX, where parameter data is passed very differently than it is for CUDA compute kernels.

* Reverted changes to `ir-entry-point-uniforms` that had made its logic depend on the compilation target. The parameter binding logic was already responsible for deciding if a given target needed to wrap up its entry-point parameters in a constant buffer, and the IR pass was respecting that layout information. The current workaround had been removing the `ConstantBuffer<T>` indirection from this IR pass for CPU/CUDA, but then reintroducing the same indirection later on in the emit step.

* Added an explicit IR pass with the task of collecting global-scope parameters of uniform/ordinary type and packaging them up into a `struct`, and then optionally packaging that `struct` up in a constant buffer. This pass bases its decisions on the IR layout information that was already computed, so it should match whatever policy choices were made at the layout level.

* Changed the "key" operand on IR `struct` layout information to not assume an `IRStructKey`. The problem here is that the global scope gets a `StructTypeLayout` to represent its members, and this is convenient (rather than having to always special-case logic that handles the global scope), but the "fields" of that struct are global variables which do not have `IRStructKey`s associated with them. The simplest solution is to use the variables themselves as the keys, which required removing the assumption in the IR encoding.

* Updated the IR layout process to compute a layout for the global scope of an entire program, and to attach that to the `IRModule` via a decoration. Updated the IR linking process to carry through that decoration to the linked output. This is necessary so that the IR pass that transforms global parameters can access the global-scope layout information.

An important concern with this approach is that the contents and layout of the monolithic `GlobalParams` structure depends on the exact set of modules that were linked (and the order in which they were specified, in some cases). This isn't really a new thing with this change, but it becomes more important as we start to think of how to generalize things to better support separate compilation and linking.

There are changes that can (and should) be made to the way that IR layouts are computed for programs (e.g., so that we compute layout per-module and then combine them rather than as a whole-program step). In this case, the problem of forming the combined/linked global layout can be moved down the IR level and not be reliant on AST-level information.

Just changing the way layout and linking interact would not change the fundamental problem that global shader parameters as they currently exist in Slang/HLSL/GLSL are not readily compatible with true separate compilation. We either need to find a solution strategy that we can apply to allow existing shaders to work with separate compilation *or* we need to incrementally work toward removing support for global-scope shader parameters in favor of explicit entry-point parameters in all cases.

* fixup: missing files

* fixup: comment the new code

24 files changed, 655 insertions, 256 deletions

diff --git a/examples/heterogeneous-hello-world/heterogeneous-hello-world.vcxproj b/examples/heterogeneous-hello-world/heterogeneous-hello-world.vcxproj
index d80fbc30b..8a53a719a 100644
--- a/examples/heterogeneous-hello-world/heterogeneous-hello-world.vcxproj
+++ b/examples/heterogeneous-hello-world/heterogeneous-hello-world.vcxproj
@@ -182,4 +182,4 @@
   <Import Project="$(VCTargetsPath)\Microsoft.Cpp.targets" />
   <ImportGroup Label="ExtensionTargets">
   </ImportGroup>
-</Project> 
\ No newline at end of file
+</Project>
\ No newline at end of file
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>
diff --git a/tests/cross-compile/unknown-image-format.slang.glsl b/tests/cross-compile/unknown-image-format.slang.glsl
index 9995bba72..329405ab6 100644
--- a/tests/cross-compile/unknown-image-format.slang.glsl
+++ b/tests/cross-compile/unknown-image-format.slang.glsl
@@ -30,36 +30,36 @@ layout(binding = 1, set = 1)
 uniform image2D gBlock_explicitFormat_0;
 
 layout(binding = 3)
-uniform image2D _S2;
+uniform image2D entryPointParams_noFormat_0;
 
 layout(rgba16f)
 layout(binding = 4)
-uniform image2D _S3;
+uniform image2D entryPointParams_explicitFormat_0;
 
 layout(location = 0)
-out vec4 _S4;
+out vec4 _S2;
 
 void main()
 {
     const vec4 result_0 = vec4(0);
 
-    float _S5 = (imageLoad((gNoFormat_0), ivec2((C_0._data.index_0))).x);
-    vec4 result_1 = result_0 + _S5;
+    float _S3 = (imageLoad((gNoFormat_0), ivec2((C_0._data.index_0))).x);
+    vec4 result_1 = result_0 + _S3;
 
-    float _S6 = (imageLoad((gExplicitFormat_0), ivec2((C_0._data.index_0))).x);
-    vec4 result_2 = result_1 + _S6;
+    float _S4 = (imageLoad((gExplicitFormat_0), ivec2((C_0._data.index_0))).x);
+    vec4 result_2 = result_1 + _S4;
 
-    vec4 _S7 = (imageLoad((gBlock_noFormat_0), ivec2((C_0._data.index_0))));
-    vec4 result_3 = result_2 + _S7;
+    vec4 _S5 = (imageLoad((gBlock_noFormat_0), ivec2((C_0._data.index_0))));
+    vec4 result_3 = result_2 + _S5;
 
-    vec4 _S8 = (imageLoad((gBlock_explicitFormat_0), ivec2((C_0._data.index_0))));
-    vec4 result_4 = result_3 + _S8;
+    vec4 _S6 = (imageLoad((gBlock_explicitFormat_0), ivec2((C_0._data.index_0))));
+    vec4 result_4 = result_3 + _S6;
 
-    vec4 _S9 = (imageLoad((_S2), ivec2((C_0._data.index_0))));
-    vec4 result_5 = result_4 + _S9;
+    vec4 _S7 = (imageLoad((entryPointParams_noFormat_0), ivec2((C_0._data.index_0))));
+    vec4 result_5 = result_4 + _S7;
 
-    vec4 _S10 = (imageLoad((_S3), ivec2((C_0._data.index_0))));
-    _S4 = result_5 + _S10;
+    vec4 _S8 = (imageLoad((entryPointParams_explicitFormat_0), ivec2((C_0._data.index_0))));
+    _S2 = result_5 + _S8;
 
     return;
 }
diff --git a/tests/diagnostics/global-uniform.slang b/tests/diagnostics/global-uniform.slang
index 6b17016f2..de4537115 100644
--- a/tests/diagnostics/global-uniform.slang
+++ b/tests/diagnostics/global-uniform.slang
@@ -2,8 +2,8 @@
 //DIAGNOSTIC_TEST:SIMPLE:-target hlsl
 //DIAGNOSTIC_TEST:COMMAND_LINE_SIMPLE:-target hlsl
 
-// Any attempt to declare a global variable that actually declares a
-// global uniform should be diagnosed as unsupported.
+// An attempt to declare a global variable that actually declares a
+// global shader parameter should be diagnosed, unless `uniform` was used.
 
 uniform float a;
 
diff --git a/tests/diagnostics/global-uniform.slang.expected b/tests/diagnostics/global-uniform.slang.expected
index a77144c4f..6f69ec0c5 100644
--- a/tests/diagnostics/global-uniform.slang.expected
+++ b/tests/diagnostics/global-uniform.slang.expected
@@ -1,8 +1,7 @@
-result code = -1
+result code = 0
 standard error = {
-tests/diagnostics/global-uniform.slang(8): error 39016: 'a' 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.
-tests/diagnostics/global-uniform.slang(10): error 39016: 'b' 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.
-tests/diagnostics/global-uniform.slang(13): error 39016: 'c' 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.
+tests/diagnostics/global-uniform.slang(10): warning 39016: 'b' 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.
+tests/diagnostics/global-uniform.slang(13): warning 39016: 'c' 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.
 }
 standard output = {
 }
diff --git a/tests/language-feature/shader-params/global-uniform-params.slang b/tests/language-feature/shader-params/global-uniform-params.slang
new file mode 100644
index 000000000..6b4e5a834
--- /dev/null
+++ b/tests/language-feature/shader-params/global-uniform-params.slang
@@ -0,0 +1,28 @@
+// global-uniform-params.slang
+
+//TEST(compute):COMPARE_COMPUTE:
+
+// Test that code can use uniform parameters
+// of "ordinary" type declared at the global scope
+
+//TEST_INPUT:cbuffer(data=[256 1]):name=$Globals
+uniform int a;
+uniform int b;
+
+int test(int val)
+{
+	return a*(val+1) + b*(val+2);
+}
+
+//TEST_INPUT:ubuffer(data=[0 0 0 0], stride=4):out,name=outputBuffer
+RWStructuredBuffer<int> outputBuffer;
+
+[numthreads(4, 1, 1)]
+[shader("compute")]
+void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID)
+{
+    uint tid = dispatchThreadID.x;
+    int inVal = tid;
+    int outVal = test(inVal);
+    outputBuffer[tid] = outVal;
+}
diff --git a/tests/language-feature/shader-params/global-uniform-params.slang.expected.txt b/tests/language-feature/shader-params/global-uniform-params.slang.expected.txt
new file mode 100644
index 000000000..4cf6581b3
--- /dev/null
+++ b/tests/language-feature/shader-params/global-uniform-params.slang.expected.txt
@@ -0,0 +1,4 @@
+102
+203
+304
+405