Allow non-struct buffer store types
For SPIR-V, wrap non-struct types in structs in the
AddSpirvBlockDecoration transform.
For MSL, wrap runtime-sized arrays in structs in the
ModuleScopeVarToEntryPointParam transform.
Bug: tint:1372
Change-Id: Icced5d77b4538e816aa9fab57a634a9f4c52fdab
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/76162
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
diff --git a/src/resolver/resolver_validation.cc b/src/resolver/resolver_validation.cc
index 1df63ba..3218e4a 100644
--- a/src/resolver/resolver_validation.cc
+++ b/src/resolver/resolver_validation.cc
@@ -459,42 +459,6 @@
return false;
}
- switch (var->StorageClass()) {
- case ast::StorageClass::kStorage: {
- // https://gpuweb.github.io/gpuweb/wgsl/#module-scope-variables
- // A variable in the storage storage class is a storage buffer variable.
- // Its store type must be a host-shareable structure type with block
- // attribute, satisfying the storage class constraints.
-
- auto* str = var->Type()->UnwrapRef()->As<sem::Struct>();
- if (!str) {
- AddError(
- "variables declared in the <storage> storage class must be of a "
- "structure type",
- decl->source);
- return false;
- }
- break;
- }
- case ast::StorageClass::kUniform: {
- // https://gpuweb.github.io/gpuweb/wgsl/#module-scope-variables
- // A variable in the uniform storage class is a uniform buffer variable.
- // Its store type must be a host-shareable structure type with block
- // attribute, satisfying the storage class constraints.
- auto* str = var->Type()->UnwrapRef()->As<sem::Struct>();
- if (!str) {
- AddError(
- "variables declared in the <uniform> storage class must be of a "
- "structure type",
- decl->source);
- return false;
- }
- break;
- }
- default:
- break;
- }
-
if (!decl->is_const) {
if (!ValidateAtomicVariable(var)) {
return false;
@@ -580,14 +544,6 @@
return false;
}
- if (auto* r = storage_ty->As<sem::Array>()) {
- if (r->IsRuntimeSized()) {
- AddError("runtime arrays may only appear as the last member of a struct",
- decl->source);
- return false;
- }
- }
-
if (auto* r = storage_ty->As<sem::MultisampledTexture>()) {
if (r->dim() != ast::TextureDimension::k2d) {
AddError("only 2d multisampled textures are supported", decl->source);
diff --git a/src/resolver/storage_class_validation_test.cc b/src/resolver/storage_class_validation_test.cc
index f836088..5d922c2 100644
--- a/src/resolver/storage_class_validation_test.cc
+++ b/src/resolver/storage_class_validation_test.cc
@@ -92,6 +92,40 @@
}
TEST_F(ResolverStorageClassValidationTest, StorageBufferBool) {
+ // var<storage> g : bool;
+ Global(Source{{56, 78}}, "g", ty.bool_(), ast::StorageClass::kStorage,
+ ast::DecorationList{
+ create<ast::BindingDecoration>(0),
+ create<ast::GroupDecoration>(0),
+ });
+
+ ASSERT_FALSE(r()->Resolve());
+
+ EXPECT_EQ(
+ r()->error(),
+ R"(56:78 error: Type 'bool' cannot be used in storage class 'storage' as it is non-host-shareable
+56:78 note: while instantiating variable g)");
+}
+
+TEST_F(ResolverStorageClassValidationTest, StorageBufferPointer) {
+ // var<storage> g : ptr<private, f32>;
+ Global(Source{{56, 78}}, "g",
+ ty.pointer(ty.f32(), ast::StorageClass::kPrivate),
+ ast::StorageClass::kStorage,
+ ast::DecorationList{
+ create<ast::BindingDecoration>(0),
+ create<ast::GroupDecoration>(0),
+ });
+
+ ASSERT_FALSE(r()->Resolve());
+
+ EXPECT_EQ(
+ r()->error(),
+ R"(56:78 error: Type 'ptr<private, f32, read_write>' cannot be used in storage class 'storage' as it is non-host-shareable
+56:78 note: while instantiating variable g)");
+}
+
+TEST_F(ResolverStorageClassValidationTest, StorageBufferIntScalar) {
// var<storage> g : i32;
Global(Source{{56, 78}}, "g", ty.i32(), ast::StorageClass::kStorage,
ast::DecorationList{
@@ -99,14 +133,10 @@
create<ast::GroupDecoration>(0),
});
- ASSERT_FALSE(r()->Resolve());
-
- EXPECT_EQ(
- r()->error(),
- R"(56:78 error: variables declared in the <storage> storage class must be of a structure type)");
+ ASSERT_TRUE(r()->Resolve()) << r()->error();
}
-TEST_F(ResolverStorageClassValidationTest, StorageBufferPointer) {
+TEST_F(ResolverStorageClassValidationTest, StorageBufferVector) {
// var<storage> g : vec4<f32>;
Global(Source{{56, 78}}, "g", ty.vec4<f32>(), ast::StorageClass::kStorage,
ast::DecorationList{
@@ -114,11 +144,7 @@
create<ast::GroupDecoration>(0),
});
- ASSERT_FALSE(r()->Resolve());
-
- EXPECT_EQ(
- r()->error(),
- R"(56:78 error: variables declared in the <storage> storage class must be of a structure type)");
+ ASSERT_TRUE(r()->Resolve()) << r()->error();
}
TEST_F(ResolverStorageClassValidationTest, StorageBufferArray) {
@@ -132,11 +158,7 @@
create<ast::GroupDecoration>(0),
});
- ASSERT_FALSE(r()->Resolve());
-
- EXPECT_EQ(
- r()->error(),
- R"(56:78 error: variables declared in the <storage> storage class must be of a structure type)");
+ ASSERT_TRUE(r()->Resolve()) << r()->error();
}
TEST_F(ResolverStorageClassValidationTest, StorageBufferBoolAlias) {
@@ -240,8 +262,10 @@
}
TEST_F(ResolverStorageClassValidationTest, UniformBufferPointer) {
- // var<uniform> g : vec4<f32>;
- Global(Source{{56, 78}}, "g", ty.vec4<f32>(), ast::StorageClass::kUniform,
+ // var<uniform> g : ptr<private, f32>;
+ Global(Source{{56, 78}}, "g",
+ ty.pointer(ty.f32(), ast::StorageClass::kPrivate),
+ ast::StorageClass::kUniform,
ast::DecorationList{
create<ast::BindingDecoration>(0),
create<ast::GroupDecoration>(0),
@@ -251,7 +275,30 @@
EXPECT_EQ(
r()->error(),
- R"(56:78 error: variables declared in the <uniform> storage class must be of a structure type)");
+ R"(56:78 error: Type 'ptr<private, f32, read_write>' cannot be used in storage class 'uniform' as it is non-host-shareable
+56:78 note: while instantiating variable g)");
+}
+
+TEST_F(ResolverStorageClassValidationTest, UniformBufferIntScalar) {
+ // var<uniform> g : i32;
+ Global(Source{{56, 78}}, "g", ty.i32(), ast::StorageClass::kUniform,
+ ast::DecorationList{
+ create<ast::BindingDecoration>(0),
+ create<ast::GroupDecoration>(0),
+ });
+
+ ASSERT_TRUE(r()->Resolve()) << r()->error();
+}
+
+TEST_F(ResolverStorageClassValidationTest, UniformBufferVector) {
+ // var<uniform> g : vec4<f32>;
+ Global(Source{{56, 78}}, "g", ty.vec4<f32>(), ast::StorageClass::kUniform,
+ ast::DecorationList{
+ create<ast::BindingDecoration>(0),
+ create<ast::GroupDecoration>(0),
+ });
+
+ ASSERT_TRUE(r()->Resolve()) << r()->error();
}
TEST_F(ResolverStorageClassValidationTest, UniformBufferArray) {
@@ -264,11 +311,7 @@
create<ast::GroupDecoration>(0),
});
- ASSERT_FALSE(r()->Resolve());
-
- EXPECT_EQ(
- r()->error(),
- R"(56:78 error: variables declared in the <uniform> storage class must be of a structure type)");
+ ASSERT_TRUE(r()->Resolve()) << r()->error();
}
TEST_F(ResolverStorageClassValidationTest, UniformBufferBoolAlias) {
diff --git a/src/transform/add_spirv_block_decoration.cc b/src/transform/add_spirv_block_decoration.cc
index e4829c9..3995c83 100644
--- a/src/transform/add_spirv_block_decoration.cc
+++ b/src/transform/add_spirv_block_decoration.cc
@@ -52,9 +52,9 @@
}
}
- // A map from a struct in the source program to a block-decorated wrapper that
+ // A map from a type in the source program to a block-decorated wrapper that
// contains it in the destination program.
- std::unordered_map<const sem::Struct*, const ast::Struct*> wrapper_structs;
+ std::unordered_map<const sem::Type*, const ast::Struct*> wrapper_structs;
// Process global variables that are buffers.
for (auto* var : ctx.src->AST().GlobalVariables()) {
@@ -64,40 +64,33 @@
continue;
}
- auto* str = sem.Get<sem::Struct>(var->type);
- if (!str) {
- // TODO(jrprice): We'll need to wrap these too, when WGSL supports this.
- TINT_ICE(Transform, ctx.dst->Diagnostics())
- << "non-struct buffer types are not yet supported";
- continue;
- }
+ auto* ty = sem.Get(var->type);
+ auto* str = ty->As<sem::Struct>();
+ if (!str || nested_structs.count(str)) {
+ const char* kMemberName = "inner";
- if (nested_structs.count(str)) {
- const char* kInnerStructMemberName = "inner";
-
- // This struct is nested somewhere else, so we need to wrap it first.
- auto* wrapper = utils::GetOrCreate(wrapper_structs, str, [&]() {
+ // This is a non-struct or a struct that is nested somewhere else, so we
+ // need to wrap it first.
+ auto* wrapper = utils::GetOrCreate(wrapper_structs, ty, [&]() {
auto* block =
ctx.dst->ASTNodes().Create<SpirvBlockDecoration>(ctx.dst->ID());
- auto wrapper_name =
- ctx.src->Symbols().NameFor(str->Declaration()->name) + "_block";
+ auto wrapper_name = ctx.src->Symbols().NameFor(var->symbol) + "_block";
auto* ret = ctx.dst->create<ast::Struct>(
ctx.dst->Symbols().New(wrapper_name),
- ast::StructMemberList{ctx.dst->Member(kInnerStructMemberName,
- CreateASTTypeFor(ctx, str))},
+ ast::StructMemberList{
+ ctx.dst->Member(kMemberName, CreateASTTypeFor(ctx, ty))},
ast::DecorationList{block});
- ctx.InsertAfter(ctx.src->AST().GlobalDeclarations(), str->Declaration(),
- ret);
+ ctx.InsertBefore(ctx.src->AST().GlobalDeclarations(), var, ret);
return ret;
});
ctx.Replace(var->type, ctx.dst->ty.Of(wrapper));
- // Insert a member accessor to get the original struct from the wrapper at
+ // Insert a member accessor to get the original type from the wrapper at
// any usage of the original variable.
for (auto* user : sem_var->Users()) {
- ctx.Replace(user->Declaration(),
- ctx.dst->MemberAccessor(ctx.Clone(var->symbol),
- kInnerStructMemberName));
+ ctx.Replace(
+ user->Declaration(),
+ ctx.dst->MemberAccessor(ctx.Clone(var->symbol), kMemberName));
}
} else {
// Add a block decoration to this struct directly.
diff --git a/src/transform/add_spirv_block_decoration_test.cc b/src/transform/add_spirv_block_decoration_test.cc
index 3780c20..74e4b33 100644
--- a/src/transform/add_spirv_block_decoration_test.cc
+++ b/src/transform/add_spirv_block_decoration_test.cc
@@ -73,7 +73,98 @@
EXPECT_EQ(expect, str(got));
}
-TEST_F(AddSpirvBlockDecorationTest, Basic) {
+TEST_F(AddSpirvBlockDecorationTest, BasicScalar) {
+ auto* src = R"(
+[[group(0), binding(0)]]
+var<uniform> u : f32;
+
+[[stage(fragment)]]
+fn main() {
+ let f = u;
+}
+)";
+ auto* expect = R"(
+[[internal(spirv_block)]]
+struct u_block {
+ inner : f32;
+};
+
+[[group(0), binding(0)]] var<uniform> u : u_block;
+
+[[stage(fragment)]]
+fn main() {
+ let f = u.inner;
+}
+)";
+
+ auto got = Run<AddSpirvBlockDecoration>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(AddSpirvBlockDecorationTest, BasicArray) {
+ auto* src = R"(
+[[group(0), binding(0)]]
+var<uniform> u : array<vec4<f32>, 4u>;
+
+[[stage(fragment)]]
+fn main() {
+ let a = u;
+}
+)";
+ auto* expect = R"(
+[[internal(spirv_block)]]
+struct u_block {
+ inner : array<vec4<f32>, 4u>;
+};
+
+[[group(0), binding(0)]] var<uniform> u : u_block;
+
+[[stage(fragment)]]
+fn main() {
+ let a = u.inner;
+}
+)";
+
+ auto got = Run<AddSpirvBlockDecoration>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(AddSpirvBlockDecorationTest, BasicArray_Alias) {
+ auto* src = R"(
+type Numbers = array<vec4<f32>, 4u>;
+
+[[group(0), binding(0)]]
+var<uniform> u : Numbers;
+
+[[stage(fragment)]]
+fn main() {
+ let a = u;
+}
+)";
+ auto* expect = R"(
+type Numbers = array<vec4<f32>, 4u>;
+
+[[internal(spirv_block)]]
+struct u_block {
+ inner : array<vec4<f32>, 4u>;
+};
+
+[[group(0), binding(0)]] var<uniform> u : u_block;
+
+[[stage(fragment)]]
+fn main() {
+ let a = u.inner;
+}
+)";
+
+ auto got = Run<AddSpirvBlockDecoration>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(AddSpirvBlockDecorationTest, BasicStruct) {
auto* src = R"(
struct S {
f : f32;
@@ -175,18 +266,18 @@
};
[[internal(spirv_block)]]
-struct Inner_block {
- inner : Inner;
-};
-
-[[internal(spirv_block)]]
struct Outer {
i : Inner;
};
[[group(0), binding(0)]] var<uniform> u0 : Outer;
-[[group(0), binding(1)]] var<uniform> u1 : Inner_block;
+[[internal(spirv_block)]]
+struct u1_block {
+ inner : Inner;
+};
+
+[[group(0), binding(1)]] var<uniform> u1 : u1_block;
[[stage(fragment)]]
fn main() {
@@ -226,18 +317,18 @@
f : f32;
};
-[[internal(spirv_block)]]
-struct Inner_block {
- inner : Inner;
-};
-
struct Outer {
i : Inner;
};
var<private> p : Outer;
-[[group(0), binding(1)]] var<uniform> u : Inner_block;
+[[internal(spirv_block)]]
+struct u_block {
+ inner : Inner;
+};
+
+[[group(0), binding(1)]] var<uniform> u : u_block;
[[stage(fragment)]]
fn main() {
@@ -283,20 +374,20 @@
};
[[internal(spirv_block)]]
-struct Inner_block {
- inner : Inner;
-};
-
-[[internal(spirv_block)]]
struct S {
i : Inner;
};
[[group(0), binding(0)]] var<uniform> u0 : S;
-[[group(0), binding(1)]] var<uniform> u1 : Inner_block;
+[[internal(spirv_block)]]
+struct u1_block {
+ inner : Inner;
+};
-[[group(0), binding(2)]] var<uniform> u2 : Inner_block;
+[[group(0), binding(1)]] var<uniform> u1 : u1_block;
+
+[[group(0), binding(2)]] var<uniform> u2 : u1_block;
[[stage(fragment)]]
fn main() {
@@ -332,11 +423,11 @@
};
[[internal(spirv_block)]]
-struct S_block {
+struct u_block {
inner : S;
};
-[[group(0), binding(0)]] var<uniform> u : S_block;
+[[group(0), binding(0)]] var<uniform> u : u_block;
[[stage(fragment)]]
fn main() {
@@ -375,13 +466,13 @@
};
[[internal(spirv_block)]]
-struct S_block {
+struct u0_block {
inner : S;
};
-[[group(0), binding(0)]] var<uniform> u0 : S_block;
+[[group(0), binding(0)]] var<uniform> u0 : u0_block;
-[[group(0), binding(1)]] var<uniform> u1 : S_block;
+[[group(0), binding(1)]] var<uniform> u1 : u0_block;
[[stage(fragment)]]
fn main() {
@@ -427,11 +518,6 @@
f : f32;
};
-[[internal(spirv_block)]]
-struct Inner_block {
- inner : Inner;
-};
-
type MyInner = Inner;
[[internal(spirv_block)]]
@@ -443,7 +529,12 @@
[[group(0), binding(0)]] var<uniform> u0 : MyOuter;
-[[group(0), binding(1)]] var<uniform> u1 : Inner_block;
+[[internal(spirv_block)]]
+struct u1_block {
+ inner : Inner;
+};
+
+[[group(0), binding(1)]] var<uniform> u1 : u1_block;
[[stage(fragment)]]
fn main() {
diff --git a/src/transform/module_scope_var_to_entry_point_param.cc b/src/transform/module_scope_var_to_entry_point_param.cc
index ba61160..960fc84 100644
--- a/src/transform/module_scope_var_to_entry_point_param.cc
+++ b/src/transform/module_scope_var_to_entry_point_param.cc
@@ -157,6 +157,7 @@
for (auto* var : func_sem->TransitivelyReferencedGlobals()) {
auto sc = var->StorageClass();
+ auto* ty = var->Type()->UnwrapRef();
if (sc == ast::StorageClass::kNone) {
continue;
}
@@ -174,13 +175,15 @@
auto new_var_symbol = ctx.dst->Sym();
// Helper to create an AST node for the store type of the variable.
- auto store_type = [&]() {
- return CreateASTTypeFor(ctx, var->Type()->UnwrapRef());
- };
+ auto store_type = [&]() { return CreateASTTypeFor(ctx, ty); };
// Track whether the new variable is a pointer or not.
bool is_pointer = false;
+ // Track whether the new variable was wrapped in a struct or not.
+ bool is_wrapped = false;
+ const char* kWrappedArrayMemberName = "arr";
+
if (is_entry_point) {
if (var->Type()->UnwrapRef()->is_handle()) {
// For a texture or sampler variable, redeclare it as an entry point
@@ -200,8 +203,23 @@
ast::DisabledValidation::kEntryPointParameter));
attributes.push_back(
ctx.dst->Disable(ast::DisabledValidation::kIgnoreStorageClass));
- auto* param_type = ctx.dst->ty.pointer(
- store_type(), sc, var->Declaration()->declared_access);
+
+ auto* param_type = store_type();
+ if (auto* arr = ty->As<sem::Array>();
+ arr && arr->IsRuntimeSized()) {
+ // Wrap runtime-sized arrays in structures, so that we can declare
+ // pointers to them. Ideally we'd just emit the array itself as a
+ // pointer, but this is not representable in Tint's AST.
+ CloneStructTypes(ty);
+ auto* wrapper = ctx.dst->Structure(
+ ctx.dst->Sym(),
+ {ctx.dst->Member(kWrappedArrayMemberName, param_type)});
+ param_type = ctx.dst->ty.Of(wrapper);
+ is_wrapped = true;
+ }
+
+ param_type = ctx.dst->ty.pointer(
+ param_type, sc, var->Declaration()->declared_access);
auto* param =
ctx.dst->Param(new_var_symbol, param_type, attributes);
ctx.InsertFront(func_ast->params, param);
@@ -283,6 +301,10 @@
expr = ctx.dst->Deref(expr);
}
+ if (is_wrapped) {
+ // Get the member from the wrapper structure.
+ expr = ctx.dst->MemberAccessor(expr, kWrappedArrayMemberName);
+ }
ctx.Replace(user->Declaration(), expr);
}
}
diff --git a/src/transform/module_scope_var_to_entry_point_param_test.cc b/src/transform/module_scope_var_to_entry_point_param_test.cc
index a1620e7..0f7d518 100644
--- a/src/transform/module_scope_var_to_entry_point_param_test.cc
+++ b/src/transform/module_scope_var_to_entry_point_param_test.cc
@@ -232,6 +232,99 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(ModuleScopeVarToEntryPointParamTest, Buffer_RuntimeArray) {
+ auto* src = R"(
+[[group(0), binding(0)]]
+var<storage> buffer : array<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = buffer[0];
+}
+)";
+
+ auto* expect = R"(
+struct tint_symbol_1 {
+ arr : array<f32>;
+};
+
+[[stage(compute), workgroup_size(1)]]
+fn main([[group(0), binding(0), internal(disable_validation__entry_point_parameter), internal(disable_validation__ignore_storage_class)]] tint_symbol : ptr<storage, tint_symbol_1>) {
+ _ = (*(tint_symbol)).arr[0];
+}
+)";
+
+ auto got = Run<ModuleScopeVarToEntryPointParam>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(ModuleScopeVarToEntryPointParamTest, Buffer_RuntimeArray_Alias) {
+ auto* src = R"(
+type myarray = array<f32>;
+
+[[group(0), binding(0)]]
+var<storage> buffer : myarray;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = buffer[0];
+}
+)";
+
+ auto* expect = R"(
+struct tint_symbol_1 {
+ arr : array<f32>;
+};
+
+type myarray = array<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main([[group(0), binding(0), internal(disable_validation__entry_point_parameter), internal(disable_validation__ignore_storage_class)]] tint_symbol : ptr<storage, tint_symbol_1>) {
+ _ = (*(tint_symbol)).arr[0];
+}
+)";
+
+ auto got = Run<ModuleScopeVarToEntryPointParam>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(ModuleScopeVarToEntryPointParamTest, Buffer_ArrayOfStruct) {
+ auto* src = R"(
+struct S {
+ f : f32;
+};
+
+[[group(0), binding(0)]]
+var<storage> buffer : array<S>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = buffer[0];
+}
+)";
+
+ auto* expect = R"(
+struct S {
+ f : f32;
+};
+
+struct tint_symbol_1 {
+ arr : array<S>;
+};
+
+[[stage(compute), workgroup_size(1)]]
+fn main([[group(0), binding(0), internal(disable_validation__entry_point_parameter), internal(disable_validation__ignore_storage_class)]] tint_symbol : ptr<storage, tint_symbol_1>) {
+ _ = (*(tint_symbol)).arr[0];
+}
+)";
+
+ auto got = Run<ModuleScopeVarToEntryPointParam>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(ModuleScopeVarToEntryPointParamTest, Buffers_FunctionCalls) {
auto* src = R"(
struct S {
diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index 76dc529..bae3044 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -2869,14 +2869,6 @@
auto* decl = var->Declaration();
auto binding_point = decl->BindingPoint();
auto* type = var->Type()->UnwrapRef();
-
- auto* str = type->As<sem::Struct>();
- if (!str) {
- // https://www.w3.org/TR/WGSL/#module-scope-variables
- TINT_ICE(Writer, diagnostics_)
- << "variables with uniform storage must be structure";
- }
-
auto name = builder_.Symbols().NameFor(decl->symbol);
line() << "cbuffer cbuffer_" << name << RegisterAndSpace('b', binding_point)
<< " {";
@@ -3513,13 +3505,7 @@
out << "ByteAddressBuffer";
return true;
case ast::StorageClass::kUniform: {
- auto* str = type->As<sem::Struct>();
- if (!str) {
- // https://www.w3.org/TR/WGSL/#module-scope-variables
- TINT_ICE(Writer, diagnostics_)
- << "variables with uniform storage must be structure";
- }
- auto array_length = (str->Size() + 15) / 16;
+ auto array_length = (type->Size() + 15) / 16;
out << "uint4 " << name << "[" << array_length << "]";
if (name_printed) {
*name_printed = true;
diff --git a/test/buffer/storage/types/array.wgsl b/test/buffer/storage/types/array.wgsl
new file mode 100644
index 0000000..6105a26
--- /dev/null
+++ b/test/buffer/storage/types/array.wgsl
@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : array<f32, 4>;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : array<f32, 4>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ out = in;
+}
diff --git a/test/buffer/storage/types/array.wgsl.expected.hlsl b/test/buffer/storage/types/array.wgsl.expected.hlsl
new file mode 100644
index 0000000..4b02ccd
--- /dev/null
+++ b/test/buffer/storage/types/array.wgsl.expected.hlsl
@@ -0,0 +1,28 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+void tint_symbol_2(RWByteAddressBuffer buffer, uint offset, float value[4]) {
+ float array[4] = value;
+ {
+ [loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+ buffer.Store((offset + (i * 4u)), asuint(array[i]));
+ }
+ }
+}
+
+typedef float tint_symbol_4_ret[4];
+tint_symbol_4_ret tint_symbol_4(ByteAddressBuffer buffer, uint offset) {
+ float arr[4] = (float[4])0;
+ {
+ [loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
+ arr[i_1] = asfloat(buffer.Load((offset + (i_1 * 4u))));
+ }
+ }
+ return arr;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ tint_symbol_2(tint_symbol_1, 0u, tint_symbol_4(tint_symbol, 0u));
+ return;
+}
diff --git a/test/buffer/storage/types/array.wgsl.expected.msl b/test/buffer/storage/types/array.wgsl.expected.msl
new file mode 100644
index 0000000..a95963f
--- /dev/null
+++ b/test/buffer/storage/types/array.wgsl.expected.msl
@@ -0,0 +1,12 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_array_wrapper {
+ /* 0x0000 */ float arr[4];
+};
+
+kernel void tint_symbol(device tint_array_wrapper* tint_symbol_1 [[buffer(1)]], const device tint_array_wrapper* tint_symbol_2 [[buffer(0)]]) {
+ *(tint_symbol_1) = *(tint_symbol_2);
+ return;
+}
+
diff --git a/test/buffer/storage/types/array.wgsl.expected.spvasm b/test/buffer/storage/types/array.wgsl.expected.spvasm
new file mode 100644
index 0000000..65f933d
--- /dev/null
+++ b/test/buffer/storage/types/array.wgsl.expected.spvasm
@@ -0,0 +1,42 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 18
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %in_block "in_block"
+ OpMemberName %in_block 0 "inner"
+ OpName %in "in"
+ OpName %out "out"
+ OpName %main "main"
+ OpDecorate %in_block Block
+ OpMemberDecorate %in_block 0 Offset 0
+ OpDecorate %_arr_float_uint_4 ArrayStride 4
+ OpDecorate %in NonWritable
+ OpDecorate %in DescriptorSet 0
+ OpDecorate %in Binding 0
+ OpDecorate %out DescriptorSet 0
+ OpDecorate %out Binding 1
+ %float = OpTypeFloat 32
+ %uint = OpTypeInt 32 0
+ %uint_4 = OpConstant %uint 4
+%_arr_float_uint_4 = OpTypeArray %float %uint_4
+ %in_block = OpTypeStruct %_arr_float_uint_4
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+ %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+ %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+ %void = OpTypeVoid
+ %9 = OpTypeFunction %void
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer__arr_float_uint_4 = OpTypePointer StorageBuffer %_arr_float_uint_4
+ %main = OpFunction %void None %9
+ %12 = OpLabel
+ %15 = OpAccessChain %_ptr_StorageBuffer__arr_float_uint_4 %out %uint_0
+ %16 = OpAccessChain %_ptr_StorageBuffer__arr_float_uint_4 %in %uint_0
+ %17 = OpLoad %_arr_float_uint_4 %16
+ OpStore %15 %17
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/storage/types/array.wgsl.expected.wgsl b/test/buffer/storage/types/array.wgsl.expected.wgsl
new file mode 100644
index 0000000..0aaacae
--- /dev/null
+++ b/test/buffer/storage/types/array.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : array<f32, 4>;
+
+[[group(0), binding(1)]] var<storage, read_write> out : array<f32, 4>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ out = in;
+}
diff --git a/test/buffer/storage/types/f32.wgsl b/test/buffer/storage/types/f32.wgsl
new file mode 100644
index 0000000..d66cdb0
--- /dev/null
+++ b/test/buffer/storage/types/f32.wgsl
@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : f32;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : f32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ out = in;
+}
diff --git a/test/buffer/storage/types/f32.wgsl.expected.hlsl b/test/buffer/storage/types/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..e6dca9e
--- /dev/null
+++ b/test/buffer/storage/types/f32.wgsl.expected.hlsl
@@ -0,0 +1,8 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+ tint_symbol_1.Store(0u, asuint(asfloat(tint_symbol.Load(0u))));
+ return;
+}
diff --git a/test/buffer/storage/types/f32.wgsl.expected.msl b/test/buffer/storage/types/f32.wgsl.expected.msl
new file mode 100644
index 0000000..033b5e2
--- /dev/null
+++ b/test/buffer/storage/types/f32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device float* tint_symbol_1 [[buffer(1)]], const device float* tint_symbol_2 [[buffer(0)]]) {
+ *(tint_symbol_1) = *(tint_symbol_2);
+ return;
+}
+
diff --git a/test/buffer/storage/types/f32.wgsl.expected.spvasm b/test/buffer/storage/types/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..dea11b4
--- /dev/null
+++ b/test/buffer/storage/types/f32.wgsl.expected.spvasm
@@ -0,0 +1,39 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %in_block "in_block"
+ OpMemberName %in_block 0 "inner"
+ OpName %in "in"
+ OpName %out "out"
+ OpName %main "main"
+ OpDecorate %in_block Block
+ OpMemberDecorate %in_block 0 Offset 0
+ OpDecorate %in NonWritable
+ OpDecorate %in DescriptorSet 0
+ OpDecorate %in Binding 0
+ OpDecorate %out DescriptorSet 0
+ OpDecorate %out Binding 1
+ %float = OpTypeFloat 32
+ %in_block = OpTypeStruct %float
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+ %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+ %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+ %void = OpTypeVoid
+ %6 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+ %main = OpFunction %void None %6
+ %9 = OpLabel
+ %13 = OpAccessChain %_ptr_StorageBuffer_float %out %uint_0
+ %14 = OpAccessChain %_ptr_StorageBuffer_float %in %uint_0
+ %15 = OpLoad %float %14
+ OpStore %13 %15
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/storage/types/f32.wgsl.expected.wgsl b/test/buffer/storage/types/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..a385a75
--- /dev/null
+++ b/test/buffer/storage/types/f32.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : f32;
+
+[[group(0), binding(1)]] var<storage, read_write> out : f32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ out = in;
+}
diff --git a/test/buffer/storage/types/i32.wgsl b/test/buffer/storage/types/i32.wgsl
new file mode 100644
index 0000000..2ffd35f
--- /dev/null
+++ b/test/buffer/storage/types/i32.wgsl
@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : i32;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : i32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ out = in;
+}
diff --git a/test/buffer/storage/types/i32.wgsl.expected.hlsl b/test/buffer/storage/types/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..3c53558
--- /dev/null
+++ b/test/buffer/storage/types/i32.wgsl.expected.hlsl
@@ -0,0 +1,8 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+ tint_symbol_1.Store(0u, asuint(asint(tint_symbol.Load(0u))));
+ return;
+}
diff --git a/test/buffer/storage/types/i32.wgsl.expected.msl b/test/buffer/storage/types/i32.wgsl.expected.msl
new file mode 100644
index 0000000..1de9e03
--- /dev/null
+++ b/test/buffer/storage/types/i32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device int* tint_symbol_1 [[buffer(1)]], const device int* tint_symbol_2 [[buffer(0)]]) {
+ *(tint_symbol_1) = *(tint_symbol_2);
+ return;
+}
+
diff --git a/test/buffer/storage/types/i32.wgsl.expected.spvasm b/test/buffer/storage/types/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..4abab36
--- /dev/null
+++ b/test/buffer/storage/types/i32.wgsl.expected.spvasm
@@ -0,0 +1,39 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %in_block "in_block"
+ OpMemberName %in_block 0 "inner"
+ OpName %in "in"
+ OpName %out "out"
+ OpName %main "main"
+ OpDecorate %in_block Block
+ OpMemberDecorate %in_block 0 Offset 0
+ OpDecorate %in NonWritable
+ OpDecorate %in DescriptorSet 0
+ OpDecorate %in Binding 0
+ OpDecorate %out DescriptorSet 0
+ OpDecorate %out Binding 1
+ %int = OpTypeInt 32 1
+ %in_block = OpTypeStruct %int
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+ %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+ %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+ %void = OpTypeVoid
+ %6 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+ %main = OpFunction %void None %6
+ %9 = OpLabel
+ %13 = OpAccessChain %_ptr_StorageBuffer_int %out %uint_0
+ %14 = OpAccessChain %_ptr_StorageBuffer_int %in %uint_0
+ %15 = OpLoad %int %14
+ OpStore %13 %15
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/storage/types/i32.wgsl.expected.wgsl b/test/buffer/storage/types/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..3424f2e
--- /dev/null
+++ b/test/buffer/storage/types/i32.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : i32;
+
+[[group(0), binding(1)]] var<storage, read_write> out : i32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ out = in;
+}
diff --git a/test/buffer/storage/types/mat2x2.wgsl b/test/buffer/storage/types/mat2x2.wgsl
new file mode 100644
index 0000000..60ff548
--- /dev/null
+++ b/test/buffer/storage/types/mat2x2.wgsl
@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : mat2x2<f32>;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : mat2x2<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ out = in;
+}
diff --git a/test/buffer/storage/types/mat2x2.wgsl.expected.hlsl b/test/buffer/storage/types/mat2x2.wgsl.expected.hlsl
new file mode 100644
index 0000000..a1fc1b6
--- /dev/null
+++ b/test/buffer/storage/types/mat2x2.wgsl.expected.hlsl
@@ -0,0 +1,17 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+void tint_symbol_2(RWByteAddressBuffer buffer, uint offset, float2x2 value) {
+ buffer.Store2((offset + 0u), asuint(value[0u]));
+ buffer.Store2((offset + 8u), asuint(value[1u]));
+}
+
+float2x2 tint_symbol_4(ByteAddressBuffer buffer, uint offset) {
+ return float2x2(asfloat(buffer.Load2((offset + 0u))), asfloat(buffer.Load2((offset + 8u))));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ tint_symbol_2(tint_symbol_1, 0u, tint_symbol_4(tint_symbol, 0u));
+ return;
+}
diff --git a/test/buffer/storage/types/mat2x2.wgsl.expected.msl b/test/buffer/storage/types/mat2x2.wgsl.expected.msl
new file mode 100644
index 0000000..c9ffff7
--- /dev/null
+++ b/test/buffer/storage/types/mat2x2.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device float2x2* tint_symbol_1 [[buffer(1)]], const device float2x2* tint_symbol_2 [[buffer(0)]]) {
+ *(tint_symbol_1) = *(tint_symbol_2);
+ return;
+}
+
diff --git a/test/buffer/storage/types/mat2x2.wgsl.expected.spvasm b/test/buffer/storage/types/mat2x2.wgsl.expected.spvasm
new file mode 100644
index 0000000..a372c43
--- /dev/null
+++ b/test/buffer/storage/types/mat2x2.wgsl.expected.spvasm
@@ -0,0 +1,43 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 18
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %in_block "in_block"
+ OpMemberName %in_block 0 "inner"
+ OpName %in "in"
+ OpName %out "out"
+ OpName %main "main"
+ OpDecorate %in_block Block
+ OpMemberDecorate %in_block 0 Offset 0
+ OpMemberDecorate %in_block 0 ColMajor
+ OpMemberDecorate %in_block 0 MatrixStride 8
+ OpDecorate %in NonWritable
+ OpDecorate %in DescriptorSet 0
+ OpDecorate %in Binding 0
+ OpDecorate %out DescriptorSet 0
+ OpDecorate %out Binding 1
+ %float = OpTypeFloat 32
+ %v2float = OpTypeVector %float 2
+%mat2v2float = OpTypeMatrix %v2float 2
+ %in_block = OpTypeStruct %mat2v2float
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+ %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+ %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+ %void = OpTypeVoid
+ %8 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_mat2v2float = OpTypePointer StorageBuffer %mat2v2float
+ %main = OpFunction %void None %8
+ %11 = OpLabel
+ %15 = OpAccessChain %_ptr_StorageBuffer_mat2v2float %out %uint_0
+ %16 = OpAccessChain %_ptr_StorageBuffer_mat2v2float %in %uint_0
+ %17 = OpLoad %mat2v2float %16
+ OpStore %15 %17
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/storage/types/mat2x2.wgsl.expected.wgsl b/test/buffer/storage/types/mat2x2.wgsl.expected.wgsl
new file mode 100644
index 0000000..8d2cbae
--- /dev/null
+++ b/test/buffer/storage/types/mat2x2.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : mat2x2<f32>;
+
+[[group(0), binding(1)]] var<storage, read_write> out : mat2x2<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ out = in;
+}
diff --git a/test/buffer/storage/types/mat2x3.wgsl b/test/buffer/storage/types/mat2x3.wgsl
new file mode 100644
index 0000000..151db83
--- /dev/null
+++ b/test/buffer/storage/types/mat2x3.wgsl
@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : mat2x3<f32>;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : mat2x3<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ out = in;
+}
diff --git a/test/buffer/storage/types/mat2x3.wgsl.expected.hlsl b/test/buffer/storage/types/mat2x3.wgsl.expected.hlsl
new file mode 100644
index 0000000..baff4c5
--- /dev/null
+++ b/test/buffer/storage/types/mat2x3.wgsl.expected.hlsl
@@ -0,0 +1,17 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+void tint_symbol_2(RWByteAddressBuffer buffer, uint offset, float2x3 value) {
+ buffer.Store3((offset + 0u), asuint(value[0u]));
+ buffer.Store3((offset + 16u), asuint(value[1u]));
+}
+
+float2x3 tint_symbol_4(ByteAddressBuffer buffer, uint offset) {
+ return float2x3(asfloat(buffer.Load3((offset + 0u))), asfloat(buffer.Load3((offset + 16u))));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ tint_symbol_2(tint_symbol_1, 0u, tint_symbol_4(tint_symbol, 0u));
+ return;
+}
diff --git a/test/buffer/storage/types/mat2x3.wgsl.expected.msl b/test/buffer/storage/types/mat2x3.wgsl.expected.msl
new file mode 100644
index 0000000..1b704d4
--- /dev/null
+++ b/test/buffer/storage/types/mat2x3.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device float2x3* tint_symbol_1 [[buffer(1)]], const device float2x3* tint_symbol_2 [[buffer(0)]]) {
+ *(tint_symbol_1) = *(tint_symbol_2);
+ return;
+}
+
diff --git a/test/buffer/storage/types/mat2x3.wgsl.expected.spvasm b/test/buffer/storage/types/mat2x3.wgsl.expected.spvasm
new file mode 100644
index 0000000..31ddb99
--- /dev/null
+++ b/test/buffer/storage/types/mat2x3.wgsl.expected.spvasm
@@ -0,0 +1,43 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 18
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %in_block "in_block"
+ OpMemberName %in_block 0 "inner"
+ OpName %in "in"
+ OpName %out "out"
+ OpName %main "main"
+ OpDecorate %in_block Block
+ OpMemberDecorate %in_block 0 Offset 0
+ OpMemberDecorate %in_block 0 ColMajor
+ OpMemberDecorate %in_block 0 MatrixStride 16
+ OpDecorate %in NonWritable
+ OpDecorate %in DescriptorSet 0
+ OpDecorate %in Binding 0
+ OpDecorate %out DescriptorSet 0
+ OpDecorate %out Binding 1
+ %float = OpTypeFloat 32
+ %v3float = OpTypeVector %float 3
+%mat2v3float = OpTypeMatrix %v3float 2
+ %in_block = OpTypeStruct %mat2v3float
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+ %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+ %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+ %void = OpTypeVoid
+ %8 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_mat2v3float = OpTypePointer StorageBuffer %mat2v3float
+ %main = OpFunction %void None %8
+ %11 = OpLabel
+ %15 = OpAccessChain %_ptr_StorageBuffer_mat2v3float %out %uint_0
+ %16 = OpAccessChain %_ptr_StorageBuffer_mat2v3float %in %uint_0
+ %17 = OpLoad %mat2v3float %16
+ OpStore %15 %17
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/storage/types/mat2x3.wgsl.expected.wgsl b/test/buffer/storage/types/mat2x3.wgsl.expected.wgsl
new file mode 100644
index 0000000..ab3a5fd
--- /dev/null
+++ b/test/buffer/storage/types/mat2x3.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : mat2x3<f32>;
+
+[[group(0), binding(1)]] var<storage, read_write> out : mat2x3<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ out = in;
+}
diff --git a/test/buffer/storage/types/mat3x2.wgsl b/test/buffer/storage/types/mat3x2.wgsl
new file mode 100644
index 0000000..d77ac81
--- /dev/null
+++ b/test/buffer/storage/types/mat3x2.wgsl
@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : mat3x2<f32>;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : mat3x2<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ out = in;
+}
diff --git a/test/buffer/storage/types/mat3x2.wgsl.expected.hlsl b/test/buffer/storage/types/mat3x2.wgsl.expected.hlsl
new file mode 100644
index 0000000..31d606b
--- /dev/null
+++ b/test/buffer/storage/types/mat3x2.wgsl.expected.hlsl
@@ -0,0 +1,18 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+void tint_symbol_2(RWByteAddressBuffer buffer, uint offset, float3x2 value) {
+ buffer.Store2((offset + 0u), asuint(value[0u]));
+ buffer.Store2((offset + 8u), asuint(value[1u]));
+ buffer.Store2((offset + 16u), asuint(value[2u]));
+}
+
+float3x2 tint_symbol_4(ByteAddressBuffer buffer, uint offset) {
+ return float3x2(asfloat(buffer.Load2((offset + 0u))), asfloat(buffer.Load2((offset + 8u))), asfloat(buffer.Load2((offset + 16u))));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ tint_symbol_2(tint_symbol_1, 0u, tint_symbol_4(tint_symbol, 0u));
+ return;
+}
diff --git a/test/buffer/storage/types/mat3x2.wgsl.expected.msl b/test/buffer/storage/types/mat3x2.wgsl.expected.msl
new file mode 100644
index 0000000..b8765f0
--- /dev/null
+++ b/test/buffer/storage/types/mat3x2.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device float3x2* tint_symbol_1 [[buffer(1)]], const device float3x2* tint_symbol_2 [[buffer(0)]]) {
+ *(tint_symbol_1) = *(tint_symbol_2);
+ return;
+}
+
diff --git a/test/buffer/storage/types/mat3x2.wgsl.expected.spvasm b/test/buffer/storage/types/mat3x2.wgsl.expected.spvasm
new file mode 100644
index 0000000..c352e18
--- /dev/null
+++ b/test/buffer/storage/types/mat3x2.wgsl.expected.spvasm
@@ -0,0 +1,43 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 18
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %in_block "in_block"
+ OpMemberName %in_block 0 "inner"
+ OpName %in "in"
+ OpName %out "out"
+ OpName %main "main"
+ OpDecorate %in_block Block
+ OpMemberDecorate %in_block 0 Offset 0
+ OpMemberDecorate %in_block 0 ColMajor
+ OpMemberDecorate %in_block 0 MatrixStride 8
+ OpDecorate %in NonWritable
+ OpDecorate %in DescriptorSet 0
+ OpDecorate %in Binding 0
+ OpDecorate %out DescriptorSet 0
+ OpDecorate %out Binding 1
+ %float = OpTypeFloat 32
+ %v2float = OpTypeVector %float 2
+%mat3v2float = OpTypeMatrix %v2float 3
+ %in_block = OpTypeStruct %mat3v2float
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+ %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+ %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+ %void = OpTypeVoid
+ %8 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_mat3v2float = OpTypePointer StorageBuffer %mat3v2float
+ %main = OpFunction %void None %8
+ %11 = OpLabel
+ %15 = OpAccessChain %_ptr_StorageBuffer_mat3v2float %out %uint_0
+ %16 = OpAccessChain %_ptr_StorageBuffer_mat3v2float %in %uint_0
+ %17 = OpLoad %mat3v2float %16
+ OpStore %15 %17
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/storage/types/mat3x2.wgsl.expected.wgsl b/test/buffer/storage/types/mat3x2.wgsl.expected.wgsl
new file mode 100644
index 0000000..4c87f4a
--- /dev/null
+++ b/test/buffer/storage/types/mat3x2.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : mat3x2<f32>;
+
+[[group(0), binding(1)]] var<storage, read_write> out : mat3x2<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ out = in;
+}
diff --git a/test/buffer/storage/types/mat4x4.wgsl b/test/buffer/storage/types/mat4x4.wgsl
new file mode 100644
index 0000000..bf18a25
--- /dev/null
+++ b/test/buffer/storage/types/mat4x4.wgsl
@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : mat4x4<f32>;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : mat4x4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ out = in;
+}
diff --git a/test/buffer/storage/types/mat4x4.wgsl.expected.hlsl b/test/buffer/storage/types/mat4x4.wgsl.expected.hlsl
new file mode 100644
index 0000000..e818c09
--- /dev/null
+++ b/test/buffer/storage/types/mat4x4.wgsl.expected.hlsl
@@ -0,0 +1,19 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+void tint_symbol_2(RWByteAddressBuffer buffer, uint offset, float4x4 value) {
+ buffer.Store4((offset + 0u), asuint(value[0u]));
+ buffer.Store4((offset + 16u), asuint(value[1u]));
+ buffer.Store4((offset + 32u), asuint(value[2u]));
+ buffer.Store4((offset + 48u), asuint(value[3u]));
+}
+
+float4x4 tint_symbol_4(ByteAddressBuffer buffer, uint offset) {
+ return float4x4(asfloat(buffer.Load4((offset + 0u))), asfloat(buffer.Load4((offset + 16u))), asfloat(buffer.Load4((offset + 32u))), asfloat(buffer.Load4((offset + 48u))));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ tint_symbol_2(tint_symbol_1, 0u, tint_symbol_4(tint_symbol, 0u));
+ return;
+}
diff --git a/test/buffer/storage/types/mat4x4.wgsl.expected.msl b/test/buffer/storage/types/mat4x4.wgsl.expected.msl
new file mode 100644
index 0000000..6b33874
--- /dev/null
+++ b/test/buffer/storage/types/mat4x4.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device float4x4* tint_symbol_1 [[buffer(1)]], const device float4x4* tint_symbol_2 [[buffer(0)]]) {
+ *(tint_symbol_1) = *(tint_symbol_2);
+ return;
+}
+
diff --git a/test/buffer/storage/types/mat4x4.wgsl.expected.spvasm b/test/buffer/storage/types/mat4x4.wgsl.expected.spvasm
new file mode 100644
index 0000000..5123414
--- /dev/null
+++ b/test/buffer/storage/types/mat4x4.wgsl.expected.spvasm
@@ -0,0 +1,43 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 18
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %in_block "in_block"
+ OpMemberName %in_block 0 "inner"
+ OpName %in "in"
+ OpName %out "out"
+ OpName %main "main"
+ OpDecorate %in_block Block
+ OpMemberDecorate %in_block 0 Offset 0
+ OpMemberDecorate %in_block 0 ColMajor
+ OpMemberDecorate %in_block 0 MatrixStride 16
+ OpDecorate %in NonWritable
+ OpDecorate %in DescriptorSet 0
+ OpDecorate %in Binding 0
+ OpDecorate %out DescriptorSet 0
+ OpDecorate %out Binding 1
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+%mat4v4float = OpTypeMatrix %v4float 4
+ %in_block = OpTypeStruct %mat4v4float
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+ %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+ %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+ %void = OpTypeVoid
+ %8 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_mat4v4float = OpTypePointer StorageBuffer %mat4v4float
+ %main = OpFunction %void None %8
+ %11 = OpLabel
+ %15 = OpAccessChain %_ptr_StorageBuffer_mat4v4float %out %uint_0
+ %16 = OpAccessChain %_ptr_StorageBuffer_mat4v4float %in %uint_0
+ %17 = OpLoad %mat4v4float %16
+ OpStore %15 %17
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/storage/types/mat4x4.wgsl.expected.wgsl b/test/buffer/storage/types/mat4x4.wgsl.expected.wgsl
new file mode 100644
index 0000000..7c1ebe9
--- /dev/null
+++ b/test/buffer/storage/types/mat4x4.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : mat4x4<f32>;
+
+[[group(0), binding(1)]] var<storage, read_write> out : mat4x4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ out = in;
+}
diff --git a/test/buffer/storage/types/runtime_array.wgsl b/test/buffer/storage/types/runtime_array.wgsl
new file mode 100644
index 0000000..f403fd3
--- /dev/null
+++ b/test/buffer/storage/types/runtime_array.wgsl
@@ -0,0 +1,14 @@
+struct S {
+ f : f32;
+};
+
+[[group(0), binding(0)]]
+var<storage, read> in : array<S>;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : array<S>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ out[0] = in[0];
+}
diff --git a/test/buffer/storage/types/runtime_array.wgsl.expected.hlsl b/test/buffer/storage/types/runtime_array.wgsl.expected.hlsl
new file mode 100644
index 0000000..035878a
--- /dev/null
+++ b/test/buffer/storage/types/runtime_array.wgsl.expected.hlsl
@@ -0,0 +1,21 @@
+struct S {
+ float f;
+};
+
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+void tint_symbol_2(RWByteAddressBuffer buffer, uint offset, S value) {
+ buffer.Store((offset + 0u), asuint(value.f));
+}
+
+S tint_symbol_4(ByteAddressBuffer buffer, uint offset) {
+ const S tint_symbol_6 = {asfloat(buffer.Load((offset + 0u)))};
+ return tint_symbol_6;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ tint_symbol_2(tint_symbol_1, (4u * uint(0)), tint_symbol_4(tint_symbol, (4u * uint(0))));
+ return;
+}
diff --git a/test/buffer/storage/types/runtime_array.wgsl.expected.msl b/test/buffer/storage/types/runtime_array.wgsl.expected.msl
new file mode 100644
index 0000000..ff6a283
--- /dev/null
+++ b/test/buffer/storage/types/runtime_array.wgsl.expected.msl
@@ -0,0 +1,18 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+ /* 0x0000 */ float f;
+};
+struct tint_symbol_2 {
+ /* 0x0000 */ S arr[1];
+};
+struct tint_symbol_4 {
+ /* 0x0000 */ S arr[1];
+};
+
+kernel void tint_symbol(device tint_symbol_2* tint_symbol_1 [[buffer(1)]], const device tint_symbol_4* tint_symbol_3 [[buffer(0)]]) {
+ (*(tint_symbol_1)).arr[0] = (*(tint_symbol_3)).arr[0];
+ return;
+}
+
diff --git a/test/buffer/storage/types/runtime_array.wgsl.expected.spvasm b/test/buffer/storage/types/runtime_array.wgsl.expected.spvasm
new file mode 100644
index 0000000..5a63e67
--- /dev/null
+++ b/test/buffer/storage/types/runtime_array.wgsl.expected.spvasm
@@ -0,0 +1,47 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 20
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %in_block "in_block"
+ OpMemberName %in_block 0 "inner"
+ OpName %S "S"
+ OpMemberName %S 0 "f"
+ OpName %in "in"
+ OpName %out "out"
+ OpName %main "main"
+ OpDecorate %in_block Block
+ OpMemberDecorate %in_block 0 Offset 0
+ OpMemberDecorate %S 0 Offset 0
+ OpDecorate %_runtimearr_S ArrayStride 4
+ OpDecorate %in NonWritable
+ OpDecorate %in DescriptorSet 0
+ OpDecorate %in Binding 0
+ OpDecorate %out DescriptorSet 0
+ OpDecorate %out Binding 1
+ %float = OpTypeFloat 32
+ %S = OpTypeStruct %float
+%_runtimearr_S = OpTypeRuntimeArray %S
+ %in_block = OpTypeStruct %_runtimearr_S
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+ %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+ %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+ %void = OpTypeVoid
+ %8 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+ %int = OpTypeInt 32 1
+ %int_0 = OpConstant %int 0
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+ %main = OpFunction %void None %8
+ %11 = OpLabel
+ %17 = OpAccessChain %_ptr_StorageBuffer_S %out %uint_0 %int_0
+ %18 = OpAccessChain %_ptr_StorageBuffer_S %in %uint_0 %int_0
+ %19 = OpLoad %S %18
+ OpStore %17 %19
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/storage/types/runtime_array.wgsl.expected.wgsl b/test/buffer/storage/types/runtime_array.wgsl.expected.wgsl
new file mode 100644
index 0000000..418c67d
--- /dev/null
+++ b/test/buffer/storage/types/runtime_array.wgsl.expected.wgsl
@@ -0,0 +1,12 @@
+struct S {
+ f : f32;
+};
+
+[[group(0), binding(0)]] var<storage, read> in : array<S>;
+
+[[group(0), binding(1)]] var<storage, read_write> out : array<S>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ out[0] = in[0];
+}
diff --git a/test/buffer/storage/types/struct.wgsl b/test/buffer/storage/types/struct.wgsl
new file mode 100644
index 0000000..42acfbc
--- /dev/null
+++ b/test/buffer/storage/types/struct.wgsl
@@ -0,0 +1,17 @@
+struct Inner {
+ f : f32;
+};
+struct S {
+ inner : Inner;
+};
+
+[[group(0), binding(0)]]
+var<storage, read> in : S;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ out = in;
+}
diff --git a/test/buffer/storage/types/struct.wgsl.expected.hlsl b/test/buffer/storage/types/struct.wgsl.expected.hlsl
new file mode 100644
index 0000000..05be600
--- /dev/null
+++ b/test/buffer/storage/types/struct.wgsl.expected.hlsl
@@ -0,0 +1,33 @@
+struct Inner {
+ float f;
+};
+struct S {
+ Inner inner;
+};
+
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+void tint_symbol_3(RWByteAddressBuffer buffer, uint offset, Inner value) {
+ buffer.Store((offset + 0u), asuint(value.f));
+}
+
+void tint_symbol_2(RWByteAddressBuffer buffer, uint offset, S value) {
+ tint_symbol_3(buffer, (offset + 0u), value.inner);
+}
+
+Inner tint_symbol_6(ByteAddressBuffer buffer, uint offset) {
+ const Inner tint_symbol_8 = {asfloat(buffer.Load((offset + 0u)))};
+ return tint_symbol_8;
+}
+
+S tint_symbol_5(ByteAddressBuffer buffer, uint offset) {
+ const S tint_symbol_9 = {tint_symbol_6(buffer, (offset + 0u))};
+ return tint_symbol_9;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ tint_symbol_2(tint_symbol_1, 0u, tint_symbol_5(tint_symbol, 0u));
+ return;
+}
diff --git a/test/buffer/storage/types/struct.wgsl.expected.msl b/test/buffer/storage/types/struct.wgsl.expected.msl
new file mode 100644
index 0000000..d5c7427
--- /dev/null
+++ b/test/buffer/storage/types/struct.wgsl.expected.msl
@@ -0,0 +1,15 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Inner {
+ /* 0x0000 */ float f;
+};
+struct S {
+ /* 0x0000 */ Inner inner;
+};
+
+kernel void tint_symbol(device S* tint_symbol_1 [[buffer(0)]], const device S* tint_symbol_2 [[buffer(1)]]) {
+ *(tint_symbol_1) = *(tint_symbol_2);
+ return;
+}
+
diff --git a/test/buffer/storage/types/struct.wgsl.expected.spvasm b/test/buffer/storage/types/struct.wgsl.expected.spvasm
new file mode 100644
index 0000000..058a3e5
--- /dev/null
+++ b/test/buffer/storage/types/struct.wgsl.expected.spvasm
@@ -0,0 +1,38 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 12
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %S "S"
+ OpMemberName %S 0 "inner"
+ OpName %Inner "Inner"
+ OpMemberName %Inner 0 "f"
+ OpName %in "in"
+ OpName %out "out"
+ OpName %main "main"
+ OpDecorate %S Block
+ OpMemberDecorate %S 0 Offset 0
+ OpMemberDecorate %Inner 0 Offset 0
+ OpDecorate %in NonWritable
+ OpDecorate %in DescriptorSet 0
+ OpDecorate %in Binding 0
+ OpDecorate %out DescriptorSet 0
+ OpDecorate %out Binding 1
+ %float = OpTypeFloat 32
+ %Inner = OpTypeStruct %float
+ %S = OpTypeStruct %Inner
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+ %in = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+ %out = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+ %void = OpTypeVoid
+ %7 = OpTypeFunction %void
+ %main = OpFunction %void None %7
+ %10 = OpLabel
+ %11 = OpLoad %S %in
+ OpStore %out %11
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/storage/types/struct.wgsl.expected.wgsl b/test/buffer/storage/types/struct.wgsl.expected.wgsl
new file mode 100644
index 0000000..f86c26b
--- /dev/null
+++ b/test/buffer/storage/types/struct.wgsl.expected.wgsl
@@ -0,0 +1,16 @@
+struct Inner {
+ f : f32;
+};
+
+struct S {
+ inner : Inner;
+};
+
+[[group(0), binding(0)]] var<storage, read> in : S;
+
+[[group(0), binding(1)]] var<storage, read_write> out : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ out = in;
+}
diff --git a/test/buffer/storage/types/u32.wgsl b/test/buffer/storage/types/u32.wgsl
new file mode 100644
index 0000000..4664d29
--- /dev/null
+++ b/test/buffer/storage/types/u32.wgsl
@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : u32;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : u32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ out = in;
+}
diff --git a/test/buffer/storage/types/u32.wgsl.expected.hlsl b/test/buffer/storage/types/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..8574daf
--- /dev/null
+++ b/test/buffer/storage/types/u32.wgsl.expected.hlsl
@@ -0,0 +1,8 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+ tint_symbol_1.Store(0u, asuint(tint_symbol.Load(0u)));
+ return;
+}
diff --git a/test/buffer/storage/types/u32.wgsl.expected.msl b/test/buffer/storage/types/u32.wgsl.expected.msl
new file mode 100644
index 0000000..bb65310
--- /dev/null
+++ b/test/buffer/storage/types/u32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device uint* tint_symbol_1 [[buffer(1)]], const device uint* tint_symbol_2 [[buffer(0)]]) {
+ *(tint_symbol_1) = *(tint_symbol_2);
+ return;
+}
+
diff --git a/test/buffer/storage/types/u32.wgsl.expected.spvasm b/test/buffer/storage/types/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..201cc29
--- /dev/null
+++ b/test/buffer/storage/types/u32.wgsl.expected.spvasm
@@ -0,0 +1,38 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 15
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %in_block "in_block"
+ OpMemberName %in_block 0 "inner"
+ OpName %in "in"
+ OpName %out "out"
+ OpName %main "main"
+ OpDecorate %in_block Block
+ OpMemberDecorate %in_block 0 Offset 0
+ OpDecorate %in NonWritable
+ OpDecorate %in DescriptorSet 0
+ OpDecorate %in Binding 0
+ OpDecorate %out DescriptorSet 0
+ OpDecorate %out Binding 1
+ %uint = OpTypeInt 32 0
+ %in_block = OpTypeStruct %uint
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+ %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+ %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+ %void = OpTypeVoid
+ %6 = OpTypeFunction %void
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+ %main = OpFunction %void None %6
+ %9 = OpLabel
+ %12 = OpAccessChain %_ptr_StorageBuffer_uint %out %uint_0
+ %13 = OpAccessChain %_ptr_StorageBuffer_uint %in %uint_0
+ %14 = OpLoad %uint %13
+ OpStore %12 %14
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/storage/types/u32.wgsl.expected.wgsl b/test/buffer/storage/types/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..46dc9d3
--- /dev/null
+++ b/test/buffer/storage/types/u32.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : u32;
+
+[[group(0), binding(1)]] var<storage, read_write> out : u32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ out = in;
+}
diff --git a/test/buffer/storage/types/vec2.wgsl b/test/buffer/storage/types/vec2.wgsl
new file mode 100644
index 0000000..9841fd4
--- /dev/null
+++ b/test/buffer/storage/types/vec2.wgsl
@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : vec2<i32>;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : vec2<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ out = in;
+}
diff --git a/test/buffer/storage/types/vec2.wgsl.expected.hlsl b/test/buffer/storage/types/vec2.wgsl.expected.hlsl
new file mode 100644
index 0000000..3c5c7d7
--- /dev/null
+++ b/test/buffer/storage/types/vec2.wgsl.expected.hlsl
@@ -0,0 +1,8 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+ tint_symbol_1.Store2(0u, asuint(asint(tint_symbol.Load2(0u))));
+ return;
+}
diff --git a/test/buffer/storage/types/vec2.wgsl.expected.msl b/test/buffer/storage/types/vec2.wgsl.expected.msl
new file mode 100644
index 0000000..c730e05
--- /dev/null
+++ b/test/buffer/storage/types/vec2.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device int2* tint_symbol_1 [[buffer(1)]], const device int2* tint_symbol_2 [[buffer(0)]]) {
+ *(tint_symbol_1) = *(tint_symbol_2);
+ return;
+}
+
diff --git a/test/buffer/storage/types/vec2.wgsl.expected.spvasm b/test/buffer/storage/types/vec2.wgsl.expected.spvasm
new file mode 100644
index 0000000..85b79d3
--- /dev/null
+++ b/test/buffer/storage/types/vec2.wgsl.expected.spvasm
@@ -0,0 +1,40 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 17
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %in_block "in_block"
+ OpMemberName %in_block 0 "inner"
+ OpName %in "in"
+ OpName %out "out"
+ OpName %main "main"
+ OpDecorate %in_block Block
+ OpMemberDecorate %in_block 0 Offset 0
+ OpDecorate %in NonWritable
+ OpDecorate %in DescriptorSet 0
+ OpDecorate %in Binding 0
+ OpDecorate %out DescriptorSet 0
+ OpDecorate %out Binding 1
+ %int = OpTypeInt 32 1
+ %v2int = OpTypeVector %int 2
+ %in_block = OpTypeStruct %v2int
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+ %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+ %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+ %void = OpTypeVoid
+ %7 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v2int = OpTypePointer StorageBuffer %v2int
+ %main = OpFunction %void None %7
+ %10 = OpLabel
+ %14 = OpAccessChain %_ptr_StorageBuffer_v2int %out %uint_0
+ %15 = OpAccessChain %_ptr_StorageBuffer_v2int %in %uint_0
+ %16 = OpLoad %v2int %15
+ OpStore %14 %16
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/storage/types/vec2.wgsl.expected.wgsl b/test/buffer/storage/types/vec2.wgsl.expected.wgsl
new file mode 100644
index 0000000..a945b32
--- /dev/null
+++ b/test/buffer/storage/types/vec2.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : vec2<i32>;
+
+[[group(0), binding(1)]] var<storage, read_write> out : vec2<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ out = in;
+}
diff --git a/test/buffer/storage/types/vec3.wgsl b/test/buffer/storage/types/vec3.wgsl
new file mode 100644
index 0000000..5335059
--- /dev/null
+++ b/test/buffer/storage/types/vec3.wgsl
@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : vec3<u32>;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : vec3<u32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ out = in;
+}
diff --git a/test/buffer/storage/types/vec3.wgsl.expected.hlsl b/test/buffer/storage/types/vec3.wgsl.expected.hlsl
new file mode 100644
index 0000000..9af25f4
--- /dev/null
+++ b/test/buffer/storage/types/vec3.wgsl.expected.hlsl
@@ -0,0 +1,8 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+ tint_symbol_1.Store3(0u, asuint(tint_symbol.Load3(0u)));
+ return;
+}
diff --git a/test/buffer/storage/types/vec3.wgsl.expected.msl b/test/buffer/storage/types/vec3.wgsl.expected.msl
new file mode 100644
index 0000000..6d55796
--- /dev/null
+++ b/test/buffer/storage/types/vec3.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device uint3* tint_symbol_1 [[buffer(1)]], const device uint3* tint_symbol_2 [[buffer(0)]]) {
+ *(tint_symbol_1) = *(tint_symbol_2);
+ return;
+}
+
diff --git a/test/buffer/storage/types/vec3.wgsl.expected.spvasm b/test/buffer/storage/types/vec3.wgsl.expected.spvasm
new file mode 100644
index 0000000..70a2daf
--- /dev/null
+++ b/test/buffer/storage/types/vec3.wgsl.expected.spvasm
@@ -0,0 +1,39 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %in_block "in_block"
+ OpMemberName %in_block 0 "inner"
+ OpName %in "in"
+ OpName %out "out"
+ OpName %main "main"
+ OpDecorate %in_block Block
+ OpMemberDecorate %in_block 0 Offset 0
+ OpDecorate %in NonWritable
+ OpDecorate %in DescriptorSet 0
+ OpDecorate %in Binding 0
+ OpDecorate %out DescriptorSet 0
+ OpDecorate %out Binding 1
+ %uint = OpTypeInt 32 0
+ %v3uint = OpTypeVector %uint 3
+ %in_block = OpTypeStruct %v3uint
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+ %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+ %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+ %void = OpTypeVoid
+ %7 = OpTypeFunction %void
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v3uint = OpTypePointer StorageBuffer %v3uint
+ %main = OpFunction %void None %7
+ %10 = OpLabel
+ %13 = OpAccessChain %_ptr_StorageBuffer_v3uint %out %uint_0
+ %14 = OpAccessChain %_ptr_StorageBuffer_v3uint %in %uint_0
+ %15 = OpLoad %v3uint %14
+ OpStore %13 %15
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/storage/types/vec3.wgsl.expected.wgsl b/test/buffer/storage/types/vec3.wgsl.expected.wgsl
new file mode 100644
index 0000000..71377b0
--- /dev/null
+++ b/test/buffer/storage/types/vec3.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : vec3<u32>;
+
+[[group(0), binding(1)]] var<storage, read_write> out : vec3<u32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ out = in;
+}
diff --git a/test/buffer/storage/types/vec4.wgsl b/test/buffer/storage/types/vec4.wgsl
new file mode 100644
index 0000000..494a89a
--- /dev/null
+++ b/test/buffer/storage/types/vec4.wgsl
@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : vec4<f32>;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : vec4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ out = in;
+}
diff --git a/test/buffer/storage/types/vec4.wgsl.expected.hlsl b/test/buffer/storage/types/vec4.wgsl.expected.hlsl
new file mode 100644
index 0000000..dd068ba
--- /dev/null
+++ b/test/buffer/storage/types/vec4.wgsl.expected.hlsl
@@ -0,0 +1,8 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+ tint_symbol_1.Store4(0u, asuint(asfloat(tint_symbol.Load4(0u))));
+ return;
+}
diff --git a/test/buffer/storage/types/vec4.wgsl.expected.msl b/test/buffer/storage/types/vec4.wgsl.expected.msl
new file mode 100644
index 0000000..4af349b
--- /dev/null
+++ b/test/buffer/storage/types/vec4.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device float4* tint_symbol_1 [[buffer(1)]], const device float4* tint_symbol_2 [[buffer(0)]]) {
+ *(tint_symbol_1) = *(tint_symbol_2);
+ return;
+}
+
diff --git a/test/buffer/storage/types/vec4.wgsl.expected.spvasm b/test/buffer/storage/types/vec4.wgsl.expected.spvasm
new file mode 100644
index 0000000..d74f298
--- /dev/null
+++ b/test/buffer/storage/types/vec4.wgsl.expected.spvasm
@@ -0,0 +1,40 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 17
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %in_block "in_block"
+ OpMemberName %in_block 0 "inner"
+ OpName %in "in"
+ OpName %out "out"
+ OpName %main "main"
+ OpDecorate %in_block Block
+ OpMemberDecorate %in_block 0 Offset 0
+ OpDecorate %in NonWritable
+ OpDecorate %in DescriptorSet 0
+ OpDecorate %in Binding 0
+ OpDecorate %out DescriptorSet 0
+ OpDecorate %out Binding 1
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+ %in_block = OpTypeStruct %v4float
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+ %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+ %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+ %void = OpTypeVoid
+ %7 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v4float = OpTypePointer StorageBuffer %v4float
+ %main = OpFunction %void None %7
+ %10 = OpLabel
+ %14 = OpAccessChain %_ptr_StorageBuffer_v4float %out %uint_0
+ %15 = OpAccessChain %_ptr_StorageBuffer_v4float %in %uint_0
+ %16 = OpLoad %v4float %15
+ OpStore %14 %16
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/storage/types/vec4.wgsl.expected.wgsl b/test/buffer/storage/types/vec4.wgsl.expected.wgsl
new file mode 100644
index 0000000..7b278bd
--- /dev/null
+++ b/test/buffer/storage/types/vec4.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : vec4<f32>;
+
+[[group(0), binding(1)]] var<storage, read_write> out : vec4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ out = in;
+}
diff --git a/test/buffer/uniform/types/array.wgsl b/test/buffer/uniform/types/array.wgsl
new file mode 100644
index 0000000..7151b53
--- /dev/null
+++ b/test/buffer/uniform/types/array.wgsl
@@ -0,0 +1,7 @@
+[[group(0), binding(0)]]
+var<uniform> u : array<vec4<f32>, 4>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let x = u;
+}
diff --git a/test/buffer/uniform/types/array.wgsl.expected.hlsl b/test/buffer/uniform/types/array.wgsl.expected.hlsl
new file mode 100644
index 0000000..34e3ef1
--- /dev/null
+++ b/test/buffer/uniform/types/array.wgsl.expected.hlsl
@@ -0,0 +1,21 @@
+cbuffer cbuffer_u : register(b0, space0) {
+ uint4 u[4];
+};
+
+typedef float4 tint_symbol_ret[4];
+tint_symbol_ret tint_symbol(uint4 buffer[4], uint offset) {
+ float4 arr[4] = (float4[4])0;
+ {
+ [loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+ const uint scalar_offset = ((offset + (i * 16u))) / 4;
+ arr[i] = asfloat(buffer[scalar_offset / 4]);
+ }
+ }
+ return arr;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ const float4 x[4] = tint_symbol(u, 0u);
+ return;
+}
diff --git a/test/buffer/uniform/types/array.wgsl.expected.msl b/test/buffer/uniform/types/array.wgsl.expected.msl
new file mode 100644
index 0000000..71c0b99
--- /dev/null
+++ b/test/buffer/uniform/types/array.wgsl.expected.msl
@@ -0,0 +1,12 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_array_wrapper {
+ /* 0x0000 */ float4 arr[4];
+};
+
+kernel void tint_symbol(const constant tint_array_wrapper* tint_symbol_1 [[buffer(0)]]) {
+ tint_array_wrapper const x = *(tint_symbol_1);
+ return;
+}
+
diff --git a/test/buffer/uniform/types/array.wgsl.expected.spvasm b/test/buffer/uniform/types/array.wgsl.expected.spvasm
new file mode 100644
index 0000000..ef26f1d
--- /dev/null
+++ b/test/buffer/uniform/types/array.wgsl.expected.spvasm
@@ -0,0 +1,37 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 17
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %u_block "u_block"
+ OpMemberName %u_block 0 "inner"
+ OpName %u "u"
+ OpName %main "main"
+ OpDecorate %u_block Block
+ OpMemberDecorate %u_block 0 Offset 0
+ OpDecorate %_arr_v4float_uint_4 ArrayStride 16
+ OpDecorate %u NonWritable
+ OpDecorate %u DescriptorSet 0
+ OpDecorate %u Binding 0
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+ %uint = OpTypeInt 32 0
+ %uint_4 = OpConstant %uint 4
+%_arr_v4float_uint_4 = OpTypeArray %v4float %uint_4
+ %u_block = OpTypeStruct %_arr_v4float_uint_4
+%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
+ %u = OpVariable %_ptr_Uniform_u_block Uniform
+ %void = OpTypeVoid
+ %9 = OpTypeFunction %void
+ %uint_0 = OpConstant %uint 0
+%_ptr_Uniform__arr_v4float_uint_4 = OpTypePointer Uniform %_arr_v4float_uint_4
+ %main = OpFunction %void None %9
+ %12 = OpLabel
+ %15 = OpAccessChain %_ptr_Uniform__arr_v4float_uint_4 %u %uint_0
+ %16 = OpLoad %_arr_v4float_uint_4 %15
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/uniform/types/array.wgsl.expected.wgsl b/test/buffer/uniform/types/array.wgsl.expected.wgsl
new file mode 100644
index 0000000..3cfa21e
--- /dev/null
+++ b/test/buffer/uniform/types/array.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<uniform> u : array<vec4<f32>, 4>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let x = u;
+}
diff --git a/test/buffer/uniform/types/f32.wgsl b/test/buffer/uniform/types/f32.wgsl
new file mode 100644
index 0000000..1f3fbfa
--- /dev/null
+++ b/test/buffer/uniform/types/f32.wgsl
@@ -0,0 +1,7 @@
+[[group(0), binding(0)]]
+var<uniform> u : f32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let x = u;
+}
diff --git a/test/buffer/uniform/types/f32.wgsl.expected.hlsl b/test/buffer/uniform/types/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..496e5c9
--- /dev/null
+++ b/test/buffer/uniform/types/f32.wgsl.expected.hlsl
@@ -0,0 +1,9 @@
+cbuffer cbuffer_u : register(b0, space0) {
+ uint4 u[1];
+};
+
+[numthreads(1, 1, 1)]
+void main() {
+ const float x = asfloat(u[0].x);
+ return;
+}
diff --git a/test/buffer/uniform/types/f32.wgsl.expected.msl b/test/buffer/uniform/types/f32.wgsl.expected.msl
new file mode 100644
index 0000000..81695e1
--- /dev/null
+++ b/test/buffer/uniform/types/f32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(const constant float* tint_symbol_1 [[buffer(0)]]) {
+ float const x = *(tint_symbol_1);
+ return;
+}
+
diff --git a/test/buffer/uniform/types/f32.wgsl.expected.spvasm b/test/buffer/uniform/types/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..bb09d33
--- /dev/null
+++ b/test/buffer/uniform/types/f32.wgsl.expected.spvasm
@@ -0,0 +1,33 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 14
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %u_block "u_block"
+ OpMemberName %u_block 0 "inner"
+ OpName %u "u"
+ OpName %main "main"
+ OpDecorate %u_block Block
+ OpMemberDecorate %u_block 0 Offset 0
+ OpDecorate %u NonWritable
+ OpDecorate %u DescriptorSet 0
+ OpDecorate %u Binding 0
+ %float = OpTypeFloat 32
+ %u_block = OpTypeStruct %float
+%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
+ %u = OpVariable %_ptr_Uniform_u_block Uniform
+ %void = OpTypeVoid
+ %5 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_float = OpTypePointer Uniform %float
+ %main = OpFunction %void None %5
+ %8 = OpLabel
+ %12 = OpAccessChain %_ptr_Uniform_float %u %uint_0
+ %13 = OpLoad %float %12
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/uniform/types/f32.wgsl.expected.wgsl b/test/buffer/uniform/types/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..6bd2751
--- /dev/null
+++ b/test/buffer/uniform/types/f32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<uniform> u : f32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let x = u;
+}
diff --git a/test/buffer/uniform/types/i32.wgsl b/test/buffer/uniform/types/i32.wgsl
new file mode 100644
index 0000000..ca165b4
--- /dev/null
+++ b/test/buffer/uniform/types/i32.wgsl
@@ -0,0 +1,7 @@
+[[group(0), binding(0)]]
+var<uniform> u : i32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let x = u;
+}
diff --git a/test/buffer/uniform/types/i32.wgsl.expected.hlsl b/test/buffer/uniform/types/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..e6ed76e
--- /dev/null
+++ b/test/buffer/uniform/types/i32.wgsl.expected.hlsl
@@ -0,0 +1,9 @@
+cbuffer cbuffer_u : register(b0, space0) {
+ uint4 u[1];
+};
+
+[numthreads(1, 1, 1)]
+void main() {
+ const int x = asint(u[0].x);
+ return;
+}
diff --git a/test/buffer/uniform/types/i32.wgsl.expected.msl b/test/buffer/uniform/types/i32.wgsl.expected.msl
new file mode 100644
index 0000000..3d4cdde
--- /dev/null
+++ b/test/buffer/uniform/types/i32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(const constant int* tint_symbol_1 [[buffer(0)]]) {
+ int const x = *(tint_symbol_1);
+ return;
+}
+
diff --git a/test/buffer/uniform/types/i32.wgsl.expected.spvasm b/test/buffer/uniform/types/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..1f7b1f7
--- /dev/null
+++ b/test/buffer/uniform/types/i32.wgsl.expected.spvasm
@@ -0,0 +1,33 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 14
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %u_block "u_block"
+ OpMemberName %u_block 0 "inner"
+ OpName %u "u"
+ OpName %main "main"
+ OpDecorate %u_block Block
+ OpMemberDecorate %u_block 0 Offset 0
+ OpDecorate %u NonWritable
+ OpDecorate %u DescriptorSet 0
+ OpDecorate %u Binding 0
+ %int = OpTypeInt 32 1
+ %u_block = OpTypeStruct %int
+%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
+ %u = OpVariable %_ptr_Uniform_u_block Uniform
+ %void = OpTypeVoid
+ %5 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_int = OpTypePointer Uniform %int
+ %main = OpFunction %void None %5
+ %8 = OpLabel
+ %12 = OpAccessChain %_ptr_Uniform_int %u %uint_0
+ %13 = OpLoad %int %12
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/uniform/types/i32.wgsl.expected.wgsl b/test/buffer/uniform/types/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..0cf7b04
--- /dev/null
+++ b/test/buffer/uniform/types/i32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<uniform> u : i32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let x = u;
+}
diff --git a/test/buffer/uniform/types/mat2x2.wgsl b/test/buffer/uniform/types/mat2x2.wgsl
new file mode 100644
index 0000000..d1cc3d3
--- /dev/null
+++ b/test/buffer/uniform/types/mat2x2.wgsl
@@ -0,0 +1,7 @@
+[[group(0), binding(0)]]
+var<uniform> u : mat2x2<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let x = u;
+}
diff --git a/test/buffer/uniform/types/mat2x2.wgsl.expected.hlsl b/test/buffer/uniform/types/mat2x2.wgsl.expected.hlsl
new file mode 100644
index 0000000..c34b4c5
--- /dev/null
+++ b/test/buffer/uniform/types/mat2x2.wgsl.expected.hlsl
@@ -0,0 +1,17 @@
+cbuffer cbuffer_u : register(b0, space0) {
+ uint4 u[1];
+};
+
+float2x2 tint_symbol(uint4 buffer[1], uint offset) {
+ const uint scalar_offset = ((offset + 0u)) / 4;
+ uint4 ubo_load = buffer[scalar_offset / 4];
+ const uint scalar_offset_1 = ((offset + 8u)) / 4;
+ uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
+ return float2x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ const float2x2 x = tint_symbol(u, 0u);
+ return;
+}
diff --git a/test/buffer/uniform/types/mat2x2.wgsl.expected.msl b/test/buffer/uniform/types/mat2x2.wgsl.expected.msl
new file mode 100644
index 0000000..226e541
--- /dev/null
+++ b/test/buffer/uniform/types/mat2x2.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(const constant float2x2* tint_symbol_1 [[buffer(0)]]) {
+ float2x2 const x = *(tint_symbol_1);
+ return;
+}
+
diff --git a/test/buffer/uniform/types/mat2x2.wgsl.expected.spvasm b/test/buffer/uniform/types/mat2x2.wgsl.expected.spvasm
new file mode 100644
index 0000000..b5fb317
--- /dev/null
+++ b/test/buffer/uniform/types/mat2x2.wgsl.expected.spvasm
@@ -0,0 +1,37 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %u_block "u_block"
+ OpMemberName %u_block 0 "inner"
+ OpName %u "u"
+ OpName %main "main"
+ OpDecorate %u_block Block
+ OpMemberDecorate %u_block 0 Offset 0
+ OpMemberDecorate %u_block 0 ColMajor
+ OpMemberDecorate %u_block 0 MatrixStride 8
+ OpDecorate %u NonWritable
+ OpDecorate %u DescriptorSet 0
+ OpDecorate %u Binding 0
+ %float = OpTypeFloat 32
+ %v2float = OpTypeVector %float 2
+%mat2v2float = OpTypeMatrix %v2float 2
+ %u_block = OpTypeStruct %mat2v2float
+%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
+ %u = OpVariable %_ptr_Uniform_u_block Uniform
+ %void = OpTypeVoid
+ %7 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_mat2v2float = OpTypePointer Uniform %mat2v2float
+ %main = OpFunction %void None %7
+ %10 = OpLabel
+ %14 = OpAccessChain %_ptr_Uniform_mat2v2float %u %uint_0
+ %15 = OpLoad %mat2v2float %14
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/uniform/types/mat2x2.wgsl.expected.wgsl b/test/buffer/uniform/types/mat2x2.wgsl.expected.wgsl
new file mode 100644
index 0000000..8c42684
--- /dev/null
+++ b/test/buffer/uniform/types/mat2x2.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<uniform> u : mat2x2<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let x = u;
+}
diff --git a/test/buffer/uniform/types/mat2x3.wgsl b/test/buffer/uniform/types/mat2x3.wgsl
new file mode 100644
index 0000000..0fb978d
--- /dev/null
+++ b/test/buffer/uniform/types/mat2x3.wgsl
@@ -0,0 +1,7 @@
+[[group(0), binding(0)]]
+var<uniform> u : mat2x3<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let x = u;
+}
diff --git a/test/buffer/uniform/types/mat2x3.wgsl.expected.hlsl b/test/buffer/uniform/types/mat2x3.wgsl.expected.hlsl
new file mode 100644
index 0000000..e92e916
--- /dev/null
+++ b/test/buffer/uniform/types/mat2x3.wgsl.expected.hlsl
@@ -0,0 +1,15 @@
+cbuffer cbuffer_u : register(b0, space0) {
+ uint4 u[2];
+};
+
+float2x3 tint_symbol(uint4 buffer[2], uint offset) {
+ const uint scalar_offset = ((offset + 0u)) / 4;
+ const uint scalar_offset_1 = ((offset + 16u)) / 4;
+ return float2x3(asfloat(buffer[scalar_offset / 4].xyz), asfloat(buffer[scalar_offset_1 / 4].xyz));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ const float2x3 x = tint_symbol(u, 0u);
+ return;
+}
diff --git a/test/buffer/uniform/types/mat2x3.wgsl.expected.msl b/test/buffer/uniform/types/mat2x3.wgsl.expected.msl
new file mode 100644
index 0000000..1c226a8
--- /dev/null
+++ b/test/buffer/uniform/types/mat2x3.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(const constant float2x3* tint_symbol_1 [[buffer(0)]]) {
+ float2x3 const x = *(tint_symbol_1);
+ return;
+}
+
diff --git a/test/buffer/uniform/types/mat2x3.wgsl.expected.spvasm b/test/buffer/uniform/types/mat2x3.wgsl.expected.spvasm
new file mode 100644
index 0000000..3a86048
--- /dev/null
+++ b/test/buffer/uniform/types/mat2x3.wgsl.expected.spvasm
@@ -0,0 +1,37 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %u_block "u_block"
+ OpMemberName %u_block 0 "inner"
+ OpName %u "u"
+ OpName %main "main"
+ OpDecorate %u_block Block
+ OpMemberDecorate %u_block 0 Offset 0
+ OpMemberDecorate %u_block 0 ColMajor
+ OpMemberDecorate %u_block 0 MatrixStride 16
+ OpDecorate %u NonWritable
+ OpDecorate %u DescriptorSet 0
+ OpDecorate %u Binding 0
+ %float = OpTypeFloat 32
+ %v3float = OpTypeVector %float 3
+%mat2v3float = OpTypeMatrix %v3float 2
+ %u_block = OpTypeStruct %mat2v3float
+%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
+ %u = OpVariable %_ptr_Uniform_u_block Uniform
+ %void = OpTypeVoid
+ %7 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_mat2v3float = OpTypePointer Uniform %mat2v3float
+ %main = OpFunction %void None %7
+ %10 = OpLabel
+ %14 = OpAccessChain %_ptr_Uniform_mat2v3float %u %uint_0
+ %15 = OpLoad %mat2v3float %14
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/uniform/types/mat2x3.wgsl.expected.wgsl b/test/buffer/uniform/types/mat2x3.wgsl.expected.wgsl
new file mode 100644
index 0000000..5a74e94
--- /dev/null
+++ b/test/buffer/uniform/types/mat2x3.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<uniform> u : mat2x3<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let x = u;
+}
diff --git a/test/buffer/uniform/types/mat3x2.wgsl b/test/buffer/uniform/types/mat3x2.wgsl
new file mode 100644
index 0000000..59b1b80
--- /dev/null
+++ b/test/buffer/uniform/types/mat3x2.wgsl
@@ -0,0 +1,7 @@
+[[group(0), binding(0)]]
+var<uniform> u : mat3x2<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let x = u;
+}
diff --git a/test/buffer/uniform/types/mat3x2.wgsl.expected.hlsl b/test/buffer/uniform/types/mat3x2.wgsl.expected.hlsl
new file mode 100644
index 0000000..44f2aa4
--- /dev/null
+++ b/test/buffer/uniform/types/mat3x2.wgsl.expected.hlsl
@@ -0,0 +1,19 @@
+cbuffer cbuffer_u : register(b0, space0) {
+ uint4 u[2];
+};
+
+float3x2 tint_symbol(uint4 buffer[2], uint offset) {
+ const uint scalar_offset = ((offset + 0u)) / 4;
+ uint4 ubo_load = buffer[scalar_offset / 4];
+ const uint scalar_offset_1 = ((offset + 8u)) / 4;
+ uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
+ const uint scalar_offset_2 = ((offset + 16u)) / 4;
+ uint4 ubo_load_2 = buffer[scalar_offset_2 / 4];
+ return float3x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)), asfloat(((scalar_offset_2 & 2) ? ubo_load_2.zw : ubo_load_2.xy)));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ const float3x2 x = tint_symbol(u, 0u);
+ return;
+}
diff --git a/test/buffer/uniform/types/mat3x2.wgsl.expected.msl b/test/buffer/uniform/types/mat3x2.wgsl.expected.msl
new file mode 100644
index 0000000..b39621d
--- /dev/null
+++ b/test/buffer/uniform/types/mat3x2.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(const constant float3x2* tint_symbol_1 [[buffer(0)]]) {
+ float3x2 const x = *(tint_symbol_1);
+ return;
+}
+
diff --git a/test/buffer/uniform/types/mat3x2.wgsl.expected.spvasm b/test/buffer/uniform/types/mat3x2.wgsl.expected.spvasm
new file mode 100644
index 0000000..550ed15
--- /dev/null
+++ b/test/buffer/uniform/types/mat3x2.wgsl.expected.spvasm
@@ -0,0 +1,37 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %u_block "u_block"
+ OpMemberName %u_block 0 "inner"
+ OpName %u "u"
+ OpName %main "main"
+ OpDecorate %u_block Block
+ OpMemberDecorate %u_block 0 Offset 0
+ OpMemberDecorate %u_block 0 ColMajor
+ OpMemberDecorate %u_block 0 MatrixStride 8
+ OpDecorate %u NonWritable
+ OpDecorate %u DescriptorSet 0
+ OpDecorate %u Binding 0
+ %float = OpTypeFloat 32
+ %v2float = OpTypeVector %float 2
+%mat3v2float = OpTypeMatrix %v2float 3
+ %u_block = OpTypeStruct %mat3v2float
+%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
+ %u = OpVariable %_ptr_Uniform_u_block Uniform
+ %void = OpTypeVoid
+ %7 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_mat3v2float = OpTypePointer Uniform %mat3v2float
+ %main = OpFunction %void None %7
+ %10 = OpLabel
+ %14 = OpAccessChain %_ptr_Uniform_mat3v2float %u %uint_0
+ %15 = OpLoad %mat3v2float %14
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/uniform/types/mat3x2.wgsl.expected.wgsl b/test/buffer/uniform/types/mat3x2.wgsl.expected.wgsl
new file mode 100644
index 0000000..5f16f76
--- /dev/null
+++ b/test/buffer/uniform/types/mat3x2.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<uniform> u : mat3x2<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let x = u;
+}
diff --git a/test/buffer/uniform/types/mat4x4.wgsl b/test/buffer/uniform/types/mat4x4.wgsl
new file mode 100644
index 0000000..7d2a57b
--- /dev/null
+++ b/test/buffer/uniform/types/mat4x4.wgsl
@@ -0,0 +1,7 @@
+[[group(0), binding(0)]]
+var<uniform> u : mat4x4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let x = u;
+}
diff --git a/test/buffer/uniform/types/mat4x4.wgsl.expected.hlsl b/test/buffer/uniform/types/mat4x4.wgsl.expected.hlsl
new file mode 100644
index 0000000..0b14de0
--- /dev/null
+++ b/test/buffer/uniform/types/mat4x4.wgsl.expected.hlsl
@@ -0,0 +1,17 @@
+cbuffer cbuffer_u : register(b0, space0) {
+ uint4 u[4];
+};
+
+float4x4 tint_symbol(uint4 buffer[4], uint offset) {
+ const uint scalar_offset = ((offset + 0u)) / 4;
+ const uint scalar_offset_1 = ((offset + 16u)) / 4;
+ const uint scalar_offset_2 = ((offset + 32u)) / 4;
+ const uint scalar_offset_3 = ((offset + 48u)) / 4;
+ return float4x4(asfloat(buffer[scalar_offset / 4]), asfloat(buffer[scalar_offset_1 / 4]), asfloat(buffer[scalar_offset_2 / 4]), asfloat(buffer[scalar_offset_3 / 4]));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ const float4x4 x = tint_symbol(u, 0u);
+ return;
+}
diff --git a/test/buffer/uniform/types/mat4x4.wgsl.expected.msl b/test/buffer/uniform/types/mat4x4.wgsl.expected.msl
new file mode 100644
index 0000000..5a16521
--- /dev/null
+++ b/test/buffer/uniform/types/mat4x4.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(const constant float4x4* tint_symbol_1 [[buffer(0)]]) {
+ float4x4 const x = *(tint_symbol_1);
+ return;
+}
+
diff --git a/test/buffer/uniform/types/mat4x4.wgsl.expected.spvasm b/test/buffer/uniform/types/mat4x4.wgsl.expected.spvasm
new file mode 100644
index 0000000..eb84aa3
--- /dev/null
+++ b/test/buffer/uniform/types/mat4x4.wgsl.expected.spvasm
@@ -0,0 +1,37 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %u_block "u_block"
+ OpMemberName %u_block 0 "inner"
+ OpName %u "u"
+ OpName %main "main"
+ OpDecorate %u_block Block
+ OpMemberDecorate %u_block 0 Offset 0
+ OpMemberDecorate %u_block 0 ColMajor
+ OpMemberDecorate %u_block 0 MatrixStride 16
+ OpDecorate %u NonWritable
+ OpDecorate %u DescriptorSet 0
+ OpDecorate %u Binding 0
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+%mat4v4float = OpTypeMatrix %v4float 4
+ %u_block = OpTypeStruct %mat4v4float
+%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
+ %u = OpVariable %_ptr_Uniform_u_block Uniform
+ %void = OpTypeVoid
+ %7 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_mat4v4float = OpTypePointer Uniform %mat4v4float
+ %main = OpFunction %void None %7
+ %10 = OpLabel
+ %14 = OpAccessChain %_ptr_Uniform_mat4v4float %u %uint_0
+ %15 = OpLoad %mat4v4float %14
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/uniform/types/mat4x4.wgsl.expected.wgsl b/test/buffer/uniform/types/mat4x4.wgsl.expected.wgsl
new file mode 100644
index 0000000..a1188d4
--- /dev/null
+++ b/test/buffer/uniform/types/mat4x4.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<uniform> u : mat4x4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let x = u;
+}
diff --git a/test/buffer/uniform/types/struct.wgsl b/test/buffer/uniform/types/struct.wgsl
new file mode 100644
index 0000000..135bcab
--- /dev/null
+++ b/test/buffer/uniform/types/struct.wgsl
@@ -0,0 +1,14 @@
+struct Inner {
+ f : f32;
+};
+struct S {
+ inner : Inner;
+};
+
+[[group(0), binding(0)]]
+var<uniform> u : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let x = u;
+}
diff --git a/test/buffer/uniform/types/struct.wgsl.expected.hlsl b/test/buffer/uniform/types/struct.wgsl.expected.hlsl
new file mode 100644
index 0000000..d15eb31
--- /dev/null
+++ b/test/buffer/uniform/types/struct.wgsl.expected.hlsl
@@ -0,0 +1,27 @@
+struct Inner {
+ float f;
+};
+struct S {
+ Inner inner;
+};
+
+cbuffer cbuffer_u : register(b0, space0) {
+ uint4 u[1];
+};
+
+Inner tint_symbol_1(uint4 buffer[1], uint offset) {
+ const uint scalar_offset = ((offset + 0u)) / 4;
+ const Inner tint_symbol_3 = {asfloat(buffer[scalar_offset / 4][scalar_offset % 4])};
+ return tint_symbol_3;
+}
+
+S tint_symbol(uint4 buffer[1], uint offset) {
+ const S tint_symbol_4 = {tint_symbol_1(buffer, (offset + 0u))};
+ return tint_symbol_4;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ const S x = tint_symbol(u, 0u);
+ return;
+}
diff --git a/test/buffer/uniform/types/struct.wgsl.expected.msl b/test/buffer/uniform/types/struct.wgsl.expected.msl
new file mode 100644
index 0000000..4e2f5c9
--- /dev/null
+++ b/test/buffer/uniform/types/struct.wgsl.expected.msl
@@ -0,0 +1,15 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Inner {
+ /* 0x0000 */ float f;
+};
+struct S {
+ /* 0x0000 */ Inner inner;
+};
+
+kernel void tint_symbol(const constant S* tint_symbol_1 [[buffer(0)]]) {
+ S const x = *(tint_symbol_1);
+ return;
+}
+
diff --git a/test/buffer/uniform/types/struct.wgsl.expected.spvasm b/test/buffer/uniform/types/struct.wgsl.expected.spvasm
new file mode 100644
index 0000000..b8d2ad8
--- /dev/null
+++ b/test/buffer/uniform/types/struct.wgsl.expected.spvasm
@@ -0,0 +1,33 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 11
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %S "S"
+ OpMemberName %S 0 "inner"
+ OpName %Inner "Inner"
+ OpMemberName %Inner 0 "f"
+ OpName %u "u"
+ OpName %main "main"
+ OpDecorate %S Block
+ OpMemberDecorate %S 0 Offset 0
+ OpMemberDecorate %Inner 0 Offset 0
+ OpDecorate %u NonWritable
+ OpDecorate %u DescriptorSet 0
+ OpDecorate %u Binding 0
+ %float = OpTypeFloat 32
+ %Inner = OpTypeStruct %float
+ %S = OpTypeStruct %Inner
+%_ptr_Uniform_S = OpTypePointer Uniform %S
+ %u = OpVariable %_ptr_Uniform_S Uniform
+ %void = OpTypeVoid
+ %6 = OpTypeFunction %void
+ %main = OpFunction %void None %6
+ %9 = OpLabel
+ %10 = OpLoad %S %u
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/uniform/types/struct.wgsl.expected.wgsl b/test/buffer/uniform/types/struct.wgsl.expected.wgsl
new file mode 100644
index 0000000..fed4de3
--- /dev/null
+++ b/test/buffer/uniform/types/struct.wgsl.expected.wgsl
@@ -0,0 +1,14 @@
+struct Inner {
+ f : f32;
+};
+
+struct S {
+ inner : Inner;
+};
+
+[[group(0), binding(0)]] var<uniform> u : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let x = u;
+}
diff --git a/test/buffer/uniform/types/u32.wgsl b/test/buffer/uniform/types/u32.wgsl
new file mode 100644
index 0000000..a1c80a3
--- /dev/null
+++ b/test/buffer/uniform/types/u32.wgsl
@@ -0,0 +1,7 @@
+[[group(0), binding(0)]]
+var<uniform> u : u32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let x = u;
+}
diff --git a/test/buffer/uniform/types/u32.wgsl.expected.hlsl b/test/buffer/uniform/types/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..f033091
--- /dev/null
+++ b/test/buffer/uniform/types/u32.wgsl.expected.hlsl
@@ -0,0 +1,9 @@
+cbuffer cbuffer_u : register(b0, space0) {
+ uint4 u[1];
+};
+
+[numthreads(1, 1, 1)]
+void main() {
+ const uint x = u[0].x;
+ return;
+}
diff --git a/test/buffer/uniform/types/u32.wgsl.expected.msl b/test/buffer/uniform/types/u32.wgsl.expected.msl
new file mode 100644
index 0000000..4d088dd
--- /dev/null
+++ b/test/buffer/uniform/types/u32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(const constant uint* tint_symbol_1 [[buffer(0)]]) {
+ uint const x = *(tint_symbol_1);
+ return;
+}
+
diff --git a/test/buffer/uniform/types/u32.wgsl.expected.spvasm b/test/buffer/uniform/types/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..b73dd01
--- /dev/null
+++ b/test/buffer/uniform/types/u32.wgsl.expected.spvasm
@@ -0,0 +1,32 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 13
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %u_block "u_block"
+ OpMemberName %u_block 0 "inner"
+ OpName %u "u"
+ OpName %main "main"
+ OpDecorate %u_block Block
+ OpMemberDecorate %u_block 0 Offset 0
+ OpDecorate %u NonWritable
+ OpDecorate %u DescriptorSet 0
+ OpDecorate %u Binding 0
+ %uint = OpTypeInt 32 0
+ %u_block = OpTypeStruct %uint
+%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
+ %u = OpVariable %_ptr_Uniform_u_block Uniform
+ %void = OpTypeVoid
+ %5 = OpTypeFunction %void
+ %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+ %main = OpFunction %void None %5
+ %8 = OpLabel
+ %11 = OpAccessChain %_ptr_Uniform_uint %u %uint_0
+ %12 = OpLoad %uint %11
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/uniform/types/u32.wgsl.expected.wgsl b/test/buffer/uniform/types/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..6751a05
--- /dev/null
+++ b/test/buffer/uniform/types/u32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<uniform> u : u32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let x = u;
+}
diff --git a/test/buffer/uniform/types/vec2.wgsl b/test/buffer/uniform/types/vec2.wgsl
new file mode 100644
index 0000000..8f2debe
--- /dev/null
+++ b/test/buffer/uniform/types/vec2.wgsl
@@ -0,0 +1,7 @@
+[[group(0), binding(0)]]
+var<uniform> u : vec2<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let x = u;
+}
diff --git a/test/buffer/uniform/types/vec2.wgsl.expected.hlsl b/test/buffer/uniform/types/vec2.wgsl.expected.hlsl
new file mode 100644
index 0000000..a36e25c
--- /dev/null
+++ b/test/buffer/uniform/types/vec2.wgsl.expected.hlsl
@@ -0,0 +1,9 @@
+cbuffer cbuffer_u : register(b0, space0) {
+ uint4 u[1];
+};
+
+[numthreads(1, 1, 1)]
+void main() {
+ const int2 x = asint(u[0].xy);
+ return;
+}
diff --git a/test/buffer/uniform/types/vec2.wgsl.expected.msl b/test/buffer/uniform/types/vec2.wgsl.expected.msl
new file mode 100644
index 0000000..3f10795
--- /dev/null
+++ b/test/buffer/uniform/types/vec2.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(const constant int2* tint_symbol_1 [[buffer(0)]]) {
+ int2 const x = *(tint_symbol_1);
+ return;
+}
+
diff --git a/test/buffer/uniform/types/vec2.wgsl.expected.spvasm b/test/buffer/uniform/types/vec2.wgsl.expected.spvasm
new file mode 100644
index 0000000..913911a
--- /dev/null
+++ b/test/buffer/uniform/types/vec2.wgsl.expected.spvasm
@@ -0,0 +1,34 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 15
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %u_block "u_block"
+ OpMemberName %u_block 0 "inner"
+ OpName %u "u"
+ OpName %main "main"
+ OpDecorate %u_block Block
+ OpMemberDecorate %u_block 0 Offset 0
+ OpDecorate %u NonWritable
+ OpDecorate %u DescriptorSet 0
+ OpDecorate %u Binding 0
+ %int = OpTypeInt 32 1
+ %v2int = OpTypeVector %int 2
+ %u_block = OpTypeStruct %v2int
+%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
+ %u = OpVariable %_ptr_Uniform_u_block Uniform
+ %void = OpTypeVoid
+ %6 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_v2int = OpTypePointer Uniform %v2int
+ %main = OpFunction %void None %6
+ %9 = OpLabel
+ %13 = OpAccessChain %_ptr_Uniform_v2int %u %uint_0
+ %14 = OpLoad %v2int %13
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/uniform/types/vec2.wgsl.expected.wgsl b/test/buffer/uniform/types/vec2.wgsl.expected.wgsl
new file mode 100644
index 0000000..61695c1
--- /dev/null
+++ b/test/buffer/uniform/types/vec2.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<uniform> u : vec2<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let x = u;
+}
diff --git a/test/buffer/uniform/types/vec3.wgsl b/test/buffer/uniform/types/vec3.wgsl
new file mode 100644
index 0000000..012d269
--- /dev/null
+++ b/test/buffer/uniform/types/vec3.wgsl
@@ -0,0 +1,7 @@
+[[group(0), binding(0)]]
+var<uniform> u : vec3<u32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let x = u;
+}
diff --git a/test/buffer/uniform/types/vec3.wgsl.expected.hlsl b/test/buffer/uniform/types/vec3.wgsl.expected.hlsl
new file mode 100644
index 0000000..8e081c8
--- /dev/null
+++ b/test/buffer/uniform/types/vec3.wgsl.expected.hlsl
@@ -0,0 +1,9 @@
+cbuffer cbuffer_u : register(b0, space0) {
+ uint4 u[1];
+};
+
+[numthreads(1, 1, 1)]
+void main() {
+ const uint3 x = u[0].xyz;
+ return;
+}
diff --git a/test/buffer/uniform/types/vec3.wgsl.expected.msl b/test/buffer/uniform/types/vec3.wgsl.expected.msl
new file mode 100644
index 0000000..407c0e7
--- /dev/null
+++ b/test/buffer/uniform/types/vec3.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(const constant uint3* tint_symbol_1 [[buffer(0)]]) {
+ uint3 const x = *(tint_symbol_1);
+ return;
+}
+
diff --git a/test/buffer/uniform/types/vec3.wgsl.expected.spvasm b/test/buffer/uniform/types/vec3.wgsl.expected.spvasm
new file mode 100644
index 0000000..52fd93a
--- /dev/null
+++ b/test/buffer/uniform/types/vec3.wgsl.expected.spvasm
@@ -0,0 +1,33 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 14
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %u_block "u_block"
+ OpMemberName %u_block 0 "inner"
+ OpName %u "u"
+ OpName %main "main"
+ OpDecorate %u_block Block
+ OpMemberDecorate %u_block 0 Offset 0
+ OpDecorate %u NonWritable
+ OpDecorate %u DescriptorSet 0
+ OpDecorate %u Binding 0
+ %uint = OpTypeInt 32 0
+ %v3uint = OpTypeVector %uint 3
+ %u_block = OpTypeStruct %v3uint
+%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
+ %u = OpVariable %_ptr_Uniform_u_block Uniform
+ %void = OpTypeVoid
+ %6 = OpTypeFunction %void
+ %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_v3uint = OpTypePointer Uniform %v3uint
+ %main = OpFunction %void None %6
+ %9 = OpLabel
+ %12 = OpAccessChain %_ptr_Uniform_v3uint %u %uint_0
+ %13 = OpLoad %v3uint %12
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/uniform/types/vec3.wgsl.expected.wgsl b/test/buffer/uniform/types/vec3.wgsl.expected.wgsl
new file mode 100644
index 0000000..d8d99e8
--- /dev/null
+++ b/test/buffer/uniform/types/vec3.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<uniform> u : vec3<u32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let x = u;
+}
diff --git a/test/buffer/uniform/types/vec4.wgsl b/test/buffer/uniform/types/vec4.wgsl
new file mode 100644
index 0000000..24d63f0
--- /dev/null
+++ b/test/buffer/uniform/types/vec4.wgsl
@@ -0,0 +1,7 @@
+[[group(0), binding(0)]]
+var<uniform> u : vec4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let x = u;
+}
diff --git a/test/buffer/uniform/types/vec4.wgsl.expected.hlsl b/test/buffer/uniform/types/vec4.wgsl.expected.hlsl
new file mode 100644
index 0000000..2ae2ec1
--- /dev/null
+++ b/test/buffer/uniform/types/vec4.wgsl.expected.hlsl
@@ -0,0 +1,9 @@
+cbuffer cbuffer_u : register(b0, space0) {
+ uint4 u[1];
+};
+
+[numthreads(1, 1, 1)]
+void main() {
+ const float4 x = asfloat(u[0]);
+ return;
+}
diff --git a/test/buffer/uniform/types/vec4.wgsl.expected.msl b/test/buffer/uniform/types/vec4.wgsl.expected.msl
new file mode 100644
index 0000000..28d31ec
--- /dev/null
+++ b/test/buffer/uniform/types/vec4.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(const constant float4* tint_symbol_1 [[buffer(0)]]) {
+ float4 const x = *(tint_symbol_1);
+ return;
+}
+
diff --git a/test/buffer/uniform/types/vec4.wgsl.expected.spvasm b/test/buffer/uniform/types/vec4.wgsl.expected.spvasm
new file mode 100644
index 0000000..f95df57
--- /dev/null
+++ b/test/buffer/uniform/types/vec4.wgsl.expected.spvasm
@@ -0,0 +1,34 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 15
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %u_block "u_block"
+ OpMemberName %u_block 0 "inner"
+ OpName %u "u"
+ OpName %main "main"
+ OpDecorate %u_block Block
+ OpMemberDecorate %u_block 0 Offset 0
+ OpDecorate %u NonWritable
+ OpDecorate %u DescriptorSet 0
+ OpDecorate %u Binding 0
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+ %u_block = OpTypeStruct %v4float
+%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
+ %u = OpVariable %_ptr_Uniform_u_block Uniform
+ %void = OpTypeVoid
+ %6 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_v4float = OpTypePointer Uniform %v4float
+ %main = OpFunction %void None %6
+ %9 = OpLabel
+ %13 = OpAccessChain %_ptr_Uniform_v4float %u %uint_0
+ %14 = OpLoad %v4float %13
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/uniform/types/vec4.wgsl.expected.wgsl b/test/buffer/uniform/types/vec4.wgsl.expected.wgsl
new file mode 100644
index 0000000..fe83de0
--- /dev/null
+++ b/test/buffer/uniform/types/vec4.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<uniform> u : vec4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let x = u;
+}
