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>

3 files changed, 79 insertions, 0 deletions

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;
+}