From cbe816f93dac1d0bd23e82e924fc90f05899688c Mon Sep 17 00:00:00 2001 From: James Price Date: Wed, 19 May 2021 10:38:18 +0000 Subject: [PATCH] writer/msl: Generate address spaces for pointers Add more E2E tests to cover pointers with different storage classes. Fixed: tint:815 Change-Id: I224a794cdf60648ce71dc9a0922d489542995be1 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/51404 Auto-Submit: James Price Commit-Queue: Ben Clayton Kokoro: Kokoro Reviewed-by: Ben Clayton --- src/writer/msl/generator_impl.cc | 18 +++++++++- test/ptr_ref/access/matrix.wgsl.expected.msl | 2 +- test/ptr_ref/access/vector.wgsl.expected.msl | 2 +- .../ptr_ref/copy/ptr_copy.spvasm.expected.msl | 4 +-- test/ptr_ref/load/local/ptr_function.wgsl | 6 ++++ .../local/ptr_function.wgsl.expected.hlsl | 1 + .../load/local/ptr_function.wgsl.expected.msl | 10 ++++++ .../local/ptr_function.wgsl.expected.spvasm | 26 ++++++++++++++ .../local/ptr_function.wgsl.expected.wgsl | 6 ++++ test/ptr_ref/load/local/ptr_private.wgsl | 7 ++++ .../load/local/ptr_private.wgsl.expected.hlsl | 1 + .../load/local/ptr_private.wgsl.expected.msl | 1 + .../local/ptr_private.wgsl.expected.spvasm | 24 +++++++++++++ .../load/local/ptr_private.wgsl.expected.wgsl | 7 ++++ test/ptr_ref/load/local/ptr_storage.wgsl | 13 +++++++ .../load/local/ptr_storage.wgsl.expected.hlsl | 1 + .../load/local/ptr_storage.wgsl.expected.msl | 13 +++++++ .../local/ptr_storage.wgsl.expected.spvasm | 34 +++++++++++++++++++ .../load/local/ptr_storage.wgsl.expected.wgsl | 12 +++++++ test/ptr_ref/load/local/ptr_uniform.wgsl | 13 +++++++ .../load/local/ptr_uniform.wgsl.expected.hlsl | 1 + .../load/local/ptr_uniform.wgsl.expected.msl | 13 +++++++ .../local/ptr_uniform.wgsl.expected.spvasm | 34 +++++++++++++++++++ .../load/local/ptr_uniform.wgsl.expected.wgsl | 12 +++++++ test/ptr_ref/load/local/ptr_workgroup.wgsl | 8 +++++ .../local/ptr_workgroup.wgsl.expected.hlsl | 1 + .../local/ptr_workgroup.wgsl.expected.msl | 1 + .../local/ptr_workgroup.wgsl.expected.spvasm | 25 ++++++++++++++ .../local/ptr_workgroup.wgsl.expected.wgsl | 8 +++++ .../load/param/ptr.spvasm.expected.msl | 2 +- test/ptr_ref/load/param/ptr.wgsl.expected.msl | 2 +- .../ptr_ref/store/local/i32.wgsl.expected.msl | 2 +- .../store/param/ptr.spvasm.expected.msl | 2 +- .../ptr_ref/store/param/ptr.wgsl.expected.msl | 2 +- 34 files changed, 304 insertions(+), 10 deletions(-) create mode 100644 test/ptr_ref/load/local/ptr_function.wgsl create mode 100644 test/ptr_ref/load/local/ptr_function.wgsl.expected.hlsl create mode 100644 test/ptr_ref/load/local/ptr_function.wgsl.expected.msl create mode 100644 test/ptr_ref/load/local/ptr_function.wgsl.expected.spvasm create mode 100644 test/ptr_ref/load/local/ptr_function.wgsl.expected.wgsl create mode 100644 test/ptr_ref/load/local/ptr_private.wgsl create mode 100644 test/ptr_ref/load/local/ptr_private.wgsl.expected.hlsl create mode 100644 test/ptr_ref/load/local/ptr_private.wgsl.expected.msl create mode 100644 test/ptr_ref/load/local/ptr_private.wgsl.expected.spvasm create mode 100644 test/ptr_ref/load/local/ptr_private.wgsl.expected.wgsl create mode 100644 test/ptr_ref/load/local/ptr_storage.wgsl create mode 100644 test/ptr_ref/load/local/ptr_storage.wgsl.expected.hlsl create mode 100644 test/ptr_ref/load/local/ptr_storage.wgsl.expected.msl create mode 100644 test/ptr_ref/load/local/ptr_storage.wgsl.expected.spvasm create mode 100644 test/ptr_ref/load/local/ptr_storage.wgsl.expected.wgsl create mode 100644 test/ptr_ref/load/local/ptr_uniform.wgsl create mode 100644 test/ptr_ref/load/local/ptr_uniform.wgsl.expected.hlsl create mode 100644 test/ptr_ref/load/local/ptr_uniform.wgsl.expected.msl create mode 100644 test/ptr_ref/load/local/ptr_uniform.wgsl.expected.spvasm create mode 100644 test/ptr_ref/load/local/ptr_uniform.wgsl.expected.wgsl create mode 100644 test/ptr_ref/load/local/ptr_workgroup.wgsl create mode 100644 test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.hlsl create mode 100644 test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.msl create mode 100644 test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.spvasm create mode 100644 test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.wgsl diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc index c4d25cf43e..ee976d0cae 100644 --- a/src/writer/msl/generator_impl.cc +++ b/src/writer/msl/generator_impl.cc @@ -1913,7 +1913,23 @@ bool GeneratorImpl::EmitType(const sem::Type* type, const std::string& name) { } out_ << mat->columns() << "x" << mat->rows(); } else if (auto* ptr = type->As()) { - // TODO(dsinclair): Storage class? + switch (ptr->StorageClass()) { + case ast::StorageClass::kFunction: + case ast::StorageClass::kPrivate: + out_ << "thread "; + break; + case ast::StorageClass::kWorkgroup: + out_ << "threadgroup "; + break; + case ast::StorageClass::kStorage: + out_ << "device "; + break; + case ast::StorageClass::kUniform: + out_ << "constant "; + break; + default: + TINT_ICE(diagnostics_) << "unhandled storage class for pointer"; + } if (!EmitType(ptr->StoreType(), "")) { return false; } diff --git a/test/ptr_ref/access/matrix.wgsl.expected.msl b/test/ptr_ref/access/matrix.wgsl.expected.msl index 5bf9bb9a15..1a95911ba5 100644 --- a/test/ptr_ref/access/matrix.wgsl.expected.msl +++ b/test/ptr_ref/access/matrix.wgsl.expected.msl @@ -3,7 +3,7 @@ using namespace metal; kernel void tint_symbol() { float3x3 m = float3x3(float3(1.0f, 2.0f, 3.0f), float3(4.0f, 5.0f, 6.0f), float3(7.0f, 8.0f, 9.0f)); - float3* const v = &(m[1]); + thread float3* const v = &(m[1]); *(v) = float3(5.0f, 5.0f, 5.0f); return; } diff --git a/test/ptr_ref/access/vector.wgsl.expected.msl b/test/ptr_ref/access/vector.wgsl.expected.msl index 88a1b99c98..fc618513db 100644 --- a/test/ptr_ref/access/vector.wgsl.expected.msl +++ b/test/ptr_ref/access/vector.wgsl.expected.msl @@ -3,7 +3,7 @@ using namespace metal; kernel void tint_symbol() { float3 v = float3(1.0f, 2.0f, 3.0f); - float* const f = &(v.y); + thread float* const f = &(v.y); *(f) = 5.0f; return; } diff --git a/test/ptr_ref/copy/ptr_copy.spvasm.expected.msl b/test/ptr_ref/copy/ptr_copy.spvasm.expected.msl index ec58031418..6577255310 100644 --- a/test/ptr_ref/copy/ptr_copy.spvasm.expected.msl +++ b/test/ptr_ref/copy/ptr_copy.spvasm.expected.msl @@ -3,8 +3,8 @@ using namespace metal; kernel void tint_symbol() { uint x_10 = 0u; - uint* const x_1 = &(x_10); - uint* const x_2 = x_1; + thread uint* const x_1 = &(x_10); + thread uint* const x_2 = x_1; return; } diff --git a/test/ptr_ref/load/local/ptr_function.wgsl b/test/ptr_ref/load/local/ptr_function.wgsl new file mode 100644 index 0000000000..4c7f82c8e3 --- /dev/null +++ b/test/ptr_ref/load/local/ptr_function.wgsl @@ -0,0 +1,6 @@ +[[stage(compute)]] +fn main() { + var i : i32 = 123; + let p : ptr = &i; + let use : i32 = *p + 1; +} diff --git a/test/ptr_ref/load/local/ptr_function.wgsl.expected.hlsl b/test/ptr_ref/load/local/ptr_function.wgsl.expected.hlsl new file mode 100644 index 0000000000..b3db42f63d --- /dev/null +++ b/test/ptr_ref/load/local/ptr_function.wgsl.expected.hlsl @@ -0,0 +1 @@ +SKIP: Failed to generate: error: pointers not supported in HLSL diff --git a/test/ptr_ref/load/local/ptr_function.wgsl.expected.msl b/test/ptr_ref/load/local/ptr_function.wgsl.expected.msl new file mode 100644 index 0000000000..ce60e0c8c5 --- /dev/null +++ b/test/ptr_ref/load/local/ptr_function.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void tint_symbol() { + int i = 123; + thread int* const p = &(i); + int const use = (*(p) + 1); + return; +} + diff --git a/test/ptr_ref/load/local/ptr_function.wgsl.expected.spvasm b/test/ptr_ref/load/local/ptr_function.wgsl.expected.spvasm new file mode 100644 index 0000000000..3cb667bcab --- /dev/null +++ b/test/ptr_ref/load/local/ptr_function.wgsl.expected.spvasm @@ -0,0 +1,26 @@ +; 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 %main "main" + OpName %i "i" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %int_123 = OpConstant %int 123 +%_ptr_Function_int = OpTypePointer Function %int + %9 = OpConstantNull %int + %int_1 = OpConstant %int 1 + %main = OpFunction %void None %1 + %4 = OpLabel + %i = OpVariable %_ptr_Function_int Function %9 + OpStore %i %int_123 + %12 = OpLoad %int %i + %14 = OpIAdd %int %12 %int_1 + OpReturn + OpFunctionEnd diff --git a/test/ptr_ref/load/local/ptr_function.wgsl.expected.wgsl b/test/ptr_ref/load/local/ptr_function.wgsl.expected.wgsl new file mode 100644 index 0000000000..be1837a0bf --- /dev/null +++ b/test/ptr_ref/load/local/ptr_function.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute)]] +fn main() { + var i : i32 = 123; + let p : ptr = &(i); + let use : i32 = (*(p) + 1); +} diff --git a/test/ptr_ref/load/local/ptr_private.wgsl b/test/ptr_ref/load/local/ptr_private.wgsl new file mode 100644 index 0000000000..bd93db3680 --- /dev/null +++ b/test/ptr_ref/load/local/ptr_private.wgsl @@ -0,0 +1,7 @@ +var i : i32 = 123; + +[[stage(compute)]] +fn main() { + let p : ptr = &i; + let use : i32 = *p + 1; +} diff --git a/test/ptr_ref/load/local/ptr_private.wgsl.expected.hlsl b/test/ptr_ref/load/local/ptr_private.wgsl.expected.hlsl new file mode 100644 index 0000000000..b3db42f63d --- /dev/null +++ b/test/ptr_ref/load/local/ptr_private.wgsl.expected.hlsl @@ -0,0 +1 @@ +SKIP: Failed to generate: error: pointers not supported in HLSL diff --git a/test/ptr_ref/load/local/ptr_private.wgsl.expected.msl b/test/ptr_ref/load/local/ptr_private.wgsl.expected.msl new file mode 100644 index 0000000000..1b2d3db8ae --- /dev/null +++ b/test/ptr_ref/load/local/ptr_private.wgsl.expected.msl @@ -0,0 +1 @@ +SKIP: TINT_UNIMPLEMENTED crbug.com/tint/726: module-scope private and workgroup variables not yet implemented diff --git a/test/ptr_ref/load/local/ptr_private.wgsl.expected.spvasm b/test/ptr_ref/load/local/ptr_private.wgsl.expected.spvasm new file mode 100644 index 0000000000..7bf6bf91d5 --- /dev/null +++ b/test/ptr_ref/load/local/ptr_private.wgsl.expected.spvasm @@ -0,0 +1,24 @@ +; 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 %i "i" + OpName %main "main" + %int = OpTypeInt 32 1 + %int_123 = OpConstant %int 123 +%_ptr_Private_int = OpTypePointer Private %int + %i = OpVariable %_ptr_Private_int Private %int_123 + %void = OpTypeVoid + %5 = OpTypeFunction %void + %int_1 = OpConstant %int 1 + %main = OpFunction %void None %5 + %8 = OpLabel + %11 = OpLoad %int %i + %13 = OpIAdd %int %11 %int_1 + OpReturn + OpFunctionEnd diff --git a/test/ptr_ref/load/local/ptr_private.wgsl.expected.wgsl b/test/ptr_ref/load/local/ptr_private.wgsl.expected.wgsl new file mode 100644 index 0000000000..a865378ae9 --- /dev/null +++ b/test/ptr_ref/load/local/ptr_private.wgsl.expected.wgsl @@ -0,0 +1,7 @@ +var i : i32 = 123; + +[[stage(compute)]] +fn main() { + let p : ptr = &(i); + let use : i32 = (*(p) + 1); +} diff --git a/test/ptr_ref/load/local/ptr_storage.wgsl b/test/ptr_ref/load/local/ptr_storage.wgsl new file mode 100644 index 0000000000..d334f3420c --- /dev/null +++ b/test/ptr_ref/load/local/ptr_storage.wgsl @@ -0,0 +1,13 @@ +[[block]] +struct S { + a : i32; +}; + +[[group(0), binding(0)]] +var v : [[access(read_write)]] S; + +[[stage(compute)]] +fn main() { + let p : ptr = &v.a; + let use : i32 = *p + 1; +} diff --git a/test/ptr_ref/load/local/ptr_storage.wgsl.expected.hlsl b/test/ptr_ref/load/local/ptr_storage.wgsl.expected.hlsl new file mode 100644 index 0000000000..c77b5fd355 --- /dev/null +++ b/test/ptr_ref/load/local/ptr_storage.wgsl.expected.hlsl @@ -0,0 +1 @@ +SKIP: error: cannot take the address of expression diff --git a/test/ptr_ref/load/local/ptr_storage.wgsl.expected.msl b/test/ptr_ref/load/local/ptr_storage.wgsl.expected.msl new file mode 100644 index 0000000000..17d0c1bf6f --- /dev/null +++ b/test/ptr_ref/load/local/ptr_storage.wgsl.expected.msl @@ -0,0 +1,13 @@ +#include + +using namespace metal; +struct S { + /* 0x0000 */ int a; +}; + +kernel void tint_symbol(device S& v [[buffer(0)]]) { + device int* const p = &(v.a); + int const use = (*(p) + 1); + return; +} + diff --git a/test/ptr_ref/load/local/ptr_storage.wgsl.expected.spvasm b/test/ptr_ref/load/local/ptr_storage.wgsl.expected.spvasm new file mode 100644 index 0000000000..e0e29bce0b --- /dev/null +++ b/test/ptr_ref/load/local/ptr_storage.wgsl.expected.spvasm @@ -0,0 +1,34 @@ +; 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 %S "S" + OpMemberName %S 0 "a" + OpName %v "v" + OpName %main "main" + OpDecorate %S Block + OpMemberDecorate %S 0 Offset 0 + OpDecorate %v DescriptorSet 0 + OpDecorate %v Binding 0 + %int = OpTypeInt 32 1 + %S = OpTypeStruct %int +%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S + %v = OpVariable %_ptr_StorageBuffer_S StorageBuffer + %void = OpTypeVoid + %5 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %uint_0 = OpConstant %uint 0 +%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int + %int_1 = OpConstant %int 1 + %main = OpFunction %void None %5 + %8 = OpLabel + %13 = OpAccessChain %_ptr_StorageBuffer_int %v %uint_0 + %15 = OpLoad %int %13 + %17 = OpIAdd %int %15 %int_1 + OpReturn + OpFunctionEnd diff --git a/test/ptr_ref/load/local/ptr_storage.wgsl.expected.wgsl b/test/ptr_ref/load/local/ptr_storage.wgsl.expected.wgsl new file mode 100644 index 0000000000..0c88535dda --- /dev/null +++ b/test/ptr_ref/load/local/ptr_storage.wgsl.expected.wgsl @@ -0,0 +1,12 @@ +[[block]] +struct S { + a : i32; +}; + +[[group(0), binding(0)]] var v : [[access(read_write)]] S; + +[[stage(compute)]] +fn main() { + let p : ptr = &(v.a); + let use : i32 = (*(p) + 1); +} diff --git a/test/ptr_ref/load/local/ptr_uniform.wgsl b/test/ptr_ref/load/local/ptr_uniform.wgsl new file mode 100644 index 0000000000..3dd38d469c --- /dev/null +++ b/test/ptr_ref/load/local/ptr_uniform.wgsl @@ -0,0 +1,13 @@ +[[block]] +struct S { + a : i32; +}; + +[[group(0), binding(0)]] +var v : S; + +[[stage(compute)]] +fn main() { + let p : ptr = &v.a; + let use : i32 = *p + 1; +} diff --git a/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.hlsl b/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.hlsl new file mode 100644 index 0000000000..b3db42f63d --- /dev/null +++ b/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.hlsl @@ -0,0 +1 @@ +SKIP: Failed to generate: error: pointers not supported in HLSL diff --git a/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.msl b/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.msl new file mode 100644 index 0000000000..9008b2fbd3 --- /dev/null +++ b/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.msl @@ -0,0 +1,13 @@ +#include + +using namespace metal; +struct S { + /* 0x0000 */ int a; +}; + +kernel void tint_symbol(constant S& v [[buffer(0)]]) { + constant int* const p = &(v.a); + int const use = (*(p) + 1); + return; +} + diff --git a/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.spvasm b/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.spvasm new file mode 100644 index 0000000000..6810fb2f1c --- /dev/null +++ b/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.spvasm @@ -0,0 +1,34 @@ +; 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 %S "S" + OpMemberName %S 0 "a" + OpName %v "v" + OpName %main "main" + OpDecorate %S Block + OpMemberDecorate %S 0 Offset 0 + OpDecorate %v DescriptorSet 0 + OpDecorate %v Binding 0 + %int = OpTypeInt 32 1 + %S = OpTypeStruct %int +%_ptr_Uniform_S = OpTypePointer Uniform %S + %v = OpVariable %_ptr_Uniform_S Uniform + %void = OpTypeVoid + %5 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_int = OpTypePointer Uniform %int + %int_1 = OpConstant %int 1 + %main = OpFunction %void None %5 + %8 = OpLabel + %13 = OpAccessChain %_ptr_Uniform_int %v %uint_0 + %15 = OpLoad %int %13 + %17 = OpIAdd %int %15 %int_1 + OpReturn + OpFunctionEnd diff --git a/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.wgsl b/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.wgsl new file mode 100644 index 0000000000..6cc99839ae --- /dev/null +++ b/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.wgsl @@ -0,0 +1,12 @@ +[[block]] +struct S { + a : i32; +}; + +[[group(0), binding(0)]] var v : S; + +[[stage(compute)]] +fn main() { + let p : ptr = &(v.a); + let use : i32 = (*(p) + 1); +} diff --git a/test/ptr_ref/load/local/ptr_workgroup.wgsl b/test/ptr_ref/load/local/ptr_workgroup.wgsl new file mode 100644 index 0000000000..3327d40db2 --- /dev/null +++ b/test/ptr_ref/load/local/ptr_workgroup.wgsl @@ -0,0 +1,8 @@ +var i : i32; + +[[stage(compute)]] +fn main() { + i = 123; + let p : ptr = &i; + let use : i32 = *p + 1; +} diff --git a/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.hlsl b/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.hlsl new file mode 100644 index 0000000000..b3db42f63d --- /dev/null +++ b/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.hlsl @@ -0,0 +1 @@ +SKIP: Failed to generate: error: pointers not supported in HLSL diff --git a/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.msl b/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.msl new file mode 100644 index 0000000000..1b2d3db8ae --- /dev/null +++ b/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.msl @@ -0,0 +1 @@ +SKIP: TINT_UNIMPLEMENTED crbug.com/tint/726: module-scope private and workgroup variables not yet implemented diff --git a/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.spvasm b/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.spvasm new file mode 100644 index 0000000000..1cda0154fd --- /dev/null +++ b/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.spvasm @@ -0,0 +1,25 @@ +; 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 %i "i" + OpName %main "main" + %int = OpTypeInt 32 1 +%_ptr_Workgroup_int = OpTypePointer Workgroup %int + %i = OpVariable %_ptr_Workgroup_int Workgroup + %void = OpTypeVoid + %4 = OpTypeFunction %void + %int_123 = OpConstant %int 123 + %int_1 = OpConstant %int 1 + %main = OpFunction %void None %4 + %7 = OpLabel + OpStore %i %int_123 + %11 = OpLoad %int %i + %13 = OpIAdd %int %11 %int_1 + OpReturn + OpFunctionEnd diff --git a/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.wgsl b/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.wgsl new file mode 100644 index 0000000000..3f9fb08425 --- /dev/null +++ b/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.wgsl @@ -0,0 +1,8 @@ +var i : i32; + +[[stage(compute)]] +fn main() { + i = 123; + let p : ptr = &(i); + let use : i32 = (*(p) + 1); +} diff --git a/test/ptr_ref/load/param/ptr.spvasm.expected.msl b/test/ptr_ref/load/param/ptr.spvasm.expected.msl index cbe0a55a64..cf1067f7c1 100644 --- a/test/ptr_ref/load/param/ptr.spvasm.expected.msl +++ b/test/ptr_ref/load/param/ptr.spvasm.expected.msl @@ -1,7 +1,7 @@ #include using namespace metal; -int func(int value, int* pointer) { +int func(int value, thread int* pointer) { int const x_9 = *(pointer); return (value + x_9); } diff --git a/test/ptr_ref/load/param/ptr.wgsl.expected.msl b/test/ptr_ref/load/param/ptr.wgsl.expected.msl index 48556e1555..20a0cc99c0 100644 --- a/test/ptr_ref/load/param/ptr.wgsl.expected.msl +++ b/test/ptr_ref/load/param/ptr.wgsl.expected.msl @@ -1,7 +1,7 @@ #include using namespace metal; -int func(int value, int* pointer) { +int func(int value, thread int* pointer) { return (value + *(pointer)); } diff --git a/test/ptr_ref/store/local/i32.wgsl.expected.msl b/test/ptr_ref/store/local/i32.wgsl.expected.msl index 6674f82ca9..3c243ed285 100644 --- a/test/ptr_ref/store/local/i32.wgsl.expected.msl +++ b/test/ptr_ref/store/local/i32.wgsl.expected.msl @@ -3,7 +3,7 @@ using namespace metal; kernel void tint_symbol() { int i = 123; - int* const p = &(i); + thread int* const p = &(i); *(p) = 123; *(p) = ((100 + 20) + 3); return; diff --git a/test/ptr_ref/store/param/ptr.spvasm.expected.msl b/test/ptr_ref/store/param/ptr.spvasm.expected.msl index 5d459da3dd..0633ba656c 100644 --- a/test/ptr_ref/store/param/ptr.spvasm.expected.msl +++ b/test/ptr_ref/store/param/ptr.spvasm.expected.msl @@ -1,7 +1,7 @@ #include using namespace metal; -void func(int value, int* pointer) { +void func(int value, thread int* pointer) { *(pointer) = value; return; } diff --git a/test/ptr_ref/store/param/ptr.wgsl.expected.msl b/test/ptr_ref/store/param/ptr.wgsl.expected.msl index 517183a5e8..9a64bcb243 100644 --- a/test/ptr_ref/store/param/ptr.wgsl.expected.msl +++ b/test/ptr_ref/store/param/ptr.wgsl.expected.msl @@ -1,7 +1,7 @@ #include using namespace metal; -void func(int value, int* pointer) { +void func(int value, thread int* pointer) { *(pointer) = value; }