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/var/round/106c0b.wgsl b/test/tint/builtins/gen/var/round/106c0b.wgsl
index 7c77f82..f4d5e6a 100644
--- a/test/tint/builtins/gen/var/round/106c0b.wgsl
+++ b/test/tint/builtins/gen/var/round/106c0b.wgsl
@@ -25,7 +25,9 @@
fn round_106c0b() {
var arg_0 = vec4<f32>(3.5f);
var res: vec4<f32> = round(arg_0);
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<f32>;
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
diff --git a/test/tint/builtins/gen/var/round/106c0b.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/round/106c0b.wgsl.expected.dxc.hlsl
index c0d5f98..37a7885 100644
--- a/test/tint/builtins/gen/var/round/106c0b.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/var/round/106c0b.wgsl.expected.dxc.hlsl
@@ -1,6 +1,9 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
void round_106c0b() {
float4 arg_0 = (3.5f).xxxx;
float4 res = round(arg_0);
+ prevent_dce.Store4(0u, asuint(res));
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/var/round/106c0b.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/round/106c0b.wgsl.expected.fxc.hlsl
index c0d5f98..37a7885 100644
--- a/test/tint/builtins/gen/var/round/106c0b.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/var/round/106c0b.wgsl.expected.fxc.hlsl
@@ -1,6 +1,9 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
void round_106c0b() {
float4 arg_0 = (3.5f).xxxx;
float4 res = round(arg_0);
+ prevent_dce.Store4(0u, asuint(res));
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/var/round/106c0b.wgsl.expected.glsl b/test/tint/builtins/gen/var/round/106c0b.wgsl.expected.glsl
index 84edebf..78273ec 100644
--- a/test/tint/builtins/gen/var/round/106c0b.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/round/106c0b.wgsl.expected.glsl
@@ -1,8 +1,13 @@
#version 310 es
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ vec4 inner;
+} prevent_dce;
+
void round_106c0b() {
vec4 arg_0 = vec4(3.5f);
vec4 res = round(arg_0);
+ prevent_dce.inner = res;
}
vec4 vertex_main() {
@@ -21,9 +26,14 @@
#version 310 es
precision mediump float;
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ vec4 inner;
+} prevent_dce;
+
void round_106c0b() {
vec4 arg_0 = vec4(3.5f);
vec4 res = round(arg_0);
+ prevent_dce.inner = res;
}
void fragment_main() {
@@ -36,9 +46,14 @@
}
#version 310 es
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ vec4 inner;
+} prevent_dce;
+
void round_106c0b() {
vec4 arg_0 = vec4(3.5f);
vec4 res = round(arg_0);
+ prevent_dce.inner = res;
}
void compute_main() {
diff --git a/test/tint/builtins/gen/var/round/106c0b.wgsl.expected.msl b/test/tint/builtins/gen/var/round/106c0b.wgsl.expected.msl
index e35922e..2ceb8d1 100644
--- a/test/tint/builtins/gen/var/round/106c0b.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/round/106c0b.wgsl.expected.msl
@@ -1,34 +1,35 @@
#include <metal_stdlib>
using namespace metal;
-void round_106c0b() {
+void round_106c0b(device float4* const tint_symbol_1) {
float4 arg_0 = float4(3.5f);
float4 res = rint(arg_0);
+ *(tint_symbol_1) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner() {
- round_106c0b();
+float4 vertex_main_inner(device float4* const tint_symbol_2) {
+ round_106c0b(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 float4* 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() {
- round_106c0b();
+fragment void fragment_main(device float4* tint_symbol_4 [[buffer(0)]]) {
+ round_106c0b(tint_symbol_4);
return;
}
-kernel void compute_main() {
- round_106c0b();
+kernel void compute_main(device float4* tint_symbol_5 [[buffer(0)]]) {
+ round_106c0b(tint_symbol_5);
return;
}
diff --git a/test/tint/builtins/gen/var/round/106c0b.wgsl.expected.spvasm b/test/tint/builtins/gen/var/round/106c0b.wgsl.expected.spvasm
index ce01c78..9fa0d67 100644
--- a/test/tint/builtins/gen/var/round/106c0b.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/round/106c0b.wgsl.expected.spvasm
@@ -1,10 +1,10 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 35
+; Bound: 43
; Schema: 0
OpCapability Shader
- %18 = OpExtInstImport "GLSL.std.450"
+ %21 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -13,6 +13,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 %round_106c0b "round_106c0b"
OpName %arg_0 "arg_0"
OpName %res "res"
@@ -22,6 +25,10 @@
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
+ OpDecorate %prevent_dce DescriptorSet 2
+ OpDecorate %prevent_dce Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
@@ -30,42 +37,51 @@
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+%prevent_dce_block = OpTypeStruct %v4float
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
%void = OpTypeVoid
- %9 = OpTypeFunction %void
+ %12 = OpTypeFunction %void
%float_3_5 = OpConstant %float 3.5
- %14 = OpConstantComposite %v4float %float_3_5 %float_3_5 %float_3_5 %float_3_5
+ %17 = OpConstantComposite %v4float %float_3_5 %float_3_5 %float_3_5 %float_3_5
%_ptr_Function_v4float = OpTypePointer Function %v4float
- %21 = OpTypeFunction %v4float
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v4float = OpTypePointer StorageBuffer %v4float
+ %29 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
-%round_106c0b = OpFunction %void None %9
- %12 = OpLabel
+%round_106c0b = OpFunction %void None %12
+ %15 = OpLabel
%arg_0 = OpVariable %_ptr_Function_v4float Function %5
%res = OpVariable %_ptr_Function_v4float Function %5
- OpStore %arg_0 %14
- %19 = OpLoad %v4float %arg_0
- %17 = OpExtInst %v4float %18 RoundEven %19
- OpStore %res %17
+ OpStore %arg_0 %17
+ %22 = OpLoad %v4float %arg_0
+ %20 = OpExtInst %v4float %21 RoundEven %22
+ OpStore %res %20
+ %27 = OpAccessChain %_ptr_StorageBuffer_v4float %prevent_dce %uint_0
+ %28 = OpLoad %v4float %res
+ OpStore %27 %28
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %21
- %23 = OpLabel
- %24 = OpFunctionCall %void %round_106c0b
+%vertex_main_inner = OpFunction %v4float None %29
+ %31 = OpLabel
+ %32 = OpFunctionCall %void %round_106c0b
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 %12
+ %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
- %30 = OpLabel
- %31 = OpFunctionCall %void %round_106c0b
+%fragment_main = OpFunction %void None %12
+ %38 = OpLabel
+ %39 = OpFunctionCall %void %round_106c0b
OpReturn
OpFunctionEnd
-%compute_main = OpFunction %void None %9
- %33 = OpLabel
- %34 = OpFunctionCall %void %round_106c0b
+%compute_main = OpFunction %void None %12
+ %41 = OpLabel
+ %42 = OpFunctionCall %void %round_106c0b
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/round/106c0b.wgsl.expected.wgsl b/test/tint/builtins/gen/var/round/106c0b.wgsl.expected.wgsl
index 129028b..e1b09dc 100644
--- a/test/tint/builtins/gen/var/round/106c0b.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/var/round/106c0b.wgsl.expected.wgsl
@@ -1,8 +1,11 @@
fn round_106c0b() {
var arg_0 = vec4<f32>(3.5f);
var res : vec4<f32> = round(arg_0);
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<f32>;
+
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
round_106c0b();
diff --git a/test/tint/builtins/gen/var/round/184d5a.wgsl b/test/tint/builtins/gen/var/round/184d5a.wgsl
index 175eca3..fc32550 100644
--- a/test/tint/builtins/gen/var/round/184d5a.wgsl
+++ b/test/tint/builtins/gen/var/round/184d5a.wgsl
@@ -26,7 +26,6 @@
const arg_0 = vec4(3.5);
var res = round(arg_0);
}
-
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
round_184d5a();
diff --git a/test/tint/builtins/gen/var/round/1c7897.wgsl b/test/tint/builtins/gen/var/round/1c7897.wgsl
index 9eac08e..97fa10e 100644
--- a/test/tint/builtins/gen/var/round/1c7897.wgsl
+++ b/test/tint/builtins/gen/var/round/1c7897.wgsl
@@ -25,7 +25,9 @@
fn round_1c7897() {
var arg_0 = vec3<f32>(3.5f);
var res: vec3<f32> = round(arg_0);
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec3<f32>;
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
diff --git a/test/tint/builtins/gen/var/round/1c7897.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/round/1c7897.wgsl.expected.dxc.hlsl
index b41ad5d..dcca5dc 100644
--- a/test/tint/builtins/gen/var/round/1c7897.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/var/round/1c7897.wgsl.expected.dxc.hlsl
@@ -1,6 +1,9 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
void round_1c7897() {
float3 arg_0 = (3.5f).xxx;
float3 res = round(arg_0);
+ prevent_dce.Store3(0u, asuint(res));
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/var/round/1c7897.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/round/1c7897.wgsl.expected.fxc.hlsl
index b41ad5d..dcca5dc 100644
--- a/test/tint/builtins/gen/var/round/1c7897.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/var/round/1c7897.wgsl.expected.fxc.hlsl
@@ -1,6 +1,9 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
void round_1c7897() {
float3 arg_0 = (3.5f).xxx;
float3 res = round(arg_0);
+ prevent_dce.Store3(0u, asuint(res));
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/var/round/1c7897.wgsl.expected.glsl b/test/tint/builtins/gen/var/round/1c7897.wgsl.expected.glsl
index 1a4453d..182cce5 100644
--- a/test/tint/builtins/gen/var/round/1c7897.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/round/1c7897.wgsl.expected.glsl
@@ -1,8 +1,13 @@
#version 310 es
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ vec3 inner;
+} prevent_dce;
+
void round_1c7897() {
vec3 arg_0 = vec3(3.5f);
vec3 res = round(arg_0);
+ prevent_dce.inner = res;
}
vec4 vertex_main() {
@@ -21,9 +26,14 @@
#version 310 es
precision mediump float;
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ vec3 inner;
+} prevent_dce;
+
void round_1c7897() {
vec3 arg_0 = vec3(3.5f);
vec3 res = round(arg_0);
+ prevent_dce.inner = res;
}
void fragment_main() {
@@ -36,9 +46,14 @@
}
#version 310 es
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ vec3 inner;
+} prevent_dce;
+
void round_1c7897() {
vec3 arg_0 = vec3(3.5f);
vec3 res = round(arg_0);
+ prevent_dce.inner = res;
}
void compute_main() {
diff --git a/test/tint/builtins/gen/var/round/1c7897.wgsl.expected.msl b/test/tint/builtins/gen/var/round/1c7897.wgsl.expected.msl
index f7efbda..f041dd1 100644
--- a/test/tint/builtins/gen/var/round/1c7897.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/round/1c7897.wgsl.expected.msl
@@ -1,34 +1,35 @@
#include <metal_stdlib>
using namespace metal;
-void round_1c7897() {
+void round_1c7897(device packed_float3* const tint_symbol_1) {
float3 arg_0 = float3(3.5f);
float3 res = rint(arg_0);
+ *(tint_symbol_1) = packed_float3(res);
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner() {
- round_1c7897();
+float4 vertex_main_inner(device packed_float3* const tint_symbol_2) {
+ round_1c7897(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 packed_float3* 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() {
- round_1c7897();
+fragment void fragment_main(device packed_float3* tint_symbol_4 [[buffer(0)]]) {
+ round_1c7897(tint_symbol_4);
return;
}
-kernel void compute_main() {
- round_1c7897();
+kernel void compute_main(device packed_float3* tint_symbol_5 [[buffer(0)]]) {
+ round_1c7897(tint_symbol_5);
return;
}
diff --git a/test/tint/builtins/gen/var/round/1c7897.wgsl.expected.spvasm b/test/tint/builtins/gen/var/round/1c7897.wgsl.expected.spvasm
index 02c1b23..101f475 100644
--- a/test/tint/builtins/gen/var/round/1c7897.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/round/1c7897.wgsl.expected.spvasm
@@ -1,10 +1,10 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 37
+; Bound: 45
; Schema: 0
OpCapability Shader
- %20 = OpExtInstImport "GLSL.std.450"
+ %23 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -13,6 +13,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 %round_1c7897 "round_1c7897"
OpName %arg_0 "arg_0"
OpName %res "res"
@@ -22,6 +25,10 @@
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
+ OpDecorate %prevent_dce DescriptorSet 2
+ OpDecorate %prevent_dce Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
@@ -30,44 +37,53 @@
%_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
+%prevent_dce_block = OpTypeStruct %v3float
+%_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_3_5 = OpConstant %float 3.5
- %15 = OpConstantComposite %v3float %float_3_5 %float_3_5 %float_3_5
+ %18 = OpConstantComposite %v3float %float_3_5 %float_3_5 %float_3_5
%_ptr_Function_v3float = OpTypePointer Function %v3float
- %18 = OpConstantNull %v3float
- %23 = OpTypeFunction %v4float
+ %21 = OpConstantNull %v3float
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v3float = OpTypePointer StorageBuffer %v3float
+ %31 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
-%round_1c7897 = OpFunction %void None %9
- %12 = OpLabel
- %arg_0 = OpVariable %_ptr_Function_v3float Function %18
- %res = OpVariable %_ptr_Function_v3float Function %18
- OpStore %arg_0 %15
- %21 = OpLoad %v3float %arg_0
- %19 = OpExtInst %v3float %20 RoundEven %21
- OpStore %res %19
+%round_1c7897 = OpFunction %void None %13
+ %16 = OpLabel
+ %arg_0 = OpVariable %_ptr_Function_v3float Function %21
+ %res = OpVariable %_ptr_Function_v3float Function %21
+ OpStore %arg_0 %18
+ %24 = OpLoad %v3float %arg_0
+ %22 = OpExtInst %v3float %23 RoundEven %24
+ OpStore %res %22
+ %29 = OpAccessChain %_ptr_StorageBuffer_v3float %prevent_dce %uint_0
+ %30 = OpLoad %v3float %res
+ OpStore %29 %30
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %23
- %25 = OpLabel
- %26 = OpFunctionCall %void %round_1c7897
+%vertex_main_inner = OpFunction %v4float None %31
+ %33 = OpLabel
+ %34 = OpFunctionCall %void %round_1c7897
OpReturnValue %5
OpFunctionEnd
-%vertex_main = OpFunction %void None %9
- %28 = OpLabel
- %29 = OpFunctionCall %v4float %vertex_main_inner
- OpStore %value %29
+%vertex_main = OpFunction %void None %13
+ %36 = OpLabel
+ %37 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %37
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
-%fragment_main = OpFunction %void None %9
- %32 = OpLabel
- %33 = OpFunctionCall %void %round_1c7897
+%fragment_main = OpFunction %void None %13
+ %40 = OpLabel
+ %41 = OpFunctionCall %void %round_1c7897
OpReturn
OpFunctionEnd
-%compute_main = OpFunction %void None %9
- %35 = OpLabel
- %36 = OpFunctionCall %void %round_1c7897
+%compute_main = OpFunction %void None %13
+ %43 = OpLabel
+ %44 = OpFunctionCall %void %round_1c7897
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/round/1c7897.wgsl.expected.wgsl b/test/tint/builtins/gen/var/round/1c7897.wgsl.expected.wgsl
index e644d19..d0b3457 100644
--- a/test/tint/builtins/gen/var/round/1c7897.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/var/round/1c7897.wgsl.expected.wgsl
@@ -1,8 +1,11 @@
fn round_1c7897() {
var arg_0 = vec3<f32>(3.5f);
var res : vec3<f32> = round(arg_0);
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec3<f32>;
+
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
round_1c7897();
diff --git a/test/tint/builtins/gen/var/round/52c84d.wgsl b/test/tint/builtins/gen/var/round/52c84d.wgsl
index 1bef0ea..2f43814 100644
--- a/test/tint/builtins/gen/var/round/52c84d.wgsl
+++ b/test/tint/builtins/gen/var/round/52c84d.wgsl
@@ -25,7 +25,9 @@
fn round_52c84d() {
var arg_0 = vec2<f32>(3.5f);
var res: vec2<f32> = round(arg_0);
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec2<f32>;
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
diff --git a/test/tint/builtins/gen/var/round/52c84d.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/round/52c84d.wgsl.expected.dxc.hlsl
index 45e93bb..ad68e69 100644
--- a/test/tint/builtins/gen/var/round/52c84d.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/var/round/52c84d.wgsl.expected.dxc.hlsl
@@ -1,6 +1,9 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
void round_52c84d() {
float2 arg_0 = (3.5f).xx;
float2 res = round(arg_0);
+ prevent_dce.Store2(0u, asuint(res));
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/var/round/52c84d.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/round/52c84d.wgsl.expected.fxc.hlsl
index 45e93bb..ad68e69 100644
--- a/test/tint/builtins/gen/var/round/52c84d.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/var/round/52c84d.wgsl.expected.fxc.hlsl
@@ -1,6 +1,9 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
void round_52c84d() {
float2 arg_0 = (3.5f).xx;
float2 res = round(arg_0);
+ prevent_dce.Store2(0u, asuint(res));
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/var/round/52c84d.wgsl.expected.glsl b/test/tint/builtins/gen/var/round/52c84d.wgsl.expected.glsl
index 0d66f20..d11ef5d 100644
--- a/test/tint/builtins/gen/var/round/52c84d.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/round/52c84d.wgsl.expected.glsl
@@ -1,8 +1,13 @@
#version 310 es
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ vec2 inner;
+} prevent_dce;
+
void round_52c84d() {
vec2 arg_0 = vec2(3.5f);
vec2 res = round(arg_0);
+ prevent_dce.inner = res;
}
vec4 vertex_main() {
@@ -21,9 +26,14 @@
#version 310 es
precision mediump float;
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ vec2 inner;
+} prevent_dce;
+
void round_52c84d() {
vec2 arg_0 = vec2(3.5f);
vec2 res = round(arg_0);
+ prevent_dce.inner = res;
}
void fragment_main() {
@@ -36,9 +46,14 @@
}
#version 310 es
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ vec2 inner;
+} prevent_dce;
+
void round_52c84d() {
vec2 arg_0 = vec2(3.5f);
vec2 res = round(arg_0);
+ prevent_dce.inner = res;
}
void compute_main() {
diff --git a/test/tint/builtins/gen/var/round/52c84d.wgsl.expected.msl b/test/tint/builtins/gen/var/round/52c84d.wgsl.expected.msl
index 3b5b707..fd1a172 100644
--- a/test/tint/builtins/gen/var/round/52c84d.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/round/52c84d.wgsl.expected.msl
@@ -1,34 +1,35 @@
#include <metal_stdlib>
using namespace metal;
-void round_52c84d() {
+void round_52c84d(device float2* const tint_symbol_1) {
float2 arg_0 = float2(3.5f);
float2 res = rint(arg_0);
+ *(tint_symbol_1) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner() {
- round_52c84d();
+float4 vertex_main_inner(device float2* const tint_symbol_2) {
+ round_52c84d(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 float2* 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() {
- round_52c84d();
+fragment void fragment_main(device float2* tint_symbol_4 [[buffer(0)]]) {
+ round_52c84d(tint_symbol_4);
return;
}
-kernel void compute_main() {
- round_52c84d();
+kernel void compute_main(device float2* tint_symbol_5 [[buffer(0)]]) {
+ round_52c84d(tint_symbol_5);
return;
}
diff --git a/test/tint/builtins/gen/var/round/52c84d.wgsl.expected.spvasm b/test/tint/builtins/gen/var/round/52c84d.wgsl.expected.spvasm
index 9abdd19..dde53eb 100644
--- a/test/tint/builtins/gen/var/round/52c84d.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/round/52c84d.wgsl.expected.spvasm
@@ -1,10 +1,10 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 37
+; Bound: 45
; Schema: 0
OpCapability Shader
- %20 = OpExtInstImport "GLSL.std.450"
+ %23 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -13,6 +13,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 %round_52c84d "round_52c84d"
OpName %arg_0 "arg_0"
OpName %res "res"
@@ -22,6 +25,10 @@
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
+ OpDecorate %prevent_dce DescriptorSet 2
+ OpDecorate %prevent_dce Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
@@ -30,44 +37,53 @@
%_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
+%prevent_dce_block = OpTypeStruct %v2float
+%_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_3_5 = OpConstant %float 3.5
- %15 = OpConstantComposite %v2float %float_3_5 %float_3_5
+ %18 = OpConstantComposite %v2float %float_3_5 %float_3_5
%_ptr_Function_v2float = OpTypePointer Function %v2float
- %18 = OpConstantNull %v2float
- %23 = OpTypeFunction %v4float
+ %21 = OpConstantNull %v2float
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v2float = OpTypePointer StorageBuffer %v2float
+ %31 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
-%round_52c84d = OpFunction %void None %9
- %12 = OpLabel
- %arg_0 = OpVariable %_ptr_Function_v2float Function %18
- %res = OpVariable %_ptr_Function_v2float Function %18
- OpStore %arg_0 %15
- %21 = OpLoad %v2float %arg_0
- %19 = OpExtInst %v2float %20 RoundEven %21
- OpStore %res %19
+%round_52c84d = OpFunction %void None %13
+ %16 = OpLabel
+ %arg_0 = OpVariable %_ptr_Function_v2float Function %21
+ %res = OpVariable %_ptr_Function_v2float Function %21
+ OpStore %arg_0 %18
+ %24 = OpLoad %v2float %arg_0
+ %22 = OpExtInst %v2float %23 RoundEven %24
+ OpStore %res %22
+ %29 = OpAccessChain %_ptr_StorageBuffer_v2float %prevent_dce %uint_0
+ %30 = OpLoad %v2float %res
+ OpStore %29 %30
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %23
- %25 = OpLabel
- %26 = OpFunctionCall %void %round_52c84d
+%vertex_main_inner = OpFunction %v4float None %31
+ %33 = OpLabel
+ %34 = OpFunctionCall %void %round_52c84d
OpReturnValue %5
OpFunctionEnd
-%vertex_main = OpFunction %void None %9
- %28 = OpLabel
- %29 = OpFunctionCall %v4float %vertex_main_inner
- OpStore %value %29
+%vertex_main = OpFunction %void None %13
+ %36 = OpLabel
+ %37 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %37
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
-%fragment_main = OpFunction %void None %9
- %32 = OpLabel
- %33 = OpFunctionCall %void %round_52c84d
+%fragment_main = OpFunction %void None %13
+ %40 = OpLabel
+ %41 = OpFunctionCall %void %round_52c84d
OpReturn
OpFunctionEnd
-%compute_main = OpFunction %void None %9
- %35 = OpLabel
- %36 = OpFunctionCall %void %round_52c84d
+%compute_main = OpFunction %void None %13
+ %43 = OpLabel
+ %44 = OpFunctionCall %void %round_52c84d
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/round/52c84d.wgsl.expected.wgsl b/test/tint/builtins/gen/var/round/52c84d.wgsl.expected.wgsl
index c017b09..86cef3a 100644
--- a/test/tint/builtins/gen/var/round/52c84d.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/var/round/52c84d.wgsl.expected.wgsl
@@ -1,8 +1,11 @@
fn round_52c84d() {
var arg_0 = vec2<f32>(3.5f);
var res : vec2<f32> = round(arg_0);
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec2<f32>;
+
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
round_52c84d();
diff --git a/test/tint/builtins/gen/var/round/773a8f.wgsl b/test/tint/builtins/gen/var/round/773a8f.wgsl
index e5c9393..0f4d254 100644
--- a/test/tint/builtins/gen/var/round/773a8f.wgsl
+++ b/test/tint/builtins/gen/var/round/773a8f.wgsl
@@ -26,7 +26,6 @@
const arg_0 = 3.5;
var res = round(arg_0);
}
-
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
round_773a8f();
diff --git a/test/tint/builtins/gen/var/round/8fdca3.wgsl b/test/tint/builtins/gen/var/round/8fdca3.wgsl
index 07c3c24..9bc9f0d 100644
--- a/test/tint/builtins/gen/var/round/8fdca3.wgsl
+++ b/test/tint/builtins/gen/var/round/8fdca3.wgsl
@@ -26,7 +26,6 @@
const arg_0 = vec2(3.5);
var res = round(arg_0);
}
-
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
round_8fdca3();
diff --git a/test/tint/builtins/gen/var/round/9078ef.wgsl b/test/tint/builtins/gen/var/round/9078ef.wgsl
index ec8ec54..acfebda 100644
--- a/test/tint/builtins/gen/var/round/9078ef.wgsl
+++ b/test/tint/builtins/gen/var/round/9078ef.wgsl
@@ -27,7 +27,9 @@
fn round_9078ef() {
var arg_0 = 3.5h;
var res: f16 = round(arg_0);
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : f16;
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
diff --git a/test/tint/builtins/gen/var/round/9078ef.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/round/9078ef.wgsl.expected.dxc.hlsl
index 3f465c9..e4fb385 100644
--- a/test/tint/builtins/gen/var/round/9078ef.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/var/round/9078ef.wgsl.expected.dxc.hlsl
@@ -1,6 +1,9 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
void round_9078ef() {
float16_t arg_0 = float16_t(3.5h);
float16_t res = round(arg_0);
+ prevent_dce.Store<float16_t>(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/var/round/9078ef.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/round/9078ef.wgsl.expected.fxc.hlsl
index ed3e148..956c07f 100644
--- a/test/tint/builtins/gen/var/round/9078ef.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/var/round/9078ef.wgsl.expected.fxc.hlsl
@@ -1,8 +1,11 @@
SKIP: FAILED
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
void round_9078ef() {
- float16_t arg_0 = float16_t(3.3984375h);
+ float16_t arg_0 = float16_t(3.5h);
float16_t res = round(arg_0);
+ prevent_dce.Store<float16_t>(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/var/round/9078ef.wgsl.expected.glsl b/test/tint/builtins/gen/var/round/9078ef.wgsl.expected.glsl
index 5fe11aa..c4f1ff0 100644
--- a/test/tint/builtins/gen/var/round/9078ef.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/round/9078ef.wgsl.expected.glsl
@@ -1,9 +1,14 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ float16_t inner;
+} prevent_dce;
+
void round_9078ef() {
float16_t arg_0 = 3.5hf;
float16_t res = round(arg_0);
+ prevent_dce.inner = res;
}
vec4 vertex_main() {
@@ -23,9 +28,14 @@
#extension GL_AMD_gpu_shader_half_float : require
precision mediump float;
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ float16_t inner;
+} prevent_dce;
+
void round_9078ef() {
float16_t arg_0 = 3.5hf;
float16_t res = round(arg_0);
+ prevent_dce.inner = res;
}
void fragment_main() {
@@ -39,9 +49,14 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ float16_t inner;
+} prevent_dce;
+
void round_9078ef() {
float16_t arg_0 = 3.5hf;
float16_t res = round(arg_0);
+ prevent_dce.inner = res;
}
void compute_main() {
diff --git a/test/tint/builtins/gen/var/round/9078ef.wgsl.expected.msl b/test/tint/builtins/gen/var/round/9078ef.wgsl.expected.msl
index 14da505..0d2fac4 100644
--- a/test/tint/builtins/gen/var/round/9078ef.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/round/9078ef.wgsl.expected.msl
@@ -1,34 +1,35 @@
#include <metal_stdlib>
using namespace metal;
-void round_9078ef() {
+void round_9078ef(device half* const tint_symbol_1) {
half arg_0 = 3.5h;
half res = rint(arg_0);
+ *(tint_symbol_1) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner() {
- round_9078ef();
+float4 vertex_main_inner(device half* const tint_symbol_2) {
+ round_9078ef(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 half* 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() {
- round_9078ef();
+fragment void fragment_main(device half* tint_symbol_4 [[buffer(0)]]) {
+ round_9078ef(tint_symbol_4);
return;
}
-kernel void compute_main() {
- round_9078ef();
+kernel void compute_main(device half* tint_symbol_5 [[buffer(0)]]) {
+ round_9078ef(tint_symbol_5);
return;
}
diff --git a/test/tint/builtins/gen/var/round/9078ef.wgsl.expected.spvasm b/test/tint/builtins/gen/var/round/9078ef.wgsl.expected.spvasm
index 1b3b587..4ebeac2 100644
--- a/test/tint/builtins/gen/var/round/9078ef.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/round/9078ef.wgsl.expected.spvasm
@@ -1,14 +1,14 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 36
+; Bound: 44
; Schema: 0
OpCapability Shader
OpCapability Float16
OpCapability UniformAndStorageBuffer16BitAccess
OpCapability StorageBuffer16BitAccess
OpCapability StorageInputOutput16
- %19 = OpExtInstImport "GLSL.std.450"
+ %22 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -17,6 +17,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 %round_9078ef "round_9078ef"
OpName %arg_0 "arg_0"
OpName %res "res"
@@ -26,6 +29,10 @@
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
+ OpDecorate %prevent_dce DescriptorSet 2
+ OpDecorate %prevent_dce Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
@@ -34,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
+%prevent_dce_block = OpTypeStruct %half
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+ %void = OpTypeVoid
+ %13 = OpTypeFunction %void
%half_0x1_cp_1 = OpConstant %half 0x1.cp+1
%_ptr_Function_half = OpTypePointer Function %half
- %17 = OpConstantNull %half
- %22 = OpTypeFunction %v4float
+ %20 = OpConstantNull %half
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_half = OpTypePointer StorageBuffer %half
+ %30 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
-%round_9078ef = OpFunction %void None %9
- %12 = OpLabel
- %arg_0 = OpVariable %_ptr_Function_half Function %17
- %res = OpVariable %_ptr_Function_half Function %17
+%round_9078ef = OpFunction %void None %13
+ %16 = OpLabel
+ %arg_0 = OpVariable %_ptr_Function_half Function %20
+ %res = OpVariable %_ptr_Function_half Function %20
OpStore %arg_0 %half_0x1_cp_1
- %20 = OpLoad %half %arg_0
- %18 = OpExtInst %half %19 RoundEven %20
- OpStore %res %18
+ %23 = OpLoad %half %arg_0
+ %21 = OpExtInst %half %22 RoundEven %23
+ OpStore %res %21
+ %28 = OpAccessChain %_ptr_StorageBuffer_half %prevent_dce %uint_0
+ %29 = OpLoad %half %res
+ OpStore %28 %29
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %22
- %24 = OpLabel
- %25 = OpFunctionCall %void %round_9078ef
+%vertex_main_inner = OpFunction %v4float None %30
+ %32 = OpLabel
+ %33 = OpFunctionCall %void %round_9078ef
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 %13
+ %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 %round_9078ef
+%fragment_main = OpFunction %void None %13
+ %39 = OpLabel
+ %40 = OpFunctionCall %void %round_9078ef
OpReturn
OpFunctionEnd
-%compute_main = OpFunction %void None %9
- %34 = OpLabel
- %35 = OpFunctionCall %void %round_9078ef
+%compute_main = OpFunction %void None %13
+ %42 = OpLabel
+ %43 = OpFunctionCall %void %round_9078ef
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/round/9078ef.wgsl.expected.wgsl b/test/tint/builtins/gen/var/round/9078ef.wgsl.expected.wgsl
index 4ba193e..b9d6661 100644
--- a/test/tint/builtins/gen/var/round/9078ef.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/var/round/9078ef.wgsl.expected.wgsl
@@ -3,8 +3,11 @@
fn round_9078ef() {
var arg_0 = 3.5h;
var res : f16 = round(arg_0);
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : f16;
+
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
round_9078ef();
diff --git a/test/tint/builtins/gen/var/round/9edc38.wgsl b/test/tint/builtins/gen/var/round/9edc38.wgsl
index 39ea9fd..83c7127 100644
--- a/test/tint/builtins/gen/var/round/9edc38.wgsl
+++ b/test/tint/builtins/gen/var/round/9edc38.wgsl
@@ -25,7 +25,9 @@
fn round_9edc38() {
var arg_0 = 3.5f;
var res: f32 = round(arg_0);
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : f32;
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
diff --git a/test/tint/builtins/gen/var/round/9edc38.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/round/9edc38.wgsl.expected.dxc.hlsl
index 6c6d0e5..5c09d16 100644
--- a/test/tint/builtins/gen/var/round/9edc38.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/var/round/9edc38.wgsl.expected.dxc.hlsl
@@ -1,6 +1,9 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
void round_9edc38() {
float arg_0 = 3.5f;
float res = round(arg_0);
+ prevent_dce.Store(0u, asuint(res));
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/var/round/9edc38.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/round/9edc38.wgsl.expected.fxc.hlsl
index 6c6d0e5..5c09d16 100644
--- a/test/tint/builtins/gen/var/round/9edc38.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/var/round/9edc38.wgsl.expected.fxc.hlsl
@@ -1,6 +1,9 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
void round_9edc38() {
float arg_0 = 3.5f;
float res = round(arg_0);
+ prevent_dce.Store(0u, asuint(res));
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/var/round/9edc38.wgsl.expected.glsl b/test/tint/builtins/gen/var/round/9edc38.wgsl.expected.glsl
index 84da2d8..8114437 100644
--- a/test/tint/builtins/gen/var/round/9edc38.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/round/9edc38.wgsl.expected.glsl
@@ -1,8 +1,13 @@
#version 310 es
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ float inner;
+} prevent_dce;
+
void round_9edc38() {
float arg_0 = 3.5f;
float res = round(arg_0);
+ prevent_dce.inner = res;
}
vec4 vertex_main() {
@@ -21,9 +26,14 @@
#version 310 es
precision mediump float;
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ float inner;
+} prevent_dce;
+
void round_9edc38() {
float arg_0 = 3.5f;
float res = round(arg_0);
+ prevent_dce.inner = res;
}
void fragment_main() {
@@ -36,9 +46,14 @@
}
#version 310 es
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ float inner;
+} prevent_dce;
+
void round_9edc38() {
float arg_0 = 3.5f;
float res = round(arg_0);
+ prevent_dce.inner = res;
}
void compute_main() {
diff --git a/test/tint/builtins/gen/var/round/9edc38.wgsl.expected.msl b/test/tint/builtins/gen/var/round/9edc38.wgsl.expected.msl
index d64a376..28152d7 100644
--- a/test/tint/builtins/gen/var/round/9edc38.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/round/9edc38.wgsl.expected.msl
@@ -1,34 +1,35 @@
#include <metal_stdlib>
using namespace metal;
-void round_9edc38() {
+void round_9edc38(device float* const tint_symbol_1) {
float arg_0 = 3.5f;
float res = rint(arg_0);
+ *(tint_symbol_1) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner() {
- round_9edc38();
+float4 vertex_main_inner(device float* const tint_symbol_2) {
+ round_9edc38(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 float* 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() {
- round_9edc38();
+fragment void fragment_main(device float* tint_symbol_4 [[buffer(0)]]) {
+ round_9edc38(tint_symbol_4);
return;
}
-kernel void compute_main() {
- round_9edc38();
+kernel void compute_main(device float* tint_symbol_5 [[buffer(0)]]) {
+ round_9edc38(tint_symbol_5);
return;
}
diff --git a/test/tint/builtins/gen/var/round/9edc38.wgsl.expected.spvasm b/test/tint/builtins/gen/var/round/9edc38.wgsl.expected.spvasm
index d817cdb..e43eb91 100644
--- a/test/tint/builtins/gen/var/round/9edc38.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/round/9edc38.wgsl.expected.spvasm
@@ -1,10 +1,10 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 34
+; Bound: 42
; Schema: 0
OpCapability Shader
- %17 = OpExtInstImport "GLSL.std.450"
+ %20 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -13,6 +13,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 %round_9edc38 "round_9edc38"
OpName %arg_0 "arg_0"
OpName %res "res"
@@ -22,6 +25,10 @@
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
+ OpDecorate %prevent_dce DescriptorSet 2
+ OpDecorate %prevent_dce Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
@@ -30,41 +37,50 @@
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+%prevent_dce_block = OpTypeStruct %float
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
%void = OpTypeVoid
- %9 = OpTypeFunction %void
+ %12 = OpTypeFunction %void
%float_3_5 = OpConstant %float 3.5
%_ptr_Function_float = OpTypePointer Function %float
- %20 = OpTypeFunction %v4float
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+ %28 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
-%round_9edc38 = OpFunction %void None %9
- %12 = OpLabel
+%round_9edc38 = OpFunction %void None %12
+ %15 = OpLabel
%arg_0 = OpVariable %_ptr_Function_float Function %8
%res = OpVariable %_ptr_Function_float Function %8
OpStore %arg_0 %float_3_5
- %18 = OpLoad %float %arg_0
- %16 = OpExtInst %float %17 RoundEven %18
- OpStore %res %16
+ %21 = OpLoad %float %arg_0
+ %19 = OpExtInst %float %20 RoundEven %21
+ OpStore %res %19
+ %26 = OpAccessChain %_ptr_StorageBuffer_float %prevent_dce %uint_0
+ %27 = OpLoad %float %res
+ OpStore %26 %27
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %20
- %22 = OpLabel
- %23 = OpFunctionCall %void %round_9edc38
+%vertex_main_inner = OpFunction %v4float None %28
+ %30 = OpLabel
+ %31 = OpFunctionCall %void %round_9edc38
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 %12
+ %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
- %29 = OpLabel
- %30 = OpFunctionCall %void %round_9edc38
+%fragment_main = OpFunction %void None %12
+ %37 = OpLabel
+ %38 = OpFunctionCall %void %round_9edc38
OpReturn
OpFunctionEnd
-%compute_main = OpFunction %void None %9
- %32 = OpLabel
- %33 = OpFunctionCall %void %round_9edc38
+%compute_main = OpFunction %void None %12
+ %40 = OpLabel
+ %41 = OpFunctionCall %void %round_9edc38
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/round/9edc38.wgsl.expected.wgsl b/test/tint/builtins/gen/var/round/9edc38.wgsl.expected.wgsl
index a2f0a1f..89d96e8 100644
--- a/test/tint/builtins/gen/var/round/9edc38.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/var/round/9edc38.wgsl.expected.wgsl
@@ -1,8 +1,11 @@
fn round_9edc38() {
var arg_0 = 3.5f;
var res : f32 = round(arg_0);
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : f32;
+
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
round_9edc38();
diff --git a/test/tint/builtins/gen/var/round/a1673d.wgsl b/test/tint/builtins/gen/var/round/a1673d.wgsl
index 09b3094..8812a77 100644
--- a/test/tint/builtins/gen/var/round/a1673d.wgsl
+++ b/test/tint/builtins/gen/var/round/a1673d.wgsl
@@ -26,7 +26,6 @@
const arg_0 = vec3(3.5);
var res = round(arg_0);
}
-
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
round_a1673d();
diff --git a/test/tint/builtins/gen/var/round/d87e84.wgsl b/test/tint/builtins/gen/var/round/d87e84.wgsl
index 8b11473..3b1c764 100644
--- a/test/tint/builtins/gen/var/round/d87e84.wgsl
+++ b/test/tint/builtins/gen/var/round/d87e84.wgsl
@@ -27,7 +27,9 @@
fn round_d87e84() {
var arg_0 = vec2<f16>(3.5h);
var res: vec2<f16> = round(arg_0);
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec2<f16>;
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
diff --git a/test/tint/builtins/gen/var/round/d87e84.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/round/d87e84.wgsl.expected.dxc.hlsl
index 9895f74b..2a01f0a 100644
--- a/test/tint/builtins/gen/var/round/d87e84.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/var/round/d87e84.wgsl.expected.dxc.hlsl
@@ -1,6 +1,9 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
void round_d87e84() {
vector<float16_t, 2> arg_0 = (float16_t(3.5h)).xx;
vector<float16_t, 2> res = round(arg_0);
+ prevent_dce.Store<vector<float16_t, 2> >(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/var/round/d87e84.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/round/d87e84.wgsl.expected.fxc.hlsl
index c09e44f..531e516 100644
--- a/test/tint/builtins/gen/var/round/d87e84.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/var/round/d87e84.wgsl.expected.fxc.hlsl
@@ -1,8 +1,11 @@
SKIP: FAILED
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
void round_d87e84() {
- vector<float16_t, 2> arg_0 = (float16_t(3.3984375h)).xx;
+ vector<float16_t, 2> arg_0 = (float16_t(3.5h)).xx;
vector<float16_t, 2> res = round(arg_0);
+ prevent_dce.Store<vector<float16_t, 2> >(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/var/round/d87e84.wgsl.expected.glsl b/test/tint/builtins/gen/var/round/d87e84.wgsl.expected.glsl
index 2d6d5a8..1f8bf23 100644
--- a/test/tint/builtins/gen/var/round/d87e84.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/round/d87e84.wgsl.expected.glsl
@@ -1,9 +1,14 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16vec2 inner;
+} prevent_dce;
+
void round_d87e84() {
f16vec2 arg_0 = f16vec2(3.5hf);
f16vec2 res = round(arg_0);
+ prevent_dce.inner = res;
}
vec4 vertex_main() {
@@ -23,9 +28,14 @@
#extension GL_AMD_gpu_shader_half_float : require
precision mediump float;
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16vec2 inner;
+} prevent_dce;
+
void round_d87e84() {
f16vec2 arg_0 = f16vec2(3.5hf);
f16vec2 res = round(arg_0);
+ prevent_dce.inner = res;
}
void fragment_main() {
@@ -39,9 +49,14 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16vec2 inner;
+} prevent_dce;
+
void round_d87e84() {
f16vec2 arg_0 = f16vec2(3.5hf);
f16vec2 res = round(arg_0);
+ prevent_dce.inner = res;
}
void compute_main() {
diff --git a/test/tint/builtins/gen/var/round/d87e84.wgsl.expected.msl b/test/tint/builtins/gen/var/round/d87e84.wgsl.expected.msl
index 2a9ee39..551e4df 100644
--- a/test/tint/builtins/gen/var/round/d87e84.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/round/d87e84.wgsl.expected.msl
@@ -1,34 +1,35 @@
#include <metal_stdlib>
using namespace metal;
-void round_d87e84() {
+void round_d87e84(device half2* const tint_symbol_1) {
half2 arg_0 = half2(3.5h);
half2 res = rint(arg_0);
+ *(tint_symbol_1) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner() {
- round_d87e84();
+float4 vertex_main_inner(device half2* const tint_symbol_2) {
+ round_d87e84(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 half2* 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() {
- round_d87e84();
+fragment void fragment_main(device half2* tint_symbol_4 [[buffer(0)]]) {
+ round_d87e84(tint_symbol_4);
return;
}
-kernel void compute_main() {
- round_d87e84();
+kernel void compute_main(device half2* tint_symbol_5 [[buffer(0)]]) {
+ round_d87e84(tint_symbol_5);
return;
}
diff --git a/test/tint/builtins/gen/var/round/d87e84.wgsl.expected.spvasm b/test/tint/builtins/gen/var/round/d87e84.wgsl.expected.spvasm
index 1696b37..74fdb39 100644
--- a/test/tint/builtins/gen/var/round/d87e84.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/round/d87e84.wgsl.expected.spvasm
@@ -1,14 +1,14 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 38
+; Bound: 46
; Schema: 0
OpCapability Shader
OpCapability Float16
OpCapability UniformAndStorageBuffer16BitAccess
OpCapability StorageBuffer16BitAccess
OpCapability StorageInputOutput16
- %21 = OpExtInstImport "GLSL.std.450"
+ %24 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -17,6 +17,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 %round_d87e84 "round_d87e84"
OpName %arg_0 "arg_0"
OpName %res "res"
@@ -26,6 +29,10 @@
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
+ OpDecorate %prevent_dce DescriptorSet 2
+ OpDecorate %prevent_dce Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
@@ -34,45 +41,54 @@
%_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
+%prevent_dce_block = OpTypeStruct %v2half
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+ %void = OpTypeVoid
+ %14 = OpTypeFunction %void
%half_0x1_cp_1 = OpConstant %half 0x1.cp+1
- %16 = OpConstantComposite %v2half %half_0x1_cp_1 %half_0x1_cp_1
+ %19 = OpConstantComposite %v2half %half_0x1_cp_1 %half_0x1_cp_1
%_ptr_Function_v2half = OpTypePointer Function %v2half
- %19 = OpConstantNull %v2half
- %24 = OpTypeFunction %v4float
+ %22 = OpConstantNull %v2half
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v2half = OpTypePointer StorageBuffer %v2half
+ %32 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
-%round_d87e84 = OpFunction %void None %9
- %12 = OpLabel
- %arg_0 = OpVariable %_ptr_Function_v2half Function %19
- %res = OpVariable %_ptr_Function_v2half Function %19
- OpStore %arg_0 %16
- %22 = OpLoad %v2half %arg_0
- %20 = OpExtInst %v2half %21 RoundEven %22
- OpStore %res %20
+%round_d87e84 = OpFunction %void None %14
+ %17 = OpLabel
+ %arg_0 = OpVariable %_ptr_Function_v2half Function %22
+ %res = OpVariable %_ptr_Function_v2half Function %22
+ OpStore %arg_0 %19
+ %25 = OpLoad %v2half %arg_0
+ %23 = OpExtInst %v2half %24 RoundEven %25
+ OpStore %res %23
+ %30 = OpAccessChain %_ptr_StorageBuffer_v2half %prevent_dce %uint_0
+ %31 = OpLoad %v2half %res
+ OpStore %30 %31
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %24
- %26 = OpLabel
- %27 = OpFunctionCall %void %round_d87e84
+%vertex_main_inner = OpFunction %v4float None %32
+ %34 = OpLabel
+ %35 = OpFunctionCall %void %round_d87e84
OpReturnValue %5
OpFunctionEnd
-%vertex_main = OpFunction %void None %9
- %29 = OpLabel
- %30 = OpFunctionCall %v4float %vertex_main_inner
- OpStore %value %30
+%vertex_main = OpFunction %void None %14
+ %37 = OpLabel
+ %38 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %38
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
-%fragment_main = OpFunction %void None %9
- %33 = OpLabel
- %34 = OpFunctionCall %void %round_d87e84
+%fragment_main = OpFunction %void None %14
+ %41 = OpLabel
+ %42 = OpFunctionCall %void %round_d87e84
OpReturn
OpFunctionEnd
-%compute_main = OpFunction %void None %9
- %36 = OpLabel
- %37 = OpFunctionCall %void %round_d87e84
+%compute_main = OpFunction %void None %14
+ %44 = OpLabel
+ %45 = OpFunctionCall %void %round_d87e84
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/round/d87e84.wgsl.expected.wgsl b/test/tint/builtins/gen/var/round/d87e84.wgsl.expected.wgsl
index f91e883..8b6e553 100644
--- a/test/tint/builtins/gen/var/round/d87e84.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/var/round/d87e84.wgsl.expected.wgsl
@@ -3,8 +3,11 @@
fn round_d87e84() {
var arg_0 = vec2<f16>(3.5h);
var res : vec2<f16> = round(arg_0);
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec2<f16>;
+
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
round_d87e84();
diff --git a/test/tint/builtins/gen/var/round/e1bba2.wgsl b/test/tint/builtins/gen/var/round/e1bba2.wgsl
index a4fbeab..dfac69f 100644
--- a/test/tint/builtins/gen/var/round/e1bba2.wgsl
+++ b/test/tint/builtins/gen/var/round/e1bba2.wgsl
@@ -27,7 +27,9 @@
fn round_e1bba2() {
var arg_0 = vec3<f16>(3.5h);
var res: vec3<f16> = round(arg_0);
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec3<f16>;
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
diff --git a/test/tint/builtins/gen/var/round/e1bba2.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/round/e1bba2.wgsl.expected.dxc.hlsl
index a92457e..544ef7f 100644
--- a/test/tint/builtins/gen/var/round/e1bba2.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/var/round/e1bba2.wgsl.expected.dxc.hlsl
@@ -1,6 +1,9 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
void round_e1bba2() {
vector<float16_t, 3> arg_0 = (float16_t(3.5h)).xxx;
vector<float16_t, 3> res = round(arg_0);
+ prevent_dce.Store<vector<float16_t, 3> >(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/var/round/e1bba2.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/round/e1bba2.wgsl.expected.fxc.hlsl
index 762bb62..84d41c4 100644
--- a/test/tint/builtins/gen/var/round/e1bba2.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/var/round/e1bba2.wgsl.expected.fxc.hlsl
@@ -1,8 +1,11 @@
SKIP: FAILED
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
void round_e1bba2() {
- vector<float16_t, 3> arg_0 = (float16_t(3.3984375h)).xxx;
+ vector<float16_t, 3> arg_0 = (float16_t(3.5h)).xxx;
vector<float16_t, 3> res = round(arg_0);
+ prevent_dce.Store<vector<float16_t, 3> >(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/var/round/e1bba2.wgsl.expected.glsl b/test/tint/builtins/gen/var/round/e1bba2.wgsl.expected.glsl
index 8d58b62..e305e2b 100644
--- a/test/tint/builtins/gen/var/round/e1bba2.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/round/e1bba2.wgsl.expected.glsl
@@ -1,9 +1,14 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16vec3 inner;
+} prevent_dce;
+
void round_e1bba2() {
f16vec3 arg_0 = f16vec3(3.5hf);
f16vec3 res = round(arg_0);
+ prevent_dce.inner = res;
}
vec4 vertex_main() {
@@ -23,9 +28,14 @@
#extension GL_AMD_gpu_shader_half_float : require
precision mediump float;
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16vec3 inner;
+} prevent_dce;
+
void round_e1bba2() {
f16vec3 arg_0 = f16vec3(3.5hf);
f16vec3 res = round(arg_0);
+ prevent_dce.inner = res;
}
void fragment_main() {
@@ -39,9 +49,14 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16vec3 inner;
+} prevent_dce;
+
void round_e1bba2() {
f16vec3 arg_0 = f16vec3(3.5hf);
f16vec3 res = round(arg_0);
+ prevent_dce.inner = res;
}
void compute_main() {
diff --git a/test/tint/builtins/gen/var/round/e1bba2.wgsl.expected.msl b/test/tint/builtins/gen/var/round/e1bba2.wgsl.expected.msl
index cc62ca3..da90b41 100644
--- a/test/tint/builtins/gen/var/round/e1bba2.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/round/e1bba2.wgsl.expected.msl
@@ -1,34 +1,35 @@
#include <metal_stdlib>
using namespace metal;
-void round_e1bba2() {
+void round_e1bba2(device packed_half3* const tint_symbol_1) {
half3 arg_0 = half3(3.5h);
half3 res = rint(arg_0);
+ *(tint_symbol_1) = packed_half3(res);
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner() {
- round_e1bba2();
+float4 vertex_main_inner(device packed_half3* const tint_symbol_2) {
+ round_e1bba2(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 packed_half3* 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() {
- round_e1bba2();
+fragment void fragment_main(device packed_half3* tint_symbol_4 [[buffer(0)]]) {
+ round_e1bba2(tint_symbol_4);
return;
}
-kernel void compute_main() {
- round_e1bba2();
+kernel void compute_main(device packed_half3* tint_symbol_5 [[buffer(0)]]) {
+ round_e1bba2(tint_symbol_5);
return;
}
diff --git a/test/tint/builtins/gen/var/round/e1bba2.wgsl.expected.spvasm b/test/tint/builtins/gen/var/round/e1bba2.wgsl.expected.spvasm
index e23c3a1..984dab4 100644
--- a/test/tint/builtins/gen/var/round/e1bba2.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/round/e1bba2.wgsl.expected.spvasm
@@ -1,14 +1,14 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 38
+; Bound: 46
; Schema: 0
OpCapability Shader
OpCapability Float16
OpCapability UniformAndStorageBuffer16BitAccess
OpCapability StorageBuffer16BitAccess
OpCapability StorageInputOutput16
- %21 = OpExtInstImport "GLSL.std.450"
+ %24 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -17,6 +17,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 %round_e1bba2 "round_e1bba2"
OpName %arg_0 "arg_0"
OpName %res "res"
@@ -26,6 +29,10 @@
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
+ OpDecorate %prevent_dce DescriptorSet 2
+ OpDecorate %prevent_dce Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
@@ -34,45 +41,54 @@
%_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
+%prevent_dce_block = OpTypeStruct %v3half
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+ %void = OpTypeVoid
+ %14 = OpTypeFunction %void
%half_0x1_cp_1 = OpConstant %half 0x1.cp+1
- %16 = OpConstantComposite %v3half %half_0x1_cp_1 %half_0x1_cp_1 %half_0x1_cp_1
+ %19 = OpConstantComposite %v3half %half_0x1_cp_1 %half_0x1_cp_1 %half_0x1_cp_1
%_ptr_Function_v3half = OpTypePointer Function %v3half
- %19 = OpConstantNull %v3half
- %24 = OpTypeFunction %v4float
+ %22 = OpConstantNull %v3half
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v3half = OpTypePointer StorageBuffer %v3half
+ %32 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
-%round_e1bba2 = OpFunction %void None %9
- %12 = OpLabel
- %arg_0 = OpVariable %_ptr_Function_v3half Function %19
- %res = OpVariable %_ptr_Function_v3half Function %19
- OpStore %arg_0 %16
- %22 = OpLoad %v3half %arg_0
- %20 = OpExtInst %v3half %21 RoundEven %22
- OpStore %res %20
+%round_e1bba2 = OpFunction %void None %14
+ %17 = OpLabel
+ %arg_0 = OpVariable %_ptr_Function_v3half Function %22
+ %res = OpVariable %_ptr_Function_v3half Function %22
+ OpStore %arg_0 %19
+ %25 = OpLoad %v3half %arg_0
+ %23 = OpExtInst %v3half %24 RoundEven %25
+ OpStore %res %23
+ %30 = OpAccessChain %_ptr_StorageBuffer_v3half %prevent_dce %uint_0
+ %31 = OpLoad %v3half %res
+ OpStore %30 %31
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %24
- %26 = OpLabel
- %27 = OpFunctionCall %void %round_e1bba2
+%vertex_main_inner = OpFunction %v4float None %32
+ %34 = OpLabel
+ %35 = OpFunctionCall %void %round_e1bba2
OpReturnValue %5
OpFunctionEnd
-%vertex_main = OpFunction %void None %9
- %29 = OpLabel
- %30 = OpFunctionCall %v4float %vertex_main_inner
- OpStore %value %30
+%vertex_main = OpFunction %void None %14
+ %37 = OpLabel
+ %38 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %38
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
-%fragment_main = OpFunction %void None %9
- %33 = OpLabel
- %34 = OpFunctionCall %void %round_e1bba2
+%fragment_main = OpFunction %void None %14
+ %41 = OpLabel
+ %42 = OpFunctionCall %void %round_e1bba2
OpReturn
OpFunctionEnd
-%compute_main = OpFunction %void None %9
- %36 = OpLabel
- %37 = OpFunctionCall %void %round_e1bba2
+%compute_main = OpFunction %void None %14
+ %44 = OpLabel
+ %45 = OpFunctionCall %void %round_e1bba2
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/round/e1bba2.wgsl.expected.wgsl b/test/tint/builtins/gen/var/round/e1bba2.wgsl.expected.wgsl
index 7cce6c4..75cf80a 100644
--- a/test/tint/builtins/gen/var/round/e1bba2.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/var/round/e1bba2.wgsl.expected.wgsl
@@ -3,8 +3,11 @@
fn round_e1bba2() {
var arg_0 = vec3<f16>(3.5h);
var res : vec3<f16> = round(arg_0);
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec3<f16>;
+
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
round_e1bba2();
diff --git a/test/tint/builtins/gen/var/round/f665b5.wgsl b/test/tint/builtins/gen/var/round/f665b5.wgsl
index 3b7fd3a..f4e36db 100644
--- a/test/tint/builtins/gen/var/round/f665b5.wgsl
+++ b/test/tint/builtins/gen/var/round/f665b5.wgsl
@@ -27,7 +27,9 @@
fn round_f665b5() {
var arg_0 = vec4<f16>(3.5h);
var res: vec4<f16> = round(arg_0);
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<f16>;
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
diff --git a/test/tint/builtins/gen/var/round/f665b5.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/round/f665b5.wgsl.expected.dxc.hlsl
index f789e47..86ebb56 100644
--- a/test/tint/builtins/gen/var/round/f665b5.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/var/round/f665b5.wgsl.expected.dxc.hlsl
@@ -1,6 +1,9 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
void round_f665b5() {
vector<float16_t, 4> arg_0 = (float16_t(3.5h)).xxxx;
vector<float16_t, 4> res = round(arg_0);
+ prevent_dce.Store<vector<float16_t, 4> >(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/var/round/f665b5.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/round/f665b5.wgsl.expected.fxc.hlsl
index fd541db..982e87b 100644
--- a/test/tint/builtins/gen/var/round/f665b5.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/var/round/f665b5.wgsl.expected.fxc.hlsl
@@ -1,8 +1,11 @@
SKIP: FAILED
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
void round_f665b5() {
- vector<float16_t, 4> arg_0 = (float16_t(3.3984375h)).xxxx;
+ vector<float16_t, 4> arg_0 = (float16_t(3.5h)).xxxx;
vector<float16_t, 4> res = round(arg_0);
+ prevent_dce.Store<vector<float16_t, 4> >(0u, res);
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/var/round/f665b5.wgsl.expected.glsl b/test/tint/builtins/gen/var/round/f665b5.wgsl.expected.glsl
index 344164d..cfd3fe7 100644
--- a/test/tint/builtins/gen/var/round/f665b5.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/round/f665b5.wgsl.expected.glsl
@@ -1,9 +1,14 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16vec4 inner;
+} prevent_dce;
+
void round_f665b5() {
f16vec4 arg_0 = f16vec4(3.5hf);
f16vec4 res = round(arg_0);
+ prevent_dce.inner = res;
}
vec4 vertex_main() {
@@ -23,9 +28,14 @@
#extension GL_AMD_gpu_shader_half_float : require
precision mediump float;
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16vec4 inner;
+} prevent_dce;
+
void round_f665b5() {
f16vec4 arg_0 = f16vec4(3.5hf);
f16vec4 res = round(arg_0);
+ prevent_dce.inner = res;
}
void fragment_main() {
@@ -39,9 +49,14 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ f16vec4 inner;
+} prevent_dce;
+
void round_f665b5() {
f16vec4 arg_0 = f16vec4(3.5hf);
f16vec4 res = round(arg_0);
+ prevent_dce.inner = res;
}
void compute_main() {
diff --git a/test/tint/builtins/gen/var/round/f665b5.wgsl.expected.msl b/test/tint/builtins/gen/var/round/f665b5.wgsl.expected.msl
index 1e6e527..e8fecd5 100644
--- a/test/tint/builtins/gen/var/round/f665b5.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/round/f665b5.wgsl.expected.msl
@@ -1,34 +1,35 @@
#include <metal_stdlib>
using namespace metal;
-void round_f665b5() {
+void round_f665b5(device half4* const tint_symbol_1) {
half4 arg_0 = half4(3.5h);
half4 res = rint(arg_0);
+ *(tint_symbol_1) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner() {
- round_f665b5();
+float4 vertex_main_inner(device half4* const tint_symbol_2) {
+ round_f665b5(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 half4* 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() {
- round_f665b5();
+fragment void fragment_main(device half4* tint_symbol_4 [[buffer(0)]]) {
+ round_f665b5(tint_symbol_4);
return;
}
-kernel void compute_main() {
- round_f665b5();
+kernel void compute_main(device half4* tint_symbol_5 [[buffer(0)]]) {
+ round_f665b5(tint_symbol_5);
return;
}
diff --git a/test/tint/builtins/gen/var/round/f665b5.wgsl.expected.spvasm b/test/tint/builtins/gen/var/round/f665b5.wgsl.expected.spvasm
index aee30a5..eb1a49a 100644
--- a/test/tint/builtins/gen/var/round/f665b5.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/round/f665b5.wgsl.expected.spvasm
@@ -1,14 +1,14 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 38
+; Bound: 46
; Schema: 0
OpCapability Shader
OpCapability Float16
OpCapability UniformAndStorageBuffer16BitAccess
OpCapability StorageBuffer16BitAccess
OpCapability StorageInputOutput16
- %21 = OpExtInstImport "GLSL.std.450"
+ %24 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -17,6 +17,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 %round_f665b5 "round_f665b5"
OpName %arg_0 "arg_0"
OpName %res "res"
@@ -26,6 +29,10 @@
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
+ OpDecorate %prevent_dce DescriptorSet 2
+ OpDecorate %prevent_dce Binding 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
@@ -34,45 +41,54 @@
%_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
+%prevent_dce_block = OpTypeStruct %v4half
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+ %void = OpTypeVoid
+ %14 = OpTypeFunction %void
%half_0x1_cp_1 = OpConstant %half 0x1.cp+1
- %16 = OpConstantComposite %v4half %half_0x1_cp_1 %half_0x1_cp_1 %half_0x1_cp_1 %half_0x1_cp_1
+ %19 = OpConstantComposite %v4half %half_0x1_cp_1 %half_0x1_cp_1 %half_0x1_cp_1 %half_0x1_cp_1
%_ptr_Function_v4half = OpTypePointer Function %v4half
- %19 = OpConstantNull %v4half
- %24 = OpTypeFunction %v4float
+ %22 = OpConstantNull %v4half
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v4half = OpTypePointer StorageBuffer %v4half
+ %32 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
-%round_f665b5 = OpFunction %void None %9
- %12 = OpLabel
- %arg_0 = OpVariable %_ptr_Function_v4half Function %19
- %res = OpVariable %_ptr_Function_v4half Function %19
- OpStore %arg_0 %16
- %22 = OpLoad %v4half %arg_0
- %20 = OpExtInst %v4half %21 RoundEven %22
- OpStore %res %20
+%round_f665b5 = OpFunction %void None %14
+ %17 = OpLabel
+ %arg_0 = OpVariable %_ptr_Function_v4half Function %22
+ %res = OpVariable %_ptr_Function_v4half Function %22
+ OpStore %arg_0 %19
+ %25 = OpLoad %v4half %arg_0
+ %23 = OpExtInst %v4half %24 RoundEven %25
+ OpStore %res %23
+ %30 = OpAccessChain %_ptr_StorageBuffer_v4half %prevent_dce %uint_0
+ %31 = OpLoad %v4half %res
+ OpStore %30 %31
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %24
- %26 = OpLabel
- %27 = OpFunctionCall %void %round_f665b5
+%vertex_main_inner = OpFunction %v4float None %32
+ %34 = OpLabel
+ %35 = OpFunctionCall %void %round_f665b5
OpReturnValue %5
OpFunctionEnd
-%vertex_main = OpFunction %void None %9
- %29 = OpLabel
- %30 = OpFunctionCall %v4float %vertex_main_inner
- OpStore %value %30
+%vertex_main = OpFunction %void None %14
+ %37 = OpLabel
+ %38 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %38
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
-%fragment_main = OpFunction %void None %9
- %33 = OpLabel
- %34 = OpFunctionCall %void %round_f665b5
+%fragment_main = OpFunction %void None %14
+ %41 = OpLabel
+ %42 = OpFunctionCall %void %round_f665b5
OpReturn
OpFunctionEnd
-%compute_main = OpFunction %void None %9
- %36 = OpLabel
- %37 = OpFunctionCall %void %round_f665b5
+%compute_main = OpFunction %void None %14
+ %44 = OpLabel
+ %45 = OpFunctionCall %void %round_f665b5
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/round/f665b5.wgsl.expected.wgsl b/test/tint/builtins/gen/var/round/f665b5.wgsl.expected.wgsl
index 4773a82..db8d1aa 100644
--- a/test/tint/builtins/gen/var/round/f665b5.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/var/round/f665b5.wgsl.expected.wgsl
@@ -3,8 +3,11 @@
fn round_f665b5() {
var arg_0 = vec4<f16>(3.5h);
var res : vec4<f16> = round(arg_0);
+ prevent_dce = res;
}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<f16>;
+
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
round_f665b5();
