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 --- source/slang/slang-check-decl.cpp | 52 ++++++++++++++++++++++ source/slang/slang-diagnostic-defs.h | 2 + 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 +- 10 files changed, 150 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 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()) + { + 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(varDecl); + if (auto declRefType = as(varDecl->getType())) + { + if (auto genericAppRefDecl = as(declRefType->getDeclRefBase())) + { + auto baseGenericRefType = genericAppRefDecl->getBase()->getDecl(); + if (auto baseTypeStruct = as(baseGenericRefType)) + { + structDecl = baseTypeStruct; + } + else if (auto genericDecl = as(baseGenericRefType)) + { + if(auto innerTypeStruct = as(genericDecl->inner)) + structDecl = innerTypeStruct; + } + } + } + if (structDecl) + { + // find if a type is non-copyable + if (structDecl->findModifier()) + 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(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 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