From eeda18d55ea2f998b5eaff583b4e1922891f8013 Mon Sep 17 00:00:00 2001 From: Ben Clayton Date: Wed, 7 Sep 2022 13:05:24 +0000 Subject: [PATCH] tint::transform::SingleEntryPoint: Preserve global 'const's These can be used in all sorts of places which are not tracked by sem::Function::TransitivelyReferencedGlobals(). As they're not emitted as variables by any backend, just preserve them. Fixed: tint:1598 Change-Id: I2696486cb2ffe8408bd5dd3090d7d600ca1d170f Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/101481 Reviewed-by: Dan Sinclair Commit-Queue: Dan Sinclair Auto-Submit: Ben Clayton --- src/tint/transform/single_entry_point.cc | 12 ++++++-- src/tint/transform/single_entry_point_test.cc | 29 +++++++++++++++++++ test/tint/bug/tint/1598.wgsl | 8 +++++ .../tint/bug/tint/1598.wgsl.expected.dxc.hlsl | 4 +++ .../tint/bug/tint/1598.wgsl.expected.fxc.hlsl | 4 +++ test/tint/bug/tint/1598.wgsl.expected.glsl | 11 +++++++ test/tint/bug/tint/1598.wgsl.expected.msl | 21 ++++++++++++++ test/tint/bug/tint/1598.wgsl.expected.spvasm | 25 ++++++++++++++++ test/tint/bug/tint/1598.wgsl.expected.wgsl | 8 +++++ 9 files changed, 119 insertions(+), 3 deletions(-) create mode 100644 test/tint/bug/tint/1598.wgsl create mode 100644 test/tint/bug/tint/1598.wgsl.expected.dxc.hlsl create mode 100644 test/tint/bug/tint/1598.wgsl.expected.fxc.hlsl create mode 100644 test/tint/bug/tint/1598.wgsl.expected.glsl create mode 100644 test/tint/bug/tint/1598.wgsl.expected.msl create mode 100644 test/tint/bug/tint/1598.wgsl.expected.spvasm create mode 100644 test/tint/bug/tint/1598.wgsl.expected.wgsl diff --git a/src/tint/transform/single_entry_point.cc b/src/tint/transform/single_entry_point.cc index 133c836e07..8d26a7f955 100644 --- a/src/tint/transform/single_entry_point.cc +++ b/src/tint/transform/single_entry_point.cc @@ -86,11 +86,17 @@ void SingleEntryPoint::Run(CloneContext& ctx, const DataMap& inputs, DataMap&) c ctx.dst->AST().AddGlobalVariable(ctx.Clone(override)); } }, - [&](const ast::Variable* v) { // var, let - if (referenced_vars.count(v)) { - ctx.dst->AST().AddGlobalVariable(ctx.Clone(v)); + [&](const ast::Var* var) { + if (referenced_vars.count(var)) { + ctx.dst->AST().AddGlobalVariable(ctx.Clone(var)); } }, + [&](const ast::Const* c) { + // Always keep 'const' declarations, as these can be used by attributes and array + // sizes, which are not tracked as transitively used by functions. They also don't + // typically get emitted by the backend unless they're actually used. + ctx.dst->AST().AddGlobalVariable(ctx.Clone(c)); + }, [&](const ast::Function* func) { if (sem.Get(func)->HasAncestorEntryPoint(entry_point->symbol)) { ctx.dst->AST().AddFunction(ctx.Clone(func)); diff --git a/src/tint/transform/single_entry_point_test.cc b/src/tint/transform/single_entry_point_test.cc index 020bd65c44..7451090a19 100644 --- a/src/tint/transform/single_entry_point_test.cc +++ b/src/tint/transform/single_entry_point_test.cc @@ -217,8 +217,14 @@ fn comp_main2() { )"; auto* expect = R"( +const a : f32 = 1.0; + +const b : f32 = 1.0; + const c : f32 = 1.0; +const d : f32 = 1.0; + @compute @workgroup_size(1) fn comp_main1() { let local_c : f32 = c; @@ -536,5 +542,28 @@ fn comp_main1() { EXPECT_EQ(expect, str(got)); } +TEST_F(SingleEntryPointTest, GlobalConstUsedAsArraySize) { + // See crbug.com/tint/1598 + auto* src = R"( +const MY_SIZE = 5u; + +type Arr = array; + +@fragment +fn main() { +} +)"; + + auto* expect = src; + + SingleEntryPoint::Config cfg("main"); + + DataMap data; + data.Add(cfg); + auto got = Run(src, data); + + EXPECT_EQ(expect, str(got)); +} + } // namespace } // namespace tint::transform diff --git a/test/tint/bug/tint/1598.wgsl b/test/tint/bug/tint/1598.wgsl new file mode 100644 index 0000000000..8c1721123e --- /dev/null +++ b/test/tint/bug/tint/1598.wgsl @@ -0,0 +1,8 @@ +const MY_SIZE = 5u; + +type Arr = array; + +@fragment +fn main() { + var a : Arr; +} diff --git a/test/tint/bug/tint/1598.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/1598.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..610b22069d --- /dev/null +++ b/test/tint/bug/tint/1598.wgsl.expected.dxc.hlsl @@ -0,0 +1,4 @@ +void main() { + int a[5] = (int[5])0; + return; +} diff --git a/test/tint/bug/tint/1598.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/1598.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..610b22069d --- /dev/null +++ b/test/tint/bug/tint/1598.wgsl.expected.fxc.hlsl @@ -0,0 +1,4 @@ +void main() { + int a[5] = (int[5])0; + return; +} diff --git a/test/tint/bug/tint/1598.wgsl.expected.glsl b/test/tint/bug/tint/1598.wgsl.expected.glsl new file mode 100644 index 0000000000..c5d32b6d51 --- /dev/null +++ b/test/tint/bug/tint/1598.wgsl.expected.glsl @@ -0,0 +1,11 @@ +#version 310 es +precision mediump float; + +void tint_symbol() { + int a[5] = int[5](0, 0, 0, 0, 0); +} + +void main() { + tint_symbol(); + return; +} diff --git a/test/tint/bug/tint/1598.wgsl.expected.msl b/test/tint/bug/tint/1598.wgsl.expected.msl new file mode 100644 index 0000000000..de53a2be77 --- /dev/null +++ b/test/tint/bug/tint/1598.wgsl.expected.msl @@ -0,0 +1,21 @@ +#include + +using namespace metal; + +template +struct tint_array { + const constant T& operator[](size_t i) const constant { return elements[i]; } + device T& operator[](size_t i) device { return elements[i]; } + const device T& operator[](size_t i) const device { return elements[i]; } + thread T& operator[](size_t i) thread { return elements[i]; } + const thread T& operator[](size_t i) const thread { return elements[i]; } + threadgroup T& operator[](size_t i) threadgroup { return elements[i]; } + const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; } + T elements[N]; +}; + +fragment void tint_symbol() { + tint_array a = {}; + return; +} + diff --git a/test/tint/bug/tint/1598.wgsl.expected.spvasm b/test/tint/bug/tint/1598.wgsl.expected.spvasm new file mode 100644 index 0000000000..a51819575d --- /dev/null +++ b/test/tint/bug/tint/1598.wgsl.expected.spvasm @@ -0,0 +1,25 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 12 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint Fragment %main "main" + OpExecutionMode %main OriginUpperLeft + OpName %main "main" + OpName %a "a" + OpDecorate %_arr_int_uint_5 ArrayStride 4 + %void = OpTypeVoid + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %uint = OpTypeInt 32 0 + %uint_5 = OpConstant %uint 5 +%_arr_int_uint_5 = OpTypeArray %int %uint_5 +%_ptr_Function__arr_int_uint_5 = OpTypePointer Function %_arr_int_uint_5 + %11 = OpConstantNull %_arr_int_uint_5 + %main = OpFunction %void None %1 + %4 = OpLabel + %a = OpVariable %_ptr_Function__arr_int_uint_5 Function %11 + OpReturn + OpFunctionEnd diff --git a/test/tint/bug/tint/1598.wgsl.expected.wgsl b/test/tint/bug/tint/1598.wgsl.expected.wgsl new file mode 100644 index 0000000000..f9fa321478 --- /dev/null +++ b/test/tint/bug/tint/1598.wgsl.expected.wgsl @@ -0,0 +1,8 @@ +const MY_SIZE = 5u; + +type Arr = array; + +@fragment +fn main() { + var a : Arr; +}