[tint][inspector] Walk globals only once in GetResourceBindings
Previously many bespoke functions walked the globals computed a list of
ResourceBinding, then results were aggregated. Instead walks globals
once and for each of them decide if/how it should be turned into a
ResourceBinding.
This will simplify the future addition of support for BindingArray that
will only need to be added in one place.
end2end tests for MSL are updated because the order of bindings is
modified by the sort in GetResourceBindings (that's used by
FlattenBindings)
Bug: 393558555
Change-Id: I795d4a60c444415bf5f69a513cbe4d318363bd42
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/229595
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Commit-Queue: Corentin Wallez <cwallez@chromium.org>
diff --git a/src/tint/lang/wgsl/inspector/inspector.cc b/src/tint/lang/wgsl/inspector/inspector.cc
index 25d44c2..b329a24 100644
--- a/src/tint/lang/wgsl/inspector/inspector.cc
+++ b/src/tint/lang/wgsl/inspector/inspector.cc
@@ -75,17 +75,6 @@
namespace tint::inspector {
namespace {
-void AppendResourceBindings(std::vector<ResourceBinding>* dest,
- const std::vector<ResourceBinding>& orig) {
- TINT_ASSERT(dest);
- if (!dest) {
- return;
- }
-
- dest->reserve(dest->size() + orig.size());
- dest->insert(dest->end(), orig.begin(), orig.end());
-}
-
std::tuple<ComponentType, CompositionType> CalculateComponentAndComposition(
const core::type::Type* type) {
// entry point in/out variables must of numeric scalar or vector types.
@@ -131,6 +120,106 @@
builtin_fn == wgsl::BuiltinFn::kTextureGather;
}
+ResourceBinding ConvertBufferToResourceBinding(const tint::sem::GlobalVariable* buffer) {
+ ResourceBinding result;
+ result.bind_group = buffer->Attributes().binding_point->group;
+ result.binding = buffer->Attributes().binding_point->binding;
+ result.variable_name = buffer->Declaration()->name->symbol.Name();
+
+ auto* unwrapped_type = buffer->Type()->UnwrapRef();
+ result.size = unwrapped_type->Size();
+ result.size_no_padding = result.size;
+ if (auto* str = unwrapped_type->As<sem::Struct>()) {
+ result.size_no_padding = str->SizeNoPadding();
+ }
+
+ if (buffer->AddressSpace() == core::AddressSpace::kStorage) {
+ if (buffer->Access() == core::Access::kReadWrite) {
+ result.resource_type = ResourceBinding::ResourceType::kStorageBuffer;
+ } else {
+ TINT_ASSERT(buffer->Access() == core::Access::kRead);
+ result.resource_type = ResourceBinding::ResourceType::kReadOnlyStorageBuffer;
+ }
+ } else {
+ TINT_ASSERT(buffer->AddressSpace() == core::AddressSpace::kUniform);
+ result.resource_type = ResourceBinding::ResourceType::kUniformBuffer;
+ }
+
+ return result;
+}
+
+ResourceBinding ConvertHandleToResourceBinding(const tint::sem::GlobalVariable* handle) {
+ ResourceBinding result;
+ result.bind_group = handle->Attributes().binding_point->group;
+ result.binding = handle->Attributes().binding_point->binding;
+ result.variable_name = handle->Declaration()->name->symbol.Name();
+
+ const core::type::Type* handle_type = handle->Type()->UnwrapRef();
+ Switch(
+ handle_type,
+
+ [&](const core::type::Sampler* sampler) {
+ if (sampler->Kind() == core::type::SamplerKind::kSampler) {
+ result.resource_type = ResourceBinding::ResourceType::kSampler;
+ } else {
+ TINT_ASSERT(sampler->Kind() == core::type::SamplerKind::kComparisonSampler);
+ result.resource_type = ResourceBinding::ResourceType::kComparisonSampler;
+ }
+ },
+
+ [&](const core::type::SampledTexture* tex) {
+ result.resource_type = ResourceBinding::ResourceType::kSampledTexture;
+ result.dim = TypeTextureDimensionToResourceBindingTextureDimension(tex->Dim());
+ result.sampled_kind = BaseTypeToSampledKind(tex->Type());
+ },
+ [&](const core::type::MultisampledTexture* tex) {
+ result.resource_type = ResourceBinding::ResourceType::kMultisampledTexture;
+ result.dim = TypeTextureDimensionToResourceBindingTextureDimension(tex->Dim());
+ result.sampled_kind = BaseTypeToSampledKind(tex->Type());
+ },
+ [&](const core::type::DepthTexture* tex) {
+ result.resource_type = ResourceBinding::ResourceType::kDepthTexture;
+ result.dim = TypeTextureDimensionToResourceBindingTextureDimension(tex->Dim());
+ },
+ [&](const core::type::DepthMultisampledTexture* tex) {
+ result.resource_type = ResourceBinding::ResourceType::kDepthMultisampledTexture;
+ result.dim = TypeTextureDimensionToResourceBindingTextureDimension(tex->Dim());
+ },
+ [&](const core::type::StorageTexture* tex) {
+ switch (tex->Access()) {
+ case core::Access::kWrite:
+ result.resource_type = ResourceBinding::ResourceType::kWriteOnlyStorageTexture;
+ break;
+ case core::Access::kReadWrite:
+ result.resource_type = ResourceBinding::ResourceType::kReadWriteStorageTexture;
+ break;
+ case core::Access::kRead:
+ result.resource_type = ResourceBinding::ResourceType::kReadOnlyStorageTexture;
+ break;
+ case core::Access::kUndefined:
+ TINT_UNREACHABLE() << "unhandled storage texture access";
+ }
+ result.dim = TypeTextureDimensionToResourceBindingTextureDimension(tex->Dim());
+ result.sampled_kind = BaseTypeToSampledKind(tex->Type());
+ result.image_format = TypeTexelFormatToResourceBindingTexelFormat(tex->TexelFormat());
+ },
+ [&](const core::type::ExternalTexture*) {
+ result.resource_type = ResourceBinding::ResourceType::kExternalTexture;
+ result.dim = ResourceBinding::TextureDimension::k2d;
+ },
+
+ [&](const core::type::InputAttachment* attachment) {
+ result.resource_type = ResourceBinding::ResourceType::kInputAttachment;
+ result.input_attachment_index = handle->Attributes().input_attachment_index.value();
+ result.sampled_kind = BaseTypeToSampledKind(attachment->Type());
+ result.dim = TypeTextureDimensionToResourceBindingTextureDimension(attachment->Dim());
+ },
+
+ TINT_ICE_ON_NO_MATCH);
+
+ return result;
+}
+
} // namespace
Inspector::Inspector(const Program& program) : program_(program) {}
@@ -367,218 +456,28 @@
}
std::vector<ResourceBinding> result;
- for (auto fn : {
- &Inspector::GetUniformBufferResourceBindings,
- &Inspector::GetStorageBufferResourceBindings,
- &Inspector::GetReadOnlyStorageBufferResourceBindings,
- &Inspector::GetSamplerResourceBindings,
- &Inspector::GetComparisonSamplerResourceBindings,
- &Inspector::GetSampledTextureResourceBindings,
- &Inspector::GetMultisampledTextureResourceBindings,
- &Inspector::GetStorageTextureResourceBindings,
- &Inspector::GetDepthTextureResourceBindings,
- &Inspector::GetDepthMultisampledTextureResourceBindings,
- &Inspector::GetExternalTextureResourceBindings,
- &Inspector::GetInputAttachmentResourceBindings,
- }) {
- AppendResourceBindings(&result, (this->*fn)(entry_point));
- }
- return result;
-}
-
-std::vector<ResourceBinding> Inspector::GetUniformBufferResourceBindings(
- const std::string& entry_point) {
- auto* func = FindEntryPointByName(entry_point);
- if (!func) {
- return {};
- }
-
- std::vector<ResourceBinding> result;
-
auto* func_sem = program_.Sem().Get(func);
- for (auto& ruv : func_sem->TransitivelyReferencedUniformVariables()) {
- auto* var = ruv.first;
- auto binding_info = ruv.second;
+ for (auto& global : func_sem->TransitivelyReferencedGlobals()) {
+ switch (global->AddressSpace()) {
+ // Resources cannot be in these address spaces.
+ case core::AddressSpace::kPrivate:
+ case core::AddressSpace::kFunction:
+ case core::AddressSpace::kWorkgroup:
+ case core::AddressSpace::kPushConstant:
+ case core::AddressSpace::kPixelLocal:
+ case core::AddressSpace::kIn:
+ case core::AddressSpace::kOut:
+ case core::AddressSpace::kUndefined:
+ continue;
- auto* unwrapped_type = var->Type()->UnwrapRef();
-
- ResourceBinding entry;
- entry.resource_type = ResourceBinding::ResourceType::kUniformBuffer;
- entry.bind_group = binding_info.group;
- entry.binding = binding_info.binding;
- entry.size = unwrapped_type->Size();
- entry.size_no_padding = entry.size;
- if (auto* str = unwrapped_type->As<sem::Struct>()) {
- entry.size_no_padding = str->SizeNoPadding();
- } else {
- entry.size_no_padding = entry.size;
+ case core::AddressSpace::kUniform:
+ case core::AddressSpace::kStorage:
+ result.push_back(ConvertBufferToResourceBinding(global));
+ break;
+ case core::AddressSpace::kHandle:
+ result.push_back(ConvertHandleToResourceBinding(global));
+ break;
}
- entry.variable_name = var->Declaration()->name->symbol.Name();
-
- result.push_back(entry);
- }
-
- return result;
-}
-
-std::vector<ResourceBinding> Inspector::GetStorageBufferResourceBindings(
- const std::string& entry_point) {
- return GetStorageBufferResourceBindingsImpl(entry_point, false);
-}
-
-std::vector<ResourceBinding> Inspector::GetReadOnlyStorageBufferResourceBindings(
- const std::string& entry_point) {
- return GetStorageBufferResourceBindingsImpl(entry_point, true);
-}
-
-std::vector<ResourceBinding> Inspector::GetSamplerResourceBindings(const std::string& entry_point) {
- auto* func = FindEntryPointByName(entry_point);
- if (!func) {
- return {};
- }
-
- std::vector<ResourceBinding> result;
-
- auto* func_sem = program_.Sem().Get(func);
- for (auto& rs : func_sem->TransitivelyReferencedSamplerVariables()) {
- auto binding_info = rs.second;
-
- ResourceBinding entry;
- entry.resource_type = ResourceBinding::ResourceType::kSampler;
- entry.bind_group = binding_info.group;
- entry.binding = binding_info.binding;
- entry.variable_name = rs.first->Declaration()->name->symbol.Name();
-
- result.push_back(entry);
- }
-
- return result;
-}
-
-std::vector<ResourceBinding> Inspector::GetComparisonSamplerResourceBindings(
- const std::string& entry_point) {
- auto* func = FindEntryPointByName(entry_point);
- if (!func) {
- return {};
- }
-
- std::vector<ResourceBinding> result;
-
- auto* func_sem = program_.Sem().Get(func);
- for (auto& rcs : func_sem->TransitivelyReferencedComparisonSamplerVariables()) {
- auto binding_info = rcs.second;
-
- ResourceBinding entry;
- entry.resource_type = ResourceBinding::ResourceType::kComparisonSampler;
- entry.bind_group = binding_info.group;
- entry.binding = binding_info.binding;
- entry.variable_name = rcs.first->Declaration()->name->symbol.Name();
-
- result.push_back(entry);
- }
-
- return result;
-}
-
-std::vector<ResourceBinding> Inspector::GetSampledTextureResourceBindings(
- const std::string& entry_point) {
- return GetSampledTextureResourceBindingsImpl(entry_point, false);
-}
-
-std::vector<ResourceBinding> Inspector::GetMultisampledTextureResourceBindings(
- const std::string& entry_point) {
- return GetSampledTextureResourceBindingsImpl(entry_point, true);
-}
-
-std::vector<ResourceBinding> Inspector::GetStorageTextureResourceBindings(
- const std::string& entry_point) {
- return GetStorageTextureResourceBindingsImpl(entry_point);
-}
-
-std::vector<ResourceBinding> Inspector::GetTextureResourceBindings(
- const std::string& entry_point,
- const tint::TypeInfo* texture_type,
- ResourceBinding::ResourceType resource_type) {
- auto* func = FindEntryPointByName(entry_point);
- if (!func) {
- return {};
- }
-
- std::vector<ResourceBinding> result;
- auto* func_sem = program_.Sem().Get(func);
- for (auto& ref : func_sem->TransitivelyReferencedVariablesOfType(texture_type)) {
- auto* var = ref.first;
- auto binding_info = ref.second;
-
- ResourceBinding entry;
- entry.resource_type = resource_type;
- entry.bind_group = binding_info.group;
- entry.binding = binding_info.binding;
- entry.variable_name = var->Declaration()->name->symbol.Name();
-
- auto* tex = var->Type()->UnwrapRef()->As<core::type::Texture>();
- entry.dim = TypeTextureDimensionToResourceBindingTextureDimension(tex->Dim());
-
- result.push_back(entry);
- }
-
- return result;
-}
-
-std::vector<ResourceBinding> Inspector::GetDepthTextureResourceBindings(
- const std::string& entry_point) {
- return GetTextureResourceBindings(entry_point, &tint::TypeInfo::Of<core::type::DepthTexture>(),
- ResourceBinding::ResourceType::kDepthTexture);
-}
-
-std::vector<ResourceBinding> Inspector::GetDepthMultisampledTextureResourceBindings(
- const std::string& entry_point) {
- return GetTextureResourceBindings(entry_point,
- &tint::TypeInfo::Of<core::type::DepthMultisampledTexture>(),
- ResourceBinding::ResourceType::kDepthMultisampledTexture);
-}
-
-std::vector<ResourceBinding> Inspector::GetExternalTextureResourceBindings(
- const std::string& entry_point) {
- return GetTextureResourceBindings(entry_point,
- &tint::TypeInfo::Of<core::type::ExternalTexture>(),
- ResourceBinding::ResourceType::kExternalTexture);
-}
-
-std::vector<ResourceBinding> Inspector::GetInputAttachmentResourceBindings(
- const std::string& entry_point) {
- auto* func = FindEntryPointByName(entry_point);
- if (!func) {
- return {};
- }
-
- std::vector<ResourceBinding> result;
- auto* func_sem = program_.Sem().Get(func);
- for (auto& ref : func_sem->TransitivelyReferencedVariablesOfType(
- &tint::TypeInfo::Of<core::type::InputAttachment>())) {
- auto* var = ref.first;
- auto binding_info = ref.second;
-
- ResourceBinding entry;
- entry.resource_type = ResourceBinding::ResourceType::kInputAttachment;
- entry.bind_group = binding_info.group;
- entry.binding = binding_info.binding;
-
- auto* sem_var = var->As<sem::GlobalVariable>();
- TINT_ASSERT(sem_var);
- TINT_ASSERT(sem_var->Attributes().input_attachment_index);
- entry.input_attachmnt_index = sem_var->Attributes().input_attachment_index.value();
-
- auto* input_attachment_type = var->Type()->UnwrapRef()->As<core::type::InputAttachment>();
- auto* base_type = input_attachment_type->Type();
- entry.sampled_kind = BaseTypeToSampledKind(base_type);
-
- entry.variable_name = var->Declaration()->name->symbol.Name();
-
- entry.dim =
- TypeTextureDimensionToResourceBindingTextureDimension(input_attachment_type->Dim());
-
- result.push_back(entry);
}
return result;
@@ -867,134 +766,6 @@
return std::nullopt;
}
-std::vector<ResourceBinding> Inspector::GetStorageBufferResourceBindingsImpl(
- const std::string& entry_point,
- bool read_only) {
- auto* func = FindEntryPointByName(entry_point);
- if (!func) {
- return {};
- }
-
- auto* func_sem = program_.Sem().Get(func);
- std::vector<ResourceBinding> result;
- for (auto& rsv : func_sem->TransitivelyReferencedStorageBufferVariables()) {
- auto* var = rsv.first;
- auto binding_info = rsv.second;
-
- if (read_only != (var->Access() == core::Access::kRead)) {
- continue;
- }
-
- auto* unwrapped_type = var->Type()->UnwrapRef();
-
- ResourceBinding entry;
- entry.resource_type = read_only ? ResourceBinding::ResourceType::kReadOnlyStorageBuffer
- : ResourceBinding::ResourceType::kStorageBuffer;
- entry.bind_group = binding_info.group;
- entry.binding = binding_info.binding;
- entry.size = unwrapped_type->Size();
- if (auto* str = unwrapped_type->As<sem::Struct>()) {
- entry.size_no_padding = str->SizeNoPadding();
- } else {
- entry.size_no_padding = entry.size;
- }
- entry.variable_name = var->Declaration()->name->symbol.Name();
-
- result.push_back(entry);
- }
-
- return result;
-}
-
-std::vector<ResourceBinding> Inspector::GetSampledTextureResourceBindingsImpl(
- const std::string& entry_point,
- bool multisampled_only) {
- auto* func = FindEntryPointByName(entry_point);
- if (!func) {
- return {};
- }
-
- std::vector<ResourceBinding> result;
- auto* func_sem = program_.Sem().Get(func);
- auto referenced_variables = multisampled_only
- ? func_sem->TransitivelyReferencedMultisampledTextureVariables()
- : func_sem->TransitivelyReferencedSampledTextureVariables();
- for (auto& ref : referenced_variables) {
- auto* var = ref.first;
- auto binding_info = ref.second;
-
- ResourceBinding entry;
- entry.resource_type = multisampled_only
- ? ResourceBinding::ResourceType::kMultisampledTexture
- : ResourceBinding::ResourceType::kSampledTexture;
- entry.bind_group = binding_info.group;
- entry.binding = binding_info.binding;
- entry.variable_name = var->Declaration()->name->symbol.Name();
-
- auto* texture_type = var->Type()->UnwrapRef()->As<core::type::Texture>();
- entry.dim = TypeTextureDimensionToResourceBindingTextureDimension(texture_type->Dim());
-
- const core::type::Type* base_type = nullptr;
- if (multisampled_only) {
- base_type = texture_type->As<core::type::MultisampledTexture>()->Type();
- } else {
- base_type = texture_type->As<core::type::SampledTexture>()->Type();
- }
- entry.sampled_kind = BaseTypeToSampledKind(base_type);
-
- result.push_back(entry);
- }
-
- return result;
-}
-
-std::vector<ResourceBinding> Inspector::GetStorageTextureResourceBindingsImpl(
- const std::string& entry_point) {
- auto* func = FindEntryPointByName(entry_point);
- if (!func) {
- return {};
- }
-
- auto* func_sem = program_.Sem().Get(func);
- std::vector<ResourceBinding> result;
- for (auto& ref :
- func_sem->TransitivelyReferencedVariablesOfType<core::type::StorageTexture>()) {
- auto* var = ref.first;
- auto binding_info = ref.second;
-
- auto* texture_type = var->Type()->UnwrapRef()->As<core::type::StorageTexture>();
-
- ResourceBinding entry;
- switch (texture_type->Access()) {
- case core::Access::kWrite:
- entry.resource_type = ResourceBinding::ResourceType::kWriteOnlyStorageTexture;
- break;
- case core::Access::kReadWrite:
- entry.resource_type = ResourceBinding::ResourceType::kReadWriteStorageTexture;
- break;
- case core::Access::kRead:
- entry.resource_type = ResourceBinding::ResourceType::kReadOnlyStorageTexture;
- break;
- case core::Access::kUndefined:
- TINT_UNREACHABLE() << "unhandled storage texture access";
- }
- entry.bind_group = binding_info.group;
- entry.binding = binding_info.binding;
- entry.variable_name = var->Declaration()->name->symbol.Name();
-
- entry.dim = TypeTextureDimensionToResourceBindingTextureDimension(texture_type->Dim());
-
- auto* base_type = texture_type->Type();
- entry.sampled_kind = BaseTypeToSampledKind(base_type);
- entry.image_format =
- TypeTexelFormatToResourceBindingTexelFormat(texture_type->TexelFormat());
-
- result.push_back(entry);
- }
-
- return result;
-}
-
std::tuple<InterpolationType, InterpolationSampling> Inspector::CalculateInterpolationData(
VectorRef<const ast::Attribute*> attributes) const {
auto* interpolation_attribute = ast::GetAttribute<ast::InterpolateAttribute>(attributes);
diff --git a/src/tint/lang/wgsl/inspector/inspector.h b/src/tint/lang/wgsl/inspector/inspector.h
index 356e06a..f248f8c 100644
--- a/src/tint/lang/wgsl/inspector/inspector.h
+++ b/src/tint/lang/wgsl/inspector/inspector.h
@@ -176,94 +176,6 @@
/// @returns the array length of the builtin clip_distances or empty when it is not used
std::optional<uint32_t> GetClipDistancesBuiltinSize(const core::type::Type* type) const;
- /// @param entry_point name of the entry point to get information about.
- /// @returns vector of all of the bindings for uniform buffers.
- std::vector<ResourceBinding> GetUniformBufferResourceBindings(const std::string& entry_point);
-
- /// @param entry_point name of the entry point to get information about.
- /// @returns vector of all of the bindings for storage buffers.
- std::vector<ResourceBinding> GetStorageBufferResourceBindings(const std::string& entry_point);
-
- /// @param entry_point name of the entry point to get information about.
- /// @returns vector of all of the bindings for read-only storage buffers.
- std::vector<ResourceBinding> GetReadOnlyStorageBufferResourceBindings(
- const std::string& entry_point);
-
- /// @param entry_point name of the entry point to get information about.
- /// @returns vector of all of the bindings for regular samplers.
- std::vector<ResourceBinding> GetSamplerResourceBindings(const std::string& entry_point);
-
- /// @param entry_point name of the entry point to get information about.
- /// @returns vector of all of the bindings for comparison samplers.
- std::vector<ResourceBinding> GetComparisonSamplerResourceBindings(
- const std::string& entry_point);
-
- /// @param entry_point name of the entry point to get information about.
- /// @returns vector of all of the bindings for sampled textures.
- std::vector<ResourceBinding> GetSampledTextureResourceBindings(const std::string& entry_point);
-
- /// @param entry_point name of the entry point to get information about.
- /// @returns vector of all of the bindings for multisampled textures.
- std::vector<ResourceBinding> GetMultisampledTextureResourceBindings(
- const std::string& entry_point);
-
- /// @param entry_point name of the entry point to get information about.
- /// @returns vector of all of the bindings for write-only storage textures.
- std::vector<ResourceBinding> GetStorageTextureResourceBindings(const std::string& entry_point);
-
- /// @param entry_point name of the entry point to get information about.
- /// @returns vector of all of the bindings for depth textures.
- std::vector<ResourceBinding> GetDepthTextureResourceBindings(const std::string& entry_point);
-
- /// @param entry_point name of the entry point to get information about.
- /// @returns vector of all of the bindings for depth textures.
- std::vector<ResourceBinding> GetDepthMultisampledTextureResourceBindings(
- const std::string& entry_point);
-
- /// Gathers all the resource bindings of the input attachment type for the given
- /// entry point.
- /// @param entry_point name of the entry point to get information about.
- /// texture type.
- /// @returns vector of all of the bindings for input attachments.
- std::vector<ResourceBinding> GetInputAttachmentResourceBindings(const std::string& entry_point);
-
- /// @param entry_point name of the entry point to get information about.
- /// @returns vector of all of the bindings for external textures.
- std::vector<ResourceBinding> GetExternalTextureResourceBindings(const std::string& entry_point);
-
- /// Gathers all the texture resource bindings of the given type for the given
- /// entry point.
- /// @param entry_point name of the entry point to get information about.
- /// @param texture_type the type of the textures to gather.
- /// @param resource_type the ResourceBinding::ResourceType for the given
- /// texture type.
- /// @returns vector of all of the bindings for depth textures.
- std::vector<ResourceBinding> GetTextureResourceBindings(
- const std::string& entry_point,
- const tint::TypeInfo* texture_type,
- ResourceBinding::ResourceType resource_type);
-
- /// @param entry_point name of the entry point to get information about.
- /// @param read_only if true get only read-only bindings, if false get
- /// write-only bindings.
- /// @returns vector of all of the bindings for the requested storage buffers.
- std::vector<ResourceBinding> GetStorageBufferResourceBindingsImpl(
- const std::string& entry_point,
- bool read_only);
-
- /// @param entry_point name of the entry point to get information about.
- /// @param multisampled_only only get multisampled textures if true, otherwise
- /// only get sampled textures.
- /// @returns vector of all of the bindings for the request storage buffers.
- std::vector<ResourceBinding> GetSampledTextureResourceBindingsImpl(
- const std::string& entry_point,
- bool multisampled_only);
-
- /// @param entry_point name of the entry point to get information about.
- /// @returns vector of all of the bindings for the requested storage textures.
- std::vector<ResourceBinding> GetStorageTextureResourceBindingsImpl(
- const std::string& entry_point);
-
/// Constructs |sampler_targets_| if it hasn't already been instantiated.
void GenerateSamplerTargets();
diff --git a/src/tint/lang/wgsl/inspector/inspector_test.cc b/src/tint/lang/wgsl/inspector/inspector_test.cc
index 45a2fdf..f668e6f 100644
--- a/src/tint/lang/wgsl/inspector/inspector_test.cc
+++ b/src/tint/lang/wgsl/inspector/inspector_test.cc
@@ -1980,15 +1980,6 @@
EXPECT_EQ(result["c"], program_->Sem().Get(c)->Attributes().override_id);
}
-/// Sorts the ResourceBindings using their bind point to make their order deterministic in tests
-/// @param resources the list of resources to sort in place
-void SortResourceBindings(std::vector<ResourceBinding>* resources) {
- std::sort(resources->begin(), resources->end(),
- [](const ResourceBinding& a, const ResourceBinding& b) {
- return std::tie(a.bind_group, a.binding) < std::tie(b.bind_group, b.binding);
- });
-}
-
TEST_F(InspectorGetResourceBindingsTest, Empty) {
MakeCallerBodyFunction("ep_func", tint::Empty,
Vector{
@@ -2076,7 +2067,6 @@
auto result = inspector.GetResourceBindings("ep_func");
ASSERT_FALSE(inspector.has_error()) << inspector.error();
ASSERT_EQ(9u, result.size());
- SortResourceBindings(&result);
EXPECT_EQ(ResourceBinding::ResourceType::kUniformBuffer, result[0].resource_type);
EXPECT_EQ(0u, result[0].bind_group);
@@ -2161,18 +2151,17 @@
auto result = inspector.GetResourceBindings("main");
ASSERT_FALSE(inspector.has_error()) << inspector.error();
ASSERT_EQ(2u, result.size());
- SortResourceBindings(&result);
EXPECT_EQ(ResourceBinding::ResourceType::kInputAttachment, result[0].resource_type);
EXPECT_EQ(0u, result[0].bind_group);
EXPECT_EQ(1u, result[0].binding);
- EXPECT_EQ(3u, result[0].input_attachmnt_index);
+ EXPECT_EQ(3u, result[0].input_attachment_index);
EXPECT_EQ(inspector::ResourceBinding::SampledKind::kFloat, result[0].sampled_kind);
EXPECT_EQ(ResourceBinding::ResourceType::kInputAttachment, result[1].resource_type);
EXPECT_EQ(4u, result[1].bind_group);
EXPECT_EQ(3u, result[1].binding);
- EXPECT_EQ(1u, result[1].input_attachmnt_index);
+ EXPECT_EQ(1u, result[1].input_attachment_index);
EXPECT_EQ(inspector::ResourceBinding::SampledKind::kSInt, result[1].sampled_kind);
}
@@ -2383,7 +2372,6 @@
auto result = inspector.GetResourceBindings("ep_func");
ASSERT_FALSE(inspector.has_error()) << inspector.error();
ASSERT_EQ(3u, result.size());
- SortResourceBindings(&result);
EXPECT_EQ(ResourceBinding::ResourceType::kUniformBuffer, result[0].resource_type);
EXPECT_EQ(0u, result[0].bind_group);
@@ -2567,7 +2555,6 @@
auto result = inspector.GetResourceBindings("ep_func");
ASSERT_FALSE(inspector.has_error()) << inspector.error();
ASSERT_EQ(3u, result.size());
- SortResourceBindings(&result);
EXPECT_EQ(ResourceBinding::ResourceType::kStorageBuffer, result[0].resource_type);
EXPECT_EQ(0u, result[0].bind_group);
@@ -2768,7 +2755,6 @@
auto result = inspector.GetResourceBindings("ep_func");
ASSERT_FALSE(inspector.has_error()) << inspector.error();
ASSERT_EQ(3u, result.size());
- SortResourceBindings(&result);
EXPECT_EQ(ResourceBinding::ResourceType::kReadOnlyStorageBuffer, result[0].resource_type);
EXPECT_EQ(0u, result[0].bind_group);
@@ -2807,9 +2793,9 @@
ASSERT_EQ(2u, result.size());
- EXPECT_EQ(ResourceBinding::ResourceType::kSampler, result[0].resource_type);
- EXPECT_EQ(0u, result[0].bind_group);
- EXPECT_EQ(0u, result[0].binding);
+ EXPECT_EQ(ResourceBinding::ResourceType::kSampler, result[1].resource_type);
+ EXPECT_EQ(0u, result[1].bind_group);
+ EXPECT_EQ(0u, result[1].binding);
}
TEST_F(InspectorGetResourceBindingsTest, Sampler_InFunction) {
@@ -2832,9 +2818,9 @@
ASSERT_FALSE(inspector.has_error()) << inspector.error();
ASSERT_EQ(2u, result.size());
- EXPECT_EQ(ResourceBinding::ResourceType::kSampler, result[0].resource_type);
- EXPECT_EQ(0u, result[0].bind_group);
- EXPECT_EQ(0u, result[0].binding);
+ EXPECT_EQ(ResourceBinding::ResourceType::kSampler, result[1].resource_type);
+ EXPECT_EQ(0u, result[1].bind_group);
+ EXPECT_EQ(0u, result[1].binding);
}
TEST_F(InspectorGetResourceBindingsTest, Sampler_Comparison) {
@@ -2856,9 +2842,9 @@
ASSERT_FALSE(inspector.has_error()) << inspector.error();
ASSERT_EQ(2u, result.size());
- EXPECT_EQ(ResourceBinding::ResourceType::kComparisonSampler, result[0].resource_type);
- EXPECT_EQ(0u, result[0].bind_group);
- EXPECT_EQ(0u, result[0].binding);
+ EXPECT_EQ(ResourceBinding::ResourceType::kComparisonSampler, result[1].resource_type);
+ EXPECT_EQ(0u, result[1].bind_group);
+ EXPECT_EQ(0u, result[1].binding);
}
TEST_P(InspectorGetResourceBindingsTest_WithSampledTextureParams, TextureSample) {
@@ -2880,7 +2866,6 @@
auto result = inspector.GetResourceBindings("ep");
ASSERT_FALSE(inspector.has_error()) << inspector.error();
ASSERT_EQ(2u, result.size());
- SortResourceBindings(&result);
EXPECT_EQ(ResourceBinding::ResourceType::kSampledTexture, result[0].resource_type);
EXPECT_EQ(0u, result[0].bind_group);
@@ -2925,7 +2910,6 @@
auto result = inspector.GetResourceBindings("ep");
ASSERT_FALSE(inspector.has_error()) << inspector.error();
ASSERT_EQ(2u, result.size());
- SortResourceBindings(&result);
EXPECT_EQ(ResourceBinding::ResourceType::kSampledTexture, result[0].resource_type);
EXPECT_EQ(0u, result[0].bind_group);
diff --git a/src/tint/lang/wgsl/inspector/resource_binding.h b/src/tint/lang/wgsl/inspector/resource_binding.h
index a62819b..be26cdc 100644
--- a/src/tint/lang/wgsl/inspector/resource_binding.h
+++ b/src/tint/lang/wgsl/inspector/resource_binding.h
@@ -109,7 +109,7 @@
/// Identifier to identify this binding within the bind group
uint32_t binding;
/// Input attachment index. Only available for input attachments.
- uint32_t input_attachmnt_index;
+ uint32_t input_attachment_index;
/// Size for this binding, in bytes, if defined.
uint64_t size;
/// Size for this binding without trailing structure padding, in bytes, if
diff --git a/src/tint/lang/wgsl/sem/function.cc b/src/tint/lang/wgsl/sem/function.cc
index e7ad174..14ffa8a 100644
--- a/src/tint/lang/wgsl/sem/function.cc
+++ b/src/tint/lang/wgsl/sem/function.cc
@@ -73,36 +73,6 @@
}
}
-Function::VariableBindings Function::TransitivelyReferencedUniformVariables() const {
- VariableBindings ret;
-
- for (auto* global : TransitivelyReferencedGlobals()) {
- if (global->AddressSpace() != core::AddressSpace::kUniform) {
- continue;
- }
-
- if (auto bp = global->Attributes().binding_point) {
- ret.push_back({global, *bp});
- }
- }
- return ret;
-}
-
-Function::VariableBindings Function::TransitivelyReferencedStorageBufferVariables() const {
- VariableBindings ret;
-
- for (auto* global : TransitivelyReferencedGlobals()) {
- if (global->AddressSpace() != core::AddressSpace::kStorage) {
- continue;
- }
-
- if (auto bp = global->Attributes().binding_point) {
- ret.push_back({global, *bp});
- }
- }
- return ret;
-}
-
std::vector<std::pair<const Variable*, const ast::BuiltinAttribute*>>
Function::TransitivelyReferencedBuiltinVariables() const {
std::vector<std::pair<const Variable*, const ast::BuiltinAttribute*>> ret;
@@ -118,36 +88,6 @@
return ret;
}
-Function::VariableBindings Function::TransitivelyReferencedSamplerVariables() const {
- return TransitivelyReferencedSamplerVariablesImpl(core::type::SamplerKind::kSampler);
-}
-
-Function::VariableBindings Function::TransitivelyReferencedComparisonSamplerVariables() const {
- return TransitivelyReferencedSamplerVariablesImpl(core::type::SamplerKind::kComparisonSampler);
-}
-
-Function::VariableBindings Function::TransitivelyReferencedSampledTextureVariables() const {
- return TransitivelyReferencedSampledTextureVariablesImpl(false);
-}
-
-Function::VariableBindings Function::TransitivelyReferencedMultisampledTextureVariables() const {
- return TransitivelyReferencedSampledTextureVariablesImpl(true);
-}
-
-Function::VariableBindings Function::TransitivelyReferencedVariablesOfType(
- const tint::TypeInfo* type) const {
- VariableBindings ret;
- for (auto* global : TransitivelyReferencedGlobals()) {
- auto* unwrapped_type = global->Type()->UnwrapRef();
- if (unwrapped_type->TypeInfo().Is(type)) {
- if (auto bp = global->Attributes().binding_point) {
- ret.push_back({global, *bp});
- }
- }
- }
- return ret;
-}
-
bool Function::HasAncestorEntryPoint(Symbol symbol) const {
for (const auto* point : ancestor_entry_points_) {
if (point->Declaration()->name->symbol == symbol) {
@@ -157,50 +97,6 @@
return false;
}
-Function::VariableBindings Function::TransitivelyReferencedSamplerVariablesImpl(
- core::type::SamplerKind kind) const {
- VariableBindings ret;
-
- for (auto* global : TransitivelyReferencedGlobals()) {
- auto* unwrapped_type = global->Type()->UnwrapRef();
- auto* sampler = unwrapped_type->As<core::type::Sampler>();
- if (sampler == nullptr || sampler->Kind() != kind) {
- continue;
- }
-
- if (auto bp = global->Attributes().binding_point) {
- ret.push_back({global, *bp});
- }
- }
- return ret;
-}
-
-Function::VariableBindings Function::TransitivelyReferencedSampledTextureVariablesImpl(
- bool multisampled) const {
- VariableBindings ret;
-
- for (auto* global : TransitivelyReferencedGlobals()) {
- auto* unwrapped_type = global->Type()->UnwrapRef();
- auto* texture = unwrapped_type->As<core::type::Texture>();
- if (texture == nullptr) {
- continue;
- }
-
- auto is_multisampled = texture->Is<core::type::MultisampledTexture>();
- auto is_sampled = texture->Is<core::type::SampledTexture>();
-
- if ((multisampled && !is_multisampled) || (!multisampled && !is_sampled)) {
- continue;
- }
-
- if (auto bp = global->Attributes().binding_point) {
- ret.push_back({global, *bp});
- }
- }
-
- return ret;
-}
-
void Function::SetDiagnosticSeverity(wgsl::DiagnosticRule rule, wgsl::DiagnosticSeverity severity) {
diagnostic_severities_.Add(rule, severity);
}
diff --git a/src/tint/lang/wgsl/sem/function.h b/src/tint/lang/wgsl/sem/function.h
index 1737dbe..646529e 100644
--- a/src/tint/lang/wgsl/sem/function.h
+++ b/src/tint/lang/wgsl/sem/function.h
@@ -63,9 +63,6 @@
/// Function holds the semantic information for function nodes.
class Function final : public Castable<Function, CallTarget> {
public:
- /// A vector of [Variable*, BindingPoint] pairs
- using VariableBindings = std::vector<std::pair<const Variable*, BindingPoint>>;
-
/// Constructor
/// @param declaration the ast::Function
explicit Function(const ast::Function* declaration);
@@ -213,50 +210,6 @@
std::vector<std::pair<const Variable*, const ast::BuiltinAttribute*>>
TransitivelyReferencedBuiltinVariables() const;
- /// Retrieves any referenced uniform variables. Note, the variables must be
- /// decorated with both binding and group attributes.
- /// @returns the referenced uniforms
- VariableBindings TransitivelyReferencedUniformVariables() const;
-
- /// Retrieves any referenced storagebuffer variables. Note, the variables
- /// must be decorated with both binding and group attributes.
- /// @returns the referenced storagebuffers
- VariableBindings TransitivelyReferencedStorageBufferVariables() const;
-
- /// Retrieves any referenced regular Sampler variables. Note, the
- /// variables must be decorated with both binding and group attributes.
- /// @returns the referenced storagebuffers
- VariableBindings TransitivelyReferencedSamplerVariables() const;
-
- /// Retrieves any referenced comparison Sampler variables. Note, the
- /// variables must be decorated with both binding and group attributes.
- /// @returns the referenced storagebuffers
- VariableBindings TransitivelyReferencedComparisonSamplerVariables() const;
-
- /// Retrieves any referenced sampled textures variables. Note, the
- /// variables must be decorated with both binding and group attributes.
- /// @returns the referenced sampled textures
- VariableBindings TransitivelyReferencedSampledTextureVariables() const;
-
- /// Retrieves any referenced multisampled textures variables. Note, the
- /// variables must be decorated with both binding and group attributes.
- /// @returns the referenced sampled textures
- VariableBindings TransitivelyReferencedMultisampledTextureVariables() const;
-
- /// Retrieves any referenced variables of the given type. Note, the variables
- /// must be decorated with both binding and group attributes.
- /// @param type the type of the variables to find
- /// @returns the referenced variables
- VariableBindings TransitivelyReferencedVariablesOfType(const tint::TypeInfo* type) const;
-
- /// Retrieves any referenced variables of the given type. Note, the variables
- /// must be decorated with both binding and group attributes.
- /// @returns the referenced variables
- template <typename T>
- VariableBindings TransitivelyReferencedVariablesOfType() const {
- return TransitivelyReferencedVariablesOfType(&tint::TypeInfo::Of<T>());
- }
-
/// Checks if the given entry point is an ancestor
/// @param sym the entry point symbol
/// @returns true if `sym` is an ancestor entry point of this function
@@ -300,9 +253,6 @@
Function(const Function&) = delete;
Function(Function&&) = delete;
- VariableBindings TransitivelyReferencedSamplerVariablesImpl(core::type::SamplerKind kind) const;
- VariableBindings TransitivelyReferencedSampledTextureVariablesImpl(bool multisampled) const;
-
const ast::Function* const declaration_;
sem::WorkgroupSize workgroup_size_;
diff --git a/test/tint/array/assign_to_storage_var.wgsl.expected.msl b/test/tint/array/assign_to_storage_var.wgsl.expected.msl
index 914a238..1dd571f 100644
--- a/test/tint/array/assign_to_storage_var.wgsl.expected.msl
+++ b/test/tint/array/assign_to_storage_var.wgsl.expected.msl
@@ -80,7 +80,7 @@
foo(ary, tint_module_vars);
}
-kernel void v_2(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_1* v_3 [[threadgroup(0)]], const constant S* src_uniform [[buffer(0)]], device S* src_storage [[buffer(2)]], device S* dst [[buffer(1)]], device S_nested* dst_nested [[buffer(3)]]) {
+kernel void v_2(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_1* v_3 [[threadgroup(0)]], const constant S* src_uniform [[buffer(1)]], device S* src_storage [[buffer(2)]], device S* dst [[buffer(0)]], device S_nested* dst_nested [[buffer(3)]]) {
thread tint_array<int4, 4> src_private = {};
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.src_private=(&src_private), .src_workgroup=(&(*v_3).tint_symbol), .src_uniform=src_uniform, .src_storage=src_storage, .dst=dst, .dst_nested=dst_nested};
main_inner(tint_local_index, tint_module_vars);
diff --git a/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.msl b/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.msl
index e5081c1..94fc405 100644
--- a/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.msl
+++ b/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.msl
@@ -113,7 +113,7 @@
(*tint_module_vars.s) = as_type<int>((as_type<uint>(v_19) + as_type<uint>(tint_f32_to_i32(arr2_vec3_f32[0u].x))));
}
-kernel void v_20(uint idx [[thread_index_in_threadgroup]], const device S_packed_vec3* sb [[buffer(1)]], device int* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
+kernel void v_20(uint idx [[thread_index_in_threadgroup]], const device S_packed_vec3* sb [[buffer(0)]], device int* s [[buffer(1)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.sb=sb, .s=s, .tint_storage_buffer_sizes=tint_storage_buffer_sizes};
main_inner(idx, tint_module_vars);
}
diff --git a/test/tint/buffer/storage/dynamic_index/read_f16.wgsl.expected.msl b/test/tint/buffer/storage/dynamic_index/read_f16.wgsl.expected.msl
index 23c0822..aeedcc9 100644
--- a/test/tint/buffer/storage/dynamic_index/read_f16.wgsl.expected.msl
+++ b/test/tint/buffer/storage/dynamic_index/read_f16.wgsl.expected.msl
@@ -171,7 +171,7 @@
(*tint_module_vars.s) = as_type<int>((as_type<uint>(v_36) + as_type<uint>(tint_f32_to_i32(arr2_vec3_f32[0u].x))));
}
-kernel void v_37(uint idx [[thread_index_in_threadgroup]], const device S_packed_vec3* sb [[buffer(1)]], device int* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
+kernel void v_37(uint idx [[thread_index_in_threadgroup]], const device S_packed_vec3* sb [[buffer(0)]], device int* s [[buffer(1)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.sb=sb, .s=s, .tint_storage_buffer_sizes=tint_storage_buffer_sizes};
main_inner(idx, tint_module_vars);
}
diff --git a/test/tint/buffer/storage/static_index/read.wgsl.expected.msl b/test/tint/buffer/storage/static_index/read.wgsl.expected.msl
index 96a7ac0..591a442 100644
--- a/test/tint/buffer/storage/static_index/read.wgsl.expected.msl
+++ b/test/tint/buffer/storage/static_index/read.wgsl.expected.msl
@@ -70,7 +70,7 @@
return tint_array<float3, 2>{float3((*from)[0u].packed), float3((*from)[1u].packed)};
}
-kernel void v(const device S_packed_vec3* sb [[buffer(1)]], device int* s [[buffer(0)]]) {
+kernel void v(const device S_packed_vec3* sb [[buffer(0)]], device int* s [[buffer(1)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.sb=sb, .s=s};
float const scalar_f32 = (*tint_module_vars.sb).scalar_f32;
int const scalar_i32 = (*tint_module_vars.sb).scalar_i32;
diff --git a/test/tint/buffer/storage/static_index/read_f16.wgsl.expected.msl b/test/tint/buffer/storage/static_index/read_f16.wgsl.expected.msl
index 2029d1d..84755ae 100644
--- a/test/tint/buffer/storage/static_index/read_f16.wgsl.expected.msl
+++ b/test/tint/buffer/storage/static_index/read_f16.wgsl.expected.msl
@@ -99,7 +99,7 @@
return tint_array<float3, 2>{float3((*from)[0u].packed), float3((*from)[1u].packed)};
}
-kernel void v(const device S_packed_vec3* sb [[buffer(1)]], device int* s [[buffer(0)]]) {
+kernel void v(const device S_packed_vec3* sb [[buffer(0)]], device int* s [[buffer(1)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.sb=sb, .s=s};
float const scalar_f32 = (*tint_module_vars.sb).scalar_f32;
int const scalar_i32 = (*tint_module_vars.sb).scalar_i32;
diff --git a/test/tint/buffer/storage/types/struct_f16.wgsl.expected.msl b/test/tint/buffer/storage/types/struct_f16.wgsl.expected.msl
index cfc0b0a..6333788 100644
--- a/test/tint/buffer/storage/types/struct_f16.wgsl.expected.msl
+++ b/test/tint/buffer/storage/types/struct_f16.wgsl.expected.msl
@@ -58,7 +58,7 @@
return S{.inner=tint_load_struct_packed_vec3_1((&(*from).inner))};
}
-kernel void v(const device S_packed_vec3* in [[buffer(1)]], device S_packed_vec3* out [[buffer(0)]]) {
+kernel void v(const device S_packed_vec3* in [[buffer(0)]], device S_packed_vec3* out [[buffer(1)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.in=in, .out=out};
S const t = tint_load_struct_packed_vec3(tint_module_vars.in);
tint_store_and_preserve_padding(tint_module_vars.out, t);
diff --git a/test/tint/buffer/storage/types/struct_f32.wgsl.expected.msl b/test/tint/buffer/storage/types/struct_f32.wgsl.expected.msl
index 8e12400..c9c37fc 100644
--- a/test/tint/buffer/storage/types/struct_f32.wgsl.expected.msl
+++ b/test/tint/buffer/storage/types/struct_f32.wgsl.expected.msl
@@ -58,7 +58,7 @@
return S{.inner=tint_load_struct_packed_vec3_1((&(*from).inner))};
}
-kernel void v(const device S_packed_vec3* in [[buffer(1)]], device S_packed_vec3* out [[buffer(0)]]) {
+kernel void v(const device S_packed_vec3* in [[buffer(0)]], device S_packed_vec3* out [[buffer(1)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.in=in, .out=out};
S const t = tint_load_struct_packed_vec3(tint_module_vars.in);
tint_store_and_preserve_padding(tint_module_vars.out, t);
diff --git a/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_fn.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_fn.wgsl.expected.msl
index 408809b..59f8ff2 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_fn.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_fn.wgsl.expected.msl
@@ -34,7 +34,7 @@
return f_1;
}
-kernel void f(const constant tint_array<float2x2, 4>* u [[buffer(0)]], device float* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<float2x2, 4>* u [[buffer(1)]], device float* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
float const v_1 = a((*tint_module_vars.u));
float const v_2 = (v_1 + b((*tint_module_vars.u)[1u]));
diff --git a/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_storage.wgsl.expected.msl
index ad55014..543c191 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_storage.wgsl.expected.msl
@@ -18,7 +18,7 @@
device tint_array<float2x2, 4>* s;
};
-kernel void f(const constant tint_array<float2x2, 4>* u [[buffer(0)]], device tint_array<float2x2, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<float2x2, 4>* u [[buffer(1)]], device tint_array<float2x2, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
(*tint_module_vars.s) = (*tint_module_vars.u);
(*tint_module_vars.s)[1u] = (*tint_module_vars.u)[2u];
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_fn.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_fn.wgsl.expected.msl
index 05e2ca2..00333e3 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_fn.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_fn.wgsl.expected.msl
@@ -50,7 +50,7 @@
return tint_array<half2x3, 4>{v_2, v_4, v_6, half2x3(half3(v_7[0u].packed), half3(v_7[1u].packed))};
}
-kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f16_array_element, 2>, 4>* u [[buffer(0)]], device half* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f16_array_element, 2>, 4>* u [[buffer(1)]], device half* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
half const v_8 = a(tint_load_array_packed_vec3(tint_module_vars.u));
tint_array<tint_packed_vec3_f16_array_element, 2> const v_9 = (*tint_module_vars.u)[1u];
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.msl
index 4a8cdd8..25d2b85 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.msl
@@ -57,7 +57,7 @@
return tint_array<half2x3, 4>{v_3, v_5, v_7, half2x3(half3(v_8[0u].packed), half3(v_8[1u].packed))};
}
-kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f16_array_element, 2>, 4>* u [[buffer(0)]], device tint_array<tint_array<tint_packed_vec3_f16_array_element, 2>, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f16_array_element, 2>, 4>* u [[buffer(1)]], device tint_array<tint_array<tint_packed_vec3_f16_array_element, 2>, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
tint_store_and_preserve_padding(tint_module_vars.s, tint_load_array_packed_vec3(tint_module_vars.u));
tint_array<tint_packed_vec3_f16_array_element, 2> const v_9 = (*tint_module_vars.u)[2u];
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_fn.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_fn.wgsl.expected.msl
index 3bfdc5c..6229846 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_fn.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_fn.wgsl.expected.msl
@@ -50,7 +50,7 @@
return tint_array<float2x3, 4>{v_2, v_4, v_6, float2x3(float3(v_7[0u].packed), float3(v_7[1u].packed))};
}
-kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 2>, 4>* u [[buffer(0)]], device float* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 2>, 4>* u [[buffer(1)]], device float* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
float const v_8 = a(tint_load_array_packed_vec3(tint_module_vars.u));
tint_array<tint_packed_vec3_f32_array_element, 2> const v_9 = (*tint_module_vars.u)[1u];
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.msl
index d26bcbb..9211b80 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.msl
@@ -57,7 +57,7 @@
return tint_array<float2x3, 4>{v_3, v_5, v_7, float2x3(float3(v_8[0u].packed), float3(v_8[1u].packed))};
}
-kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 2>, 4>* u [[buffer(0)]], device tint_array<tint_array<tint_packed_vec3_f32_array_element, 2>, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 2>, 4>* u [[buffer(1)]], device tint_array<tint_array<tint_packed_vec3_f32_array_element, 2>, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
tint_store_and_preserve_padding(tint_module_vars.s, tint_load_array_packed_vec3(tint_module_vars.u));
tint_array<tint_packed_vec3_f32_array_element, 2> const v_9 = (*tint_module_vars.u)[2u];
diff --git a/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_fn.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_fn.wgsl.expected.msl
index cd09bcc..197396c 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_fn.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_fn.wgsl.expected.msl
@@ -34,7 +34,7 @@
return f_1;
}
-kernel void f(const constant tint_array<half2x4, 4>* u [[buffer(0)]], device half* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<half2x4, 4>* u [[buffer(1)]], device half* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
half const v_1 = a((*tint_module_vars.u));
half const v_2 = (v_1 + b((*tint_module_vars.u)[1u]));
diff --git a/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_storage.wgsl.expected.msl
index 16daafc..60178d6 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_storage.wgsl.expected.msl
@@ -18,7 +18,7 @@
device tint_array<half2x4, 4>* s;
};
-kernel void f(const constant tint_array<half2x4, 4>* u [[buffer(0)]], device tint_array<half2x4, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<half2x4, 4>* u [[buffer(1)]], device tint_array<half2x4, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
(*tint_module_vars.s) = (*tint_module_vars.u);
(*tint_module_vars.s)[1u] = (*tint_module_vars.u)[2u];
diff --git a/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_fn.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_fn.wgsl.expected.msl
index 7fd6b3b..479edbc 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_fn.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_fn.wgsl.expected.msl
@@ -34,7 +34,7 @@
return f_1;
}
-kernel void f(const constant tint_array<float2x4, 4>* u [[buffer(0)]], device float* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<float2x4, 4>* u [[buffer(1)]], device float* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
float const v_1 = a((*tint_module_vars.u));
float const v_2 = (v_1 + b((*tint_module_vars.u)[1u]));
diff --git a/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_storage.wgsl.expected.msl
index 288e8b1..f0a7ee8 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_storage.wgsl.expected.msl
@@ -18,7 +18,7 @@
device tint_array<float2x4, 4>* s;
};
-kernel void f(const constant tint_array<float2x4, 4>* u [[buffer(0)]], device tint_array<float2x4, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<float2x4, 4>* u [[buffer(1)]], device tint_array<float2x4, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
(*tint_module_vars.s) = (*tint_module_vars.u);
(*tint_module_vars.s)[1u] = (*tint_module_vars.u)[2u];
diff --git a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_fn.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_fn.wgsl.expected.msl
index 8687dc5..0aa59c2 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_fn.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_fn.wgsl.expected.msl
@@ -50,7 +50,7 @@
return tint_array<float3x3, 4>{v_2, v_4, v_6, float3x3(float3(v_7[0u].packed), float3(v_7[1u].packed), float3(v_7[2u].packed))};
}
-kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 4>* u [[buffer(0)]], device float* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 4>* u [[buffer(1)]], device float* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
float const v_8 = a(tint_load_array_packed_vec3(tint_module_vars.u));
tint_array<tint_packed_vec3_f32_array_element, 3> const v_9 = (*tint_module_vars.u)[1u];
diff --git a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.msl
index 15903ff..d0bb699 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.msl
@@ -58,7 +58,7 @@
return tint_array<float3x3, 4>{v_3, v_5, v_7, float3x3(float3(v_8[0u].packed), float3(v_8[1u].packed), float3(v_8[2u].packed))};
}
-kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 4>* u [[buffer(0)]], device tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 4>* u [[buffer(1)]], device tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
tint_store_and_preserve_padding(tint_module_vars.s, tint_load_array_packed_vec3(tint_module_vars.u));
tint_array<tint_packed_vec3_f32_array_element, 3> const v_9 = (*tint_module_vars.u)[2u];
diff --git a/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_fn.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_fn.wgsl.expected.msl
index ab08384..d89f461 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_fn.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_fn.wgsl.expected.msl
@@ -34,7 +34,7 @@
return f_1;
}
-kernel void f(const constant tint_array<float3x4, 4>* u [[buffer(0)]], device float* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<float3x4, 4>* u [[buffer(1)]], device float* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
float const v_1 = a((*tint_module_vars.u));
float const v_2 = (v_1 + b((*tint_module_vars.u)[1u]));
diff --git a/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_storage.wgsl.expected.msl
index ceffe27..e3ec4ff 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_storage.wgsl.expected.msl
@@ -18,7 +18,7 @@
device tint_array<float3x4, 4>* s;
};
-kernel void f(const constant tint_array<float3x4, 4>* u [[buffer(0)]], device tint_array<float3x4, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<float3x4, 4>* u [[buffer(1)]], device tint_array<float3x4, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
(*tint_module_vars.s) = (*tint_module_vars.u);
(*tint_module_vars.s)[1u] = (*tint_module_vars.u)[2u];
diff --git a/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_fn.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_fn.wgsl.expected.msl
index f86bd5d..f9fbd3d 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_fn.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_fn.wgsl.expected.msl
@@ -34,7 +34,7 @@
return f_1;
}
-kernel void f(const constant tint_array<half4x2, 4>* u [[buffer(0)]], device half* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<half4x2, 4>* u [[buffer(1)]], device half* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
half const v_1 = a((*tint_module_vars.u));
half const v_2 = (v_1 + b((*tint_module_vars.u)[1u]));
diff --git a/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_storage.wgsl.expected.msl
index 6d4e3d1..94c67a4 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_storage.wgsl.expected.msl
@@ -18,7 +18,7 @@
device tint_array<half4x2, 4>* s;
};
-kernel void f(const constant tint_array<half4x2, 4>* u [[buffer(0)]], device tint_array<half4x2, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<half4x2, 4>* u [[buffer(1)]], device tint_array<half4x2, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
(*tint_module_vars.s) = (*tint_module_vars.u);
(*tint_module_vars.s)[1u] = (*tint_module_vars.u)[2u];
diff --git a/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_fn.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_fn.wgsl.expected.msl
index 126e347..eece13e 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_fn.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_fn.wgsl.expected.msl
@@ -34,7 +34,7 @@
return f_1;
}
-kernel void f(const constant tint_array<float4x2, 4>* u [[buffer(0)]], device float* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<float4x2, 4>* u [[buffer(1)]], device float* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
float const v_1 = a((*tint_module_vars.u));
float const v_2 = (v_1 + b((*tint_module_vars.u)[1u]));
diff --git a/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_storage.wgsl.expected.msl
index 61f49f0..f1477ab 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_storage.wgsl.expected.msl
@@ -18,7 +18,7 @@
device tint_array<float4x2, 4>* s;
};
-kernel void f(const constant tint_array<float4x2, 4>* u [[buffer(0)]], device tint_array<float4x2, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<float4x2, 4>* u [[buffer(1)]], device tint_array<float4x2, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
(*tint_module_vars.s) = (*tint_module_vars.u);
(*tint_module_vars.s)[1u] = (*tint_module_vars.u)[2u];
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_fn.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_fn.wgsl.expected.msl
index bd05237..7436493 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_fn.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_fn.wgsl.expected.msl
@@ -50,7 +50,7 @@
return tint_array<half4x3, 4>{v_2, v_4, v_6, half4x3(half3(v_7[0u].packed), half3(v_7[1u].packed), half3(v_7[2u].packed), half3(v_7[3u].packed))};
}
-kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f16_array_element, 4>, 4>* u [[buffer(0)]], device half* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f16_array_element, 4>, 4>* u [[buffer(1)]], device half* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
half const v_8 = a(tint_load_array_packed_vec3(tint_module_vars.u));
tint_array<tint_packed_vec3_f16_array_element, 4> const v_9 = (*tint_module_vars.u)[1u];
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.msl
index f02a3de..5c3b827 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.msl
@@ -59,7 +59,7 @@
return tint_array<half4x3, 4>{v_3, v_5, v_7, half4x3(half3(v_8[0u].packed), half3(v_8[1u].packed), half3(v_8[2u].packed), half3(v_8[3u].packed))};
}
-kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f16_array_element, 4>, 4>* u [[buffer(0)]], device tint_array<tint_array<tint_packed_vec3_f16_array_element, 4>, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f16_array_element, 4>, 4>* u [[buffer(1)]], device tint_array<tint_array<tint_packed_vec3_f16_array_element, 4>, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
tint_store_and_preserve_padding(tint_module_vars.s, tint_load_array_packed_vec3(tint_module_vars.u));
tint_array<tint_packed_vec3_f16_array_element, 4> const v_9 = (*tint_module_vars.u)[2u];
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_fn.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_fn.wgsl.expected.msl
index 8ce6622..92bad27 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_fn.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_fn.wgsl.expected.msl
@@ -50,7 +50,7 @@
return tint_array<float4x3, 4>{v_2, v_4, v_6, float4x3(float3(v_7[0u].packed), float3(v_7[1u].packed), float3(v_7[2u].packed), float3(v_7[3u].packed))};
}
-kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 4>, 4>* u [[buffer(0)]], device float* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 4>, 4>* u [[buffer(1)]], device float* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
float const v_8 = a(tint_load_array_packed_vec3(tint_module_vars.u));
tint_array<tint_packed_vec3_f32_array_element, 4> const v_9 = (*tint_module_vars.u)[1u];
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.msl
index a02df21..becd496 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.msl
@@ -59,7 +59,7 @@
return tint_array<float4x3, 4>{v_3, v_5, v_7, float4x3(float3(v_8[0u].packed), float3(v_8[1u].packed), float3(v_8[2u].packed), float3(v_8[3u].packed))};
}
-kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 4>, 4>* u [[buffer(0)]], device tint_array<tint_array<tint_packed_vec3_f32_array_element, 4>, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 4>, 4>* u [[buffer(1)]], device tint_array<tint_array<tint_packed_vec3_f32_array_element, 4>, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
tint_store_and_preserve_padding(tint_module_vars.s, tint_load_array_packed_vec3(tint_module_vars.u));
tint_array<tint_packed_vec3_f32_array_element, 4> const v_9 = (*tint_module_vars.u)[2u];
diff --git a/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_fn.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_fn.wgsl.expected.msl
index 5034b53..77ab440 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_fn.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_fn.wgsl.expected.msl
@@ -34,7 +34,7 @@
return f_1;
}
-kernel void f(const constant tint_array<half4x4, 4>* u [[buffer(0)]], device half* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<half4x4, 4>* u [[buffer(1)]], device half* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
half const v_1 = a((*tint_module_vars.u));
half const v_2 = (v_1 + b((*tint_module_vars.u)[1u]));
diff --git a/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_storage.wgsl.expected.msl
index ae57bd1..bab5fc8 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_storage.wgsl.expected.msl
@@ -18,7 +18,7 @@
device tint_array<half4x4, 4>* s;
};
-kernel void f(const constant tint_array<half4x4, 4>* u [[buffer(0)]], device tint_array<half4x4, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<half4x4, 4>* u [[buffer(1)]], device tint_array<half4x4, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
(*tint_module_vars.s) = (*tint_module_vars.u);
(*tint_module_vars.s)[1u] = (*tint_module_vars.u)[2u];
diff --git a/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_fn.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_fn.wgsl.expected.msl
index 6ac6867..1844372 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_fn.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_fn.wgsl.expected.msl
@@ -34,7 +34,7 @@
return f_1;
}
-kernel void f(const constant tint_array<float4x4, 4>* u [[buffer(0)]], device float* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<float4x4, 4>* u [[buffer(1)]], device float* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
float const v_1 = a((*tint_module_vars.u));
float const v_2 = (v_1 + b((*tint_module_vars.u)[1u]));
diff --git a/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_storage.wgsl.expected.msl
index a8935af..e21f52c 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_storage.wgsl.expected.msl
@@ -18,7 +18,7 @@
device tint_array<float4x4, 4>* s;
};
-kernel void f(const constant tint_array<float4x4, 4>* u [[buffer(0)]], device tint_array<float4x4, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<float4x4, 4>* u [[buffer(1)]], device tint_array<float4x4, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
(*tint_module_vars.s) = (*tint_module_vars.u);
(*tint_module_vars.s)[1u] = (*tint_module_vars.u)[2u];
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.msl
index f7c76c5..c506a89 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.msl
@@ -50,7 +50,7 @@
}
}
-kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]], device tint_array<S, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<S, 4>* u [[buffer(1)]], device tint_array<S, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
tint_store_and_preserve_padding(tint_module_vars.s, (*tint_module_vars.u));
tint_store_and_preserve_padding_1((&(*tint_module_vars.s)[1u]), (*tint_module_vars.u)[2u]);
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.msl
index b182233..b4873dd 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.msl
@@ -51,7 +51,7 @@
}
}
-kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]], device tint_array<S, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<S, 4>* u [[buffer(1)]], device tint_array<S, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
tint_store_and_preserve_padding(tint_module_vars.s, (*tint_module_vars.u));
tint_store_and_preserve_padding_1((&(*tint_module_vars.s)[1u]), (*tint_module_vars.u)[2u]);
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.msl
index d083fe4..6922711 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.msl
@@ -81,7 +81,7 @@
return tint_array<S, 4>{v_5, v_6, v_7, tint_load_struct_packed_vec3((&(*from)[3u]))};
}
-kernel void f(const constant tint_array<S_packed_vec3, 4>* u [[buffer(0)]], device tint_array<S_packed_vec3, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<S_packed_vec3, 4>* u [[buffer(1)]], device tint_array<S_packed_vec3, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
tint_store_and_preserve_padding(tint_module_vars.s, tint_load_array_packed_vec3(tint_module_vars.u));
tint_store_and_preserve_padding_1((&(*tint_module_vars.s)[1u]), tint_load_struct_packed_vec3((&(*tint_module_vars.u)[2u])));
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.msl
index 6d4708b..d803efb 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.msl
@@ -81,7 +81,7 @@
return tint_array<S, 4>{v_5, v_6, v_7, tint_load_struct_packed_vec3((&(*from)[3u]))};
}
-kernel void f(const constant tint_array<S_packed_vec3, 4>* u [[buffer(0)]], device tint_array<S_packed_vec3, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<S_packed_vec3, 4>* u [[buffer(1)]], device tint_array<S_packed_vec3, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
tint_store_and_preserve_padding(tint_module_vars.s, tint_load_array_packed_vec3(tint_module_vars.u));
tint_store_and_preserve_padding_1((&(*tint_module_vars.s)[1u]), tint_load_struct_packed_vec3((&(*tint_module_vars.u)[2u])));
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.msl
index d605359..222192e 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.msl
@@ -51,7 +51,7 @@
}
}
-kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]], device tint_array<S, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<S, 4>* u [[buffer(1)]], device tint_array<S, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
tint_store_and_preserve_padding(tint_module_vars.s, (*tint_module_vars.u));
tint_store_and_preserve_padding_1((&(*tint_module_vars.s)[1u]), (*tint_module_vars.u)[2u]);
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.msl
index 5738334..6b18fc3 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.msl
@@ -51,7 +51,7 @@
}
}
-kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]], device tint_array<S, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<S, 4>* u [[buffer(1)]], device tint_array<S, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
tint_store_and_preserve_padding(tint_module_vars.s, (*tint_module_vars.u));
tint_store_and_preserve_padding_1((&(*tint_module_vars.s)[1u]), (*tint_module_vars.u)[2u]);
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.msl
index 371f30d..4b681a8 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.msl
@@ -50,7 +50,7 @@
}
}
-kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]], device tint_array<S, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<S, 4>* u [[buffer(1)]], device tint_array<S, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
tint_store_and_preserve_padding(tint_module_vars.s, (*tint_module_vars.u));
tint_store_and_preserve_padding_1((&(*tint_module_vars.s)[1u]), (*tint_module_vars.u)[2u]);
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.msl
index 10dd8fe..9bbb5ce 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.msl
@@ -51,7 +51,7 @@
}
}
-kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]], device tint_array<S, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<S, 4>* u [[buffer(1)]], device tint_array<S, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
tint_store_and_preserve_padding(tint_module_vars.s, (*tint_module_vars.u));
tint_store_and_preserve_padding_1((&(*tint_module_vars.s)[1u]), (*tint_module_vars.u)[2u]);
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.msl
index bffb357..a84320e 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.msl
@@ -82,7 +82,7 @@
return tint_array<S, 4>{v_5, v_6, v_7, tint_load_struct_packed_vec3((&(*from)[3u]))};
}
-kernel void f(const constant tint_array<S_packed_vec3, 4>* u [[buffer(0)]], device tint_array<S_packed_vec3, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<S_packed_vec3, 4>* u [[buffer(1)]], device tint_array<S_packed_vec3, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
tint_store_and_preserve_padding(tint_module_vars.s, tint_load_array_packed_vec3(tint_module_vars.u));
tint_store_and_preserve_padding_1((&(*tint_module_vars.s)[1u]), tint_load_struct_packed_vec3((&(*tint_module_vars.u)[2u])));
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.msl
index b539217..c000cdb 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.msl
@@ -81,7 +81,7 @@
return tint_array<S, 4>{v_5, v_6, v_7, tint_load_struct_packed_vec3((&(*from)[3u]))};
}
-kernel void f(const constant tint_array<S_packed_vec3, 4>* u [[buffer(0)]], device tint_array<S_packed_vec3, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<S_packed_vec3, 4>* u [[buffer(1)]], device tint_array<S_packed_vec3, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
tint_store_and_preserve_padding(tint_module_vars.s, tint_load_array_packed_vec3(tint_module_vars.u));
tint_store_and_preserve_padding_1((&(*tint_module_vars.s)[1u]), tint_load_struct_packed_vec3((&(*tint_module_vars.u)[2u])));
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.msl
index 0b63816..5ccd2b6 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.msl
@@ -51,7 +51,7 @@
}
}
-kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]], device tint_array<S, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<S, 4>* u [[buffer(1)]], device tint_array<S, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
tint_store_and_preserve_padding(tint_module_vars.s, (*tint_module_vars.u));
tint_store_and_preserve_padding_1((&(*tint_module_vars.s)[1u]), (*tint_module_vars.u)[2u]);
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.msl
index c18e93c..f73f1a5 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.msl
@@ -50,7 +50,7 @@
}
}
-kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]], device tint_array<S, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<S, 4>* u [[buffer(1)]], device tint_array<S, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
tint_store_and_preserve_padding(tint_module_vars.s, (*tint_module_vars.u));
tint_store_and_preserve_padding_1((&(*tint_module_vars.s)[1u]), (*tint_module_vars.u)[2u]);
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.msl
index 123b5cc..f6d5f36 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.msl
@@ -50,7 +50,7 @@
}
}
-kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]], device tint_array<S, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<S, 4>* u [[buffer(1)]], device tint_array<S, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
tint_store_and_preserve_padding(tint_module_vars.s, (*tint_module_vars.u));
tint_store_and_preserve_padding_1((&(*tint_module_vars.s)[1u]), (*tint_module_vars.u)[2u]);
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.msl
index 60d1bc8..9fcc539 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.msl
@@ -51,7 +51,7 @@
}
}
-kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]], device tint_array<S, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<S, 4>* u [[buffer(1)]], device tint_array<S, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
tint_store_and_preserve_padding(tint_module_vars.s, (*tint_module_vars.u));
tint_store_and_preserve_padding_1((&(*tint_module_vars.s)[1u]), (*tint_module_vars.u)[2u]);
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.msl
index 9afbba2..3bfb11d 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.msl
@@ -83,7 +83,7 @@
return tint_array<S, 4>{v_5, v_6, v_7, tint_load_struct_packed_vec3((&(*from)[3u]))};
}
-kernel void f(const constant tint_array<S_packed_vec3, 4>* u [[buffer(0)]], device tint_array<S_packed_vec3, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<S_packed_vec3, 4>* u [[buffer(1)]], device tint_array<S_packed_vec3, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
tint_store_and_preserve_padding(tint_module_vars.s, tint_load_array_packed_vec3(tint_module_vars.u));
tint_store_and_preserve_padding_1((&(*tint_module_vars.s)[1u]), tint_load_struct_packed_vec3((&(*tint_module_vars.u)[2u])));
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.msl
index 77dc976..360c786 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.msl
@@ -83,7 +83,7 @@
return tint_array<S, 4>{v_5, v_6, v_7, tint_load_struct_packed_vec3((&(*from)[3u]))};
}
-kernel void f(const constant tint_array<S_packed_vec3, 4>* u [[buffer(0)]], device tint_array<S_packed_vec3, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<S_packed_vec3, 4>* u [[buffer(1)]], device tint_array<S_packed_vec3, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
tint_store_and_preserve_padding(tint_module_vars.s, tint_load_array_packed_vec3(tint_module_vars.u));
tint_store_and_preserve_padding_1((&(*tint_module_vars.s)[1u]), tint_load_struct_packed_vec3((&(*tint_module_vars.u)[2u])));
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.msl
index ee49f0d..e5b9104 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.msl
@@ -51,7 +51,7 @@
}
}
-kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]], device tint_array<S, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<S, 4>* u [[buffer(1)]], device tint_array<S, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
tint_store_and_preserve_padding(tint_module_vars.s, (*tint_module_vars.u));
tint_store_and_preserve_padding_1((&(*tint_module_vars.s)[1u]), (*tint_module_vars.u)[2u]);
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.msl
index a66d52d..d5255047 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.msl
@@ -51,7 +51,7 @@
}
}
-kernel void f(const constant tint_array<S, 4>* u [[buffer(0)]], device tint_array<S, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<S, 4>* u [[buffer(1)]], device tint_array<S, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
tint_store_and_preserve_padding(tint_module_vars.s, (*tint_module_vars.u));
tint_store_and_preserve_padding_1((&(*tint_module_vars.s)[1u]), (*tint_module_vars.u)[2u]);
diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x2_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/unnested/mat2x2_f16/to_storage.wgsl.expected.msl
index 8544309..7c62628 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x2_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x2_f16/to_storage.wgsl.expected.msl
@@ -6,7 +6,7 @@
device half2x2* s;
};
-kernel void f(const constant half2x2* u [[buffer(0)]], device half2x2* s [[buffer(1)]]) {
+kernel void f(const constant half2x2* u [[buffer(1)]], device half2x2* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
(*tint_module_vars.s) = (*tint_module_vars.u);
(*tint_module_vars.s)[1u] = (*tint_module_vars.u)[0u];
diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x2_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/unnested/mat2x2_f32/to_storage.wgsl.expected.msl
index 5ac2f2f..42b2a62 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x2_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x2_f32/to_storage.wgsl.expected.msl
@@ -6,7 +6,7 @@
device float2x2* s;
};
-kernel void f(const constant float2x2* u [[buffer(0)]], device float2x2* s [[buffer(1)]]) {
+kernel void f(const constant float2x2* u [[buffer(1)]], device float2x2* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
(*tint_module_vars.s) = (*tint_module_vars.u);
(*tint_module_vars.s)[1u] = (*tint_module_vars.u)[0u];
diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x3_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/unnested/mat2x3_f16/to_storage.wgsl.expected.msl
index 1929fc8..508b21f 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x3_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x3_f16/to_storage.wgsl.expected.msl
@@ -28,7 +28,7 @@
(*target)[1u].packed = packed_half3(value_param[1u]);
}
-kernel void f(const constant tint_array<tint_packed_vec3_f16_array_element, 2>* u [[buffer(0)]], device tint_array<tint_packed_vec3_f16_array_element, 2>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<tint_packed_vec3_f16_array_element, 2>* u [[buffer(1)]], device tint_array<tint_packed_vec3_f16_array_element, 2>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
tint_array<tint_packed_vec3_f16_array_element, 2> const v = (*tint_module_vars.u);
tint_store_and_preserve_padding(tint_module_vars.s, half2x3(half3(v[0u].packed), half3(v[1u].packed)));
diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x3_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/unnested/mat2x3_f32/to_storage.wgsl.expected.msl
index a78e9b1..c765428 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x3_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x3_f32/to_storage.wgsl.expected.msl
@@ -28,7 +28,7 @@
(*target)[1u].packed = packed_float3(value_param[1u]);
}
-kernel void f(const constant tint_array<tint_packed_vec3_f32_array_element, 2>* u [[buffer(0)]], device tint_array<tint_packed_vec3_f32_array_element, 2>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<tint_packed_vec3_f32_array_element, 2>* u [[buffer(1)]], device tint_array<tint_packed_vec3_f32_array_element, 2>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
tint_array<tint_packed_vec3_f32_array_element, 2> const v = (*tint_module_vars.u);
tint_store_and_preserve_padding(tint_module_vars.s, float2x3(float3(v[0u].packed), float3(v[1u].packed)));
diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x4_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/unnested/mat2x4_f16/to_storage.wgsl.expected.msl
index 4240081..2059a89 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x4_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x4_f16/to_storage.wgsl.expected.msl
@@ -6,7 +6,7 @@
device half2x4* s;
};
-kernel void f(const constant half2x4* u [[buffer(0)]], device half2x4* s [[buffer(1)]]) {
+kernel void f(const constant half2x4* u [[buffer(1)]], device half2x4* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
(*tint_module_vars.s) = (*tint_module_vars.u);
(*tint_module_vars.s)[1u] = (*tint_module_vars.u)[0u];
diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x4_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/unnested/mat2x4_f32/to_storage.wgsl.expected.msl
index 2146611..cb0a424 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x4_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x4_f32/to_storage.wgsl.expected.msl
@@ -6,7 +6,7 @@
device float2x4* s;
};
-kernel void f(const constant float2x4* u [[buffer(0)]], device float2x4* s [[buffer(1)]]) {
+kernel void f(const constant float2x4* u [[buffer(1)]], device float2x4* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
(*tint_module_vars.s) = (*tint_module_vars.u);
(*tint_module_vars.s)[1u] = (*tint_module_vars.u)[0u];
diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x2_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/unnested/mat3x2_f16/to_storage.wgsl.expected.msl
index de5386a..f9e6e99 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x2_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x2_f16/to_storage.wgsl.expected.msl
@@ -6,7 +6,7 @@
device half3x2* s;
};
-kernel void f(const constant half3x2* u [[buffer(0)]], device half3x2* s [[buffer(1)]]) {
+kernel void f(const constant half3x2* u [[buffer(1)]], device half3x2* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
(*tint_module_vars.s) = (*tint_module_vars.u);
(*tint_module_vars.s)[1u] = (*tint_module_vars.u)[0u];
diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x2_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/unnested/mat3x2_f32/to_storage.wgsl.expected.msl
index bae1216..8950157 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x2_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x2_f32/to_storage.wgsl.expected.msl
@@ -6,7 +6,7 @@
device float3x2* s;
};
-kernel void f(const constant float3x2* u [[buffer(0)]], device float3x2* s [[buffer(1)]]) {
+kernel void f(const constant float3x2* u [[buffer(1)]], device float3x2* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
(*tint_module_vars.s) = (*tint_module_vars.u);
(*tint_module_vars.s)[1u] = (*tint_module_vars.u)[0u];
diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x3_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/unnested/mat3x3_f16/to_storage.wgsl.expected.msl
index 4c9b52d..c2f25b7 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x3_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x3_f16/to_storage.wgsl.expected.msl
@@ -29,7 +29,7 @@
(*target)[2u].packed = packed_half3(value_param[2u]);
}
-kernel void f(const constant tint_array<tint_packed_vec3_f16_array_element, 3>* u [[buffer(0)]], device tint_array<tint_packed_vec3_f16_array_element, 3>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<tint_packed_vec3_f16_array_element, 3>* u [[buffer(1)]], device tint_array<tint_packed_vec3_f16_array_element, 3>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
tint_array<tint_packed_vec3_f16_array_element, 3> const v = (*tint_module_vars.u);
tint_store_and_preserve_padding(tint_module_vars.s, half3x3(half3(v[0u].packed), half3(v[1u].packed), half3(v[2u].packed)));
diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x3_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/unnested/mat3x3_f32/to_storage.wgsl.expected.msl
index 4947863..261a433 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x3_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x3_f32/to_storage.wgsl.expected.msl
@@ -29,7 +29,7 @@
(*target)[2u].packed = packed_float3(value_param[2u]);
}
-kernel void f(const constant tint_array<tint_packed_vec3_f32_array_element, 3>* u [[buffer(0)]], device tint_array<tint_packed_vec3_f32_array_element, 3>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<tint_packed_vec3_f32_array_element, 3>* u [[buffer(1)]], device tint_array<tint_packed_vec3_f32_array_element, 3>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
tint_array<tint_packed_vec3_f32_array_element, 3> const v = (*tint_module_vars.u);
tint_store_and_preserve_padding(tint_module_vars.s, float3x3(float3(v[0u].packed), float3(v[1u].packed), float3(v[2u].packed)));
diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x4_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/unnested/mat3x4_f16/to_storage.wgsl.expected.msl
index 1849ad5..f086548 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x4_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x4_f16/to_storage.wgsl.expected.msl
@@ -6,7 +6,7 @@
device half3x4* s;
};
-kernel void f(const constant half3x4* u [[buffer(0)]], device half3x4* s [[buffer(1)]]) {
+kernel void f(const constant half3x4* u [[buffer(1)]], device half3x4* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
(*tint_module_vars.s) = (*tint_module_vars.u);
(*tint_module_vars.s)[1u] = (*tint_module_vars.u)[0u];
diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x4_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/unnested/mat3x4_f32/to_storage.wgsl.expected.msl
index 8a77f60..0d592a8 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x4_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x4_f32/to_storage.wgsl.expected.msl
@@ -6,7 +6,7 @@
device float3x4* s;
};
-kernel void f(const constant float3x4* u [[buffer(0)]], device float3x4* s [[buffer(1)]]) {
+kernel void f(const constant float3x4* u [[buffer(1)]], device float3x4* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
(*tint_module_vars.s) = (*tint_module_vars.u);
(*tint_module_vars.s)[1u] = (*tint_module_vars.u)[0u];
diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x2_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/unnested/mat4x2_f16/to_storage.wgsl.expected.msl
index 64764c6..e7c32a0 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x2_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x2_f16/to_storage.wgsl.expected.msl
@@ -6,7 +6,7 @@
device half4x2* s;
};
-kernel void f(const constant half4x2* u [[buffer(0)]], device half4x2* s [[buffer(1)]]) {
+kernel void f(const constant half4x2* u [[buffer(1)]], device half4x2* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
(*tint_module_vars.s) = (*tint_module_vars.u);
(*tint_module_vars.s)[1u] = (*tint_module_vars.u)[0u];
diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x2_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/unnested/mat4x2_f32/to_storage.wgsl.expected.msl
index 045bcf9..c4e4536 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x2_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x2_f32/to_storage.wgsl.expected.msl
@@ -6,7 +6,7 @@
device float4x2* s;
};
-kernel void f(const constant float4x2* u [[buffer(0)]], device float4x2* s [[buffer(1)]]) {
+kernel void f(const constant float4x2* u [[buffer(1)]], device float4x2* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
(*tint_module_vars.s) = (*tint_module_vars.u);
(*tint_module_vars.s)[1u] = (*tint_module_vars.u)[0u];
diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x3_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/unnested/mat4x3_f16/to_storage.wgsl.expected.msl
index 3acb9b7..86ff3ca 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x3_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x3_f16/to_storage.wgsl.expected.msl
@@ -30,7 +30,7 @@
(*target)[3u].packed = packed_half3(value_param[3u]);
}
-kernel void f(const constant tint_array<tint_packed_vec3_f16_array_element, 4>* u [[buffer(0)]], device tint_array<tint_packed_vec3_f16_array_element, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<tint_packed_vec3_f16_array_element, 4>* u [[buffer(1)]], device tint_array<tint_packed_vec3_f16_array_element, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
tint_array<tint_packed_vec3_f16_array_element, 4> const v = (*tint_module_vars.u);
tint_store_and_preserve_padding(tint_module_vars.s, half4x3(half3(v[0u].packed), half3(v[1u].packed), half3(v[2u].packed), half3(v[3u].packed)));
diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x3_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/unnested/mat4x3_f32/to_storage.wgsl.expected.msl
index 071fdde..c8cd797 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x3_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x3_f32/to_storage.wgsl.expected.msl
@@ -30,7 +30,7 @@
(*target)[3u].packed = packed_float3(value_param[3u]);
}
-kernel void f(const constant tint_array<tint_packed_vec3_f32_array_element, 4>* u [[buffer(0)]], device tint_array<tint_packed_vec3_f32_array_element, 4>* s [[buffer(1)]]) {
+kernel void f(const constant tint_array<tint_packed_vec3_f32_array_element, 4>* u [[buffer(1)]], device tint_array<tint_packed_vec3_f32_array_element, 4>* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
tint_array<tint_packed_vec3_f32_array_element, 4> const v = (*tint_module_vars.u);
tint_store_and_preserve_padding(tint_module_vars.s, float4x3(float3(v[0u].packed), float3(v[1u].packed), float3(v[2u].packed), float3(v[3u].packed)));
diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x4_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/unnested/mat4x4_f16/to_storage.wgsl.expected.msl
index 34ef1e4..36db13f 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x4_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x4_f16/to_storage.wgsl.expected.msl
@@ -6,7 +6,7 @@
device half4x4* s;
};
-kernel void f(const constant half4x4* u [[buffer(0)]], device half4x4* s [[buffer(1)]]) {
+kernel void f(const constant half4x4* u [[buffer(1)]], device half4x4* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
(*tint_module_vars.s) = (*tint_module_vars.u);
(*tint_module_vars.s)[1u] = (*tint_module_vars.u)[0u];
diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x4_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/unnested/mat4x4_f32/to_storage.wgsl.expected.msl
index 3f05c7b..79c964f 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x4_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x4_f32/to_storage.wgsl.expected.msl
@@ -6,7 +6,7 @@
device float4x4* s;
};
-kernel void f(const constant float4x4* u [[buffer(0)]], device float4x4* s [[buffer(1)]]) {
+kernel void f(const constant float4x4* u [[buffer(1)]], device float4x4* s [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .s=s};
(*tint_module_vars.s) = (*tint_module_vars.u);
(*tint_module_vars.s)[1u] = (*tint_module_vars.u)[0u];
diff --git a/test/tint/bug/chromium/1434271.wgsl.expected.msl b/test/tint/bug/chromium/1434271.wgsl.expected.msl
index 708dd0e..1ebc213 100644
--- a/test/tint/bug/chromium/1434271.wgsl.expected.msl
+++ b/test/tint/bug/chromium/1434271.wgsl.expected.msl
@@ -249,7 +249,7 @@
}
}
-kernel void export_level(uint3 coord [[thread_position_in_grid]], const constant UBO* ubo [[buffer(0)]], const device Buffer* buf_in [[buffer(2)]], device Buffer* buf_out [[buffer(1)]], texture2d<float, access::write> tex_out [[texture(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
+kernel void export_level(uint3 coord [[thread_position_in_grid]], const constant UBO* ubo [[buffer(0)]], const device Buffer* buf_in [[buffer(1)]], device Buffer* buf_out [[buffer(2)]], texture2d<float, access::write> tex_out [[texture(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.ubo=ubo, .buf_in=buf_in, .buf_out=buf_out, .tex_out=tex_out, .tint_storage_buffer_sizes=tint_storage_buffer_sizes};
export_level_inner(coord, tint_module_vars);
}
diff --git a/test/tint/bug/fxc/dyn_array_idx/read/function.wgsl.expected.msl b/test/tint/bug/fxc/dyn_array_idx/read/function.wgsl.expected.msl
index 0e55b68..da9fbb1 100644
--- a/test/tint/bug/fxc/dyn_array_idx/read/function.wgsl.expected.msl
+++ b/test/tint/bug/fxc/dyn_array_idx/read/function.wgsl.expected.msl
@@ -30,7 +30,7 @@
tint_array<int, 64> data;
};
-kernel void f(const constant UBO* ubo [[buffer(0)]], device Result* result [[buffer(1)]]) {
+kernel void f(const constant UBO* ubo [[buffer(1)]], device Result* result [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.ubo=ubo, .result=result};
S s = {};
(*tint_module_vars.result).out = s.data[min(uint((*tint_module_vars.ubo).dynamic_idx), 63u)];
diff --git a/test/tint/bug/fxc/dyn_array_idx/read/private.wgsl.expected.msl b/test/tint/bug/fxc/dyn_array_idx/read/private.wgsl.expected.msl
index defda94..27ad461 100644
--- a/test/tint/bug/fxc/dyn_array_idx/read/private.wgsl.expected.msl
+++ b/test/tint/bug/fxc/dyn_array_idx/read/private.wgsl.expected.msl
@@ -31,7 +31,7 @@
thread S* s;
};
-kernel void f(const constant UBO* ubo [[buffer(0)]], device Result* result [[buffer(1)]]) {
+kernel void f(const constant UBO* ubo [[buffer(1)]], device Result* result [[buffer(0)]]) {
thread S s = {};
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.ubo=ubo, .result=result, .s=(&s)};
(*tint_module_vars.result).out = (*tint_module_vars.s).data[min(uint((*tint_module_vars.ubo).dynamic_idx), 63u)];
diff --git a/test/tint/bug/fxc/dyn_array_idx/read/storage.wgsl.expected.msl b/test/tint/bug/fxc/dyn_array_idx/read/storage.wgsl.expected.msl
index e7b5866..269b84b 100644
--- a/test/tint/bug/fxc/dyn_array_idx/read/storage.wgsl.expected.msl
+++ b/test/tint/bug/fxc/dyn_array_idx/read/storage.wgsl.expected.msl
@@ -31,7 +31,7 @@
device SSBO* ssbo;
};
-kernel void f(const constant UBO* ubo [[buffer(0)]], device Result* result [[buffer(1)]], device SSBO* ssbo [[buffer(2)]]) {
+kernel void f(const constant UBO* ubo [[buffer(2)]], device Result* result [[buffer(0)]], device SSBO* ssbo [[buffer(1)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.ubo=ubo, .result=result, .ssbo=ssbo};
(*tint_module_vars.result).out = (*tint_module_vars.ssbo).data[min(uint((*tint_module_vars.ubo).dynamic_idx), 3u)];
}
diff --git a/test/tint/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.msl b/test/tint/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.msl
index 0a816e3..f4199c7 100644
--- a/test/tint/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.msl
+++ b/test/tint/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.msl
@@ -28,7 +28,7 @@
device Result* result;
};
-kernel void f(const constant UBO* ubo [[buffer(0)]], device Result* result [[buffer(1)]]) {
+kernel void f(const constant UBO* ubo [[buffer(1)]], device Result* result [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.ubo=ubo, .result=result};
(*tint_module_vars.result).out = (*tint_module_vars.ubo).data[min(uint((*tint_module_vars.ubo).dynamic_idx), 3u)].x;
}
diff --git a/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.msl b/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.msl
index 3546f28..ff3f8d5 100644
--- a/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.msl
+++ b/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.msl
@@ -55,7 +55,7 @@
(*tint_module_vars.result).out = (*tint_module_vars.s).data[min(uint((*tint_module_vars.ubo).dynamic_idx), 63u)];
}
-kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant UBO* ubo [[buffer(0)]], device Result* result [[buffer(1)]], threadgroup tint_symbol_1* v_2 [[threadgroup(0)]]) {
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant UBO* ubo [[buffer(1)]], device Result* result [[buffer(0)]], threadgroup tint_symbol_1* v_2 [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.ubo=ubo, .result=result, .s=(&(*v_2).tint_symbol)};
f_inner(tint_local_index, tint_module_vars);
}
diff --git a/test/tint/bug/fxc/dyn_array_idx/write/storage.wgsl.expected.msl b/test/tint/bug/fxc/dyn_array_idx/write/storage.wgsl.expected.msl
index 27293b7..a2b176b 100644
--- a/test/tint/bug/fxc/dyn_array_idx/write/storage.wgsl.expected.msl
+++ b/test/tint/bug/fxc/dyn_array_idx/write/storage.wgsl.expected.msl
@@ -31,7 +31,7 @@
device SSBO* ssbo;
};
-kernel void f(const constant UBO* ubo [[buffer(0)]], device Result* result [[buffer(2)]], device SSBO* ssbo [[buffer(1)]]) {
+kernel void f(const constant UBO* ubo [[buffer(1)]], device Result* result [[buffer(2)]], device SSBO* ssbo [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.ubo=ubo, .result=result, .ssbo=ssbo};
(*tint_module_vars.ssbo).data[min(uint((*tint_module_vars.ubo).dynamic_idx), 3u)] = 1;
(*tint_module_vars.result).out = (*tint_module_vars.ssbo).data[3u];
diff --git a/test/tint/bug/fxc/indexed_assign_to_array_in_struct/1206.wgsl.expected.msl b/test/tint/bug/fxc/indexed_assign_to_array_in_struct/1206.wgsl.expected.msl
index 9e08465..d19719d 100644
--- a/test/tint/bug/fxc/indexed_assign_to_array_in_struct/1206.wgsl.expected.msl
+++ b/test/tint/bug/fxc/indexed_assign_to_array_in_struct/1206.wgsl.expected.msl
@@ -57,7 +57,7 @@
return Particle{.position=v, .lifetime=(*from).lifetime, .color=(*from).color, .velocity=float3((*from).velocity)};
}
-kernel void v_1(const device Particles_packed_vec3* particles [[buffer(1)]], const constant Simulation* sim [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
+kernel void v_1(const device Particles_packed_vec3* particles [[buffer(0)]], const constant Simulation* sim [[buffer(1)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.particles=particles, .sim=sim, .tint_storage_buffer_sizes=tint_storage_buffer_sizes};
uint const v_2 = ((((*tint_module_vars.tint_storage_buffer_sizes)[0u].x - 0u) / 176u) - 1u);
Particle particle = tint_load_struct_packed_vec3((&(*tint_module_vars.particles).p[min(uint(0), v_2)]));
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.msl b/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.msl
index a7654ca..d9d538f 100644
--- a/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.msl
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.msl
@@ -6,7 +6,7 @@
device packed_float3* v1;
};
-kernel void v(const constant uint* i [[buffer(0)]], device packed_float3* v1 [[buffer(1)]]) {
+kernel void v(const constant uint* i [[buffer(1)]], device packed_float3* v1 [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.i=i, .v1=v1};
(*tint_module_vars.v1)[min((*tint_module_vars.i), 2u)] = 1.0f;
}
diff --git a/test/tint/bug/tint/1121.wgsl.expected.msl b/test/tint/bug/tint/1121.wgsl.expected.msl
index 689b765..5c5eb41 100644
--- a/test/tint/bug/tint/1121.wgsl.expected.msl
+++ b/test/tint/bug/tint/1121.wgsl.expected.msl
@@ -181,7 +181,7 @@
}
}
-kernel void v_9(uint3 GlobalInvocationID [[thread_position_in_grid]], device LightsBuffer_packed_vec3* lightsBuffer [[buffer(2)]], device Tiles* tileLightId [[buffer(3)]], const constant Config* config [[buffer(0)]], const constant Uniforms* uniforms [[buffer(1)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
+kernel void v_9(uint3 GlobalInvocationID [[thread_position_in_grid]], device LightsBuffer_packed_vec3* lightsBuffer [[buffer(1)]], device Tiles* tileLightId [[buffer(3)]], const constant Config* config [[buffer(0)]], const constant Uniforms* uniforms [[buffer(2)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.lightsBuffer=lightsBuffer, .tileLightId=tileLightId, .config=config, .uniforms=uniforms, .tint_storage_buffer_sizes=tint_storage_buffer_sizes};
main_inner(GlobalInvocationID, tint_module_vars);
}
diff --git a/test/tint/bug/tint/1739.wgsl.expected.msl b/test/tint/bug/tint/1739.wgsl.expected.msl
index 3bf7a14..2bce7a0 100644
--- a/test/tint/bug/tint/1739.wgsl.expected.msl
+++ b/test/tint/bug/tint/1739.wgsl.expected.msl
@@ -111,7 +111,7 @@
return tint_ExternalTextureParams{.numPlanes=v_10, .doYuvToRgbConversionOnly=v_11, .yuvToRgbConversionMatrix=v_12, .gammaDecodeParams=v_13, .gammaEncodeParams=v_14, .gamutConversionMatrix=v_16, .sampleTransform=(*from).sampleTransform, .loadTransform=(*from).loadTransform, .samplePlane0RectMin=(*from).samplePlane0RectMin, .samplePlane0RectMax=(*from).samplePlane0RectMax, .samplePlane1RectMin=(*from).samplePlane1RectMin, .samplePlane1RectMax=(*from).samplePlane1RectMax, .apparentSize=(*from).apparentSize, .plane1CoordFactor=(*from).plane1CoordFactor};
}
-kernel void v_17(texture2d<float, access::sample> t_plane0 [[texture(1)]], texture2d<float, access::sample> t_plane1 [[texture(2)]], const constant tint_ExternalTextureParams_packed_vec3* t_params [[buffer(3)]], texture2d<float, access::write> outImage [[texture(0)]]) {
+kernel void v_17(texture2d<float, access::sample> t_plane0 [[texture(0)]], texture2d<float, access::sample> t_plane1 [[texture(2)]], const constant tint_ExternalTextureParams_packed_vec3* t_params [[buffer(3)]], texture2d<float, access::write> outImage [[texture(1)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.t_plane0=t_plane0, .t_plane1=t_plane1, .t_params=t_params, .outImage=outImage};
tint_ExternalTextureParams const v_18 = tint_load_struct_packed_vec3(tint_module_vars.t_params);
float4 red = tint_TextureLoadExternal(tint_module_vars.t_plane0, tint_module_vars.t_plane1, v_18, min(uint2(int2(10)), ((v_18.apparentSize + uint2(1u)) - uint2(1u))));
diff --git a/test/tint/bug/tint/2010.spvasm.expected.msl b/test/tint/bug/tint/2010.spvasm.expected.msl
index 848e381..507a221 100644
--- a/test/tint/bug/tint/2010.spvasm.expected.msl
+++ b/test/tint/bug/tint/2010.spvasm.expected.msl
@@ -186,7 +186,7 @@
main_1(tint_module_vars);
}
-kernel void v_4(uint3 x_3_param [[thread_position_in_threadgroup]], uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_5* v_5 [[threadgroup(0)]], const constant S_2* x_6 [[buffer(0)]], const device S_3* x_9 [[buffer(2)]], device S_4* x_12 [[buffer(1)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
+kernel void v_4(uint3 x_3_param [[thread_position_in_threadgroup]], uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_5* v_5 [[threadgroup(0)]], const constant S_2* x_6 [[buffer(0)]], const device S_3* x_9 [[buffer(1)]], device S_4* x_12 [[buffer(2)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
thread uint3 x_3 = 0u;
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.x_28=(&(*v_5).tint_symbol), .x_34=(&(*v_5).tint_symbol_1), .x_35=(&(*v_5).tint_symbol_2), .x_36=(&(*v_5).tint_symbol_3), .x_37=(&(*v_5).tint_symbol_4), .x_3=(&x_3), .x_6=x_6, .x_9=x_9, .x_12=x_12, .tint_storage_buffer_sizes=tint_storage_buffer_sizes};
main_inner(x_3_param, tint_local_index, tint_module_vars);
diff --git a/test/tint/bug/tint/744.wgsl.expected.msl b/test/tint/bug/tint/744.wgsl.expected.msl
index c40e9db..f1aeca3 100644
--- a/test/tint/bug/tint/744.wgsl.expected.msl
+++ b/test/tint/bug/tint/744.wgsl.expected.msl
@@ -56,7 +56,7 @@
(*tint_module_vars.resultMatrix).numbers[min(index, ((((*tint_module_vars.tint_storage_buffer_sizes)[0u].z - 0u) / 4u) - 1u))] = result;
}
-kernel void v(uint3 global_id [[thread_position_in_grid]], const device Matrix* firstMatrix [[buffer(2)]], const device Matrix* secondMatrix [[buffer(3)]], device Matrix* resultMatrix [[buffer(1)]], const constant Uniforms* uniforms [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
+kernel void v(uint3 global_id [[thread_position_in_grid]], const device Matrix* firstMatrix [[buffer(1)]], const device Matrix* secondMatrix [[buffer(2)]], device Matrix* resultMatrix [[buffer(3)]], const constant Uniforms* uniforms [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.firstMatrix=firstMatrix, .secondMatrix=secondMatrix, .resultMatrix=resultMatrix, .uniforms=uniforms, .tint_storage_buffer_sizes=tint_storage_buffer_sizes};
main_inner(global_id, tint_module_vars);
}
diff --git a/test/tint/bug/tint/914.wgsl.expected.msl b/test/tint/bug/tint/914.wgsl.expected.msl
index 1c33e60..5a619af 100644
--- a/test/tint/bug/tint/914.wgsl.expected.msl
+++ b/test/tint/bug/tint/914.wgsl.expected.msl
@@ -287,7 +287,7 @@
}
}
-kernel void v_6(uint3 local_id [[thread_position_in_threadgroup]], uint3 global_id [[thread_position_in_grid]], uint tint_local_index [[thread_index_in_threadgroup]], const device Matrix* firstMatrix [[buffer(2)]], const device Matrix* secondMatrix [[buffer(3)]], device Matrix* resultMatrix [[buffer(1)]], const constant Uniforms* uniforms [[buffer(0)]], threadgroup tint_symbol_2* v_7 [[threadgroup(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
+kernel void v_6(uint3 local_id [[thread_position_in_threadgroup]], uint3 global_id [[thread_position_in_grid]], uint tint_local_index [[thread_index_in_threadgroup]], const device Matrix* firstMatrix [[buffer(1)]], const device Matrix* secondMatrix [[buffer(2)]], device Matrix* resultMatrix [[buffer(3)]], const constant Uniforms* uniforms [[buffer(0)]], threadgroup tint_symbol_2* v_7 [[threadgroup(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.firstMatrix=firstMatrix, .secondMatrix=secondMatrix, .resultMatrix=resultMatrix, .uniforms=uniforms, .mm_Asub=(&(*v_7).tint_symbol), .mm_Bsub=(&(*v_7).tint_symbol_1), .tint_storage_buffer_sizes=tint_storage_buffer_sizes};
main_inner(local_id, global_id, tint_local_index, tint_module_vars);
}
diff --git a/test/tint/bug/tint/959.wgsl.expected.msl b/test/tint/bug/tint/959.wgsl.expected.msl
index 7a1b235..0f61e2c 100644
--- a/test/tint/bug/tint/959.wgsl.expected.msl
+++ b/test/tint/bug/tint/959.wgsl.expected.msl
@@ -56,6 +56,6 @@
sampler s15;
};
-fragment void v(const device S* b0 [[buffer(8)]], const device S* b1 [[buffer(9)]], const device S* b2 [[buffer(10)]], const device S* b3 [[buffer(11)]], const device S* b4 [[buffer(12)]], const device S* b5 [[buffer(13)]], const device S* b6 [[buffer(14)]], const device S* b7 [[buffer(15)]], const constant S* b8 [[buffer(0)]], const constant S* b9 [[buffer(1)]], const constant S* b10 [[buffer(2)]], const constant S* b11 [[buffer(3)]], const constant S* b12 [[buffer(4)]], const constant S* b13 [[buffer(5)]], const constant S* b14 [[buffer(6)]], const constant S* b15 [[buffer(7)]], texture2d<float, access::sample> t0 [[texture(0)]], texture2d<float, access::sample> t1 [[texture(1)]], texture2d<float, access::sample> t2 [[texture(2)]], texture2d<float, access::sample> t3 [[texture(3)]], texture2d<float, access::sample> t4 [[texture(4)]], texture2d<float, access::sample> t5 [[texture(5)]], texture2d<float, access::sample> t6 [[texture(6)]], texture2d<float, access::sample> t7 [[texture(7)]], depth2d<float, access::sample> t8 [[texture(8)]], depth2d<float, access::sample> t9 [[texture(9)]], depth2d<float, access::sample> t10 [[texture(10)]], depth2d<float, access::sample> t11 [[texture(11)]], depth2d<float, access::sample> t12 [[texture(12)]], depth2d<float, access::sample> t13 [[texture(13)]], depth2d<float, access::sample> t14 [[texture(14)]], depth2d<float, access::sample> t15 [[texture(15)]], sampler s0 [[sampler(0)]], sampler s1 [[sampler(1)]], sampler s2 [[sampler(2)]], sampler s3 [[sampler(3)]], sampler s4 [[sampler(4)]], sampler s5 [[sampler(5)]], sampler s6 [[sampler(6)]], sampler s7 [[sampler(7)]], sampler s8 [[sampler(8)]], sampler s9 [[sampler(9)]], sampler s10 [[sampler(10)]], sampler s11 [[sampler(11)]], sampler s12 [[sampler(12)]], sampler s13 [[sampler(13)]], sampler s14 [[sampler(14)]], sampler s15 [[sampler(15)]]) {
+fragment void v(const device S* b0 [[buffer(0)]], const device S* b1 [[buffer(1)]], const device S* b2 [[buffer(2)]], const device S* b3 [[buffer(3)]], const device S* b4 [[buffer(4)]], const device S* b5 [[buffer(5)]], const device S* b6 [[buffer(6)]], const device S* b7 [[buffer(7)]], const constant S* b8 [[buffer(8)]], const constant S* b9 [[buffer(9)]], const constant S* b10 [[buffer(10)]], const constant S* b11 [[buffer(11)]], const constant S* b12 [[buffer(12)]], const constant S* b13 [[buffer(13)]], const constant S* b14 [[buffer(14)]], const constant S* b15 [[buffer(15)]], texture2d<float, access::sample> t0 [[texture(0)]], texture2d<float, access::sample> t1 [[texture(1)]], texture2d<float, access::sample> t2 [[texture(2)]], texture2d<float, access::sample> t3 [[texture(3)]], texture2d<float, access::sample> t4 [[texture(4)]], texture2d<float, access::sample> t5 [[texture(5)]], texture2d<float, access::sample> t6 [[texture(6)]], texture2d<float, access::sample> t7 [[texture(7)]], depth2d<float, access::sample> t8 [[texture(8)]], depth2d<float, access::sample> t9 [[texture(9)]], depth2d<float, access::sample> t10 [[texture(10)]], depth2d<float, access::sample> t11 [[texture(11)]], depth2d<float, access::sample> t12 [[texture(12)]], depth2d<float, access::sample> t13 [[texture(13)]], depth2d<float, access::sample> t14 [[texture(14)]], depth2d<float, access::sample> t15 [[texture(15)]], sampler s0 [[sampler(0)]], sampler s1 [[sampler(1)]], sampler s2 [[sampler(2)]], sampler s3 [[sampler(3)]], sampler s4 [[sampler(4)]], sampler s5 [[sampler(5)]], sampler s6 [[sampler(6)]], sampler s7 [[sampler(7)]], sampler s8 [[sampler(8)]], sampler s9 [[sampler(9)]], sampler s10 [[sampler(10)]], sampler s11 [[sampler(11)]], sampler s12 [[sampler(12)]], sampler s13 [[sampler(13)]], sampler s14 [[sampler(14)]], sampler s15 [[sampler(15)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.b0=b0, .b1=b1, .b2=b2, .b3=b3, .b4=b4, .b5=b5, .b6=b6, .b7=b7, .b8=b8, .b9=b9, .b10=b10, .b11=b11, .b12=b12, .b13=b13, .b14=b14, .b15=b15, .t0=t0, .t1=t1, .t2=t2, .t3=t3, .t4=t4, .t5=t5, .t6=t6, .t7=t7, .t8=t8, .t9=t9, .t10=t10, .t11=t11, .t12=t12, .t13=t13, .t14=t14, .t15=t15, .s0=s0, .s1=s1, .s2=s2, .s3=s3, .s4=s4, .s5=s5, .s6=s6, .s7=s7, .s8=s8, .s9=s9, .s10=s10, .s11=s11, .s12=s12, .s13=s13, .s14=s14, .s15=s15};
}
diff --git a/test/tint/bug/tint/993.wgsl.expected.msl b/test/tint/bug/tint/993.wgsl.expected.msl
index 43dacfb..a61a0ec 100644
--- a/test/tint/bug/tint/993.wgsl.expected.msl
+++ b/test/tint/bug/tint/993.wgsl.expected.msl
@@ -35,7 +35,7 @@
return atomic_load_explicit((&(*tint_module_vars.s).data[min((0u + uint((*tint_module_vars.constants).zero)), 2u)]), memory_order_relaxed);
}
-kernel void v(const constant Constants* constants [[buffer(0)]], device Result* result [[buffer(1)]], device TestData* s [[buffer(2)]]) {
+kernel void v(const constant Constants* constants [[buffer(2)]], device Result* result [[buffer(0)]], device TestData* s [[buffer(1)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.constants=constants, .result=result, .s=s};
(*tint_module_vars.result).value = uint(runTest(tint_module_vars));
}
diff --git a/test/tint/samples/compute_boids.wgsl.expected.msl b/test/tint/samples/compute_boids.wgsl.expected.msl
index ed383a4..3697735 100644
--- a/test/tint/samples/compute_boids.wgsl.expected.msl
+++ b/test/tint/samples/compute_boids.wgsl.expected.msl
@@ -163,7 +163,7 @@
(*tint_module_vars.particlesB).particles[min(index, 4u)].vel = vVel;
}
-kernel void comp_main(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], const constant SimParams* params [[buffer(0)]], device Particles* particlesA [[buffer(1)]], device Particles* particlesB [[buffer(2)]]) {
+kernel void comp_main(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], const constant SimParams* params [[buffer(1)]], device Particles* particlesA [[buffer(0)]], device Particles* particlesB [[buffer(2)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.params=params, .particlesA=particlesA, .particlesB=particlesB};
comp_main_inner(gl_GlobalInvocationID, tint_module_vars);
}
diff --git a/test/tint/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array.wgsl.expected.msl b/test/tint/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array.wgsl.expected.msl
index a7e3131..e79671a 100644
--- a/test/tint/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array.wgsl.expected.msl
+++ b/test/tint/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array.wgsl.expected.msl
@@ -31,7 +31,7 @@
const constant tint_array<uint4, 1>* tint_storage_buffer_sizes;
};
-kernel void v_1(const constant Uniforms* uniforms [[buffer(0)]], device OuterS* s1 [[buffer(1)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
+kernel void v_1(const constant Uniforms* uniforms [[buffer(1)]], device OuterS* s1 [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.uniforms=uniforms, .s1=s1, .tint_storage_buffer_sizes=tint_storage_buffer_sizes};
InnerS v = {};
(*tint_module_vars.s1).a1[min((*tint_module_vars.uniforms).i, ((((*tint_module_vars.tint_storage_buffer_sizes)[0u].x - 0u) / 4u) - 1u))] = v;
diff --git a/test/tint/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array_struct_array.wgsl.expected.msl b/test/tint/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array_struct_array.wgsl.expected.msl
index c9f23db..bbce5fe 100644
--- a/test/tint/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array_struct_array.wgsl.expected.msl
+++ b/test/tint/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array_struct_array.wgsl.expected.msl
@@ -36,7 +36,7 @@
const constant tint_array<uint4, 1>* tint_storage_buffer_sizes;
};
-kernel void v_1(const constant Uniforms* uniforms [[buffer(0)]], device OuterS* s [[buffer(1)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
+kernel void v_1(const constant Uniforms* uniforms [[buffer(1)]], device OuterS* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.uniforms=uniforms, .s=s, .tint_storage_buffer_sizes=tint_storage_buffer_sizes};
InnerS v = {};
(*tint_module_vars.s).a1[min((*tint_module_vars.uniforms).i, ((((*tint_module_vars.tint_storage_buffer_sizes)[0u].x - 0u) / 32u) - 1u))].a2[min((*tint_module_vars.uniforms).j, 7u)] = v;
