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 <dsinclair@chromium.org> Commit-Queue: Dan Sinclair <dsinclair@chromium.org> Auto-Submit: Ben Clayton <bclayton@google.com>
This commit is contained in:
parent
688c6dc808
commit
eeda18d55e
|
@ -86,11 +86,17 @@ void SingleEntryPoint::Run(CloneContext& ctx, const DataMap& inputs, DataMap&) c
|
||||||
ctx.dst->AST().AddGlobalVariable(ctx.Clone(override));
|
ctx.dst->AST().AddGlobalVariable(ctx.Clone(override));
|
||||||
}
|
}
|
||||||
},
|
},
|
||||||
[&](const ast::Variable* v) { // var, let
|
[&](const ast::Var* var) {
|
||||||
if (referenced_vars.count(v)) {
|
if (referenced_vars.count(var)) {
|
||||||
ctx.dst->AST().AddGlobalVariable(ctx.Clone(v));
|
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) {
|
[&](const ast::Function* func) {
|
||||||
if (sem.Get(func)->HasAncestorEntryPoint(entry_point->symbol)) {
|
if (sem.Get(func)->HasAncestorEntryPoint(entry_point->symbol)) {
|
||||||
ctx.dst->AST().AddFunction(ctx.Clone(func));
|
ctx.dst->AST().AddFunction(ctx.Clone(func));
|
||||||
|
|
|
@ -217,8 +217,14 @@ fn comp_main2() {
|
||||||
)";
|
)";
|
||||||
|
|
||||||
auto* expect = R"(
|
auto* expect = R"(
|
||||||
|
const a : f32 = 1.0;
|
||||||
|
|
||||||
|
const b : f32 = 1.0;
|
||||||
|
|
||||||
const c : f32 = 1.0;
|
const c : f32 = 1.0;
|
||||||
|
|
||||||
|
const d : f32 = 1.0;
|
||||||
|
|
||||||
@compute @workgroup_size(1)
|
@compute @workgroup_size(1)
|
||||||
fn comp_main1() {
|
fn comp_main1() {
|
||||||
let local_c : f32 = c;
|
let local_c : f32 = c;
|
||||||
|
@ -536,5 +542,28 @@ fn comp_main1() {
|
||||||
EXPECT_EQ(expect, str(got));
|
EXPECT_EQ(expect, str(got));
|
||||||
}
|
}
|
||||||
|
|
||||||
|
TEST_F(SingleEntryPointTest, GlobalConstUsedAsArraySize) {
|
||||||
|
// See crbug.com/tint/1598
|
||||||
|
auto* src = R"(
|
||||||
|
const MY_SIZE = 5u;
|
||||||
|
|
||||||
|
type Arr = array<i32, MY_SIZE>;
|
||||||
|
|
||||||
|
@fragment
|
||||||
|
fn main() {
|
||||||
|
}
|
||||||
|
)";
|
||||||
|
|
||||||
|
auto* expect = src;
|
||||||
|
|
||||||
|
SingleEntryPoint::Config cfg("main");
|
||||||
|
|
||||||
|
DataMap data;
|
||||||
|
data.Add<SingleEntryPoint::Config>(cfg);
|
||||||
|
auto got = Run<SingleEntryPoint>(src, data);
|
||||||
|
|
||||||
|
EXPECT_EQ(expect, str(got));
|
||||||
|
}
|
||||||
|
|
||||||
} // namespace
|
} // namespace
|
||||||
} // namespace tint::transform
|
} // namespace tint::transform
|
||||||
|
|
|
@ -0,0 +1,8 @@
|
||||||
|
const MY_SIZE = 5u;
|
||||||
|
|
||||||
|
type Arr = array<i32, MY_SIZE>;
|
||||||
|
|
||||||
|
@fragment
|
||||||
|
fn main() {
|
||||||
|
var a : Arr;
|
||||||
|
}
|
|
@ -0,0 +1,4 @@
|
||||||
|
void main() {
|
||||||
|
int a[5] = (int[5])0;
|
||||||
|
return;
|
||||||
|
}
|
|
@ -0,0 +1,4 @@
|
||||||
|
void main() {
|
||||||
|
int a[5] = (int[5])0;
|
||||||
|
return;
|
||||||
|
}
|
|
@ -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;
|
||||||
|
}
|
|
@ -0,0 +1,21 @@
|
||||||
|
#include <metal_stdlib>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
|
||||||
|
template<typename T, size_t N>
|
||||||
|
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<int, 5> a = {};
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
|
@ -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
|
|
@ -0,0 +1,8 @@
|
||||||
|
const MY_SIZE = 5u;
|
||||||
|
|
||||||
|
type Arr = array<i32, MY_SIZE>;
|
||||||
|
|
||||||
|
@fragment
|
||||||
|
fn main() {
|
||||||
|
var a : Arr;
|
||||||
|
}
|
Loading…
Reference in New Issue