Do not zero-initialize groupshared and rayquery variables (#4838)

* Do not zero-initialize groupshared and rayquery variables

Fixes: #4824

`-zero-initialize` option will explicitly not:
1. Set any groupshared values to defaults
2. Set any rayQuery object to a default state (currently invalid code generation)

* grammer

* disallow groupshared initializers

disallow groupshared initializers & adjust tests accordingly

* remove disallowed groupshared-init expression

* do not default init if non-copyable

---------

Co-authored-by: Yong He <yonghe@outlook.com>

10 files changed, 150 insertions, 42 deletions

diff --git a/source/slang/slang-check-decl.cpp b/source/slang/slang-check-decl.cpp
index 7211565dd..2e5e13360 100644
--- a/source/slang/slang-check-decl.cpp
+++ b/source/slang/slang-check-decl.cpp
@@ -1922,10 +1922,59 @@ namespace Slang
         checkVisibility(classDecl);
     }
 
+    bool DiagnoseIsAllowedInitExpr(VarDeclBase* varDecl, DiagnosticSink* sink)
+    {
+        // find groupshared modifier
+        if (varDecl->findModifier<HLSLGroupSharedModifier>())
+        {
+            if (sink && varDecl->initExpr)
+                sink->diagnose(varDecl, Diagnostics::cannotHaveInitializer, varDecl, "groupshared");
+            return false;
+        }
+
+        return true;
+    }
+
+    bool isDefaultInitializable(VarDeclBase* varDecl)
+    {
+        if (!DiagnoseIsAllowedInitExpr(varDecl, nullptr))
+            return false;
+
+        // Find struct and modifiers associated with varDecl
+        StructDecl* structDecl = as<StructDecl>(varDecl);
+        if (auto declRefType = as<DeclRefType>(varDecl->getType()))
+        {
+            if (auto genericAppRefDecl = as<GenericAppDeclRef>(declRefType->getDeclRefBase()))
+            {
+                auto baseGenericRefType = genericAppRefDecl->getBase()->getDecl();
+                if (auto baseTypeStruct = as<StructDecl>(baseGenericRefType))
+                {
+                    structDecl = baseTypeStruct;
+                }
+                else if (auto genericDecl = as<GenericDecl>(baseGenericRefType))
+                {
+                    if(auto innerTypeStruct = as<StructDecl>(genericDecl->inner))
+                        structDecl = innerTypeStruct;
+                }
+            }
+        }
+        if (structDecl)
+        {
+            // find if a type is non-copyable
+            if (structDecl->findModifier<NonCopyableTypeAttribute>())
+                return false;
+        }
+
+        return true;
+    }
+
     static Expr* constructDefaultInitExprForVar(SemanticsVisitor* visitor, VarDeclBase* varDecl)
     {
         if (!varDecl->type || !varDecl->type.type)
             return nullptr;
+        
+        if (!isDefaultInitializable(varDecl))
+            return nullptr;
 
         ConstructorDecl* defaultCtor = nullptr;
         auto declRefType = as<DeclRefType>(varDecl->type.type);
@@ -1951,8 +2000,11 @@ namespace Slang
             return defaultCall;
         }
     }
+
     void SemanticsDeclBodyVisitor::checkVarDeclCommon(VarDeclBase* varDecl)
     {
+        DiagnoseIsAllowedInitExpr(varDecl, getSink());
+
         // if zero initialize is true, set everything to a default
         if (getOptionSet().hasOption(CompilerOptionName::ZeroInitialize) 
             && !varDecl->initExpr
diff --git a/source/slang/slang-diagnostic-defs.h b/source/slang/slang-diagnostic-defs.h
index 44e8aa13d..03e8efbc0 100644
--- a/source/slang/slang-diagnostic-defs.h
+++ b/source/slang/slang-diagnostic-defs.h
@@ -520,6 +520,8 @@ DIAGNOSTIC(30504, Error, cannotUseInitializerListForType, "cannot use initialize
 // 3062x: variables
 DIAGNOSTIC(30620, Error, varWithoutTypeMustHaveInitializer, "a variable declaration without an initial-value expression must be given an explicit type")
 DIAGNOSTIC(30622, Error, ambiguousDefaultInitializerForType, "more than one default initializer was found for type '$0'")
+DIAGNOSTIC(30623, Error, cannotHaveInitializer, "'$0' cannot have an initializer because it is $1")
+
 
 // 307xx: parameters
 DIAGNOSTIC(30700, Error, outputParameterCannotHaveDefaultValue, "an 'out' or 'inout' parameter cannot have a default-value expression")
diff --git a/tests/compute/groupshared-init.slang b/tests/compute/groupshared-init.slang
deleted file mode 100644
index 5a9758826..000000000
--- a/tests/compute/groupshared-init.slang
+++ /dev/null
@@ -1,28 +0,0 @@
-//TEST(compute):COMPARE_COMPUTE(filecheck-buffer=CHECK):-cpu -output-using-type
-//TEST(compute):COMPARE_COMPUTE(filecheck-buffer=CHECK):-dx12 -output-using-type -use-dxil
-//TEST(compute):COMPARE_COMPUTE(filecheck-buffer=CHECK):-vk -output-using-type
-//TEST(compute):COMPARE_COMPUTE(filecheck-buffer=CHECK):-vk -output-using-type -emit-spirv-directly
-
-// CHECK: type: uint32_t
-// CHECK-NEXT: 1231
-// CHECK-NEXT: 1232
-// CHECK-NEXT: 1233
-// CHECK-NEXT: 1234
-
-// This is a basic test for Slang compute shader.
-
-//TEST_INPUT:ubuffer(data=[0 0 0 0], stride=4):out,name=outputBuffer
-RWStructuredBuffer<uint> outputBuffer;
-
-groupshared uint myGroupSharedValue = foo();
-
-uint foo()
-{
-  return 1231;
-}
-
-[numthreads(4, 1, 1)]
-void computeMain(uint i : SV_GroupIndex)
-{
-    outputBuffer[i] = i + myGroupSharedValue;
-}
diff --git a/tests/diagnostics/uninitialized-globals.slang b/tests/diagnostics/uninitialized-globals.slang
index 635db1e29..730b0343f 100644
--- a/tests/diagnostics/uninitialized-globals.slang
+++ b/tests/diagnostics/uninitialized-globals.slang
@@ -1,15 +1,8 @@
 //TEST:SIMPLE(filecheck=CHK): -allow-glsl -target spirv
 
 // Using groupshared variables
-groupshared float4 gsConstexpr = float4(1.0f);
 groupshared float4 gsUninitialized;
 
-// OK
-float use_constexpr_initialized_gs()
-{
-    return gsConstexpr.x;
-}
-
 float use_undefined_gs()
 {
     //CHK-DAG: ([[# @LINE + 1]]): warning 41017: use of uninitialized global variable 'gsUninitialized'
diff --git a/tests/hlsl-intrinsic/atomic/atomic-intrinsics-64bit.slang b/tests/hlsl-intrinsic/atomic/atomic-intrinsics-64bit.slang
index aa05f9750..aee2b40e4 100644
--- a/tests/hlsl-intrinsic/atomic/atomic-intrinsics-64bit.slang
+++ b/tests/hlsl-intrinsic/atomic/atomic-intrinsics-64bit.slang
@@ -17,10 +17,10 @@ RWByteAddressBuffer uBuf;
 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0]):name=iBuf
 RWByteAddressBuffer iBuf;
 
-groupshared float    f32Shared[4] = { 0.f, 0.f, 0.f, 0.f };
-groupshared uint64_t u64Shared[4] = { 0, 0, 0, 0 };
-groupshared int64_t  i64Shared[4] = { 0, 0, 0, 0 };
-groupshared uint64_t indexAlloc = 0;
+groupshared float f32Shared[4];
+groupshared uint64_t u64Shared[4];
+groupshared int64_t i64Shared[4];
+groupshared uint64_t indexAlloc;
 
 //TEST_INPUT: ubuffer(data=[0 0 0 0], stride=4):out,name outputBuffer
 RWStructuredBuffer<int> outputBuffer;
@@ -28,6 +28,12 @@ RWStructuredBuffer<int> outputBuffer;
 [numthreads(4, 1, 1)]
 void computeMain(uint groupIndex : SV_GroupIndex, int3 dispatchThreadID: SV_DispatchThreadID)
 {
+    f32Shared = { 0.f, 0.f, 0.f, 0.f };
+    u64Shared = { 0, 0, 0, 0 };
+    i64Shared = { 0, 0, 0, 0 };
+    indexAlloc = 0;
+    GroupMemoryBarrierWithGroupSync();
+
     int idx = dispatchThreadID.x;
     bool result = true;
 
diff --git a/tests/hlsl-intrinsic/atomic/atomic-intrinsics.slang b/tests/hlsl-intrinsic/atomic/atomic-intrinsics.slang
index c118b9d23..390b66be2 100644
--- a/tests/hlsl-intrinsic/atomic/atomic-intrinsics.slang
+++ b/tests/hlsl-intrinsic/atomic/atomic-intrinsics.slang
@@ -11,8 +11,8 @@ RWStructuredBuffer<uint> uintBuffer;
 //TEST_INPUT:ubuffer(data=[0 1 2 3], stride=4):name=intBuffer
 RWStructuredBuffer<int> intBuffer;
 
-groupshared uint shareMemUI[4] = { 0, 0, 0, 0 };
-groupshared int shareMemI[4] = { 0, 0, 0, 0 };
+groupshared uint shareMemUI[4];
+groupshared int shareMemI[4];
 
 //TEST_INPUT: ubuffer(data=[0 0 0 0], stride=4):out,name outputBuffer
 RWStructuredBuffer<float> outputBuffer;
@@ -20,6 +20,9 @@ RWStructuredBuffer<float> outputBuffer;
 [numthreads(4, 1, 1)]
 void computeMain(uint groupIndex : SV_GroupIndex, int3 dispatchThreadID: SV_DispatchThreadID)
 {
+    shareMemUI = { 0, 0, 0, 0 };
+    shareMemI = { 0, 0, 0, 0 };
+    AllMemoryBarrierWithGroupSync();
     int idx = dispatchThreadID.x;
     float val = 0.0f;
 
diff --git a/tests/language-feature/shared-memory-initializer-error.slang b/tests/language-feature/shared-memory-initializer-error.slang
new file mode 100644
index 000000000..ceb361176
--- /dev/null
+++ b/tests/language-feature/shared-memory-initializer-error.slang
@@ -0,0 +1,14 @@
+//TEST:SIMPLE(filecheck=HLSL): -target hlsl -stage compute -entry computeMain
+//TEST:SIMPLE(filecheck=GLSL): -target glsl -stage compute -entry computeMain
+
+RWStructuredBuffer<uint> outputBuffer;
+
+// GLSL: error 30623
+// HLSL: error 30623
+groupshared uint globalMem = 1;
+
+[numthreads(1, 1, 1)]
+void computeMain(int3 dispatchThreadID: SV_DispatchThreadID)
+{
+    outputBuffer[0] = globalMem ;
+}
diff --git a/tests/language-feature/zero-initialize/rayquery.slang b/tests/language-feature/zero-initialize/rayquery.slang
new file mode 100644
index 000000000..6c48d3c65
--- /dev/null
+++ b/tests/language-feature/zero-initialize/rayquery.slang
@@ -0,0 +1,47 @@
+//TEST:SIMPLE(filecheck=HLSL): -target hlsl -stage compute -entry computeMain -zero-initialize
+//TEST:SIMPLE(filecheck=GLSL): -target glsl -stage compute -entry computeMain -zero-initialize
+
+// HLSL-NOT: RayQuery{{.*}} {{.*}} =
+// GLSL-NOT: rayQueryEXT {{.*}} =
+
+uniform RaytracingAccelerationStructure scene;
+RWStructuredBuffer<float> outputBuffer;
+
+bool traceRayClosestHit(
+    float3 rayOrigin,
+    float3 rayDir,
+    out float t)
+{
+    RayDesc ray;
+    ray.Origin = rayOrigin;
+    ray.TMin = 0.01f;
+    ray.Direction = rayDir;
+    ray.TMax = 1e4f;
+    RayQuery<RAY_FLAG_NONE> q;
+    let rayFlags = RAY_FLAG_NONE;
+
+    q.TraceRayInline(
+        scene,
+        rayFlags,
+        0xff,
+        ray);
+
+    q.Proceed();
+    if(q.CommittedStatus() == COMMITTED_TRIANGLE_HIT)
+    {
+        t = q.CommittedRayT();
+        return true;
+    }
+    unused(t);
+    return false;
+}
+
+[shader("compute")]
+[numthreads(1,1,1)]
+void computeMain(
+    uint3 threadIdx : SV_DispatchThreadID)
+{
+    float t = 0.0;
+    traceRayClosestHit(float3(0.1, 0.1, 0.0), float3(0,0,1), t);
+    outputBuffer[threadIdx.x] = t;
+}
\ No newline at end of file
diff --git a/tests/language-feature/zero-initialize/shared-memory.slang b/tests/language-feature/zero-initialize/shared-memory.slang
new file mode 100644
index 000000000..39243f796
--- /dev/null
+++ b/tests/language-feature/zero-initialize/shared-memory.slang
@@ -0,0 +1,18 @@
+//TEST:SIMPLE(filecheck=HLSL): -target hlsl -stage compute -entry computeMain -zero-initialize
+//TEST:SIMPLE(filecheck=GLSL): -target glsl -stage compute -entry computeMain -zero-initialize
+
+RWStructuredBuffer<uint> outputBuffer;
+
+// GLSL-NOT: error 30623
+// HLSL-NOT: error 30623
+
+// GLSL-NOT: globalMem{{.*}} =
+// HLSL-NOT: globalMem{{.*}} =
+
+groupshared uint globalMem;
+
+[numthreads(1, 1, 1)]
+void computeMain(int3 dispatchThreadID: SV_DispatchThreadID)
+{
+    outputBuffer[0] = globalMem;
+}
diff --git a/tests/metal/groupshared-threadlocal-same-parameter.slang b/tests/metal/groupshared-threadlocal-same-parameter.slang
index ae4cac8df..1249d4703 100644
--- a/tests/metal/groupshared-threadlocal-same-parameter.slang
+++ b/tests/metal/groupshared-threadlocal-same-parameter.slang
@@ -6,7 +6,7 @@
 //TEST_INPUT:ubuffer(data=[0 0], stride=4):out,name outputBuffer
 RWStructuredBuffer<uint> outputBuffer;
 
-static groupshared uint g_values[2] = { 1, 0 };
+static groupshared uint g_values[2];
 static uint g_altValues[2] = { 2, 3 };
 
 static groupshared uint g_valuesReturned[2];
@@ -27,6 +27,7 @@ uint[2] maybeGroupSharedReturn(uint id)
 [numthreads(2, 1, 1)]
 void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID)
 {
+    g_values = { 1, 0 };
     AllMemoryBarrierWithGroupSync();
     uint tid = dispatchThreadID.x;
     if (tid == 0)