test/tint/builtins/gen: Use returned value
It appears that FXC and DXC do some validation post dead-code-elimination.
These tests have been updated so that the return value is assigned to a storage buffer, ensuring that all validation is performed.
Many DXC tests are affected by https://github.com/microsoft/DirectXShaderCompiler/issues/5082, which have been SKIP'ed.
Fixed: tint:1859
Change-Id: I0001a9a9821846cd0855c3d8ce2bec79ab8e64ef
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/122662
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: James Price <jrprice@google.com>
Auto-Submit: Ben Clayton <bclayton@google.com>
diff --git a/test/tint/builtins/gen/literal/transpose/06794e.wgsl b/test/tint/builtins/gen/literal/transpose/06794e.wgsl
index 53a99be..c9e2f62 100644
--- a/test/tint/builtins/gen/literal/transpose/06794e.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/06794e.wgsl
@@ -26,7 +26,9 @@
// fn transpose(mat<3, 3, f16>) -> mat<3, 3, f16>
fn transpose_06794e() {
var res: mat3x3<f16> = transpose(mat3x3<f16>(1.h, 1.h, 1.h, 1.h, 1.h, 1.h, 1.h, 1.h, 1.h));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat3x3<f16>;
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
diff --git a/test/tint/builtins/gen/literal/transpose/06794e.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/transpose/06794e.wgsl.expected.dxc.hlsl
index 3977313..eb13a3a 100644
--- a/test/tint/builtins/gen/literal/transpose/06794e.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/06794e.wgsl.expected.dxc.hlsl
@@ -1,5 +1,14 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, matrix<float16_t, 3, 3> value) {
+ prevent_dce.Store<vector<float16_t, 3> >((offset + 0u), value[0u]);
+ prevent_dce.Store<vector<float16_t, 3> >((offset + 8u), value[1u]);
+ prevent_dce.Store<vector<float16_t, 3> >((offset + 16u), value[2u]);
+}
+
void transpose_06794e() {
matrix<float16_t, 3, 3> res = matrix<float16_t, 3, 3>((float16_t(1.0h)).xxx, (float16_t(1.0h)).xxx, (float16_t(1.0h)).xxx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/06794e.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/transpose/06794e.wgsl.expected.fxc.hlsl
index 5d50e02..80d5ee1 100644
--- a/test/tint/builtins/gen/literal/transpose/06794e.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/06794e.wgsl.expected.fxc.hlsl
@@ -1,7 +1,16 @@
SKIP: FAILED
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, matrix<float16_t, 3, 3> value) {
+ prevent_dce.Store<vector<float16_t, 3> >((offset + 0u), value[0u]);
+ prevent_dce.Store<vector<float16_t, 3> >((offset + 8u), value[1u]);
+ prevent_dce.Store<vector<float16_t, 3> >((offset + 16u), value[2u]);
+}
+
void transpose_06794e() {
matrix<float16_t, 3, 3> res = matrix<float16_t, 3, 3>((float16_t(1.0h)).xxx, (float16_t(1.0h)).xxx, (float16_t(1.0h)).xxx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/06794e.wgsl.expected.glsl b/test/tint/builtins/gen/literal/transpose/06794e.wgsl.expected.glsl
index bceb985..7b06a26 100644
--- a/test/tint/builtins/gen/literal/transpose/06794e.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/transpose/06794e.wgsl.expected.glsl
@@ -1,8 +1,19 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16mat3 inner;
+} prevent_dce;
+
+void assign_and_preserve_padding_prevent_dce(f16mat3 value) {
+ prevent_dce.inner[0] = value[0u];
+ prevent_dce.inner[1] = value[1u];
+ prevent_dce.inner[2] = value[2u];
+}
+
void transpose_06794e() {
f16mat3 res = f16mat3(f16vec3(1.0hf), f16vec3(1.0hf), f16vec3(1.0hf));
+ assign_and_preserve_padding_prevent_dce(res);
}
vec4 vertex_main() {
@@ -22,8 +33,19 @@
#extension GL_AMD_gpu_shader_half_float : require
precision mediump float;
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16mat3 inner;
+} prevent_dce;
+
+void assign_and_preserve_padding_prevent_dce(f16mat3 value) {
+ prevent_dce.inner[0] = value[0u];
+ prevent_dce.inner[1] = value[1u];
+ prevent_dce.inner[2] = value[2u];
+}
+
void transpose_06794e() {
f16mat3 res = f16mat3(f16vec3(1.0hf), f16vec3(1.0hf), f16vec3(1.0hf));
+ assign_and_preserve_padding_prevent_dce(res);
}
void fragment_main() {
@@ -37,8 +59,19 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16mat3 inner;
+} prevent_dce;
+
+void assign_and_preserve_padding_prevent_dce(f16mat3 value) {
+ prevent_dce.inner[0] = value[0u];
+ prevent_dce.inner[1] = value[1u];
+ prevent_dce.inner[2] = value[2u];
+}
+
void transpose_06794e() {
f16mat3 res = f16mat3(f16vec3(1.0hf), f16vec3(1.0hf), f16vec3(1.0hf));
+ assign_and_preserve_padding_prevent_dce(res);
}
void compute_main() {
diff --git a/test/tint/builtins/gen/literal/transpose/06794e.wgsl.expected.msl b/test/tint/builtins/gen/literal/transpose/06794e.wgsl.expected.msl
index e11011f..20b2ccf 100644
--- a/test/tint/builtins/gen/literal/transpose/06794e.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/transpose/06794e.wgsl.expected.msl
@@ -1,33 +1,58 @@
#include <metal_stdlib>
using namespace metal;
-void transpose_06794e() {
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+struct tint_packed_vec3_f16_array_element {
+ /* 0x0000 */ packed_half3 elements;
+ /* 0x0006 */ tint_array<int8_t, 2> tint_pad;
+};
+
+void assign_and_preserve_padding(device tint_array<tint_packed_vec3_f16_array_element, 3>* const dest, half3x3 value) {
+ (*(dest))[0u].elements = packed_half3(value[0u]);
+ (*(dest))[1u].elements = packed_half3(value[1u]);
+ (*(dest))[2u].elements = packed_half3(value[2u]);
+}
+
+void transpose_06794e(device tint_array<tint_packed_vec3_f16_array_element, 3>* const tint_symbol_1) {
half3x3 res = half3x3(half3(1.0h), half3(1.0h), half3(1.0h));
+ assign_and_preserve_padding(tint_symbol_1, res);
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner() {
- transpose_06794e();
+float4 vertex_main_inner(device tint_array<tint_packed_vec3_f16_array_element, 3>* const tint_symbol_2) {
+ transpose_06794e(tint_symbol_2);
return float4(0.0f);
}
-vertex tint_symbol vertex_main() {
- float4 const inner_result = vertex_main_inner();
+vertex tint_symbol vertex_main(device tint_array<tint_packed_vec3_f16_array_element, 3>* tint_symbol_3 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_3);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main() {
- transpose_06794e();
+fragment void fragment_main(device tint_array<tint_packed_vec3_f16_array_element, 3>* tint_symbol_4 [[buffer(0)]]) {
+ transpose_06794e(tint_symbol_4);
return;
}
-kernel void compute_main() {
- transpose_06794e();
+kernel void compute_main(device tint_array<tint_packed_vec3_f16_array_element, 3>* tint_symbol_5 [[buffer(0)]]) {
+ transpose_06794e(tint_symbol_5);
return;
}
diff --git a/test/tint/builtins/gen/literal/transpose/06794e.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/transpose/06794e.wgsl.expected.spvasm
index 03c5725..a227368 100644
--- a/test/tint/builtins/gen/literal/transpose/06794e.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/transpose/06794e.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 36
+; Bound: 61
; Schema: 0
OpCapability Shader
OpCapability Float16
@@ -9,66 +9,107 @@
OpCapability StorageBuffer16BitAccess
OpCapability StorageInputOutput16
OpMemoryModel Logical GLSL450
- OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+ OpEntryPoint Vertex %vertex_main "vertex_main" %value_1 %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
OpEntryPoint GLCompute %compute_main "compute_main"
OpExecutionMode %fragment_main OriginUpperLeft
OpExecutionMode %compute_main LocalSize 1 1 1
- OpName %value "value"
+ OpName %value_1 "value_1"
OpName %vertex_point_size "vertex_point_size"
+ OpName %prevent_dce_block "prevent_dce_block"
+ OpMemberName %prevent_dce_block 0 "inner"
+ OpName %prevent_dce "prevent_dce"
+ OpName %assign_and_preserve_padding_prevent_dce "assign_and_preserve_padding_prevent_dce"
+ OpName %value "value"
OpName %transpose_06794e "transpose_06794e"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
OpName %vertex_main "vertex_main"
OpName %fragment_main "fragment_main"
OpName %compute_main "compute_main"
- OpDecorate %value BuiltIn Position
+ OpDecorate %value_1 BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
+ OpDecorate %prevent_dce_block Block
+ OpMemberDecorate %prevent_dce_block 0 Offset 0
+ OpMemberDecorate %prevent_dce_block 0 ColMajor
+ OpMemberDecorate %prevent_dce_block 0 MatrixStride 8
+ OpDecorate %prevent_dce DescriptorSet 2
+ OpDecorate %prevent_dce Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
%5 = OpConstantNull %v4float
- %value = OpVariable %_ptr_Output_v4float Output %5
+ %value_1 = OpVariable %_ptr_Output_v4float Output %5
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
- %void = OpTypeVoid
- %9 = OpTypeFunction %void
%half = OpTypeFloat 16
%v3half = OpTypeVector %half 3
%mat3v3half = OpTypeMatrix %v3half 3
+%prevent_dce_block = OpTypeStruct %mat3v3half
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+ %void = OpTypeVoid
+ %15 = OpTypeFunction %void %mat3v3half
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+ %int = OpTypeInt 32 1
+ %23 = OpConstantNull %int
+%_ptr_StorageBuffer_v3half = OpTypePointer StorageBuffer %v3half
+ %26 = OpConstantNull %uint
+ %int_1 = OpConstant %int 1
+ %uint_1 = OpConstant %uint 1
+ %int_2 = OpConstant %int 2
+ %uint_2 = OpConstant %uint 2
+ %36 = OpTypeFunction %void
%half_0x1p_0 = OpConstant %half 0x1p+0
- %17 = OpConstantComposite %v3half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
- %18 = OpConstantComposite %mat3v3half %17 %17 %17
+ %40 = OpConstantComposite %v3half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
+ %41 = OpConstantComposite %mat3v3half %40 %40 %40
%_ptr_Function_mat3v3half = OpTypePointer Function %mat3v3half
- %21 = OpConstantNull %mat3v3half
- %22 = OpTypeFunction %v4float
+ %44 = OpConstantNull %mat3v3half
+ %47 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
-%transpose_06794e = OpFunction %void None %9
- %12 = OpLabel
- %res = OpVariable %_ptr_Function_mat3v3half Function %21
- OpStore %res %18
+%assign_and_preserve_padding_prevent_dce = OpFunction %void None %15
+ %value = OpFunctionParameter %mat3v3half
+ %19 = OpLabel
+ %25 = OpAccessChain %_ptr_StorageBuffer_v3half %prevent_dce %uint_0 %23
+ %27 = OpCompositeExtract %v3half %value 0
+ OpStore %25 %27
+ %29 = OpAccessChain %_ptr_StorageBuffer_v3half %prevent_dce %uint_0 %int_1
+ %31 = OpCompositeExtract %v3half %value 1
+ OpStore %29 %31
+ %33 = OpAccessChain %_ptr_StorageBuffer_v3half %prevent_dce %uint_0 %int_2
+ %35 = OpCompositeExtract %v3half %value 2
+ OpStore %33 %35
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %22
- %24 = OpLabel
- %25 = OpFunctionCall %void %transpose_06794e
+%transpose_06794e = OpFunction %void None %36
+ %38 = OpLabel
+ %res = OpVariable %_ptr_Function_mat3v3half Function %44
+ OpStore %res %41
+ %46 = OpLoad %mat3v3half %res
+ %45 = OpFunctionCall %void %assign_and_preserve_padding_prevent_dce %46
+ OpReturn
+ OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %47
+ %49 = OpLabel
+ %50 = OpFunctionCall %void %transpose_06794e
OpReturnValue %5
OpFunctionEnd
-%vertex_main = OpFunction %void None %9
- %27 = OpLabel
- %28 = OpFunctionCall %v4float %vertex_main_inner
- OpStore %value %28
+%vertex_main = OpFunction %void None %36
+ %52 = OpLabel
+ %53 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value_1 %53
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
-%fragment_main = OpFunction %void None %9
- %31 = OpLabel
- %32 = OpFunctionCall %void %transpose_06794e
+%fragment_main = OpFunction %void None %36
+ %56 = OpLabel
+ %57 = OpFunctionCall %void %transpose_06794e
OpReturn
OpFunctionEnd
-%compute_main = OpFunction %void None %9
- %34 = OpLabel
- %35 = OpFunctionCall %void %transpose_06794e
+%compute_main = OpFunction %void None %36
+ %59 = OpLabel
+ %60 = OpFunctionCall %void %transpose_06794e
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/transpose/06794e.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/transpose/06794e.wgsl.expected.wgsl
index b1ba096..5025574 100644
--- a/test/tint/builtins/gen/literal/transpose/06794e.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/06794e.wgsl.expected.wgsl
@@ -2,8 +2,11 @@
fn transpose_06794e() {
var res : mat3x3<f16> = transpose(mat3x3<f16>(1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat3x3<f16>;
+
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
transpose_06794e();
diff --git a/test/tint/builtins/gen/literal/transpose/2585cd.wgsl b/test/tint/builtins/gen/literal/transpose/2585cd.wgsl
index 944e58d..95d65e9 100644
--- a/test/tint/builtins/gen/literal/transpose/2585cd.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/2585cd.wgsl
@@ -24,7 +24,9 @@
// fn transpose(mat<4, 3, f32>) -> mat<3, 4, f32>
fn transpose_2585cd() {
var res: mat3x4<f32> = transpose(mat4x3<f32>(1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat3x4<f32>;
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
diff --git a/test/tint/builtins/gen/literal/transpose/2585cd.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/transpose/2585cd.wgsl.expected.dxc.hlsl
index d86773b..00be678 100644
--- a/test/tint/builtins/gen/literal/transpose/2585cd.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/2585cd.wgsl.expected.dxc.hlsl
@@ -1,5 +1,14 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, float3x4 value) {
+ prevent_dce.Store4((offset + 0u), asuint(value[0u]));
+ prevent_dce.Store4((offset + 16u), asuint(value[1u]));
+ prevent_dce.Store4((offset + 32u), asuint(value[2u]));
+}
+
void transpose_2585cd() {
float3x4 res = float3x4((1.0f).xxxx, (1.0f).xxxx, (1.0f).xxxx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/2585cd.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/transpose/2585cd.wgsl.expected.fxc.hlsl
index d86773b..00be678 100644
--- a/test/tint/builtins/gen/literal/transpose/2585cd.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/2585cd.wgsl.expected.fxc.hlsl
@@ -1,5 +1,14 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, float3x4 value) {
+ prevent_dce.Store4((offset + 0u), asuint(value[0u]));
+ prevent_dce.Store4((offset + 16u), asuint(value[1u]));
+ prevent_dce.Store4((offset + 32u), asuint(value[2u]));
+}
+
void transpose_2585cd() {
float3x4 res = float3x4((1.0f).xxxx, (1.0f).xxxx, (1.0f).xxxx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/2585cd.wgsl.expected.glsl b/test/tint/builtins/gen/literal/transpose/2585cd.wgsl.expected.glsl
index 60836cf..6d632f3 100644
--- a/test/tint/builtins/gen/literal/transpose/2585cd.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/transpose/2585cd.wgsl.expected.glsl
@@ -1,7 +1,12 @@
#version 310 es
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ mat3x4 inner;
+} prevent_dce;
+
void transpose_2585cd() {
mat3x4 res = mat3x4(vec4(1.0f), vec4(1.0f), vec4(1.0f));
+ prevent_dce.inner = res;
}
vec4 vertex_main() {
@@ -20,8 +25,13 @@
#version 310 es
precision mediump float;
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ mat3x4 inner;
+} prevent_dce;
+
void transpose_2585cd() {
mat3x4 res = mat3x4(vec4(1.0f), vec4(1.0f), vec4(1.0f));
+ prevent_dce.inner = res;
}
void fragment_main() {
@@ -34,8 +44,13 @@
}
#version 310 es
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ mat3x4 inner;
+} prevent_dce;
+
void transpose_2585cd() {
mat3x4 res = mat3x4(vec4(1.0f), vec4(1.0f), vec4(1.0f));
+ prevent_dce.inner = res;
}
void compute_main() {
diff --git a/test/tint/builtins/gen/literal/transpose/2585cd.wgsl.expected.msl b/test/tint/builtins/gen/literal/transpose/2585cd.wgsl.expected.msl
index f0936d5..e2fd1af 100644
--- a/test/tint/builtins/gen/literal/transpose/2585cd.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/transpose/2585cd.wgsl.expected.msl
@@ -1,33 +1,34 @@
#include <metal_stdlib>
using namespace metal;
-void transpose_2585cd() {
+void transpose_2585cd(device float3x4* const tint_symbol_1) {
float3x4 res = float3x4(float4(1.0f), float4(1.0f), float4(1.0f));
+ *(tint_symbol_1) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner() {
- transpose_2585cd();
+float4 vertex_main_inner(device float3x4* const tint_symbol_2) {
+ transpose_2585cd(tint_symbol_2);
return float4(0.0f);
}
-vertex tint_symbol vertex_main() {
- float4 const inner_result = vertex_main_inner();
+vertex tint_symbol vertex_main(device float3x4* tint_symbol_3 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_3);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main() {
- transpose_2585cd();
+fragment void fragment_main(device float3x4* tint_symbol_4 [[buffer(0)]]) {
+ transpose_2585cd(tint_symbol_4);
return;
}
-kernel void compute_main() {
- transpose_2585cd();
+kernel void compute_main(device float3x4* tint_symbol_5 [[buffer(0)]]) {
+ transpose_2585cd(tint_symbol_5);
return;
}
diff --git a/test/tint/builtins/gen/literal/transpose/2585cd.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/transpose/2585cd.wgsl.expected.spvasm
index 4925227..a1ea658 100644
--- a/test/tint/builtins/gen/literal/transpose/2585cd.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/transpose/2585cd.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 33
+; Bound: 41
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -12,6 +12,9 @@
OpExecutionMode %compute_main LocalSize 1 1 1
OpName %value "value"
OpName %vertex_point_size "vertex_point_size"
+ OpName %prevent_dce_block "prevent_dce_block"
+ OpMemberName %prevent_dce_block 0 "inner"
+ OpName %prevent_dce "prevent_dce"
OpName %transpose_2585cd "transpose_2585cd"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
@@ -20,6 +23,12 @@
OpName %compute_main "compute_main"
OpDecorate %value BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
+ OpDecorate %prevent_dce_block Block
+ OpMemberDecorate %prevent_dce_block 0 Offset 0
+ OpMemberDecorate %prevent_dce_block 0 ColMajor
+ OpMemberDecorate %prevent_dce_block 0 MatrixStride 16
+ OpDecorate %prevent_dce DescriptorSet 2
+ OpDecorate %prevent_dce Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
@@ -28,40 +37,49 @@
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
- %void = OpTypeVoid
- %9 = OpTypeFunction %void
%mat3v4float = OpTypeMatrix %v4float 3
+%prevent_dce_block = OpTypeStruct %mat3v4float
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+ %void = OpTypeVoid
+ %13 = OpTypeFunction %void
%float_1 = OpConstant %float 1
- %15 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
- %16 = OpConstantComposite %mat3v4float %15 %15 %15
+ %18 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
+ %19 = OpConstantComposite %mat3v4float %18 %18 %18
%_ptr_Function_mat3v4float = OpTypePointer Function %mat3v4float
- %19 = OpConstantNull %mat3v4float
- %20 = OpTypeFunction %v4float
-%transpose_2585cd = OpFunction %void None %9
- %12 = OpLabel
- %res = OpVariable %_ptr_Function_mat3v4float Function %19
- OpStore %res %16
+ %22 = OpConstantNull %mat3v4float
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_mat3v4float = OpTypePointer StorageBuffer %mat3v4float
+ %28 = OpTypeFunction %v4float
+%transpose_2585cd = OpFunction %void None %13
+ %16 = OpLabel
+ %res = OpVariable %_ptr_Function_mat3v4float Function %22
+ OpStore %res %19
+ %26 = OpAccessChain %_ptr_StorageBuffer_mat3v4float %prevent_dce %uint_0
+ %27 = OpLoad %mat3v4float %res
+ OpStore %26 %27
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %20
- %22 = OpLabel
- %23 = OpFunctionCall %void %transpose_2585cd
+%vertex_main_inner = OpFunction %v4float None %28
+ %30 = OpLabel
+ %31 = OpFunctionCall %void %transpose_2585cd
OpReturnValue %5
OpFunctionEnd
-%vertex_main = OpFunction %void None %9
- %25 = OpLabel
- %26 = OpFunctionCall %v4float %vertex_main_inner
- OpStore %value %26
+%vertex_main = OpFunction %void None %13
+ %33 = OpLabel
+ %34 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %34
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
-%fragment_main = OpFunction %void None %9
- %28 = OpLabel
- %29 = OpFunctionCall %void %transpose_2585cd
+%fragment_main = OpFunction %void None %13
+ %36 = OpLabel
+ %37 = OpFunctionCall %void %transpose_2585cd
OpReturn
OpFunctionEnd
-%compute_main = OpFunction %void None %9
- %31 = OpLabel
- %32 = OpFunctionCall %void %transpose_2585cd
+%compute_main = OpFunction %void None %13
+ %39 = OpLabel
+ %40 = OpFunctionCall %void %transpose_2585cd
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/transpose/2585cd.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/transpose/2585cd.wgsl.expected.wgsl
index ca39b02..f254d65 100644
--- a/test/tint/builtins/gen/literal/transpose/2585cd.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/2585cd.wgsl.expected.wgsl
@@ -1,7 +1,10 @@
fn transpose_2585cd() {
var res : mat3x4<f32> = transpose(mat4x3<f32>(1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat3x4<f32>;
+
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
transpose_2585cd();
diff --git a/test/tint/builtins/gen/literal/transpose/31d679.wgsl b/test/tint/builtins/gen/literal/transpose/31d679.wgsl
index 3f15b9d..cec2b32 100644
--- a/test/tint/builtins/gen/literal/transpose/31d679.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/31d679.wgsl
@@ -24,7 +24,9 @@
// fn transpose(mat<2, 2, f32>) -> mat<2, 2, f32>
fn transpose_31d679() {
var res: mat2x2<f32> = transpose(mat2x2<f32>(1.f, 1.f, 1.f, 1.f));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat2x2<f32>;
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
diff --git a/test/tint/builtins/gen/literal/transpose/31d679.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/transpose/31d679.wgsl.expected.dxc.hlsl
index efb6810..3972cd4 100644
--- a/test/tint/builtins/gen/literal/transpose/31d679.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/31d679.wgsl.expected.dxc.hlsl
@@ -1,5 +1,13 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, float2x2 value) {
+ prevent_dce.Store2((offset + 0u), asuint(value[0u]));
+ prevent_dce.Store2((offset + 8u), asuint(value[1u]));
+}
+
void transpose_31d679() {
float2x2 res = float2x2((1.0f).xx, (1.0f).xx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/31d679.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/transpose/31d679.wgsl.expected.fxc.hlsl
index efb6810..3972cd4 100644
--- a/test/tint/builtins/gen/literal/transpose/31d679.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/31d679.wgsl.expected.fxc.hlsl
@@ -1,5 +1,13 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, float2x2 value) {
+ prevent_dce.Store2((offset + 0u), asuint(value[0u]));
+ prevent_dce.Store2((offset + 8u), asuint(value[1u]));
+}
+
void transpose_31d679() {
float2x2 res = float2x2((1.0f).xx, (1.0f).xx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/31d679.wgsl.expected.glsl b/test/tint/builtins/gen/literal/transpose/31d679.wgsl.expected.glsl
index 6e858e6..5948d0b 100644
--- a/test/tint/builtins/gen/literal/transpose/31d679.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/transpose/31d679.wgsl.expected.glsl
@@ -1,7 +1,12 @@
#version 310 es
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ mat2 inner;
+} prevent_dce;
+
void transpose_31d679() {
mat2 res = mat2(vec2(1.0f), vec2(1.0f));
+ prevent_dce.inner = res;
}
vec4 vertex_main() {
@@ -20,8 +25,13 @@
#version 310 es
precision mediump float;
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ mat2 inner;
+} prevent_dce;
+
void transpose_31d679() {
mat2 res = mat2(vec2(1.0f), vec2(1.0f));
+ prevent_dce.inner = res;
}
void fragment_main() {
@@ -34,8 +44,13 @@
}
#version 310 es
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ mat2 inner;
+} prevent_dce;
+
void transpose_31d679() {
mat2 res = mat2(vec2(1.0f), vec2(1.0f));
+ prevent_dce.inner = res;
}
void compute_main() {
diff --git a/test/tint/builtins/gen/literal/transpose/31d679.wgsl.expected.msl b/test/tint/builtins/gen/literal/transpose/31d679.wgsl.expected.msl
index 6170477..27eae81 100644
--- a/test/tint/builtins/gen/literal/transpose/31d679.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/transpose/31d679.wgsl.expected.msl
@@ -1,33 +1,34 @@
#include <metal_stdlib>
using namespace metal;
-void transpose_31d679() {
+void transpose_31d679(device float2x2* const tint_symbol_1) {
float2x2 res = float2x2(float2(1.0f), float2(1.0f));
+ *(tint_symbol_1) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner() {
- transpose_31d679();
+float4 vertex_main_inner(device float2x2* const tint_symbol_2) {
+ transpose_31d679(tint_symbol_2);
return float4(0.0f);
}
-vertex tint_symbol vertex_main() {
- float4 const inner_result = vertex_main_inner();
+vertex tint_symbol vertex_main(device float2x2* tint_symbol_3 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_3);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main() {
- transpose_31d679();
+fragment void fragment_main(device float2x2* tint_symbol_4 [[buffer(0)]]) {
+ transpose_31d679(tint_symbol_4);
return;
}
-kernel void compute_main() {
- transpose_31d679();
+kernel void compute_main(device float2x2* tint_symbol_5 [[buffer(0)]]) {
+ transpose_31d679(tint_symbol_5);
return;
}
diff --git a/test/tint/builtins/gen/literal/transpose/31d679.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/transpose/31d679.wgsl.expected.spvasm
index 0c2a9ca..e6b5c2f 100644
--- a/test/tint/builtins/gen/literal/transpose/31d679.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/transpose/31d679.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 34
+; Bound: 42
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -12,6 +12,9 @@
OpExecutionMode %compute_main LocalSize 1 1 1
OpName %value "value"
OpName %vertex_point_size "vertex_point_size"
+ OpName %prevent_dce_block "prevent_dce_block"
+ OpMemberName %prevent_dce_block 0 "inner"
+ OpName %prevent_dce "prevent_dce"
OpName %transpose_31d679 "transpose_31d679"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
@@ -20,6 +23,12 @@
OpName %compute_main "compute_main"
OpDecorate %value BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
+ OpDecorate %prevent_dce_block Block
+ OpMemberDecorate %prevent_dce_block 0 Offset 0
+ OpMemberDecorate %prevent_dce_block 0 ColMajor
+ OpMemberDecorate %prevent_dce_block 0 MatrixStride 8
+ OpDecorate %prevent_dce DescriptorSet 2
+ OpDecorate %prevent_dce Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
@@ -28,41 +37,50 @@
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
- %void = OpTypeVoid
- %9 = OpTypeFunction %void
%v2float = OpTypeVector %float 2
%mat2v2float = OpTypeMatrix %v2float 2
+%prevent_dce_block = OpTypeStruct %mat2v2float
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+ %void = OpTypeVoid
+ %14 = OpTypeFunction %void
%float_1 = OpConstant %float 1
- %16 = OpConstantComposite %v2float %float_1 %float_1
- %17 = OpConstantComposite %mat2v2float %16 %16
+ %19 = OpConstantComposite %v2float %float_1 %float_1
+ %20 = OpConstantComposite %mat2v2float %19 %19
%_ptr_Function_mat2v2float = OpTypePointer Function %mat2v2float
- %20 = OpConstantNull %mat2v2float
- %21 = OpTypeFunction %v4float
-%transpose_31d679 = OpFunction %void None %9
- %12 = OpLabel
- %res = OpVariable %_ptr_Function_mat2v2float Function %20
- OpStore %res %17
+ %23 = OpConstantNull %mat2v2float
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_mat2v2float = OpTypePointer StorageBuffer %mat2v2float
+ %29 = OpTypeFunction %v4float
+%transpose_31d679 = OpFunction %void None %14
+ %17 = OpLabel
+ %res = OpVariable %_ptr_Function_mat2v2float Function %23
+ OpStore %res %20
+ %27 = OpAccessChain %_ptr_StorageBuffer_mat2v2float %prevent_dce %uint_0
+ %28 = OpLoad %mat2v2float %res
+ OpStore %27 %28
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %21
- %23 = OpLabel
- %24 = OpFunctionCall %void %transpose_31d679
+%vertex_main_inner = OpFunction %v4float None %29
+ %31 = OpLabel
+ %32 = OpFunctionCall %void %transpose_31d679
OpReturnValue %5
OpFunctionEnd
-%vertex_main = OpFunction %void None %9
- %26 = OpLabel
- %27 = OpFunctionCall %v4float %vertex_main_inner
- OpStore %value %27
+%vertex_main = OpFunction %void None %14
+ %34 = OpLabel
+ %35 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %35
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
-%fragment_main = OpFunction %void None %9
- %29 = OpLabel
- %30 = OpFunctionCall %void %transpose_31d679
+%fragment_main = OpFunction %void None %14
+ %37 = OpLabel
+ %38 = OpFunctionCall %void %transpose_31d679
OpReturn
OpFunctionEnd
-%compute_main = OpFunction %void None %9
- %32 = OpLabel
- %33 = OpFunctionCall %void %transpose_31d679
+%compute_main = OpFunction %void None %14
+ %40 = OpLabel
+ %41 = OpFunctionCall %void %transpose_31d679
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/transpose/31d679.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/transpose/31d679.wgsl.expected.wgsl
index 32267a1..e2d0c07 100644
--- a/test/tint/builtins/gen/literal/transpose/31d679.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/31d679.wgsl.expected.wgsl
@@ -1,7 +1,10 @@
fn transpose_31d679() {
var res : mat2x2<f32> = transpose(mat2x2<f32>(1.0f, 1.0f, 1.0f, 1.0f));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat2x2<f32>;
+
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
transpose_31d679();
diff --git a/test/tint/builtins/gen/literal/transpose/31e37e.wgsl b/test/tint/builtins/gen/literal/transpose/31e37e.wgsl
index ec41dcb..5e27ba1 100644
--- a/test/tint/builtins/gen/literal/transpose/31e37e.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/31e37e.wgsl
@@ -24,7 +24,9 @@
// fn transpose(mat<4, 2, f32>) -> mat<2, 4, f32>
fn transpose_31e37e() {
var res: mat2x4<f32> = transpose(mat4x2<f32>(1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat2x4<f32>;
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
diff --git a/test/tint/builtins/gen/literal/transpose/31e37e.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/transpose/31e37e.wgsl.expected.dxc.hlsl
index 387bd69..9c0f8bb 100644
--- a/test/tint/builtins/gen/literal/transpose/31e37e.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/31e37e.wgsl.expected.dxc.hlsl
@@ -1,5 +1,13 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, float2x4 value) {
+ prevent_dce.Store4((offset + 0u), asuint(value[0u]));
+ prevent_dce.Store4((offset + 16u), asuint(value[1u]));
+}
+
void transpose_31e37e() {
float2x4 res = float2x4((1.0f).xxxx, (1.0f).xxxx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/31e37e.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/transpose/31e37e.wgsl.expected.fxc.hlsl
index 387bd69..9c0f8bb 100644
--- a/test/tint/builtins/gen/literal/transpose/31e37e.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/31e37e.wgsl.expected.fxc.hlsl
@@ -1,5 +1,13 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, float2x4 value) {
+ prevent_dce.Store4((offset + 0u), asuint(value[0u]));
+ prevent_dce.Store4((offset + 16u), asuint(value[1u]));
+}
+
void transpose_31e37e() {
float2x4 res = float2x4((1.0f).xxxx, (1.0f).xxxx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/31e37e.wgsl.expected.glsl b/test/tint/builtins/gen/literal/transpose/31e37e.wgsl.expected.glsl
index 031e2ab..c1dd309 100644
--- a/test/tint/builtins/gen/literal/transpose/31e37e.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/transpose/31e37e.wgsl.expected.glsl
@@ -1,7 +1,12 @@
#version 310 es
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ mat2x4 inner;
+} prevent_dce;
+
void transpose_31e37e() {
mat2x4 res = mat2x4(vec4(1.0f), vec4(1.0f));
+ prevent_dce.inner = res;
}
vec4 vertex_main() {
@@ -20,8 +25,13 @@
#version 310 es
precision mediump float;
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ mat2x4 inner;
+} prevent_dce;
+
void transpose_31e37e() {
mat2x4 res = mat2x4(vec4(1.0f), vec4(1.0f));
+ prevent_dce.inner = res;
}
void fragment_main() {
@@ -34,8 +44,13 @@
}
#version 310 es
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ mat2x4 inner;
+} prevent_dce;
+
void transpose_31e37e() {
mat2x4 res = mat2x4(vec4(1.0f), vec4(1.0f));
+ prevent_dce.inner = res;
}
void compute_main() {
diff --git a/test/tint/builtins/gen/literal/transpose/31e37e.wgsl.expected.msl b/test/tint/builtins/gen/literal/transpose/31e37e.wgsl.expected.msl
index 8cbf0ed..6c5f1d8 100644
--- a/test/tint/builtins/gen/literal/transpose/31e37e.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/transpose/31e37e.wgsl.expected.msl
@@ -1,33 +1,34 @@
#include <metal_stdlib>
using namespace metal;
-void transpose_31e37e() {
+void transpose_31e37e(device float2x4* const tint_symbol_1) {
float2x4 res = float2x4(float4(1.0f), float4(1.0f));
+ *(tint_symbol_1) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner() {
- transpose_31e37e();
+float4 vertex_main_inner(device float2x4* const tint_symbol_2) {
+ transpose_31e37e(tint_symbol_2);
return float4(0.0f);
}
-vertex tint_symbol vertex_main() {
- float4 const inner_result = vertex_main_inner();
+vertex tint_symbol vertex_main(device float2x4* tint_symbol_3 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_3);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main() {
- transpose_31e37e();
+fragment void fragment_main(device float2x4* tint_symbol_4 [[buffer(0)]]) {
+ transpose_31e37e(tint_symbol_4);
return;
}
-kernel void compute_main() {
- transpose_31e37e();
+kernel void compute_main(device float2x4* tint_symbol_5 [[buffer(0)]]) {
+ transpose_31e37e(tint_symbol_5);
return;
}
diff --git a/test/tint/builtins/gen/literal/transpose/31e37e.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/transpose/31e37e.wgsl.expected.spvasm
index 3f0ece2..d2bdc93 100644
--- a/test/tint/builtins/gen/literal/transpose/31e37e.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/transpose/31e37e.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 33
+; Bound: 41
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -12,6 +12,9 @@
OpExecutionMode %compute_main LocalSize 1 1 1
OpName %value "value"
OpName %vertex_point_size "vertex_point_size"
+ OpName %prevent_dce_block "prevent_dce_block"
+ OpMemberName %prevent_dce_block 0 "inner"
+ OpName %prevent_dce "prevent_dce"
OpName %transpose_31e37e "transpose_31e37e"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
@@ -20,6 +23,12 @@
OpName %compute_main "compute_main"
OpDecorate %value BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
+ OpDecorate %prevent_dce_block Block
+ OpMemberDecorate %prevent_dce_block 0 Offset 0
+ OpMemberDecorate %prevent_dce_block 0 ColMajor
+ OpMemberDecorate %prevent_dce_block 0 MatrixStride 16
+ OpDecorate %prevent_dce DescriptorSet 2
+ OpDecorate %prevent_dce Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
@@ -28,40 +37,49 @@
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
- %void = OpTypeVoid
- %9 = OpTypeFunction %void
%mat2v4float = OpTypeMatrix %v4float 2
+%prevent_dce_block = OpTypeStruct %mat2v4float
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+ %void = OpTypeVoid
+ %13 = OpTypeFunction %void
%float_1 = OpConstant %float 1
- %15 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
- %16 = OpConstantComposite %mat2v4float %15 %15
+ %18 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
+ %19 = OpConstantComposite %mat2v4float %18 %18
%_ptr_Function_mat2v4float = OpTypePointer Function %mat2v4float
- %19 = OpConstantNull %mat2v4float
- %20 = OpTypeFunction %v4float
-%transpose_31e37e = OpFunction %void None %9
- %12 = OpLabel
- %res = OpVariable %_ptr_Function_mat2v4float Function %19
- OpStore %res %16
+ %22 = OpConstantNull %mat2v4float
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_mat2v4float = OpTypePointer StorageBuffer %mat2v4float
+ %28 = OpTypeFunction %v4float
+%transpose_31e37e = OpFunction %void None %13
+ %16 = OpLabel
+ %res = OpVariable %_ptr_Function_mat2v4float Function %22
+ OpStore %res %19
+ %26 = OpAccessChain %_ptr_StorageBuffer_mat2v4float %prevent_dce %uint_0
+ %27 = OpLoad %mat2v4float %res
+ OpStore %26 %27
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %20
- %22 = OpLabel
- %23 = OpFunctionCall %void %transpose_31e37e
+%vertex_main_inner = OpFunction %v4float None %28
+ %30 = OpLabel
+ %31 = OpFunctionCall %void %transpose_31e37e
OpReturnValue %5
OpFunctionEnd
-%vertex_main = OpFunction %void None %9
- %25 = OpLabel
- %26 = OpFunctionCall %v4float %vertex_main_inner
- OpStore %value %26
+%vertex_main = OpFunction %void None %13
+ %33 = OpLabel
+ %34 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %34
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
-%fragment_main = OpFunction %void None %9
- %28 = OpLabel
- %29 = OpFunctionCall %void %transpose_31e37e
+%fragment_main = OpFunction %void None %13
+ %36 = OpLabel
+ %37 = OpFunctionCall %void %transpose_31e37e
OpReturn
OpFunctionEnd
-%compute_main = OpFunction %void None %9
- %31 = OpLabel
- %32 = OpFunctionCall %void %transpose_31e37e
+%compute_main = OpFunction %void None %13
+ %39 = OpLabel
+ %40 = OpFunctionCall %void %transpose_31e37e
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/transpose/31e37e.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/transpose/31e37e.wgsl.expected.wgsl
index 31e0e1c..06a0782 100644
--- a/test/tint/builtins/gen/literal/transpose/31e37e.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/31e37e.wgsl.expected.wgsl
@@ -1,7 +1,10 @@
fn transpose_31e37e() {
var res : mat2x4<f32> = transpose(mat4x2<f32>(1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat2x4<f32>;
+
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
transpose_31e37e();
diff --git a/test/tint/builtins/gen/literal/transpose/32dd64.wgsl b/test/tint/builtins/gen/literal/transpose/32dd64.wgsl
index 3efd9cc..2fa4efe 100644
--- a/test/tint/builtins/gen/literal/transpose/32dd64.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/32dd64.wgsl
@@ -25,7 +25,6 @@
fn transpose_32dd64() {
var res = transpose(mat3x4(1., 1., 1., 1., 1., 1., 1., 1., 1., 1., 1., 1.));
}
-
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
transpose_32dd64();
diff --git a/test/tint/builtins/gen/literal/transpose/4ce359.wgsl b/test/tint/builtins/gen/literal/transpose/4ce359.wgsl
index 40502b7..396f7bb 100644
--- a/test/tint/builtins/gen/literal/transpose/4ce359.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/4ce359.wgsl
@@ -24,7 +24,9 @@
// fn transpose(mat<2, 4, f32>) -> mat<4, 2, f32>
fn transpose_4ce359() {
var res: mat4x2<f32> = transpose(mat2x4<f32>(1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat4x2<f32>;
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
diff --git a/test/tint/builtins/gen/literal/transpose/4ce359.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/transpose/4ce359.wgsl.expected.dxc.hlsl
index a3c33b4..7bc6aa4 100644
--- a/test/tint/builtins/gen/literal/transpose/4ce359.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/4ce359.wgsl.expected.dxc.hlsl
@@ -1,5 +1,15 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, float4x2 value) {
+ prevent_dce.Store2((offset + 0u), asuint(value[0u]));
+ prevent_dce.Store2((offset + 8u), asuint(value[1u]));
+ prevent_dce.Store2((offset + 16u), asuint(value[2u]));
+ prevent_dce.Store2((offset + 24u), asuint(value[3u]));
+}
+
void transpose_4ce359() {
float4x2 res = float4x2((1.0f).xx, (1.0f).xx, (1.0f).xx, (1.0f).xx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/4ce359.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/transpose/4ce359.wgsl.expected.fxc.hlsl
index a3c33b4..7bc6aa4 100644
--- a/test/tint/builtins/gen/literal/transpose/4ce359.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/4ce359.wgsl.expected.fxc.hlsl
@@ -1,5 +1,15 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, float4x2 value) {
+ prevent_dce.Store2((offset + 0u), asuint(value[0u]));
+ prevent_dce.Store2((offset + 8u), asuint(value[1u]));
+ prevent_dce.Store2((offset + 16u), asuint(value[2u]));
+ prevent_dce.Store2((offset + 24u), asuint(value[3u]));
+}
+
void transpose_4ce359() {
float4x2 res = float4x2((1.0f).xx, (1.0f).xx, (1.0f).xx, (1.0f).xx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/4ce359.wgsl.expected.glsl b/test/tint/builtins/gen/literal/transpose/4ce359.wgsl.expected.glsl
index 101eec8..6c40d55 100644
--- a/test/tint/builtins/gen/literal/transpose/4ce359.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/transpose/4ce359.wgsl.expected.glsl
@@ -1,7 +1,12 @@
#version 310 es
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ mat4x2 inner;
+} prevent_dce;
+
void transpose_4ce359() {
mat4x2 res = mat4x2(vec2(1.0f), vec2(1.0f), vec2(1.0f), vec2(1.0f));
+ prevent_dce.inner = res;
}
vec4 vertex_main() {
@@ -20,8 +25,13 @@
#version 310 es
precision mediump float;
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ mat4x2 inner;
+} prevent_dce;
+
void transpose_4ce359() {
mat4x2 res = mat4x2(vec2(1.0f), vec2(1.0f), vec2(1.0f), vec2(1.0f));
+ prevent_dce.inner = res;
}
void fragment_main() {
@@ -34,8 +44,13 @@
}
#version 310 es
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ mat4x2 inner;
+} prevent_dce;
+
void transpose_4ce359() {
mat4x2 res = mat4x2(vec2(1.0f), vec2(1.0f), vec2(1.0f), vec2(1.0f));
+ prevent_dce.inner = res;
}
void compute_main() {
diff --git a/test/tint/builtins/gen/literal/transpose/4ce359.wgsl.expected.msl b/test/tint/builtins/gen/literal/transpose/4ce359.wgsl.expected.msl
index 150821f..2cf1364 100644
--- a/test/tint/builtins/gen/literal/transpose/4ce359.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/transpose/4ce359.wgsl.expected.msl
@@ -1,33 +1,34 @@
#include <metal_stdlib>
using namespace metal;
-void transpose_4ce359() {
+void transpose_4ce359(device float4x2* const tint_symbol_1) {
float4x2 res = float4x2(float2(1.0f), float2(1.0f), float2(1.0f), float2(1.0f));
+ *(tint_symbol_1) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner() {
- transpose_4ce359();
+float4 vertex_main_inner(device float4x2* const tint_symbol_2) {
+ transpose_4ce359(tint_symbol_2);
return float4(0.0f);
}
-vertex tint_symbol vertex_main() {
- float4 const inner_result = vertex_main_inner();
+vertex tint_symbol vertex_main(device float4x2* tint_symbol_3 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_3);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main() {
- transpose_4ce359();
+fragment void fragment_main(device float4x2* tint_symbol_4 [[buffer(0)]]) {
+ transpose_4ce359(tint_symbol_4);
return;
}
-kernel void compute_main() {
- transpose_4ce359();
+kernel void compute_main(device float4x2* tint_symbol_5 [[buffer(0)]]) {
+ transpose_4ce359(tint_symbol_5);
return;
}
diff --git a/test/tint/builtins/gen/literal/transpose/4ce359.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/transpose/4ce359.wgsl.expected.spvasm
index 96c31a9..63a3fa8 100644
--- a/test/tint/builtins/gen/literal/transpose/4ce359.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/transpose/4ce359.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 34
+; Bound: 42
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -12,6 +12,9 @@
OpExecutionMode %compute_main LocalSize 1 1 1
OpName %value "value"
OpName %vertex_point_size "vertex_point_size"
+ OpName %prevent_dce_block "prevent_dce_block"
+ OpMemberName %prevent_dce_block 0 "inner"
+ OpName %prevent_dce "prevent_dce"
OpName %transpose_4ce359 "transpose_4ce359"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
@@ -20,6 +23,12 @@
OpName %compute_main "compute_main"
OpDecorate %value BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
+ OpDecorate %prevent_dce_block Block
+ OpMemberDecorate %prevent_dce_block 0 Offset 0
+ OpMemberDecorate %prevent_dce_block 0 ColMajor
+ OpMemberDecorate %prevent_dce_block 0 MatrixStride 8
+ OpDecorate %prevent_dce DescriptorSet 2
+ OpDecorate %prevent_dce Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
@@ -28,41 +37,50 @@
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
- %void = OpTypeVoid
- %9 = OpTypeFunction %void
%v2float = OpTypeVector %float 2
%mat4v2float = OpTypeMatrix %v2float 4
+%prevent_dce_block = OpTypeStruct %mat4v2float
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+ %void = OpTypeVoid
+ %14 = OpTypeFunction %void
%float_1 = OpConstant %float 1
- %16 = OpConstantComposite %v2float %float_1 %float_1
- %17 = OpConstantComposite %mat4v2float %16 %16 %16 %16
+ %19 = OpConstantComposite %v2float %float_1 %float_1
+ %20 = OpConstantComposite %mat4v2float %19 %19 %19 %19
%_ptr_Function_mat4v2float = OpTypePointer Function %mat4v2float
- %20 = OpConstantNull %mat4v2float
- %21 = OpTypeFunction %v4float
-%transpose_4ce359 = OpFunction %void None %9
- %12 = OpLabel
- %res = OpVariable %_ptr_Function_mat4v2float Function %20
- OpStore %res %17
+ %23 = OpConstantNull %mat4v2float
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_mat4v2float = OpTypePointer StorageBuffer %mat4v2float
+ %29 = OpTypeFunction %v4float
+%transpose_4ce359 = OpFunction %void None %14
+ %17 = OpLabel
+ %res = OpVariable %_ptr_Function_mat4v2float Function %23
+ OpStore %res %20
+ %27 = OpAccessChain %_ptr_StorageBuffer_mat4v2float %prevent_dce %uint_0
+ %28 = OpLoad %mat4v2float %res
+ OpStore %27 %28
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %21
- %23 = OpLabel
- %24 = OpFunctionCall %void %transpose_4ce359
+%vertex_main_inner = OpFunction %v4float None %29
+ %31 = OpLabel
+ %32 = OpFunctionCall %void %transpose_4ce359
OpReturnValue %5
OpFunctionEnd
-%vertex_main = OpFunction %void None %9
- %26 = OpLabel
- %27 = OpFunctionCall %v4float %vertex_main_inner
- OpStore %value %27
+%vertex_main = OpFunction %void None %14
+ %34 = OpLabel
+ %35 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %35
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
-%fragment_main = OpFunction %void None %9
- %29 = OpLabel
- %30 = OpFunctionCall %void %transpose_4ce359
+%fragment_main = OpFunction %void None %14
+ %37 = OpLabel
+ %38 = OpFunctionCall %void %transpose_4ce359
OpReturn
OpFunctionEnd
-%compute_main = OpFunction %void None %9
- %32 = OpLabel
- %33 = OpFunctionCall %void %transpose_4ce359
+%compute_main = OpFunction %void None %14
+ %40 = OpLabel
+ %41 = OpFunctionCall %void %transpose_4ce359
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/transpose/4ce359.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/transpose/4ce359.wgsl.expected.wgsl
index 47e4a28..32e6a5e 100644
--- a/test/tint/builtins/gen/literal/transpose/4ce359.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/4ce359.wgsl.expected.wgsl
@@ -1,7 +1,10 @@
fn transpose_4ce359() {
var res : mat4x2<f32> = transpose(mat2x4<f32>(1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat4x2<f32>;
+
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
transpose_4ce359();
diff --git a/test/tint/builtins/gen/literal/transpose/4dc9a1.wgsl b/test/tint/builtins/gen/literal/transpose/4dc9a1.wgsl
index d9e4eed..e65e895 100644
--- a/test/tint/builtins/gen/literal/transpose/4dc9a1.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/4dc9a1.wgsl
@@ -24,7 +24,9 @@
// fn transpose(mat<2, 3, f32>) -> mat<3, 2, f32>
fn transpose_4dc9a1() {
var res: mat3x2<f32> = transpose(mat2x3<f32>(1.f, 1.f, 1.f, 1.f, 1.f, 1.f));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat3x2<f32>;
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
diff --git a/test/tint/builtins/gen/literal/transpose/4dc9a1.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/transpose/4dc9a1.wgsl.expected.dxc.hlsl
index 2f3ba72..8e37587 100644
--- a/test/tint/builtins/gen/literal/transpose/4dc9a1.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/4dc9a1.wgsl.expected.dxc.hlsl
@@ -1,5 +1,14 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, float3x2 value) {
+ prevent_dce.Store2((offset + 0u), asuint(value[0u]));
+ prevent_dce.Store2((offset + 8u), asuint(value[1u]));
+ prevent_dce.Store2((offset + 16u), asuint(value[2u]));
+}
+
void transpose_4dc9a1() {
float3x2 res = float3x2((1.0f).xx, (1.0f).xx, (1.0f).xx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/4dc9a1.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/transpose/4dc9a1.wgsl.expected.fxc.hlsl
index 2f3ba72..8e37587 100644
--- a/test/tint/builtins/gen/literal/transpose/4dc9a1.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/4dc9a1.wgsl.expected.fxc.hlsl
@@ -1,5 +1,14 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, float3x2 value) {
+ prevent_dce.Store2((offset + 0u), asuint(value[0u]));
+ prevent_dce.Store2((offset + 8u), asuint(value[1u]));
+ prevent_dce.Store2((offset + 16u), asuint(value[2u]));
+}
+
void transpose_4dc9a1() {
float3x2 res = float3x2((1.0f).xx, (1.0f).xx, (1.0f).xx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/4dc9a1.wgsl.expected.glsl b/test/tint/builtins/gen/literal/transpose/4dc9a1.wgsl.expected.glsl
index 81feae0..70de138 100644
--- a/test/tint/builtins/gen/literal/transpose/4dc9a1.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/transpose/4dc9a1.wgsl.expected.glsl
@@ -1,7 +1,12 @@
#version 310 es
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ mat3x2 inner;
+} prevent_dce;
+
void transpose_4dc9a1() {
mat3x2 res = mat3x2(vec2(1.0f), vec2(1.0f), vec2(1.0f));
+ prevent_dce.inner = res;
}
vec4 vertex_main() {
@@ -20,8 +25,13 @@
#version 310 es
precision mediump float;
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ mat3x2 inner;
+} prevent_dce;
+
void transpose_4dc9a1() {
mat3x2 res = mat3x2(vec2(1.0f), vec2(1.0f), vec2(1.0f));
+ prevent_dce.inner = res;
}
void fragment_main() {
@@ -34,8 +44,13 @@
}
#version 310 es
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ mat3x2 inner;
+} prevent_dce;
+
void transpose_4dc9a1() {
mat3x2 res = mat3x2(vec2(1.0f), vec2(1.0f), vec2(1.0f));
+ prevent_dce.inner = res;
}
void compute_main() {
diff --git a/test/tint/builtins/gen/literal/transpose/4dc9a1.wgsl.expected.msl b/test/tint/builtins/gen/literal/transpose/4dc9a1.wgsl.expected.msl
index b1d4d7a..1aa279e 100644
--- a/test/tint/builtins/gen/literal/transpose/4dc9a1.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/transpose/4dc9a1.wgsl.expected.msl
@@ -1,33 +1,34 @@
#include <metal_stdlib>
using namespace metal;
-void transpose_4dc9a1() {
+void transpose_4dc9a1(device float3x2* const tint_symbol_1) {
float3x2 res = float3x2(float2(1.0f), float2(1.0f), float2(1.0f));
+ *(tint_symbol_1) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner() {
- transpose_4dc9a1();
+float4 vertex_main_inner(device float3x2* const tint_symbol_2) {
+ transpose_4dc9a1(tint_symbol_2);
return float4(0.0f);
}
-vertex tint_symbol vertex_main() {
- float4 const inner_result = vertex_main_inner();
+vertex tint_symbol vertex_main(device float3x2* tint_symbol_3 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_3);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main() {
- transpose_4dc9a1();
+fragment void fragment_main(device float3x2* tint_symbol_4 [[buffer(0)]]) {
+ transpose_4dc9a1(tint_symbol_4);
return;
}
-kernel void compute_main() {
- transpose_4dc9a1();
+kernel void compute_main(device float3x2* tint_symbol_5 [[buffer(0)]]) {
+ transpose_4dc9a1(tint_symbol_5);
return;
}
diff --git a/test/tint/builtins/gen/literal/transpose/4dc9a1.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/transpose/4dc9a1.wgsl.expected.spvasm
index 725a125..77ec602 100644
--- a/test/tint/builtins/gen/literal/transpose/4dc9a1.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/transpose/4dc9a1.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 34
+; Bound: 42
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -12,6 +12,9 @@
OpExecutionMode %compute_main LocalSize 1 1 1
OpName %value "value"
OpName %vertex_point_size "vertex_point_size"
+ OpName %prevent_dce_block "prevent_dce_block"
+ OpMemberName %prevent_dce_block 0 "inner"
+ OpName %prevent_dce "prevent_dce"
OpName %transpose_4dc9a1 "transpose_4dc9a1"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
@@ -20,6 +23,12 @@
OpName %compute_main "compute_main"
OpDecorate %value BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
+ OpDecorate %prevent_dce_block Block
+ OpMemberDecorate %prevent_dce_block 0 Offset 0
+ OpMemberDecorate %prevent_dce_block 0 ColMajor
+ OpMemberDecorate %prevent_dce_block 0 MatrixStride 8
+ OpDecorate %prevent_dce DescriptorSet 2
+ OpDecorate %prevent_dce Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
@@ -28,41 +37,50 @@
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
- %void = OpTypeVoid
- %9 = OpTypeFunction %void
%v2float = OpTypeVector %float 2
%mat3v2float = OpTypeMatrix %v2float 3
+%prevent_dce_block = OpTypeStruct %mat3v2float
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+ %void = OpTypeVoid
+ %14 = OpTypeFunction %void
%float_1 = OpConstant %float 1
- %16 = OpConstantComposite %v2float %float_1 %float_1
- %17 = OpConstantComposite %mat3v2float %16 %16 %16
+ %19 = OpConstantComposite %v2float %float_1 %float_1
+ %20 = OpConstantComposite %mat3v2float %19 %19 %19
%_ptr_Function_mat3v2float = OpTypePointer Function %mat3v2float
- %20 = OpConstantNull %mat3v2float
- %21 = OpTypeFunction %v4float
-%transpose_4dc9a1 = OpFunction %void None %9
- %12 = OpLabel
- %res = OpVariable %_ptr_Function_mat3v2float Function %20
- OpStore %res %17
+ %23 = OpConstantNull %mat3v2float
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_mat3v2float = OpTypePointer StorageBuffer %mat3v2float
+ %29 = OpTypeFunction %v4float
+%transpose_4dc9a1 = OpFunction %void None %14
+ %17 = OpLabel
+ %res = OpVariable %_ptr_Function_mat3v2float Function %23
+ OpStore %res %20
+ %27 = OpAccessChain %_ptr_StorageBuffer_mat3v2float %prevent_dce %uint_0
+ %28 = OpLoad %mat3v2float %res
+ OpStore %27 %28
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %21
- %23 = OpLabel
- %24 = OpFunctionCall %void %transpose_4dc9a1
+%vertex_main_inner = OpFunction %v4float None %29
+ %31 = OpLabel
+ %32 = OpFunctionCall %void %transpose_4dc9a1
OpReturnValue %5
OpFunctionEnd
-%vertex_main = OpFunction %void None %9
- %26 = OpLabel
- %27 = OpFunctionCall %v4float %vertex_main_inner
- OpStore %value %27
+%vertex_main = OpFunction %void None %14
+ %34 = OpLabel
+ %35 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %35
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
-%fragment_main = OpFunction %void None %9
- %29 = OpLabel
- %30 = OpFunctionCall %void %transpose_4dc9a1
+%fragment_main = OpFunction %void None %14
+ %37 = OpLabel
+ %38 = OpFunctionCall %void %transpose_4dc9a1
OpReturn
OpFunctionEnd
-%compute_main = OpFunction %void None %9
- %32 = OpLabel
- %33 = OpFunctionCall %void %transpose_4dc9a1
+%compute_main = OpFunction %void None %14
+ %40 = OpLabel
+ %41 = OpFunctionCall %void %transpose_4dc9a1
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/transpose/4dc9a1.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/transpose/4dc9a1.wgsl.expected.wgsl
index 5db00a5..494d81f 100644
--- a/test/tint/builtins/gen/literal/transpose/4dc9a1.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/4dc9a1.wgsl.expected.wgsl
@@ -1,7 +1,10 @@
fn transpose_4dc9a1() {
var res : mat3x2<f32> = transpose(mat2x3<f32>(1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat3x2<f32>;
+
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
transpose_4dc9a1();
diff --git a/test/tint/builtins/gen/literal/transpose/553e90.wgsl b/test/tint/builtins/gen/literal/transpose/553e90.wgsl
index 5eb0ace..ac448b0 100644
--- a/test/tint/builtins/gen/literal/transpose/553e90.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/553e90.wgsl
@@ -25,7 +25,6 @@
fn transpose_553e90() {
var res = transpose(mat4x2(1., 1., 1., 1., 1., 1., 1., 1.));
}
-
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
transpose_553e90();
diff --git a/test/tint/builtins/gen/literal/transpose/5c133c.wgsl b/test/tint/builtins/gen/literal/transpose/5c133c.wgsl
index 487aad89..e477412 100644
--- a/test/tint/builtins/gen/literal/transpose/5c133c.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/5c133c.wgsl
@@ -25,7 +25,6 @@
fn transpose_5c133c() {
var res = transpose(mat4x3(1., 1., 1., 1., 1., 1., 1., 1., 1., 1., 1., 1.));
}
-
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
transpose_5c133c();
diff --git a/test/tint/builtins/gen/literal/transpose/5edd96.wgsl b/test/tint/builtins/gen/literal/transpose/5edd96.wgsl
index ca694d5..f526477 100644
--- a/test/tint/builtins/gen/literal/transpose/5edd96.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/5edd96.wgsl
@@ -26,7 +26,9 @@
// fn transpose(mat<4, 2, f16>) -> mat<2, 4, f16>
fn transpose_5edd96() {
var res: mat2x4<f16> = transpose(mat4x2<f16>(1.h, 1.h, 1.h, 1.h, 1.h, 1.h, 1.h, 1.h));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat2x4<f16>;
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
diff --git a/test/tint/builtins/gen/literal/transpose/5edd96.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/transpose/5edd96.wgsl.expected.dxc.hlsl
index 736166a..5ebd512 100644
--- a/test/tint/builtins/gen/literal/transpose/5edd96.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/5edd96.wgsl.expected.dxc.hlsl
@@ -1,5 +1,13 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, matrix<float16_t, 2, 4> value) {
+ prevent_dce.Store<vector<float16_t, 4> >((offset + 0u), value[0u]);
+ prevent_dce.Store<vector<float16_t, 4> >((offset + 8u), value[1u]);
+}
+
void transpose_5edd96() {
matrix<float16_t, 2, 4> res = matrix<float16_t, 2, 4>((float16_t(1.0h)).xxxx, (float16_t(1.0h)).xxxx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/5edd96.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/transpose/5edd96.wgsl.expected.fxc.hlsl
index 1011d3c..8b775e9 100644
--- a/test/tint/builtins/gen/literal/transpose/5edd96.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/5edd96.wgsl.expected.fxc.hlsl
@@ -1,7 +1,15 @@
SKIP: FAILED
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, matrix<float16_t, 2, 4> value) {
+ prevent_dce.Store<vector<float16_t, 4> >((offset + 0u), value[0u]);
+ prevent_dce.Store<vector<float16_t, 4> >((offset + 8u), value[1u]);
+}
+
void transpose_5edd96() {
matrix<float16_t, 2, 4> res = matrix<float16_t, 2, 4>((float16_t(1.0h)).xxxx, (float16_t(1.0h)).xxxx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/5edd96.wgsl.expected.glsl b/test/tint/builtins/gen/literal/transpose/5edd96.wgsl.expected.glsl
index f2c3448..2968eda 100644
--- a/test/tint/builtins/gen/literal/transpose/5edd96.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/transpose/5edd96.wgsl.expected.glsl
@@ -1,8 +1,13 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16mat2x4 inner;
+} prevent_dce;
+
void transpose_5edd96() {
f16mat2x4 res = f16mat2x4(f16vec4(1.0hf), f16vec4(1.0hf));
+ prevent_dce.inner = res;
}
vec4 vertex_main() {
@@ -22,8 +27,13 @@
#extension GL_AMD_gpu_shader_half_float : require
precision mediump float;
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16mat2x4 inner;
+} prevent_dce;
+
void transpose_5edd96() {
f16mat2x4 res = f16mat2x4(f16vec4(1.0hf), f16vec4(1.0hf));
+ prevent_dce.inner = res;
}
void fragment_main() {
@@ -37,8 +47,13 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16mat2x4 inner;
+} prevent_dce;
+
void transpose_5edd96() {
f16mat2x4 res = f16mat2x4(f16vec4(1.0hf), f16vec4(1.0hf));
+ prevent_dce.inner = res;
}
void compute_main() {
diff --git a/test/tint/builtins/gen/literal/transpose/5edd96.wgsl.expected.msl b/test/tint/builtins/gen/literal/transpose/5edd96.wgsl.expected.msl
index 9e612d6..75e2d46 100644
--- a/test/tint/builtins/gen/literal/transpose/5edd96.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/transpose/5edd96.wgsl.expected.msl
@@ -1,33 +1,34 @@
#include <metal_stdlib>
using namespace metal;
-void transpose_5edd96() {
+void transpose_5edd96(device half2x4* const tint_symbol_1) {
half2x4 res = half2x4(half4(1.0h), half4(1.0h));
+ *(tint_symbol_1) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner() {
- transpose_5edd96();
+float4 vertex_main_inner(device half2x4* const tint_symbol_2) {
+ transpose_5edd96(tint_symbol_2);
return float4(0.0f);
}
-vertex tint_symbol vertex_main() {
- float4 const inner_result = vertex_main_inner();
+vertex tint_symbol vertex_main(device half2x4* tint_symbol_3 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_3);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main() {
- transpose_5edd96();
+fragment void fragment_main(device half2x4* tint_symbol_4 [[buffer(0)]]) {
+ transpose_5edd96(tint_symbol_4);
return;
}
-kernel void compute_main() {
- transpose_5edd96();
+kernel void compute_main(device half2x4* tint_symbol_5 [[buffer(0)]]) {
+ transpose_5edd96(tint_symbol_5);
return;
}
diff --git a/test/tint/builtins/gen/literal/transpose/5edd96.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/transpose/5edd96.wgsl.expected.spvasm
index 603fe36..11b39f6 100644
--- a/test/tint/builtins/gen/literal/transpose/5edd96.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/transpose/5edd96.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 36
+; Bound: 44
; Schema: 0
OpCapability Shader
OpCapability Float16
@@ -16,6 +16,9 @@
OpExecutionMode %compute_main LocalSize 1 1 1
OpName %value "value"
OpName %vertex_point_size "vertex_point_size"
+ OpName %prevent_dce_block "prevent_dce_block"
+ OpMemberName %prevent_dce_block 0 "inner"
+ OpName %prevent_dce "prevent_dce"
OpName %transpose_5edd96 "transpose_5edd96"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
@@ -24,6 +27,12 @@
OpName %compute_main "compute_main"
OpDecorate %value BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
+ OpDecorate %prevent_dce_block Block
+ OpMemberDecorate %prevent_dce_block 0 Offset 0
+ OpMemberDecorate %prevent_dce_block 0 ColMajor
+ OpMemberDecorate %prevent_dce_block 0 MatrixStride 8
+ OpDecorate %prevent_dce DescriptorSet 2
+ OpDecorate %prevent_dce Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
@@ -32,43 +41,52 @@
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
- %void = OpTypeVoid
- %9 = OpTypeFunction %void
%half = OpTypeFloat 16
%v4half = OpTypeVector %half 4
%mat2v4half = OpTypeMatrix %v4half 2
+%prevent_dce_block = OpTypeStruct %mat2v4half
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+ %void = OpTypeVoid
+ %15 = OpTypeFunction %void
%half_0x1p_0 = OpConstant %half 0x1p+0
- %17 = OpConstantComposite %v4half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
- %18 = OpConstantComposite %mat2v4half %17 %17
+ %20 = OpConstantComposite %v4half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
+ %21 = OpConstantComposite %mat2v4half %20 %20
%_ptr_Function_mat2v4half = OpTypePointer Function %mat2v4half
- %21 = OpConstantNull %mat2v4half
- %22 = OpTypeFunction %v4float
+ %24 = OpConstantNull %mat2v4half
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_mat2v4half = OpTypePointer StorageBuffer %mat2v4half
+ %30 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
-%transpose_5edd96 = OpFunction %void None %9
- %12 = OpLabel
- %res = OpVariable %_ptr_Function_mat2v4half Function %21
- OpStore %res %18
+%transpose_5edd96 = OpFunction %void None %15
+ %18 = OpLabel
+ %res = OpVariable %_ptr_Function_mat2v4half Function %24
+ OpStore %res %21
+ %28 = OpAccessChain %_ptr_StorageBuffer_mat2v4half %prevent_dce %uint_0
+ %29 = OpLoad %mat2v4half %res
+ OpStore %28 %29
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %22
- %24 = OpLabel
- %25 = OpFunctionCall %void %transpose_5edd96
+%vertex_main_inner = OpFunction %v4float None %30
+ %32 = OpLabel
+ %33 = OpFunctionCall %void %transpose_5edd96
OpReturnValue %5
OpFunctionEnd
-%vertex_main = OpFunction %void None %9
- %27 = OpLabel
- %28 = OpFunctionCall %v4float %vertex_main_inner
- OpStore %value %28
+%vertex_main = OpFunction %void None %15
+ %35 = OpLabel
+ %36 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %36
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
-%fragment_main = OpFunction %void None %9
- %31 = OpLabel
- %32 = OpFunctionCall %void %transpose_5edd96
+%fragment_main = OpFunction %void None %15
+ %39 = OpLabel
+ %40 = OpFunctionCall %void %transpose_5edd96
OpReturn
OpFunctionEnd
-%compute_main = OpFunction %void None %9
- %34 = OpLabel
- %35 = OpFunctionCall %void %transpose_5edd96
+%compute_main = OpFunction %void None %15
+ %42 = OpLabel
+ %43 = OpFunctionCall %void %transpose_5edd96
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/transpose/5edd96.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/transpose/5edd96.wgsl.expected.wgsl
index 3edbfa2..5ff72d3 100644
--- a/test/tint/builtins/gen/literal/transpose/5edd96.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/5edd96.wgsl.expected.wgsl
@@ -2,8 +2,11 @@
fn transpose_5edd96() {
var res : mat2x4<f16> = transpose(mat4x2<f16>(1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat2x4<f16>;
+
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
transpose_5edd96();
diff --git a/test/tint/builtins/gen/literal/transpose/5f36bf.wgsl b/test/tint/builtins/gen/literal/transpose/5f36bf.wgsl
index f5992a4..109d13c 100644
--- a/test/tint/builtins/gen/literal/transpose/5f36bf.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/5f36bf.wgsl
@@ -26,7 +26,9 @@
// fn transpose(mat<4, 3, f16>) -> mat<3, 4, f16>
fn transpose_5f36bf() {
var res: mat3x4<f16> = transpose(mat4x3<f16>(1.h, 1.h, 1.h, 1.h, 1.h, 1.h, 1.h, 1.h, 1.h, 1.h, 1.h, 1.h));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat3x4<f16>;
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
diff --git a/test/tint/builtins/gen/literal/transpose/5f36bf.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/transpose/5f36bf.wgsl.expected.dxc.hlsl
index 6bef090..8fa5fe7 100644
--- a/test/tint/builtins/gen/literal/transpose/5f36bf.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/5f36bf.wgsl.expected.dxc.hlsl
@@ -1,5 +1,14 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, matrix<float16_t, 3, 4> value) {
+ prevent_dce.Store<vector<float16_t, 4> >((offset + 0u), value[0u]);
+ prevent_dce.Store<vector<float16_t, 4> >((offset + 8u), value[1u]);
+ prevent_dce.Store<vector<float16_t, 4> >((offset + 16u), value[2u]);
+}
+
void transpose_5f36bf() {
matrix<float16_t, 3, 4> res = matrix<float16_t, 3, 4>((float16_t(1.0h)).xxxx, (float16_t(1.0h)).xxxx, (float16_t(1.0h)).xxxx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/5f36bf.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/transpose/5f36bf.wgsl.expected.fxc.hlsl
index 6336c52..6d2e6f4 100644
--- a/test/tint/builtins/gen/literal/transpose/5f36bf.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/5f36bf.wgsl.expected.fxc.hlsl
@@ -1,7 +1,16 @@
SKIP: FAILED
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, matrix<float16_t, 3, 4> value) {
+ prevent_dce.Store<vector<float16_t, 4> >((offset + 0u), value[0u]);
+ prevent_dce.Store<vector<float16_t, 4> >((offset + 8u), value[1u]);
+ prevent_dce.Store<vector<float16_t, 4> >((offset + 16u), value[2u]);
+}
+
void transpose_5f36bf() {
matrix<float16_t, 3, 4> res = matrix<float16_t, 3, 4>((float16_t(1.0h)).xxxx, (float16_t(1.0h)).xxxx, (float16_t(1.0h)).xxxx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/5f36bf.wgsl.expected.glsl b/test/tint/builtins/gen/literal/transpose/5f36bf.wgsl.expected.glsl
index 6e78f07..7bf364b 100644
--- a/test/tint/builtins/gen/literal/transpose/5f36bf.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/transpose/5f36bf.wgsl.expected.glsl
@@ -1,8 +1,13 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16mat3x4 inner;
+} prevent_dce;
+
void transpose_5f36bf() {
f16mat3x4 res = f16mat3x4(f16vec4(1.0hf), f16vec4(1.0hf), f16vec4(1.0hf));
+ prevent_dce.inner = res;
}
vec4 vertex_main() {
@@ -22,8 +27,13 @@
#extension GL_AMD_gpu_shader_half_float : require
precision mediump float;
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16mat3x4 inner;
+} prevent_dce;
+
void transpose_5f36bf() {
f16mat3x4 res = f16mat3x4(f16vec4(1.0hf), f16vec4(1.0hf), f16vec4(1.0hf));
+ prevent_dce.inner = res;
}
void fragment_main() {
@@ -37,8 +47,13 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16mat3x4 inner;
+} prevent_dce;
+
void transpose_5f36bf() {
f16mat3x4 res = f16mat3x4(f16vec4(1.0hf), f16vec4(1.0hf), f16vec4(1.0hf));
+ prevent_dce.inner = res;
}
void compute_main() {
diff --git a/test/tint/builtins/gen/literal/transpose/5f36bf.wgsl.expected.msl b/test/tint/builtins/gen/literal/transpose/5f36bf.wgsl.expected.msl
index 7e141f9..8d19bfe 100644
--- a/test/tint/builtins/gen/literal/transpose/5f36bf.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/transpose/5f36bf.wgsl.expected.msl
@@ -1,33 +1,34 @@
#include <metal_stdlib>
using namespace metal;
-void transpose_5f36bf() {
+void transpose_5f36bf(device half3x4* const tint_symbol_1) {
half3x4 res = half3x4(half4(1.0h), half4(1.0h), half4(1.0h));
+ *(tint_symbol_1) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner() {
- transpose_5f36bf();
+float4 vertex_main_inner(device half3x4* const tint_symbol_2) {
+ transpose_5f36bf(tint_symbol_2);
return float4(0.0f);
}
-vertex tint_symbol vertex_main() {
- float4 const inner_result = vertex_main_inner();
+vertex tint_symbol vertex_main(device half3x4* tint_symbol_3 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_3);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main() {
- transpose_5f36bf();
+fragment void fragment_main(device half3x4* tint_symbol_4 [[buffer(0)]]) {
+ transpose_5f36bf(tint_symbol_4);
return;
}
-kernel void compute_main() {
- transpose_5f36bf();
+kernel void compute_main(device half3x4* tint_symbol_5 [[buffer(0)]]) {
+ transpose_5f36bf(tint_symbol_5);
return;
}
diff --git a/test/tint/builtins/gen/literal/transpose/5f36bf.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/transpose/5f36bf.wgsl.expected.spvasm
index 08bc2fd..5e7050c 100644
--- a/test/tint/builtins/gen/literal/transpose/5f36bf.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/transpose/5f36bf.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 36
+; Bound: 44
; Schema: 0
OpCapability Shader
OpCapability Float16
@@ -16,6 +16,9 @@
OpExecutionMode %compute_main LocalSize 1 1 1
OpName %value "value"
OpName %vertex_point_size "vertex_point_size"
+ OpName %prevent_dce_block "prevent_dce_block"
+ OpMemberName %prevent_dce_block 0 "inner"
+ OpName %prevent_dce "prevent_dce"
OpName %transpose_5f36bf "transpose_5f36bf"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
@@ -24,6 +27,12 @@
OpName %compute_main "compute_main"
OpDecorate %value BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
+ OpDecorate %prevent_dce_block Block
+ OpMemberDecorate %prevent_dce_block 0 Offset 0
+ OpMemberDecorate %prevent_dce_block 0 ColMajor
+ OpMemberDecorate %prevent_dce_block 0 MatrixStride 8
+ OpDecorate %prevent_dce DescriptorSet 2
+ OpDecorate %prevent_dce Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
@@ -32,43 +41,52 @@
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
- %void = OpTypeVoid
- %9 = OpTypeFunction %void
%half = OpTypeFloat 16
%v4half = OpTypeVector %half 4
%mat3v4half = OpTypeMatrix %v4half 3
+%prevent_dce_block = OpTypeStruct %mat3v4half
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+ %void = OpTypeVoid
+ %15 = OpTypeFunction %void
%half_0x1p_0 = OpConstant %half 0x1p+0
- %17 = OpConstantComposite %v4half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
- %18 = OpConstantComposite %mat3v4half %17 %17 %17
+ %20 = OpConstantComposite %v4half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
+ %21 = OpConstantComposite %mat3v4half %20 %20 %20
%_ptr_Function_mat3v4half = OpTypePointer Function %mat3v4half
- %21 = OpConstantNull %mat3v4half
- %22 = OpTypeFunction %v4float
+ %24 = OpConstantNull %mat3v4half
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_mat3v4half = OpTypePointer StorageBuffer %mat3v4half
+ %30 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
-%transpose_5f36bf = OpFunction %void None %9
- %12 = OpLabel
- %res = OpVariable %_ptr_Function_mat3v4half Function %21
- OpStore %res %18
+%transpose_5f36bf = OpFunction %void None %15
+ %18 = OpLabel
+ %res = OpVariable %_ptr_Function_mat3v4half Function %24
+ OpStore %res %21
+ %28 = OpAccessChain %_ptr_StorageBuffer_mat3v4half %prevent_dce %uint_0
+ %29 = OpLoad %mat3v4half %res
+ OpStore %28 %29
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %22
- %24 = OpLabel
- %25 = OpFunctionCall %void %transpose_5f36bf
+%vertex_main_inner = OpFunction %v4float None %30
+ %32 = OpLabel
+ %33 = OpFunctionCall %void %transpose_5f36bf
OpReturnValue %5
OpFunctionEnd
-%vertex_main = OpFunction %void None %9
- %27 = OpLabel
- %28 = OpFunctionCall %v4float %vertex_main_inner
- OpStore %value %28
+%vertex_main = OpFunction %void None %15
+ %35 = OpLabel
+ %36 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %36
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
-%fragment_main = OpFunction %void None %9
- %31 = OpLabel
- %32 = OpFunctionCall %void %transpose_5f36bf
+%fragment_main = OpFunction %void None %15
+ %39 = OpLabel
+ %40 = OpFunctionCall %void %transpose_5f36bf
OpReturn
OpFunctionEnd
-%compute_main = OpFunction %void None %9
- %34 = OpLabel
- %35 = OpFunctionCall %void %transpose_5f36bf
+%compute_main = OpFunction %void None %15
+ %42 = OpLabel
+ %43 = OpFunctionCall %void %transpose_5f36bf
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/transpose/5f36bf.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/transpose/5f36bf.wgsl.expected.wgsl
index 374302e..eba7377 100644
--- a/test/tint/builtins/gen/literal/transpose/5f36bf.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/5f36bf.wgsl.expected.wgsl
@@ -2,8 +2,11 @@
fn transpose_5f36bf() {
var res : mat3x4<f16> = transpose(mat4x3<f16>(1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat3x4<f16>;
+
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
transpose_5f36bf();
diff --git a/test/tint/builtins/gen/literal/transpose/66fce8.wgsl b/test/tint/builtins/gen/literal/transpose/66fce8.wgsl
index 51e7da6..7eed3a2 100644
--- a/test/tint/builtins/gen/literal/transpose/66fce8.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/66fce8.wgsl
@@ -25,7 +25,6 @@
fn transpose_66fce8() {
var res = transpose(mat3x3(1., 1., 1., 1., 1., 1., 1., 1., 1.));
}
-
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
transpose_66fce8();
diff --git a/test/tint/builtins/gen/literal/transpose/70ca11.wgsl b/test/tint/builtins/gen/literal/transpose/70ca11.wgsl
index 7ef92e5..430acf3 100644
--- a/test/tint/builtins/gen/literal/transpose/70ca11.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/70ca11.wgsl
@@ -25,7 +25,6 @@
fn transpose_70ca11() {
var res = transpose(mat2x3(1., 1., 1., 1., 1., 1.));
}
-
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
transpose_70ca11();
diff --git a/test/tint/builtins/gen/literal/transpose/7be8b2.wgsl b/test/tint/builtins/gen/literal/transpose/7be8b2.wgsl
index bd62bb3..eee4741 100644
--- a/test/tint/builtins/gen/literal/transpose/7be8b2.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/7be8b2.wgsl
@@ -26,7 +26,9 @@
// fn transpose(mat<2, 2, f16>) -> mat<2, 2, f16>
fn transpose_7be8b2() {
var res: mat2x2<f16> = transpose(mat2x2<f16>(1.h, 1.h, 1.h, 1.h));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat2x2<f16>;
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
diff --git a/test/tint/builtins/gen/literal/transpose/7be8b2.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/transpose/7be8b2.wgsl.expected.dxc.hlsl
index 08cac08..fa38c79 100644
--- a/test/tint/builtins/gen/literal/transpose/7be8b2.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/7be8b2.wgsl.expected.dxc.hlsl
@@ -1,5 +1,13 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, matrix<float16_t, 2, 2> value) {
+ prevent_dce.Store<vector<float16_t, 2> >((offset + 0u), value[0u]);
+ prevent_dce.Store<vector<float16_t, 2> >((offset + 4u), value[1u]);
+}
+
void transpose_7be8b2() {
matrix<float16_t, 2, 2> res = matrix<float16_t, 2, 2>((float16_t(1.0h)).xx, (float16_t(1.0h)).xx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/7be8b2.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/transpose/7be8b2.wgsl.expected.fxc.hlsl
index ed432d6..f9f1c02 100644
--- a/test/tint/builtins/gen/literal/transpose/7be8b2.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/7be8b2.wgsl.expected.fxc.hlsl
@@ -1,7 +1,15 @@
SKIP: FAILED
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, matrix<float16_t, 2, 2> value) {
+ prevent_dce.Store<vector<float16_t, 2> >((offset + 0u), value[0u]);
+ prevent_dce.Store<vector<float16_t, 2> >((offset + 4u), value[1u]);
+}
+
void transpose_7be8b2() {
matrix<float16_t, 2, 2> res = matrix<float16_t, 2, 2>((float16_t(1.0h)).xx, (float16_t(1.0h)).xx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/7be8b2.wgsl.expected.glsl b/test/tint/builtins/gen/literal/transpose/7be8b2.wgsl.expected.glsl
index db2a179..76b1843 100644
--- a/test/tint/builtins/gen/literal/transpose/7be8b2.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/transpose/7be8b2.wgsl.expected.glsl
@@ -1,8 +1,13 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16mat2 inner;
+} prevent_dce;
+
void transpose_7be8b2() {
f16mat2 res = f16mat2(f16vec2(1.0hf), f16vec2(1.0hf));
+ prevent_dce.inner = res;
}
vec4 vertex_main() {
@@ -22,8 +27,13 @@
#extension GL_AMD_gpu_shader_half_float : require
precision mediump float;
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16mat2 inner;
+} prevent_dce;
+
void transpose_7be8b2() {
f16mat2 res = f16mat2(f16vec2(1.0hf), f16vec2(1.0hf));
+ prevent_dce.inner = res;
}
void fragment_main() {
@@ -37,8 +47,13 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16mat2 inner;
+} prevent_dce;
+
void transpose_7be8b2() {
f16mat2 res = f16mat2(f16vec2(1.0hf), f16vec2(1.0hf));
+ prevent_dce.inner = res;
}
void compute_main() {
diff --git a/test/tint/builtins/gen/literal/transpose/7be8b2.wgsl.expected.msl b/test/tint/builtins/gen/literal/transpose/7be8b2.wgsl.expected.msl
index 90e5d8a..aa428f6 100644
--- a/test/tint/builtins/gen/literal/transpose/7be8b2.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/transpose/7be8b2.wgsl.expected.msl
@@ -1,33 +1,34 @@
#include <metal_stdlib>
using namespace metal;
-void transpose_7be8b2() {
+void transpose_7be8b2(device half2x2* const tint_symbol_1) {
half2x2 res = half2x2(half2(1.0h), half2(1.0h));
+ *(tint_symbol_1) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner() {
- transpose_7be8b2();
+float4 vertex_main_inner(device half2x2* const tint_symbol_2) {
+ transpose_7be8b2(tint_symbol_2);
return float4(0.0f);
}
-vertex tint_symbol vertex_main() {
- float4 const inner_result = vertex_main_inner();
+vertex tint_symbol vertex_main(device half2x2* tint_symbol_3 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_3);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main() {
- transpose_7be8b2();
+fragment void fragment_main(device half2x2* tint_symbol_4 [[buffer(0)]]) {
+ transpose_7be8b2(tint_symbol_4);
return;
}
-kernel void compute_main() {
- transpose_7be8b2();
+kernel void compute_main(device half2x2* tint_symbol_5 [[buffer(0)]]) {
+ transpose_7be8b2(tint_symbol_5);
return;
}
diff --git a/test/tint/builtins/gen/literal/transpose/7be8b2.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/transpose/7be8b2.wgsl.expected.spvasm
index 05c185c..fd02dd4 100644
--- a/test/tint/builtins/gen/literal/transpose/7be8b2.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/transpose/7be8b2.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 36
+; Bound: 44
; Schema: 0
OpCapability Shader
OpCapability Float16
@@ -16,6 +16,9 @@
OpExecutionMode %compute_main LocalSize 1 1 1
OpName %value "value"
OpName %vertex_point_size "vertex_point_size"
+ OpName %prevent_dce_block "prevent_dce_block"
+ OpMemberName %prevent_dce_block 0 "inner"
+ OpName %prevent_dce "prevent_dce"
OpName %transpose_7be8b2 "transpose_7be8b2"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
@@ -24,6 +27,12 @@
OpName %compute_main "compute_main"
OpDecorate %value BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
+ OpDecorate %prevent_dce_block Block
+ OpMemberDecorate %prevent_dce_block 0 Offset 0
+ OpMemberDecorate %prevent_dce_block 0 ColMajor
+ OpMemberDecorate %prevent_dce_block 0 MatrixStride 4
+ OpDecorate %prevent_dce DescriptorSet 2
+ OpDecorate %prevent_dce Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
@@ -32,43 +41,52 @@
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
- %void = OpTypeVoid
- %9 = OpTypeFunction %void
%half = OpTypeFloat 16
%v2half = OpTypeVector %half 2
%mat2v2half = OpTypeMatrix %v2half 2
+%prevent_dce_block = OpTypeStruct %mat2v2half
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+ %void = OpTypeVoid
+ %15 = OpTypeFunction %void
%half_0x1p_0 = OpConstant %half 0x1p+0
- %17 = OpConstantComposite %v2half %half_0x1p_0 %half_0x1p_0
- %18 = OpConstantComposite %mat2v2half %17 %17
+ %20 = OpConstantComposite %v2half %half_0x1p_0 %half_0x1p_0
+ %21 = OpConstantComposite %mat2v2half %20 %20
%_ptr_Function_mat2v2half = OpTypePointer Function %mat2v2half
- %21 = OpConstantNull %mat2v2half
- %22 = OpTypeFunction %v4float
+ %24 = OpConstantNull %mat2v2half
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_mat2v2half = OpTypePointer StorageBuffer %mat2v2half
+ %30 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
-%transpose_7be8b2 = OpFunction %void None %9
- %12 = OpLabel
- %res = OpVariable %_ptr_Function_mat2v2half Function %21
- OpStore %res %18
+%transpose_7be8b2 = OpFunction %void None %15
+ %18 = OpLabel
+ %res = OpVariable %_ptr_Function_mat2v2half Function %24
+ OpStore %res %21
+ %28 = OpAccessChain %_ptr_StorageBuffer_mat2v2half %prevent_dce %uint_0
+ %29 = OpLoad %mat2v2half %res
+ OpStore %28 %29
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %22
- %24 = OpLabel
- %25 = OpFunctionCall %void %transpose_7be8b2
+%vertex_main_inner = OpFunction %v4float None %30
+ %32 = OpLabel
+ %33 = OpFunctionCall %void %transpose_7be8b2
OpReturnValue %5
OpFunctionEnd
-%vertex_main = OpFunction %void None %9
- %27 = OpLabel
- %28 = OpFunctionCall %v4float %vertex_main_inner
- OpStore %value %28
+%vertex_main = OpFunction %void None %15
+ %35 = OpLabel
+ %36 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %36
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
-%fragment_main = OpFunction %void None %9
- %31 = OpLabel
- %32 = OpFunctionCall %void %transpose_7be8b2
+%fragment_main = OpFunction %void None %15
+ %39 = OpLabel
+ %40 = OpFunctionCall %void %transpose_7be8b2
OpReturn
OpFunctionEnd
-%compute_main = OpFunction %void None %9
- %34 = OpLabel
- %35 = OpFunctionCall %void %transpose_7be8b2
+%compute_main = OpFunction %void None %15
+ %42 = OpLabel
+ %43 = OpFunctionCall %void %transpose_7be8b2
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/transpose/7be8b2.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/transpose/7be8b2.wgsl.expected.wgsl
index 9b4281f..2985afa 100644
--- a/test/tint/builtins/gen/literal/transpose/7be8b2.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/7be8b2.wgsl.expected.wgsl
@@ -2,8 +2,11 @@
fn transpose_7be8b2() {
var res : mat2x2<f16> = transpose(mat2x2<f16>(1.0h, 1.0h, 1.0h, 1.0h));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat2x2<f16>;
+
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
transpose_7be8b2();
diff --git a/test/tint/builtins/gen/literal/transpose/7eb2c5.wgsl b/test/tint/builtins/gen/literal/transpose/7eb2c5.wgsl
index 84f0709..e55fc8a 100644
--- a/test/tint/builtins/gen/literal/transpose/7eb2c5.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/7eb2c5.wgsl
@@ -25,7 +25,6 @@
fn transpose_7eb2c5() {
var res = transpose(mat2x2(1., 1., 1., 1.));
}
-
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
transpose_7eb2c5();
diff --git a/test/tint/builtins/gen/literal/transpose/844869.wgsl b/test/tint/builtins/gen/literal/transpose/844869.wgsl
index 0548c1b..60f4010 100644
--- a/test/tint/builtins/gen/literal/transpose/844869.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/844869.wgsl
@@ -26,7 +26,9 @@
// fn transpose(mat<4, 4, f16>) -> mat<4, 4, f16>
fn transpose_844869() {
var res: mat4x4<f16> = transpose(mat4x4<f16>(1.h, 1.h, 1.h, 1.h, 1.h, 1.h, 1.h, 1.h, 1.h, 1.h, 1.h, 1.h, 1.h, 1.h, 1.h, 1.h));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat4x4<f16>;
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
diff --git a/test/tint/builtins/gen/literal/transpose/844869.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/transpose/844869.wgsl.expected.dxc.hlsl
index ec32137..f38fa90 100644
--- a/test/tint/builtins/gen/literal/transpose/844869.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/844869.wgsl.expected.dxc.hlsl
@@ -1,5 +1,15 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, matrix<float16_t, 4, 4> value) {
+ prevent_dce.Store<vector<float16_t, 4> >((offset + 0u), value[0u]);
+ prevent_dce.Store<vector<float16_t, 4> >((offset + 8u), value[1u]);
+ prevent_dce.Store<vector<float16_t, 4> >((offset + 16u), value[2u]);
+ prevent_dce.Store<vector<float16_t, 4> >((offset + 24u), value[3u]);
+}
+
void transpose_844869() {
matrix<float16_t, 4, 4> res = matrix<float16_t, 4, 4>((float16_t(1.0h)).xxxx, (float16_t(1.0h)).xxxx, (float16_t(1.0h)).xxxx, (float16_t(1.0h)).xxxx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/844869.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/transpose/844869.wgsl.expected.fxc.hlsl
index 2a2a2b4..ec0bb48 100644
--- a/test/tint/builtins/gen/literal/transpose/844869.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/844869.wgsl.expected.fxc.hlsl
@@ -1,7 +1,17 @@
SKIP: FAILED
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, matrix<float16_t, 4, 4> value) {
+ prevent_dce.Store<vector<float16_t, 4> >((offset + 0u), value[0u]);
+ prevent_dce.Store<vector<float16_t, 4> >((offset + 8u), value[1u]);
+ prevent_dce.Store<vector<float16_t, 4> >((offset + 16u), value[2u]);
+ prevent_dce.Store<vector<float16_t, 4> >((offset + 24u), value[3u]);
+}
+
void transpose_844869() {
matrix<float16_t, 4, 4> res = matrix<float16_t, 4, 4>((float16_t(1.0h)).xxxx, (float16_t(1.0h)).xxxx, (float16_t(1.0h)).xxxx, (float16_t(1.0h)).xxxx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/844869.wgsl.expected.glsl b/test/tint/builtins/gen/literal/transpose/844869.wgsl.expected.glsl
index 6fd102a..f077991 100644
--- a/test/tint/builtins/gen/literal/transpose/844869.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/transpose/844869.wgsl.expected.glsl
@@ -1,8 +1,13 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16mat4 inner;
+} prevent_dce;
+
void transpose_844869() {
f16mat4 res = f16mat4(f16vec4(1.0hf), f16vec4(1.0hf), f16vec4(1.0hf), f16vec4(1.0hf));
+ prevent_dce.inner = res;
}
vec4 vertex_main() {
@@ -22,8 +27,13 @@
#extension GL_AMD_gpu_shader_half_float : require
precision mediump float;
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16mat4 inner;
+} prevent_dce;
+
void transpose_844869() {
f16mat4 res = f16mat4(f16vec4(1.0hf), f16vec4(1.0hf), f16vec4(1.0hf), f16vec4(1.0hf));
+ prevent_dce.inner = res;
}
void fragment_main() {
@@ -37,8 +47,13 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16mat4 inner;
+} prevent_dce;
+
void transpose_844869() {
f16mat4 res = f16mat4(f16vec4(1.0hf), f16vec4(1.0hf), f16vec4(1.0hf), f16vec4(1.0hf));
+ prevent_dce.inner = res;
}
void compute_main() {
diff --git a/test/tint/builtins/gen/literal/transpose/844869.wgsl.expected.msl b/test/tint/builtins/gen/literal/transpose/844869.wgsl.expected.msl
index fe21910..9d757f4 100644
--- a/test/tint/builtins/gen/literal/transpose/844869.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/transpose/844869.wgsl.expected.msl
@@ -1,33 +1,34 @@
#include <metal_stdlib>
using namespace metal;
-void transpose_844869() {
+void transpose_844869(device half4x4* const tint_symbol_1) {
half4x4 res = half4x4(half4(1.0h), half4(1.0h), half4(1.0h), half4(1.0h));
+ *(tint_symbol_1) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner() {
- transpose_844869();
+float4 vertex_main_inner(device half4x4* const tint_symbol_2) {
+ transpose_844869(tint_symbol_2);
return float4(0.0f);
}
-vertex tint_symbol vertex_main() {
- float4 const inner_result = vertex_main_inner();
+vertex tint_symbol vertex_main(device half4x4* tint_symbol_3 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_3);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main() {
- transpose_844869();
+fragment void fragment_main(device half4x4* tint_symbol_4 [[buffer(0)]]) {
+ transpose_844869(tint_symbol_4);
return;
}
-kernel void compute_main() {
- transpose_844869();
+kernel void compute_main(device half4x4* tint_symbol_5 [[buffer(0)]]) {
+ transpose_844869(tint_symbol_5);
return;
}
diff --git a/test/tint/builtins/gen/literal/transpose/844869.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/transpose/844869.wgsl.expected.spvasm
index 3364395..1fe7f8e 100644
--- a/test/tint/builtins/gen/literal/transpose/844869.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/transpose/844869.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 36
+; Bound: 44
; Schema: 0
OpCapability Shader
OpCapability Float16
@@ -16,6 +16,9 @@
OpExecutionMode %compute_main LocalSize 1 1 1
OpName %value "value"
OpName %vertex_point_size "vertex_point_size"
+ OpName %prevent_dce_block "prevent_dce_block"
+ OpMemberName %prevent_dce_block 0 "inner"
+ OpName %prevent_dce "prevent_dce"
OpName %transpose_844869 "transpose_844869"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
@@ -24,6 +27,12 @@
OpName %compute_main "compute_main"
OpDecorate %value BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
+ OpDecorate %prevent_dce_block Block
+ OpMemberDecorate %prevent_dce_block 0 Offset 0
+ OpMemberDecorate %prevent_dce_block 0 ColMajor
+ OpMemberDecorate %prevent_dce_block 0 MatrixStride 8
+ OpDecorate %prevent_dce DescriptorSet 2
+ OpDecorate %prevent_dce Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
@@ -32,43 +41,52 @@
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
- %void = OpTypeVoid
- %9 = OpTypeFunction %void
%half = OpTypeFloat 16
%v4half = OpTypeVector %half 4
%mat4v4half = OpTypeMatrix %v4half 4
+%prevent_dce_block = OpTypeStruct %mat4v4half
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+ %void = OpTypeVoid
+ %15 = OpTypeFunction %void
%half_0x1p_0 = OpConstant %half 0x1p+0
- %17 = OpConstantComposite %v4half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
- %18 = OpConstantComposite %mat4v4half %17 %17 %17 %17
+ %20 = OpConstantComposite %v4half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
+ %21 = OpConstantComposite %mat4v4half %20 %20 %20 %20
%_ptr_Function_mat4v4half = OpTypePointer Function %mat4v4half
- %21 = OpConstantNull %mat4v4half
- %22 = OpTypeFunction %v4float
+ %24 = OpConstantNull %mat4v4half
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_mat4v4half = OpTypePointer StorageBuffer %mat4v4half
+ %30 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
-%transpose_844869 = OpFunction %void None %9
- %12 = OpLabel
- %res = OpVariable %_ptr_Function_mat4v4half Function %21
- OpStore %res %18
+%transpose_844869 = OpFunction %void None %15
+ %18 = OpLabel
+ %res = OpVariable %_ptr_Function_mat4v4half Function %24
+ OpStore %res %21
+ %28 = OpAccessChain %_ptr_StorageBuffer_mat4v4half %prevent_dce %uint_0
+ %29 = OpLoad %mat4v4half %res
+ OpStore %28 %29
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %22
- %24 = OpLabel
- %25 = OpFunctionCall %void %transpose_844869
+%vertex_main_inner = OpFunction %v4float None %30
+ %32 = OpLabel
+ %33 = OpFunctionCall %void %transpose_844869
OpReturnValue %5
OpFunctionEnd
-%vertex_main = OpFunction %void None %9
- %27 = OpLabel
- %28 = OpFunctionCall %v4float %vertex_main_inner
- OpStore %value %28
+%vertex_main = OpFunction %void None %15
+ %35 = OpLabel
+ %36 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %36
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
-%fragment_main = OpFunction %void None %9
- %31 = OpLabel
- %32 = OpFunctionCall %void %transpose_844869
+%fragment_main = OpFunction %void None %15
+ %39 = OpLabel
+ %40 = OpFunctionCall %void %transpose_844869
OpReturn
OpFunctionEnd
-%compute_main = OpFunction %void None %9
- %34 = OpLabel
- %35 = OpFunctionCall %void %transpose_844869
+%compute_main = OpFunction %void None %15
+ %42 = OpLabel
+ %43 = OpFunctionCall %void %transpose_844869
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/transpose/844869.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/transpose/844869.wgsl.expected.wgsl
index 8f36149..e5637e9 100644
--- a/test/tint/builtins/gen/literal/transpose/844869.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/844869.wgsl.expected.wgsl
@@ -2,8 +2,11 @@
fn transpose_844869() {
var res : mat4x4<f16> = transpose(mat4x4<f16>(1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat4x4<f16>;
+
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
transpose_844869();
diff --git a/test/tint/builtins/gen/literal/transpose/84a763.wgsl b/test/tint/builtins/gen/literal/transpose/84a763.wgsl
index a965485..0851133 100644
--- a/test/tint/builtins/gen/literal/transpose/84a763.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/84a763.wgsl
@@ -25,7 +25,6 @@
fn transpose_84a763() {
var res = transpose(mat2x4(1., 1., 1., 1., 1., 1., 1., 1.));
}
-
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
transpose_84a763();
diff --git a/test/tint/builtins/gen/literal/transpose/854336.wgsl b/test/tint/builtins/gen/literal/transpose/854336.wgsl
index 68d9306..30adea2 100644
--- a/test/tint/builtins/gen/literal/transpose/854336.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/854336.wgsl
@@ -24,7 +24,9 @@
// fn transpose(mat<3, 3, f32>) -> mat<3, 3, f32>
fn transpose_854336() {
var res: mat3x3<f32> = transpose(mat3x3<f32>(1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat3x3<f32>;
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
diff --git a/test/tint/builtins/gen/literal/transpose/854336.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/transpose/854336.wgsl.expected.dxc.hlsl
index b094732..ec5e8e19 100644
--- a/test/tint/builtins/gen/literal/transpose/854336.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/854336.wgsl.expected.dxc.hlsl
@@ -1,5 +1,14 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, float3x3 value) {
+ prevent_dce.Store3((offset + 0u), asuint(value[0u]));
+ prevent_dce.Store3((offset + 16u), asuint(value[1u]));
+ prevent_dce.Store3((offset + 32u), asuint(value[2u]));
+}
+
void transpose_854336() {
float3x3 res = float3x3((1.0f).xxx, (1.0f).xxx, (1.0f).xxx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/854336.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/transpose/854336.wgsl.expected.fxc.hlsl
index b094732..ec5e8e19 100644
--- a/test/tint/builtins/gen/literal/transpose/854336.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/854336.wgsl.expected.fxc.hlsl
@@ -1,5 +1,14 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, float3x3 value) {
+ prevent_dce.Store3((offset + 0u), asuint(value[0u]));
+ prevent_dce.Store3((offset + 16u), asuint(value[1u]));
+ prevent_dce.Store3((offset + 32u), asuint(value[2u]));
+}
+
void transpose_854336() {
float3x3 res = float3x3((1.0f).xxx, (1.0f).xxx, (1.0f).xxx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/854336.wgsl.expected.glsl b/test/tint/builtins/gen/literal/transpose/854336.wgsl.expected.glsl
index 5092a2e..c33db6b 100644
--- a/test/tint/builtins/gen/literal/transpose/854336.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/transpose/854336.wgsl.expected.glsl
@@ -1,7 +1,18 @@
#version 310 es
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ mat3 inner;
+} prevent_dce;
+
+void assign_and_preserve_padding_prevent_dce(mat3 value) {
+ prevent_dce.inner[0] = value[0u];
+ prevent_dce.inner[1] = value[1u];
+ prevent_dce.inner[2] = value[2u];
+}
+
void transpose_854336() {
mat3 res = mat3(vec3(1.0f), vec3(1.0f), vec3(1.0f));
+ assign_and_preserve_padding_prevent_dce(res);
}
vec4 vertex_main() {
@@ -20,8 +31,19 @@
#version 310 es
precision mediump float;
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ mat3 inner;
+} prevent_dce;
+
+void assign_and_preserve_padding_prevent_dce(mat3 value) {
+ prevent_dce.inner[0] = value[0u];
+ prevent_dce.inner[1] = value[1u];
+ prevent_dce.inner[2] = value[2u];
+}
+
void transpose_854336() {
mat3 res = mat3(vec3(1.0f), vec3(1.0f), vec3(1.0f));
+ assign_and_preserve_padding_prevent_dce(res);
}
void fragment_main() {
@@ -34,8 +56,19 @@
}
#version 310 es
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ mat3 inner;
+} prevent_dce;
+
+void assign_and_preserve_padding_prevent_dce(mat3 value) {
+ prevent_dce.inner[0] = value[0u];
+ prevent_dce.inner[1] = value[1u];
+ prevent_dce.inner[2] = value[2u];
+}
+
void transpose_854336() {
mat3 res = mat3(vec3(1.0f), vec3(1.0f), vec3(1.0f));
+ assign_and_preserve_padding_prevent_dce(res);
}
void compute_main() {
diff --git a/test/tint/builtins/gen/literal/transpose/854336.wgsl.expected.msl b/test/tint/builtins/gen/literal/transpose/854336.wgsl.expected.msl
index 13b5c65..ca7a6d1 100644
--- a/test/tint/builtins/gen/literal/transpose/854336.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/transpose/854336.wgsl.expected.msl
@@ -1,33 +1,58 @@
#include <metal_stdlib>
using namespace metal;
-void transpose_854336() {
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+struct tint_packed_vec3_f32_array_element {
+ /* 0x0000 */ packed_float3 elements;
+ /* 0x000c */ tint_array<int8_t, 4> tint_pad;
+};
+
+void assign_and_preserve_padding(device tint_array<tint_packed_vec3_f32_array_element, 3>* const dest, float3x3 value) {
+ (*(dest))[0u].elements = packed_float3(value[0u]);
+ (*(dest))[1u].elements = packed_float3(value[1u]);
+ (*(dest))[2u].elements = packed_float3(value[2u]);
+}
+
+void transpose_854336(device tint_array<tint_packed_vec3_f32_array_element, 3>* const tint_symbol_1) {
float3x3 res = float3x3(float3(1.0f), float3(1.0f), float3(1.0f));
+ assign_and_preserve_padding(tint_symbol_1, res);
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner() {
- transpose_854336();
+float4 vertex_main_inner(device tint_array<tint_packed_vec3_f32_array_element, 3>* const tint_symbol_2) {
+ transpose_854336(tint_symbol_2);
return float4(0.0f);
}
-vertex tint_symbol vertex_main() {
- float4 const inner_result = vertex_main_inner();
+vertex tint_symbol vertex_main(device tint_array<tint_packed_vec3_f32_array_element, 3>* tint_symbol_3 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_3);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main() {
- transpose_854336();
+fragment void fragment_main(device tint_array<tint_packed_vec3_f32_array_element, 3>* tint_symbol_4 [[buffer(0)]]) {
+ transpose_854336(tint_symbol_4);
return;
}
-kernel void compute_main() {
- transpose_854336();
+kernel void compute_main(device tint_array<tint_packed_vec3_f32_array_element, 3>* tint_symbol_5 [[buffer(0)]]) {
+ transpose_854336(tint_symbol_5);
return;
}
diff --git a/test/tint/builtins/gen/literal/transpose/854336.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/transpose/854336.wgsl.expected.spvasm
index 3d93f98..02fb272 100644
--- a/test/tint/builtins/gen/literal/transpose/854336.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/transpose/854336.wgsl.expected.spvasm
@@ -1,68 +1,109 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 34
+; Bound: 59
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
- OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+ OpEntryPoint Vertex %vertex_main "vertex_main" %value_1 %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
OpEntryPoint GLCompute %compute_main "compute_main"
OpExecutionMode %fragment_main OriginUpperLeft
OpExecutionMode %compute_main LocalSize 1 1 1
- OpName %value "value"
+ OpName %value_1 "value_1"
OpName %vertex_point_size "vertex_point_size"
+ OpName %prevent_dce_block "prevent_dce_block"
+ OpMemberName %prevent_dce_block 0 "inner"
+ OpName %prevent_dce "prevent_dce"
+ OpName %assign_and_preserve_padding_prevent_dce "assign_and_preserve_padding_prevent_dce"
+ OpName %value "value"
OpName %transpose_854336 "transpose_854336"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
OpName %vertex_main "vertex_main"
OpName %fragment_main "fragment_main"
OpName %compute_main "compute_main"
- OpDecorate %value BuiltIn Position
+ OpDecorate %value_1 BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
+ OpDecorate %prevent_dce_block Block
+ OpMemberDecorate %prevent_dce_block 0 Offset 0
+ OpMemberDecorate %prevent_dce_block 0 ColMajor
+ OpMemberDecorate %prevent_dce_block 0 MatrixStride 16
+ OpDecorate %prevent_dce DescriptorSet 2
+ OpDecorate %prevent_dce Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
%5 = OpConstantNull %v4float
- %value = OpVariable %_ptr_Output_v4float Output %5
+ %value_1 = OpVariable %_ptr_Output_v4float Output %5
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
- %void = OpTypeVoid
- %9 = OpTypeFunction %void
%v3float = OpTypeVector %float 3
%mat3v3float = OpTypeMatrix %v3float 3
+%prevent_dce_block = OpTypeStruct %mat3v3float
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+ %void = OpTypeVoid
+ %14 = OpTypeFunction %void %mat3v3float
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+ %int = OpTypeInt 32 1
+ %22 = OpConstantNull %int
+%_ptr_StorageBuffer_v3float = OpTypePointer StorageBuffer %v3float
+ %25 = OpConstantNull %uint
+ %int_1 = OpConstant %int 1
+ %uint_1 = OpConstant %uint 1
+ %int_2 = OpConstant %int 2
+ %uint_2 = OpConstant %uint 2
+ %35 = OpTypeFunction %void
%float_1 = OpConstant %float 1
- %16 = OpConstantComposite %v3float %float_1 %float_1 %float_1
- %17 = OpConstantComposite %mat3v3float %16 %16 %16
+ %39 = OpConstantComposite %v3float %float_1 %float_1 %float_1
+ %40 = OpConstantComposite %mat3v3float %39 %39 %39
%_ptr_Function_mat3v3float = OpTypePointer Function %mat3v3float
- %20 = OpConstantNull %mat3v3float
- %21 = OpTypeFunction %v4float
-%transpose_854336 = OpFunction %void None %9
- %12 = OpLabel
- %res = OpVariable %_ptr_Function_mat3v3float Function %20
- OpStore %res %17
+ %43 = OpConstantNull %mat3v3float
+ %46 = OpTypeFunction %v4float
+%assign_and_preserve_padding_prevent_dce = OpFunction %void None %14
+ %value = OpFunctionParameter %mat3v3float
+ %18 = OpLabel
+ %24 = OpAccessChain %_ptr_StorageBuffer_v3float %prevent_dce %uint_0 %22
+ %26 = OpCompositeExtract %v3float %value 0
+ OpStore %24 %26
+ %28 = OpAccessChain %_ptr_StorageBuffer_v3float %prevent_dce %uint_0 %int_1
+ %30 = OpCompositeExtract %v3float %value 1
+ OpStore %28 %30
+ %32 = OpAccessChain %_ptr_StorageBuffer_v3float %prevent_dce %uint_0 %int_2
+ %34 = OpCompositeExtract %v3float %value 2
+ OpStore %32 %34
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %21
- %23 = OpLabel
- %24 = OpFunctionCall %void %transpose_854336
+%transpose_854336 = OpFunction %void None %35
+ %37 = OpLabel
+ %res = OpVariable %_ptr_Function_mat3v3float Function %43
+ OpStore %res %40
+ %45 = OpLoad %mat3v3float %res
+ %44 = OpFunctionCall %void %assign_and_preserve_padding_prevent_dce %45
+ OpReturn
+ OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %46
+ %48 = OpLabel
+ %49 = OpFunctionCall %void %transpose_854336
OpReturnValue %5
OpFunctionEnd
-%vertex_main = OpFunction %void None %9
- %26 = OpLabel
- %27 = OpFunctionCall %v4float %vertex_main_inner
- OpStore %value %27
+%vertex_main = OpFunction %void None %35
+ %51 = OpLabel
+ %52 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value_1 %52
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
-%fragment_main = OpFunction %void None %9
- %29 = OpLabel
- %30 = OpFunctionCall %void %transpose_854336
+%fragment_main = OpFunction %void None %35
+ %54 = OpLabel
+ %55 = OpFunctionCall %void %transpose_854336
OpReturn
OpFunctionEnd
-%compute_main = OpFunction %void None %9
- %32 = OpLabel
- %33 = OpFunctionCall %void %transpose_854336
+%compute_main = OpFunction %void None %35
+ %57 = OpLabel
+ %58 = OpFunctionCall %void %transpose_854336
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/transpose/854336.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/transpose/854336.wgsl.expected.wgsl
index f5aaa29..da84300 100644
--- a/test/tint/builtins/gen/literal/transpose/854336.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/854336.wgsl.expected.wgsl
@@ -1,7 +1,10 @@
fn transpose_854336() {
var res : mat3x3<f32> = transpose(mat3x3<f32>(1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat3x3<f32>;
+
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
transpose_854336();
diff --git a/test/tint/builtins/gen/literal/transpose/8c06ce.wgsl b/test/tint/builtins/gen/literal/transpose/8c06ce.wgsl
index b755bbd..1d2affa 100644
--- a/test/tint/builtins/gen/literal/transpose/8c06ce.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/8c06ce.wgsl
@@ -26,7 +26,9 @@
// fn transpose(mat<3, 4, f16>) -> mat<4, 3, f16>
fn transpose_8c06ce() {
var res: mat4x3<f16> = transpose(mat3x4<f16>(1.h, 1.h, 1.h, 1.h, 1.h, 1.h, 1.h, 1.h, 1.h, 1.h, 1.h, 1.h));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat4x3<f16>;
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
diff --git a/test/tint/builtins/gen/literal/transpose/8c06ce.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/transpose/8c06ce.wgsl.expected.dxc.hlsl
index fbbd86d..2fb3562 100644
--- a/test/tint/builtins/gen/literal/transpose/8c06ce.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/8c06ce.wgsl.expected.dxc.hlsl
@@ -1,5 +1,15 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, matrix<float16_t, 4, 3> value) {
+ prevent_dce.Store<vector<float16_t, 3> >((offset + 0u), value[0u]);
+ prevent_dce.Store<vector<float16_t, 3> >((offset + 8u), value[1u]);
+ prevent_dce.Store<vector<float16_t, 3> >((offset + 16u), value[2u]);
+ prevent_dce.Store<vector<float16_t, 3> >((offset + 24u), value[3u]);
+}
+
void transpose_8c06ce() {
matrix<float16_t, 4, 3> res = matrix<float16_t, 4, 3>((float16_t(1.0h)).xxx, (float16_t(1.0h)).xxx, (float16_t(1.0h)).xxx, (float16_t(1.0h)).xxx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/8c06ce.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/transpose/8c06ce.wgsl.expected.fxc.hlsl
index 34277a0..11937a1 100644
--- a/test/tint/builtins/gen/literal/transpose/8c06ce.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/8c06ce.wgsl.expected.fxc.hlsl
@@ -1,7 +1,17 @@
SKIP: FAILED
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, matrix<float16_t, 4, 3> value) {
+ prevent_dce.Store<vector<float16_t, 3> >((offset + 0u), value[0u]);
+ prevent_dce.Store<vector<float16_t, 3> >((offset + 8u), value[1u]);
+ prevent_dce.Store<vector<float16_t, 3> >((offset + 16u), value[2u]);
+ prevent_dce.Store<vector<float16_t, 3> >((offset + 24u), value[3u]);
+}
+
void transpose_8c06ce() {
matrix<float16_t, 4, 3> res = matrix<float16_t, 4, 3>((float16_t(1.0h)).xxx, (float16_t(1.0h)).xxx, (float16_t(1.0h)).xxx, (float16_t(1.0h)).xxx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/8c06ce.wgsl.expected.glsl b/test/tint/builtins/gen/literal/transpose/8c06ce.wgsl.expected.glsl
index 439264c..8e9363a 100644
--- a/test/tint/builtins/gen/literal/transpose/8c06ce.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/transpose/8c06ce.wgsl.expected.glsl
@@ -1,8 +1,20 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16mat4x3 inner;
+} prevent_dce;
+
+void assign_and_preserve_padding_prevent_dce(f16mat4x3 value) {
+ prevent_dce.inner[0] = value[0u];
+ prevent_dce.inner[1] = value[1u];
+ prevent_dce.inner[2] = value[2u];
+ prevent_dce.inner[3] = value[3u];
+}
+
void transpose_8c06ce() {
f16mat4x3 res = f16mat4x3(f16vec3(1.0hf), f16vec3(1.0hf), f16vec3(1.0hf), f16vec3(1.0hf));
+ assign_and_preserve_padding_prevent_dce(res);
}
vec4 vertex_main() {
@@ -22,8 +34,20 @@
#extension GL_AMD_gpu_shader_half_float : require
precision mediump float;
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16mat4x3 inner;
+} prevent_dce;
+
+void assign_and_preserve_padding_prevent_dce(f16mat4x3 value) {
+ prevent_dce.inner[0] = value[0u];
+ prevent_dce.inner[1] = value[1u];
+ prevent_dce.inner[2] = value[2u];
+ prevent_dce.inner[3] = value[3u];
+}
+
void transpose_8c06ce() {
f16mat4x3 res = f16mat4x3(f16vec3(1.0hf), f16vec3(1.0hf), f16vec3(1.0hf), f16vec3(1.0hf));
+ assign_and_preserve_padding_prevent_dce(res);
}
void fragment_main() {
@@ -37,8 +61,20 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16mat4x3 inner;
+} prevent_dce;
+
+void assign_and_preserve_padding_prevent_dce(f16mat4x3 value) {
+ prevent_dce.inner[0] = value[0u];
+ prevent_dce.inner[1] = value[1u];
+ prevent_dce.inner[2] = value[2u];
+ prevent_dce.inner[3] = value[3u];
+}
+
void transpose_8c06ce() {
f16mat4x3 res = f16mat4x3(f16vec3(1.0hf), f16vec3(1.0hf), f16vec3(1.0hf), f16vec3(1.0hf));
+ assign_and_preserve_padding_prevent_dce(res);
}
void compute_main() {
diff --git a/test/tint/builtins/gen/literal/transpose/8c06ce.wgsl.expected.msl b/test/tint/builtins/gen/literal/transpose/8c06ce.wgsl.expected.msl
index 3144e11..8715fbe 100644
--- a/test/tint/builtins/gen/literal/transpose/8c06ce.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/transpose/8c06ce.wgsl.expected.msl
@@ -1,33 +1,59 @@
#include <metal_stdlib>
using namespace metal;
-void transpose_8c06ce() {
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+struct tint_packed_vec3_f16_array_element {
+ /* 0x0000 */ packed_half3 elements;
+ /* 0x0006 */ tint_array<int8_t, 2> tint_pad;
+};
+
+void assign_and_preserve_padding(device tint_array<tint_packed_vec3_f16_array_element, 4>* const dest, half4x3 value) {
+ (*(dest))[0u].elements = packed_half3(value[0u]);
+ (*(dest))[1u].elements = packed_half3(value[1u]);
+ (*(dest))[2u].elements = packed_half3(value[2u]);
+ (*(dest))[3u].elements = packed_half3(value[3u]);
+}
+
+void transpose_8c06ce(device tint_array<tint_packed_vec3_f16_array_element, 4>* const tint_symbol_1) {
half4x3 res = half4x3(half3(1.0h), half3(1.0h), half3(1.0h), half3(1.0h));
+ assign_and_preserve_padding(tint_symbol_1, res);
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner() {
- transpose_8c06ce();
+float4 vertex_main_inner(device tint_array<tint_packed_vec3_f16_array_element, 4>* const tint_symbol_2) {
+ transpose_8c06ce(tint_symbol_2);
return float4(0.0f);
}
-vertex tint_symbol vertex_main() {
- float4 const inner_result = vertex_main_inner();
+vertex tint_symbol vertex_main(device tint_array<tint_packed_vec3_f16_array_element, 4>* tint_symbol_3 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_3);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main() {
- transpose_8c06ce();
+fragment void fragment_main(device tint_array<tint_packed_vec3_f16_array_element, 4>* tint_symbol_4 [[buffer(0)]]) {
+ transpose_8c06ce(tint_symbol_4);
return;
}
-kernel void compute_main() {
- transpose_8c06ce();
+kernel void compute_main(device tint_array<tint_packed_vec3_f16_array_element, 4>* tint_symbol_5 [[buffer(0)]]) {
+ transpose_8c06ce(tint_symbol_5);
return;
}
diff --git a/test/tint/builtins/gen/literal/transpose/8c06ce.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/transpose/8c06ce.wgsl.expected.spvasm
index ef9708c..29ee10e 100644
--- a/test/tint/builtins/gen/literal/transpose/8c06ce.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/transpose/8c06ce.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 36
+; Bound: 65
; Schema: 0
OpCapability Shader
OpCapability Float16
@@ -9,66 +9,112 @@
OpCapability StorageBuffer16BitAccess
OpCapability StorageInputOutput16
OpMemoryModel Logical GLSL450
- OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+ OpEntryPoint Vertex %vertex_main "vertex_main" %value_1 %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
OpEntryPoint GLCompute %compute_main "compute_main"
OpExecutionMode %fragment_main OriginUpperLeft
OpExecutionMode %compute_main LocalSize 1 1 1
- OpName %value "value"
+ OpName %value_1 "value_1"
OpName %vertex_point_size "vertex_point_size"
+ OpName %prevent_dce_block "prevent_dce_block"
+ OpMemberName %prevent_dce_block 0 "inner"
+ OpName %prevent_dce "prevent_dce"
+ OpName %assign_and_preserve_padding_prevent_dce "assign_and_preserve_padding_prevent_dce"
+ OpName %value "value"
OpName %transpose_8c06ce "transpose_8c06ce"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
OpName %vertex_main "vertex_main"
OpName %fragment_main "fragment_main"
OpName %compute_main "compute_main"
- OpDecorate %value BuiltIn Position
+ OpDecorate %value_1 BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
+ OpDecorate %prevent_dce_block Block
+ OpMemberDecorate %prevent_dce_block 0 Offset 0
+ OpMemberDecorate %prevent_dce_block 0 ColMajor
+ OpMemberDecorate %prevent_dce_block 0 MatrixStride 8
+ OpDecorate %prevent_dce DescriptorSet 2
+ OpDecorate %prevent_dce Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
%5 = OpConstantNull %v4float
- %value = OpVariable %_ptr_Output_v4float Output %5
+ %value_1 = OpVariable %_ptr_Output_v4float Output %5
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
- %void = OpTypeVoid
- %9 = OpTypeFunction %void
%half = OpTypeFloat 16
%v3half = OpTypeVector %half 3
%mat4v3half = OpTypeMatrix %v3half 4
+%prevent_dce_block = OpTypeStruct %mat4v3half
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+ %void = OpTypeVoid
+ %15 = OpTypeFunction %void %mat4v3half
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+ %int = OpTypeInt 32 1
+ %23 = OpConstantNull %int
+%_ptr_StorageBuffer_v3half = OpTypePointer StorageBuffer %v3half
+ %26 = OpConstantNull %uint
+ %int_1 = OpConstant %int 1
+ %uint_1 = OpConstant %uint 1
+ %int_2 = OpConstant %int 2
+ %uint_2 = OpConstant %uint 2
+ %int_3 = OpConstant %int 3
+ %uint_3 = OpConstant %uint 3
+ %40 = OpTypeFunction %void
%half_0x1p_0 = OpConstant %half 0x1p+0
- %17 = OpConstantComposite %v3half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
- %18 = OpConstantComposite %mat4v3half %17 %17 %17 %17
+ %44 = OpConstantComposite %v3half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
+ %45 = OpConstantComposite %mat4v3half %44 %44 %44 %44
%_ptr_Function_mat4v3half = OpTypePointer Function %mat4v3half
- %21 = OpConstantNull %mat4v3half
- %22 = OpTypeFunction %v4float
+ %48 = OpConstantNull %mat4v3half
+ %51 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
-%transpose_8c06ce = OpFunction %void None %9
- %12 = OpLabel
- %res = OpVariable %_ptr_Function_mat4v3half Function %21
- OpStore %res %18
+%assign_and_preserve_padding_prevent_dce = OpFunction %void None %15
+ %value = OpFunctionParameter %mat4v3half
+ %19 = OpLabel
+ %25 = OpAccessChain %_ptr_StorageBuffer_v3half %prevent_dce %uint_0 %23
+ %27 = OpCompositeExtract %v3half %value 0
+ OpStore %25 %27
+ %29 = OpAccessChain %_ptr_StorageBuffer_v3half %prevent_dce %uint_0 %int_1
+ %31 = OpCompositeExtract %v3half %value 1
+ OpStore %29 %31
+ %33 = OpAccessChain %_ptr_StorageBuffer_v3half %prevent_dce %uint_0 %int_2
+ %35 = OpCompositeExtract %v3half %value 2
+ OpStore %33 %35
+ %37 = OpAccessChain %_ptr_StorageBuffer_v3half %prevent_dce %uint_0 %int_3
+ %39 = OpCompositeExtract %v3half %value 3
+ OpStore %37 %39
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %22
- %24 = OpLabel
- %25 = OpFunctionCall %void %transpose_8c06ce
+%transpose_8c06ce = OpFunction %void None %40
+ %42 = OpLabel
+ %res = OpVariable %_ptr_Function_mat4v3half Function %48
+ OpStore %res %45
+ %50 = OpLoad %mat4v3half %res
+ %49 = OpFunctionCall %void %assign_and_preserve_padding_prevent_dce %50
+ OpReturn
+ OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %51
+ %53 = OpLabel
+ %54 = OpFunctionCall %void %transpose_8c06ce
OpReturnValue %5
OpFunctionEnd
-%vertex_main = OpFunction %void None %9
- %27 = OpLabel
- %28 = OpFunctionCall %v4float %vertex_main_inner
- OpStore %value %28
+%vertex_main = OpFunction %void None %40
+ %56 = OpLabel
+ %57 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value_1 %57
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
-%fragment_main = OpFunction %void None %9
- %31 = OpLabel
- %32 = OpFunctionCall %void %transpose_8c06ce
+%fragment_main = OpFunction %void None %40
+ %60 = OpLabel
+ %61 = OpFunctionCall %void %transpose_8c06ce
OpReturn
OpFunctionEnd
-%compute_main = OpFunction %void None %9
- %34 = OpLabel
- %35 = OpFunctionCall %void %transpose_8c06ce
+%compute_main = OpFunction %void None %40
+ %63 = OpLabel
+ %64 = OpFunctionCall %void %transpose_8c06ce
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/transpose/8c06ce.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/transpose/8c06ce.wgsl.expected.wgsl
index f2b4dc4..1335df4 100644
--- a/test/tint/builtins/gen/literal/transpose/8c06ce.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/8c06ce.wgsl.expected.wgsl
@@ -2,8 +2,11 @@
fn transpose_8c06ce() {
var res : mat4x3<f16> = transpose(mat3x4<f16>(1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat4x3<f16>;
+
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
transpose_8c06ce();
diff --git a/test/tint/builtins/gen/literal/transpose/ace596.wgsl b/test/tint/builtins/gen/literal/transpose/ace596.wgsl
index b55c252..b75a104 100644
--- a/test/tint/builtins/gen/literal/transpose/ace596.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/ace596.wgsl
@@ -25,7 +25,6 @@
fn transpose_ace596() {
var res = transpose(mat3x2(1., 1., 1., 1., 1., 1.));
}
-
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
transpose_ace596();
diff --git a/test/tint/builtins/gen/literal/transpose/b9ad1f.wgsl b/test/tint/builtins/gen/literal/transpose/b9ad1f.wgsl
index d2e4c2b..1a750f6 100644
--- a/test/tint/builtins/gen/literal/transpose/b9ad1f.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/b9ad1f.wgsl
@@ -26,7 +26,9 @@
// fn transpose(mat<3, 2, f16>) -> mat<2, 3, f16>
fn transpose_b9ad1f() {
var res: mat2x3<f16> = transpose(mat3x2<f16>(1.h, 1.h, 1.h, 1.h, 1.h, 1.h));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat2x3<f16>;
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
diff --git a/test/tint/builtins/gen/literal/transpose/b9ad1f.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/transpose/b9ad1f.wgsl.expected.dxc.hlsl
index 126d8cb..7eb6c0c 100644
--- a/test/tint/builtins/gen/literal/transpose/b9ad1f.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/b9ad1f.wgsl.expected.dxc.hlsl
@@ -1,5 +1,13 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, matrix<float16_t, 2, 3> value) {
+ prevent_dce.Store<vector<float16_t, 3> >((offset + 0u), value[0u]);
+ prevent_dce.Store<vector<float16_t, 3> >((offset + 8u), value[1u]);
+}
+
void transpose_b9ad1f() {
matrix<float16_t, 2, 3> res = matrix<float16_t, 2, 3>((float16_t(1.0h)).xxx, (float16_t(1.0h)).xxx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/b9ad1f.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/transpose/b9ad1f.wgsl.expected.fxc.hlsl
index 7b35fe7..df8512e 100644
--- a/test/tint/builtins/gen/literal/transpose/b9ad1f.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/b9ad1f.wgsl.expected.fxc.hlsl
@@ -1,7 +1,15 @@
SKIP: FAILED
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, matrix<float16_t, 2, 3> value) {
+ prevent_dce.Store<vector<float16_t, 3> >((offset + 0u), value[0u]);
+ prevent_dce.Store<vector<float16_t, 3> >((offset + 8u), value[1u]);
+}
+
void transpose_b9ad1f() {
matrix<float16_t, 2, 3> res = matrix<float16_t, 2, 3>((float16_t(1.0h)).xxx, (float16_t(1.0h)).xxx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/b9ad1f.wgsl.expected.glsl b/test/tint/builtins/gen/literal/transpose/b9ad1f.wgsl.expected.glsl
index 82e02e7..e59d4ad 100644
--- a/test/tint/builtins/gen/literal/transpose/b9ad1f.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/transpose/b9ad1f.wgsl.expected.glsl
@@ -1,8 +1,18 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16mat2x3 inner;
+} prevent_dce;
+
+void assign_and_preserve_padding_prevent_dce(f16mat2x3 value) {
+ prevent_dce.inner[0] = value[0u];
+ prevent_dce.inner[1] = value[1u];
+}
+
void transpose_b9ad1f() {
f16mat2x3 res = f16mat2x3(f16vec3(1.0hf), f16vec3(1.0hf));
+ assign_and_preserve_padding_prevent_dce(res);
}
vec4 vertex_main() {
@@ -22,8 +32,18 @@
#extension GL_AMD_gpu_shader_half_float : require
precision mediump float;
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16mat2x3 inner;
+} prevent_dce;
+
+void assign_and_preserve_padding_prevent_dce(f16mat2x3 value) {
+ prevent_dce.inner[0] = value[0u];
+ prevent_dce.inner[1] = value[1u];
+}
+
void transpose_b9ad1f() {
f16mat2x3 res = f16mat2x3(f16vec3(1.0hf), f16vec3(1.0hf));
+ assign_and_preserve_padding_prevent_dce(res);
}
void fragment_main() {
@@ -37,8 +57,18 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16mat2x3 inner;
+} prevent_dce;
+
+void assign_and_preserve_padding_prevent_dce(f16mat2x3 value) {
+ prevent_dce.inner[0] = value[0u];
+ prevent_dce.inner[1] = value[1u];
+}
+
void transpose_b9ad1f() {
f16mat2x3 res = f16mat2x3(f16vec3(1.0hf), f16vec3(1.0hf));
+ assign_and_preserve_padding_prevent_dce(res);
}
void compute_main() {
diff --git a/test/tint/builtins/gen/literal/transpose/b9ad1f.wgsl.expected.msl b/test/tint/builtins/gen/literal/transpose/b9ad1f.wgsl.expected.msl
index 9d97e55..9752e51 100644
--- a/test/tint/builtins/gen/literal/transpose/b9ad1f.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/transpose/b9ad1f.wgsl.expected.msl
@@ -1,33 +1,57 @@
#include <metal_stdlib>
using namespace metal;
-void transpose_b9ad1f() {
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+struct tint_packed_vec3_f16_array_element {
+ /* 0x0000 */ packed_half3 elements;
+ /* 0x0006 */ tint_array<int8_t, 2> tint_pad;
+};
+
+void assign_and_preserve_padding(device tint_array<tint_packed_vec3_f16_array_element, 2>* const dest, half2x3 value) {
+ (*(dest))[0u].elements = packed_half3(value[0u]);
+ (*(dest))[1u].elements = packed_half3(value[1u]);
+}
+
+void transpose_b9ad1f(device tint_array<tint_packed_vec3_f16_array_element, 2>* const tint_symbol_1) {
half2x3 res = half2x3(half3(1.0h), half3(1.0h));
+ assign_and_preserve_padding(tint_symbol_1, res);
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner() {
- transpose_b9ad1f();
+float4 vertex_main_inner(device tint_array<tint_packed_vec3_f16_array_element, 2>* const tint_symbol_2) {
+ transpose_b9ad1f(tint_symbol_2);
return float4(0.0f);
}
-vertex tint_symbol vertex_main() {
- float4 const inner_result = vertex_main_inner();
+vertex tint_symbol vertex_main(device tint_array<tint_packed_vec3_f16_array_element, 2>* tint_symbol_3 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_3);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main() {
- transpose_b9ad1f();
+fragment void fragment_main(device tint_array<tint_packed_vec3_f16_array_element, 2>* tint_symbol_4 [[buffer(0)]]) {
+ transpose_b9ad1f(tint_symbol_4);
return;
}
-kernel void compute_main() {
- transpose_b9ad1f();
+kernel void compute_main(device tint_array<tint_packed_vec3_f16_array_element, 2>* tint_symbol_5 [[buffer(0)]]) {
+ transpose_b9ad1f(tint_symbol_5);
return;
}
diff --git a/test/tint/builtins/gen/literal/transpose/b9ad1f.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/transpose/b9ad1f.wgsl.expected.spvasm
index 63db7d8..08aa33c 100644
--- a/test/tint/builtins/gen/literal/transpose/b9ad1f.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/transpose/b9ad1f.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 36
+; Bound: 57
; Schema: 0
OpCapability Shader
OpCapability Float16
@@ -9,66 +9,102 @@
OpCapability StorageBuffer16BitAccess
OpCapability StorageInputOutput16
OpMemoryModel Logical GLSL450
- OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+ OpEntryPoint Vertex %vertex_main "vertex_main" %value_1 %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
OpEntryPoint GLCompute %compute_main "compute_main"
OpExecutionMode %fragment_main OriginUpperLeft
OpExecutionMode %compute_main LocalSize 1 1 1
- OpName %value "value"
+ OpName %value_1 "value_1"
OpName %vertex_point_size "vertex_point_size"
+ OpName %prevent_dce_block "prevent_dce_block"
+ OpMemberName %prevent_dce_block 0 "inner"
+ OpName %prevent_dce "prevent_dce"
+ OpName %assign_and_preserve_padding_prevent_dce "assign_and_preserve_padding_prevent_dce"
+ OpName %value "value"
OpName %transpose_b9ad1f "transpose_b9ad1f"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
OpName %vertex_main "vertex_main"
OpName %fragment_main "fragment_main"
OpName %compute_main "compute_main"
- OpDecorate %value BuiltIn Position
+ OpDecorate %value_1 BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
+ OpDecorate %prevent_dce_block Block
+ OpMemberDecorate %prevent_dce_block 0 Offset 0
+ OpMemberDecorate %prevent_dce_block 0 ColMajor
+ OpMemberDecorate %prevent_dce_block 0 MatrixStride 8
+ OpDecorate %prevent_dce DescriptorSet 2
+ OpDecorate %prevent_dce Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
%5 = OpConstantNull %v4float
- %value = OpVariable %_ptr_Output_v4float Output %5
+ %value_1 = OpVariable %_ptr_Output_v4float Output %5
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
- %void = OpTypeVoid
- %9 = OpTypeFunction %void
%half = OpTypeFloat 16
%v3half = OpTypeVector %half 3
%mat2v3half = OpTypeMatrix %v3half 2
+%prevent_dce_block = OpTypeStruct %mat2v3half
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+ %void = OpTypeVoid
+ %15 = OpTypeFunction %void %mat2v3half
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+ %int = OpTypeInt 32 1
+ %23 = OpConstantNull %int
+%_ptr_StorageBuffer_v3half = OpTypePointer StorageBuffer %v3half
+ %26 = OpConstantNull %uint
+ %int_1 = OpConstant %int 1
+ %uint_1 = OpConstant %uint 1
+ %32 = OpTypeFunction %void
%half_0x1p_0 = OpConstant %half 0x1p+0
- %17 = OpConstantComposite %v3half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
- %18 = OpConstantComposite %mat2v3half %17 %17
+ %36 = OpConstantComposite %v3half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
+ %37 = OpConstantComposite %mat2v3half %36 %36
%_ptr_Function_mat2v3half = OpTypePointer Function %mat2v3half
- %21 = OpConstantNull %mat2v3half
- %22 = OpTypeFunction %v4float
+ %40 = OpConstantNull %mat2v3half
+ %43 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
-%transpose_b9ad1f = OpFunction %void None %9
- %12 = OpLabel
- %res = OpVariable %_ptr_Function_mat2v3half Function %21
- OpStore %res %18
+%assign_and_preserve_padding_prevent_dce = OpFunction %void None %15
+ %value = OpFunctionParameter %mat2v3half
+ %19 = OpLabel
+ %25 = OpAccessChain %_ptr_StorageBuffer_v3half %prevent_dce %uint_0 %23
+ %27 = OpCompositeExtract %v3half %value 0
+ OpStore %25 %27
+ %29 = OpAccessChain %_ptr_StorageBuffer_v3half %prevent_dce %uint_0 %int_1
+ %31 = OpCompositeExtract %v3half %value 1
+ OpStore %29 %31
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %22
- %24 = OpLabel
- %25 = OpFunctionCall %void %transpose_b9ad1f
+%transpose_b9ad1f = OpFunction %void None %32
+ %34 = OpLabel
+ %res = OpVariable %_ptr_Function_mat2v3half Function %40
+ OpStore %res %37
+ %42 = OpLoad %mat2v3half %res
+ %41 = OpFunctionCall %void %assign_and_preserve_padding_prevent_dce %42
+ OpReturn
+ OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %43
+ %45 = OpLabel
+ %46 = OpFunctionCall %void %transpose_b9ad1f
OpReturnValue %5
OpFunctionEnd
-%vertex_main = OpFunction %void None %9
- %27 = OpLabel
- %28 = OpFunctionCall %v4float %vertex_main_inner
- OpStore %value %28
+%vertex_main = OpFunction %void None %32
+ %48 = OpLabel
+ %49 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value_1 %49
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
-%fragment_main = OpFunction %void None %9
- %31 = OpLabel
- %32 = OpFunctionCall %void %transpose_b9ad1f
+%fragment_main = OpFunction %void None %32
+ %52 = OpLabel
+ %53 = OpFunctionCall %void %transpose_b9ad1f
OpReturn
OpFunctionEnd
-%compute_main = OpFunction %void None %9
- %34 = OpLabel
- %35 = OpFunctionCall %void %transpose_b9ad1f
+%compute_main = OpFunction %void None %32
+ %55 = OpLabel
+ %56 = OpFunctionCall %void %transpose_b9ad1f
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/transpose/b9ad1f.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/transpose/b9ad1f.wgsl.expected.wgsl
index 3d06432..8265d11 100644
--- a/test/tint/builtins/gen/literal/transpose/b9ad1f.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/b9ad1f.wgsl.expected.wgsl
@@ -2,8 +2,11 @@
fn transpose_b9ad1f() {
var res : mat2x3<f16> = transpose(mat3x2<f16>(1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat2x3<f16>;
+
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
transpose_b9ad1f();
diff --git a/test/tint/builtins/gen/literal/transpose/c1b600.wgsl b/test/tint/builtins/gen/literal/transpose/c1b600.wgsl
index 888b2f4..e540291 100644
--- a/test/tint/builtins/gen/literal/transpose/c1b600.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/c1b600.wgsl
@@ -24,7 +24,9 @@
// fn transpose(mat<4, 4, f32>) -> mat<4, 4, f32>
fn transpose_c1b600() {
var res: mat4x4<f32> = transpose(mat4x4<f32>(1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat4x4<f32>;
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
diff --git a/test/tint/builtins/gen/literal/transpose/c1b600.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/transpose/c1b600.wgsl.expected.dxc.hlsl
index d06f52d..5411696 100644
--- a/test/tint/builtins/gen/literal/transpose/c1b600.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/c1b600.wgsl.expected.dxc.hlsl
@@ -1,5 +1,15 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, float4x4 value) {
+ prevent_dce.Store4((offset + 0u), asuint(value[0u]));
+ prevent_dce.Store4((offset + 16u), asuint(value[1u]));
+ prevent_dce.Store4((offset + 32u), asuint(value[2u]));
+ prevent_dce.Store4((offset + 48u), asuint(value[3u]));
+}
+
void transpose_c1b600() {
float4x4 res = float4x4((1.0f).xxxx, (1.0f).xxxx, (1.0f).xxxx, (1.0f).xxxx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/c1b600.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/transpose/c1b600.wgsl.expected.fxc.hlsl
index d06f52d..5411696 100644
--- a/test/tint/builtins/gen/literal/transpose/c1b600.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/c1b600.wgsl.expected.fxc.hlsl
@@ -1,5 +1,15 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, float4x4 value) {
+ prevent_dce.Store4((offset + 0u), asuint(value[0u]));
+ prevent_dce.Store4((offset + 16u), asuint(value[1u]));
+ prevent_dce.Store4((offset + 32u), asuint(value[2u]));
+ prevent_dce.Store4((offset + 48u), asuint(value[3u]));
+}
+
void transpose_c1b600() {
float4x4 res = float4x4((1.0f).xxxx, (1.0f).xxxx, (1.0f).xxxx, (1.0f).xxxx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/c1b600.wgsl.expected.glsl b/test/tint/builtins/gen/literal/transpose/c1b600.wgsl.expected.glsl
index 088adc7..697f42c 100644
--- a/test/tint/builtins/gen/literal/transpose/c1b600.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/transpose/c1b600.wgsl.expected.glsl
@@ -1,7 +1,12 @@
#version 310 es
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ mat4 inner;
+} prevent_dce;
+
void transpose_c1b600() {
mat4 res = mat4(vec4(1.0f), vec4(1.0f), vec4(1.0f), vec4(1.0f));
+ prevent_dce.inner = res;
}
vec4 vertex_main() {
@@ -20,8 +25,13 @@
#version 310 es
precision mediump float;
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ mat4 inner;
+} prevent_dce;
+
void transpose_c1b600() {
mat4 res = mat4(vec4(1.0f), vec4(1.0f), vec4(1.0f), vec4(1.0f));
+ prevent_dce.inner = res;
}
void fragment_main() {
@@ -34,8 +44,13 @@
}
#version 310 es
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ mat4 inner;
+} prevent_dce;
+
void transpose_c1b600() {
mat4 res = mat4(vec4(1.0f), vec4(1.0f), vec4(1.0f), vec4(1.0f));
+ prevent_dce.inner = res;
}
void compute_main() {
diff --git a/test/tint/builtins/gen/literal/transpose/c1b600.wgsl.expected.msl b/test/tint/builtins/gen/literal/transpose/c1b600.wgsl.expected.msl
index eedd9e4..719f723 100644
--- a/test/tint/builtins/gen/literal/transpose/c1b600.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/transpose/c1b600.wgsl.expected.msl
@@ -1,33 +1,34 @@
#include <metal_stdlib>
using namespace metal;
-void transpose_c1b600() {
+void transpose_c1b600(device float4x4* const tint_symbol_1) {
float4x4 res = float4x4(float4(1.0f), float4(1.0f), float4(1.0f), float4(1.0f));
+ *(tint_symbol_1) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner() {
- transpose_c1b600();
+float4 vertex_main_inner(device float4x4* const tint_symbol_2) {
+ transpose_c1b600(tint_symbol_2);
return float4(0.0f);
}
-vertex tint_symbol vertex_main() {
- float4 const inner_result = vertex_main_inner();
+vertex tint_symbol vertex_main(device float4x4* tint_symbol_3 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_3);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main() {
- transpose_c1b600();
+fragment void fragment_main(device float4x4* tint_symbol_4 [[buffer(0)]]) {
+ transpose_c1b600(tint_symbol_4);
return;
}
-kernel void compute_main() {
- transpose_c1b600();
+kernel void compute_main(device float4x4* tint_symbol_5 [[buffer(0)]]) {
+ transpose_c1b600(tint_symbol_5);
return;
}
diff --git a/test/tint/builtins/gen/literal/transpose/c1b600.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/transpose/c1b600.wgsl.expected.spvasm
index 8b8bf4e..803e526 100644
--- a/test/tint/builtins/gen/literal/transpose/c1b600.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/transpose/c1b600.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 33
+; Bound: 41
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -12,6 +12,9 @@
OpExecutionMode %compute_main LocalSize 1 1 1
OpName %value "value"
OpName %vertex_point_size "vertex_point_size"
+ OpName %prevent_dce_block "prevent_dce_block"
+ OpMemberName %prevent_dce_block 0 "inner"
+ OpName %prevent_dce "prevent_dce"
OpName %transpose_c1b600 "transpose_c1b600"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
@@ -20,6 +23,12 @@
OpName %compute_main "compute_main"
OpDecorate %value BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
+ OpDecorate %prevent_dce_block Block
+ OpMemberDecorate %prevent_dce_block 0 Offset 0
+ OpMemberDecorate %prevent_dce_block 0 ColMajor
+ OpMemberDecorate %prevent_dce_block 0 MatrixStride 16
+ OpDecorate %prevent_dce DescriptorSet 2
+ OpDecorate %prevent_dce Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
@@ -28,40 +37,49 @@
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
- %void = OpTypeVoid
- %9 = OpTypeFunction %void
%mat4v4float = OpTypeMatrix %v4float 4
+%prevent_dce_block = OpTypeStruct %mat4v4float
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+ %void = OpTypeVoid
+ %13 = OpTypeFunction %void
%float_1 = OpConstant %float 1
- %15 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
- %16 = OpConstantComposite %mat4v4float %15 %15 %15 %15
+ %18 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
+ %19 = OpConstantComposite %mat4v4float %18 %18 %18 %18
%_ptr_Function_mat4v4float = OpTypePointer Function %mat4v4float
- %19 = OpConstantNull %mat4v4float
- %20 = OpTypeFunction %v4float
-%transpose_c1b600 = OpFunction %void None %9
- %12 = OpLabel
- %res = OpVariable %_ptr_Function_mat4v4float Function %19
- OpStore %res %16
+ %22 = OpConstantNull %mat4v4float
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_mat4v4float = OpTypePointer StorageBuffer %mat4v4float
+ %28 = OpTypeFunction %v4float
+%transpose_c1b600 = OpFunction %void None %13
+ %16 = OpLabel
+ %res = OpVariable %_ptr_Function_mat4v4float Function %22
+ OpStore %res %19
+ %26 = OpAccessChain %_ptr_StorageBuffer_mat4v4float %prevent_dce %uint_0
+ %27 = OpLoad %mat4v4float %res
+ OpStore %26 %27
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %20
- %22 = OpLabel
- %23 = OpFunctionCall %void %transpose_c1b600
+%vertex_main_inner = OpFunction %v4float None %28
+ %30 = OpLabel
+ %31 = OpFunctionCall %void %transpose_c1b600
OpReturnValue %5
OpFunctionEnd
-%vertex_main = OpFunction %void None %9
- %25 = OpLabel
- %26 = OpFunctionCall %v4float %vertex_main_inner
- OpStore %value %26
+%vertex_main = OpFunction %void None %13
+ %33 = OpLabel
+ %34 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %34
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
-%fragment_main = OpFunction %void None %9
- %28 = OpLabel
- %29 = OpFunctionCall %void %transpose_c1b600
+%fragment_main = OpFunction %void None %13
+ %36 = OpLabel
+ %37 = OpFunctionCall %void %transpose_c1b600
OpReturn
OpFunctionEnd
-%compute_main = OpFunction %void None %9
- %31 = OpLabel
- %32 = OpFunctionCall %void %transpose_c1b600
+%compute_main = OpFunction %void None %13
+ %39 = OpLabel
+ %40 = OpFunctionCall %void %transpose_c1b600
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/transpose/c1b600.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/transpose/c1b600.wgsl.expected.wgsl
index bf9bbea..2f63028 100644
--- a/test/tint/builtins/gen/literal/transpose/c1b600.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/c1b600.wgsl.expected.wgsl
@@ -1,7 +1,10 @@
fn transpose_c1b600() {
var res : mat4x4<f32> = transpose(mat4x4<f32>(1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat4x4<f32>;
+
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
transpose_c1b600();
diff --git a/test/tint/builtins/gen/literal/transpose/d6faec.wgsl b/test/tint/builtins/gen/literal/transpose/d6faec.wgsl
index ebaa669..c61c862 100644
--- a/test/tint/builtins/gen/literal/transpose/d6faec.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/d6faec.wgsl
@@ -26,7 +26,9 @@
// fn transpose(mat<2, 3, f16>) -> mat<3, 2, f16>
fn transpose_d6faec() {
var res: mat3x2<f16> = transpose(mat2x3<f16>(1.h, 1.h, 1.h, 1.h, 1.h, 1.h));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat3x2<f16>;
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
diff --git a/test/tint/builtins/gen/literal/transpose/d6faec.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/transpose/d6faec.wgsl.expected.dxc.hlsl
index 85b78fd..1196440 100644
--- a/test/tint/builtins/gen/literal/transpose/d6faec.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/d6faec.wgsl.expected.dxc.hlsl
@@ -1,5 +1,14 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, matrix<float16_t, 3, 2> value) {
+ prevent_dce.Store<vector<float16_t, 2> >((offset + 0u), value[0u]);
+ prevent_dce.Store<vector<float16_t, 2> >((offset + 4u), value[1u]);
+ prevent_dce.Store<vector<float16_t, 2> >((offset + 8u), value[2u]);
+}
+
void transpose_d6faec() {
matrix<float16_t, 3, 2> res = matrix<float16_t, 3, 2>((float16_t(1.0h)).xx, (float16_t(1.0h)).xx, (float16_t(1.0h)).xx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/d6faec.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/transpose/d6faec.wgsl.expected.fxc.hlsl
index 961226e..5852f9e 100644
--- a/test/tint/builtins/gen/literal/transpose/d6faec.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/d6faec.wgsl.expected.fxc.hlsl
@@ -1,7 +1,16 @@
SKIP: FAILED
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, matrix<float16_t, 3, 2> value) {
+ prevent_dce.Store<vector<float16_t, 2> >((offset + 0u), value[0u]);
+ prevent_dce.Store<vector<float16_t, 2> >((offset + 4u), value[1u]);
+ prevent_dce.Store<vector<float16_t, 2> >((offset + 8u), value[2u]);
+}
+
void transpose_d6faec() {
matrix<float16_t, 3, 2> res = matrix<float16_t, 3, 2>((float16_t(1.0h)).xx, (float16_t(1.0h)).xx, (float16_t(1.0h)).xx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/d6faec.wgsl.expected.glsl b/test/tint/builtins/gen/literal/transpose/d6faec.wgsl.expected.glsl
index 6028fb6a..d4dff64 100644
--- a/test/tint/builtins/gen/literal/transpose/d6faec.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/transpose/d6faec.wgsl.expected.glsl
@@ -1,8 +1,13 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16mat3x2 inner;
+} prevent_dce;
+
void transpose_d6faec() {
f16mat3x2 res = f16mat3x2(f16vec2(1.0hf), f16vec2(1.0hf), f16vec2(1.0hf));
+ prevent_dce.inner = res;
}
vec4 vertex_main() {
@@ -22,8 +27,13 @@
#extension GL_AMD_gpu_shader_half_float : require
precision mediump float;
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16mat3x2 inner;
+} prevent_dce;
+
void transpose_d6faec() {
f16mat3x2 res = f16mat3x2(f16vec2(1.0hf), f16vec2(1.0hf), f16vec2(1.0hf));
+ prevent_dce.inner = res;
}
void fragment_main() {
@@ -37,8 +47,13 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16mat3x2 inner;
+} prevent_dce;
+
void transpose_d6faec() {
f16mat3x2 res = f16mat3x2(f16vec2(1.0hf), f16vec2(1.0hf), f16vec2(1.0hf));
+ prevent_dce.inner = res;
}
void compute_main() {
diff --git a/test/tint/builtins/gen/literal/transpose/d6faec.wgsl.expected.msl b/test/tint/builtins/gen/literal/transpose/d6faec.wgsl.expected.msl
index 8b55e1a..1904769 100644
--- a/test/tint/builtins/gen/literal/transpose/d6faec.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/transpose/d6faec.wgsl.expected.msl
@@ -1,33 +1,34 @@
#include <metal_stdlib>
using namespace metal;
-void transpose_d6faec() {
+void transpose_d6faec(device half3x2* const tint_symbol_1) {
half3x2 res = half3x2(half2(1.0h), half2(1.0h), half2(1.0h));
+ *(tint_symbol_1) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner() {
- transpose_d6faec();
+float4 vertex_main_inner(device half3x2* const tint_symbol_2) {
+ transpose_d6faec(tint_symbol_2);
return float4(0.0f);
}
-vertex tint_symbol vertex_main() {
- float4 const inner_result = vertex_main_inner();
+vertex tint_symbol vertex_main(device half3x2* tint_symbol_3 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_3);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main() {
- transpose_d6faec();
+fragment void fragment_main(device half3x2* tint_symbol_4 [[buffer(0)]]) {
+ transpose_d6faec(tint_symbol_4);
return;
}
-kernel void compute_main() {
- transpose_d6faec();
+kernel void compute_main(device half3x2* tint_symbol_5 [[buffer(0)]]) {
+ transpose_d6faec(tint_symbol_5);
return;
}
diff --git a/test/tint/builtins/gen/literal/transpose/d6faec.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/transpose/d6faec.wgsl.expected.spvasm
index d5c2f0a..2c380d7 100644
--- a/test/tint/builtins/gen/literal/transpose/d6faec.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/transpose/d6faec.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 36
+; Bound: 44
; Schema: 0
OpCapability Shader
OpCapability Float16
@@ -16,6 +16,9 @@
OpExecutionMode %compute_main LocalSize 1 1 1
OpName %value "value"
OpName %vertex_point_size "vertex_point_size"
+ OpName %prevent_dce_block "prevent_dce_block"
+ OpMemberName %prevent_dce_block 0 "inner"
+ OpName %prevent_dce "prevent_dce"
OpName %transpose_d6faec "transpose_d6faec"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
@@ -24,6 +27,12 @@
OpName %compute_main "compute_main"
OpDecorate %value BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
+ OpDecorate %prevent_dce_block Block
+ OpMemberDecorate %prevent_dce_block 0 Offset 0
+ OpMemberDecorate %prevent_dce_block 0 ColMajor
+ OpMemberDecorate %prevent_dce_block 0 MatrixStride 4
+ OpDecorate %prevent_dce DescriptorSet 2
+ OpDecorate %prevent_dce Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
@@ -32,43 +41,52 @@
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
- %void = OpTypeVoid
- %9 = OpTypeFunction %void
%half = OpTypeFloat 16
%v2half = OpTypeVector %half 2
%mat3v2half = OpTypeMatrix %v2half 3
+%prevent_dce_block = OpTypeStruct %mat3v2half
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+ %void = OpTypeVoid
+ %15 = OpTypeFunction %void
%half_0x1p_0 = OpConstant %half 0x1p+0
- %17 = OpConstantComposite %v2half %half_0x1p_0 %half_0x1p_0
- %18 = OpConstantComposite %mat3v2half %17 %17 %17
+ %20 = OpConstantComposite %v2half %half_0x1p_0 %half_0x1p_0
+ %21 = OpConstantComposite %mat3v2half %20 %20 %20
%_ptr_Function_mat3v2half = OpTypePointer Function %mat3v2half
- %21 = OpConstantNull %mat3v2half
- %22 = OpTypeFunction %v4float
+ %24 = OpConstantNull %mat3v2half
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_mat3v2half = OpTypePointer StorageBuffer %mat3v2half
+ %30 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
-%transpose_d6faec = OpFunction %void None %9
- %12 = OpLabel
- %res = OpVariable %_ptr_Function_mat3v2half Function %21
- OpStore %res %18
+%transpose_d6faec = OpFunction %void None %15
+ %18 = OpLabel
+ %res = OpVariable %_ptr_Function_mat3v2half Function %24
+ OpStore %res %21
+ %28 = OpAccessChain %_ptr_StorageBuffer_mat3v2half %prevent_dce %uint_0
+ %29 = OpLoad %mat3v2half %res
+ OpStore %28 %29
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %22
- %24 = OpLabel
- %25 = OpFunctionCall %void %transpose_d6faec
+%vertex_main_inner = OpFunction %v4float None %30
+ %32 = OpLabel
+ %33 = OpFunctionCall %void %transpose_d6faec
OpReturnValue %5
OpFunctionEnd
-%vertex_main = OpFunction %void None %9
- %27 = OpLabel
- %28 = OpFunctionCall %v4float %vertex_main_inner
- OpStore %value %28
+%vertex_main = OpFunction %void None %15
+ %35 = OpLabel
+ %36 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %36
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
-%fragment_main = OpFunction %void None %9
- %31 = OpLabel
- %32 = OpFunctionCall %void %transpose_d6faec
+%fragment_main = OpFunction %void None %15
+ %39 = OpLabel
+ %40 = OpFunctionCall %void %transpose_d6faec
OpReturn
OpFunctionEnd
-%compute_main = OpFunction %void None %9
- %34 = OpLabel
- %35 = OpFunctionCall %void %transpose_d6faec
+%compute_main = OpFunction %void None %15
+ %42 = OpLabel
+ %43 = OpFunctionCall %void %transpose_d6faec
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/transpose/d6faec.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/transpose/d6faec.wgsl.expected.wgsl
index 2bf363a..671562d 100644
--- a/test/tint/builtins/gen/literal/transpose/d6faec.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/d6faec.wgsl.expected.wgsl
@@ -2,8 +2,11 @@
fn transpose_d6faec() {
var res : mat3x2<f16> = transpose(mat2x3<f16>(1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat3x2<f16>;
+
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
transpose_d6faec();
diff --git a/test/tint/builtins/gen/literal/transpose/d8f8ba.wgsl b/test/tint/builtins/gen/literal/transpose/d8f8ba.wgsl
index 2d38478..15e0379 100644
--- a/test/tint/builtins/gen/literal/transpose/d8f8ba.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/d8f8ba.wgsl
@@ -24,7 +24,9 @@
// fn transpose(mat<3, 4, f32>) -> mat<4, 3, f32>
fn transpose_d8f8ba() {
var res: mat4x3<f32> = transpose(mat3x4<f32>(1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat4x3<f32>;
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
diff --git a/test/tint/builtins/gen/literal/transpose/d8f8ba.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/transpose/d8f8ba.wgsl.expected.dxc.hlsl
index 2c8bd5d..2efae76 100644
--- a/test/tint/builtins/gen/literal/transpose/d8f8ba.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/d8f8ba.wgsl.expected.dxc.hlsl
@@ -1,5 +1,15 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, float4x3 value) {
+ prevent_dce.Store3((offset + 0u), asuint(value[0u]));
+ prevent_dce.Store3((offset + 16u), asuint(value[1u]));
+ prevent_dce.Store3((offset + 32u), asuint(value[2u]));
+ prevent_dce.Store3((offset + 48u), asuint(value[3u]));
+}
+
void transpose_d8f8ba() {
float4x3 res = float4x3((1.0f).xxx, (1.0f).xxx, (1.0f).xxx, (1.0f).xxx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/d8f8ba.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/transpose/d8f8ba.wgsl.expected.fxc.hlsl
index 2c8bd5d..2efae76 100644
--- a/test/tint/builtins/gen/literal/transpose/d8f8ba.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/d8f8ba.wgsl.expected.fxc.hlsl
@@ -1,5 +1,15 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, float4x3 value) {
+ prevent_dce.Store3((offset + 0u), asuint(value[0u]));
+ prevent_dce.Store3((offset + 16u), asuint(value[1u]));
+ prevent_dce.Store3((offset + 32u), asuint(value[2u]));
+ prevent_dce.Store3((offset + 48u), asuint(value[3u]));
+}
+
void transpose_d8f8ba() {
float4x3 res = float4x3((1.0f).xxx, (1.0f).xxx, (1.0f).xxx, (1.0f).xxx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/d8f8ba.wgsl.expected.glsl b/test/tint/builtins/gen/literal/transpose/d8f8ba.wgsl.expected.glsl
index 58df856..afe4b08 100644
--- a/test/tint/builtins/gen/literal/transpose/d8f8ba.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/transpose/d8f8ba.wgsl.expected.glsl
@@ -1,7 +1,19 @@
#version 310 es
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ mat4x3 inner;
+} prevent_dce;
+
+void assign_and_preserve_padding_prevent_dce(mat4x3 value) {
+ prevent_dce.inner[0] = value[0u];
+ prevent_dce.inner[1] = value[1u];
+ prevent_dce.inner[2] = value[2u];
+ prevent_dce.inner[3] = value[3u];
+}
+
void transpose_d8f8ba() {
mat4x3 res = mat4x3(vec3(1.0f), vec3(1.0f), vec3(1.0f), vec3(1.0f));
+ assign_and_preserve_padding_prevent_dce(res);
}
vec4 vertex_main() {
@@ -20,8 +32,20 @@
#version 310 es
precision mediump float;
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ mat4x3 inner;
+} prevent_dce;
+
+void assign_and_preserve_padding_prevent_dce(mat4x3 value) {
+ prevent_dce.inner[0] = value[0u];
+ prevent_dce.inner[1] = value[1u];
+ prevent_dce.inner[2] = value[2u];
+ prevent_dce.inner[3] = value[3u];
+}
+
void transpose_d8f8ba() {
mat4x3 res = mat4x3(vec3(1.0f), vec3(1.0f), vec3(1.0f), vec3(1.0f));
+ assign_and_preserve_padding_prevent_dce(res);
}
void fragment_main() {
@@ -34,8 +58,20 @@
}
#version 310 es
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ mat4x3 inner;
+} prevent_dce;
+
+void assign_and_preserve_padding_prevent_dce(mat4x3 value) {
+ prevent_dce.inner[0] = value[0u];
+ prevent_dce.inner[1] = value[1u];
+ prevent_dce.inner[2] = value[2u];
+ prevent_dce.inner[3] = value[3u];
+}
+
void transpose_d8f8ba() {
mat4x3 res = mat4x3(vec3(1.0f), vec3(1.0f), vec3(1.0f), vec3(1.0f));
+ assign_and_preserve_padding_prevent_dce(res);
}
void compute_main() {
diff --git a/test/tint/builtins/gen/literal/transpose/d8f8ba.wgsl.expected.msl b/test/tint/builtins/gen/literal/transpose/d8f8ba.wgsl.expected.msl
index fa01e0a..e016106 100644
--- a/test/tint/builtins/gen/literal/transpose/d8f8ba.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/transpose/d8f8ba.wgsl.expected.msl
@@ -1,33 +1,59 @@
#include <metal_stdlib>
using namespace metal;
-void transpose_d8f8ba() {
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+struct tint_packed_vec3_f32_array_element {
+ /* 0x0000 */ packed_float3 elements;
+ /* 0x000c */ tint_array<int8_t, 4> tint_pad;
+};
+
+void assign_and_preserve_padding(device tint_array<tint_packed_vec3_f32_array_element, 4>* const dest, float4x3 value) {
+ (*(dest))[0u].elements = packed_float3(value[0u]);
+ (*(dest))[1u].elements = packed_float3(value[1u]);
+ (*(dest))[2u].elements = packed_float3(value[2u]);
+ (*(dest))[3u].elements = packed_float3(value[3u]);
+}
+
+void transpose_d8f8ba(device tint_array<tint_packed_vec3_f32_array_element, 4>* const tint_symbol_1) {
float4x3 res = float4x3(float3(1.0f), float3(1.0f), float3(1.0f), float3(1.0f));
+ assign_and_preserve_padding(tint_symbol_1, res);
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner() {
- transpose_d8f8ba();
+float4 vertex_main_inner(device tint_array<tint_packed_vec3_f32_array_element, 4>* const tint_symbol_2) {
+ transpose_d8f8ba(tint_symbol_2);
return float4(0.0f);
}
-vertex tint_symbol vertex_main() {
- float4 const inner_result = vertex_main_inner();
+vertex tint_symbol vertex_main(device tint_array<tint_packed_vec3_f32_array_element, 4>* tint_symbol_3 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_3);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main() {
- transpose_d8f8ba();
+fragment void fragment_main(device tint_array<tint_packed_vec3_f32_array_element, 4>* tint_symbol_4 [[buffer(0)]]) {
+ transpose_d8f8ba(tint_symbol_4);
return;
}
-kernel void compute_main() {
- transpose_d8f8ba();
+kernel void compute_main(device tint_array<tint_packed_vec3_f32_array_element, 4>* tint_symbol_5 [[buffer(0)]]) {
+ transpose_d8f8ba(tint_symbol_5);
return;
}
diff --git a/test/tint/builtins/gen/literal/transpose/d8f8ba.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/transpose/d8f8ba.wgsl.expected.spvasm
index 61e237e..9884f48 100644
--- a/test/tint/builtins/gen/literal/transpose/d8f8ba.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/transpose/d8f8ba.wgsl.expected.spvasm
@@ -1,68 +1,114 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 34
+; Bound: 63
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
- OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+ OpEntryPoint Vertex %vertex_main "vertex_main" %value_1 %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
OpEntryPoint GLCompute %compute_main "compute_main"
OpExecutionMode %fragment_main OriginUpperLeft
OpExecutionMode %compute_main LocalSize 1 1 1
- OpName %value "value"
+ OpName %value_1 "value_1"
OpName %vertex_point_size "vertex_point_size"
+ OpName %prevent_dce_block "prevent_dce_block"
+ OpMemberName %prevent_dce_block 0 "inner"
+ OpName %prevent_dce "prevent_dce"
+ OpName %assign_and_preserve_padding_prevent_dce "assign_and_preserve_padding_prevent_dce"
+ OpName %value "value"
OpName %transpose_d8f8ba "transpose_d8f8ba"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
OpName %vertex_main "vertex_main"
OpName %fragment_main "fragment_main"
OpName %compute_main "compute_main"
- OpDecorate %value BuiltIn Position
+ OpDecorate %value_1 BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
+ OpDecorate %prevent_dce_block Block
+ OpMemberDecorate %prevent_dce_block 0 Offset 0
+ OpMemberDecorate %prevent_dce_block 0 ColMajor
+ OpMemberDecorate %prevent_dce_block 0 MatrixStride 16
+ OpDecorate %prevent_dce DescriptorSet 2
+ OpDecorate %prevent_dce Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
%5 = OpConstantNull %v4float
- %value = OpVariable %_ptr_Output_v4float Output %5
+ %value_1 = OpVariable %_ptr_Output_v4float Output %5
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
- %void = OpTypeVoid
- %9 = OpTypeFunction %void
%v3float = OpTypeVector %float 3
%mat4v3float = OpTypeMatrix %v3float 4
+%prevent_dce_block = OpTypeStruct %mat4v3float
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+ %void = OpTypeVoid
+ %14 = OpTypeFunction %void %mat4v3float
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+ %int = OpTypeInt 32 1
+ %22 = OpConstantNull %int
+%_ptr_StorageBuffer_v3float = OpTypePointer StorageBuffer %v3float
+ %25 = OpConstantNull %uint
+ %int_1 = OpConstant %int 1
+ %uint_1 = OpConstant %uint 1
+ %int_2 = OpConstant %int 2
+ %uint_2 = OpConstant %uint 2
+ %int_3 = OpConstant %int 3
+ %uint_3 = OpConstant %uint 3
+ %39 = OpTypeFunction %void
%float_1 = OpConstant %float 1
- %16 = OpConstantComposite %v3float %float_1 %float_1 %float_1
- %17 = OpConstantComposite %mat4v3float %16 %16 %16 %16
+ %43 = OpConstantComposite %v3float %float_1 %float_1 %float_1
+ %44 = OpConstantComposite %mat4v3float %43 %43 %43 %43
%_ptr_Function_mat4v3float = OpTypePointer Function %mat4v3float
- %20 = OpConstantNull %mat4v3float
- %21 = OpTypeFunction %v4float
-%transpose_d8f8ba = OpFunction %void None %9
- %12 = OpLabel
- %res = OpVariable %_ptr_Function_mat4v3float Function %20
- OpStore %res %17
+ %47 = OpConstantNull %mat4v3float
+ %50 = OpTypeFunction %v4float
+%assign_and_preserve_padding_prevent_dce = OpFunction %void None %14
+ %value = OpFunctionParameter %mat4v3float
+ %18 = OpLabel
+ %24 = OpAccessChain %_ptr_StorageBuffer_v3float %prevent_dce %uint_0 %22
+ %26 = OpCompositeExtract %v3float %value 0
+ OpStore %24 %26
+ %28 = OpAccessChain %_ptr_StorageBuffer_v3float %prevent_dce %uint_0 %int_1
+ %30 = OpCompositeExtract %v3float %value 1
+ OpStore %28 %30
+ %32 = OpAccessChain %_ptr_StorageBuffer_v3float %prevent_dce %uint_0 %int_2
+ %34 = OpCompositeExtract %v3float %value 2
+ OpStore %32 %34
+ %36 = OpAccessChain %_ptr_StorageBuffer_v3float %prevent_dce %uint_0 %int_3
+ %38 = OpCompositeExtract %v3float %value 3
+ OpStore %36 %38
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %21
- %23 = OpLabel
- %24 = OpFunctionCall %void %transpose_d8f8ba
+%transpose_d8f8ba = OpFunction %void None %39
+ %41 = OpLabel
+ %res = OpVariable %_ptr_Function_mat4v3float Function %47
+ OpStore %res %44
+ %49 = OpLoad %mat4v3float %res
+ %48 = OpFunctionCall %void %assign_and_preserve_padding_prevent_dce %49
+ OpReturn
+ OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %50
+ %52 = OpLabel
+ %53 = OpFunctionCall %void %transpose_d8f8ba
OpReturnValue %5
OpFunctionEnd
-%vertex_main = OpFunction %void None %9
- %26 = OpLabel
- %27 = OpFunctionCall %v4float %vertex_main_inner
- OpStore %value %27
+%vertex_main = OpFunction %void None %39
+ %55 = OpLabel
+ %56 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value_1 %56
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
-%fragment_main = OpFunction %void None %9
- %29 = OpLabel
- %30 = OpFunctionCall %void %transpose_d8f8ba
+%fragment_main = OpFunction %void None %39
+ %58 = OpLabel
+ %59 = OpFunctionCall %void %transpose_d8f8ba
OpReturn
OpFunctionEnd
-%compute_main = OpFunction %void None %9
- %32 = OpLabel
- %33 = OpFunctionCall %void %transpose_d8f8ba
+%compute_main = OpFunction %void None %39
+ %61 = OpLabel
+ %62 = OpFunctionCall %void %transpose_d8f8ba
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/transpose/d8f8ba.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/transpose/d8f8ba.wgsl.expected.wgsl
index c08b152..cca9ee2 100644
--- a/test/tint/builtins/gen/literal/transpose/d8f8ba.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/d8f8ba.wgsl.expected.wgsl
@@ -1,7 +1,10 @@
fn transpose_d8f8ba() {
var res : mat4x3<f32> = transpose(mat3x4<f32>(1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat4x3<f32>;
+
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
transpose_d8f8ba();
diff --git a/test/tint/builtins/gen/literal/transpose/dc671a.wgsl b/test/tint/builtins/gen/literal/transpose/dc671a.wgsl
index 7ce8237..f2feaa7 100644
--- a/test/tint/builtins/gen/literal/transpose/dc671a.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/dc671a.wgsl
@@ -25,7 +25,6 @@
fn transpose_dc671a() {
var res = transpose(mat4x4(1., 1., 1., 1., 1., 1., 1., 1., 1., 1., 1., 1., 1., 1., 1., 1.));
}
-
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
transpose_dc671a();
diff --git a/test/tint/builtins/gen/literal/transpose/ed4bdc.wgsl b/test/tint/builtins/gen/literal/transpose/ed4bdc.wgsl
index 101c03d..567ae33 100644
--- a/test/tint/builtins/gen/literal/transpose/ed4bdc.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/ed4bdc.wgsl
@@ -24,7 +24,9 @@
// fn transpose(mat<3, 2, f32>) -> mat<2, 3, f32>
fn transpose_ed4bdc() {
var res: mat2x3<f32> = transpose(mat3x2<f32>(1.f, 1.f, 1.f, 1.f, 1.f, 1.f));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat2x3<f32>;
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
diff --git a/test/tint/builtins/gen/literal/transpose/ed4bdc.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/transpose/ed4bdc.wgsl.expected.dxc.hlsl
index 47e3175..ffef122 100644
--- a/test/tint/builtins/gen/literal/transpose/ed4bdc.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/ed4bdc.wgsl.expected.dxc.hlsl
@@ -1,5 +1,13 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, float2x3 value) {
+ prevent_dce.Store3((offset + 0u), asuint(value[0u]));
+ prevent_dce.Store3((offset + 16u), asuint(value[1u]));
+}
+
void transpose_ed4bdc() {
float2x3 res = float2x3((1.0f).xxx, (1.0f).xxx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/ed4bdc.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/transpose/ed4bdc.wgsl.expected.fxc.hlsl
index 47e3175..ffef122 100644
--- a/test/tint/builtins/gen/literal/transpose/ed4bdc.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/ed4bdc.wgsl.expected.fxc.hlsl
@@ -1,5 +1,13 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, float2x3 value) {
+ prevent_dce.Store3((offset + 0u), asuint(value[0u]));
+ prevent_dce.Store3((offset + 16u), asuint(value[1u]));
+}
+
void transpose_ed4bdc() {
float2x3 res = float2x3((1.0f).xxx, (1.0f).xxx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/ed4bdc.wgsl.expected.glsl b/test/tint/builtins/gen/literal/transpose/ed4bdc.wgsl.expected.glsl
index 3179df8..f3e9088 100644
--- a/test/tint/builtins/gen/literal/transpose/ed4bdc.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/transpose/ed4bdc.wgsl.expected.glsl
@@ -1,7 +1,17 @@
#version 310 es
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ mat2x3 inner;
+} prevent_dce;
+
+void assign_and_preserve_padding_prevent_dce(mat2x3 value) {
+ prevent_dce.inner[0] = value[0u];
+ prevent_dce.inner[1] = value[1u];
+}
+
void transpose_ed4bdc() {
mat2x3 res = mat2x3(vec3(1.0f), vec3(1.0f));
+ assign_and_preserve_padding_prevent_dce(res);
}
vec4 vertex_main() {
@@ -20,8 +30,18 @@
#version 310 es
precision mediump float;
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ mat2x3 inner;
+} prevent_dce;
+
+void assign_and_preserve_padding_prevent_dce(mat2x3 value) {
+ prevent_dce.inner[0] = value[0u];
+ prevent_dce.inner[1] = value[1u];
+}
+
void transpose_ed4bdc() {
mat2x3 res = mat2x3(vec3(1.0f), vec3(1.0f));
+ assign_and_preserve_padding_prevent_dce(res);
}
void fragment_main() {
@@ -34,8 +54,18 @@
}
#version 310 es
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ mat2x3 inner;
+} prevent_dce;
+
+void assign_and_preserve_padding_prevent_dce(mat2x3 value) {
+ prevent_dce.inner[0] = value[0u];
+ prevent_dce.inner[1] = value[1u];
+}
+
void transpose_ed4bdc() {
mat2x3 res = mat2x3(vec3(1.0f), vec3(1.0f));
+ assign_and_preserve_padding_prevent_dce(res);
}
void compute_main() {
diff --git a/test/tint/builtins/gen/literal/transpose/ed4bdc.wgsl.expected.msl b/test/tint/builtins/gen/literal/transpose/ed4bdc.wgsl.expected.msl
index 1f9c1a8..1be5ac3 100644
--- a/test/tint/builtins/gen/literal/transpose/ed4bdc.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/transpose/ed4bdc.wgsl.expected.msl
@@ -1,33 +1,57 @@
#include <metal_stdlib>
using namespace metal;
-void transpose_ed4bdc() {
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+struct tint_packed_vec3_f32_array_element {
+ /* 0x0000 */ packed_float3 elements;
+ /* 0x000c */ tint_array<int8_t, 4> tint_pad;
+};
+
+void assign_and_preserve_padding(device tint_array<tint_packed_vec3_f32_array_element, 2>* const dest, float2x3 value) {
+ (*(dest))[0u].elements = packed_float3(value[0u]);
+ (*(dest))[1u].elements = packed_float3(value[1u]);
+}
+
+void transpose_ed4bdc(device tint_array<tint_packed_vec3_f32_array_element, 2>* const tint_symbol_1) {
float2x3 res = float2x3(float3(1.0f), float3(1.0f));
+ assign_and_preserve_padding(tint_symbol_1, res);
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner() {
- transpose_ed4bdc();
+float4 vertex_main_inner(device tint_array<tint_packed_vec3_f32_array_element, 2>* const tint_symbol_2) {
+ transpose_ed4bdc(tint_symbol_2);
return float4(0.0f);
}
-vertex tint_symbol vertex_main() {
- float4 const inner_result = vertex_main_inner();
+vertex tint_symbol vertex_main(device tint_array<tint_packed_vec3_f32_array_element, 2>* tint_symbol_3 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_3);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main() {
- transpose_ed4bdc();
+fragment void fragment_main(device tint_array<tint_packed_vec3_f32_array_element, 2>* tint_symbol_4 [[buffer(0)]]) {
+ transpose_ed4bdc(tint_symbol_4);
return;
}
-kernel void compute_main() {
- transpose_ed4bdc();
+kernel void compute_main(device tint_array<tint_packed_vec3_f32_array_element, 2>* tint_symbol_5 [[buffer(0)]]) {
+ transpose_ed4bdc(tint_symbol_5);
return;
}
diff --git a/test/tint/builtins/gen/literal/transpose/ed4bdc.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/transpose/ed4bdc.wgsl.expected.spvasm
index 434bdee..fa85b25 100644
--- a/test/tint/builtins/gen/literal/transpose/ed4bdc.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/transpose/ed4bdc.wgsl.expected.spvasm
@@ -1,68 +1,104 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 34
+; Bound: 55
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
- OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+ OpEntryPoint Vertex %vertex_main "vertex_main" %value_1 %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
OpEntryPoint GLCompute %compute_main "compute_main"
OpExecutionMode %fragment_main OriginUpperLeft
OpExecutionMode %compute_main LocalSize 1 1 1
- OpName %value "value"
+ OpName %value_1 "value_1"
OpName %vertex_point_size "vertex_point_size"
+ OpName %prevent_dce_block "prevent_dce_block"
+ OpMemberName %prevent_dce_block 0 "inner"
+ OpName %prevent_dce "prevent_dce"
+ OpName %assign_and_preserve_padding_prevent_dce "assign_and_preserve_padding_prevent_dce"
+ OpName %value "value"
OpName %transpose_ed4bdc "transpose_ed4bdc"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
OpName %vertex_main "vertex_main"
OpName %fragment_main "fragment_main"
OpName %compute_main "compute_main"
- OpDecorate %value BuiltIn Position
+ OpDecorate %value_1 BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
+ OpDecorate %prevent_dce_block Block
+ OpMemberDecorate %prevent_dce_block 0 Offset 0
+ OpMemberDecorate %prevent_dce_block 0 ColMajor
+ OpMemberDecorate %prevent_dce_block 0 MatrixStride 16
+ OpDecorate %prevent_dce DescriptorSet 2
+ OpDecorate %prevent_dce Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
%5 = OpConstantNull %v4float
- %value = OpVariable %_ptr_Output_v4float Output %5
+ %value_1 = OpVariable %_ptr_Output_v4float Output %5
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
- %void = OpTypeVoid
- %9 = OpTypeFunction %void
%v3float = OpTypeVector %float 3
%mat2v3float = OpTypeMatrix %v3float 2
+%prevent_dce_block = OpTypeStruct %mat2v3float
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+ %void = OpTypeVoid
+ %14 = OpTypeFunction %void %mat2v3float
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+ %int = OpTypeInt 32 1
+ %22 = OpConstantNull %int
+%_ptr_StorageBuffer_v3float = OpTypePointer StorageBuffer %v3float
+ %25 = OpConstantNull %uint
+ %int_1 = OpConstant %int 1
+ %uint_1 = OpConstant %uint 1
+ %31 = OpTypeFunction %void
%float_1 = OpConstant %float 1
- %16 = OpConstantComposite %v3float %float_1 %float_1 %float_1
- %17 = OpConstantComposite %mat2v3float %16 %16
+ %35 = OpConstantComposite %v3float %float_1 %float_1 %float_1
+ %36 = OpConstantComposite %mat2v3float %35 %35
%_ptr_Function_mat2v3float = OpTypePointer Function %mat2v3float
- %20 = OpConstantNull %mat2v3float
- %21 = OpTypeFunction %v4float
-%transpose_ed4bdc = OpFunction %void None %9
- %12 = OpLabel
- %res = OpVariable %_ptr_Function_mat2v3float Function %20
- OpStore %res %17
+ %39 = OpConstantNull %mat2v3float
+ %42 = OpTypeFunction %v4float
+%assign_and_preserve_padding_prevent_dce = OpFunction %void None %14
+ %value = OpFunctionParameter %mat2v3float
+ %18 = OpLabel
+ %24 = OpAccessChain %_ptr_StorageBuffer_v3float %prevent_dce %uint_0 %22
+ %26 = OpCompositeExtract %v3float %value 0
+ OpStore %24 %26
+ %28 = OpAccessChain %_ptr_StorageBuffer_v3float %prevent_dce %uint_0 %int_1
+ %30 = OpCompositeExtract %v3float %value 1
+ OpStore %28 %30
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %21
- %23 = OpLabel
- %24 = OpFunctionCall %void %transpose_ed4bdc
+%transpose_ed4bdc = OpFunction %void None %31
+ %33 = OpLabel
+ %res = OpVariable %_ptr_Function_mat2v3float Function %39
+ OpStore %res %36
+ %41 = OpLoad %mat2v3float %res
+ %40 = OpFunctionCall %void %assign_and_preserve_padding_prevent_dce %41
+ OpReturn
+ OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %42
+ %44 = OpLabel
+ %45 = OpFunctionCall %void %transpose_ed4bdc
OpReturnValue %5
OpFunctionEnd
-%vertex_main = OpFunction %void None %9
- %26 = OpLabel
- %27 = OpFunctionCall %v4float %vertex_main_inner
- OpStore %value %27
+%vertex_main = OpFunction %void None %31
+ %47 = OpLabel
+ %48 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value_1 %48
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
-%fragment_main = OpFunction %void None %9
- %29 = OpLabel
- %30 = OpFunctionCall %void %transpose_ed4bdc
+%fragment_main = OpFunction %void None %31
+ %50 = OpLabel
+ %51 = OpFunctionCall %void %transpose_ed4bdc
OpReturn
OpFunctionEnd
-%compute_main = OpFunction %void None %9
- %32 = OpLabel
- %33 = OpFunctionCall %void %transpose_ed4bdc
+%compute_main = OpFunction %void None %31
+ %53 = OpLabel
+ %54 = OpFunctionCall %void %transpose_ed4bdc
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/transpose/ed4bdc.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/transpose/ed4bdc.wgsl.expected.wgsl
index f8d73e8..f9da29e 100644
--- a/test/tint/builtins/gen/literal/transpose/ed4bdc.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/ed4bdc.wgsl.expected.wgsl
@@ -1,7 +1,10 @@
fn transpose_ed4bdc() {
var res : mat2x3<f32> = transpose(mat3x2<f32>(1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat2x3<f32>;
+
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
transpose_ed4bdc();
diff --git a/test/tint/builtins/gen/literal/transpose/faeb05.wgsl b/test/tint/builtins/gen/literal/transpose/faeb05.wgsl
index 71346b0..0551fa8 100644
--- a/test/tint/builtins/gen/literal/transpose/faeb05.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/faeb05.wgsl
@@ -26,7 +26,9 @@
// fn transpose(mat<2, 4, f16>) -> mat<4, 2, f16>
fn transpose_faeb05() {
var res: mat4x2<f16> = transpose(mat2x4<f16>(1.h, 1.h, 1.h, 1.h, 1.h, 1.h, 1.h, 1.h));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat4x2<f16>;
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
diff --git a/test/tint/builtins/gen/literal/transpose/faeb05.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/transpose/faeb05.wgsl.expected.dxc.hlsl
index 9be9430..2f55e37 100644
--- a/test/tint/builtins/gen/literal/transpose/faeb05.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/faeb05.wgsl.expected.dxc.hlsl
@@ -1,5 +1,15 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, matrix<float16_t, 4, 2> value) {
+ prevent_dce.Store<vector<float16_t, 2> >((offset + 0u), value[0u]);
+ prevent_dce.Store<vector<float16_t, 2> >((offset + 4u), value[1u]);
+ prevent_dce.Store<vector<float16_t, 2> >((offset + 8u), value[2u]);
+ prevent_dce.Store<vector<float16_t, 2> >((offset + 12u), value[3u]);
+}
+
void transpose_faeb05() {
matrix<float16_t, 4, 2> res = matrix<float16_t, 4, 2>((float16_t(1.0h)).xx, (float16_t(1.0h)).xx, (float16_t(1.0h)).xx, (float16_t(1.0h)).xx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/faeb05.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/transpose/faeb05.wgsl.expected.fxc.hlsl
index 32820e1..ab5a3c4 100644
--- a/test/tint/builtins/gen/literal/transpose/faeb05.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/transpose/faeb05.wgsl.expected.fxc.hlsl
@@ -1,7 +1,17 @@
SKIP: FAILED
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
+void prevent_dce_store(uint offset, matrix<float16_t, 4, 2> value) {
+ prevent_dce.Store<vector<float16_t, 2> >((offset + 0u), value[0u]);
+ prevent_dce.Store<vector<float16_t, 2> >((offset + 4u), value[1u]);
+ prevent_dce.Store<vector<float16_t, 2> >((offset + 8u), value[2u]);
+ prevent_dce.Store<vector<float16_t, 2> >((offset + 12u), value[3u]);
+}
+
void transpose_faeb05() {
matrix<float16_t, 4, 2> res = matrix<float16_t, 4, 2>((float16_t(1.0h)).xx, (float16_t(1.0h)).xx, (float16_t(1.0h)).xx, (float16_t(1.0h)).xx);
+ prevent_dce_store(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/transpose/faeb05.wgsl.expected.glsl b/test/tint/builtins/gen/literal/transpose/faeb05.wgsl.expected.glsl
index 07ae2dd..6725813 100644
--- a/test/tint/builtins/gen/literal/transpose/faeb05.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/transpose/faeb05.wgsl.expected.glsl
@@ -1,8 +1,13 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16mat4x2 inner;
+} prevent_dce;
+
void transpose_faeb05() {
f16mat4x2 res = f16mat4x2(f16vec2(1.0hf), f16vec2(1.0hf), f16vec2(1.0hf), f16vec2(1.0hf));
+ prevent_dce.inner = res;
}
vec4 vertex_main() {
@@ -22,8 +27,13 @@
#extension GL_AMD_gpu_shader_half_float : require
precision mediump float;
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16mat4x2 inner;
+} prevent_dce;
+
void transpose_faeb05() {
f16mat4x2 res = f16mat4x2(f16vec2(1.0hf), f16vec2(1.0hf), f16vec2(1.0hf), f16vec2(1.0hf));
+ prevent_dce.inner = res;
}
void fragment_main() {
@@ -37,8 +47,13 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16mat4x2 inner;
+} prevent_dce;
+
void transpose_faeb05() {
f16mat4x2 res = f16mat4x2(f16vec2(1.0hf), f16vec2(1.0hf), f16vec2(1.0hf), f16vec2(1.0hf));
+ prevent_dce.inner = res;
}
void compute_main() {
diff --git a/test/tint/builtins/gen/literal/transpose/faeb05.wgsl.expected.msl b/test/tint/builtins/gen/literal/transpose/faeb05.wgsl.expected.msl
index 845b1bc..3aa5ea3 100644
--- a/test/tint/builtins/gen/literal/transpose/faeb05.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/transpose/faeb05.wgsl.expected.msl
@@ -1,33 +1,34 @@
#include <metal_stdlib>
using namespace metal;
-void transpose_faeb05() {
+void transpose_faeb05(device half4x2* const tint_symbol_1) {
half4x2 res = half4x2(half2(1.0h), half2(1.0h), half2(1.0h), half2(1.0h));
+ *(tint_symbol_1) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner() {
- transpose_faeb05();
+float4 vertex_main_inner(device half4x2* const tint_symbol_2) {
+ transpose_faeb05(tint_symbol_2);
return float4(0.0f);
}
-vertex tint_symbol vertex_main() {
- float4 const inner_result = vertex_main_inner();
+vertex tint_symbol vertex_main(device half4x2* tint_symbol_3 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_3);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main() {
- transpose_faeb05();
+fragment void fragment_main(device half4x2* tint_symbol_4 [[buffer(0)]]) {
+ transpose_faeb05(tint_symbol_4);
return;
}
-kernel void compute_main() {
- transpose_faeb05();
+kernel void compute_main(device half4x2* tint_symbol_5 [[buffer(0)]]) {
+ transpose_faeb05(tint_symbol_5);
return;
}
diff --git a/test/tint/builtins/gen/literal/transpose/faeb05.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/transpose/faeb05.wgsl.expected.spvasm
index 0d70272..e18ed7b 100644
--- a/test/tint/builtins/gen/literal/transpose/faeb05.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/transpose/faeb05.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 36
+; Bound: 44
; Schema: 0
OpCapability Shader
OpCapability Float16
@@ -16,6 +16,9 @@
OpExecutionMode %compute_main LocalSize 1 1 1
OpName %value "value"
OpName %vertex_point_size "vertex_point_size"
+ OpName %prevent_dce_block "prevent_dce_block"
+ OpMemberName %prevent_dce_block 0 "inner"
+ OpName %prevent_dce "prevent_dce"
OpName %transpose_faeb05 "transpose_faeb05"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
@@ -24,6 +27,12 @@
OpName %compute_main "compute_main"
OpDecorate %value BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
+ OpDecorate %prevent_dce_block Block
+ OpMemberDecorate %prevent_dce_block 0 Offset 0
+ OpMemberDecorate %prevent_dce_block 0 ColMajor
+ OpMemberDecorate %prevent_dce_block 0 MatrixStride 4
+ OpDecorate %prevent_dce DescriptorSet 2
+ OpDecorate %prevent_dce Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
@@ -32,43 +41,52 @@
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
- %void = OpTypeVoid
- %9 = OpTypeFunction %void
%half = OpTypeFloat 16
%v2half = OpTypeVector %half 2
%mat4v2half = OpTypeMatrix %v2half 4
+%prevent_dce_block = OpTypeStruct %mat4v2half
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+ %void = OpTypeVoid
+ %15 = OpTypeFunction %void
%half_0x1p_0 = OpConstant %half 0x1p+0
- %17 = OpConstantComposite %v2half %half_0x1p_0 %half_0x1p_0
- %18 = OpConstantComposite %mat4v2half %17 %17 %17 %17
+ %20 = OpConstantComposite %v2half %half_0x1p_0 %half_0x1p_0
+ %21 = OpConstantComposite %mat4v2half %20 %20 %20 %20
%_ptr_Function_mat4v2half = OpTypePointer Function %mat4v2half
- %21 = OpConstantNull %mat4v2half
- %22 = OpTypeFunction %v4float
+ %24 = OpConstantNull %mat4v2half
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_mat4v2half = OpTypePointer StorageBuffer %mat4v2half
+ %30 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
-%transpose_faeb05 = OpFunction %void None %9
- %12 = OpLabel
- %res = OpVariable %_ptr_Function_mat4v2half Function %21
- OpStore %res %18
+%transpose_faeb05 = OpFunction %void None %15
+ %18 = OpLabel
+ %res = OpVariable %_ptr_Function_mat4v2half Function %24
+ OpStore %res %21
+ %28 = OpAccessChain %_ptr_StorageBuffer_mat4v2half %prevent_dce %uint_0
+ %29 = OpLoad %mat4v2half %res
+ OpStore %28 %29
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %22
- %24 = OpLabel
- %25 = OpFunctionCall %void %transpose_faeb05
+%vertex_main_inner = OpFunction %v4float None %30
+ %32 = OpLabel
+ %33 = OpFunctionCall %void %transpose_faeb05
OpReturnValue %5
OpFunctionEnd
-%vertex_main = OpFunction %void None %9
- %27 = OpLabel
- %28 = OpFunctionCall %v4float %vertex_main_inner
- OpStore %value %28
+%vertex_main = OpFunction %void None %15
+ %35 = OpLabel
+ %36 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %36
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
-%fragment_main = OpFunction %void None %9
- %31 = OpLabel
- %32 = OpFunctionCall %void %transpose_faeb05
+%fragment_main = OpFunction %void None %15
+ %39 = OpLabel
+ %40 = OpFunctionCall %void %transpose_faeb05
OpReturn
OpFunctionEnd
-%compute_main = OpFunction %void None %9
- %34 = OpLabel
- %35 = OpFunctionCall %void %transpose_faeb05
+%compute_main = OpFunction %void None %15
+ %42 = OpLabel
+ %43 = OpFunctionCall %void %transpose_faeb05
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/transpose/faeb05.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/transpose/faeb05.wgsl.expected.wgsl
index 364abc3..7cff494 100644
--- a/test/tint/builtins/gen/literal/transpose/faeb05.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/transpose/faeb05.wgsl.expected.wgsl
@@ -2,8 +2,11 @@
fn transpose_faeb05() {
var res : mat4x2<f16> = transpose(mat2x4<f16>(1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h, 1.0h));
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : mat4x2<f16>;
+
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
transpose_faeb05();
