From 8a96c7893148bd45ab95165de4c026a702adaa51 Mon Sep 17 00:00:00 2001 From: Ben Clayton Date: Thu, 15 Jul 2021 20:29:09 +0000 Subject: [PATCH] transform: Fixes for DecomposeMemoryAccess CloneContext::Replace(T* what, T* with) is bug-prone, as complex transforms may want to clone `what` multiple times, or not at all. In both cases, this will likely result in an ICE as either the replacement will be reachable multiple times, or not at all. The CTS test: webgpu:shader,execution,robust_access:linear_memory:storageClass="storage";storageMode="read_write";access="read";atomic=true;baseType="i32" Was triggering this brokenness with DecomposeMemoryAccess's use of CloneContext::Replace(T*, T*). Switch the usage of CloneContext::Replace(T*, T*) to the new function form. As std::function is copyable, it cannot hold a captured std::unique_ptr. This prevented the Replace() lambdas from capturing the necessary `BufferAccess` data, as this held a `std::unique_ptr`. To fix this, use a `BlockAllocator` for Offsets, and use raw pointers instead. Because the function passed to Replace() is called just before the node is cloned, insertion of new functions will occur just before the currently evaluated module-scope entity. This allows us to remove the "insert_after" arguments to LoadFunc(), StoreFunc(), and AtomicFunc(). We can also kill the icky InsertGlobal() and TypeDeclOf() helpers. Bug: tint:993 Change-Id: I60972bc13a2fa819a163ee2671f61e82d0e68d2a Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/58222 Kokoro: Kokoro Reviewed-by: James Price Commit-Queue: Ben Clayton --- src/transform/decompose_memory_access.cc | 419 ++++++++---------- src/transform/decompose_memory_access_test.cc | 36 +- .../generator_impl_member_accessor_test.cc | 6 +- .../assign_to_function_var.wgsl.expected.hlsl | 38 +- .../assign_to_private_var.wgsl.expected.hlsl | 38 +- .../assign_to_storage_var.wgsl.expected.hlsl | 82 ++-- ...assign_to_workgroup_var.wgsl.expected.hlsl | 38 +- .../dynamic_index/read.wgsl.expected.hlsl | 12 +- .../dynamic_index/write.wgsl.expected.hlsl | 12 +- .../static_index/read.wgsl.expected.hlsl | 4 +- .../static_index/write.wgsl.expected.hlsl | 4 +- .../dynamic_index/read.wgsl.expected.hlsl | 12 +- .../static_index/read.wgsl.expected.hlsl | 8 +- test/bug/tint/403.wgsl.expected.hlsl | 28 +- test/bug/tint/870.spvasm.expected.hlsl | 4 +- test/bug/tint/922.wgsl.expected.hlsl | 49 +- test/bug/tint/993.wgsl | 24 + test/bug/tint/993.wgsl.expected.hlsl | 23 + test/bug/tint/993.wgsl.expected.msl | 34 ++ test/bug/tint/993.wgsl.expected.spvasm | 73 +++ test/bug/tint/993.wgsl.expected.wgsl | 29 ++ test/samples/cube.wgsl.expected.hlsl | 16 +- ...d_struct_storage_buffer.wgsl.expected.hlsl | 12 +- 23 files changed, 573 insertions(+), 428 deletions(-) create mode 100644 test/bug/tint/993.wgsl create mode 100644 test/bug/tint/993.wgsl.expected.hlsl create mode 100644 test/bug/tint/993.wgsl.expected.msl create mode 100644 test/bug/tint/993.wgsl.expected.spvasm create mode 100644 test/bug/tint/993.wgsl.expected.wgsl diff --git a/src/transform/decompose_memory_access.cc b/src/transform/decompose_memory_access.cc index 1210c720cc..a43c4c0820 100644 --- a/src/transform/decompose_memory_access.cc +++ b/src/transform/decompose_memory_access.cc @@ -26,6 +26,7 @@ #include "src/ast/scalar_constructor_expression.h" #include "src/ast/type_name.h" #include "src/ast/unary_op.h" +#include "src/block_allocator.h" #include "src/program_builder.h" #include "src/sem/array.h" #include "src/sem/atomic_type.h" @@ -50,7 +51,7 @@ namespace { /// offsets for storage and uniform buffer accesses. struct Offset : Castable { /// @returns builds and returns the ast::Expression in `ctx.dst` - virtual ast::Expression* Build(CloneContext& ctx) = 0; + virtual ast::Expression* Build(CloneContext& ctx) const = 0; }; /// OffsetExpr is an implementation of Offset that clones and casts the given @@ -60,7 +61,7 @@ struct OffsetExpr : Offset { explicit OffsetExpr(ast::Expression* e) : expr(e) {} - ast::Expression* Build(CloneContext& ctx) override { + ast::Expression* Build(CloneContext& ctx) const override { auto* type = ctx.src->Sem().Get(expr)->Type()->UnwrapRef(); auto* res = ctx.Clone(expr); if (!type->Is()) { @@ -77,7 +78,7 @@ struct OffsetLiteral : Castable { explicit OffsetLiteral(uint32_t lit) : literal(lit) {} - ast::Expression* Build(CloneContext& ctx) override { + ast::Expression* Build(CloneContext& ctx) const override { return ctx.dst->Expr(literal); } }; @@ -86,103 +87,20 @@ struct OffsetLiteral : Castable { /// two Offsets. struct OffsetBinOp : Offset { ast::BinaryOp op; - std::unique_ptr lhs; - std::unique_ptr rhs; + Offset const* lhs = nullptr; + Offset const* rhs = nullptr; - ast::Expression* Build(CloneContext& ctx) override { + ast::Expression* Build(CloneContext& ctx) const override { return ctx.dst->create(op, lhs->Build(ctx), rhs->Build(ctx)); } }; -/// @returns an Offset for the given literal value -std::unique_ptr ToOffset(uint32_t offset) { - return std::make_unique(offset); -} - -/// @returns an Offset for the given ast::Expression -std::unique_ptr ToOffset(ast::Expression* expr) { - if (auto* scalar = expr->As()) { - if (auto* u32 = scalar->literal()->As()) { - return std::make_unique(u32->value()); - } else if (auto* i32 = scalar->literal()->As()) { - if (i32->value() > 0) { - return std::make_unique(i32->value()); - } - } - } - return std::make_unique(expr); -} - -/// @returns the given offset (pass-through) -std::unique_ptr ToOffset(std::unique_ptr offset) { - return offset; -} - -/// @return an Offset that is a sum of lhs and rhs, performing basic constant -/// folding if possible -template -std::unique_ptr Add(LHS&& lhs_, RHS&& rhs_) { - std::unique_ptr lhs = ToOffset(std::forward(lhs_)); - std::unique_ptr rhs = ToOffset(std::forward(rhs_)); - auto* lhs_lit = lhs->As(); - auto* rhs_lit = rhs->As(); - if (lhs_lit && lhs_lit->literal == 0) { - return rhs; - } - if (rhs_lit && rhs_lit->literal == 0) { - return lhs; - } - if (lhs_lit && rhs_lit) { - if (static_cast(lhs_lit->literal) + - static_cast(rhs_lit->literal) <= - 0xffffffff) { - return std::make_unique(lhs_lit->literal + - rhs_lit->literal); - } - } - auto out = std::make_unique(); - out->op = ast::BinaryOp::kAdd; - out->lhs = std::move(lhs); - out->rhs = std::move(rhs); - return out; -} - -/// @return an Offset that is the multiplication of lhs and rhs, performing -/// basic constant folding if possible -template -std::unique_ptr Mul(LHS&& lhs_, RHS&& rhs_) { - std::unique_ptr lhs = ToOffset(std::forward(lhs_)); - std::unique_ptr rhs = ToOffset(std::forward(rhs_)); - auto* lhs_lit = lhs->As(); - auto* rhs_lit = rhs->As(); - if (lhs_lit && lhs_lit->literal == 0) { - return std::make_unique(0); - } - if (rhs_lit && rhs_lit->literal == 0) { - return std::make_unique(0); - } - if (lhs_lit && lhs_lit->literal == 1) { - return rhs; - } - if (rhs_lit && rhs_lit->literal == 1) { - return lhs; - } - if (lhs_lit && rhs_lit) { - return std::make_unique(lhs_lit->literal * rhs_lit->literal); - } - auto out = std::make_unique(); - out->op = ast::BinaryOp::kMultiply; - out->lhs = std::move(lhs); - out->rhs = std::move(rhs); - return out; -} - /// LoadStoreKey is the unordered map key to a load or store intrinsic. struct LoadStoreKey { ast::StorageClass const storage_class; // buffer storage class - sem::Type const* buf_ty; // buffer type - sem::Type const* el_ty; // element type + sem::Type const* buf_ty = nullptr; // buffer type + sem::Type const* el_ty = nullptr; // element type bool operator==(const LoadStoreKey& rhs) const { return storage_class == rhs.storage_class && buf_ty == rhs.buf_ty && el_ty == rhs.el_ty; @@ -196,9 +114,9 @@ struct LoadStoreKey { /// AtomicKey is the unordered map key to an atomic intrinsic. struct AtomicKey { - sem::Type const* buf_ty; // buffer type - sem::Type const* el_ty; // element type - sem::IntrinsicType const op; // atomic op + sem::Type const* buf_ty = nullptr; // buffer type + sem::Type const* el_ty = nullptr; // element type + sem::IntrinsicType const op; // atomic op bool operator==(const AtomicKey& rhs) const { return buf_ty == rhs.buf_ty && el_ty == rhs.el_ty && op == rhs.op; } @@ -367,39 +285,10 @@ DecomposeMemoryAccess::Intrinsic* IntrinsicAtomicFor(ProgramBuilder* builder, builder->ID(), op, ast::StorageClass::kStorage, type); } -/// Inserts `node` before `insert_after` in the global declarations of -/// `ctx.dst`. If `insert_after` is nullptr, then `node` is inserted at the top -/// of the module. -void InsertGlobal(CloneContext& ctx, - const Cloneable* insert_after, - Cloneable* node) { - auto& globals = ctx.src->AST().GlobalDeclarations(); - if (insert_after) { - ctx.InsertAfter(globals, insert_after, node); - } else { - ctx.InsertBefore(globals, *globals.begin(), node); - } -} - -/// @returns the unwrapped, user-declared type of ty. -const ast::TypeDecl* TypeDeclOf(const sem::Type* ty) { - while (true) { - if (auto* ref = ty->As()) { - ty = ref->StoreType(); - continue; - } - if (auto* str = ty->As()) { - return str->Declaration(); - } - // Not a declared type - return nullptr; - } -} - /// BufferAccess describes a single storage or uniform buffer access struct BufferAccess { sem::Expression const* var = nullptr; // Storage buffer variable - std::unique_ptr offset; // The byte offset on var + Offset const* offset = nullptr; // The byte offset on var sem::Type const* type = nullptr; // The type of the access operator bool() const { return var; } // Returns true if valid }; @@ -430,14 +319,105 @@ struct DecomposeMemoryAccess::State { std::unordered_map atomic_funcs; /// List of storage or uniform buffer writes std::vector stores; + /// Allocations for offsets + BlockAllocator offsets_; + + /// @param offset the offset value to wrap in an Offset + /// @returns an Offset for the given literal value + const Offset* ToOffset(uint32_t offset) { + return offsets_.Create(offset); + } + + /// @param expr the expression to convert to an Offset + /// @returns an Offset for the given ast::Expression + const Offset* ToOffset(ast::Expression* expr) { + if (auto* scalar = expr->As()) { + if (auto* u32 = scalar->literal()->As()) { + return offsets_.Create(u32->value()); + } else if (auto* i32 = scalar->literal()->As()) { + if (i32->value() > 0) { + return offsets_.Create(i32->value()); + } + } + } + return offsets_.Create(expr); + } + + /// @param offset the Offset that is returned + /// @returns the given offset (pass-through) + const Offset* ToOffset(const Offset* offset) { return offset; } + + /// @param lhs_ the left-hand side of the add expression + /// @param rhs_ the right-hand side of the add expression + /// @return an Offset that is a sum of lhs and rhs, performing basic constant + /// folding if possible + template + const Offset* Add(LHS&& lhs_, RHS&& rhs_) { + auto* lhs = ToOffset(std::forward(lhs_)); + auto* rhs = ToOffset(std::forward(rhs_)); + auto* lhs_lit = tint::As(lhs); + auto* rhs_lit = tint::As(rhs); + if (lhs_lit && lhs_lit->literal == 0) { + return rhs; + } + if (rhs_lit && rhs_lit->literal == 0) { + return lhs; + } + if (lhs_lit && rhs_lit) { + if (static_cast(lhs_lit->literal) + + static_cast(rhs_lit->literal) <= + 0xffffffff) { + return offsets_.Create(lhs_lit->literal + + rhs_lit->literal); + } + } + auto* out = offsets_.Create(); + out->op = ast::BinaryOp::kAdd; + out->lhs = lhs; + out->rhs = rhs; + return out; + } + + /// @param lhs_ the left-hand side of the multiply expression + /// @param rhs_ the right-hand side of the multiply expression + /// @return an Offset that is the multiplication of lhs and rhs, performing + /// basic constant folding if possible + template + const Offset* Mul(LHS&& lhs_, RHS&& rhs_) { + auto* lhs = ToOffset(std::forward(lhs_)); + auto* rhs = ToOffset(std::forward(rhs_)); + auto* lhs_lit = tint::As(lhs); + auto* rhs_lit = tint::As(rhs); + if (lhs_lit && lhs_lit->literal == 0) { + return offsets_.Create(0); + } + if (rhs_lit && rhs_lit->literal == 0) { + return offsets_.Create(0); + } + if (lhs_lit && lhs_lit->literal == 1) { + return rhs; + } + if (rhs_lit && rhs_lit->literal == 1) { + return lhs; + } + if (lhs_lit && rhs_lit) { + return offsets_.Create(lhs_lit->literal * + rhs_lit->literal); + } + auto* out = offsets_.Create(); + out->op = ast::BinaryOp::kMultiply; + out->lhs = lhs; + out->rhs = rhs; + return out; + } /// AddAccess() adds the `expr -> access` map item to #accesses, and `expr` /// to #expression_order. /// @param expr the expression that performs the access /// @param access the access - void AddAccess(ast::Expression* expr, BufferAccess&& access) { + void AddAccess(ast::Expression* expr, const BufferAccess& access) { TINT_ASSERT(Transform, access.type); - accesses.emplace(expr, std::move(access)); + accesses.emplace(expr, access); expression_order.emplace_back(expr); } @@ -451,7 +431,7 @@ struct DecomposeMemoryAccess::State { if (lhs_it == accesses.end()) { return {}; } - auto access = std::move(lhs_it->second); + auto access = lhs_it->second; accesses.erase(node); return access; } @@ -461,13 +441,11 @@ struct DecomposeMemoryAccess::State { /// The emitted function has the signature: /// `fn load(buf : buf_ty, offset : u32) -> el_ty` /// @param ctx the CloneContext - /// @param insert_after the user-declared type to insert the function after /// @param buf_ty the storage or uniform buffer type /// @param el_ty the storage or uniform buffer element type /// @param var_user the variable user /// @return the name of the function that performs the load Symbol LoadFunc(CloneContext& ctx, - const ast::TypeDecl* insert_after, const sem::Type* buf_ty, const sem::Type* el_ty, const sem::VariableUser* var_user) { @@ -509,8 +487,7 @@ struct DecomposeMemoryAccess::State { ast::ExpressionList values; if (auto* mat_ty = el_ty->As()) { auto* vec_ty = mat_ty->ColumnType(); - Symbol load = - LoadFunc(ctx, insert_after, buf_ty, vec_ty, var_user); + Symbol load = LoadFunc(ctx, buf_ty, vec_ty, var_user); for (uint32_t i = 0; i < mat_ty->columns(); i++) { auto* offset = ctx.dst->Add("offset", i * MatrixColumnStride(mat_ty)); @@ -519,14 +496,14 @@ struct DecomposeMemoryAccess::State { } else if (auto* str = el_ty->As()) { for (auto* member : str->Members()) { auto* offset = ctx.dst->Add("offset", member->Offset()); - Symbol load = LoadFunc(ctx, insert_after, buf_ty, - member->Type()->UnwrapRef(), var_user); + Symbol load = LoadFunc(ctx, buf_ty, member->Type()->UnwrapRef(), + var_user); values.emplace_back(ctx.dst->Call(load, "buffer", offset)); } } else if (auto* arr = el_ty->As()) { for (uint32_t i = 0; i < arr->Count(); i++) { auto* offset = ctx.dst->Add("offset", arr->Stride() * i); - Symbol load = LoadFunc(ctx, insert_after, buf_ty, + Symbol load = LoadFunc(ctx, buf_ty, arr->ElemType()->UnwrapRef(), var_user); values.emplace_back(ctx.dst->Call(load, "buffer", offset)); } @@ -539,7 +516,7 @@ struct DecomposeMemoryAccess::State { CreateASTTypeFor(&ctx, el_ty), values))), ast::DecorationList{}, ast::DecorationList{}); } - InsertGlobal(ctx, insert_after, func); + ctx.dst->AST().AddFunction(func); return func->symbol(); }); } @@ -549,13 +526,11 @@ struct DecomposeMemoryAccess::State { /// The function has the signature: /// `fn store(buf : buf_ty, offset : u32, value : el_ty)` /// @param ctx the CloneContext - /// @param insert_after the user-declared type to insert the function after /// @param buf_ty the storage buffer type /// @param el_ty the storage buffer element type /// @param var_user the variable user /// @return the name of the function that performs the store Symbol StoreFunc(CloneContext& ctx, - const ast::TypeDecl* insert_after, const sem::Type* buf_ty, const sem::Type* el_ty, const sem::VariableUser* var_user) { @@ -597,8 +572,7 @@ struct DecomposeMemoryAccess::State { ast::StatementList body; if (auto* mat_ty = el_ty->As()) { auto* vec_ty = mat_ty->ColumnType(); - Symbol store = - StoreFunc(ctx, insert_after, buf_ty, vec_ty, var_user); + Symbol store = StoreFunc(ctx, buf_ty, vec_ty, var_user); for (uint32_t i = 0; i < mat_ty->columns(); i++) { auto* offset = ctx.dst->Add("offset", i * MatrixColumnStride(mat_ty)); @@ -611,7 +585,7 @@ struct DecomposeMemoryAccess::State { auto* offset = ctx.dst->Add("offset", member->Offset()); auto* access = ctx.dst->MemberAccessor( "value", ctx.Clone(member->Declaration()->symbol())); - Symbol store = StoreFunc(ctx, insert_after, buf_ty, + Symbol store = StoreFunc(ctx, buf_ty, member->Type()->UnwrapRef(), var_user); auto* call = ctx.dst->Call(store, "buffer", offset, access); body.emplace_back(ctx.dst->create(call)); @@ -621,9 +595,8 @@ struct DecomposeMemoryAccess::State { auto* offset = ctx.dst->Add("offset", arr->Stride() * i); auto* access = ctx.dst->IndexAccessor("value", ctx.dst->Expr(i)); - Symbol store = - StoreFunc(ctx, insert_after, buf_ty, - arr->ElemType()->UnwrapRef(), var_user); + Symbol store = StoreFunc( + ctx, buf_ty, arr->ElemType()->UnwrapRef(), var_user); auto* call = ctx.dst->Call(store, "buffer", offset, access); body.emplace_back(ctx.dst->create(call)); } @@ -634,7 +607,7 @@ struct DecomposeMemoryAccess::State { ast::DecorationList{}); } - InsertGlobal(ctx, insert_after, func); + ctx.dst->AST().AddFunction(func); return func->symbol(); }); } @@ -644,14 +617,12 @@ struct DecomposeMemoryAccess::State { /// the signature: // `fn atomic_op(buf : buf_ty, offset : u32, ...) -> T` /// @param ctx the CloneContext - /// @param insert_after the user-declared type to insert the function after /// @param buf_ty the storage buffer type /// @param el_ty the storage buffer element type /// @param intrinsic the atomic intrinsic /// @param var_user the variable user /// @return the name of the function that performs the load Symbol AtomicFunc(CloneContext& ctx, - const ast::TypeDecl* insert_after, const sem::Type* buf_ty, const sem::Type* el_ty, const sem::Intrinsic* intrinsic, @@ -700,7 +671,7 @@ struct DecomposeMemoryAccess::State { }, ast::DecorationList{}); - InsertGlobal(ctx, insert_after, func); + ctx.dst->AST().AddFunction(func); return func->symbol(); }); } @@ -825,7 +796,7 @@ void DecomposeMemoryAccess::Run(CloneContext& ctx, const DataMap&, DataMap&) { // Variable to a storage or uniform buffer state.AddAccess(ident, { var, - ToOffset(0u), + state.ToOffset(0u), var->Type()->UnwrapRef(), }); } @@ -840,14 +811,13 @@ void DecomposeMemoryAccess::Run(CloneContext& ctx, const DataMap&, DataMap&) { if (swizzle->Indices().size() == 1) { if (auto access = state.TakeAccess(accessor->structure())) { auto* vec_ty = access.type->As(); - auto offset = - Mul(ScalarSize(vec_ty->type()), swizzle->Indices()[0]); - state.AddAccess( - accessor, { - access.var, - Add(std::move(access.offset), std::move(offset)), - vec_ty->type()->UnwrapRef(), - }); + auto* offset = + state.Mul(ScalarSize(vec_ty->type()), swizzle->Indices()[0]); + state.AddAccess(accessor, { + access.var, + state.Add(access.offset, offset), + vec_ty->type()->UnwrapRef(), + }); } } } else { @@ -855,12 +825,11 @@ void DecomposeMemoryAccess::Run(CloneContext& ctx, const DataMap&, DataMap&) { auto* str_ty = access.type->As(); auto* member = str_ty->FindMember(accessor->member()->symbol()); auto offset = member->Offset(); - state.AddAccess(accessor, - { - access.var, - Add(std::move(access.offset), std::move(offset)), - member->Type()->UnwrapRef(), - }); + state.AddAccess(accessor, { + access.var, + state.Add(access.offset, offset), + member->Type()->UnwrapRef(), + }); } } continue; @@ -870,33 +839,32 @@ void DecomposeMemoryAccess::Run(CloneContext& ctx, const DataMap&, DataMap&) { if (auto access = state.TakeAccess(accessor->array())) { // X[Y] if (auto* arr = access.type->As()) { - auto offset = Mul(arr->Stride(), accessor->idx_expr()); - state.AddAccess(accessor, - { - access.var, - Add(std::move(access.offset), std::move(offset)), - arr->ElemType()->UnwrapRef(), - }); + auto* offset = state.Mul(arr->Stride(), accessor->idx_expr()); + state.AddAccess(accessor, { + access.var, + state.Add(access.offset, offset), + arr->ElemType()->UnwrapRef(), + }); continue; } if (auto* vec_ty = access.type->As()) { - auto offset = Mul(ScalarSize(vec_ty->type()), accessor->idx_expr()); - state.AddAccess(accessor, - { - access.var, - Add(std::move(access.offset), std::move(offset)), - vec_ty->type()->UnwrapRef(), - }); + auto* offset = + state.Mul(ScalarSize(vec_ty->type()), accessor->idx_expr()); + state.AddAccess(accessor, { + access.var, + state.Add(access.offset, offset), + vec_ty->type()->UnwrapRef(), + }); continue; } if (auto* mat_ty = access.type->As()) { - auto offset = Mul(MatrixColumnStride(mat_ty), accessor->idx_expr()); - state.AddAccess(accessor, - { - access.var, - Add(std::move(access.offset), std::move(offset)), - mat_ty->ColumnType(), - }); + auto* offset = + state.Mul(MatrixColumnStride(mat_ty), accessor->idx_expr()); + state.AddAccess(accessor, { + access.var, + state.Add(access.offset, offset), + mat_ty->ColumnType(), + }); continue; } } @@ -908,7 +876,7 @@ void DecomposeMemoryAccess::Run(CloneContext& ctx, const DataMap&, DataMap&) { if (auto access = state.TakeAccess(op->expr())) { // HLSL does not support pointers, so just take the access from the // reference and place it on the pointer. - state.AddAccess(op, std::move(access)); + state.AddAccess(op, access); continue; } } @@ -918,7 +886,7 @@ void DecomposeMemoryAccess::Run(CloneContext& ctx, const DataMap&, DataMap&) { // X = Y // Move the LHS access to a store. if (auto lhs = state.TakeAccess(assign->lhs())) { - state.stores.emplace_back(Store{assign, std::move(lhs)}); + state.stores.emplace_back(Store{assign, lhs}); } } @@ -934,23 +902,22 @@ void DecomposeMemoryAccess::Run(CloneContext& ctx, const DataMap&, DataMap&) { if (intrinsic->IsAtomic()) { if (auto access = state.TakeAccess(call_expr->params()[0])) { // atomic___(X) + ctx.Replace(call_expr, [=, &ctx, &state] { + auto* buf = access.var->Declaration(); + auto* offset = access.offset->Build(ctx); + auto* buf_ty = access.var->Type()->UnwrapRef(); + auto* el_ty = access.type->UnwrapRef()->As()->Type(); + Symbol func = + state.AtomicFunc(ctx, buf_ty, el_ty, intrinsic, + access.var->As()); - auto* buf = access.var->Declaration(); - auto* offset = access.offset->Build(ctx); - auto* buf_ty = access.var->Type()->UnwrapRef(); - auto* el_ty = access.type->UnwrapRef()->As()->Type(); - auto* insert_after = TypeDeclOf(access.var->Type()); - Symbol func = - state.AtomicFunc(ctx, insert_after, buf_ty, el_ty, intrinsic, - access.var->As()); - - ast::ExpressionList args{ctx.Clone(buf), offset}; - for (size_t i = 1; i < call_expr->params().size(); i++) { - auto* arg = call_expr->params()[i]; - args.emplace_back(ctx.Clone(arg)); - } - - ctx.Replace(call_expr, ctx.dst->Call(func, args)); + ast::ExpressionList args{ctx.Clone(buf), offset}; + for (size_t i = 1; i < call_expr->params().size(); i++) { + auto* arg = call_expr->params()[i]; + args.emplace_back(ctx.Clone(arg)); + } + return ctx.dst->Call(func, args); + }); } } } @@ -964,36 +931,32 @@ void DecomposeMemoryAccess::Run(CloneContext& ctx, const DataMap&, DataMap&) { if (access_it == state.accesses.end()) { continue; } - - auto access = std::move(access_it->second); - - auto* buf = access.var->Declaration(); - auto* offset = access.offset->Build(ctx); - auto* buf_ty = access.var->Type()->UnwrapRef(); - auto* el_ty = access.type->UnwrapRef(); - auto* insert_after = TypeDeclOf(access.var->Type()); - Symbol func = state.LoadFunc(ctx, insert_after, buf_ty, el_ty, - access.var->As()); - - auto* load = ctx.dst->Call(func, ctx.Clone(buf), offset); - - ctx.Replace(expr, load); + BufferAccess access = access_it->second; + ctx.Replace(expr, [=, &ctx, &state] { + auto* buf = access.var->Declaration(); + auto* offset = access.offset->Build(ctx); + auto* buf_ty = access.var->Type()->UnwrapRef(); + auto* el_ty = access.type->UnwrapRef(); + Symbol func = state.LoadFunc(ctx, buf_ty, el_ty, + access.var->As()); + return ctx.dst->Call(func, ctx.CloneWithoutTransform(buf), offset); + }); } // And replace all storage and uniform buffer assignments with stores - for (auto& store : state.stores) { - auto* buf = store.target.var->Declaration(); - auto* offset = store.target.offset->Build(ctx); - auto* buf_ty = store.target.var->Type()->UnwrapRef(); - auto* el_ty = store.target.type->UnwrapRef(); - auto* value = store.assignment->rhs(); - auto* insert_after = TypeDeclOf(store.target.var->Type()); - Symbol func = state.StoreFunc(ctx, insert_after, buf_ty, el_ty, - store.target.var->As()); - - auto* call = ctx.dst->Call(func, ctx.Clone(buf), offset, ctx.Clone(value)); - - ctx.Replace(store.assignment, ctx.dst->create(call)); + for (auto store : state.stores) { + ctx.Replace(store.assignment, [=, &ctx, &state] { + auto* buf = store.target.var->Declaration(); + auto* offset = store.target.offset->Build(ctx); + auto* buf_ty = store.target.var->Type()->UnwrapRef(); + auto* el_ty = store.target.type->UnwrapRef(); + auto* value = store.assignment->rhs(); + Symbol func = state.StoreFunc(ctx, buf_ty, el_ty, + store.target.var->As()); + auto* call = ctx.dst->Call(func, ctx.CloneWithoutTransform(buf), offset, + ctx.Clone(value)); + return ctx.dst->create(call); + }); } ctx.Clone(); diff --git a/src/transform/decompose_memory_access_test.cc b/src/transform/decompose_memory_access_test.cc index 06069625ed..5e9e60da2e 100644 --- a/src/transform/decompose_memory_access_test.cc +++ b/src/transform/decompose_memory_access_test.cc @@ -106,6 +106,8 @@ struct SB { v : array, 2>; }; +[[group(0), binding(0)]] var sb : SB; + [[internal(intrinsic_load_storage_i32), internal(disable_validation__function_has_no_body)]] fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> i32 @@ -182,8 +184,6 @@ fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_p return array, 2>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u))); } -[[group(0), binding(0)]] var sb : SB; - [[stage(compute), workgroup_size(1)]] fn main() { var a : i32 = tint_symbol(sb, 0u); @@ -300,6 +300,8 @@ struct UB { v : array, 2>; }; +[[group(0), binding(0)]] var ub : UB; + [[internal(intrinsic_load_uniform_i32), internal(disable_validation__function_has_no_body)]] fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> i32 @@ -376,8 +378,6 @@ fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_p return array, 2>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u))); } -[[group(0), binding(0)]] var ub : UB; - [[stage(compute), workgroup_size(1)]] fn main() { var a : i32 = tint_symbol(ub, 0u); @@ -494,6 +494,8 @@ struct SB { v : array, 2>; }; +[[group(0), binding(0)]] var sb : SB; + [[internal(intrinsic_store_storage_i32), internal(disable_validation__function_has_no_body)]] fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : i32) @@ -589,8 +591,6 @@ fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_p tint_symbol_8(buffer, (offset + 16u), value[1u]); } -[[group(0), binding(0)]] var sb : SB; - [[stage(compute), workgroup_size(1)]] fn main() { tint_symbol(sb, 0u, i32()); @@ -686,6 +686,8 @@ struct SB { v : array, 2>; }; +[[group(0), binding(0)]] var sb : SB; + [[internal(intrinsic_load_storage_i32), internal(disable_validation__function_has_no_body)]] fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> i32 @@ -766,8 +768,6 @@ fn tint_symbol_22([[internal(disable_validation__ignore_constructible_function_p return SB(tint_symbol(buffer, (offset + 0u)), tint_symbol_1(buffer, (offset + 4u)), tint_symbol_2(buffer, (offset + 8u)), tint_symbol_3(buffer, (offset + 16u)), tint_symbol_4(buffer, (offset + 24u)), tint_symbol_5(buffer, (offset + 32u)), tint_symbol_6(buffer, (offset + 48u)), tint_symbol_7(buffer, (offset + 64u)), tint_symbol_8(buffer, (offset + 80u)), tint_symbol_9(buffer, (offset + 96u)), tint_symbol_10(buffer, (offset + 112u)), tint_symbol_11(buffer, (offset + 128u)), tint_symbol_12(buffer, (offset + 144u)), tint_symbol_13(buffer, (offset + 160u)), tint_symbol_14(buffer, (offset + 192u)), tint_symbol_15(buffer, (offset + 224u)), tint_symbol_16(buffer, (offset + 256u)), tint_symbol_17(buffer, (offset + 304u)), tint_symbol_18(buffer, (offset + 352u)), tint_symbol_19(buffer, (offset + 384u)), tint_symbol_20(buffer, (offset + 448u)), tint_symbol_21(buffer, (offset + 512u))); } -[[group(0), binding(0)]] var sb : SB; - [[stage(compute), workgroup_size(1)]] fn main() { var x : SB = tint_symbol_22(sb, 0u); @@ -842,6 +842,8 @@ struct SB { v : array, 2>; }; +[[group(0), binding(0)]] var sb : SB; + [[internal(intrinsic_store_storage_i32), internal(disable_validation__function_has_no_body)]] fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : i32) @@ -962,8 +964,6 @@ fn tint_symbol_22([[internal(disable_validation__ignore_constructible_function_p tint_symbol_21(buffer, (offset + 512u), value.v); } -[[group(0), binding(0)]] var sb : SB; - [[stage(compute), workgroup_size(1)]] fn main() { tint_symbol_22(sb, 0u, SB()); @@ -1031,11 +1031,11 @@ struct SB { b : [[stride(256)]] array; }; +[[group(0), binding(0)]] var sb : SB; + [[internal(intrinsic_load_storage_f32), internal(disable_validation__function_has_no_body)]] fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> f32 -[[group(0), binding(0)]] var sb : SB; - [[stage(compute), workgroup_size(1)]] fn main() { var x : f32 = tint_symbol(sb, 1224u); @@ -1099,11 +1099,11 @@ struct SB { b : [[stride(256)]] array; }; +[[group(0), binding(0)]] var sb : SB; + [[internal(intrinsic_load_storage_f32), internal(disable_validation__function_has_no_body)]] fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> f32 -[[group(0), binding(0)]] var sb : SB; - [[stage(compute), workgroup_size(1)]] fn main() { var i : i32 = 4; @@ -1186,11 +1186,11 @@ struct SB { b : A2_Array; }; +[[group(0), binding(0)]] var sb : SB; + [[internal(intrinsic_load_storage_f32), internal(disable_validation__function_has_no_body)]] fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> f32 -[[group(0), binding(0)]] var sb : SB; - [[stage(compute), workgroup_size(1)]] fn main() { var i : i32 = 4; @@ -1250,6 +1250,8 @@ struct SB { b : atomic; }; +[[group(0), binding(0)]] var sb : SB; + [[internal(intrinsic_atomic_store_storage_i32), internal(disable_validation__function_has_no_body)]] fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32) @@ -1310,8 +1312,6 @@ fn tint_symbol_18([[internal(disable_validation__ignore_constructible_function_p [[internal(intrinsic_atomic_compare_exchange_weak_storage_u32), internal(disable_validation__function_has_no_body)]] fn tint_symbol_19([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32, param_2 : u32) -> vec2 -[[group(0), binding(0)]] var sb : SB; - [[stage(compute), workgroup_size(1)]] fn main() { tint_symbol(sb, 16u, 123); diff --git a/src/writer/hlsl/generator_impl_member_accessor_test.cc b/src/writer/hlsl/generator_impl_member_accessor_test.cc index 1205e93317..10d7cb91d0 100644 --- a/src/writer/hlsl/generator_impl_member_accessor_test.cc +++ b/src/writer/hlsl/generator_impl_member_accessor_test.cc @@ -341,13 +341,13 @@ TEST_F(HlslGeneratorImplTest_MemberAccessor, StorageBuffer_Store_Matrix_Empty) { ASSERT_TRUE(gen.Generate()) << gen.error(); auto* expected = - R"(void tint_symbol_1(RWByteAddressBuffer buffer, uint offset, float2x3 value) { + R"(RWByteAddressBuffer data : register(u0, space1); + +void tint_symbol_1(RWByteAddressBuffer buffer, uint offset, float2x3 value) { buffer.Store3((offset + 0u), asuint(value[0u])); buffer.Store3((offset + 16u), asuint(value[1u])); } -RWByteAddressBuffer data : register(u0, space1); - void main() { tint_symbol_1(data, 16u, float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)); return; diff --git a/test/array/assign_to_function_var.wgsl.expected.hlsl b/test/array/assign_to_function_var.wgsl.expected.hlsl index 7fc6d52eff..2c7d6be159 100644 --- a/test/array/assign_to_function_var.wgsl.expected.hlsl +++ b/test/array/assign_to_function_var.wgsl.expected.hlsl @@ -10,22 +10,6 @@ struct S { tint_padded_array_element arr[4]; }; -typedef tint_padded_array_element tint_symbol_2_ret[4]; -tint_symbol_2_ret tint_symbol_2(uint4 buffer[4], uint offset) { - const uint scalar_offset = ((offset + 0u)) / 4; - const uint scalar_offset_1 = ((offset + 16u)) / 4; - const uint scalar_offset_2 = ((offset + 32u)) / 4; - const uint scalar_offset_3 = ((offset + 48u)) / 4; - const tint_padded_array_element tint_symbol_5[4] = {{asint(buffer[scalar_offset / 4][scalar_offset % 4])}, {asint(buffer[scalar_offset_1 / 4][scalar_offset_1 % 4])}, {asint(buffer[scalar_offset_2 / 4][scalar_offset_2 % 4])}, {asint(buffer[scalar_offset_3 / 4][scalar_offset_3 % 4])}}; - return tint_symbol_5; -} - -typedef tint_padded_array_element tint_symbol_4_ret[4]; -tint_symbol_4_ret tint_symbol_4(RWByteAddressBuffer buffer, uint offset) { - const tint_padded_array_element tint_symbol_6[4] = {{asint(buffer.Load((offset + 0u)))}, {asint(buffer.Load((offset + 16u)))}, {asint(buffer.Load((offset + 32u)))}, {asint(buffer.Load((offset + 48u)))}}; - return tint_symbol_6; -} - static tint_padded_array_element src_private[4] = (tint_padded_array_element[4])0; groupshared tint_padded_array_element src_workgroup[4]; cbuffer cbuffer_src_uniform : register(b0, space0) { @@ -35,12 +19,28 @@ RWByteAddressBuffer src_storage : register(u1, space0); typedef tint_padded_array_element ret_arr_ret[4]; ret_arr_ret ret_arr() { - const tint_padded_array_element tint_symbol_7[4] = (tint_padded_array_element[4])0; - return tint_symbol_7; + const tint_padded_array_element tint_symbol_5[4] = (tint_padded_array_element[4])0; + return tint_symbol_5; } S ret_struct_arr() { - const S tint_symbol_8 = (S)0; + const S tint_symbol_6 = (S)0; + return tint_symbol_6; +} + +typedef tint_padded_array_element tint_symbol_2_ret[4]; +tint_symbol_2_ret tint_symbol_2(uint4 buffer[4], uint offset) { + const uint scalar_offset = ((offset + 0u)) / 4; + const uint scalar_offset_1 = ((offset + 16u)) / 4; + const uint scalar_offset_2 = ((offset + 32u)) / 4; + const uint scalar_offset_3 = ((offset + 48u)) / 4; + const tint_padded_array_element tint_symbol_7[4] = {{asint(buffer[scalar_offset / 4][scalar_offset % 4])}, {asint(buffer[scalar_offset_1 / 4][scalar_offset_1 % 4])}, {asint(buffer[scalar_offset_2 / 4][scalar_offset_2 % 4])}, {asint(buffer[scalar_offset_3 / 4][scalar_offset_3 % 4])}}; + return tint_symbol_7; +} + +typedef tint_padded_array_element tint_symbol_4_ret[4]; +tint_symbol_4_ret tint_symbol_4(RWByteAddressBuffer buffer, uint offset) { + const tint_padded_array_element tint_symbol_8[4] = {{asint(buffer.Load((offset + 0u)))}, {asint(buffer.Load((offset + 16u)))}, {asint(buffer.Load((offset + 32u)))}, {asint(buffer.Load((offset + 48u)))}}; return tint_symbol_8; } diff --git a/test/array/assign_to_private_var.wgsl.expected.hlsl b/test/array/assign_to_private_var.wgsl.expected.hlsl index 55e7cc6d61..68ea5c3f47 100644 --- a/test/array/assign_to_private_var.wgsl.expected.hlsl +++ b/test/array/assign_to_private_var.wgsl.expected.hlsl @@ -10,22 +10,6 @@ struct S { tint_padded_array_element arr[4]; }; -typedef tint_padded_array_element tint_symbol_2_ret[4]; -tint_symbol_2_ret tint_symbol_2(uint4 buffer[4], uint offset) { - const uint scalar_offset = ((offset + 0u)) / 4; - const uint scalar_offset_1 = ((offset + 16u)) / 4; - const uint scalar_offset_2 = ((offset + 32u)) / 4; - const uint scalar_offset_3 = ((offset + 48u)) / 4; - const tint_padded_array_element tint_symbol_5[4] = {{asint(buffer[scalar_offset / 4][scalar_offset % 4])}, {asint(buffer[scalar_offset_1 / 4][scalar_offset_1 % 4])}, {asint(buffer[scalar_offset_2 / 4][scalar_offset_2 % 4])}, {asint(buffer[scalar_offset_3 / 4][scalar_offset_3 % 4])}}; - return tint_symbol_5; -} - -typedef tint_padded_array_element tint_symbol_4_ret[4]; -tint_symbol_4_ret tint_symbol_4(RWByteAddressBuffer buffer, uint offset) { - const tint_padded_array_element tint_symbol_6[4] = {{asint(buffer.Load((offset + 0u)))}, {asint(buffer.Load((offset + 16u)))}, {asint(buffer.Load((offset + 32u)))}, {asint(buffer.Load((offset + 48u)))}}; - return tint_symbol_6; -} - static tint_padded_array_element src_private[4] = (tint_padded_array_element[4])0; groupshared tint_padded_array_element src_workgroup[4]; cbuffer cbuffer_src_uniform : register(b0, space0) { @@ -37,12 +21,28 @@ static int dst_nested[4][3][2] = (int[4][3][2])0; typedef tint_padded_array_element ret_arr_ret[4]; ret_arr_ret ret_arr() { - const tint_padded_array_element tint_symbol_7[4] = (tint_padded_array_element[4])0; - return tint_symbol_7; + const tint_padded_array_element tint_symbol_5[4] = (tint_padded_array_element[4])0; + return tint_symbol_5; } S ret_struct_arr() { - const S tint_symbol_8 = (S)0; + const S tint_symbol_6 = (S)0; + return tint_symbol_6; +} + +typedef tint_padded_array_element tint_symbol_2_ret[4]; +tint_symbol_2_ret tint_symbol_2(uint4 buffer[4], uint offset) { + const uint scalar_offset = ((offset + 0u)) / 4; + const uint scalar_offset_1 = ((offset + 16u)) / 4; + const uint scalar_offset_2 = ((offset + 32u)) / 4; + const uint scalar_offset_3 = ((offset + 48u)) / 4; + const tint_padded_array_element tint_symbol_7[4] = {{asint(buffer[scalar_offset / 4][scalar_offset % 4])}, {asint(buffer[scalar_offset_1 / 4][scalar_offset_1 % 4])}, {asint(buffer[scalar_offset_2 / 4][scalar_offset_2 % 4])}, {asint(buffer[scalar_offset_3 / 4][scalar_offset_3 % 4])}}; + return tint_symbol_7; +} + +typedef tint_padded_array_element tint_symbol_4_ret[4]; +tint_symbol_4_ret tint_symbol_4(RWByteAddressBuffer buffer, uint offset) { + const tint_padded_array_element tint_symbol_8[4] = {{asint(buffer.Load((offset + 0u)))}, {asint(buffer.Load((offset + 16u)))}, {asint(buffer.Load((offset + 32u)))}, {asint(buffer.Load((offset + 48u)))}}; return tint_symbol_8; } diff --git a/test/array/assign_to_storage_var.wgsl.expected.hlsl b/test/array/assign_to_storage_var.wgsl.expected.hlsl index bf9232b6c7..ee58066595 100644 --- a/test/array/assign_to_storage_var.wgsl.expected.hlsl +++ b/test/array/assign_to_storage_var.wgsl.expected.hlsl @@ -10,29 +10,49 @@ struct S { tint_padded_array_element arr[4]; }; -typedef tint_padded_array_element tint_symbol_2_ret[4]; -tint_symbol_2_ret tint_symbol_2(uint4 buffer[4], uint offset) { - const uint scalar_offset = ((offset + 0u)) / 4; - const uint scalar_offset_1 = ((offset + 16u)) / 4; - const uint scalar_offset_2 = ((offset + 32u)) / 4; - const uint scalar_offset_3 = ((offset + 48u)) / 4; - const tint_padded_array_element tint_symbol_11[4] = {{asint(buffer[scalar_offset / 4][scalar_offset % 4])}, {asint(buffer[scalar_offset_1 / 4][scalar_offset_1 % 4])}, {asint(buffer[scalar_offset_2 / 4][scalar_offset_2 % 4])}, {asint(buffer[scalar_offset_3 / 4][scalar_offset_3 % 4])}}; +static tint_padded_array_element src_private[4] = (tint_padded_array_element[4])0; +groupshared tint_padded_array_element src_workgroup[4]; +cbuffer cbuffer_src_uniform : register(b0, space0) { + uint4 src_uniform[4]; +}; +RWByteAddressBuffer src_storage : register(u1, space0); +RWByteAddressBuffer tint_symbol : register(u2, space0); +RWByteAddressBuffer dst_nested : register(u3, space0); + +typedef tint_padded_array_element ret_arr_ret[4]; +ret_arr_ret ret_arr() { + const tint_padded_array_element tint_symbol_11[4] = (tint_padded_array_element[4])0; return tint_symbol_11; } -typedef tint_padded_array_element tint_symbol_4_ret[4]; -tint_symbol_4_ret tint_symbol_4(RWByteAddressBuffer buffer, uint offset) { - const tint_padded_array_element tint_symbol_12[4] = {{asint(buffer.Load((offset + 0u)))}, {asint(buffer.Load((offset + 16u)))}, {asint(buffer.Load((offset + 32u)))}, {asint(buffer.Load((offset + 48u)))}}; +S ret_struct_arr() { + const S tint_symbol_12 = (S)0; return tint_symbol_12; } -void tint_symbol_6(RWByteAddressBuffer buffer, uint offset, tint_padded_array_element value[4]) { +void tint_symbol_2(RWByteAddressBuffer buffer, uint offset, tint_padded_array_element value[4]) { buffer.Store((offset + 0u), asuint(value[0u].el)); buffer.Store((offset + 16u), asuint(value[1u].el)); buffer.Store((offset + 32u), asuint(value[2u].el)); buffer.Store((offset + 48u), asuint(value[3u].el)); } +typedef tint_padded_array_element tint_symbol_4_ret[4]; +tint_symbol_4_ret tint_symbol_4(uint4 buffer[4], uint offset) { + const uint scalar_offset = ((offset + 0u)) / 4; + const uint scalar_offset_1 = ((offset + 16u)) / 4; + const uint scalar_offset_2 = ((offset + 32u)) / 4; + const uint scalar_offset_3 = ((offset + 48u)) / 4; + const tint_padded_array_element tint_symbol_13[4] = {{asint(buffer[scalar_offset / 4][scalar_offset % 4])}, {asint(buffer[scalar_offset_1 / 4][scalar_offset_1 % 4])}, {asint(buffer[scalar_offset_2 / 4][scalar_offset_2 % 4])}, {asint(buffer[scalar_offset_3 / 4][scalar_offset_3 % 4])}}; + return tint_symbol_13; +} + +typedef tint_padded_array_element tint_symbol_6_ret[4]; +tint_symbol_6_ret tint_symbol_6(RWByteAddressBuffer buffer, uint offset) { + const tint_padded_array_element tint_symbol_14[4] = {{asint(buffer.Load((offset + 0u)))}, {asint(buffer.Load((offset + 16u)))}, {asint(buffer.Load((offset + 32u)))}, {asint(buffer.Load((offset + 48u)))}}; + return tint_symbol_14; +} + void tint_symbol_8(RWByteAddressBuffer buffer, uint offset, int value[2]) { buffer.Store((offset + 0u), asuint(value[0u])); buffer.Store((offset + 4u), asuint(value[1u])); @@ -51,40 +71,20 @@ void tint_symbol_10(RWByteAddressBuffer buffer, uint offset, int value[4][3][2]) tint_symbol_9(buffer, (offset + 72u), value[3u]); } -static tint_padded_array_element src_private[4] = (tint_padded_array_element[4])0; -groupshared tint_padded_array_element src_workgroup[4]; -cbuffer cbuffer_src_uniform : register(b0, space0) { - uint4 src_uniform[4]; -}; -RWByteAddressBuffer src_storage : register(u1, space0); -RWByteAddressBuffer tint_symbol : register(u2, space0); -RWByteAddressBuffer dst_nested : register(u3, space0); - -typedef tint_padded_array_element ret_arr_ret[4]; -ret_arr_ret ret_arr() { - const tint_padded_array_element tint_symbol_13[4] = (tint_padded_array_element[4])0; - return tint_symbol_13; -} - -S ret_struct_arr() { - const S tint_symbol_14 = (S)0; - return tint_symbol_14; -} - void foo(tint_padded_array_element src_param[4]) { tint_padded_array_element src_function[4] = (tint_padded_array_element[4])0; const tint_padded_array_element tint_symbol_15[4] = {{1}, {2}, {3}, {3}}; - tint_symbol_6(tint_symbol, 0u, tint_symbol_15); - tint_symbol_6(tint_symbol, 0u, src_param); - tint_symbol_6(tint_symbol, 0u, ret_arr()); + tint_symbol_2(tint_symbol, 0u, tint_symbol_15); + tint_symbol_2(tint_symbol, 0u, src_param); + tint_symbol_2(tint_symbol, 0u, ret_arr()); const tint_padded_array_element src_let[4] = (tint_padded_array_element[4])0; - tint_symbol_6(tint_symbol, 0u, src_let); - tint_symbol_6(tint_symbol, 0u, src_function); - tint_symbol_6(tint_symbol, 0u, src_private); - tint_symbol_6(tint_symbol, 0u, src_workgroup); - tint_symbol_6(tint_symbol, 0u, ret_struct_arr().arr); - tint_symbol_6(tint_symbol, 0u, tint_symbol_2(src_uniform, 0u)); - tint_symbol_6(tint_symbol, 0u, tint_symbol_4(src_storage, 0u)); + tint_symbol_2(tint_symbol, 0u, src_let); + tint_symbol_2(tint_symbol, 0u, src_function); + tint_symbol_2(tint_symbol, 0u, src_private); + tint_symbol_2(tint_symbol, 0u, src_workgroup); + tint_symbol_2(tint_symbol, 0u, ret_struct_arr().arr); + tint_symbol_2(tint_symbol, 0u, tint_symbol_4(src_uniform, 0u)); + tint_symbol_2(tint_symbol, 0u, tint_symbol_6(src_storage, 0u)); int src_nested[4][3][2] = (int[4][3][2])0; tint_symbol_10(dst_nested, 0u, src_nested); } diff --git a/test/array/assign_to_workgroup_var.wgsl.expected.hlsl b/test/array/assign_to_workgroup_var.wgsl.expected.hlsl index 65d3fc916f..48c5fae141 100644 --- a/test/array/assign_to_workgroup_var.wgsl.expected.hlsl +++ b/test/array/assign_to_workgroup_var.wgsl.expected.hlsl @@ -10,22 +10,6 @@ struct S { tint_padded_array_element arr[4]; }; -typedef tint_padded_array_element tint_symbol_2_ret[4]; -tint_symbol_2_ret tint_symbol_2(uint4 buffer[4], uint offset) { - const uint scalar_offset = ((offset + 0u)) / 4; - const uint scalar_offset_1 = ((offset + 16u)) / 4; - const uint scalar_offset_2 = ((offset + 32u)) / 4; - const uint scalar_offset_3 = ((offset + 48u)) / 4; - const tint_padded_array_element tint_symbol_5[4] = {{asint(buffer[scalar_offset / 4][scalar_offset % 4])}, {asint(buffer[scalar_offset_1 / 4][scalar_offset_1 % 4])}, {asint(buffer[scalar_offset_2 / 4][scalar_offset_2 % 4])}, {asint(buffer[scalar_offset_3 / 4][scalar_offset_3 % 4])}}; - return tint_symbol_5; -} - -typedef tint_padded_array_element tint_symbol_4_ret[4]; -tint_symbol_4_ret tint_symbol_4(RWByteAddressBuffer buffer, uint offset) { - const tint_padded_array_element tint_symbol_6[4] = {{asint(buffer.Load((offset + 0u)))}, {asint(buffer.Load((offset + 16u)))}, {asint(buffer.Load((offset + 32u)))}, {asint(buffer.Load((offset + 48u)))}}; - return tint_symbol_6; -} - static tint_padded_array_element src_private[4] = (tint_padded_array_element[4])0; groupshared tint_padded_array_element src_workgroup[4]; cbuffer cbuffer_src_uniform : register(b0, space0) { @@ -37,12 +21,28 @@ groupshared int dst_nested[4][3][2]; typedef tint_padded_array_element ret_arr_ret[4]; ret_arr_ret ret_arr() { - const tint_padded_array_element tint_symbol_7[4] = (tint_padded_array_element[4])0; - return tint_symbol_7; + const tint_padded_array_element tint_symbol_5[4] = (tint_padded_array_element[4])0; + return tint_symbol_5; } S ret_struct_arr() { - const S tint_symbol_8 = (S)0; + const S tint_symbol_6 = (S)0; + return tint_symbol_6; +} + +typedef tint_padded_array_element tint_symbol_2_ret[4]; +tint_symbol_2_ret tint_symbol_2(uint4 buffer[4], uint offset) { + const uint scalar_offset = ((offset + 0u)) / 4; + const uint scalar_offset_1 = ((offset + 16u)) / 4; + const uint scalar_offset_2 = ((offset + 32u)) / 4; + const uint scalar_offset_3 = ((offset + 48u)) / 4; + const tint_padded_array_element tint_symbol_7[4] = {{asint(buffer[scalar_offset / 4][scalar_offset % 4])}, {asint(buffer[scalar_offset_1 / 4][scalar_offset_1 % 4])}, {asint(buffer[scalar_offset_2 / 4][scalar_offset_2 % 4])}, {asint(buffer[scalar_offset_3 / 4][scalar_offset_3 % 4])}}; + return tint_symbol_7; +} + +typedef tint_padded_array_element tint_symbol_4_ret[4]; +tint_symbol_4_ret tint_symbol_4(RWByteAddressBuffer buffer, uint offset) { + const tint_padded_array_element tint_symbol_8[4] = {{asint(buffer.Load((offset + 0u)))}, {asint(buffer.Load((offset + 16u)))}, {asint(buffer.Load((offset + 32u)))}, {asint(buffer.Load((offset + 48u)))}}; return tint_symbol_8; } diff --git a/test/buffer/storage/dynamic_index/read.wgsl.expected.hlsl b/test/buffer/storage/dynamic_index/read.wgsl.expected.hlsl index 89fc6b972e..9b9390af18 100644 --- a/test/buffer/storage/dynamic_index/read.wgsl.expected.hlsl +++ b/test/buffer/storage/dynamic_index/read.wgsl.expected.hlsl @@ -1,3 +1,9 @@ +ByteAddressBuffer s : register(t0, space0); + +struct tint_symbol_1 { + uint idx : SV_GroupIndex; +}; + float2x3 tint_symbol_8(ByteAddressBuffer buffer, uint offset) { return float2x3(asfloat(buffer.Load3((offset + 0u))), asfloat(buffer.Load3((offset + 16u)))); } @@ -12,12 +18,6 @@ tint_symbol_12_ret tint_symbol_12(ByteAddressBuffer buffer, uint offset) { return tint_symbol_13; } -ByteAddressBuffer s : register(t0, space0); - -struct tint_symbol_1 { - uint idx : SV_GroupIndex; -}; - [numthreads(1, 1, 1)] void main(tint_symbol_1 tint_symbol) { const uint idx = tint_symbol.idx; diff --git a/test/buffer/storage/dynamic_index/write.wgsl.expected.hlsl b/test/buffer/storage/dynamic_index/write.wgsl.expected.hlsl index 9d035582db..74aeaf2068 100644 --- a/test/buffer/storage/dynamic_index/write.wgsl.expected.hlsl +++ b/test/buffer/storage/dynamic_index/write.wgsl.expected.hlsl @@ -1,3 +1,9 @@ +RWByteAddressBuffer s : register(u0, space0); + +struct tint_symbol_1 { + uint idx : SV_GroupIndex; +}; + void tint_symbol_8(RWByteAddressBuffer buffer, uint offset, float2x3 value) { buffer.Store3((offset + 0u), asuint(value[0u])); buffer.Store3((offset + 16u), asuint(value[1u])); @@ -16,12 +22,6 @@ void tint_symbol_12(RWByteAddressBuffer buffer, uint offset, int4 value[4]) { buffer.Store4((offset + 48u), asuint(value[3u])); } -RWByteAddressBuffer s : register(u0, space0); - -struct tint_symbol_1 { - uint idx : SV_GroupIndex; -}; - [numthreads(1, 1, 1)] void main(tint_symbol_1 tint_symbol) { const uint idx = tint_symbol.idx; diff --git a/test/buffer/storage/static_index/read.wgsl.expected.hlsl b/test/buffer/storage/static_index/read.wgsl.expected.hlsl index 95ed2723a0..c14caff8f4 100644 --- a/test/buffer/storage/static_index/read.wgsl.expected.hlsl +++ b/test/buffer/storage/static_index/read.wgsl.expected.hlsl @@ -5,6 +5,8 @@ struct tint_padded_array_element { Inner el; }; +ByteAddressBuffer s : register(t0, space0); + float2x3 tint_symbol_6(ByteAddressBuffer buffer, uint offset) { return float2x3(asfloat(buffer.Load3((offset + 0u))), asfloat(buffer.Load3((offset + 16u)))); } @@ -24,8 +26,6 @@ tint_symbol_10_ret tint_symbol_10(ByteAddressBuffer buffer, uint offset) { return tint_symbol_12; } -ByteAddressBuffer s : register(t0, space0); - [numthreads(1, 1, 1)] void main() { const int3 a = asint(s.Load3(0u)); diff --git a/test/buffer/storage/static_index/write.wgsl.expected.hlsl b/test/buffer/storage/static_index/write.wgsl.expected.hlsl index 52c36c6646..7e569cdaa2 100644 --- a/test/buffer/storage/static_index/write.wgsl.expected.hlsl +++ b/test/buffer/storage/static_index/write.wgsl.expected.hlsl @@ -5,6 +5,8 @@ struct tint_padded_array_element { Inner el; }; +RWByteAddressBuffer s : register(u0, space0); + void tint_symbol_6(RWByteAddressBuffer buffer, uint offset, float2x3 value) { buffer.Store3((offset + 0u), asuint(value[0u])); buffer.Store3((offset + 16u), asuint(value[1u])); @@ -27,8 +29,6 @@ void tint_symbol_10(RWByteAddressBuffer buffer, uint offset, tint_padded_array_e tint_symbol_9(buffer, (offset + 48u), value[3u].el); } -RWByteAddressBuffer s : register(u0, space0); - [numthreads(1, 1, 1)] void main() { s.Store3(0u, asuint(int3(0, 0, 0))); diff --git a/test/buffer/uniform/dynamic_index/read.wgsl.expected.hlsl b/test/buffer/uniform/dynamic_index/read.wgsl.expected.hlsl index 378e877731..22133cc826 100644 --- a/test/buffer/uniform/dynamic_index/read.wgsl.expected.hlsl +++ b/test/buffer/uniform/dynamic_index/read.wgsl.expected.hlsl @@ -1,9 +1,3 @@ -float2x3 tint_symbol_9(uint4 buffer[96], uint offset) { - const uint scalar_offset = ((offset + 0u)) / 4; - const uint scalar_offset_1 = ((offset + 16u)) / 4; - return float2x3(asfloat(buffer[scalar_offset / 4].xyz), asfloat(buffer[scalar_offset_1 / 4].xyz)); -} - cbuffer cbuffer_s : register(b0, space0) { uint4 s[96]; }; @@ -12,6 +6,12 @@ struct tint_symbol_1 { uint idx : SV_GroupIndex; }; +float2x3 tint_symbol_9(uint4 buffer[96], uint offset) { + const uint scalar_offset = ((offset + 0u)) / 4; + const uint scalar_offset_1 = ((offset + 16u)) / 4; + return float2x3(asfloat(buffer[scalar_offset / 4].xyz), asfloat(buffer[scalar_offset_1 / 4].xyz)); +} + [numthreads(1, 1, 1)] void main(tint_symbol_1 tint_symbol) { const uint idx = tint_symbol.idx; diff --git a/test/buffer/uniform/static_index/read.wgsl.expected.hlsl b/test/buffer/uniform/static_index/read.wgsl.expected.hlsl index 7acef2bd7b..fa6d133619 100644 --- a/test/buffer/uniform/static_index/read.wgsl.expected.hlsl +++ b/test/buffer/uniform/static_index/read.wgsl.expected.hlsl @@ -5,6 +5,10 @@ struct tint_padded_array_element { Inner el; }; +cbuffer cbuffer_s : register(b0, space0) { + uint4 s[13]; +}; + float2x3 tint_symbol_7(uint4 buffer[13], uint offset) { const uint scalar_offset = ((offset + 0u)) / 4; const uint scalar_offset_1 = ((offset + 16u)) / 4; @@ -33,10 +37,6 @@ tint_symbol_11_ret tint_symbol_11(uint4 buffer[13], uint offset) { return tint_symbol_13; } -cbuffer cbuffer_s : register(b0, space0) { - uint4 s[13]; -}; - [numthreads(1, 1, 1)] void main() { const int3 a = asint(s[0].xyz); diff --git a/test/bug/tint/403.wgsl.expected.hlsl b/test/bug/tint/403.wgsl.expected.hlsl index 51b6d4e072..ef0848ccce 100644 --- a/test/bug/tint/403.wgsl.expected.hlsl +++ b/test/bug/tint/403.wgsl.expected.hlsl @@ -1,3 +1,17 @@ +cbuffer cbuffer_x_20 : register(b0, space0) { + uint4 x_20[1]; +}; +cbuffer cbuffer_x_26 : register(b0, space1) { + uint4 x_26[1]; +}; + +struct tint_symbol_1 { + uint gl_VertexIndex : SV_VertexID; +}; +struct tint_symbol_2 { + float4 value : SV_Position; +}; + float2x2 tint_symbol_4(uint4 buffer[1], uint offset) { const uint scalar_offset = ((offset + 0u)) / 4; uint4 ubo_load = buffer[scalar_offset / 4]; @@ -14,20 +28,6 @@ float2x2 tint_symbol_6(uint4 buffer[1], uint offset) { return float2x2(asfloat(((scalar_offset_2 & 2) ? ubo_load_2.zw : ubo_load_2.xy)), asfloat(((scalar_offset_3 & 2) ? ubo_load_3.zw : ubo_load_3.xy))); } -cbuffer cbuffer_x_20 : register(b0, space0) { - uint4 x_20[1]; -}; -cbuffer cbuffer_x_26 : register(b0, space1) { - uint4 x_26[1]; -}; - -struct tint_symbol_1 { - uint gl_VertexIndex : SV_VertexID; -}; -struct tint_symbol_2 { - float4 value : SV_Position; -}; - tint_symbol_2 main(tint_symbol_1 tint_symbol) { const uint gl_VertexIndex = tint_symbol.gl_VertexIndex; float2 indexable[3] = (float2[3])0; diff --git a/test/bug/tint/870.spvasm.expected.hlsl b/test/bug/tint/870.spvasm.expected.hlsl index db135e64bf..b3b5071b52 100644 --- a/test/bug/tint/870.spvasm.expected.hlsl +++ b/test/bug/tint/870.spvasm.expected.hlsl @@ -1,11 +1,11 @@ +ByteAddressBuffer sspp962805860buildInformation : register(t2, space0); + typedef int tint_symbol_1_ret[6]; tint_symbol_1_ret tint_symbol_1(ByteAddressBuffer buffer, uint offset) { const int tint_symbol_2[6] = {asint(buffer.Load((offset + 0u))), asint(buffer.Load((offset + 4u))), asint(buffer.Load((offset + 8u))), asint(buffer.Load((offset + 12u))), asint(buffer.Load((offset + 16u))), asint(buffer.Load((offset + 20u)))}; return tint_symbol_2; } -ByteAddressBuffer sspp962805860buildInformation : register(t2, space0); - void main_1() { int orientation[6] = (int[6])0; const int x_23[6] = tint_symbol_1(sspp962805860buildInformation, 36u); diff --git a/test/bug/tint/922.wgsl.expected.hlsl b/test/bug/tint/922.wgsl.expected.hlsl index 60eb71643d..e0b2049f07 100644 --- a/test/bug/tint/922.wgsl.expected.hlsl +++ b/test/bug/tint/922.wgsl.expected.hlsl @@ -13,31 +13,6 @@ struct Mat4x2_ { float4 mx; float4 my; }; - -Mat4x4_ tint_symbol_7(uint4 buffer[4], uint offset) { - const uint scalar_offset = ((offset + 0u)) / 4; - const uint scalar_offset_1 = ((offset + 16u)) / 4; - const uint scalar_offset_2 = ((offset + 32u)) / 4; - const uint scalar_offset_3 = ((offset + 48u)) / 4; - const Mat4x4_ tint_symbol_10 = {asfloat(buffer[scalar_offset / 4]), asfloat(buffer[scalar_offset_1 / 4]), asfloat(buffer[scalar_offset_2 / 4]), asfloat(buffer[scalar_offset_3 / 4])}; - return tint_symbol_10; -} - -Mat4x2_ tint_symbol_9(uint4 buffer[3], uint offset) { - const uint scalar_offset_4 = ((offset + 0u)) / 4; - const uint scalar_offset_5 = ((offset + 16u)) / 4; - const Mat4x2_ tint_symbol_11 = {asfloat(buffer[scalar_offset_4 / 4]), asfloat(buffer[scalar_offset_5 / 4])}; - return tint_symbol_11; -} - -Mat4x3_ tint_symbol_5(uint4 buffer[96], uint offset) { - const uint scalar_offset_6 = ((offset + 0u)) / 4; - const uint scalar_offset_7 = ((offset + 16u)) / 4; - const uint scalar_offset_8 = ((offset + 32u)) / 4; - const Mat4x3_ tint_symbol_12 = {asfloat(buffer[scalar_offset_6 / 4]), asfloat(buffer[scalar_offset_7 / 4]), asfloat(buffer[scalar_offset_8 / 4])}; - return tint_symbol_12; -} - struct VertexOutput { float4 v_Color; float2 v_TexCoord; @@ -179,6 +154,30 @@ Mat4x3_ _Mat4x3_1(Mat4x4_ m20) { return o4; } +Mat4x3_ tint_symbol_5(uint4 buffer[96], uint offset) { + const uint scalar_offset = ((offset + 0u)) / 4; + const uint scalar_offset_1 = ((offset + 16u)) / 4; + const uint scalar_offset_2 = ((offset + 32u)) / 4; + const Mat4x3_ tint_symbol_10 = {asfloat(buffer[scalar_offset / 4]), asfloat(buffer[scalar_offset_1 / 4]), asfloat(buffer[scalar_offset_2 / 4])}; + return tint_symbol_10; +} + +Mat4x4_ tint_symbol_7(uint4 buffer[4], uint offset) { + const uint scalar_offset_3 = ((offset + 0u)) / 4; + const uint scalar_offset_4 = ((offset + 16u)) / 4; + const uint scalar_offset_5 = ((offset + 32u)) / 4; + const uint scalar_offset_6 = ((offset + 48u)) / 4; + const Mat4x4_ tint_symbol_11 = {asfloat(buffer[scalar_offset_3 / 4]), asfloat(buffer[scalar_offset_4 / 4]), asfloat(buffer[scalar_offset_5 / 4]), asfloat(buffer[scalar_offset_6 / 4])}; + return tint_symbol_11; +} + +Mat4x2_ tint_symbol_9(uint4 buffer[3], uint offset) { + const uint scalar_offset_7 = ((offset + 0u)) / 4; + const uint scalar_offset_8 = ((offset + 16u)) / 4; + const Mat4x2_ tint_symbol_12 = {asfloat(buffer[scalar_offset_7 / 4]), asfloat(buffer[scalar_offset_8 / 4])}; + return tint_symbol_12; +} + void main1() { Mat4x3_ t_PosMtx = (Mat4x3_)0; float2 t_TexSpaceCoord = float2(0.0f, 0.0f); diff --git a/test/bug/tint/993.wgsl b/test/bug/tint/993.wgsl new file mode 100644 index 0000000000..d9167f892b --- /dev/null +++ b/test/bug/tint/993.wgsl @@ -0,0 +1,24 @@ + +[[block]] struct Constants { + zero: u32; +}; +[[group(1), binding(0)]] var constants: Constants; + +[[block]] struct Result { + value: u32; +}; +[[group(1), binding(1)]] var result: Result; + +[[block]] struct TestData { + data: array,3>; +}; +[[group(0), binding(0)]] var s: TestData; + +fn runTest() -> i32 { + return atomicLoad(&s.data[(0u) + u32(constants.zero)]); +} + +[[stage(compute), workgroup_size(1)]] +fn main() { + result.value = u32(runTest()); +} \ No newline at end of file diff --git a/test/bug/tint/993.wgsl.expected.hlsl b/test/bug/tint/993.wgsl.expected.hlsl new file mode 100644 index 0000000000..2b16fd28cf --- /dev/null +++ b/test/bug/tint/993.wgsl.expected.hlsl @@ -0,0 +1,23 @@ +int atomicLoad_1(RWByteAddressBuffer buffer, uint offset) { + int value = 0; + buffer.InterlockedOr(offset, 0, value); + return value; +} + +cbuffer cbuffer_constants : register(b0, space1) { + uint4 constants[1]; +}; + +RWByteAddressBuffer result : register(u1, space1); + +RWByteAddressBuffer s : register(u0, space0); + +int runTest() { + return atomicLoad_1(s, (4u * (0u + uint(constants[0].x)))); +} + +[numthreads(1, 1, 1)] +void main() { + result.Store(0u, asuint(uint(runTest()))); + return; +} diff --git a/test/bug/tint/993.wgsl.expected.msl b/test/bug/tint/993.wgsl.expected.msl new file mode 100644 index 0000000000..50f64ad556 --- /dev/null +++ b/test/bug/tint/993.wgsl.expected.msl @@ -0,0 +1,34 @@ +SKIP: FAILED + +#include + +using namespace metal; +struct Constants { + /* 0x0000 */ uint zero; +}; +struct Result { + /* 0x0000 */ uint value; +}; +struct tint_array_wrapper { + /* 0x0000 */ atomic_int arr[3]; +}; +struct TestData { + /* 0x0000 */ tint_array_wrapper data; +}; + +int runTest(constant Constants& constants, device TestData& s) { + return atomic_load_explicit(&(s.data.arr[(0u + uint(constants.zero))]), memory_order_relaxed); +} + +kernel void tint_symbol(constant Constants& constants [[buffer(0)]], device Result& result [[buffer(1)]], device TestData& s [[buffer(0)]]) { + result.value = uint(runTest(constants, s)); + return; +} + +Compilation failed: + +program_source:21:124: error: cannot reserve 'buffer' resource location at index 0 +kernel void tint_symbol(constant Constants& constants [[buffer(0)]], device Result& result [[buffer(1)]], device TestData& s [[buffer(0)]]) { + ^ + + diff --git a/test/bug/tint/993.wgsl.expected.spvasm b/test/bug/tint/993.wgsl.expected.spvasm new file mode 100644 index 0000000000..31b94d8c8c --- /dev/null +++ b/test/bug/tint/993.wgsl.expected.spvasm @@ -0,0 +1,73 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 36 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %Constants "Constants" + OpMemberName %Constants 0 "zero" + OpName %constants "constants" + OpName %Result "Result" + OpMemberName %Result 0 "value" + OpName %result "result" + OpName %TestData "TestData" + OpMemberName %TestData 0 "data" + OpName %s "s" + OpName %runTest "runTest" + OpName %main "main" + OpDecorate %Constants Block + OpMemberDecorate %Constants 0 Offset 0 + OpDecorate %constants NonWritable + OpDecorate %constants DescriptorSet 1 + OpDecorate %constants Binding 0 + OpDecorate %Result Block + OpMemberDecorate %Result 0 Offset 0 + OpDecorate %result NonReadable + OpDecorate %result DescriptorSet 1 + OpDecorate %result Binding 1 + OpDecorate %TestData Block + OpMemberDecorate %TestData 0 Offset 0 + OpDecorate %_arr_int_uint_3 ArrayStride 4 + OpDecorate %s DescriptorSet 0 + OpDecorate %s Binding 0 + %uint = OpTypeInt 32 0 + %Constants = OpTypeStruct %uint +%_ptr_Uniform_Constants = OpTypePointer Uniform %Constants + %constants = OpVariable %_ptr_Uniform_Constants Uniform + %Result = OpTypeStruct %uint +%_ptr_StorageBuffer_Result = OpTypePointer StorageBuffer %Result + %result = OpVariable %_ptr_StorageBuffer_Result StorageBuffer + %int = OpTypeInt 32 1 + %uint_3 = OpConstant %uint 3 +%_arr_int_uint_3 = OpTypeArray %int %uint_3 + %TestData = OpTypeStruct %_arr_int_uint_3 +%_ptr_StorageBuffer_TestData = OpTypePointer StorageBuffer %TestData + %s = OpVariable %_ptr_StorageBuffer_TestData StorageBuffer + %14 = OpTypeFunction %int + %uint_1 = OpConstant %uint 1 + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint +%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int + %void = OpTypeVoid + %28 = OpTypeFunction %void +%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint + %runTest = OpFunction %int None %14 + %16 = OpLabel + %23 = OpAccessChain %_ptr_Uniform_uint %constants %uint_0 + %24 = OpLoad %uint %23 + %25 = OpIAdd %uint %uint_0 %24 + %27 = OpAccessChain %_ptr_StorageBuffer_int %s %uint_0 %25 + %17 = OpAtomicLoad %int %27 %uint_1 %uint_0 + OpReturnValue %17 + OpFunctionEnd + %main = OpFunction %void None %28 + %31 = OpLabel + %33 = OpAccessChain %_ptr_StorageBuffer_uint %result %uint_0 + %35 = OpFunctionCall %int %runTest + %34 = OpBitcast %uint %35 + OpStore %33 %34 + OpReturn + OpFunctionEnd diff --git a/test/bug/tint/993.wgsl.expected.wgsl b/test/bug/tint/993.wgsl.expected.wgsl new file mode 100644 index 0000000000..b5bf11b2b0 --- /dev/null +++ b/test/bug/tint/993.wgsl.expected.wgsl @@ -0,0 +1,29 @@ +[[block]] +struct Constants { + zero : u32; +}; + +[[group(1), binding(0)]] var constants : Constants; + +[[block]] +struct Result { + value : u32; +}; + +[[group(1), binding(1)]] var result : Result; + +[[block]] +struct TestData { + data : array, 3>; +}; + +[[group(0), binding(0)]] var s : TestData; + +fn runTest() -> i32 { + return atomicLoad(&(s.data[(0u + u32(constants.zero))])); +} + +[[stage(compute), workgroup_size(1)]] +fn main() { + result.value = u32(runTest()); +} diff --git a/test/samples/cube.wgsl.expected.hlsl b/test/samples/cube.wgsl.expected.hlsl index 851f523d7a..9791b42083 100644 --- a/test/samples/cube.wgsl.expected.hlsl +++ b/test/samples/cube.wgsl.expected.hlsl @@ -1,11 +1,3 @@ -float4x4 tint_symbol_7(uint4 buffer[4], uint offset) { - const uint scalar_offset = ((offset + 0u)) / 4; - const uint scalar_offset_1 = ((offset + 16u)) / 4; - const uint scalar_offset_2 = ((offset + 32u)) / 4; - const uint scalar_offset_3 = ((offset + 48u)) / 4; - return float4x4(asfloat(buffer[scalar_offset / 4]), asfloat(buffer[scalar_offset_1 / 4]), asfloat(buffer[scalar_offset_2 / 4]), asfloat(buffer[scalar_offset_3 / 4])); -} - cbuffer cbuffer_uniforms : register(b0, space0) { uint4 uniforms[4]; }; @@ -27,6 +19,14 @@ struct tint_symbol_2 { float4 Position : SV_Position; }; +float4x4 tint_symbol_7(uint4 buffer[4], uint offset) { + const uint scalar_offset = ((offset + 0u)) / 4; + const uint scalar_offset_1 = ((offset + 16u)) / 4; + const uint scalar_offset_2 = ((offset + 32u)) / 4; + const uint scalar_offset_3 = ((offset + 48u)) / 4; + return float4x4(asfloat(buffer[scalar_offset / 4]), asfloat(buffer[scalar_offset_1 / 4]), asfloat(buffer[scalar_offset_2 / 4]), asfloat(buffer[scalar_offset_3 / 4])); +} + tint_symbol_2 vtx_main(tint_symbol_1 tint_symbol) { const VertexInput input = {tint_symbol.cur_position, tint_symbol.color}; VertexOutput output = (VertexOutput)0; diff --git a/test/shader_io/shared_struct_storage_buffer.wgsl.expected.hlsl b/test/shader_io/shared_struct_storage_buffer.wgsl.expected.hlsl index 1c8ad7dfb8..af4440eae8 100644 --- a/test/shader_io/shared_struct_storage_buffer.wgsl.expected.hlsl +++ b/test/shader_io/shared_struct_storage_buffer.wgsl.expected.hlsl @@ -4,12 +4,6 @@ struct S { float4 v; }; -void tint_symbol_5(RWByteAddressBuffer buffer, uint offset, S value) { - buffer.Store((offset + 0u), asuint(value.f)); - buffer.Store((offset + 4u), asuint(value.u)); - buffer.Store4((offset + 128u), asuint(value.v)); -} - RWByteAddressBuffer output : register(u0, space0); struct tint_symbol_1 { @@ -18,6 +12,12 @@ struct tint_symbol_1 { float4 v : SV_Position; }; +void tint_symbol_5(RWByteAddressBuffer buffer, uint offset, S value) { + buffer.Store((offset + 0u), asuint(value.f)); + buffer.Store((offset + 4u), asuint(value.u)); + buffer.Store4((offset + 128u), asuint(value.v)); +} + void frag_main(tint_symbol_1 tint_symbol) { const S input = {tint_symbol.f, tint_symbol.u, tint_symbol.v}; const float f = input.f;