From 7395e29e709148bfe175bd19d921bac4397b22f4 Mon Sep 17 00:00:00 2001 From: James Price Date: Wed, 19 Jan 2022 15:55:56 +0000 Subject: [PATCH] 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 Kokoro: Kokoro --- src/resolver/resolver_validation.cc | 44 ----- src/resolver/storage_class_validation_test.cc | 91 ++++++++--- src/transform/add_spirv_block_decoration.cc | 41 ++--- .../add_spirv_block_decoration_test.cc | 153 ++++++++++++++---- .../module_scope_var_to_entry_point_param.cc | 32 +++- ...ule_scope_var_to_entry_point_param_test.cc | 93 +++++++++++ src/writer/hlsl/generator_impl.cc | 16 +- test/buffer/storage/types/array.wgsl | 10 ++ .../storage/types/array.wgsl.expected.hlsl | 28 ++++ .../storage/types/array.wgsl.expected.msl | 12 ++ .../storage/types/array.wgsl.expected.spvasm | 42 +++++ .../storage/types/array.wgsl.expected.wgsl | 8 + test/buffer/storage/types/f32.wgsl | 10 ++ .../storage/types/f32.wgsl.expected.hlsl | 8 + .../storage/types/f32.wgsl.expected.msl | 8 + .../storage/types/f32.wgsl.expected.spvasm | 39 +++++ .../storage/types/f32.wgsl.expected.wgsl | 8 + test/buffer/storage/types/i32.wgsl | 10 ++ .../storage/types/i32.wgsl.expected.hlsl | 8 + .../storage/types/i32.wgsl.expected.msl | 8 + .../storage/types/i32.wgsl.expected.spvasm | 39 +++++ .../storage/types/i32.wgsl.expected.wgsl | 8 + test/buffer/storage/types/mat2x2.wgsl | 10 ++ .../storage/types/mat2x2.wgsl.expected.hlsl | 17 ++ .../storage/types/mat2x2.wgsl.expected.msl | 8 + .../storage/types/mat2x2.wgsl.expected.spvasm | 43 +++++ .../storage/types/mat2x2.wgsl.expected.wgsl | 8 + test/buffer/storage/types/mat2x3.wgsl | 10 ++ .../storage/types/mat2x3.wgsl.expected.hlsl | 17 ++ .../storage/types/mat2x3.wgsl.expected.msl | 8 + .../storage/types/mat2x3.wgsl.expected.spvasm | 43 +++++ .../storage/types/mat2x3.wgsl.expected.wgsl | 8 + test/buffer/storage/types/mat3x2.wgsl | 10 ++ .../storage/types/mat3x2.wgsl.expected.hlsl | 18 +++ .../storage/types/mat3x2.wgsl.expected.msl | 8 + .../storage/types/mat3x2.wgsl.expected.spvasm | 43 +++++ .../storage/types/mat3x2.wgsl.expected.wgsl | 8 + test/buffer/storage/types/mat4x4.wgsl | 10 ++ .../storage/types/mat4x4.wgsl.expected.hlsl | 19 +++ .../storage/types/mat4x4.wgsl.expected.msl | 8 + .../storage/types/mat4x4.wgsl.expected.spvasm | 43 +++++ .../storage/types/mat4x4.wgsl.expected.wgsl | 8 + test/buffer/storage/types/runtime_array.wgsl | 14 ++ .../types/runtime_array.wgsl.expected.hlsl | 21 +++ .../types/runtime_array.wgsl.expected.msl | 18 +++ .../types/runtime_array.wgsl.expected.spvasm | 47 ++++++ .../types/runtime_array.wgsl.expected.wgsl | 12 ++ test/buffer/storage/types/struct.wgsl | 17 ++ .../storage/types/struct.wgsl.expected.hlsl | 33 ++++ .../storage/types/struct.wgsl.expected.msl | 15 ++ .../storage/types/struct.wgsl.expected.spvasm | 38 +++++ .../storage/types/struct.wgsl.expected.wgsl | 16 ++ test/buffer/storage/types/u32.wgsl | 10 ++ .../storage/types/u32.wgsl.expected.hlsl | 8 + .../storage/types/u32.wgsl.expected.msl | 8 + .../storage/types/u32.wgsl.expected.spvasm | 38 +++++ .../storage/types/u32.wgsl.expected.wgsl | 8 + test/buffer/storage/types/vec2.wgsl | 10 ++ .../storage/types/vec2.wgsl.expected.hlsl | 8 + .../storage/types/vec2.wgsl.expected.msl | 8 + .../storage/types/vec2.wgsl.expected.spvasm | 40 +++++ .../storage/types/vec2.wgsl.expected.wgsl | 8 + test/buffer/storage/types/vec3.wgsl | 10 ++ .../storage/types/vec3.wgsl.expected.hlsl | 8 + .../storage/types/vec3.wgsl.expected.msl | 8 + .../storage/types/vec3.wgsl.expected.spvasm | 39 +++++ .../storage/types/vec3.wgsl.expected.wgsl | 8 + test/buffer/storage/types/vec4.wgsl | 10 ++ .../storage/types/vec4.wgsl.expected.hlsl | 8 + .../storage/types/vec4.wgsl.expected.msl | 8 + .../storage/types/vec4.wgsl.expected.spvasm | 40 +++++ .../storage/types/vec4.wgsl.expected.wgsl | 8 + test/buffer/uniform/types/array.wgsl | 7 + .../uniform/types/array.wgsl.expected.hlsl | 21 +++ .../uniform/types/array.wgsl.expected.msl | 12 ++ .../uniform/types/array.wgsl.expected.spvasm | 37 +++++ .../uniform/types/array.wgsl.expected.wgsl | 6 + test/buffer/uniform/types/f32.wgsl | 7 + .../uniform/types/f32.wgsl.expected.hlsl | 9 ++ .../uniform/types/f32.wgsl.expected.msl | 8 + .../uniform/types/f32.wgsl.expected.spvasm | 33 ++++ .../uniform/types/f32.wgsl.expected.wgsl | 6 + test/buffer/uniform/types/i32.wgsl | 7 + .../uniform/types/i32.wgsl.expected.hlsl | 9 ++ .../uniform/types/i32.wgsl.expected.msl | 8 + .../uniform/types/i32.wgsl.expected.spvasm | 33 ++++ .../uniform/types/i32.wgsl.expected.wgsl | 6 + test/buffer/uniform/types/mat2x2.wgsl | 7 + .../uniform/types/mat2x2.wgsl.expected.hlsl | 17 ++ .../uniform/types/mat2x2.wgsl.expected.msl | 8 + .../uniform/types/mat2x2.wgsl.expected.spvasm | 37 +++++ .../uniform/types/mat2x2.wgsl.expected.wgsl | 6 + test/buffer/uniform/types/mat2x3.wgsl | 7 + .../uniform/types/mat2x3.wgsl.expected.hlsl | 15 ++ .../uniform/types/mat2x3.wgsl.expected.msl | 8 + .../uniform/types/mat2x3.wgsl.expected.spvasm | 37 +++++ .../uniform/types/mat2x3.wgsl.expected.wgsl | 6 + test/buffer/uniform/types/mat3x2.wgsl | 7 + .../uniform/types/mat3x2.wgsl.expected.hlsl | 19 +++ .../uniform/types/mat3x2.wgsl.expected.msl | 8 + .../uniform/types/mat3x2.wgsl.expected.spvasm | 37 +++++ .../uniform/types/mat3x2.wgsl.expected.wgsl | 6 + test/buffer/uniform/types/mat4x4.wgsl | 7 + .../uniform/types/mat4x4.wgsl.expected.hlsl | 17 ++ .../uniform/types/mat4x4.wgsl.expected.msl | 8 + .../uniform/types/mat4x4.wgsl.expected.spvasm | 37 +++++ .../uniform/types/mat4x4.wgsl.expected.wgsl | 6 + test/buffer/uniform/types/struct.wgsl | 14 ++ .../uniform/types/struct.wgsl.expected.hlsl | 27 ++++ .../uniform/types/struct.wgsl.expected.msl | 15 ++ .../uniform/types/struct.wgsl.expected.spvasm | 33 ++++ .../uniform/types/struct.wgsl.expected.wgsl | 14 ++ test/buffer/uniform/types/u32.wgsl | 7 + .../uniform/types/u32.wgsl.expected.hlsl | 9 ++ .../uniform/types/u32.wgsl.expected.msl | 8 + .../uniform/types/u32.wgsl.expected.spvasm | 32 ++++ .../uniform/types/u32.wgsl.expected.wgsl | 6 + test/buffer/uniform/types/vec2.wgsl | 7 + .../uniform/types/vec2.wgsl.expected.hlsl | 9 ++ .../uniform/types/vec2.wgsl.expected.msl | 8 + .../uniform/types/vec2.wgsl.expected.spvasm | 34 ++++ .../uniform/types/vec2.wgsl.expected.wgsl | 6 + test/buffer/uniform/types/vec3.wgsl | 7 + .../uniform/types/vec3.wgsl.expected.hlsl | 9 ++ .../uniform/types/vec3.wgsl.expected.msl | 8 + .../uniform/types/vec3.wgsl.expected.spvasm | 33 ++++ .../uniform/types/vec3.wgsl.expected.wgsl | 6 + test/buffer/uniform/types/vec4.wgsl | 7 + .../uniform/types/vec4.wgsl.expected.hlsl | 9 ++ .../uniform/types/vec4.wgsl.expected.msl | 8 + .../uniform/types/vec4.wgsl.expected.spvasm | 34 ++++ .../uniform/types/vec4.wgsl.expected.wgsl | 6 + 132 files changed, 2309 insertions(+), 143 deletions(-) create mode 100644 test/buffer/storage/types/array.wgsl create mode 100644 test/buffer/storage/types/array.wgsl.expected.hlsl create mode 100644 test/buffer/storage/types/array.wgsl.expected.msl create mode 100644 test/buffer/storage/types/array.wgsl.expected.spvasm create mode 100644 test/buffer/storage/types/array.wgsl.expected.wgsl create mode 100644 test/buffer/storage/types/f32.wgsl create mode 100644 test/buffer/storage/types/f32.wgsl.expected.hlsl create mode 100644 test/buffer/storage/types/f32.wgsl.expected.msl create mode 100644 test/buffer/storage/types/f32.wgsl.expected.spvasm create mode 100644 test/buffer/storage/types/f32.wgsl.expected.wgsl create mode 100644 test/buffer/storage/types/i32.wgsl create mode 100644 test/buffer/storage/types/i32.wgsl.expected.hlsl create mode 100644 test/buffer/storage/types/i32.wgsl.expected.msl create mode 100644 test/buffer/storage/types/i32.wgsl.expected.spvasm create mode 100644 test/buffer/storage/types/i32.wgsl.expected.wgsl create mode 100644 test/buffer/storage/types/mat2x2.wgsl create mode 100644 test/buffer/storage/types/mat2x2.wgsl.expected.hlsl create mode 100644 test/buffer/storage/types/mat2x2.wgsl.expected.msl create mode 100644 test/buffer/storage/types/mat2x2.wgsl.expected.spvasm create mode 100644 test/buffer/storage/types/mat2x2.wgsl.expected.wgsl create mode 100644 test/buffer/storage/types/mat2x3.wgsl create mode 100644 test/buffer/storage/types/mat2x3.wgsl.expected.hlsl create mode 100644 test/buffer/storage/types/mat2x3.wgsl.expected.msl create mode 100644 test/buffer/storage/types/mat2x3.wgsl.expected.spvasm create mode 100644 test/buffer/storage/types/mat2x3.wgsl.expected.wgsl create mode 100644 test/buffer/storage/types/mat3x2.wgsl create mode 100644 test/buffer/storage/types/mat3x2.wgsl.expected.hlsl create mode 100644 test/buffer/storage/types/mat3x2.wgsl.expected.msl create mode 100644 test/buffer/storage/types/mat3x2.wgsl.expected.spvasm create mode 100644 test/buffer/storage/types/mat3x2.wgsl.expected.wgsl create mode 100644 test/buffer/storage/types/mat4x4.wgsl create mode 100644 test/buffer/storage/types/mat4x4.wgsl.expected.hlsl create mode 100644 test/buffer/storage/types/mat4x4.wgsl.expected.msl create mode 100644 test/buffer/storage/types/mat4x4.wgsl.expected.spvasm create mode 100644 test/buffer/storage/types/mat4x4.wgsl.expected.wgsl create mode 100644 test/buffer/storage/types/runtime_array.wgsl create mode 100644 test/buffer/storage/types/runtime_array.wgsl.expected.hlsl create mode 100644 test/buffer/storage/types/runtime_array.wgsl.expected.msl create mode 100644 test/buffer/storage/types/runtime_array.wgsl.expected.spvasm create mode 100644 test/buffer/storage/types/runtime_array.wgsl.expected.wgsl create mode 100644 test/buffer/storage/types/struct.wgsl create mode 100644 test/buffer/storage/types/struct.wgsl.expected.hlsl create mode 100644 test/buffer/storage/types/struct.wgsl.expected.msl create mode 100644 test/buffer/storage/types/struct.wgsl.expected.spvasm create mode 100644 test/buffer/storage/types/struct.wgsl.expected.wgsl create mode 100644 test/buffer/storage/types/u32.wgsl create mode 100644 test/buffer/storage/types/u32.wgsl.expected.hlsl create mode 100644 test/buffer/storage/types/u32.wgsl.expected.msl create mode 100644 test/buffer/storage/types/u32.wgsl.expected.spvasm create mode 100644 test/buffer/storage/types/u32.wgsl.expected.wgsl create mode 100644 test/buffer/storage/types/vec2.wgsl create mode 100644 test/buffer/storage/types/vec2.wgsl.expected.hlsl create mode 100644 test/buffer/storage/types/vec2.wgsl.expected.msl create mode 100644 test/buffer/storage/types/vec2.wgsl.expected.spvasm create mode 100644 test/buffer/storage/types/vec2.wgsl.expected.wgsl create mode 100644 test/buffer/storage/types/vec3.wgsl create mode 100644 test/buffer/storage/types/vec3.wgsl.expected.hlsl create mode 100644 test/buffer/storage/types/vec3.wgsl.expected.msl create mode 100644 test/buffer/storage/types/vec3.wgsl.expected.spvasm create mode 100644 test/buffer/storage/types/vec3.wgsl.expected.wgsl create mode 100644 test/buffer/storage/types/vec4.wgsl create mode 100644 test/buffer/storage/types/vec4.wgsl.expected.hlsl create mode 100644 test/buffer/storage/types/vec4.wgsl.expected.msl create mode 100644 test/buffer/storage/types/vec4.wgsl.expected.spvasm create mode 100644 test/buffer/storage/types/vec4.wgsl.expected.wgsl create mode 100644 test/buffer/uniform/types/array.wgsl create mode 100644 test/buffer/uniform/types/array.wgsl.expected.hlsl create mode 100644 test/buffer/uniform/types/array.wgsl.expected.msl create mode 100644 test/buffer/uniform/types/array.wgsl.expected.spvasm create mode 100644 test/buffer/uniform/types/array.wgsl.expected.wgsl create mode 100644 test/buffer/uniform/types/f32.wgsl create mode 100644 test/buffer/uniform/types/f32.wgsl.expected.hlsl create mode 100644 test/buffer/uniform/types/f32.wgsl.expected.msl create mode 100644 test/buffer/uniform/types/f32.wgsl.expected.spvasm create mode 100644 test/buffer/uniform/types/f32.wgsl.expected.wgsl create mode 100644 test/buffer/uniform/types/i32.wgsl create mode 100644 test/buffer/uniform/types/i32.wgsl.expected.hlsl create mode 100644 test/buffer/uniform/types/i32.wgsl.expected.msl create mode 100644 test/buffer/uniform/types/i32.wgsl.expected.spvasm create mode 100644 test/buffer/uniform/types/i32.wgsl.expected.wgsl create mode 100644 test/buffer/uniform/types/mat2x2.wgsl create mode 100644 test/buffer/uniform/types/mat2x2.wgsl.expected.hlsl create mode 100644 test/buffer/uniform/types/mat2x2.wgsl.expected.msl create mode 100644 test/buffer/uniform/types/mat2x2.wgsl.expected.spvasm create mode 100644 test/buffer/uniform/types/mat2x2.wgsl.expected.wgsl create mode 100644 test/buffer/uniform/types/mat2x3.wgsl create mode 100644 test/buffer/uniform/types/mat2x3.wgsl.expected.hlsl create mode 100644 test/buffer/uniform/types/mat2x3.wgsl.expected.msl create mode 100644 test/buffer/uniform/types/mat2x3.wgsl.expected.spvasm create mode 100644 test/buffer/uniform/types/mat2x3.wgsl.expected.wgsl create mode 100644 test/buffer/uniform/types/mat3x2.wgsl create mode 100644 test/buffer/uniform/types/mat3x2.wgsl.expected.hlsl create mode 100644 test/buffer/uniform/types/mat3x2.wgsl.expected.msl create mode 100644 test/buffer/uniform/types/mat3x2.wgsl.expected.spvasm create mode 100644 test/buffer/uniform/types/mat3x2.wgsl.expected.wgsl create mode 100644 test/buffer/uniform/types/mat4x4.wgsl create mode 100644 test/buffer/uniform/types/mat4x4.wgsl.expected.hlsl create mode 100644 test/buffer/uniform/types/mat4x4.wgsl.expected.msl create mode 100644 test/buffer/uniform/types/mat4x4.wgsl.expected.spvasm create mode 100644 test/buffer/uniform/types/mat4x4.wgsl.expected.wgsl create mode 100644 test/buffer/uniform/types/struct.wgsl create mode 100644 test/buffer/uniform/types/struct.wgsl.expected.hlsl create mode 100644 test/buffer/uniform/types/struct.wgsl.expected.msl create mode 100644 test/buffer/uniform/types/struct.wgsl.expected.spvasm create mode 100644 test/buffer/uniform/types/struct.wgsl.expected.wgsl create mode 100644 test/buffer/uniform/types/u32.wgsl create mode 100644 test/buffer/uniform/types/u32.wgsl.expected.hlsl create mode 100644 test/buffer/uniform/types/u32.wgsl.expected.msl create mode 100644 test/buffer/uniform/types/u32.wgsl.expected.spvasm create mode 100644 test/buffer/uniform/types/u32.wgsl.expected.wgsl create mode 100644 test/buffer/uniform/types/vec2.wgsl create mode 100644 test/buffer/uniform/types/vec2.wgsl.expected.hlsl create mode 100644 test/buffer/uniform/types/vec2.wgsl.expected.msl create mode 100644 test/buffer/uniform/types/vec2.wgsl.expected.spvasm create mode 100644 test/buffer/uniform/types/vec2.wgsl.expected.wgsl create mode 100644 test/buffer/uniform/types/vec3.wgsl create mode 100644 test/buffer/uniform/types/vec3.wgsl.expected.hlsl create mode 100644 test/buffer/uniform/types/vec3.wgsl.expected.msl create mode 100644 test/buffer/uniform/types/vec3.wgsl.expected.spvasm create mode 100644 test/buffer/uniform/types/vec3.wgsl.expected.wgsl create mode 100644 test/buffer/uniform/types/vec4.wgsl create mode 100644 test/buffer/uniform/types/vec4.wgsl.expected.hlsl create mode 100644 test/buffer/uniform/types/vec4.wgsl.expected.msl create mode 100644 test/buffer/uniform/types/vec4.wgsl.expected.spvasm create mode 100644 test/buffer/uniform/types/vec4.wgsl.expected.wgsl diff --git a/src/resolver/resolver_validation.cc b/src/resolver/resolver_validation.cc index 1df63ba371..3218e4af49 100644 --- a/src/resolver/resolver_validation.cc +++ b/src/resolver/resolver_validation.cc @@ -459,42 +459,6 @@ bool Resolver::ValidateGlobalVariable(const sem::Variable* var) { 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(); - if (!str) { - AddError( - "variables declared in the 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(); - if (!str) { - AddError( - "variables declared in the 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 @@ bool Resolver::ValidateVariable(const sem::Variable* var) { return false; } - if (auto* r = storage_ty->As()) { - 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()) { 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 f8360882aa..5d922c2f50 100644 --- a/src/resolver/storage_class_validation_test.cc +++ b/src/resolver/storage_class_validation_test.cc @@ -92,6 +92,40 @@ note: while analysing structure member S.m } TEST_F(ResolverStorageClassValidationTest, StorageBufferBool) { + // var g : bool; + Global(Source{{56, 78}}, "g", ty.bool_(), ast::StorageClass::kStorage, + ast::DecorationList{ + create(0), + create(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 g : ptr; + Global(Source{{56, 78}}, "g", + ty.pointer(ty.f32(), ast::StorageClass::kPrivate), + ast::StorageClass::kStorage, + ast::DecorationList{ + create(0), + create(0), + }); + + ASSERT_FALSE(r()->Resolve()); + + EXPECT_EQ( + r()->error(), + R"(56:78 error: Type 'ptr' 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 g : i32; Global(Source{{56, 78}}, "g", ty.i32(), ast::StorageClass::kStorage, ast::DecorationList{ @@ -99,14 +133,10 @@ TEST_F(ResolverStorageClassValidationTest, StorageBufferBool) { create(0), }); - ASSERT_FALSE(r()->Resolve()); - - EXPECT_EQ( - r()->error(), - R"(56:78 error: variables declared in the storage class must be of a structure type)"); + ASSERT_TRUE(r()->Resolve()) << r()->error(); } -TEST_F(ResolverStorageClassValidationTest, StorageBufferPointer) { +TEST_F(ResolverStorageClassValidationTest, StorageBufferVector) { // var g : vec4; Global(Source{{56, 78}}, "g", ty.vec4(), ast::StorageClass::kStorage, ast::DecorationList{ @@ -114,11 +144,7 @@ TEST_F(ResolverStorageClassValidationTest, StorageBufferPointer) { create(0), }); - ASSERT_FALSE(r()->Resolve()); - - EXPECT_EQ( - r()->error(), - R"(56:78 error: variables declared in the storage class must be of a structure type)"); + ASSERT_TRUE(r()->Resolve()) << r()->error(); } TEST_F(ResolverStorageClassValidationTest, StorageBufferArray) { @@ -132,11 +158,7 @@ TEST_F(ResolverStorageClassValidationTest, StorageBufferArray) { create(0), }); - ASSERT_FALSE(r()->Resolve()); - - EXPECT_EQ( - r()->error(), - R"(56:78 error: variables declared in the 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, UniformBufferBool) { } TEST_F(ResolverStorageClassValidationTest, UniformBufferPointer) { - // var g : vec4; - Global(Source{{56, 78}}, "g", ty.vec4(), ast::StorageClass::kUniform, + // var g : ptr; + Global(Source{{56, 78}}, "g", + ty.pointer(ty.f32(), ast::StorageClass::kPrivate), + ast::StorageClass::kUniform, ast::DecorationList{ create(0), create(0), @@ -251,7 +275,30 @@ TEST_F(ResolverStorageClassValidationTest, UniformBufferPointer) { EXPECT_EQ( r()->error(), - R"(56:78 error: variables declared in the storage class must be of a structure type)"); + R"(56:78 error: Type 'ptr' 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 g : i32; + Global(Source{{56, 78}}, "g", ty.i32(), ast::StorageClass::kUniform, + ast::DecorationList{ + create(0), + create(0), + }); + + ASSERT_TRUE(r()->Resolve()) << r()->error(); +} + +TEST_F(ResolverStorageClassValidationTest, UniformBufferVector) { + // var g : vec4; + Global(Source{{56, 78}}, "g", ty.vec4(), ast::StorageClass::kUniform, + ast::DecorationList{ + create(0), + create(0), + }); + + ASSERT_TRUE(r()->Resolve()) << r()->error(); } TEST_F(ResolverStorageClassValidationTest, UniformBufferArray) { @@ -264,11 +311,7 @@ TEST_F(ResolverStorageClassValidationTest, UniformBufferArray) { create(0), }); - ASSERT_FALSE(r()->Resolve()); - - EXPECT_EQ( - r()->error(), - R"(56:78 error: variables declared in the 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 e4829c924b..3995c839e1 100644 --- a/src/transform/add_spirv_block_decoration.cc +++ b/src/transform/add_spirv_block_decoration.cc @@ -52,9 +52,9 @@ void AddSpirvBlockDecoration::Run(CloneContext& ctx, const DataMap&, DataMap&) { } } - // 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 wrapper_structs; + std::unordered_map wrapper_structs; // Process global variables that are buffers. for (auto* var : ctx.src->AST().GlobalVariables()) { @@ -64,40 +64,33 @@ void AddSpirvBlockDecoration::Run(CloneContext& ctx, const DataMap&, DataMap&) { continue; } - auto* str = sem.Get(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(); + 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(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( 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 3780c209b4..74e4b33386 100644 --- a/src/transform/add_spirv_block_decoration_test.cc +++ b/src/transform/add_spirv_block_decoration_test.cc @@ -73,7 +73,98 @@ fn main() -> S { EXPECT_EQ(expect, str(got)); } -TEST_F(AddSpirvBlockDecorationTest, Basic) { +TEST_F(AddSpirvBlockDecorationTest, BasicScalar) { + auto* src = R"( +[[group(0), binding(0)]] +var 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 u : u_block; + +[[stage(fragment)]] +fn main() { + let f = u.inner; +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(AddSpirvBlockDecorationTest, BasicArray) { + auto* src = R"( +[[group(0), binding(0)]] +var u : array, 4u>; + +[[stage(fragment)]] +fn main() { + let a = u; +} +)"; + auto* expect = R"( +[[internal(spirv_block)]] +struct u_block { + inner : array, 4u>; +}; + +[[group(0), binding(0)]] var u : u_block; + +[[stage(fragment)]] +fn main() { + let a = u.inner; +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(AddSpirvBlockDecorationTest, BasicArray_Alias) { + auto* src = R"( +type Numbers = array, 4u>; + +[[group(0), binding(0)]] +var u : Numbers; + +[[stage(fragment)]] +fn main() { + let a = u; +} +)"; + auto* expect = R"( +type Numbers = array, 4u>; + +[[internal(spirv_block)]] +struct u_block { + inner : array, 4u>; +}; + +[[group(0), binding(0)]] var u : u_block; + +[[stage(fragment)]] +fn main() { + let a = u.inner; +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(AddSpirvBlockDecorationTest, BasicStruct) { auto* src = R"( struct S { f : f32; @@ -174,11 +265,6 @@ struct Inner { f : f32; }; -[[internal(spirv_block)]] -struct Inner_block { - inner : Inner; -}; - [[internal(spirv_block)]] struct Outer { i : Inner; @@ -186,7 +272,12 @@ struct Outer { [[group(0), binding(0)]] var u0 : Outer; -[[group(0), binding(1)]] var u1 : Inner_block; +[[internal(spirv_block)]] +struct u1_block { + inner : Inner; +}; + +[[group(0), binding(1)]] var u1 : u1_block; [[stage(fragment)]] fn main() { @@ -226,18 +317,18 @@ struct Inner { f : f32; }; -[[internal(spirv_block)]] -struct Inner_block { - inner : Inner; -}; - struct Outer { i : Inner; }; var p : Outer; -[[group(0), binding(1)]] var u : Inner_block; +[[internal(spirv_block)]] +struct u_block { + inner : Inner; +}; + +[[group(0), binding(1)]] var u : u_block; [[stage(fragment)]] fn main() { @@ -282,11 +373,6 @@ struct Inner { f : f32; }; -[[internal(spirv_block)]] -struct Inner_block { - inner : Inner; -}; - [[internal(spirv_block)]] struct S { i : Inner; @@ -294,9 +380,14 @@ struct S { [[group(0), binding(0)]] var u0 : S; -[[group(0), binding(1)]] var u1 : Inner_block; +[[internal(spirv_block)]] +struct u1_block { + inner : Inner; +}; -[[group(0), binding(2)]] var u2 : Inner_block; +[[group(0), binding(1)]] var u1 : u1_block; + +[[group(0), binding(2)]] var u2 : u1_block; [[stage(fragment)]] fn main() { @@ -332,11 +423,11 @@ struct S { }; [[internal(spirv_block)]] -struct S_block { +struct u_block { inner : S; }; -[[group(0), binding(0)]] var u : S_block; +[[group(0), binding(0)]] var u : u_block; [[stage(fragment)]] fn main() { @@ -375,13 +466,13 @@ struct S { }; [[internal(spirv_block)]] -struct S_block { +struct u0_block { inner : S; }; -[[group(0), binding(0)]] var u0 : S_block; +[[group(0), binding(0)]] var u0 : u0_block; -[[group(0), binding(1)]] var u1 : S_block; +[[group(0), binding(1)]] var u1 : u0_block; [[stage(fragment)]] fn main() { @@ -427,11 +518,6 @@ struct Inner { f : f32; }; -[[internal(spirv_block)]] -struct Inner_block { - inner : Inner; -}; - type MyInner = Inner; [[internal(spirv_block)]] @@ -443,7 +529,12 @@ type MyOuter = Outer; [[group(0), binding(0)]] var u0 : MyOuter; -[[group(0), binding(1)]] var u1 : Inner_block; +[[internal(spirv_block)]] +struct u1_block { + inner : Inner; +}; + +[[group(0), binding(1)]] var 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 ba61160940..960fc84601 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 @@ struct ModuleScopeVarToEntryPointParam::State { 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 @@ struct ModuleScopeVarToEntryPointParam::State { 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 @@ struct ModuleScopeVarToEntryPointParam::State { 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(); + 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 @@ struct ModuleScopeVarToEntryPointParam::State { 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 a1620e787b..0f7d518a37 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 @@ fn main([[group(0), binding(0), internal(disable_validation__entry_point_paramet EXPECT_EQ(expect, str(got)); } +TEST_F(ModuleScopeVarToEntryPointParamTest, Buffer_RuntimeArray) { + auto* src = R"( +[[group(0), binding(0)]] +var buffer : array; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = buffer[0]; +} +)"; + + auto* expect = R"( +struct tint_symbol_1 { + arr : array; +}; + +[[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) { + _ = (*(tint_symbol)).arr[0]; +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(ModuleScopeVarToEntryPointParamTest, Buffer_RuntimeArray_Alias) { + auto* src = R"( +type myarray = array; + +[[group(0), binding(0)]] +var buffer : myarray; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = buffer[0]; +} +)"; + + auto* expect = R"( +struct tint_symbol_1 { + arr : array; +}; + +type myarray = array; + +[[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) { + _ = (*(tint_symbol)).arr[0]; +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(ModuleScopeVarToEntryPointParamTest, Buffer_ArrayOfStruct) { + auto* src = R"( +struct S { + f : f32; +}; + +[[group(0), binding(0)]] +var buffer : array; + +[[stage(compute), workgroup_size(1)]] +fn main() { + _ = buffer[0]; +} +)"; + + auto* expect = R"( +struct S { + f : f32; +}; + +struct tint_symbol_1 { + arr : array; +}; + +[[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) { + _ = (*(tint_symbol)).arr[0]; +} +)"; + + auto got = Run(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 76dc529d61..bae3044f39 100644 --- a/src/writer/hlsl/generator_impl.cc +++ b/src/writer/hlsl/generator_impl.cc @@ -2869,14 +2869,6 @@ bool GeneratorImpl::EmitUniformVariable(const sem::Variable* var) { auto* decl = var->Declaration(); auto binding_point = decl->BindingPoint(); auto* type = var->Type()->UnwrapRef(); - - auto* str = type->As(); - 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 @@ bool GeneratorImpl::EmitType(std::ostream& out, out << "ByteAddressBuffer"; return true; case ast::StorageClass::kUniform: { - auto* str = type->As(); - 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 0000000000..6105a26727 --- /dev/null +++ b/test/buffer/storage/types/array.wgsl @@ -0,0 +1,10 @@ +[[group(0), binding(0)]] +var in : array; + +[[group(0), binding(1)]] +var out : array; + +[[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 0000000000..4b02ccdd25 --- /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 0000000000..a95963fecf --- /dev/null +++ b/test/buffer/storage/types/array.wgsl.expected.msl @@ -0,0 +1,12 @@ +#include + +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 0000000000..65f933d554 --- /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 0000000000..0aaacaedd2 --- /dev/null +++ b/test/buffer/storage/types/array.wgsl.expected.wgsl @@ -0,0 +1,8 @@ +[[group(0), binding(0)]] var in : array; + +[[group(0), binding(1)]] var out : array; + +[[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 0000000000..d66cdb0cea --- /dev/null +++ b/test/buffer/storage/types/f32.wgsl @@ -0,0 +1,10 @@ +[[group(0), binding(0)]] +var in : f32; + +[[group(0), binding(1)]] +var 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 0000000000..e6dca9ed16 --- /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 0000000000..033b5e2feb --- /dev/null +++ b/test/buffer/storage/types/f32.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +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 0000000000..dea11b4b0e --- /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 0000000000..a385a75d5e --- /dev/null +++ b/test/buffer/storage/types/f32.wgsl.expected.wgsl @@ -0,0 +1,8 @@ +[[group(0), binding(0)]] var in : f32; + +[[group(0), binding(1)]] var 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 0000000000..2ffd35fe0f --- /dev/null +++ b/test/buffer/storage/types/i32.wgsl @@ -0,0 +1,10 @@ +[[group(0), binding(0)]] +var in : i32; + +[[group(0), binding(1)]] +var 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 0000000000..3c53558620 --- /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 0000000000..1de9e037ca --- /dev/null +++ b/test/buffer/storage/types/i32.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +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 0000000000..4abab36eac --- /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 0000000000..3424f2e6f1 --- /dev/null +++ b/test/buffer/storage/types/i32.wgsl.expected.wgsl @@ -0,0 +1,8 @@ +[[group(0), binding(0)]] var in : i32; + +[[group(0), binding(1)]] var 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 0000000000..60ff548adb --- /dev/null +++ b/test/buffer/storage/types/mat2x2.wgsl @@ -0,0 +1,10 @@ +[[group(0), binding(0)]] +var in : mat2x2; + +[[group(0), binding(1)]] +var out : mat2x2; + +[[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 0000000000..a1fc1b643f --- /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 0000000000..c9ffff792a --- /dev/null +++ b/test/buffer/storage/types/mat2x2.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +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 0000000000..a372c43cfa --- /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 0000000000..8d2cbae323 --- /dev/null +++ b/test/buffer/storage/types/mat2x2.wgsl.expected.wgsl @@ -0,0 +1,8 @@ +[[group(0), binding(0)]] var in : mat2x2; + +[[group(0), binding(1)]] var out : mat2x2; + +[[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 0000000000..151db83808 --- /dev/null +++ b/test/buffer/storage/types/mat2x3.wgsl @@ -0,0 +1,10 @@ +[[group(0), binding(0)]] +var in : mat2x3; + +[[group(0), binding(1)]] +var out : mat2x3; + +[[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 0000000000..baff4c5904 --- /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 0000000000..1b704d4a15 --- /dev/null +++ b/test/buffer/storage/types/mat2x3.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +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 0000000000..31ddb997df --- /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 0000000000..ab3a5fd966 --- /dev/null +++ b/test/buffer/storage/types/mat2x3.wgsl.expected.wgsl @@ -0,0 +1,8 @@ +[[group(0), binding(0)]] var in : mat2x3; + +[[group(0), binding(1)]] var out : mat2x3; + +[[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 0000000000..d77ac81d22 --- /dev/null +++ b/test/buffer/storage/types/mat3x2.wgsl @@ -0,0 +1,10 @@ +[[group(0), binding(0)]] +var in : mat3x2; + +[[group(0), binding(1)]] +var out : mat3x2; + +[[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 0000000000..31d606bf03 --- /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 0000000000..b8765f08ab --- /dev/null +++ b/test/buffer/storage/types/mat3x2.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +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 0000000000..c352e184c9 --- /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 0000000000..4c87f4ac1e --- /dev/null +++ b/test/buffer/storage/types/mat3x2.wgsl.expected.wgsl @@ -0,0 +1,8 @@ +[[group(0), binding(0)]] var in : mat3x2; + +[[group(0), binding(1)]] var out : mat3x2; + +[[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 0000000000..bf18a256f9 --- /dev/null +++ b/test/buffer/storage/types/mat4x4.wgsl @@ -0,0 +1,10 @@ +[[group(0), binding(0)]] +var in : mat4x4; + +[[group(0), binding(1)]] +var out : mat4x4; + +[[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 0000000000..e818c0922d --- /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 0000000000..6b33874fa0 --- /dev/null +++ b/test/buffer/storage/types/mat4x4.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +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 0000000000..512341416c --- /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 0000000000..7c1ebe9e7a --- /dev/null +++ b/test/buffer/storage/types/mat4x4.wgsl.expected.wgsl @@ -0,0 +1,8 @@ +[[group(0), binding(0)]] var in : mat4x4; + +[[group(0), binding(1)]] var out : mat4x4; + +[[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 0000000000..f403fd377d --- /dev/null +++ b/test/buffer/storage/types/runtime_array.wgsl @@ -0,0 +1,14 @@ +struct S { + f : f32; +}; + +[[group(0), binding(0)]] +var in : array; + +[[group(0), binding(1)]] +var out : array; + +[[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 0000000000..035878ad15 --- /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 0000000000..ff6a283788 --- /dev/null +++ b/test/buffer/storage/types/runtime_array.wgsl.expected.msl @@ -0,0 +1,18 @@ +#include + +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 0000000000..5a63e67be4 --- /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 0000000000..418c67d64f --- /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 in : array; + +[[group(0), binding(1)]] var out : array; + +[[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 0000000000..42acfbce5e --- /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 in : S; + +[[group(0), binding(1)]] +var 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 0000000000..05be6006f9 --- /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 0000000000..d5c7427870 --- /dev/null +++ b/test/buffer/storage/types/struct.wgsl.expected.msl @@ -0,0 +1,15 @@ +#include + +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 0000000000..058a3e5427 --- /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 0000000000..f86c26b429 --- /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 in : S; + +[[group(0), binding(1)]] var 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 0000000000..4664d29d6c --- /dev/null +++ b/test/buffer/storage/types/u32.wgsl @@ -0,0 +1,10 @@ +[[group(0), binding(0)]] +var in : u32; + +[[group(0), binding(1)]] +var 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 0000000000..8574dafdc8 --- /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 0000000000..bb65310eec --- /dev/null +++ b/test/buffer/storage/types/u32.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +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 0000000000..201cc294b3 --- /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 0000000000..46dc9d383e --- /dev/null +++ b/test/buffer/storage/types/u32.wgsl.expected.wgsl @@ -0,0 +1,8 @@ +[[group(0), binding(0)]] var in : u32; + +[[group(0), binding(1)]] var 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 0000000000..9841fd4181 --- /dev/null +++ b/test/buffer/storage/types/vec2.wgsl @@ -0,0 +1,10 @@ +[[group(0), binding(0)]] +var in : vec2; + +[[group(0), binding(1)]] +var out : vec2; + +[[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 0000000000..3c5c7d7ebe --- /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 0000000000..c730e05c48 --- /dev/null +++ b/test/buffer/storage/types/vec2.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +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 0000000000..85b79d3146 --- /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 0000000000..a945b3228b --- /dev/null +++ b/test/buffer/storage/types/vec2.wgsl.expected.wgsl @@ -0,0 +1,8 @@ +[[group(0), binding(0)]] var in : vec2; + +[[group(0), binding(1)]] var out : vec2; + +[[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 0000000000..5335059754 --- /dev/null +++ b/test/buffer/storage/types/vec3.wgsl @@ -0,0 +1,10 @@ +[[group(0), binding(0)]] +var in : vec3; + +[[group(0), binding(1)]] +var out : vec3; + +[[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 0000000000..9af25f4d63 --- /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 0000000000..6d55796494 --- /dev/null +++ b/test/buffer/storage/types/vec3.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +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 0000000000..70a2dafdb0 --- /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 0000000000..71377b0d4d --- /dev/null +++ b/test/buffer/storage/types/vec3.wgsl.expected.wgsl @@ -0,0 +1,8 @@ +[[group(0), binding(0)]] var in : vec3; + +[[group(0), binding(1)]] var out : vec3; + +[[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 0000000000..494a89abdc --- /dev/null +++ b/test/buffer/storage/types/vec4.wgsl @@ -0,0 +1,10 @@ +[[group(0), binding(0)]] +var in : vec4; + +[[group(0), binding(1)]] +var out : vec4; + +[[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 0000000000..dd068ba25a --- /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 0000000000..4af349b7cf --- /dev/null +++ b/test/buffer/storage/types/vec4.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +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 0000000000..d74f2981da --- /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 0000000000..7b278bd0fc --- /dev/null +++ b/test/buffer/storage/types/vec4.wgsl.expected.wgsl @@ -0,0 +1,8 @@ +[[group(0), binding(0)]] var in : vec4; + +[[group(0), binding(1)]] var out : vec4; + +[[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 0000000000..7151b53956 --- /dev/null +++ b/test/buffer/uniform/types/array.wgsl @@ -0,0 +1,7 @@ +[[group(0), binding(0)]] +var u : array, 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 0000000000..34e3ef1c6a --- /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 0000000000..71c0b99b0f --- /dev/null +++ b/test/buffer/uniform/types/array.wgsl.expected.msl @@ -0,0 +1,12 @@ +#include + +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 0000000000..ef26f1dffb --- /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 0000000000..3cfa21e32d --- /dev/null +++ b/test/buffer/uniform/types/array.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[group(0), binding(0)]] var u : array, 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 0000000000..1f3fbfa74c --- /dev/null +++ b/test/buffer/uniform/types/f32.wgsl @@ -0,0 +1,7 @@ +[[group(0), binding(0)]] +var 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 0000000000..496e5c9fd4 --- /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 0000000000..81695e10d4 --- /dev/null +++ b/test/buffer/uniform/types/f32.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +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 0000000000..bb09d33562 --- /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 0000000000..6bd2751aaf --- /dev/null +++ b/test/buffer/uniform/types/f32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[group(0), binding(0)]] var 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 0000000000..ca165b4601 --- /dev/null +++ b/test/buffer/uniform/types/i32.wgsl @@ -0,0 +1,7 @@ +[[group(0), binding(0)]] +var 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 0000000000..e6ed76ea43 --- /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 0000000000..3d4cdde979 --- /dev/null +++ b/test/buffer/uniform/types/i32.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +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 0000000000..1f7b1f7701 --- /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 0000000000..0cf7b04a61 --- /dev/null +++ b/test/buffer/uniform/types/i32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[group(0), binding(0)]] var 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 0000000000..d1cc3d3d93 --- /dev/null +++ b/test/buffer/uniform/types/mat2x2.wgsl @@ -0,0 +1,7 @@ +[[group(0), binding(0)]] +var u : mat2x2; + +[[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 0000000000..c34b4c53df --- /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 0000000000..226e54180d --- /dev/null +++ b/test/buffer/uniform/types/mat2x2.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +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 0000000000..b5fb3179b8 --- /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 0000000000..8c42684b4f --- /dev/null +++ b/test/buffer/uniform/types/mat2x2.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[group(0), binding(0)]] var u : mat2x2; + +[[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 0000000000..0fb978d61f --- /dev/null +++ b/test/buffer/uniform/types/mat2x3.wgsl @@ -0,0 +1,7 @@ +[[group(0), binding(0)]] +var u : mat2x3; + +[[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 0000000000..e92e916135 --- /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 0000000000..1c226a8283 --- /dev/null +++ b/test/buffer/uniform/types/mat2x3.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +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 0000000000..3a86048f06 --- /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 0000000000..5a74e9478a --- /dev/null +++ b/test/buffer/uniform/types/mat2x3.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[group(0), binding(0)]] var u : mat2x3; + +[[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 0000000000..59b1b8059c --- /dev/null +++ b/test/buffer/uniform/types/mat3x2.wgsl @@ -0,0 +1,7 @@ +[[group(0), binding(0)]] +var u : mat3x2; + +[[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 0000000000..44f2aa4431 --- /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 0000000000..b39621d687 --- /dev/null +++ b/test/buffer/uniform/types/mat3x2.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +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 0000000000..550ed15496 --- /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 0000000000..5f16f764b7 --- /dev/null +++ b/test/buffer/uniform/types/mat3x2.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[group(0), binding(0)]] var u : mat3x2; + +[[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 0000000000..7d2a57b0b1 --- /dev/null +++ b/test/buffer/uniform/types/mat4x4.wgsl @@ -0,0 +1,7 @@ +[[group(0), binding(0)]] +var u : mat4x4; + +[[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 0000000000..0b14de0372 --- /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 0000000000..5a16521c42 --- /dev/null +++ b/test/buffer/uniform/types/mat4x4.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +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 0000000000..eb84aa3485 --- /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 0000000000..a1188d409b --- /dev/null +++ b/test/buffer/uniform/types/mat4x4.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[group(0), binding(0)]] var u : mat4x4; + +[[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 0000000000..135bcab6fd --- /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 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 0000000000..d15eb31450 --- /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 0000000000..4e2f5c9797 --- /dev/null +++ b/test/buffer/uniform/types/struct.wgsl.expected.msl @@ -0,0 +1,15 @@ +#include + +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 0000000000..b8d2ad8fce --- /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 0000000000..fed4de3d78 --- /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 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 0000000000..a1c80a397b --- /dev/null +++ b/test/buffer/uniform/types/u32.wgsl @@ -0,0 +1,7 @@ +[[group(0), binding(0)]] +var 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 0000000000..f033091a5a --- /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 0000000000..4d088dd484 --- /dev/null +++ b/test/buffer/uniform/types/u32.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +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 0000000000..b73dd01a1f --- /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 0000000000..6751a05e52 --- /dev/null +++ b/test/buffer/uniform/types/u32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[group(0), binding(0)]] var 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 0000000000..8f2debec27 --- /dev/null +++ b/test/buffer/uniform/types/vec2.wgsl @@ -0,0 +1,7 @@ +[[group(0), binding(0)]] +var u : vec2; + +[[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 0000000000..a36e25c3a9 --- /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 0000000000..3f1079516e --- /dev/null +++ b/test/buffer/uniform/types/vec2.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +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 0000000000..913911a522 --- /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 0000000000..61695c1daf --- /dev/null +++ b/test/buffer/uniform/types/vec2.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[group(0), binding(0)]] var u : vec2; + +[[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 0000000000..012d269fa4 --- /dev/null +++ b/test/buffer/uniform/types/vec3.wgsl @@ -0,0 +1,7 @@ +[[group(0), binding(0)]] +var u : vec3; + +[[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 0000000000..8e081c8a73 --- /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 0000000000..407c0e7545 --- /dev/null +++ b/test/buffer/uniform/types/vec3.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +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 0000000000..52fd93a432 --- /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 0000000000..d8d99e8ea4 --- /dev/null +++ b/test/buffer/uniform/types/vec3.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[group(0), binding(0)]] var u : vec3; + +[[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 0000000000..24d63f09fc --- /dev/null +++ b/test/buffer/uniform/types/vec4.wgsl @@ -0,0 +1,7 @@ +[[group(0), binding(0)]] +var u : vec4; + +[[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 0000000000..2ae2ec1970 --- /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 0000000000..28d31ecb39 --- /dev/null +++ b/test/buffer/uniform/types/vec4.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +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 0000000000..f95df57838 --- /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 0000000000..fe83de0469 --- /dev/null +++ b/test/buffer/uniform/types/vec4.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[group(0), binding(0)]] var u : vec4; + +[[stage(compute), workgroup_size(1)]] +fn main() { + let x = u; +}