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 --- tests/compute/groupshared-init.slang | 28 ------------- tests/diagnostics/uninitialized-globals.slang | 7 ---- .../atomic/atomic-intrinsics-64bit.slang | 14 +++++-- .../hlsl-intrinsic/atomic/atomic-intrinsics.slang | 7 +++- .../shared-memory-initializer-error.slang | 14 +++++++ .../zero-initialize/rayquery.slang | 47 ++++++++++++++++++++++ .../zero-initialize/shared-memory.slang | 18 +++++++++ .../groupshared-threadlocal-same-parameter.slang | 3 +- 8 files changed, 96 insertions(+), 42 deletions(-) delete mode 100644 tests/compute/groupshared-init.slang create mode 100644 tests/language-feature/shared-memory-initializer-error.slang create mode 100644 tests/language-feature/zero-initialize/rayquery.slang create mode 100644 tests/language-feature/zero-initialize/shared-memory.slang (limited to 'tests') 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 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 outputBuffer; @@ -28,6 +28,12 @@ RWStructuredBuffer 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 uintBuffer; //TEST_INPUT:ubuffer(data=[0 1 2 3], stride=4):name=intBuffer RWStructuredBuffer 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 outputBuffer; @@ -20,6 +20,9 @@ RWStructuredBuffer 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 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 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 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 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 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) -- cgit v1.2.3