Fix structure builtin emission
Fix the Renamer to preserve builtin structure member names.
Fix the HLSL writer to emit the modf / frexp result type even if there is no private / function storage usage of the types.
Fixed: chromium:1236161
Change-Id: I93b9d92980682f9a9cb090d07b04e4c3f6a2f705
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/60922
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
diff --git a/src/transform/renamer.cc b/src/transform/renamer.cc
index 712f70e..956d407 100644
--- a/src/transform/renamer.cc
+++ b/src/transform/renamer.cc
@@ -896,7 +896,8 @@
// Disable auto-cloning of symbols, since we want to rename them.
CloneContext ctx(&out, in, false);
- // Swizzles and intrinsic calls need to keep their symbols preserved.
+ // Swizzles, intrinsic calls and builtin structure members need to keep their
+ // symbols preserved.
std::unordered_set<ast::IdentifierExpression*> preserve;
for (auto* node : in->ASTNodes().Objects()) {
if (auto* member = node->As<ast::MemberAccessorExpression>()) {
@@ -908,6 +909,12 @@
}
if (sem->Is<sem::Swizzle>()) {
preserve.emplace(member->member());
+ } else if (auto* str_expr = in->Sem().Get(member->structure())) {
+ if (auto* ty = str_expr->Type()->UnwrapRef()->As<sem::Struct>()) {
+ if (ty->Declaration() == nullptr) { // Builtin structure
+ preserve.emplace(member->member());
+ }
+ }
}
} else if (auto* call = node->As<ast::CallExpression>()) {
auto* sem = in->Sem().Get(call);
diff --git a/src/transform/renamer_test.cc b/src/transform/renamer_test.cc
index cfd77d5..9b9ff3f 100644
--- a/src/transform/renamer_test.cc
+++ b/src/transform/renamer_test.cc
@@ -149,6 +149,41 @@
EXPECT_THAT(data->remappings, ContainerEq(expected_remappings));
}
+TEST_F(RenamerTest, PreserveBuiltinTypes) {
+ auto* src = R"(
+[[stage(compute), workgroup_size(1)]]
+fn entry() {
+ var a = modf(1.0).whole;
+ var b = modf(1.0).fract;
+ var c = frexp(1.0).sig;
+ var d = frexp(1.0).exp;
+}
+)";
+
+ auto* expect = R"(
+[[stage(compute), workgroup_size(1)]]
+fn tint_symbol() {
+ var tint_symbol_1 = modf(1.0).whole;
+ var tint_symbol_2 = modf(1.0).fract;
+ var tint_symbol_3 = frexp(1.0).sig;
+ var tint_symbol_4 = frexp(1.0).exp;
+}
+)";
+
+ auto got = Run<Renamer>(src);
+
+ EXPECT_EQ(expect, str(got));
+
+ auto* data = got.data.Get<Renamer::Data>();
+
+ ASSERT_NE(data, nullptr);
+ Renamer::Data::Remappings expected_remappings = {
+ {"entry", "tint_symbol"}, {"a", "tint_symbol_1"}, {"b", "tint_symbol_2"},
+ {"c", "tint_symbol_3"}, {"d", "tint_symbol_4"},
+ };
+ EXPECT_THAT(data->remappings, ContainerEq(expected_remappings));
+}
+
TEST_F(RenamerTest, AttemptSymbolCollision) {
auto* src = R"(
[[stage(vertex)]]
diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index db40ef3..c36db0d 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -139,8 +139,20 @@
return false;
}
} else if (auto* str = decl->As<ast::Struct>()) {
- if (!EmitStructType(current_buffer_, builder_.Sem().Get(str))) {
- return false;
+ auto* ty = builder_.Sem().Get(str);
+ auto storage_class_uses = ty->StorageClassUsage();
+ if (storage_class_uses.size() !=
+ (storage_class_uses.count(ast::StorageClass::kStorage) +
+ storage_class_uses.count(ast::StorageClass::kUniform))) {
+ // The structure is used as something other than a storage buffer or
+ // uniform buffer, so it needs to be emitted.
+ // Storage buffer are read and written to via a ByteAddressBuffer
+ // instead of true structure.
+ // Structures used as uniform buffer are read from an array of vectors
+ // instead of true structure.
+ if (!EmitStructType(current_buffer_, ty)) {
+ return false;
+ }
}
} else if (auto* func = decl->As<ast::Function>()) {
if (func->IsEntryPoint()) {
@@ -3135,19 +3147,6 @@
}
bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) {
- auto storage_class_uses = str->StorageClassUsage();
- if (storage_class_uses.size() ==
- (storage_class_uses.count(ast::StorageClass::kStorage) +
- storage_class_uses.count(ast::StorageClass::kUniform))) {
- // The only use of the structure is as a storage buffer and / or uniform
- // buffer.
- // Structures used as storage buffer are read and written to via a
- // ByteAddressBuffer instead of true structure.
- // Structures used as uniform buffer are read from an array of vectors
- // instead of true structure.
- return true;
- }
-
line(b) << "struct " << StructName(str) << " {";
{
ScopedIndent si(b);
diff --git a/src/writer/hlsl/generator_impl_type_test.cc b/src/writer/hlsl/generator_impl_type_test.cc
index b9a8fc6..90df7cd 100644
--- a/src/writer/hlsl/generator_impl_type_test.cc
+++ b/src/writer/hlsl/generator_impl_type_test.cc
@@ -198,10 +198,8 @@
GeneratorImpl& gen = Build();
- TextGenerator::TextBuffer buf;
- auto* sem_s = program->TypeOf(s)->As<sem::Struct>();
- ASSERT_TRUE(gen.EmitStructType(&buf, sem_s)) << gen.error();
- EXPECT_EQ(buf.String(), "");
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), "RWByteAddressBuffer g : register(u0, space0);\n");
}
TEST_F(HlslGeneratorImplTest_Type, EmitType_Struct) {
diff --git a/test/bug/chromium/1236161.wgsl b/test/bug/chromium/1236161.wgsl
new file mode 100644
index 0000000..1a4f3c9
--- /dev/null
+++ b/test/bug/chromium/1236161.wgsl
@@ -0,0 +1,3 @@
+// flags: --transform renamer
+
+fn i(){let s=modf(1.).whole;}
diff --git a/test/bug/chromium/1236161.wgsl.expected.hlsl b/test/bug/chromium/1236161.wgsl.expected.hlsl
new file mode 100644
index 0000000..a0d44d9
--- /dev/null
+++ b/test/bug/chromium/1236161.wgsl.expected.hlsl
@@ -0,0 +1,19 @@
+struct modf_result {
+ float fract;
+ float whole;
+};
+modf_result tint_modf(float param_0) {
+ float whole;
+ float fract = modf(param_0, whole);
+ modf_result result = {fract, whole};
+ return result;
+}
+
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+ return;
+}
+
+void i() {
+ const float s = tint_modf(1.0f).whole;
+}
diff --git a/test/bug/chromium/1236161.wgsl.expected.msl b/test/bug/chromium/1236161.wgsl.expected.msl
new file mode 100644
index 0000000..4eb36e6
--- /dev/null
+++ b/test/bug/chromium/1236161.wgsl.expected.msl
@@ -0,0 +1,18 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+struct modf_result {
+ float fract;
+ float whole;
+};
+modf_result tint_modf(float param_0) {
+ float whole;
+ float fract = modf(param_0, whole);
+ return {fract, whole};
+}
+
+void i() {
+ float const s = tint_modf(1.0f).whole;
+}
+
diff --git a/test/bug/chromium/1236161.wgsl.expected.spvasm b/test/bug/chromium/1236161.wgsl.expected.spvasm
new file mode 100644
index 0000000..39a6a55
--- /dev/null
+++ b/test/bug/chromium/1236161.wgsl.expected.spvasm
@@ -0,0 +1,32 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 13
+; Schema: 0
+ OpCapability Shader
+ %10 = OpExtInstImport "GLSL.std.450"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+ OpExecutionMode %unused_entry_point LocalSize 1 1 1
+ OpName %unused_entry_point "unused_entry_point"
+ OpName %tint_symbol "tint_symbol"
+ OpName %_modf_result "_modf_result"
+ OpMemberName %_modf_result 0 "fract"
+ OpMemberName %_modf_result 1 "whole"
+ OpMemberDecorate %_modf_result 0 Offset 0
+ OpMemberDecorate %_modf_result 1 Offset 4
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %float = OpTypeFloat 32
+%_modf_result = OpTypeStruct %float %float
+ %float_1 = OpConstant %float 1
+%unused_entry_point = OpFunction %void None %1
+ %4 = OpLabel
+ OpReturn
+ OpFunctionEnd
+%tint_symbol = OpFunction %void None %1
+ %6 = OpLabel
+ %7 = OpExtInst %_modf_result %10 ModfStruct %float_1
+ %12 = OpCompositeExtract %float %7 1
+ OpReturn
+ OpFunctionEnd
diff --git a/test/bug/chromium/1236161.wgsl.expected.wgsl b/test/bug/chromium/1236161.wgsl.expected.wgsl
new file mode 100644
index 0000000..ec28ed3
--- /dev/null
+++ b/test/bug/chromium/1236161.wgsl.expected.wgsl
@@ -0,0 +1,3 @@
+fn tint_symbol() {
+ let tint_symbol_1 = modf(1.0).whole;
+}
