From 45b76418f9da2248b069f2058c6a1d52b05a8c74 Mon Sep 17 00:00:00 2001
From: ArielG-NV <159081215+ArielG-NV@users.noreply.github.com>
Date: Wed, 14 Aug 2024 13:05:57 -0400
Subject: 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>
---
 .../zero-initialize/rayquery.slang                 | 47 ++++++++++++++++++++++
 .../zero-initialize/shared-memory.slang            | 18 +++++++++
 2 files changed, 65 insertions(+)
 create mode 100644 tests/language-feature/zero-initialize/rayquery.slang
 create mode 100644 tests/language-feature/zero-initialize/shared-memory.slang

(limited to 'tests/language-feature/zero-initialize')

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;
+}
-- 
cgit v1.2.3