Address issues with GLSL style global in/out vars (#6669) (#6998)

* Address issues with GLSL style global in/out vars (#6669)

Asserts and segfaults were observed trying to compile a simple
vertex shader like:

````
in int2 inPos;

[shader("vertex")]
main(uniform int2 test1, int2 test2, out float4 pos: SV_Position)
void main()
{
    // Bogus use of all input vars to prevent optimizing out.
    pos = float4(inPos.x, test1.x, test2.y, 0);
}
````

Further investigation found that while replacing "uniform int2 test1"
with "int2 test1" allowed for successful compilation, the resulting
output shader would have overlapping location qualifiers. For example,
compiling the above with "int2 test1" to glsl might give:

````
...
layout(location = 0) in ivec2 test1_0;
layout(location = 1) in ivec2 test2_0;
layout(location = 0) in ivec2 translatedGlobalParams_inPos_0;
...
````

This was because Slang does not actually support mixing GLSL style global
in/out vars and entry point params. However, this is never checked for
or noted in documentation. Slang source also assumes input shaders do not
mix these and these assumptions ultimately led to the observed asserts
and seg faults when using uniform entry point params.

This change makes updates to throw an error when the compiler detects that
it is trying to translate global in/out variables into entry point params
when an entry point already contains parameters, allowing for compilation
to fail gracefully.

Certain tests have been updated to avoid mixing GLSL style global in/out
vars and entry point params. This was mostly for tests that were using
functions like WaveGetLaneIndex which use global in vars for certain
platforms (see __builtinWaveLaneIndex).

* Address issues with GLSL style global in/out vars - updates 1 (#6669)

Update addresses review feedback to support mixing GLSL-flavored global
in/out vars and entrypoint parameters when either all global in/out vars
or all entry point params have a system value binding semantic.

* Address issues with GLSL style global in/out vars - updates 2 (#6669)

This update attempts to actually allow mixing GLSL style global in
vars and entry point vars.

Change attempts to recalculate offsets when adding the global input
vars into the recreated entry point params layout.

Additional updates were made to:

-resolve further issues uncovered with entry point uniform params.

-Address improper use of SV_DispatchThreadID in wave-get-lane-index.slang
 for metal. "thread_position_in_grid" is not supported for signed integer
 scalars or vectors.

-Fix a spirv casting conflict due to the implementation of
 gl_PrimitiveID.get conflicting with PrimitiveIndex().

-Add a call to remove a global var in replaceUsesOfGlobalVar(). The global
 var is already replaced in this function and keeping it around can prevent
 it from being cleaned up by DCE if it still has decorations.

* format code

---------

Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>

7 files changed, 260 insertions, 73 deletions

diff --git a/source/slang/glsl.meta.slang b/source/slang/glsl.meta.slang
index 88c90a777..9b0a28a31 100644
--- a/source/slang/glsl.meta.slang
+++ b/source/slang/glsl.meta.slang
@@ -4764,7 +4764,12 @@ public property uint3 gl_LaunchSizeEXT
     }
 }
 
-internal in int __gl_PrimitiveID : SV_PrimitiveID;
+// While the glsl spec defines gl_PrimitiveID as an "in int",
+// we need to use uint here and cast to an int below. This is due
+// __gl_PrimitiveID being picked up as a global variable regardless
+// of the stage used below. Using int __gl_PrimitiveID leads to a
+// casting conflict due to the spirv implementation of PrimitiveIndex().
+internal in uint __gl_PrimitiveID : SV_PrimitiveID;
 
 public property int gl_PrimitiveID 
 {
@@ -4779,7 +4784,7 @@ public property int gl_PrimitiveID
             setupExtForRayTracingBuiltIn();
             return PrimitiveIndex();
         default:
-            return __gl_PrimitiveID;   
+            return __intCast<int32_t>(__gl_PrimitiveID);
         }
     }
 }
diff --git a/source/slang/slang-ir-entry-point-uniforms.cpp b/source/slang/slang-ir-entry-point-uniforms.cpp
index 12dcd187a..c5e91c54d 100644
--- a/source/slang/slang-ir-entry-point-uniforms.cpp
+++ b/source/slang/slang-ir-entry-point-uniforms.cpp
@@ -101,43 +101,76 @@ namespace Slang
 // whether a parameter represents a varying input rather than
 // a uniform parameter.
 
-
-// In order to determine whether a parameter is varying based on its
-// layout, we need to know which resource kinds represent varying
-// shader parameters.
+// Setup some flags that will be used below to check for varying
+// resource kinds. Ordinary varying input/output + Ray-tracing shader
+// input/output kinds will be considered varying.
 //
-bool isVaryingResourceKind(LayoutResourceKind kind)
+// Note: The set of cases that are considered
+// varying here would need to be extended if we
+// add more fine-grained resource kinds (e.g.,
+// if we ever add an explicit resource kind
+// for geometry shader output streams).
+static const LayoutResourceKindFlags flags =
+    LayoutResourceKindFlag::make(LayoutResourceKind::VaryingInput) |
+    LayoutResourceKindFlag::make(LayoutResourceKind::VaryingOutput) |
+    LayoutResourceKindFlag::make(LayoutResourceKind::CallablePayload) |
+    LayoutResourceKindFlag::make(LayoutResourceKind::HitAttributes) |
+    LayoutResourceKindFlag::make(LayoutResourceKind::RayPayload);
+
+bool isVaryingParameter(IRTypeLayout* typeLayout)
 {
-    switch (kind)
+    if (!typeLayout)
+        return false;
+
+    // If we're looking at a struct type layout, we need to check each of the fields.
+    // If any of the fields is not a varying resource kind, then we consider the
+    // whole struct to be uniform (and thus not varying).
+    if (auto structTypeLayout = as<IRStructTypeLayout>(typeLayout))
+    {
+        for (auto fieldAttr : structTypeLayout->getFieldLayoutAttrs())
+        {
+            if (!isVaryingParameter(fieldAttr->getLayout()))
+                return false;
+        }
+        // If we made it here, all struct fields were varying.
+        return true;
+    }
+
+    // If we're looking at a parameter group type layout, we should return false
+    // as these should always be associtated with LayoutResourceKind::Uniform.
+    else if (as<IRParameterGroupTypeLayout>(typeLayout))
     {
-    default:
         return false;
+    }
 
-        // Note: The set of cases that are considered
-        // varying here would need to be extended if we
-        // add more fine-grained resource kinds (e.g.,
-        // if we ever add an explicit resource kind
-        // for geometry shader output streams).
-        //
-        // Ordinary varying input/output:
-    case LayoutResourceKind::VaryingInput:
-    case LayoutResourceKind::VaryingOutput:
-        //
-        // Ray-tracing shader input/output:
-    case LayoutResourceKind::CallablePayload:
-    case LayoutResourceKind::HitAttributes:
-    case LayoutResourceKind::RayPayload:
+    // If we're looking at an array type layout, we need to check the element's
+    // type layout.
+    else if (auto arrayTypeLayout = as<IRArrayTypeLayout>(typeLayout))
+    {
+        return isVaryingParameter(arrayTypeLayout->getElementTypeLayout());
+    }
+
+    // If we didn't match a type layout above, try finding a kind from sizeAttrs.
+    else
+    {
+        // If all the element type's sizeAttrs have a varying kind matching our
+        // flags above, return true.
+        for (auto sizeAttr : typeLayout->getSizeAttrs())
+        {
+            if ((flags & LayoutResourceKindFlag::make(sizeAttr->getResourceKind())) == 0)
+                return false;
+        }
         return true;
     }
 }
 
-bool isVaryingParameter(IRTypeLayout* typeLayout)
+bool isVaryingParameter(IRVarLayout* varLayout)
 {
     // If *any* of the resources consumed by the parameter type
     // is *not* a varying resource kind, then we consider the
     // whole parameter to be uniform (and thus not varying).
     //
-    // Note that this means that an empty type will always
+    // Note that this means that some empty types will always
     // be considered varying, even if it had been explicitly
     // marked `uniform`.
     //
@@ -149,19 +182,25 @@ bool isVaryingParameter(IRTypeLayout* typeLayout)
     // reosurce kind, so they show up as empty. Simply
     // adding `LayoutResourceKind`s for system-value inputs
     // and outputs would allow for simpler logic here.
-    //
-    for (auto sizeAttr : typeLayout->getSizeAttrs())
-    {
-        if (!isVaryingResourceKind(sizeAttr->getResourceKind()))
-            return false;
-    }
-    return true;
-}
 
-bool isVaryingParameter(IRVarLayout* varLayout)
-{
     if (!varLayout)
         return false;
+
+    // System-value parameters currently do not have kinds setup.
+    // As a WAR, if we see a system-value attr just assume varying
+    // for now.
+    if (varLayout->findAttr<IRSystemValueSemanticAttr>())
+    {
+        return true;
+    }
+
+    if (varLayout->usesResourceFromKinds(flags))
+    {
+        return true;
+    }
+
+    // If the base cases above failed, we need to check if we are dealing with
+    // an IRVarLayout that could have "nested" IRVarLayouts, such as an IRStructTypeLayout.
     return isVaryingParameter(varLayout->getTypeLayout());
 }
 
diff --git a/source/slang/slang-ir-explicit-global-context.cpp b/source/slang/slang-ir-explicit-global-context.cpp
index f235ba3e4..9fae4ce15 100644
--- a/source/slang/slang-ir-explicit-global-context.cpp
+++ b/source/slang/slang-ir-explicit-global-context.cpp
@@ -661,6 +661,12 @@ struct IntroduceExplicitGlobalContextPass
                 ptr = builder.emitLoad(ptr);
             use->set(ptr);
         }
+
+        // We've replaced all uses of the global var so we should remove it.
+        // We leave decorations on the global var above, so if we do not remove it
+        // here, the global var will not be processed for elimination in
+        // eliminateDeadCode.
+        globalVar->removeAndDeallocate();
     }
 
     IRInst* findOrCreateContextPtrForInst(IRInst* inst)
diff --git a/source/slang/slang-ir-glsl-legalize.cpp b/source/slang/slang-ir-glsl-legalize.cpp
index c9791880f..fa7e69fa4 100644
--- a/source/slang/slang-ir-glsl-legalize.cpp
+++ b/source/slang/slang-ir-glsl-legalize.cpp
@@ -3575,28 +3575,40 @@ void legalizeEntryPointParameterForGLSL(
 
                 auto key = dec->getOperand(1);
                 IRInst* realGlobalVar = nullptr;
-                if (globalValue.flavor != ScalarizedVal::Flavor::tuple)
-                    continue;
-                if (auto tupleVal = as<ScalarizedTupleValImpl>(globalValue.impl))
+
+                // When we relate a "global variable" to a "global parameter" using
+                // kIROp_GlobalVariableShadowingGlobalParameterDecoration, the globalValue flavor
+                // is dependent on the global parameter's type. Struct types for example will relate
+                // to the tuple flavor, while vector types will be related to the address flavor.
+                if (globalValue.flavor == ScalarizedVal::Flavor::tuple)
                 {
-                    for (auto elem : tupleVal->elements)
+                    if (auto tupleVal = as<ScalarizedTupleValImpl>(globalValue.impl))
                     {
-                        if (elem.key == key)
+                        for (auto elem : tupleVal->elements)
                         {
-                            realGlobalVar = elem.val.irValue;
-                            if (!realGlobalVar &&
-                                ScalarizedVal::Flavor::typeAdapter == elem.val.flavor)
+                            if (elem.key == key)
                             {
-                                if (auto typeAdapterVal =
-                                        as<ScalarizedTypeAdapterValImpl>(elem.val.impl))
+                                realGlobalVar = elem.val.irValue;
+                                if (!realGlobalVar &&
+                                    ScalarizedVal::Flavor::typeAdapter == elem.val.flavor)
                                 {
-                                    realGlobalVar = typeAdapterVal->val.irValue;
+                                    if (auto typeAdapterVal =
+                                            as<ScalarizedTypeAdapterValImpl>(elem.val.impl))
+                                    {
+                                        realGlobalVar = typeAdapterVal->val.irValue;
+                                    }
                                 }
+                                break;
                             }
-                            break;
                         }
                     }
                 }
+                else if (globalValue.flavor == ScalarizedVal::Flavor::address)
+                {
+                    realGlobalVar = globalValue.irValue;
+                }
+                else
+                    continue;
                 SLANG_ASSERT(realGlobalVar);
 
                 // Remove all stores into the global var introduced during
diff --git a/source/slang/slang-ir-legalize-varying-params.cpp b/source/slang/slang-ir-legalize-varying-params.cpp
index f8a9839d2..e5b3de90f 100644
--- a/source/slang/slang-ir-legalize-varying-params.cpp
+++ b/source/slang/slang-ir-legalize-varying-params.cpp
@@ -1891,6 +1891,10 @@ private:
                 structTypeLayout = as<IRStructTypeLayout>(varLayout->getTypeLayout());
             Index fieldIndex = 0;
             List<IRInst*> fieldParams;
+            // TODO: We currently lose some decorations from the struct that should possibly be
+            // transfered
+            //       to the new params here, like
+            //       kIROp_GlobalVariableShadowingGlobalParameterDecoration.
             for (auto field : structType->getFields())
             {
                 auto fieldParam = builder.emitParam(field->getFieldType());
diff --git a/source/slang/slang-ir-translate-global-varying-var.cpp b/source/slang/slang-ir-translate-global-varying-var.cpp
index 80f5c42c3..57b277d25 100644
--- a/source/slang/slang-ir-translate-global-varying-var.cpp
+++ b/source/slang/slang-ir-translate-global-varying-var.cpp
@@ -75,8 +75,80 @@ struct GlobalVarTranslationContext
             builder.setInsertBefore(entryPointFunc);
             auto inputStructType = builder.createStructType();
             IRStructTypeLayout::Builder inputStructTypeLayoutBuilder(&builder);
+
+            // Go through the existing parameter layouts and add them to the
+            // replacement struct type first.
+            //
+            // We'll also need to fix up the offsets of the global varing inputs
+            // being added if we have any varing input entry point params. Set
+            // these to some defaults for now.
+            UInt nextOffset = 0;
+            IRVarLayout* lastFieldVarLayout = nullptr;
+
+            // If we have an existing entry point struct layout, we need to go through it and
+            // add any of the struct field layout attributes as fields in the new struct type
+            // layout that we are building to combine global varing ins and entry point ins.
+            IRLayoutDecoration* entryPointLayoutDecor =
+                entryPointFunc->findDecoration<IRLayoutDecoration>();
+            IREntryPointLayout* entryPointLayout = nullptr;
+            IRStructTypeLayout* entryPointParamsStructLayout = nullptr;
+
+            if (entryPointLayoutDecor)
+            {
+                entryPointLayout = as<IREntryPointLayout>(entryPointLayoutDecor->getLayout());
+                if (entryPointLayout)
+                {
+                    // The parameter layout for an entry point will either be a structure
+                    // type layout, or a parameter group type layout.
+                    entryPointParamsStructLayout = getScopeStructLayout(entryPointLayout);
+                    if (entryPointParamsStructLayout)
+                    {
+                        for (auto attr : entryPointParamsStructLayout->getFieldLayoutAttrs())
+                        {
+                            // Add existing parameter layouts to the replacement struct type layout.
+                            IRVarLayout* fieldLayout = attr->getLayout();
+                            IRInst* key = attr->getFieldKey();
+                            inputStructTypeLayoutBuilder.addField(key, fieldLayout);
+
+                            // Save the last VaryingInput type. This will be used to help
+                            // recalculate offsets for the global varying inputs.
+                            if (fieldLayout->usesResourceKind(LayoutResourceKind::VaryingInput))
+                                lastFieldVarLayout = fieldLayout;
+                        }
+
+                        // Calculate the nextOffset if necessary. If entry params had no varying
+                        // inputs nextOffset will just be 0.
+                        if (lastFieldVarLayout)
+                        {
+                            // Find the size and offset of the last varying "in" kind
+                            // param in the entry point param struct. Adding these will give us
+                            // the starting offset that should be used for the first global.
+                            if (auto sizeAttr = lastFieldVarLayout->getTypeLayout()->findSizeAttr(
+                                    LayoutResourceKind::VaryingInput))
+                            {
+                                size_t finiteSize = sizeAttr->getFiniteSize();
+                                if (auto offsetAttr = lastFieldVarLayout->findOffsetAttr(
+                                        LayoutResourceKind::VaryingInput))
+                                {
+                                    UInt lastFieldOffset = offsetAttr->getOffset();
+                                    nextOffset = finiteSize + lastFieldOffset;
+                                }
+                            }
+                        }
+                    }
+                }
+            }
+
+            // Add the global vars to the replacement struct and add new params to
+            // the entry point func based on the global input vars.
+            List<IRInst*> newParams;
             UInt inputVarIndex = 0;
             List<IRStructKey*> inputKeys;
+
+            // Setup the location where we will insert the new params
+            auto firstBlock = entryPointFunc->getFirstBlock();
+            builder.setInsertInto(firstBlock);
+
             for (auto input : inputVars)
             {
                 auto inputType = cast<IRPtrTypeBase>(input->getDataType())->getValueType();
@@ -118,14 +190,19 @@ struct GlobalVarTranslationContext
                             LayoutResourceKind::VaryingInput,
                             LayoutSize(1));
                     }
+
+                    // Start off the offset as nextOffset. If the global "in"s have existing
+                    // offsetAttr's, we will add the offsets from those as well.
+                    auto resInfo =
+                        varLayoutBuilder.findOrAddResourceInfo(LayoutResourceKind::VaryingInput);
+                    resInfo->offset = nextOffset;
+
                     if (auto layoutDecor = findVarLayout(input))
                     {
                         if (auto offsetAttr =
                                 layoutDecor->findOffsetAttr(LayoutResourceKind::VaryingInput))
                         {
-                            varLayoutBuilder
-                                .findOrAddResourceInfo(LayoutResourceKind::VaryingInput)
-                                ->offset = (UInt)offsetAttr->getOffset();
+                            resInfo->offset += (UInt)offsetAttr->getOffset();
                         }
                     }
                     if (entryPointDecor->getProfile().getStage() == Stage::Fragment)
@@ -138,41 +215,60 @@ struct GlobalVarTranslationContext
                     }
                     inputVarIndex++;
                 }
-                inputStructTypeLayoutBuilder.addField(key, varLayoutBuilder.build());
+                auto varLayout = varLayoutBuilder.build();
+                inputStructTypeLayoutBuilder.addField(key, varLayout);
                 input->transferDecorationsTo(key);
+
+                // Emit a new param here to represent the global input var.
+                auto inputParam = builder.emitParam(
+                    builder.getPtrType(kIROp_ConstRefType, inputType, AddressSpace::Input));
+
+                // Copy the global input vars original decorations onto the new param.
+                // We need to do this to ensure that we can do things like get system
+                // value info in later passes like when we legalize entry points for WGSL.
+                IRCloneEnv cloneEnv;
+                cloneInstDecorationsAndChildren(&cloneEnv, builder.getModule(), key, inputParam);
+
+                // Add the layout to the new param
+                builder.addLayoutDecoration(inputParam, varLayout);
+
+                // Add the param to our list of new params. This will allow us
+                // to connect the old "global variable" to a "global parameter"
+                // later on.
+                newParams.add(inputParam);
             }
+
+            // Build the new layout that will replace the old entryPointLayout.
             auto paramTypeLayout = inputStructTypeLayoutBuilder.build();
             IRVarLayout::Builder paramVarLayoutBuilder(&builder, paramTypeLayout);
             paramLayout = paramVarLayoutBuilder.build();
 
-            // Add an entry point parameter for all the inputs.
-            auto firstBlock = entryPointFunc->getFirstBlock();
-            builder.setInsertInto(firstBlock);
-            auto inputParam = builder.emitParam(
-                builder.getPtrType(kIROp_ConstRefType, inputStructType, AddressSpace::Input));
-            builder.addLayoutDecoration(inputParam, paramLayout);
-
             // Initialize all global variables in the order of struct member declaration.
             for (Index i = inputVars.getCount() - 1; i >= 0; i--)
             {
                 auto input = inputVars[i];
                 setInsertBeforeOrdinaryInst(&builder, firstBlock->getFirstOrdinaryInst());
-                auto inputType = cast<IRPtrTypeBase>(input->getDataType())->getValueType();
+                // TODO: Does the below TODO apply if we no longer use emitFieldExtract?
+#if 0
                 // TODO: This could be more efficient as a Load(FieldAddress(inputParam, i))
                 // operation instead of a FieldExtract(Load(inputParam)).
                 builder.emitStore(
                     input,
                     builder
                         .emitFieldExtract(inputType, builder.emitLoad(inputParam), inputKeys[i]));
-                // Relate "global variable" to a "global parameter" for use later in compilation
-                // to resolve a "global variable" shadowing a "global parameter" relationship.
+#endif
+                builder.emitStore(input, builder.emitLoad(newParams[i]));
+
+                // Relate the old "global variable" to a "global parameter" for use later in
+                // compilation to resolve a "global variable" shadowing a "global parameter"
+                // relationship.
                 builder.addGlobalVariableShadowingGlobalParameterDecoration(
-                    inputParam,
+                    newParams[i],
                     input,
                     inputKeys[i]);
             }
 
-            // For each entry point, introduce a new parameter to represent each input parameter,
+            // For each entry point, introduce a new parameter to represent each output parameter,
             // and return all outputs via a struct value.
             if (hasOutput)
             {
@@ -256,18 +352,43 @@ struct GlobalVarTranslationContext
                     }
                 }
             }
-            if (auto entryPointLayoutDecor = entryPointFunc->findDecoration<IRLayoutDecoration>())
+            if (entryPointLayout)
             {
-                if (auto entryPointLayout =
-                        as<IREntryPointLayout>(entryPointLayoutDecor->getLayout()))
+                if (paramLayout)
                 {
-                    if (paramLayout)
-                        builder.replaceOperand(entryPointLayout->getOperands(), paramLayout);
-                    if (resultVarLayout)
+                    // We try to respect the original entry point type layout when replacing it
+                    // below. The IRParameterGroupTypeLayout type layout is used when there are
+                    // things like uniform entry point params, otherwise the IRStructTypeLayout type
+                    // layout will normally represent the entry point type layout.
+                    if (auto paramGroupTypeLayout =
+                            as<IRParameterGroupTypeLayout>(entryPointParamsStructLayout))
+                    {
+                        // Build a new IRParameterGroupTypeLayout to replace the old one
+                        // representing the entryPointLayout. The ContainerVarLayout shouldn't have
+                        // changed, so resue that. Replace the rest with the new paramTypeLayout and
+                        // paramLayout that we calculated above
+                        IRParameterGroupTypeLayout::Builder paramGroupTypeLayoutBuilder(&builder);
+                        paramGroupTypeLayoutBuilder.setContainerVarLayout(
+                            paramGroupTypeLayout->getContainerVarLayout());
+                        paramGroupTypeLayoutBuilder.setElementVarLayout(paramLayout);
+                        paramGroupTypeLayoutBuilder.setOffsetElementTypeLayout(paramTypeLayout);
+                        builder.replaceOperand(
+                            entryPointLayout->getParamsLayout()->getOperands(),
+                            paramGroupTypeLayoutBuilder.build());
+                    }
+                    else if (as<IRStructTypeLayout>(entryPointParamsStructLayout))
+                    {
                         builder.replaceOperand(
-                            entryPointLayout->getOperands() + 1,
-                            resultVarLayout);
+                            entryPointLayout->getParamsLayout()->getOperands(),
+                            paramTypeLayout);
+                    }
+                    else
+                    {
+                        SLANG_UNEXPECTED("uhandled global-scope binding layout");
+                    }
                 }
+                if (resultVarLayout)
+                    builder.replaceOperand(entryPointLayout->getOperands() + 1, resultVarLayout);
             }
             // Update func type for the entry point.
             List<IRType*> paramTypes;
diff --git a/tests/hlsl-intrinsic/wave-get-lane-index.slang b/tests/hlsl-intrinsic/wave-get-lane-index.slang
index e1a9262a9..b059ff42d 100644
--- a/tests/hlsl-intrinsic/wave-get-lane-index.slang
+++ b/tests/hlsl-intrinsic/wave-get-lane-index.slang
@@ -11,9 +11,9 @@
 RWStructuredBuffer<int> outputBuffer;
 
 [numthreads(4, 1, 1)]
-void computeMain(int3 dispatchThreadID : SV_DispatchThreadID)
+void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID)
 {
-    int idx = dispatchThreadID.x;
+    uint idx = dispatchThreadID.x;
     uint laneId = WaveGetLaneIndex();
     // The laneCount will be dependent on target hardware. It seems a count of 1 is valid in spec. 
     // For now we'll just check it's not 0.