diff --git a/src/transform/renamer.cc b/src/transform/renamer.cc index 712f70ed84..956d4076ce 100644 --- a/src/transform/renamer.cc +++ b/src/transform/renamer.cc @@ -896,7 +896,8 @@ Output Renamer::Run(const Program* in, const DataMap& inputs) { // 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 preserve; for (auto* node : in->ASTNodes().Objects()) { if (auto* member = node->As()) { @@ -908,6 +909,12 @@ Output Renamer::Run(const Program* in, const DataMap& inputs) { } if (sem->Is()) { preserve.emplace(member->member()); + } else if (auto* str_expr = in->Sem().Get(member->structure())) { + if (auto* ty = str_expr->Type()->UnwrapRef()->As()) { + if (ty->Declaration() == nullptr) { // Builtin structure + preserve.emplace(member->member()); + } + } } } else if (auto* call = node->As()) { auto* sem = in->Sem().Get(call); diff --git a/src/transform/renamer_test.cc b/src/transform/renamer_test.cc index cfd77d519d..9b9ff3f274 100644 --- a/src/transform/renamer_test.cc +++ b/src/transform/renamer_test.cc @@ -149,6 +149,41 @@ fn tint_symbol() -> [[builtin(position)]] vec4 { 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(src); + + EXPECT_EQ(expect, str(got)); + + auto* data = got.data.Get(); + + 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 db40ef3281..c36db0d95a 100644 --- a/src/writer/hlsl/generator_impl.cc +++ b/src/writer/hlsl/generator_impl.cc @@ -139,8 +139,20 @@ bool GeneratorImpl::Generate() { return false; } } else if (auto* str = decl->As()) { - 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()) { if (func->IsEntryPoint()) { @@ -3135,19 +3147,6 @@ bool GeneratorImpl::EmitTypeAndName(std::ostream& out, } 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 b9a8fc6d5e..90df7cd86e 100644 --- a/src/writer/hlsl/generator_impl_type_test.cc +++ b/src/writer/hlsl/generator_impl_type_test.cc @@ -198,10 +198,8 @@ TEST_F(HlslGeneratorImplTest_Type, EmitType_StructDecl_OmittedIfStorageBuffer) { GeneratorImpl& gen = Build(); - TextGenerator::TextBuffer buf; - auto* sem_s = program->TypeOf(s)->As(); - 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 0000000000..1a4f3c9962 --- /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 0000000000..a0d44d90a7 --- /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 0000000000..4eb36e602b --- /dev/null +++ b/test/bug/chromium/1236161.wgsl.expected.msl @@ -0,0 +1,18 @@ +#include + +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 0000000000..39a6a55c80 --- /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 0000000000..ec28ed3111 --- /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; +}