Make OFI transform push_constant-friendly.
Modify the OffsetFirstIndex transform to handle a pre-existing
push_constant block in the AST.
Bug: dawn:2185
Change-Id: If328138bab53cb94d4250f4d656ccccd524e1bbb
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/171860
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Stephen White <senorblanco@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
diff --git a/src/tint/cmd/tint/main.cc b/src/tint/cmd/tint/main.cc
index 38d695b..81a9e5b 100644
--- a/src/tint/cmd/tint/main.cc
+++ b/src/tint/cmd/tint/main.cc
@@ -1067,6 +1067,10 @@
gen_options.texture_builtins_from_uniform = std::move(textureBuiltinsFromUniform);
+ // Place the first_instance push constant member after user-defined push constants (if any).
+ gen_options.first_instance_offset =
+ inspector.GetEntryPoint(entry_point_name).push_constant_size;
+
auto result = tint::glsl::writer::Generate(prg, gen_options, entry_point_name);
if (result != tint::Success) {
tint::cmd::PrintWGSL(std::cerr, prg);
diff --git a/src/tint/lang/wgsl/ast/transform/offset_first_index.cc b/src/tint/lang/wgsl/ast/transform/offset_first_index.cc
index eb64f60..04bdd9a 100644
--- a/src/tint/lang/wgsl/ast/transform/offset_first_index.cc
+++ b/src/tint/lang/wgsl/ast/transform/offset_first_index.cc
@@ -136,23 +136,29 @@
return SkipTransform;
}
- // Abort on any use of push constants in the module.
+ Vector<const ast::StructMember*, 8> members;
+
+ const ast::Variable* push_constants_var = nullptr;
+
+ // Find first push_constant.
for (auto* global : src.AST().GlobalVariables()) {
if (auto* var = global->As<ast::Var>()) {
auto* v = src.Sem().Get(var);
- if (TINT_UNLIKELY(v->AddressSpace() == core::AddressSpace::kPushConstant)) {
- TINT_ICE()
- << "OffsetFirstIndex doesn't know how to handle module that already use push "
- "constants (yet)";
- return resolver::Resolve(b);
+ if (v->AddressSpace() == core::AddressSpace::kPushConstant) {
+ push_constants_var = var;
+ auto* str = v->Type()->UnwrapRef()->As<sem::Struct>();
+ if (!str) {
+ TINT_ICE() << "expected var<push_constant> type to be struct. Was "
+ "AddBlockAttribute run?";
+ }
+ for (auto* member : str->Members()) {
+ members.Push(ctx.CloneWithoutTransform(member->Declaration()));
+ }
}
}
}
- b.Enable(wgsl::Extension::kChromiumExperimentalPushConstant);
-
// Add push constant members and calculate byte offsets
- tint::Vector<const StructMember*, 8> members;
if (has_vertex_index) {
members.Push(b.Member(kFirstVertexName, b.ty.u32(),
Vector{b.MemberOffset(AInt(*cfg->first_vertex_offset))}));
@@ -161,10 +167,37 @@
members.Push(b.Member(kFirstInstanceName, b.ty.u32(),
Vector{b.MemberOffset(AInt(*cfg->first_instance_offset))}));
}
- auto struct_ = b.Structure(b.Symbols().New("PushConstants"), std::move(members));
- // Create a global to hold the uniform buffer
- Symbol buffer_name = b.Symbols().New("push_constants");
- b.GlobalVar(buffer_name, b.ty.Of(struct_), core::AddressSpace::kPushConstant);
+
+ auto new_struct = b.Structure(b.Symbols().New("PushConstants"), std::move(members));
+
+ Symbol buffer_name;
+
+ // If this is the first use of push constants, create a global to hold them.
+ if (!push_constants_var) {
+ b.Enable(wgsl::Extension::kChromiumExperimentalPushConstant);
+
+ buffer_name = b.Symbols().New("push_constants");
+ b.GlobalVar(buffer_name, b.ty.Of(new_struct), core::AddressSpace::kPushConstant);
+ } else {
+ buffer_name = ctx.Clone(push_constants_var->name->symbol);
+ }
+
+ // Replace all variable users of the old struct with the new struct.
+ if (push_constants_var) {
+ ctx.ReplaceAll([&](const ast::Variable* var) -> const ast::Variable* {
+ if (var->type == push_constants_var->type) {
+ if (var->As<ast::Parameter>()) {
+ return ctx.dst->Param(ctx.Clone(var->name->symbol), b.ty.Of(new_struct),
+ ctx.Clone(var->attributes));
+ } else {
+ return ctx.dst->Var(ctx.Clone(var->name->symbol), b.ty.Of(new_struct),
+ ctx.Clone(var->attributes),
+ core::AddressSpace::kPushConstant);
+ }
+ }
+ return nullptr;
+ });
+ }
// Fix up all references to the builtins with the offsets
ctx.ReplaceAll([&](const Expression* expr) -> const Expression* {
@@ -172,14 +205,14 @@
if (auto* user = sem->UnwrapLoad()->As<sem::VariableUser>()) {
auto it = builtin_vars.find(user->Variable());
if (it != builtin_vars.end()) {
- return ctx.dst->Add(ctx.CloneWithoutTransform(expr),
+ return ctx.dst->Add(b.Bitcast(b.ty.u32(), ctx.CloneWithoutTransform(expr)),
ctx.dst->MemberAccessor(buffer_name, it->second));
}
}
if (auto* access = sem->As<sem::StructMemberAccess>()) {
auto it = builtin_members.find(access->Member());
if (it != builtin_members.end()) {
- return ctx.dst->Add(ctx.CloneWithoutTransform(expr),
+ return ctx.dst->Add(b.Bitcast(b.ty.u32(), ctx.CloneWithoutTransform(expr)),
ctx.dst->MemberAccessor(buffer_name, it->second));
}
}
diff --git a/src/tint/lang/wgsl/ast/transform/offset_first_index.h b/src/tint/lang/wgsl/ast/transform/offset_first_index.h
index aa37559..72d2e96 100644
--- a/src/tint/lang/wgsl/ast/transform/offset_first_index.h
+++ b/src/tint/lang/wgsl/ast/transform/offset_first_index.h
@@ -35,7 +35,7 @@
/// Adds firstVertex/Instance (injected via push constants) to
/// vertex/instance index builtins.
///
-/// This transform assumes that Name transform has been run before.
+/// This transform assumes that the Name and AddBlockAttribute transforms have been run before.
///
/// Some shading languages start vertex and instance numbering at 0,
/// regardless of the firstVertex/firstInstance value specified. This transform
diff --git a/src/tint/lang/wgsl/ast/transform/offset_first_index_test.cc b/src/tint/lang/wgsl/ast/transform/offset_first_index_test.cc
index 37a79a9..7d026ed 100644
--- a/src/tint/lang/wgsl/ast/transform/offset_first_index_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/offset_first_index_test.cc
@@ -154,7 +154,7 @@
@vertex
fn entry(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32> {
- test((vert_idx + push_constants.first_vertex));
+ test((bitcast<u32>(vert_idx) + push_constants.first_vertex));
return vec4<f32>();
}
)";
@@ -191,7 +191,7 @@
@vertex
fn entry(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32> {
- test((vert_idx + push_constants.first_vertex));
+ test((bitcast<u32>(vert_idx) + push_constants.first_vertex));
return vec4<f32>();
}
@@ -238,7 +238,7 @@
@vertex
fn entry(@builtin(instance_index) inst_idx : u32) -> @builtin(position) vec4<f32> {
- test((inst_idx + push_constants.first_instance));
+ test((bitcast<u32>(inst_idx) + push_constants.first_instance));
return vec4<f32>();
}
)";
@@ -277,7 +277,7 @@
@vertex
fn entry(@builtin(instance_index) inst_idx : u32) -> @builtin(position) vec4<f32> {
- test((inst_idx + push_constants.first_instance));
+ test((bitcast<u32>(inst_idx) + push_constants.first_instance));
return vec4<f32>();
}
@@ -336,7 +336,7 @@
@vertex
fn entry(inputs : Inputs) -> @builtin(position) vec4<f32> {
- test((inputs.instance_idx + push_constants.first_instance), (inputs.vert_idx + push_constants.first_vertex));
+ test((bitcast<u32>(inputs.instance_idx) + push_constants.first_instance), (bitcast<u32>(inputs.vert_idx) + push_constants.first_vertex));
return vec4<f32>();
}
)";
@@ -380,7 +380,7 @@
@vertex
fn entry(inputs : Inputs) -> @builtin(position) vec4<f32> {
- test((inputs.instance_idx + push_constants.first_instance), (inputs.vert_idx + push_constants.first_vertex));
+ test((bitcast<u32>(inputs.instance_idx) + push_constants.first_instance), (bitcast<u32>(inputs.vert_idx) + push_constants.first_vertex));
return vec4<f32>();
}
@@ -444,7 +444,7 @@
@vertex
fn entry(inputs : Inputs) -> @builtin(position) vec4<f32> {
- test((inputs.instance_idx + push_constants.first_instance), inputs.vert_idx);
+ test((bitcast<u32>(inputs.instance_idx) + push_constants.first_instance), inputs.vert_idx);
return vec4<f32>();
}
)";
@@ -497,7 +497,7 @@
@vertex
fn entry(inputs : Inputs) -> @builtin(position) vec4<f32> {
- test(inputs.instance_idx, (inputs.vert_idx + push_constants.first_vertex));
+ test(inputs.instance_idx, (bitcast<u32>(inputs.vert_idx) + push_constants.first_vertex));
return vec4<f32>();
}
)";
@@ -572,7 +572,7 @@
@vertex
fn entry(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32> {
- return vec4<f32>((f32((vert_idx + push_constants.first_vertex)) + p.f));
+ return vec4<f32>((f32((bitcast<u32>(vert_idx) + push_constants.first_vertex)) + p.f));
}
)";
@@ -609,7 +609,7 @@
@vertex
fn entry(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32> {
- return vec4<f32>(f32((vert_idx + push_constants_1.first_vertex)));
+ return vec4<f32>(f32((bitcast<u32>(vert_idx) + push_constants_1.first_vertex)));
}
)";
@@ -657,7 +657,7 @@
@vertex
fn entry(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32> {
- func2((vert_idx + push_constants.first_vertex));
+ func2((bitcast<u32>(vert_idx) + push_constants.first_vertex));
return vec4<f32>();
}
)";
@@ -698,7 +698,7 @@
@vertex
fn entry(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32> {
- func2((vert_idx + push_constants.first_vertex));
+ func2((bitcast<u32>(vert_idx) + push_constants.first_vertex));
return vec4<f32>();
}
@@ -761,19 +761,19 @@
@vertex
fn entry_a(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32> {
- func((vert_idx + push_constants.first_vertex));
+ func((bitcast<u32>(vert_idx) + push_constants.first_vertex));
return vec4<f32>();
}
@vertex
fn entry_b(@builtin(vertex_index) vert_idx : u32, @builtin(instance_index) inst_idx : u32) -> @builtin(position) vec4<f32> {
- func(((vert_idx + push_constants.first_vertex) + (inst_idx + push_constants.first_instance)));
+ func(((bitcast<u32>(vert_idx) + push_constants.first_vertex) + (bitcast<u32>(inst_idx) + push_constants.first_instance)));
return vec4<f32>();
}
@vertex
fn entry_c(@builtin(instance_index) inst_idx : u32) -> @builtin(position) vec4<f32> {
- func((inst_idx + push_constants.first_instance));
+ func((bitcast<u32>(inst_idx) + push_constants.first_instance));
return vec4<f32>();
}
)";
@@ -824,19 +824,19 @@
@vertex
fn entry_a(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32> {
- func((vert_idx + push_constants.first_vertex));
+ func((bitcast<u32>(vert_idx) + push_constants.first_vertex));
return vec4<f32>();
}
@vertex
fn entry_b(@builtin(vertex_index) vert_idx : u32, @builtin(instance_index) inst_idx : u32) -> @builtin(position) vec4<f32> {
- func(((vert_idx + push_constants.first_vertex) + (inst_idx + push_constants.first_instance)));
+ func(((bitcast<u32>(vert_idx) + push_constants.first_vertex) + (bitcast<u32>(inst_idx) + push_constants.first_instance)));
return vec4<f32>();
}
@vertex
fn entry_c(@builtin(instance_index) inst_idx : u32) -> @builtin(position) vec4<f32> {
- func((inst_idx + push_constants.first_instance));
+ func((bitcast<u32>(inst_idx) + push_constants.first_instance));
return vec4<f32>();
}
diff --git a/src/tint/lang/wgsl/inspector/entry_point.h b/src/tint/lang/wgsl/inspector/entry_point.h
index 24164d0..3709ade 100644
--- a/src/tint/lang/wgsl/inspector/entry_point.h
+++ b/src/tint/lang/wgsl/inspector/entry_point.h
@@ -165,6 +165,8 @@
std::optional<WorkgroupSize> workgroup_size;
/// The total size in bytes of all Workgroup storage-class storage accessed via the entry point.
uint32_t workgroup_storage_size = 0;
+ /// The total size in bytes of all push_constant variables accessed by the entry point.
+ uint32_t push_constant_size = 0;
/// List of the input variable accessed via this entry point.
std::vector<StageVariable> input_variables;
/// List of the output variable accessed via this entry point.
diff --git a/src/tint/lang/wgsl/inspector/inspector.cc b/src/tint/lang/wgsl/inspector/inspector.cc
index a69fb67..8efa350 100644
--- a/src/tint/lang/wgsl/inspector/inspector.cc
+++ b/src/tint/lang/wgsl/inspector/inspector.cc
@@ -174,6 +174,8 @@
}
}
+ entry_point.push_constant_size = ComputePushConstantSize(func);
+
for (auto* param : sem->Parameters()) {
AddEntryPointInOutVariables(param->Declaration()->name->symbol.Name(),
param->Declaration()->name->symbol.Name(), param->Type(),
@@ -927,6 +929,18 @@
return total_size;
}
+uint32_t Inspector::ComputePushConstantSize(const ast::Function* func) const {
+ uint32_t size = 0;
+ auto* func_sem = program_.Sem().Get(func);
+ for (const sem::Variable* var : func_sem->TransitivelyReferencedGlobals()) {
+ if (var->AddressSpace() == core::AddressSpace::kPushConstant) {
+ size += var->Type()->UnwrapRef()->Size();
+ }
+ }
+
+ return size;
+}
+
std::vector<PixelLocalMemberType> Inspector::ComputePixelLocalMemberTypes(
const ast::Function* func) const {
auto* func_sem = program_.Sem().Get(func);
diff --git a/src/tint/lang/wgsl/inspector/inspector.h b/src/tint/lang/wgsl/inspector/inspector.h
index af839df..736b2da 100644
--- a/src/tint/lang/wgsl/inspector/inspector.h
+++ b/src/tint/lang/wgsl/inspector/inspector.h
@@ -264,6 +264,10 @@
/// @returns the total size in bytes of all Workgroup storage-class storage accessed via func.
uint32_t ComputeWorkgroupStorageSize(const ast::Function* func) const;
+ /// @param func the root function of the callgraph to consider for the computation.
+ /// @returns the total size in bytes of all push_constant variables accessed via func.
+ uint32_t ComputePushConstantSize(const ast::Function* func) const;
+
/// @param func the root function of the callgraph to consider for the computation
/// @returns the list of member types for the `pixel_local` variable accessed via func, if any.
std::vector<PixelLocalMemberType> ComputePixelLocalMemberTypes(const ast::Function* func) const;
diff --git a/src/tint/lang/wgsl/inspector/inspector_test.cc b/src/tint/lang/wgsl/inspector/inspector_test.cc
index e3d97bb..8d1a4b3 100644
--- a/src/tint/lang/wgsl/inspector/inspector_test.cc
+++ b/src/tint/lang/wgsl/inspector/inspector_test.cc
@@ -258,6 +258,85 @@
EXPECT_EQ(1u, workgroup_size->z);
}
+// Test that push_constant_size is zero if there are no push constants.
+TEST_F(InspectorGetEntryPointTest, PushConstantSizeNone) {
+ MakeEmptyBodyFunction("foo", Vector{
+ Stage(ast::PipelineStage::kFragment),
+ });
+
+ Inspector& inspector = Build();
+
+ auto result = inspector.GetEntryPoints();
+ ASSERT_FALSE(inspector.has_error()) << inspector.error();
+
+ ASSERT_EQ(1u, result.size());
+ EXPECT_EQ(0u, result[0].push_constant_size);
+}
+
+// Test that push_constant_size is 4 (bytes) if there is a single F32 push constant.
+TEST_F(InspectorGetEntryPointTest, PushConstantSizeOneWord) {
+ Enable(wgsl::Extension::kChromiumExperimentalPushConstant);
+ GlobalVar("pc", core::AddressSpace::kPushConstant, ty.f32());
+ MakePlainGlobalReferenceBodyFunction("foo", "pc", ty.f32(),
+ Vector{
+ Stage(ast::PipelineStage::kFragment),
+ });
+
+ Inspector& inspector = Build();
+
+ auto result = inspector.GetEntryPoints();
+ ASSERT_FALSE(inspector.has_error()) << inspector.error();
+
+ ASSERT_EQ(1u, result.size());
+ EXPECT_EQ(4u, result[0].push_constant_size);
+}
+
+// Test that push_constant_size is 12 (bytes) if there is a struct containing one
+// each of i32, f32 and u32.
+TEST_F(InspectorGetEntryPointTest, PushConstantSizeThreeWords) {
+ Enable(wgsl::Extension::kChromiumExperimentalPushConstant);
+ auto* pc_struct_type =
+ MakeStructType("PushConstantStruct", Vector{ty.i32(), ty.f32(), ty.u32()});
+ GlobalVar("pc", core::AddressSpace::kPushConstant, ty.Of(pc_struct_type));
+ MakePlainGlobalReferenceBodyFunction("foo", "pc", ty.Of(pc_struct_type),
+ Vector{
+ Stage(ast::PipelineStage::kFragment),
+ });
+
+ Inspector& inspector = Build();
+
+ auto result = inspector.GetEntryPoints();
+ ASSERT_FALSE(inspector.has_error()) << inspector.error();
+
+ ASSERT_EQ(1u, result.size());
+ EXPECT_EQ(12u, result[0].push_constant_size);
+}
+
+// Test that push_constant_size is 4 (bytes) if there are two push constants,
+// one used by the entry point containing an f32, and one unused by the entry
+// point containing a struct of size 12 bytes.
+TEST_F(InspectorGetEntryPointTest, PushConstantSizeTwoConstants) {
+ Enable(wgsl::Extension::kChromiumExperimentalPushConstant);
+ auto* unused_struct_type =
+ MakeStructType("PushConstantStruct", Vector{ty.i32(), ty.f32(), ty.u32()});
+ GlobalVar("unused", core::AddressSpace::kPushConstant, ty.Of(unused_struct_type));
+ GlobalVar("pc", core::AddressSpace::kPushConstant, ty.f32());
+ MakePlainGlobalReferenceBodyFunction("foo", "pc", ty.f32(),
+ Vector{
+ Stage(ast::PipelineStage::kFragment),
+ });
+
+ Inspector& inspector = Build();
+
+ auto result = inspector.GetEntryPoints();
+ ASSERT_FALSE(inspector.has_error()) << inspector.error();
+
+ ASSERT_EQ(1u, result.size());
+
+ // Check that the result only includes the single f32 push constant.
+ EXPECT_EQ(4u, result[0].push_constant_size);
+}
+
TEST_F(InspectorGetEntryPointTest, NonDefaultWorkgroupSize) {
MakeEmptyBodyFunction("foo", Vector{
Stage(ast::PipelineStage::kCompute),
diff --git a/test/tint/bug/chromium/1251009.wgsl.expected.glsl b/test/tint/bug/chromium/1251009.wgsl.expected.glsl
index 5ae88f4..f4292a7 100644
--- a/test/tint/bug/chromium/1251009.wgsl.expected.glsl
+++ b/test/tint/bug/chromium/1251009.wgsl.expected.glsl
@@ -4,6 +4,11 @@
layout(location = 1) in uint loc1_1;
layout(location = 2) in uint loc1_2;
layout(location = 3) in vec4 loc3_1;
+struct PushConstants {
+ uint first_instance;
+};
+
+layout(location=0) uniform PushConstants push_constants;
struct VertexInputs0 {
uint vertex_index;
int loc0;
@@ -15,7 +20,7 @@
};
vec4 tint_symbol(VertexInputs0 inputs0, uint loc1, uint instance_index, VertexInputs1 inputs1) {
- uint foo = (inputs0.vertex_index + instance_index);
+ uint foo = (inputs0.vertex_index + (instance_index + push_constants.first_instance));
return vec4(0.0f);
}
diff --git a/test/tint/bug/tint/824.wgsl.expected.glsl b/test/tint/bug/tint/824.wgsl.expected.glsl
index 074047c..b2dc49c 100644
--- a/test/tint/bug/tint/824.wgsl.expected.glsl
+++ b/test/tint/bug/tint/824.wgsl.expected.glsl
@@ -1,6 +1,11 @@
#version 310 es
layout(location = 0) out vec4 color_1;
+struct PushConstants {
+ uint first_instance;
+};
+
+layout(location=0) uniform PushConstants push_constants;
struct Output {
vec4 Position;
vec4 color;
@@ -8,11 +13,11 @@
Output tint_symbol(uint VertexIndex, uint InstanceIndex) {
vec2 zv[4] = vec2[4](vec2(0.20000000298023223877f), vec2(0.30000001192092895508f), vec2(-0.10000000149011611938f), vec2(1.10000002384185791016f));
- float z = zv[InstanceIndex].x;
+ float z = zv[(InstanceIndex + push_constants.first_instance)].x;
Output tint_symbol_1 = Output(vec4(0.0f, 0.0f, 0.0f, 0.0f), vec4(0.0f, 0.0f, 0.0f, 0.0f));
tint_symbol_1.Position = vec4(0.5f, 0.5f, z, 1.0f);
vec4 colors[4] = vec4[4](vec4(1.0f, 0.0f, 0.0f, 1.0f), vec4(0.0f, 1.0f, 0.0f, 1.0f), vec4(0.0f, 0.0f, 1.0f, 1.0f), vec4(1.0f));
- tint_symbol_1.color = colors[InstanceIndex];
+ tint_symbol_1.color = colors[(InstanceIndex + push_constants.first_instance)];
return tint_symbol_1;
}
diff --git a/test/tint/types/functions/shader_io/vertex_input_builtins.wgsl.expected.glsl b/test/tint/types/functions/shader_io/vertex_input_builtins.wgsl.expected.glsl
index e5e32ae..eaa2763 100644
--- a/test/tint/types/functions/shader_io/vertex_input_builtins.wgsl.expected.glsl
+++ b/test/tint/types/functions/shader_io/vertex_input_builtins.wgsl.expected.glsl
@@ -1,7 +1,12 @@
#version 310 es
+struct PushConstants {
+ uint first_instance;
+};
+
+layout(location=0) uniform PushConstants push_constants;
vec4 tint_symbol(uint vertex_index, uint instance_index) {
- uint foo = (vertex_index + instance_index);
+ uint foo = (vertex_index + (instance_index + push_constants.first_instance));
return vec4(0.0f);
}
diff --git a/test/tint/types/functions/shader_io/vertex_input_builtins_struct.wgsl.expected.glsl b/test/tint/types/functions/shader_io/vertex_input_builtins_struct.wgsl.expected.glsl
index 7daad03..4f64082 100644
--- a/test/tint/types/functions/shader_io/vertex_input_builtins_struct.wgsl.expected.glsl
+++ b/test/tint/types/functions/shader_io/vertex_input_builtins_struct.wgsl.expected.glsl
@@ -1,12 +1,17 @@
#version 310 es
+struct PushConstants {
+ uint first_instance;
+};
+
+layout(location=0) uniform PushConstants push_constants;
struct VertexInputs {
uint vertex_index;
uint instance_index;
};
vec4 tint_symbol(VertexInputs inputs) {
- uint foo = (inputs.vertex_index + inputs.instance_index);
+ uint foo = (inputs.vertex_index + (inputs.instance_index + push_constants.first_instance));
return vec4(0.0f);
}
diff --git a/test/tint/types/functions/shader_io/vertex_input_mixed.wgsl.expected.glsl b/test/tint/types/functions/shader_io/vertex_input_mixed.wgsl.expected.glsl
index 2470a13..2208401 100644
--- a/test/tint/types/functions/shader_io/vertex_input_mixed.wgsl.expected.glsl
+++ b/test/tint/types/functions/shader_io/vertex_input_mixed.wgsl.expected.glsl
@@ -4,6 +4,11 @@
layout(location = 1) in uint loc1_1;
layout(location = 2) in float loc2_1;
layout(location = 3) in vec4 loc3_1;
+struct PushConstants {
+ uint first_instance;
+};
+
+layout(location=0) uniform PushConstants push_constants;
struct VertexInputs0 {
uint vertex_index;
int loc0;
@@ -15,7 +20,7 @@
};
vec4 tint_symbol(VertexInputs0 inputs0, uint loc1, uint instance_index, VertexInputs1 inputs1) {
- uint foo = (inputs0.vertex_index + instance_index);
+ uint foo = (inputs0.vertex_index + (instance_index + push_constants.first_instance));
int i = inputs0.loc0;
uint u = loc1;
float f = inputs1.loc2;
diff --git a/test/tint/types/functions/shader_io/vertex_input_mixed_f16.wgsl.expected.glsl b/test/tint/types/functions/shader_io/vertex_input_mixed_f16.wgsl.expected.glsl
index 8d16b68..1cd4f6a 100644
--- a/test/tint/types/functions/shader_io/vertex_input_mixed_f16.wgsl.expected.glsl
+++ b/test/tint/types/functions/shader_io/vertex_input_mixed_f16.wgsl.expected.glsl
@@ -7,6 +7,11 @@
layout(location = 3) in vec4 loc3_1;
layout(location = 5) in f16vec3 loc5_1;
layout(location = 4) in float16_t loc4_1;
+struct PushConstants {
+ uint first_instance;
+};
+
+layout(location=0) uniform PushConstants push_constants;
struct VertexInputs0 {
uint vertex_index;
int loc0;
@@ -19,7 +24,7 @@
};
vec4 tint_symbol(VertexInputs0 inputs0, uint loc1, uint instance_index, VertexInputs1 inputs1, float16_t loc4) {
- uint foo = (inputs0.vertex_index + instance_index);
+ uint foo = (inputs0.vertex_index + (instance_index + push_constants.first_instance));
int i = inputs0.loc0;
uint u = loc1;
float f = inputs1.loc2;
diff --git a/test/tint/var/uses/instance_index.wgsl b/test/tint/var/uses/instance_index.wgsl
new file mode 100644
index 0000000..6305dac
--- /dev/null
+++ b/test/tint/var/uses/instance_index.wgsl
@@ -0,0 +1,3 @@
+@vertex fn main(@builtin(instance_index) b : u32) -> @builtin(position) vec4<f32> {
+ return vec4<f32>(f32(b));
+}
diff --git a/test/tint/var/uses/instance_index.wgsl.expected.glsl b/test/tint/var/uses/instance_index.wgsl.expected.glsl
new file mode 100644
index 0000000..c543334
--- /dev/null
+++ b/test/tint/var/uses/instance_index.wgsl.expected.glsl
@@ -0,0 +1,19 @@
+#version 310 es
+
+struct PushConstants {
+ uint first_instance;
+};
+
+layout(location=0) uniform PushConstants push_constants;
+vec4 tint_symbol(uint b) {
+ return vec4(float((b + push_constants.first_instance)));
+}
+
+void main() {
+ gl_PointSize = 1.0;
+ vec4 inner_result = tint_symbol(uint(gl_InstanceID));
+ gl_Position = inner_result;
+ gl_Position.y = -(gl_Position.y);
+ gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+ return;
+}
diff --git a/test/tint/var/uses/push_constant.wgsl.expected.msl b/test/tint/var/uses/push_constant.wgsl.expected.msl
index 925fc2a..c106eee 100644
--- a/test/tint/var/uses/push_constant.wgsl.expected.msl
+++ b/test/tint/var/uses/push_constant.wgsl.expected.msl
@@ -43,6 +43,3 @@
Failed to generate: error: unhandled module-scope address space (push_constant)
error: unhandled module-scope address space (push_constant)
error: unhandled module-scope address space (push_constant)
-error: unhandled module-scope address space (push_constant)
-error: unhandled module-scope address space (push_constant)
-error: unhandled module-scope address space (push_constant)
diff --git a/test/tint/var/uses/push_constant_and_instance_index.wgsl b/test/tint/var/uses/push_constant_and_instance_index.wgsl
new file mode 100644
index 0000000..5f9d6c4
--- /dev/null
+++ b/test/tint/var/uses/push_constant_and_instance_index.wgsl
@@ -0,0 +1,7 @@
+enable chromium_experimental_push_constant;
+
+var<push_constant> a : f32;
+
+@vertex fn main(@builtin(instance_index) b : u32) -> @builtin(position) vec4<f32> {
+ return vec4<f32>(a + f32(b));
+}
diff --git a/test/tint/var/uses/push_constant_and_instance_index.wgsl.expected.dxc.hlsl b/test/tint/var/uses/push_constant_and_instance_index.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..b85ac2b
--- /dev/null
+++ b/test/tint/var/uses/push_constant_and_instance_index.wgsl.expected.dxc.hlsl
@@ -0,0 +1,13 @@
+SKIP: FAILED
+
+
+enable chromium_experimental_push_constant;
+
+var<push_constant> a : f32;
+
+@vertex
+fn main(@builtin(instance_index) b : u32) -> @builtin(position) vec4<f32> {
+ return vec4<f32>((a + f32(b)));
+}
+
+Failed to generate: error: unhandled address space push_constant
diff --git a/test/tint/var/uses/push_constant_and_instance_index.wgsl.expected.fxc.hlsl b/test/tint/var/uses/push_constant_and_instance_index.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..b85ac2b
--- /dev/null
+++ b/test/tint/var/uses/push_constant_and_instance_index.wgsl.expected.fxc.hlsl
@@ -0,0 +1,13 @@
+SKIP: FAILED
+
+
+enable chromium_experimental_push_constant;
+
+var<push_constant> a : f32;
+
+@vertex
+fn main(@builtin(instance_index) b : u32) -> @builtin(position) vec4<f32> {
+ return vec4<f32>((a + f32(b)));
+}
+
+Failed to generate: error: unhandled address space push_constant
diff --git a/test/tint/var/uses/push_constant_and_instance_index.wgsl.expected.glsl b/test/tint/var/uses/push_constant_and_instance_index.wgsl.expected.glsl
new file mode 100644
index 0000000..7989941
--- /dev/null
+++ b/test/tint/var/uses/push_constant_and_instance_index.wgsl.expected.glsl
@@ -0,0 +1,20 @@
+#version 310 es
+
+struct PushConstants {
+ float inner;
+ uint first_instance;
+};
+
+layout(location=0) uniform PushConstants a;
+vec4 tint_symbol(uint b) {
+ return vec4((a.inner + float((b + a.first_instance))));
+}
+
+void main() {
+ gl_PointSize = 1.0;
+ vec4 inner_result = tint_symbol(uint(gl_InstanceID));
+ gl_Position = inner_result;
+ gl_Position.y = -(gl_Position.y);
+ gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+ return;
+}
diff --git a/test/tint/var/uses/push_constant_and_instance_index.wgsl.expected.msl b/test/tint/var/uses/push_constant_and_instance_index.wgsl.expected.msl
new file mode 100644
index 0000000..db4c340
--- /dev/null
+++ b/test/tint/var/uses/push_constant_and_instance_index.wgsl.expected.msl
@@ -0,0 +1,9 @@
+SKIP: FAILED
+
+../../src/tint/lang/msl/writer/ast_raise/module_scope_var_to_entry_point_param.cc:220 internal compiler error: unhandled module-scope address space (push_constant)
+********************************************************************
+* The tint shader compiler has encountered an unexpected error. *
+* *
+* Please help us fix this issue by submitting a bug report at *
+* crbug.com/tint with the source program that triggered the bug. *
+********************************************************************
diff --git a/test/tint/var/uses/push_constant_and_instance_index.wgsl.expected.spvasm b/test/tint/var/uses/push_constant_and_instance_index.wgsl.expected.spvasm
new file mode 100644
index 0000000..2592150
--- /dev/null
+++ b/test/tint/var/uses/push_constant_and_instance_index.wgsl.expected.spvasm
@@ -0,0 +1,60 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 33
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Vertex %main "main" %b_1 %value %vertex_point_size
+ OpName %b_1 "b_1"
+ OpName %value "value"
+ OpName %vertex_point_size "vertex_point_size"
+ OpName %a_block "a_block"
+ OpMemberName %a_block 0 "inner"
+ OpName %a "a"
+ OpName %main_inner "main_inner"
+ OpName %b "b"
+ OpName %main "main"
+ OpDecorate %b_1 BuiltIn InstanceIndex
+ OpDecorate %value BuiltIn Position
+ OpDecorate %vertex_point_size BuiltIn PointSize
+ OpDecorate %a_block Block
+ OpMemberDecorate %a_block 0 Offset 0
+ %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+ %b_1 = OpVariable %_ptr_Input_uint Input
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+ %8 = OpConstantNull %v4float
+ %value = OpVariable %_ptr_Output_v4float Output %8
+%_ptr_Output_float = OpTypePointer Output %float
+ %11 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %11
+ %a_block = OpTypeStruct %float
+%_ptr_PushConstant_a_block = OpTypePointer PushConstant %a_block
+ %a = OpVariable %_ptr_PushConstant_a_block PushConstant
+ %15 = OpTypeFunction %v4float %uint
+ %uint_0 = OpConstant %uint 0
+%_ptr_PushConstant_float = OpTypePointer PushConstant %float
+ %void = OpTypeVoid
+ %26 = OpTypeFunction %void
+ %float_1 = OpConstant %float 1
+ %main_inner = OpFunction %v4float None %15
+ %b = OpFunctionParameter %uint
+ %18 = OpLabel
+ %21 = OpAccessChain %_ptr_PushConstant_float %a %uint_0
+ %22 = OpLoad %float %21
+ %23 = OpConvertUToF %float %b
+ %24 = OpFAdd %float %22 %23
+ %25 = OpCompositeConstruct %v4float %24 %24 %24 %24
+ OpReturnValue %25
+ OpFunctionEnd
+ %main = OpFunction %void None %26
+ %29 = OpLabel
+ %31 = OpLoad %uint %b_1
+ %30 = OpFunctionCall %v4float %main_inner %31
+ OpStore %value %30
+ OpStore %vertex_point_size %float_1
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/var/uses/push_constant_and_instance_index.wgsl.expected.wgsl b/test/tint/var/uses/push_constant_and_instance_index.wgsl.expected.wgsl
new file mode 100644
index 0000000..3b1e555
--- /dev/null
+++ b/test/tint/var/uses/push_constant_and_instance_index.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+enable chromium_experimental_push_constant;
+
+var<push_constant> a : f32;
+
+@vertex
+fn main(@builtin(instance_index) b : u32) -> @builtin(position) vec4<f32> {
+ return vec4<f32>((a + f32(b)));
+}
