Update test/tint/access to store results
This CL updates the `test/tint/access` tests to make sure the results
are stored back into a storage variable.
Change-Id: If9b003ed8eb7106345962f6f93c5ae5462d0a923
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/164541
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Auto-Submit: dan sinclair <dsinclair@chromium.org>
diff --git a/test/tint/access/let/matrix.wgsl b/test/tint/access/let/matrix.wgsl
index 3d96fba..e8a783c 100644
--- a/test/tint/access/let/matrix.wgsl
+++ b/test/tint/access/let/matrix.wgsl
@@ -1,6 +1,10 @@
+@group(0) @binding(0) var<storage, read_write> s: f32;
+
@compute @workgroup_size(1)
fn main() {
let m : mat3x3<f32> = mat3x3<f32>(vec3<f32>(1., 2., 3.), vec3<f32>(4., 5., 6.), vec3<f32>(7., 8., 9.));
let v : vec3<f32> = m[1];
let f : f32 = v[1];
+
+ s = f;
}
diff --git a/test/tint/access/let/matrix.wgsl.expected.dxc.hlsl b/test/tint/access/let/matrix.wgsl.expected.dxc.hlsl
index 38129cf..10d9f08 100644
--- a/test/tint/access/let/matrix.wgsl.expected.dxc.hlsl
+++ b/test/tint/access/let/matrix.wgsl.expected.dxc.hlsl
@@ -1,7 +1,10 @@
+RWByteAddressBuffer s : register(u0);
+
[numthreads(1, 1, 1)]
void main() {
const float3x3 m = float3x3(float3(1.0f, 2.0f, 3.0f), float3(4.0f, 5.0f, 6.0f), float3(7.0f, 8.0f, 9.0f));
const float3 v = m[1];
const float f = v[1];
+ s.Store(0u, asuint(f));
return;
}
diff --git a/test/tint/access/let/matrix.wgsl.expected.fxc.hlsl b/test/tint/access/let/matrix.wgsl.expected.fxc.hlsl
index 38129cf..10d9f08 100644
--- a/test/tint/access/let/matrix.wgsl.expected.fxc.hlsl
+++ b/test/tint/access/let/matrix.wgsl.expected.fxc.hlsl
@@ -1,7 +1,10 @@
+RWByteAddressBuffer s : register(u0);
+
[numthreads(1, 1, 1)]
void main() {
const float3x3 m = float3x3(float3(1.0f, 2.0f, 3.0f), float3(4.0f, 5.0f, 6.0f), float3(7.0f, 8.0f, 9.0f));
const float3 v = m[1];
const float f = v[1];
+ s.Store(0u, asuint(f));
return;
}
diff --git a/test/tint/access/let/matrix.wgsl.expected.glsl b/test/tint/access/let/matrix.wgsl.expected.glsl
index 0add115..5296123 100644
--- a/test/tint/access/let/matrix.wgsl.expected.glsl
+++ b/test/tint/access/let/matrix.wgsl.expected.glsl
@@ -1,9 +1,14 @@
#version 310 es
+layout(binding = 0, std430) buffer s_block_ssbo {
+ float inner;
+} s;
+
void tint_symbol() {
mat3 m = mat3(vec3(1.0f, 2.0f, 3.0f), vec3(4.0f, 5.0f, 6.0f), vec3(7.0f, 8.0f, 9.0f));
vec3 v = m[1];
float f = v[1];
+ s.inner = f;
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
diff --git a/test/tint/access/let/matrix.wgsl.expected.msl b/test/tint/access/let/matrix.wgsl.expected.msl
index f8f945d..6f2e4ee 100644
--- a/test/tint/access/let/matrix.wgsl.expected.msl
+++ b/test/tint/access/let/matrix.wgsl.expected.msl
@@ -1,10 +1,11 @@
#include <metal_stdlib>
using namespace metal;
-kernel void tint_symbol() {
+kernel void tint_symbol(device float* tint_symbol_1 [[buffer(0)]]) {
float3x3 const 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];
float const f = v[1];
+ *(tint_symbol_1) = f;
return;
}
diff --git a/test/tint/access/let/matrix.wgsl.expected.spvasm b/test/tint/access/let/matrix.wgsl.expected.spvasm
index cbc5aca..f537d63 100644
--- a/test/tint/access/let/matrix.wgsl.expected.spvasm
+++ b/test/tint/access/let/matrix.wgsl.expected.spvasm
@@ -1,36 +1,51 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 25
+; Bound: 32
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
+ OpName %s_block "s_block"
+ OpMemberName %s_block 0 "inner"
+ OpName %s "s"
OpName %main "main"
- %void = OpTypeVoid
- %1 = OpTypeFunction %void
+ OpDecorate %s_block Block
+ OpMemberDecorate %s_block 0 Offset 0
+ OpDecorate %s DescriptorSet 0
+ OpDecorate %s Binding 0
%float = OpTypeFloat 32
+ %s_block = OpTypeStruct %float
+%_ptr_StorageBuffer_s_block = OpTypePointer StorageBuffer %s_block
+ %s = OpVariable %_ptr_StorageBuffer_s_block StorageBuffer
+ %void = OpTypeVoid
+ %5 = OpTypeFunction %void
%v3float = OpTypeVector %float 3
%mat3v3float = OpTypeMatrix %v3float 3
%float_1 = OpConstant %float 1
%float_2 = OpConstant %float 2
%float_3 = OpConstant %float 3
- %11 = OpConstantComposite %v3float %float_1 %float_2 %float_3
+ %14 = OpConstantComposite %v3float %float_1 %float_2 %float_3
%float_4 = OpConstant %float 4
%float_5 = OpConstant %float 5
%float_6 = OpConstant %float 6
- %15 = OpConstantComposite %v3float %float_4 %float_5 %float_6
+ %18 = OpConstantComposite %v3float %float_4 %float_5 %float_6
%float_7 = OpConstant %float 7
%float_8 = OpConstant %float 8
%float_9 = OpConstant %float 9
- %19 = OpConstantComposite %v3float %float_7 %float_8 %float_9
- %20 = OpConstantComposite %mat3v3float %11 %15 %19
+ %22 = OpConstantComposite %v3float %float_7 %float_8 %float_9
+ %23 = OpConstantComposite %mat3v3float %14 %18 %22
%int = OpTypeInt 32 1
%int_1 = OpConstant %int 1
- %main = OpFunction %void None %1
- %4 = OpLabel
- %23 = OpCompositeExtract %v3float %20 1
- %24 = OpCompositeExtract %float %23 1
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+ %main = OpFunction %void None %5
+ %8 = OpLabel
+ %26 = OpCompositeExtract %v3float %23 1
+ %27 = OpCompositeExtract %float %26 1
+ %31 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_0
+ OpStore %31 %27
OpReturn
OpFunctionEnd
diff --git a/test/tint/access/let/matrix.wgsl.expected.wgsl b/test/tint/access/let/matrix.wgsl.expected.wgsl
index bdfd207..e33c47d 100644
--- a/test/tint/access/let/matrix.wgsl.expected.wgsl
+++ b/test/tint/access/let/matrix.wgsl.expected.wgsl
@@ -1,6 +1,9 @@
+@group(0) @binding(0) var<storage, read_write> s : f32;
+
@compute @workgroup_size(1)
fn main() {
let m : mat3x3<f32> = mat3x3<f32>(vec3<f32>(1.0, 2.0, 3.0), vec3<f32>(4.0, 5.0, 6.0), vec3<f32>(7.0, 8.0, 9.0));
let v : vec3<f32> = m[1];
let f : f32 = v[1];
+ s = f;
}
diff --git a/test/tint/access/let/vector.wgsl b/test/tint/access/let/vector.wgsl
index 87604a9..a6b8f53 100644
--- a/test/tint/access/let/vector.wgsl
+++ b/test/tint/access/let/vector.wgsl
@@ -1,7 +1,11 @@
+@group(0) @binding(0) var<storage, read_write> s: vec3f;
+
@compute @workgroup_size(1)
fn main() {
let v : vec3<f32> = vec3<f32>(1., 2., 3.);
let scalar : f32 = v.y;
let swizzle2 : vec2<f32> = v.xz;
let swizzle3 : vec3<f32> = v.xzy;
+
+ s = vec3(scalar) + vec3(swizzle2, 1) + swizzle3;
}
diff --git a/test/tint/access/let/vector.wgsl.expected.dxc.hlsl b/test/tint/access/let/vector.wgsl.expected.dxc.hlsl
index d86c8ed..89c501a 100644
--- a/test/tint/access/let/vector.wgsl.expected.dxc.hlsl
+++ b/test/tint/access/let/vector.wgsl.expected.dxc.hlsl
@@ -1,8 +1,11 @@
+RWByteAddressBuffer s : register(u0);
+
[numthreads(1, 1, 1)]
void main() {
const float3 v = float3(1.0f, 2.0f, 3.0f);
const float scalar = v.y;
const float2 swizzle2 = v.xz;
const float3 swizzle3 = v.xzy;
+ s.Store3(0u, asuint(((float3((scalar).xxx) + float3(swizzle2, 1.0f)) + swizzle3)));
return;
}
diff --git a/test/tint/access/let/vector.wgsl.expected.fxc.hlsl b/test/tint/access/let/vector.wgsl.expected.fxc.hlsl
index d86c8ed..89c501a 100644
--- a/test/tint/access/let/vector.wgsl.expected.fxc.hlsl
+++ b/test/tint/access/let/vector.wgsl.expected.fxc.hlsl
@@ -1,8 +1,11 @@
+RWByteAddressBuffer s : register(u0);
+
[numthreads(1, 1, 1)]
void main() {
const float3 v = float3(1.0f, 2.0f, 3.0f);
const float scalar = v.y;
const float2 swizzle2 = v.xz;
const float3 swizzle3 = v.xzy;
+ s.Store3(0u, asuint(((float3((scalar).xxx) + float3(swizzle2, 1.0f)) + swizzle3)));
return;
}
diff --git a/test/tint/access/let/vector.wgsl.expected.glsl b/test/tint/access/let/vector.wgsl.expected.glsl
index 4f43d1f..45fe702 100644
--- a/test/tint/access/let/vector.wgsl.expected.glsl
+++ b/test/tint/access/let/vector.wgsl.expected.glsl
@@ -1,10 +1,15 @@
#version 310 es
+layout(binding = 0, std430) buffer s_block_ssbo {
+ vec3 inner;
+} s;
+
void tint_symbol() {
vec3 v = vec3(1.0f, 2.0f, 3.0f);
float scalar = v.y;
vec2 swizzle2 = v.xz;
vec3 swizzle3 = v.xzy;
+ s.inner = ((vec3(scalar) + vec3(swizzle2, 1.0f)) + swizzle3);
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
diff --git a/test/tint/access/let/vector.wgsl.expected.msl b/test/tint/access/let/vector.wgsl.expected.msl
index 2c7a0d5..e0842c6 100644
--- a/test/tint/access/let/vector.wgsl.expected.msl
+++ b/test/tint/access/let/vector.wgsl.expected.msl
@@ -1,11 +1,12 @@
#include <metal_stdlib>
using namespace metal;
-kernel void tint_symbol() {
+kernel void tint_symbol(device packed_float3* tint_symbol_1 [[buffer(0)]]) {
float3 const v = float3(1.0f, 2.0f, 3.0f);
float const scalar = v[1];
float2 const swizzle2 = v.xz;
float3 const swizzle3 = v.xzy;
+ *(tint_symbol_1) = packed_float3(((float3(scalar) + float3(swizzle2, 1.0f)) + swizzle3));
return;
}
diff --git a/test/tint/access/let/vector.wgsl.expected.spvasm b/test/tint/access/let/vector.wgsl.expected.spvasm
index 9de59d2..e4d25ff 100644
--- a/test/tint/access/let/vector.wgsl.expected.spvasm
+++ b/test/tint/access/let/vector.wgsl.expected.spvasm
@@ -1,26 +1,47 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 15
+; Bound: 28
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
+ OpName %s_block "s_block"
+ OpMemberName %s_block 0 "inner"
+ OpName %s "s"
OpName %main "main"
- %void = OpTypeVoid
- %1 = OpTypeFunction %void
+ OpDecorate %s_block Block
+ OpMemberDecorate %s_block 0 Offset 0
+ OpDecorate %s DescriptorSet 0
+ OpDecorate %s Binding 0
%float = OpTypeFloat 32
%v3float = OpTypeVector %float 3
+ %s_block = OpTypeStruct %v3float
+%_ptr_StorageBuffer_s_block = OpTypePointer StorageBuffer %s_block
+ %s = OpVariable %_ptr_StorageBuffer_s_block StorageBuffer
+ %void = OpTypeVoid
+ %6 = OpTypeFunction %void
%float_1 = OpConstant %float 1
%float_2 = OpConstant %float 2
%float_3 = OpConstant %float 3
- %10 = OpConstantComposite %v3float %float_1 %float_2 %float_3
+ %13 = OpConstantComposite %v3float %float_1 %float_2 %float_3
%v2float = OpTypeVector %float 2
- %main = OpFunction %void None %1
- %4 = OpLabel
- %11 = OpCompositeExtract %float %10 1
- %13 = OpVectorShuffle %v2float %10 %10 0 2
- %14 = OpVectorShuffle %v3float %10 %10 0 2 1
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v3float = OpTypePointer StorageBuffer %v3float
+ %main = OpFunction %void None %6
+ %9 = OpLabel
+ %14 = OpCompositeExtract %float %13 1
+ %16 = OpVectorShuffle %v2float %13 %13 0 2
+ %17 = OpVectorShuffle %v3float %13 %13 0 2 1
+ %21 = OpAccessChain %_ptr_StorageBuffer_v3float %s %uint_0
+ %22 = OpCompositeConstruct %v3float %14 %14 %14
+ %23 = OpCompositeExtract %float %16 0
+ %24 = OpCompositeExtract %float %16 1
+ %25 = OpCompositeConstruct %v3float %23 %24 %float_1
+ %26 = OpFAdd %v3float %22 %25
+ %27 = OpFAdd %v3float %26 %17
+ OpStore %21 %27
OpReturn
OpFunctionEnd
diff --git a/test/tint/access/let/vector.wgsl.expected.wgsl b/test/tint/access/let/vector.wgsl.expected.wgsl
index 6646867..aa937a7 100644
--- a/test/tint/access/let/vector.wgsl.expected.wgsl
+++ b/test/tint/access/let/vector.wgsl.expected.wgsl
@@ -1,7 +1,10 @@
+@group(0) @binding(0) var<storage, read_write> s : vec3f;
+
@compute @workgroup_size(1)
fn main() {
let v : vec3<f32> = vec3<f32>(1.0, 2.0, 3.0);
let scalar : f32 = v.y;
let swizzle2 : vec2<f32> = v.xz;
let swizzle3 : vec3<f32> = v.xzy;
+ s = ((vec3(scalar) + vec3(swizzle2, 1)) + swizzle3);
}
diff --git a/test/tint/access/var/matrix.wgsl b/test/tint/access/var/matrix.wgsl
index 92db7c0..2cef97d 100644
--- a/test/tint/access/var/matrix.wgsl
+++ b/test/tint/access/var/matrix.wgsl
@@ -1,6 +1,10 @@
+@group(0) @binding(0) var<storage, read_write> s: f32;
+
@compute @workgroup_size(1)
fn main() {
var m : mat3x3<f32>;
let v : vec3<f32> = m[1];
let f : f32 = v[1];
+
+ s = f;
}
diff --git a/test/tint/access/var/matrix.wgsl.expected.dxc.hlsl b/test/tint/access/var/matrix.wgsl.expected.dxc.hlsl
index b402c70..8653170 100644
--- a/test/tint/access/var/matrix.wgsl.expected.dxc.hlsl
+++ b/test/tint/access/var/matrix.wgsl.expected.dxc.hlsl
@@ -1,7 +1,10 @@
+RWByteAddressBuffer s : register(u0);
+
[numthreads(1, 1, 1)]
void main() {
float3x3 m = float3x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
const float3 v = m[1];
const float f = v[1];
+ s.Store(0u, asuint(f));
return;
}
diff --git a/test/tint/access/var/matrix.wgsl.expected.fxc.hlsl b/test/tint/access/var/matrix.wgsl.expected.fxc.hlsl
index b402c70..8653170 100644
--- a/test/tint/access/var/matrix.wgsl.expected.fxc.hlsl
+++ b/test/tint/access/var/matrix.wgsl.expected.fxc.hlsl
@@ -1,7 +1,10 @@
+RWByteAddressBuffer s : register(u0);
+
[numthreads(1, 1, 1)]
void main() {
float3x3 m = float3x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
const float3 v = m[1];
const float f = v[1];
+ s.Store(0u, asuint(f));
return;
}
diff --git a/test/tint/access/var/matrix.wgsl.expected.glsl b/test/tint/access/var/matrix.wgsl.expected.glsl
index 2c0d9cf..7f85a9d 100644
--- a/test/tint/access/var/matrix.wgsl.expected.glsl
+++ b/test/tint/access/var/matrix.wgsl.expected.glsl
@@ -1,9 +1,14 @@
#version 310 es
+layout(binding = 0, std430) buffer s_block_ssbo {
+ float inner;
+} s;
+
void tint_symbol() {
mat3 m = mat3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
vec3 v = m[1];
float f = v[1];
+ s.inner = f;
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
diff --git a/test/tint/access/var/matrix.wgsl.expected.msl b/test/tint/access/var/matrix.wgsl.expected.msl
index 95fd2d6..7b13fc3 100644
--- a/test/tint/access/var/matrix.wgsl.expected.msl
+++ b/test/tint/access/var/matrix.wgsl.expected.msl
@@ -1,10 +1,11 @@
#include <metal_stdlib>
using namespace metal;
-kernel void tint_symbol() {
+kernel void tint_symbol(device float* tint_symbol_1 [[buffer(0)]]) {
float3x3 m = float3x3(0.0f);
float3 const v = m[1];
float const f = v[1];
+ *(tint_symbol_1) = f;
return;
}
diff --git a/test/tint/access/var/matrix.wgsl.expected.spvasm b/test/tint/access/var/matrix.wgsl.expected.spvasm
index 9d735ff..39846c6 100644
--- a/test/tint/access/var/matrix.wgsl.expected.spvasm
+++ b/test/tint/access/var/matrix.wgsl.expected.spvasm
@@ -1,29 +1,44 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 17
+; Bound: 24
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
+ OpName %s_block "s_block"
+ OpMemberName %s_block 0 "inner"
+ OpName %s "s"
OpName %main "main"
OpName %m "m"
- %void = OpTypeVoid
- %1 = OpTypeFunction %void
+ OpDecorate %s_block Block
+ OpMemberDecorate %s_block 0 Offset 0
+ OpDecorate %s DescriptorSet 0
+ OpDecorate %s Binding 0
%float = OpTypeFloat 32
+ %s_block = OpTypeStruct %float
+%_ptr_StorageBuffer_s_block = OpTypePointer StorageBuffer %s_block
+ %s = OpVariable %_ptr_StorageBuffer_s_block StorageBuffer
+ %void = OpTypeVoid
+ %5 = OpTypeFunction %void
%v3float = OpTypeVector %float 3
%mat3v3float = OpTypeMatrix %v3float 3
%_ptr_Function_mat3v3float = OpTypePointer Function %mat3v3float
- %10 = OpConstantNull %mat3v3float
+ %13 = OpConstantNull %mat3v3float
%int = OpTypeInt 32 1
%int_1 = OpConstant %int 1
%_ptr_Function_v3float = OpTypePointer Function %v3float
- %main = OpFunction %void None %1
- %4 = OpLabel
- %m = OpVariable %_ptr_Function_mat3v3float Function %10
- %14 = OpAccessChain %_ptr_Function_v3float %m %int_1
- %15 = OpLoad %v3float %14
- %16 = OpCompositeExtract %float %15 1
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+ %main = OpFunction %void None %5
+ %8 = OpLabel
+ %m = OpVariable %_ptr_Function_mat3v3float Function %13
+ %17 = OpAccessChain %_ptr_Function_v3float %m %int_1
+ %18 = OpLoad %v3float %17
+ %19 = OpCompositeExtract %float %18 1
+ %23 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_0
+ OpStore %23 %19
OpReturn
OpFunctionEnd
diff --git a/test/tint/access/var/matrix.wgsl.expected.wgsl b/test/tint/access/var/matrix.wgsl.expected.wgsl
index 92db7c0..5fde10b 100644
--- a/test/tint/access/var/matrix.wgsl.expected.wgsl
+++ b/test/tint/access/var/matrix.wgsl.expected.wgsl
@@ -1,6 +1,9 @@
+@group(0) @binding(0) var<storage, read_write> s : f32;
+
@compute @workgroup_size(1)
fn main() {
var m : mat3x3<f32>;
let v : vec3<f32> = m[1];
let f : f32 = v[1];
+ s = f;
}
diff --git a/test/tint/access/var/vector.wgsl b/test/tint/access/var/vector.wgsl
index 8cc17f8..6d217b0 100644
--- a/test/tint/access/var/vector.wgsl
+++ b/test/tint/access/var/vector.wgsl
@@ -1,7 +1,11 @@
+@group(0) @binding(0) var<storage, read_write> s: vec3f;
+
@compute @workgroup_size(1)
fn main() {
var v : vec3<f32>;
let scalar : f32 = v.y;
let swizzle2 : vec2<f32> = v.xz;
let swizzle3 : vec3<f32> = v.xzy;
+
+ s = vec3(scalar) + vec3(swizzle2, 1) + swizzle3;
}
diff --git a/test/tint/access/var/vector.wgsl.expected.dxc.hlsl b/test/tint/access/var/vector.wgsl.expected.dxc.hlsl
index 8b9f393..bb92af0 100644
--- a/test/tint/access/var/vector.wgsl.expected.dxc.hlsl
+++ b/test/tint/access/var/vector.wgsl.expected.dxc.hlsl
@@ -1,8 +1,11 @@
+RWByteAddressBuffer s : register(u0);
+
[numthreads(1, 1, 1)]
void main() {
float3 v = float3(0.0f, 0.0f, 0.0f);
const float scalar = v.y;
const float2 swizzle2 = v.xz;
const float3 swizzle3 = v.xzy;
+ s.Store3(0u, asuint(((float3((scalar).xxx) + float3(swizzle2, 1.0f)) + swizzle3)));
return;
}
diff --git a/test/tint/access/var/vector.wgsl.expected.fxc.hlsl b/test/tint/access/var/vector.wgsl.expected.fxc.hlsl
index 8b9f393..bb92af0 100644
--- a/test/tint/access/var/vector.wgsl.expected.fxc.hlsl
+++ b/test/tint/access/var/vector.wgsl.expected.fxc.hlsl
@@ -1,8 +1,11 @@
+RWByteAddressBuffer s : register(u0);
+
[numthreads(1, 1, 1)]
void main() {
float3 v = float3(0.0f, 0.0f, 0.0f);
const float scalar = v.y;
const float2 swizzle2 = v.xz;
const float3 swizzle3 = v.xzy;
+ s.Store3(0u, asuint(((float3((scalar).xxx) + float3(swizzle2, 1.0f)) + swizzle3)));
return;
}
diff --git a/test/tint/access/var/vector.wgsl.expected.glsl b/test/tint/access/var/vector.wgsl.expected.glsl
index bfb9f6b..6268525 100644
--- a/test/tint/access/var/vector.wgsl.expected.glsl
+++ b/test/tint/access/var/vector.wgsl.expected.glsl
@@ -1,10 +1,15 @@
#version 310 es
+layout(binding = 0, std430) buffer s_block_ssbo {
+ vec3 inner;
+} s;
+
void tint_symbol() {
vec3 v = vec3(0.0f, 0.0f, 0.0f);
float scalar = v.y;
vec2 swizzle2 = v.xz;
vec3 swizzle3 = v.xzy;
+ s.inner = ((vec3(scalar) + vec3(swizzle2, 1.0f)) + swizzle3);
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
diff --git a/test/tint/access/var/vector.wgsl.expected.msl b/test/tint/access/var/vector.wgsl.expected.msl
index 270c842..7020eed 100644
--- a/test/tint/access/var/vector.wgsl.expected.msl
+++ b/test/tint/access/var/vector.wgsl.expected.msl
@@ -1,11 +1,12 @@
#include <metal_stdlib>
using namespace metal;
-kernel void tint_symbol() {
+kernel void tint_symbol(device packed_float3* tint_symbol_1 [[buffer(0)]]) {
float3 v = 0.0f;
float const scalar = v[1];
float2 const swizzle2 = v.xz;
float3 const swizzle3 = v.xzy;
+ *(tint_symbol_1) = packed_float3(((float3(scalar) + float3(swizzle2, 1.0f)) + swizzle3));
return;
}
diff --git a/test/tint/access/var/vector.wgsl.expected.spvasm b/test/tint/access/var/vector.wgsl.expected.spvasm
index 74418bb..33e465a 100644
--- a/test/tint/access/var/vector.wgsl.expected.spvasm
+++ b/test/tint/access/var/vector.wgsl.expected.spvasm
@@ -1,32 +1,53 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 20
+; Bound: 33
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
+ OpName %s_block "s_block"
+ OpMemberName %s_block 0 "inner"
+ OpName %s "s"
OpName %main "main"
OpName %v "v"
- %void = OpTypeVoid
- %1 = OpTypeFunction %void
+ OpDecorate %s_block Block
+ OpMemberDecorate %s_block 0 Offset 0
+ OpDecorate %s DescriptorSet 0
+ OpDecorate %s Binding 0
%float = OpTypeFloat 32
%v3float = OpTypeVector %float 3
+ %s_block = OpTypeStruct %v3float
+%_ptr_StorageBuffer_s_block = OpTypePointer StorageBuffer %s_block
+ %s = OpVariable %_ptr_StorageBuffer_s_block StorageBuffer
+ %void = OpTypeVoid
+ %6 = OpTypeFunction %void
%_ptr_Function_v3float = OpTypePointer Function %v3float
- %9 = OpConstantNull %v3float
+ %12 = OpConstantNull %v3float
%uint = OpTypeInt 32 0
%uint_1 = OpConstant %uint 1
%_ptr_Function_float = OpTypePointer Function %float
%v2float = OpTypeVector %float 2
- %main = OpFunction %void None %1
- %4 = OpLabel
- %v = OpVariable %_ptr_Function_v3float Function %9
- %13 = OpAccessChain %_ptr_Function_float %v %uint_1
- %14 = OpLoad %float %13
- %15 = OpLoad %v3float %v
- %17 = OpVectorShuffle %v2float %15 %15 0 2
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v3float = OpTypePointer StorageBuffer %v3float
+ %float_1 = OpConstant %float 1
+ %main = OpFunction %void None %6
+ %9 = OpLabel
+ %v = OpVariable %_ptr_Function_v3float Function %12
+ %16 = OpAccessChain %_ptr_Function_float %v %uint_1
+ %17 = OpLoad %float %16
%18 = OpLoad %v3float %v
- %19 = OpVectorShuffle %v3float %18 %18 0 2 1
+ %20 = OpVectorShuffle %v2float %18 %18 0 2
+ %21 = OpLoad %v3float %v
+ %22 = OpVectorShuffle %v3float %21 %21 0 2 1
+ %25 = OpAccessChain %_ptr_StorageBuffer_v3float %s %uint_0
+ %26 = OpCompositeConstruct %v3float %17 %17 %17
+ %27 = OpCompositeExtract %float %20 0
+ %28 = OpCompositeExtract %float %20 1
+ %30 = OpCompositeConstruct %v3float %27 %28 %float_1
+ %31 = OpFAdd %v3float %26 %30
+ %32 = OpFAdd %v3float %31 %22
+ OpStore %25 %32
OpReturn
OpFunctionEnd
diff --git a/test/tint/access/var/vector.wgsl.expected.wgsl b/test/tint/access/var/vector.wgsl.expected.wgsl
index 8cc17f8..37ac31d 100644
--- a/test/tint/access/var/vector.wgsl.expected.wgsl
+++ b/test/tint/access/var/vector.wgsl.expected.wgsl
@@ -1,7 +1,10 @@
+@group(0) @binding(0) var<storage, read_write> s : vec3f;
+
@compute @workgroup_size(1)
fn main() {
var v : vec3<f32>;
let scalar : f32 = v.y;
let swizzle2 : vec2<f32> = v.xz;
let swizzle3 : vec3<f32> = v.xzy;
+ s = ((vec3(scalar) + vec3(swizzle2, 1)) + swizzle3);
}
