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>

test/tint/builtins/gen/gen.wgsl.tmpl

@@ -145,8 +145,23 @@ fn {{$permutation}}() {
 {{-     if $i -}}, {{end}}{{$args.Get $i -}}
 {{-   end -}}
   );
+{{- /* Ensure the call isn't dead-code eliminated */ -}}
+{{-   if and $overload.ReturnType (IsDeclarable $overload.ReturnType)}}
+{{-     if IsHostShareable $overload.ReturnType}}
+  prevent_dce = res;
+{{-     else}}
+  prevent_dce = select(0, 1, all(res == {{template "Type" $overload.ReturnType}}()));
+{{-     end}}
+{{-   end}}
 }
-{{/*new line*/ -}}
+
+{{-   if and $overload.ReturnType (IsDeclarable $overload.ReturnType)}}
+{{-     if IsHostShareable $overload.ReturnType}}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : {{template "Type" $overload.ReturnType}};
+{{-     else}}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : i32;
+{{      end}}
+{{    end}}
 
 {{- if $overload.CanBeUsedInStage.Vertex }}
 @vertex

test/tint/builtins/gen/literal/abs/002533.wgsl

@@ -24,7 +24,9 @@
 // fn abs(vec<4, f32>) -> vec<4, f32>
 fn abs_002533() {
   var res: vec4<f32> = abs(vec4<f32>(1.f));
+  prevent_dce = res;
 }
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<f32>;
 
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {

test/tint/builtins/gen/literal/abs/002533.wgsl.expected.dxc.hlsl

@@ -1,5 +1,8 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
 void abs_002533() {
   float4 res = (1.0f).xxxx;
+  prevent_dce.Store4(0u, asuint(res));
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/002533.wgsl.expected.fxc.hlsl

@@ -1,5 +1,8 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
 void abs_002533() {
   float4 res = (1.0f).xxxx;
+  prevent_dce.Store4(0u, asuint(res));
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/002533.wgsl.expected.glsl

@@ -1,7 +1,12 @@
 #version 310 es
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  vec4 inner;
+} prevent_dce;
+
 void abs_002533() {
   vec4 res = vec4(1.0f);
+  prevent_dce.inner = res;
 }
 
 vec4 vertex_main() {
@@ -20,8 +25,13 @@ void main() {
 #version 310 es
 precision mediump float;
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  vec4 inner;
+} prevent_dce;
+
 void abs_002533() {
   vec4 res = vec4(1.0f);
+  prevent_dce.inner = res;
 }
 
 void fragment_main() {
@@ -34,8 +44,13 @@ void main() {
 }
 #version 310 es
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  vec4 inner;
+} prevent_dce;
+
 void abs_002533() {
   vec4 res = vec4(1.0f);
+  prevent_dce.inner = res;
 }
 
 void compute_main() {

test/tint/builtins/gen/literal/abs/002533.wgsl.expected.msl

@@ -1,33 +1,34 @@
 #include <metal_stdlib>
 
 using namespace metal;
-void abs_002533() {
+void abs_002533(device float4* const tint_symbol_1) {
   float4 res = float4(1.0f);
+  *(tint_symbol_1) = res;
 }
 
 struct tint_symbol {
   float4 value [[position]];
 };
 
-float4 vertex_main_inner() {
+float4 vertex_main_inner(device float4* const tint_symbol_2) {
-  abs_002533();
+  abs_002533(tint_symbol_2);
   return float4(0.0f);
 }
 
-vertex tint_symbol vertex_main() {
+vertex tint_symbol vertex_main(device float4* tint_symbol_3 [[buffer(0)]]) {
-  float4 const inner_result = vertex_main_inner();
+  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() {
+fragment void fragment_main(device float4* tint_symbol_4 [[buffer(0)]]) {
-  abs_002533();
+  abs_002533(tint_symbol_4);
   return;
 }
 
-kernel void compute_main() {
+kernel void compute_main(device float4* tint_symbol_5 [[buffer(0)]]) {
-  abs_002533();
+  abs_002533(tint_symbol_5);
   return;
 }
 

test/tint/builtins/gen/literal/abs/002533.wgsl.expected.spvasm

@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 30
+; Bound: 38
 ; 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 %abs_002533 "abs_002533"
                OpName %res "res"
                OpName %vertex_main_inner "vertex_main_inner"
@@ -20,6 +23,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
@@ -28,37 +35,46 @@
 %_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_1 = OpConstant %float 1
-         %14 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
+         %17 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
 %_ptr_Function_v4float = OpTypePointer Function %v4float
-         %17 = OpTypeFunction %v4float
+       %uint = OpTypeInt 32 0
- %abs_002533 = OpFunction %void None %9
+     %uint_0 = OpConstant %uint 0
-         %12 = OpLabel
+%_ptr_StorageBuffer_v4float = OpTypePointer StorageBuffer %v4float
+         %25 = OpTypeFunction %v4float
+ %abs_002533 = OpFunction %void None %12
+         %15 = OpLabel
         %res = OpVariable %_ptr_Function_v4float Function %5
-               OpStore %res %14
+               OpStore %res %17
+         %23 = OpAccessChain %_ptr_StorageBuffer_v4float %prevent_dce %uint_0
+         %24 = OpLoad %v4float %res
+               OpStore %23 %24
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %17
+%vertex_main_inner = OpFunction %v4float None %25
-         %19 = OpLabel
+         %27 = OpLabel
-         %20 = OpFunctionCall %void %abs_002533
+         %28 = OpFunctionCall %void %abs_002533
                OpReturnValue %5
                OpFunctionEnd
-%vertex_main = OpFunction %void None %9
+%vertex_main = OpFunction %void None %12
-         %22 = OpLabel
+         %30 = OpLabel
-         %23 = OpFunctionCall %v4float %vertex_main_inner
+         %31 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %23
+               OpStore %value %31
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
-%fragment_main = OpFunction %void None %9
+%fragment_main = OpFunction %void None %12
-         %25 = OpLabel
+         %33 = OpLabel
-         %26 = OpFunctionCall %void %abs_002533
+         %34 = OpFunctionCall %void %abs_002533
                OpReturn
                OpFunctionEnd
-%compute_main = OpFunction %void None %9
+%compute_main = OpFunction %void None %12
-         %28 = OpLabel
+         %36 = OpLabel
-         %29 = OpFunctionCall %void %abs_002533
+         %37 = OpFunctionCall %void %abs_002533
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/abs/002533.wgsl.expected.wgsl

@@ -1,7 +1,10 @@
 fn abs_002533() {
   var res : vec4<f32> = abs(vec4<f32>(1.0f));
+  prevent_dce = res;
 }
 
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<f32>;
+
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   abs_002533();

test/tint/builtins/gen/literal/abs/005174.wgsl

@@ -24,7 +24,9 @@
 // fn abs(vec<3, f32>) -> vec<3, f32>
 fn abs_005174() {
   var res: vec3<f32> = abs(vec3<f32>(1.f));
+  prevent_dce = res;
 }
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec3<f32>;
 
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {

test/tint/builtins/gen/literal/abs/005174.wgsl.expected.dxc.hlsl

@@ -1,5 +1,8 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
 void abs_005174() {
   float3 res = (1.0f).xxx;
+  prevent_dce.Store3(0u, asuint(res));
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/005174.wgsl.expected.fxc.hlsl

@@ -1,5 +1,8 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
 void abs_005174() {
   float3 res = (1.0f).xxx;
+  prevent_dce.Store3(0u, asuint(res));
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/005174.wgsl.expected.glsl

@@ -1,7 +1,12 @@
 #version 310 es
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  vec3 inner;
+} prevent_dce;
+
 void abs_005174() {
   vec3 res = vec3(1.0f);
+  prevent_dce.inner = res;
 }
 
 vec4 vertex_main() {
@@ -20,8 +25,13 @@ void main() {
 #version 310 es
 precision mediump float;
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  vec3 inner;
+} prevent_dce;
+
 void abs_005174() {
   vec3 res = vec3(1.0f);
+  prevent_dce.inner = res;
 }
 
 void fragment_main() {
@@ -34,8 +44,13 @@ void main() {
 }
 #version 310 es
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  vec3 inner;
+} prevent_dce;
+
 void abs_005174() {
   vec3 res = vec3(1.0f);
+  prevent_dce.inner = res;
 }
 
 void compute_main() {

test/tint/builtins/gen/literal/abs/005174.wgsl.expected.msl

@@ -1,33 +1,34 @@
 #include <metal_stdlib>
 
 using namespace metal;
-void abs_005174() {
+void abs_005174(device packed_float3* const tint_symbol_1) {
   float3 res = float3(1.0f);
+  *(tint_symbol_1) = packed_float3(res);
 }
 
 struct tint_symbol {
   float4 value [[position]];
 };
 
-float4 vertex_main_inner() {
+float4 vertex_main_inner(device packed_float3* const tint_symbol_2) {
-  abs_005174();
+  abs_005174(tint_symbol_2);
   return float4(0.0f);
 }
 
-vertex tint_symbol vertex_main() {
+vertex tint_symbol vertex_main(device packed_float3* tint_symbol_3 [[buffer(0)]]) {
-  float4 const inner_result = vertex_main_inner();
+  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() {
+fragment void fragment_main(device packed_float3* tint_symbol_4 [[buffer(0)]]) {
-  abs_005174();
+  abs_005174(tint_symbol_4);
   return;
 }
 
-kernel void compute_main() {
+kernel void compute_main(device packed_float3* tint_symbol_5 [[buffer(0)]]) {
-  abs_005174();
+  abs_005174(tint_symbol_5);
   return;
 }
 

test/tint/builtins/gen/literal/abs/005174.wgsl.expected.spvasm

@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 32
+; Bound: 40
 ; 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 %abs_005174 "abs_005174"
                OpName %res "res"
                OpName %vertex_main_inner "vertex_main_inner"
@@ -20,6 +23,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
@@ -28,39 +35,48 @@
 %_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_1 = OpConstant %float 1
-         %15 = OpConstantComposite %v3float %float_1 %float_1 %float_1
+         %18 = OpConstantComposite %v3float %float_1 %float_1 %float_1
 %_ptr_Function_v3float = OpTypePointer Function %v3float
-         %18 = OpConstantNull %v3float
+         %21 = OpConstantNull %v3float
-         %19 = OpTypeFunction %v4float
+       %uint = OpTypeInt 32 0
- %abs_005174 = OpFunction %void None %9
+     %uint_0 = OpConstant %uint 0
-         %12 = OpLabel
+%_ptr_StorageBuffer_v3float = OpTypePointer StorageBuffer %v3float
-        %res = OpVariable %_ptr_Function_v3float Function %18
+         %27 = OpTypeFunction %v4float
-               OpStore %res %15
+ %abs_005174 = OpFunction %void None %13
+         %16 = OpLabel
+        %res = OpVariable %_ptr_Function_v3float Function %21
+               OpStore %res %18
+         %25 = OpAccessChain %_ptr_StorageBuffer_v3float %prevent_dce %uint_0
+         %26 = OpLoad %v3float %res
+               OpStore %25 %26
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %19
+%vertex_main_inner = OpFunction %v4float None %27
-         %21 = OpLabel
+         %29 = OpLabel
-         %22 = OpFunctionCall %void %abs_005174
+         %30 = OpFunctionCall %void %abs_005174
                OpReturnValue %5
                OpFunctionEnd
-%vertex_main = OpFunction %void None %9
+%vertex_main = OpFunction %void None %13
-         %24 = OpLabel
+         %32 = OpLabel
-         %25 = OpFunctionCall %v4float %vertex_main_inner
+         %33 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %25
+               OpStore %value %33
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
-%fragment_main = OpFunction %void None %9
+%fragment_main = OpFunction %void None %13
-         %27 = OpLabel
+         %35 = OpLabel
-         %28 = OpFunctionCall %void %abs_005174
+         %36 = OpFunctionCall %void %abs_005174
                OpReturn
                OpFunctionEnd
-%compute_main = OpFunction %void None %9
+%compute_main = OpFunction %void None %13
-         %30 = OpLabel
+         %38 = OpLabel
-         %31 = OpFunctionCall %void %abs_005174
+         %39 = OpFunctionCall %void %abs_005174
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/abs/005174.wgsl.expected.wgsl

@@ -1,7 +1,10 @@
 fn abs_005174() {
   var res : vec3<f32> = abs(vec3<f32>(1.0f));
+  prevent_dce = res;
 }
 
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec3<f32>;
+
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   abs_005174();

test/tint/builtins/gen/literal/abs/1ce782.wgsl

@@ -24,7 +24,9 @@
 // fn abs(vec<4, u32>) -> vec<4, u32>
 fn abs_1ce782() {
   var res: vec4<u32> = abs(vec4<u32>(1u));
+  prevent_dce = res;
 }
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
 
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {

test/tint/builtins/gen/literal/abs/1ce782.wgsl.expected.dxc.hlsl

@@ -1,5 +1,8 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
 void abs_1ce782() {
   uint4 res = (1u).xxxx;
+  prevent_dce.Store4(0u, asuint(res));
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/1ce782.wgsl.expected.fxc.hlsl

@@ -1,5 +1,8 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
 void abs_1ce782() {
   uint4 res = (1u).xxxx;
+  prevent_dce.Store4(0u, asuint(res));
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/1ce782.wgsl.expected.glsl

@@ -1,7 +1,12 @@
 #version 310 es
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  uvec4 inner;
+} prevent_dce;
+
 void abs_1ce782() {
   uvec4 res = uvec4(1u);
+  prevent_dce.inner = res;
 }
 
 vec4 vertex_main() {
@@ -20,8 +25,13 @@ void main() {
 #version 310 es
 precision mediump float;
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  uvec4 inner;
+} prevent_dce;
+
 void abs_1ce782() {
   uvec4 res = uvec4(1u);
+  prevent_dce.inner = res;
 }
 
 void fragment_main() {
@@ -34,8 +44,13 @@ void main() {
 }
 #version 310 es
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  uvec4 inner;
+} prevent_dce;
+
 void abs_1ce782() {
   uvec4 res = uvec4(1u);
+  prevent_dce.inner = res;
 }
 
 void compute_main() {

test/tint/builtins/gen/literal/abs/1ce782.wgsl.expected.msl

@@ -1,33 +1,34 @@
 #include <metal_stdlib>
 
 using namespace metal;
-void abs_1ce782() {
+void abs_1ce782(device uint4* const tint_symbol_1) {
   uint4 res = uint4(1u);
+  *(tint_symbol_1) = res;
 }
 
 struct tint_symbol {
   float4 value [[position]];
 };
 
-float4 vertex_main_inner() {
+float4 vertex_main_inner(device uint4* const tint_symbol_2) {
-  abs_1ce782();
+  abs_1ce782(tint_symbol_2);
   return float4(0.0f);
 }
 
-vertex tint_symbol vertex_main() {
+vertex tint_symbol vertex_main(device uint4* tint_symbol_3 [[buffer(0)]]) {
-  float4 const inner_result = vertex_main_inner();
+  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() {
+fragment void fragment_main(device uint4* tint_symbol_4 [[buffer(0)]]) {
-  abs_1ce782();
+  abs_1ce782(tint_symbol_4);
   return;
 }
 
-kernel void compute_main() {
+kernel void compute_main(device uint4* tint_symbol_5 [[buffer(0)]]) {
-  abs_1ce782();
+  abs_1ce782(tint_symbol_5);
   return;
 }
 

test/tint/builtins/gen/literal/abs/1ce782.wgsl.expected.spvasm

@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 34
+; 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 %abs_1ce782 "abs_1ce782"
                OpName %res "res"
                OpName %vertex_main_inner "vertex_main_inner"
@@ -20,6 +23,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
@@ -28,41 +35,49 @@
 %_ptr_Output_float = OpTypePointer Output %float
           %8 = OpConstantNull %float
 %vertex_point_size = OpVariable %_ptr_Output_float Output %8
-       %void = OpTypeVoid
-          %9 = OpTypeFunction %void
        %uint = OpTypeInt 32 0
      %v4uint = OpTypeVector %uint 4
+%prevent_dce_block = OpTypeStruct %v4uint
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+       %void = OpTypeVoid
+         %14 = OpTypeFunction %void
      %uint_1 = OpConstant %uint 1
-         %16 = OpConstantComposite %v4uint %uint_1 %uint_1 %uint_1 %uint_1
+         %19 = OpConstantComposite %v4uint %uint_1 %uint_1 %uint_1 %uint_1
 %_ptr_Function_v4uint = OpTypePointer Function %v4uint
-         %19 = OpConstantNull %v4uint
+         %22 = OpConstantNull %v4uint
-         %20 = OpTypeFunction %v4float
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v4uint = OpTypePointer StorageBuffer %v4uint
+         %27 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
- %abs_1ce782 = OpFunction %void None %9
+ %abs_1ce782 = OpFunction %void None %14
-         %12 = OpLabel
+         %17 = OpLabel
-        %res = OpVariable %_ptr_Function_v4uint Function %19
+        %res = OpVariable %_ptr_Function_v4uint Function %22
-               OpStore %res %16
+               OpStore %res %19
+         %25 = OpAccessChain %_ptr_StorageBuffer_v4uint %prevent_dce %uint_0
+         %26 = OpLoad %v4uint %res
+               OpStore %25 %26
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %20
+%vertex_main_inner = OpFunction %v4float None %27
-         %22 = OpLabel
+         %29 = OpLabel
-         %23 = OpFunctionCall %void %abs_1ce782
+         %30 = OpFunctionCall %void %abs_1ce782
                OpReturnValue %5
                OpFunctionEnd
-%vertex_main = OpFunction %void None %9
+%vertex_main = OpFunction %void None %14
-         %25 = OpLabel
+         %32 = OpLabel
-         %26 = OpFunctionCall %v4float %vertex_main_inner
+         %33 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %26
+               OpStore %value %33
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
-%fragment_main = OpFunction %void None %9
+%fragment_main = OpFunction %void None %14
-         %29 = OpLabel
+         %36 = OpLabel
-         %30 = OpFunctionCall %void %abs_1ce782
+         %37 = OpFunctionCall %void %abs_1ce782
                OpReturn
                OpFunctionEnd
-%compute_main = OpFunction %void None %9
+%compute_main = OpFunction %void None %14
-         %32 = OpLabel
+         %39 = OpLabel
-         %33 = OpFunctionCall %void %abs_1ce782
+         %40 = OpFunctionCall %void %abs_1ce782
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/abs/1ce782.wgsl.expected.wgsl

@@ -1,7 +1,10 @@
 fn abs_1ce782() {
   var res : vec4<u32> = abs(vec4<u32>(1u));
+  prevent_dce = res;
 }
 
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
+
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   abs_1ce782();

test/tint/builtins/gen/literal/abs/1e9d53.wgsl

@@ -24,7 +24,9 @@
 // fn abs(vec<2, f32>) -> vec<2, f32>
 fn abs_1e9d53() {
   var res: vec2<f32> = abs(vec2<f32>(1.f));
+  prevent_dce = res;
 }
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec2<f32>;
 
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {

test/tint/builtins/gen/literal/abs/1e9d53.wgsl.expected.dxc.hlsl

@@ -1,5 +1,8 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
 void abs_1e9d53() {
   float2 res = (1.0f).xx;
+  prevent_dce.Store2(0u, asuint(res));
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/1e9d53.wgsl.expected.fxc.hlsl

@@ -1,5 +1,8 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
 void abs_1e9d53() {
   float2 res = (1.0f).xx;
+  prevent_dce.Store2(0u, asuint(res));
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/1e9d53.wgsl.expected.glsl

@@ -1,7 +1,12 @@
 #version 310 es
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  vec2 inner;
+} prevent_dce;
+
 void abs_1e9d53() {
   vec2 res = vec2(1.0f);
+  prevent_dce.inner = res;
 }
 
 vec4 vertex_main() {
@@ -20,8 +25,13 @@ void main() {
 #version 310 es
 precision mediump float;
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  vec2 inner;
+} prevent_dce;
+
 void abs_1e9d53() {
   vec2 res = vec2(1.0f);
+  prevent_dce.inner = res;
 }
 
 void fragment_main() {
@@ -34,8 +44,13 @@ void main() {
 }
 #version 310 es
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  vec2 inner;
+} prevent_dce;
+
 void abs_1e9d53() {
   vec2 res = vec2(1.0f);
+  prevent_dce.inner = res;
 }
 
 void compute_main() {

test/tint/builtins/gen/literal/abs/1e9d53.wgsl.expected.msl

@@ -1,33 +1,34 @@
 #include <metal_stdlib>
 
 using namespace metal;
-void abs_1e9d53() {
+void abs_1e9d53(device float2* const tint_symbol_1) {
   float2 res = float2(1.0f);
+  *(tint_symbol_1) = res;
 }
 
 struct tint_symbol {
   float4 value [[position]];
 };
 
-float4 vertex_main_inner() {
+float4 vertex_main_inner(device float2* const tint_symbol_2) {
-  abs_1e9d53();
+  abs_1e9d53(tint_symbol_2);
   return float4(0.0f);
 }
 
-vertex tint_symbol vertex_main() {
+vertex tint_symbol vertex_main(device float2* tint_symbol_3 [[buffer(0)]]) {
-  float4 const inner_result = vertex_main_inner();
+  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() {
+fragment void fragment_main(device float2* tint_symbol_4 [[buffer(0)]]) {
-  abs_1e9d53();
+  abs_1e9d53(tint_symbol_4);
   return;
 }
 
-kernel void compute_main() {
+kernel void compute_main(device float2* tint_symbol_5 [[buffer(0)]]) {
-  abs_1e9d53();
+  abs_1e9d53(tint_symbol_5);
   return;
 }
 

test/tint/builtins/gen/literal/abs/1e9d53.wgsl.expected.spvasm

@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 32
+; Bound: 40
 ; 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 %abs_1e9d53 "abs_1e9d53"
                OpName %res "res"
                OpName %vertex_main_inner "vertex_main_inner"
@@ -20,6 +23,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
@@ -28,39 +35,48 @@
 %_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_1 = OpConstant %float 1
-         %15 = OpConstantComposite %v2float %float_1 %float_1
+         %18 = OpConstantComposite %v2float %float_1 %float_1
 %_ptr_Function_v2float = OpTypePointer Function %v2float
-         %18 = OpConstantNull %v2float
+         %21 = OpConstantNull %v2float
-         %19 = OpTypeFunction %v4float
+       %uint = OpTypeInt 32 0
- %abs_1e9d53 = OpFunction %void None %9
+     %uint_0 = OpConstant %uint 0
-         %12 = OpLabel
+%_ptr_StorageBuffer_v2float = OpTypePointer StorageBuffer %v2float
-        %res = OpVariable %_ptr_Function_v2float Function %18
+         %27 = OpTypeFunction %v4float
-               OpStore %res %15
+ %abs_1e9d53 = OpFunction %void None %13
+         %16 = OpLabel
+        %res = OpVariable %_ptr_Function_v2float Function %21
+               OpStore %res %18
+         %25 = OpAccessChain %_ptr_StorageBuffer_v2float %prevent_dce %uint_0
+         %26 = OpLoad %v2float %res
+               OpStore %25 %26
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %19
+%vertex_main_inner = OpFunction %v4float None %27
-         %21 = OpLabel
+         %29 = OpLabel
-         %22 = OpFunctionCall %void %abs_1e9d53
+         %30 = OpFunctionCall %void %abs_1e9d53
                OpReturnValue %5
                OpFunctionEnd
-%vertex_main = OpFunction %void None %9
+%vertex_main = OpFunction %void None %13
-         %24 = OpLabel
+         %32 = OpLabel
-         %25 = OpFunctionCall %v4float %vertex_main_inner
+         %33 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %25
+               OpStore %value %33
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
-%fragment_main = OpFunction %void None %9
+%fragment_main = OpFunction %void None %13
-         %27 = OpLabel
+         %35 = OpLabel
-         %28 = OpFunctionCall %void %abs_1e9d53
+         %36 = OpFunctionCall %void %abs_1e9d53
                OpReturn
                OpFunctionEnd
-%compute_main = OpFunction %void None %9
+%compute_main = OpFunction %void None %13
-         %30 = OpLabel
+         %38 = OpLabel
-         %31 = OpFunctionCall %void %abs_1e9d53
+         %39 = OpFunctionCall %void %abs_1e9d53
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/abs/1e9d53.wgsl.expected.wgsl

@@ -1,7 +1,10 @@
 fn abs_1e9d53() {
   var res : vec2<f32> = abs(vec2<f32>(1.0f));
+  prevent_dce = res;
 }
 
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec2<f32>;
+
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   abs_1e9d53();

test/tint/builtins/gen/literal/abs/2f861b.wgsl

@@ -25,7 +25,6 @@
 fn abs_2f861b() {
   var res = abs(vec3(1.));
 }
-
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   abs_2f861b();

test/tint/builtins/gen/literal/abs/421ca3.wgsl

@@ -26,7 +26,9 @@ enable f16;
 // fn abs(vec<3, f16>) -> vec<3, f16>
 fn abs_421ca3() {
   var res: vec3<f16> = abs(vec3<f16>(1.h));
+  prevent_dce = res;
 }
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec3<f16>;
 
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {

test/tint/builtins/gen/literal/abs/421ca3.wgsl.expected.dxc.hlsl

@@ -1,5 +1,8 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
 void abs_421ca3() {
   vector<float16_t, 3> res = (float16_t(1.0h)).xxx;
+  prevent_dce.Store<vector<float16_t, 3> >(0u, res);
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/421ca3.wgsl.expected.fxc.hlsl

@@ -1,7 +1,10 @@
 SKIP: FAILED
 
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
 void abs_421ca3() {
   vector<float16_t, 3> res = (float16_t(1.0h)).xxx;
+  prevent_dce.Store<vector<float16_t, 3> >(0u, res);
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/421ca3.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 {
+  f16vec3 inner;
+} prevent_dce;
+
 void abs_421ca3() {
   f16vec3 res = f16vec3(1.0hf);
+  prevent_dce.inner = res;
 }
 
 vec4 vertex_main() {
@@ -22,8 +27,13 @@ void main() {
 #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 abs_421ca3() {
   f16vec3 res = f16vec3(1.0hf);
+  prevent_dce.inner = res;
 }
 
 void fragment_main() {
@@ -37,8 +47,13 @@ void main() {
 #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 abs_421ca3() {
   f16vec3 res = f16vec3(1.0hf);
+  prevent_dce.inner = res;
 }
 
 void compute_main() {

test/tint/builtins/gen/literal/abs/421ca3.wgsl.expected.msl

@@ -1,33 +1,34 @@
 #include <metal_stdlib>
 
 using namespace metal;
-void abs_421ca3() {
+void abs_421ca3(device packed_half3* const tint_symbol_1) {
   half3 res = half3(1.0h);
+  *(tint_symbol_1) = packed_half3(res);
 }
 
 struct tint_symbol {
   float4 value [[position]];
 };
 
-float4 vertex_main_inner() {
+float4 vertex_main_inner(device packed_half3* const tint_symbol_2) {
-  abs_421ca3();
+  abs_421ca3(tint_symbol_2);
   return float4(0.0f);
 }
 
-vertex tint_symbol vertex_main() {
+vertex tint_symbol vertex_main(device packed_half3* tint_symbol_3 [[buffer(0)]]) {
-  float4 const inner_result = vertex_main_inner();
+  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() {
+fragment void fragment_main(device packed_half3* tint_symbol_4 [[buffer(0)]]) {
-  abs_421ca3();
+  abs_421ca3(tint_symbol_4);
   return;
 }
 
-kernel void compute_main() {
+kernel void compute_main(device packed_half3* tint_symbol_5 [[buffer(0)]]) {
-  abs_421ca3();
+  abs_421ca3(tint_symbol_5);
   return;
 }
 

test/tint/builtins/gen/literal/abs/421ca3.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
                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 %abs_421ca3 "abs_421ca3"
                OpName %res "res"
                OpName %vertex_main_inner "vertex_main_inner"
@@ -24,6 +27,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
@@ -32,41 +39,50 @@
 %_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_0x1p_0 = OpConstant %half 0x1p+0
-         %16 = OpConstantComposite %v3half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
+         %19 = OpConstantComposite %v3half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
 %_ptr_Function_v3half = OpTypePointer Function %v3half
-         %19 = OpConstantNull %v3half
+         %22 = OpConstantNull %v3half
-         %20 = OpTypeFunction %v4float
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v3half = OpTypePointer StorageBuffer %v3half
+         %28 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
- %abs_421ca3 = OpFunction %void None %9
+ %abs_421ca3 = OpFunction %void None %14
-         %12 = OpLabel
+         %17 = OpLabel
-        %res = OpVariable %_ptr_Function_v3half Function %19
+        %res = OpVariable %_ptr_Function_v3half Function %22
-               OpStore %res %16
+               OpStore %res %19
+         %26 = OpAccessChain %_ptr_StorageBuffer_v3half %prevent_dce %uint_0
+         %27 = OpLoad %v3half %res
+               OpStore %26 %27
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %20
+%vertex_main_inner = OpFunction %v4float None %28
-         %22 = OpLabel
+         %30 = OpLabel
-         %23 = OpFunctionCall %void %abs_421ca3
+         %31 = OpFunctionCall %void %abs_421ca3
                OpReturnValue %5
                OpFunctionEnd
-%vertex_main = OpFunction %void None %9
+%vertex_main = OpFunction %void None %14
-         %25 = OpLabel
+         %33 = OpLabel
-         %26 = OpFunctionCall %v4float %vertex_main_inner
+         %34 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %26
+               OpStore %value %34
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
-%fragment_main = OpFunction %void None %9
+%fragment_main = OpFunction %void None %14
-         %29 = OpLabel
+         %37 = OpLabel
-         %30 = OpFunctionCall %void %abs_421ca3
+         %38 = OpFunctionCall %void %abs_421ca3
                OpReturn
                OpFunctionEnd
-%compute_main = OpFunction %void None %9
+%compute_main = OpFunction %void None %14
-         %32 = OpLabel
+         %40 = OpLabel
-         %33 = OpFunctionCall %void %abs_421ca3
+         %41 = OpFunctionCall %void %abs_421ca3
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/abs/421ca3.wgsl.expected.wgsl

@@ -2,8 +2,11 @@ enable f16;
 
 fn abs_421ca3() {
   var res : vec3<f16> = abs(vec3<f16>(1.0h));
+  prevent_dce = res;
 }
 
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec3<f16>;
+
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   abs_421ca3();

test/tint/builtins/gen/literal/abs/467cd1.wgsl

@@ -24,7 +24,9 @@
 // fn abs(u32) -> u32
 fn abs_467cd1() {
   var res: u32 = abs(1u);
+  prevent_dce = res;
 }
+@group(2) @binding(0) var<storage, read_write> prevent_dce : u32;
 
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {

test/tint/builtins/gen/literal/abs/467cd1.wgsl.expected.dxc.hlsl

@@ -1,5 +1,8 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
 void abs_467cd1() {
   uint res = 1u;
+  prevent_dce.Store(0u, asuint(res));
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/467cd1.wgsl.expected.fxc.hlsl

@@ -1,5 +1,8 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
 void abs_467cd1() {
   uint res = 1u;
+  prevent_dce.Store(0u, asuint(res));
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/467cd1.wgsl.expected.glsl

@@ -1,7 +1,12 @@
 #version 310 es
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  uint inner;
+} prevent_dce;
+
 void abs_467cd1() {
   uint res = 1u;
+  prevent_dce.inner = res;
 }
 
 vec4 vertex_main() {
@@ -20,8 +25,13 @@ void main() {
 #version 310 es
 precision mediump float;
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  uint inner;
+} prevent_dce;
+
 void abs_467cd1() {
   uint res = 1u;
+  prevent_dce.inner = res;
 }
 
 void fragment_main() {
@@ -34,8 +44,13 @@ void main() {
 }
 #version 310 es
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  uint inner;
+} prevent_dce;
+
 void abs_467cd1() {
   uint res = 1u;
+  prevent_dce.inner = res;
 }
 
 void compute_main() {

test/tint/builtins/gen/literal/abs/467cd1.wgsl.expected.msl

@@ -1,33 +1,34 @@
 #include <metal_stdlib>
 
 using namespace metal;
-void abs_467cd1() {
+void abs_467cd1(device uint* const tint_symbol_1) {
   uint res = 1u;
+  *(tint_symbol_1) = res;
 }
 
 struct tint_symbol {
   float4 value [[position]];
 };
 
-float4 vertex_main_inner() {
+float4 vertex_main_inner(device uint* const tint_symbol_2) {
-  abs_467cd1();
+  abs_467cd1(tint_symbol_2);
   return float4(0.0f);
 }
 
-vertex tint_symbol vertex_main() {
+vertex tint_symbol vertex_main(device uint* tint_symbol_3 [[buffer(0)]]) {
-  float4 const inner_result = vertex_main_inner();
+  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() {
+fragment void fragment_main(device uint* tint_symbol_4 [[buffer(0)]]) {
-  abs_467cd1();
+  abs_467cd1(tint_symbol_4);
   return;
 }
 
-kernel void compute_main() {
+kernel void compute_main(device uint* tint_symbol_5 [[buffer(0)]]) {
-  abs_467cd1();
+  abs_467cd1(tint_symbol_5);
   return;
 }
 

test/tint/builtins/gen/literal/abs/467cd1.wgsl.expected.spvasm

@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 32
+; Bound: 39
 ; 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 %abs_467cd1 "abs_467cd1"
                OpName %res "res"
                OpName %vertex_main_inner "vertex_main_inner"
@@ -20,6 +23,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
@@ -28,39 +35,47 @@
 %_ptr_Output_float = OpTypePointer Output %float
           %8 = OpConstantNull %float
 %vertex_point_size = OpVariable %_ptr_Output_float Output %8
-       %void = OpTypeVoid
-          %9 = OpTypeFunction %void
        %uint = OpTypeInt 32 0
+%prevent_dce_block = OpTypeStruct %uint
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+       %void = OpTypeVoid
+         %13 = OpTypeFunction %void
      %uint_1 = OpConstant %uint 1
 %_ptr_Function_uint = OpTypePointer Function %uint
-         %17 = OpConstantNull %uint
+         %20 = OpConstantNull %uint
-         %18 = OpTypeFunction %v4float
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+         %25 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
- %abs_467cd1 = OpFunction %void None %9
+ %abs_467cd1 = OpFunction %void None %13
-         %12 = OpLabel
+         %16 = OpLabel
-        %res = OpVariable %_ptr_Function_uint Function %17
+        %res = OpVariable %_ptr_Function_uint Function %20
                OpStore %res %uint_1
+         %23 = OpAccessChain %_ptr_StorageBuffer_uint %prevent_dce %uint_0
+         %24 = OpLoad %uint %res
+               OpStore %23 %24
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %18
+%vertex_main_inner = OpFunction %v4float None %25
-         %20 = OpLabel
+         %27 = OpLabel
-         %21 = OpFunctionCall %void %abs_467cd1
+         %28 = OpFunctionCall %void %abs_467cd1
                OpReturnValue %5
                OpFunctionEnd
-%vertex_main = OpFunction %void None %9
+%vertex_main = OpFunction %void None %13
-         %23 = OpLabel
+         %30 = OpLabel
-         %24 = OpFunctionCall %v4float %vertex_main_inner
+         %31 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %24
+               OpStore %value %31
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
-%fragment_main = OpFunction %void None %9
+%fragment_main = OpFunction %void None %13
-         %27 = OpLabel
+         %34 = OpLabel
-         %28 = OpFunctionCall %void %abs_467cd1
+         %35 = OpFunctionCall %void %abs_467cd1
                OpReturn
                OpFunctionEnd
-%compute_main = OpFunction %void None %9
+%compute_main = OpFunction %void None %13
-         %30 = OpLabel
+         %37 = OpLabel
-         %31 = OpFunctionCall %void %abs_467cd1
+         %38 = OpFunctionCall %void %abs_467cd1
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/abs/467cd1.wgsl.expected.wgsl

@@ -1,7 +1,10 @@
 fn abs_467cd1() {
   var res : u32 = abs(1u);
+  prevent_dce = res;
 }
 
+@group(2) @binding(0) var<storage, read_write> prevent_dce : u32;
+
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   abs_467cd1();

test/tint/builtins/gen/literal/abs/4ad288.wgsl

@@ -24,7 +24,9 @@
 // fn abs(i32) -> i32
 fn abs_4ad288() {
   var res: i32 = abs(1i);
+  prevent_dce = res;
 }
+@group(2) @binding(0) var<storage, read_write> prevent_dce : i32;
 
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {

test/tint/builtins/gen/literal/abs/4ad288.wgsl.expected.dxc.hlsl

@@ -1,5 +1,8 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
 void abs_4ad288() {
   int res = 1;
+  prevent_dce.Store(0u, asuint(res));
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/4ad288.wgsl.expected.fxc.hlsl

@@ -1,5 +1,8 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
 void abs_4ad288() {
   int res = 1;
+  prevent_dce.Store(0u, asuint(res));
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/4ad288.wgsl.expected.glsl

@@ -1,7 +1,12 @@
 #version 310 es
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  int inner;
+} prevent_dce;
+
 void abs_4ad288() {
   int res = 1;
+  prevent_dce.inner = res;
 }
 
 vec4 vertex_main() {
@@ -20,8 +25,13 @@ void main() {
 #version 310 es
 precision mediump float;
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  int inner;
+} prevent_dce;
+
 void abs_4ad288() {
   int res = 1;
+  prevent_dce.inner = res;
 }
 
 void fragment_main() {
@@ -34,8 +44,13 @@ void main() {
 }
 #version 310 es
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  int inner;
+} prevent_dce;
+
 void abs_4ad288() {
   int res = 1;
+  prevent_dce.inner = res;
 }
 
 void compute_main() {

test/tint/builtins/gen/literal/abs/4ad288.wgsl.expected.msl

@@ -1,33 +1,34 @@
 #include <metal_stdlib>
 
 using namespace metal;
-void abs_4ad288() {
+void abs_4ad288(device int* const tint_symbol_1) {
   int res = 1;
+  *(tint_symbol_1) = res;
 }
 
 struct tint_symbol {
   float4 value [[position]];
 };
 
-float4 vertex_main_inner() {
+float4 vertex_main_inner(device int* const tint_symbol_2) {
-  abs_4ad288();
+  abs_4ad288(tint_symbol_2);
   return float4(0.0f);
 }
 
-vertex tint_symbol vertex_main() {
+vertex tint_symbol vertex_main(device int* tint_symbol_3 [[buffer(0)]]) {
-  float4 const inner_result = vertex_main_inner();
+  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() {
+fragment void fragment_main(device int* tint_symbol_4 [[buffer(0)]]) {
-  abs_4ad288();
+  abs_4ad288(tint_symbol_4);
   return;
 }
 
-kernel void compute_main() {
+kernel void compute_main(device int* tint_symbol_5 [[buffer(0)]]) {
-  abs_4ad288();
+  abs_4ad288(tint_symbol_5);
   return;
 }
 

test/tint/builtins/gen/literal/abs/4ad288.wgsl.expected.spvasm

@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 32
+; Bound: 40
 ; 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 %abs_4ad288 "abs_4ad288"
                OpName %res "res"
                OpName %vertex_main_inner "vertex_main_inner"
@@ -20,6 +23,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
@@ -28,39 +35,48 @@
 %_ptr_Output_float = OpTypePointer Output %float
           %8 = OpConstantNull %float
 %vertex_point_size = OpVariable %_ptr_Output_float Output %8
-       %void = OpTypeVoid
-          %9 = OpTypeFunction %void
         %int = OpTypeInt 32 1
+%prevent_dce_block = OpTypeStruct %int
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+       %void = OpTypeVoid
+         %13 = OpTypeFunction %void
       %int_1 = OpConstant %int 1
 %_ptr_Function_int = OpTypePointer Function %int
-         %17 = OpConstantNull %int
+         %20 = OpConstantNull %int
-         %18 = OpTypeFunction %v4float
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+         %26 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
- %abs_4ad288 = OpFunction %void None %9
+ %abs_4ad288 = OpFunction %void None %13
-         %12 = OpLabel
+         %16 = OpLabel
-        %res = OpVariable %_ptr_Function_int Function %17
+        %res = OpVariable %_ptr_Function_int Function %20
                OpStore %res %int_1
+         %24 = OpAccessChain %_ptr_StorageBuffer_int %prevent_dce %uint_0
+         %25 = OpLoad %int %res
+               OpStore %24 %25
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %18
+%vertex_main_inner = OpFunction %v4float None %26
-         %20 = OpLabel
+         %28 = OpLabel
-         %21 = OpFunctionCall %void %abs_4ad288
+         %29 = OpFunctionCall %void %abs_4ad288
                OpReturnValue %5
                OpFunctionEnd
-%vertex_main = OpFunction %void None %9
+%vertex_main = OpFunction %void None %13
-         %23 = OpLabel
+         %31 = OpLabel
-         %24 = OpFunctionCall %v4float %vertex_main_inner
+         %32 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %24
+               OpStore %value %32
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
-%fragment_main = OpFunction %void None %9
+%fragment_main = OpFunction %void None %13
-         %27 = OpLabel
+         %35 = OpLabel
-         %28 = OpFunctionCall %void %abs_4ad288
+         %36 = OpFunctionCall %void %abs_4ad288
                OpReturn
                OpFunctionEnd
-%compute_main = OpFunction %void None %9
+%compute_main = OpFunction %void None %13
-         %30 = OpLabel
+         %38 = OpLabel
-         %31 = OpFunctionCall %void %abs_4ad288
+         %39 = OpFunctionCall %void %abs_4ad288
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/abs/4ad288.wgsl.expected.wgsl

@@ -1,7 +1,10 @@
 fn abs_4ad288() {
   var res : i32 = abs(1i);
+  prevent_dce = res;
 }
 
+@group(2) @binding(0) var<storage, read_write> prevent_dce : i32;
+
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   abs_4ad288();

test/tint/builtins/gen/literal/abs/538d29.wgsl

@@ -26,7 +26,9 @@ enable f16;
 // fn abs(vec<4, f16>) -> vec<4, f16>
 fn abs_538d29() {
   var res: vec4<f16> = abs(vec4<f16>(1.h));
+  prevent_dce = res;
 }
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<f16>;
 
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {

test/tint/builtins/gen/literal/abs/538d29.wgsl.expected.dxc.hlsl

@@ -1,5 +1,8 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
 void abs_538d29() {
   vector<float16_t, 4> res = (float16_t(1.0h)).xxxx;
+  prevent_dce.Store<vector<float16_t, 4> >(0u, res);
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/538d29.wgsl.expected.fxc.hlsl

@@ -1,7 +1,10 @@
 SKIP: FAILED
 
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
 void abs_538d29() {
   vector<float16_t, 4> res = (float16_t(1.0h)).xxxx;
+  prevent_dce.Store<vector<float16_t, 4> >(0u, res);
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/538d29.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 {
+  f16vec4 inner;
+} prevent_dce;
+
 void abs_538d29() {
   f16vec4 res = f16vec4(1.0hf);
+  prevent_dce.inner = res;
 }
 
 vec4 vertex_main() {
@@ -22,8 +27,13 @@ void main() {
 #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 abs_538d29() {
   f16vec4 res = f16vec4(1.0hf);
+  prevent_dce.inner = res;
 }
 
 void fragment_main() {
@@ -37,8 +47,13 @@ void main() {
 #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 abs_538d29() {
   f16vec4 res = f16vec4(1.0hf);
+  prevent_dce.inner = res;
 }
 
 void compute_main() {

test/tint/builtins/gen/literal/abs/538d29.wgsl.expected.msl

@@ -1,33 +1,34 @@
 #include <metal_stdlib>
 
 using namespace metal;
-void abs_538d29() {
+void abs_538d29(device half4* const tint_symbol_1) {
   half4 res = half4(1.0h);
+  *(tint_symbol_1) = res;
 }
 
 struct tint_symbol {
   float4 value [[position]];
 };
 
-float4 vertex_main_inner() {
+float4 vertex_main_inner(device half4* const tint_symbol_2) {
-  abs_538d29();
+  abs_538d29(tint_symbol_2);
   return float4(0.0f);
 }
 
-vertex tint_symbol vertex_main() {
+vertex tint_symbol vertex_main(device half4* tint_symbol_3 [[buffer(0)]]) {
-  float4 const inner_result = vertex_main_inner();
+  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() {
+fragment void fragment_main(device half4* tint_symbol_4 [[buffer(0)]]) {
-  abs_538d29();
+  abs_538d29(tint_symbol_4);
   return;
 }
 
-kernel void compute_main() {
+kernel void compute_main(device half4* tint_symbol_5 [[buffer(0)]]) {
-  abs_538d29();
+  abs_538d29(tint_symbol_5);
   return;
 }
 

test/tint/builtins/gen/literal/abs/538d29.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
                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 %abs_538d29 "abs_538d29"
                OpName %res "res"
                OpName %vertex_main_inner "vertex_main_inner"
@@ -24,6 +27,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
@@ -32,41 +39,50 @@
 %_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_0x1p_0 = OpConstant %half 0x1p+0
-         %16 = OpConstantComposite %v4half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
+         %19 = OpConstantComposite %v4half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
 %_ptr_Function_v4half = OpTypePointer Function %v4half
-         %19 = OpConstantNull %v4half
+         %22 = OpConstantNull %v4half
-         %20 = OpTypeFunction %v4float
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v4half = OpTypePointer StorageBuffer %v4half
+         %28 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
- %abs_538d29 = OpFunction %void None %9
+ %abs_538d29 = OpFunction %void None %14
-         %12 = OpLabel
+         %17 = OpLabel
-        %res = OpVariable %_ptr_Function_v4half Function %19
+        %res = OpVariable %_ptr_Function_v4half Function %22
-               OpStore %res %16
+               OpStore %res %19
+         %26 = OpAccessChain %_ptr_StorageBuffer_v4half %prevent_dce %uint_0
+         %27 = OpLoad %v4half %res
+               OpStore %26 %27
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %20
+%vertex_main_inner = OpFunction %v4float None %28
-         %22 = OpLabel
+         %30 = OpLabel
-         %23 = OpFunctionCall %void %abs_538d29
+         %31 = OpFunctionCall %void %abs_538d29
                OpReturnValue %5
                OpFunctionEnd
-%vertex_main = OpFunction %void None %9
+%vertex_main = OpFunction %void None %14
-         %25 = OpLabel
+         %33 = OpLabel
-         %26 = OpFunctionCall %v4float %vertex_main_inner
+         %34 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %26
+               OpStore %value %34
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
-%fragment_main = OpFunction %void None %9
+%fragment_main = OpFunction %void None %14
-         %29 = OpLabel
+         %37 = OpLabel
-         %30 = OpFunctionCall %void %abs_538d29
+         %38 = OpFunctionCall %void %abs_538d29
                OpReturn
                OpFunctionEnd
-%compute_main = OpFunction %void None %9
+%compute_main = OpFunction %void None %14
-         %32 = OpLabel
+         %40 = OpLabel
-         %33 = OpFunctionCall %void %abs_538d29
+         %41 = OpFunctionCall %void %abs_538d29
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/abs/538d29.wgsl.expected.wgsl

@@ -2,8 +2,11 @@ enable f16;
 
 fn abs_538d29() {
   var res : vec4<f16> = abs(vec4<f16>(1.0h));
+  prevent_dce = res;
 }
 
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<f16>;
+
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   abs_538d29();

test/tint/builtins/gen/literal/abs/577d6e.wgsl

@@ -25,7 +25,6 @@
 fn abs_577d6e() {
   var res = abs(vec2(1));
 }
-
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   abs_577d6e();

test/tint/builtins/gen/literal/abs/5a8af1.wgsl

@@ -25,7 +25,6 @@
 fn abs_5a8af1() {
   var res = abs(1);
 }
-
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   abs_5a8af1();

test/tint/builtins/gen/literal/abs/5ad50a.wgsl

@@ -24,7 +24,9 @@
 // fn abs(vec<3, i32>) -> vec<3, i32>
 fn abs_5ad50a() {
   var res: vec3<i32> = abs(vec3<i32>(1i));
+  prevent_dce = res;
 }
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec3<i32>;
 
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {

test/tint/builtins/gen/literal/abs/5ad50a.wgsl.expected.dxc.hlsl

@@ -1,5 +1,8 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
 void abs_5ad50a() {
   int3 res = (1).xxx;
+  prevent_dce.Store3(0u, asuint(res));
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/5ad50a.wgsl.expected.fxc.hlsl

@@ -1,5 +1,8 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
 void abs_5ad50a() {
   int3 res = (1).xxx;
+  prevent_dce.Store3(0u, asuint(res));
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/5ad50a.wgsl.expected.glsl

@@ -1,7 +1,12 @@
 #version 310 es
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  ivec3 inner;
+} prevent_dce;
+
 void abs_5ad50a() {
   ivec3 res = ivec3(1);
+  prevent_dce.inner = res;
 }
 
 vec4 vertex_main() {
@@ -20,8 +25,13 @@ void main() {
 #version 310 es
 precision mediump float;
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  ivec3 inner;
+} prevent_dce;
+
 void abs_5ad50a() {
   ivec3 res = ivec3(1);
+  prevent_dce.inner = res;
 }
 
 void fragment_main() {
@@ -34,8 +44,13 @@ void main() {
 }
 #version 310 es
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  ivec3 inner;
+} prevent_dce;
+
 void abs_5ad50a() {
   ivec3 res = ivec3(1);
+  prevent_dce.inner = res;
 }
 
 void compute_main() {

test/tint/builtins/gen/literal/abs/5ad50a.wgsl.expected.msl

@@ -1,33 +1,34 @@
 #include <metal_stdlib>
 
 using namespace metal;
-void abs_5ad50a() {
+void abs_5ad50a(device packed_int3* const tint_symbol_1) {
   int3 res = int3(1);
+  *(tint_symbol_1) = packed_int3(res);
 }
 
 struct tint_symbol {
   float4 value [[position]];
 };
 
-float4 vertex_main_inner() {
+float4 vertex_main_inner(device packed_int3* const tint_symbol_2) {
-  abs_5ad50a();
+  abs_5ad50a(tint_symbol_2);
   return float4(0.0f);
 }
 
-vertex tint_symbol vertex_main() {
+vertex tint_symbol vertex_main(device packed_int3* tint_symbol_3 [[buffer(0)]]) {
-  float4 const inner_result = vertex_main_inner();
+  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() {
+fragment void fragment_main(device packed_int3* tint_symbol_4 [[buffer(0)]]) {
-  abs_5ad50a();
+  abs_5ad50a(tint_symbol_4);
   return;
 }
 
-kernel void compute_main() {
+kernel void compute_main(device packed_int3* tint_symbol_5 [[buffer(0)]]) {
-  abs_5ad50a();
+  abs_5ad50a(tint_symbol_5);
   return;
 }
 

test/tint/builtins/gen/literal/abs/5ad50a.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 %abs_5ad50a "abs_5ad50a"
                OpName %res "res"
                OpName %vertex_main_inner "vertex_main_inner"
@@ -20,6 +23,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
@@ -28,41 +35,50 @@
 %_ptr_Output_float = OpTypePointer Output %float
           %8 = OpConstantNull %float
 %vertex_point_size = OpVariable %_ptr_Output_float Output %8
-       %void = OpTypeVoid
-          %9 = OpTypeFunction %void
         %int = OpTypeInt 32 1
       %v3int = OpTypeVector %int 3
+%prevent_dce_block = OpTypeStruct %v3int
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+       %void = OpTypeVoid
+         %14 = OpTypeFunction %void
       %int_1 = OpConstant %int 1
-         %16 = OpConstantComposite %v3int %int_1 %int_1 %int_1
+         %19 = OpConstantComposite %v3int %int_1 %int_1 %int_1
 %_ptr_Function_v3int = OpTypePointer Function %v3int
-         %19 = OpConstantNull %v3int
+         %22 = OpConstantNull %v3int
-         %20 = OpTypeFunction %v4float
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v3int = OpTypePointer StorageBuffer %v3int
+         %28 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
- %abs_5ad50a = OpFunction %void None %9
+ %abs_5ad50a = OpFunction %void None %14
-         %12 = OpLabel
+         %17 = OpLabel
-        %res = OpVariable %_ptr_Function_v3int Function %19
+        %res = OpVariable %_ptr_Function_v3int Function %22
-               OpStore %res %16
+               OpStore %res %19
+         %26 = OpAccessChain %_ptr_StorageBuffer_v3int %prevent_dce %uint_0
+         %27 = OpLoad %v3int %res
+               OpStore %26 %27
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %20
+%vertex_main_inner = OpFunction %v4float None %28
-         %22 = OpLabel
+         %30 = OpLabel
-         %23 = OpFunctionCall %void %abs_5ad50a
+         %31 = OpFunctionCall %void %abs_5ad50a
                OpReturnValue %5
                OpFunctionEnd
-%vertex_main = OpFunction %void None %9
+%vertex_main = OpFunction %void None %14
-         %25 = OpLabel
+         %33 = OpLabel
-         %26 = OpFunctionCall %v4float %vertex_main_inner
+         %34 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %26
+               OpStore %value %34
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
-%fragment_main = OpFunction %void None %9
+%fragment_main = OpFunction %void None %14
-         %29 = OpLabel
+         %37 = OpLabel
-         %30 = OpFunctionCall %void %abs_5ad50a
+         %38 = OpFunctionCall %void %abs_5ad50a
                OpReturn
                OpFunctionEnd
-%compute_main = OpFunction %void None %9
+%compute_main = OpFunction %void None %14
-         %32 = OpLabel
+         %40 = OpLabel
-         %33 = OpFunctionCall %void %abs_5ad50a
+         %41 = OpFunctionCall %void %abs_5ad50a
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/abs/5ad50a.wgsl.expected.wgsl

@@ -1,7 +1,10 @@
 fn abs_5ad50a() {
   var res : vec3<i32> = abs(vec3<i32>(1i));
+  prevent_dce = res;
 }
 
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec3<i32>;
+
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   abs_5ad50a();

test/tint/builtins/gen/literal/abs/5ae4fe.wgsl

@@ -26,7 +26,9 @@ enable f16;
 // fn abs(vec<2, f16>) -> vec<2, f16>
 fn abs_5ae4fe() {
   var res: vec2<f16> = abs(vec2<f16>(1.h));
+  prevent_dce = res;
 }
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec2<f16>;
 
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {

test/tint/builtins/gen/literal/abs/5ae4fe.wgsl.expected.dxc.hlsl

@@ -1,5 +1,8 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
 void abs_5ae4fe() {
   vector<float16_t, 2> res = (float16_t(1.0h)).xx;
+  prevent_dce.Store<vector<float16_t, 2> >(0u, res);
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/5ae4fe.wgsl.expected.fxc.hlsl

@@ -1,7 +1,10 @@
 SKIP: FAILED
 
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
 void abs_5ae4fe() {
   vector<float16_t, 2> res = (float16_t(1.0h)).xx;
+  prevent_dce.Store<vector<float16_t, 2> >(0u, res);
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/5ae4fe.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 {
+  f16vec2 inner;
+} prevent_dce;
+
 void abs_5ae4fe() {
   f16vec2 res = f16vec2(1.0hf);
+  prevent_dce.inner = res;
 }
 
 vec4 vertex_main() {
@@ -22,8 +27,13 @@ void main() {
 #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 abs_5ae4fe() {
   f16vec2 res = f16vec2(1.0hf);
+  prevent_dce.inner = res;
 }
 
 void fragment_main() {
@@ -37,8 +47,13 @@ void main() {
 #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 abs_5ae4fe() {
   f16vec2 res = f16vec2(1.0hf);
+  prevent_dce.inner = res;
 }
 
 void compute_main() {

test/tint/builtins/gen/literal/abs/5ae4fe.wgsl.expected.msl

@@ -1,33 +1,34 @@
 #include <metal_stdlib>
 
 using namespace metal;
-void abs_5ae4fe() {
+void abs_5ae4fe(device half2* const tint_symbol_1) {
   half2 res = half2(1.0h);
+  *(tint_symbol_1) = res;
 }
 
 struct tint_symbol {
   float4 value [[position]];
 };
 
-float4 vertex_main_inner() {
+float4 vertex_main_inner(device half2* const tint_symbol_2) {
-  abs_5ae4fe();
+  abs_5ae4fe(tint_symbol_2);
   return float4(0.0f);
 }
 
-vertex tint_symbol vertex_main() {
+vertex tint_symbol vertex_main(device half2* tint_symbol_3 [[buffer(0)]]) {
-  float4 const inner_result = vertex_main_inner();
+  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() {
+fragment void fragment_main(device half2* tint_symbol_4 [[buffer(0)]]) {
-  abs_5ae4fe();
+  abs_5ae4fe(tint_symbol_4);
   return;
 }
 
-kernel void compute_main() {
+kernel void compute_main(device half2* tint_symbol_5 [[buffer(0)]]) {
-  abs_5ae4fe();
+  abs_5ae4fe(tint_symbol_5);
   return;
 }
 

test/tint/builtins/gen/literal/abs/5ae4fe.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
                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 %abs_5ae4fe "abs_5ae4fe"
                OpName %res "res"
                OpName %vertex_main_inner "vertex_main_inner"
@@ -24,6 +27,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
@@ -32,41 +39,50 @@
 %_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_0x1p_0 = OpConstant %half 0x1p+0
-         %16 = OpConstantComposite %v2half %half_0x1p_0 %half_0x1p_0
+         %19 = OpConstantComposite %v2half %half_0x1p_0 %half_0x1p_0
 %_ptr_Function_v2half = OpTypePointer Function %v2half
-         %19 = OpConstantNull %v2half
+         %22 = OpConstantNull %v2half
-         %20 = OpTypeFunction %v4float
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v2half = OpTypePointer StorageBuffer %v2half
+         %28 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
- %abs_5ae4fe = OpFunction %void None %9
+ %abs_5ae4fe = OpFunction %void None %14
-         %12 = OpLabel
+         %17 = OpLabel
-        %res = OpVariable %_ptr_Function_v2half Function %19
+        %res = OpVariable %_ptr_Function_v2half Function %22
-               OpStore %res %16
+               OpStore %res %19
+         %26 = OpAccessChain %_ptr_StorageBuffer_v2half %prevent_dce %uint_0
+         %27 = OpLoad %v2half %res
+               OpStore %26 %27
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %20
+%vertex_main_inner = OpFunction %v4float None %28
-         %22 = OpLabel
+         %30 = OpLabel
-         %23 = OpFunctionCall %void %abs_5ae4fe
+         %31 = OpFunctionCall %void %abs_5ae4fe
                OpReturnValue %5
                OpFunctionEnd
-%vertex_main = OpFunction %void None %9
+%vertex_main = OpFunction %void None %14
-         %25 = OpLabel
+         %33 = OpLabel
-         %26 = OpFunctionCall %v4float %vertex_main_inner
+         %34 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %26
+               OpStore %value %34
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
-%fragment_main = OpFunction %void None %9
+%fragment_main = OpFunction %void None %14
-         %29 = OpLabel
+         %37 = OpLabel
-         %30 = OpFunctionCall %void %abs_5ae4fe
+         %38 = OpFunctionCall %void %abs_5ae4fe
                OpReturn
                OpFunctionEnd
-%compute_main = OpFunction %void None %9
+%compute_main = OpFunction %void None %14
-         %32 = OpLabel
+         %40 = OpLabel
-         %33 = OpFunctionCall %void %abs_5ae4fe
+         %41 = OpFunctionCall %void %abs_5ae4fe
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/abs/5ae4fe.wgsl.expected.wgsl

@@ -2,8 +2,11 @@ enable f16;
 
 fn abs_5ae4fe() {
   var res : vec2<f16> = abs(vec2<f16>(1.0h));
+  prevent_dce = res;
 }
 
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec2<f16>;
+
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   abs_5ae4fe();

test/tint/builtins/gen/literal/abs/7326de.wgsl

@@ -24,7 +24,9 @@
 // fn abs(vec<3, u32>) -> vec<3, u32>
 fn abs_7326de() {
   var res: vec3<u32> = abs(vec3<u32>(1u));
+  prevent_dce = res;
 }
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec3<u32>;
 
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {

test/tint/builtins/gen/literal/abs/7326de.wgsl.expected.dxc.hlsl

@@ -1,5 +1,8 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
 void abs_7326de() {
   uint3 res = (1u).xxx;
+  prevent_dce.Store3(0u, asuint(res));
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/7326de.wgsl.expected.fxc.hlsl

@@ -1,5 +1,8 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
 void abs_7326de() {
   uint3 res = (1u).xxx;
+  prevent_dce.Store3(0u, asuint(res));
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/7326de.wgsl.expected.glsl

@@ -1,7 +1,12 @@
 #version 310 es
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  uvec3 inner;
+} prevent_dce;
+
 void abs_7326de() {
   uvec3 res = uvec3(1u);
+  prevent_dce.inner = res;
 }
 
 vec4 vertex_main() {
@@ -20,8 +25,13 @@ void main() {
 #version 310 es
 precision mediump float;
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  uvec3 inner;
+} prevent_dce;
+
 void abs_7326de() {
   uvec3 res = uvec3(1u);
+  prevent_dce.inner = res;
 }
 
 void fragment_main() {
@@ -34,8 +44,13 @@ void main() {
 }
 #version 310 es
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  uvec3 inner;
+} prevent_dce;
+
 void abs_7326de() {
   uvec3 res = uvec3(1u);
+  prevent_dce.inner = res;
 }
 
 void compute_main() {

test/tint/builtins/gen/literal/abs/7326de.wgsl.expected.msl

@@ -1,33 +1,34 @@
 #include <metal_stdlib>
 
 using namespace metal;
-void abs_7326de() {
+void abs_7326de(device packed_uint3* const tint_symbol_1) {
   uint3 res = uint3(1u);
+  *(tint_symbol_1) = packed_uint3(res);
 }
 
 struct tint_symbol {
   float4 value [[position]];
 };
 
-float4 vertex_main_inner() {
+float4 vertex_main_inner(device packed_uint3* const tint_symbol_2) {
-  abs_7326de();
+  abs_7326de(tint_symbol_2);
   return float4(0.0f);
 }
 
-vertex tint_symbol vertex_main() {
+vertex tint_symbol vertex_main(device packed_uint3* tint_symbol_3 [[buffer(0)]]) {
-  float4 const inner_result = vertex_main_inner();
+  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() {
+fragment void fragment_main(device packed_uint3* tint_symbol_4 [[buffer(0)]]) {
-  abs_7326de();
+  abs_7326de(tint_symbol_4);
   return;
 }
 
-kernel void compute_main() {
+kernel void compute_main(device packed_uint3* tint_symbol_5 [[buffer(0)]]) {
-  abs_7326de();
+  abs_7326de(tint_symbol_5);
   return;
 }
 

test/tint/builtins/gen/literal/abs/7326de.wgsl.expected.spvasm

@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 34
+; 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 %abs_7326de "abs_7326de"
                OpName %res "res"
                OpName %vertex_main_inner "vertex_main_inner"
@@ -20,6 +23,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
@@ -28,41 +35,49 @@
 %_ptr_Output_float = OpTypePointer Output %float
           %8 = OpConstantNull %float
 %vertex_point_size = OpVariable %_ptr_Output_float Output %8
-       %void = OpTypeVoid
-          %9 = OpTypeFunction %void
        %uint = OpTypeInt 32 0
      %v3uint = OpTypeVector %uint 3
+%prevent_dce_block = OpTypeStruct %v3uint
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+       %void = OpTypeVoid
+         %14 = OpTypeFunction %void
      %uint_1 = OpConstant %uint 1
-         %16 = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1
+         %19 = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1
 %_ptr_Function_v3uint = OpTypePointer Function %v3uint
-         %19 = OpConstantNull %v3uint
+         %22 = OpConstantNull %v3uint
-         %20 = OpTypeFunction %v4float
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v3uint = OpTypePointer StorageBuffer %v3uint
+         %27 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
- %abs_7326de = OpFunction %void None %9
+ %abs_7326de = OpFunction %void None %14
-         %12 = OpLabel
+         %17 = OpLabel
-        %res = OpVariable %_ptr_Function_v3uint Function %19
+        %res = OpVariable %_ptr_Function_v3uint Function %22
-               OpStore %res %16
+               OpStore %res %19
+         %25 = OpAccessChain %_ptr_StorageBuffer_v3uint %prevent_dce %uint_0
+         %26 = OpLoad %v3uint %res
+               OpStore %25 %26
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %20
+%vertex_main_inner = OpFunction %v4float None %27
-         %22 = OpLabel
+         %29 = OpLabel
-         %23 = OpFunctionCall %void %abs_7326de
+         %30 = OpFunctionCall %void %abs_7326de
                OpReturnValue %5
                OpFunctionEnd
-%vertex_main = OpFunction %void None %9
+%vertex_main = OpFunction %void None %14
-         %25 = OpLabel
+         %32 = OpLabel
-         %26 = OpFunctionCall %v4float %vertex_main_inner
+         %33 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %26
+               OpStore %value %33
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
-%fragment_main = OpFunction %void None %9
+%fragment_main = OpFunction %void None %14
-         %29 = OpLabel
+         %36 = OpLabel
-         %30 = OpFunctionCall %void %abs_7326de
+         %37 = OpFunctionCall %void %abs_7326de
                OpReturn
                OpFunctionEnd
-%compute_main = OpFunction %void None %9
+%compute_main = OpFunction %void None %14
-         %32 = OpLabel
+         %39 = OpLabel
-         %33 = OpFunctionCall %void %abs_7326de
+         %40 = OpFunctionCall %void %abs_7326de
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/abs/7326de.wgsl.expected.wgsl

@@ -1,7 +1,10 @@
 fn abs_7326de() {
   var res : vec3<u32> = abs(vec3<u32>(1u));
+  prevent_dce = res;
 }
 
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec3<u32>;
+
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   abs_7326de();

test/tint/builtins/gen/literal/abs/7f28e6.wgsl

@@ -24,7 +24,9 @@
 // fn abs(vec<2, u32>) -> vec<2, u32>
 fn abs_7f28e6() {
   var res: vec2<u32> = abs(vec2<u32>(1u));
+  prevent_dce = res;
 }
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec2<u32>;
 
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {

test/tint/builtins/gen/literal/abs/7f28e6.wgsl.expected.dxc.hlsl

@@ -1,5 +1,8 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
 void abs_7f28e6() {
   uint2 res = (1u).xx;
+  prevent_dce.Store2(0u, asuint(res));
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/7f28e6.wgsl.expected.fxc.hlsl

@@ -1,5 +1,8 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
 void abs_7f28e6() {
   uint2 res = (1u).xx;
+  prevent_dce.Store2(0u, asuint(res));
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/7f28e6.wgsl.expected.glsl

@@ -1,7 +1,12 @@
 #version 310 es
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  uvec2 inner;
+} prevent_dce;
+
 void abs_7f28e6() {
   uvec2 res = uvec2(1u);
+  prevent_dce.inner = res;
 }
 
 vec4 vertex_main() {
@@ -20,8 +25,13 @@ void main() {
 #version 310 es
 precision mediump float;
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  uvec2 inner;
+} prevent_dce;
+
 void abs_7f28e6() {
   uvec2 res = uvec2(1u);
+  prevent_dce.inner = res;
 }
 
 void fragment_main() {
@@ -34,8 +44,13 @@ void main() {
 }
 #version 310 es
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  uvec2 inner;
+} prevent_dce;
+
 void abs_7f28e6() {
   uvec2 res = uvec2(1u);
+  prevent_dce.inner = res;
 }
 
 void compute_main() {

test/tint/builtins/gen/literal/abs/7f28e6.wgsl.expected.msl

@@ -1,33 +1,34 @@
 #include <metal_stdlib>
 
 using namespace metal;
-void abs_7f28e6() {
+void abs_7f28e6(device uint2* const tint_symbol_1) {
   uint2 res = uint2(1u);
+  *(tint_symbol_1) = res;
 }
 
 struct tint_symbol {
   float4 value [[position]];
 };
 
-float4 vertex_main_inner() {
+float4 vertex_main_inner(device uint2* const tint_symbol_2) {
-  abs_7f28e6();
+  abs_7f28e6(tint_symbol_2);
   return float4(0.0f);
 }
 
-vertex tint_symbol vertex_main() {
+vertex tint_symbol vertex_main(device uint2* tint_symbol_3 [[buffer(0)]]) {
-  float4 const inner_result = vertex_main_inner();
+  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() {
+fragment void fragment_main(device uint2* tint_symbol_4 [[buffer(0)]]) {
-  abs_7f28e6();
+  abs_7f28e6(tint_symbol_4);
   return;
 }
 
-kernel void compute_main() {
+kernel void compute_main(device uint2* tint_symbol_5 [[buffer(0)]]) {
-  abs_7f28e6();
+  abs_7f28e6(tint_symbol_5);
   return;
 }
 

test/tint/builtins/gen/literal/abs/7f28e6.wgsl.expected.spvasm

@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 34
+; 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 %abs_7f28e6 "abs_7f28e6"
                OpName %res "res"
                OpName %vertex_main_inner "vertex_main_inner"
@@ -20,6 +23,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
@@ -28,41 +35,49 @@
 %_ptr_Output_float = OpTypePointer Output %float
           %8 = OpConstantNull %float
 %vertex_point_size = OpVariable %_ptr_Output_float Output %8
-       %void = OpTypeVoid
-          %9 = OpTypeFunction %void
        %uint = OpTypeInt 32 0
      %v2uint = OpTypeVector %uint 2
+%prevent_dce_block = OpTypeStruct %v2uint
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+       %void = OpTypeVoid
+         %14 = OpTypeFunction %void
      %uint_1 = OpConstant %uint 1
-         %16 = OpConstantComposite %v2uint %uint_1 %uint_1
+         %19 = OpConstantComposite %v2uint %uint_1 %uint_1
 %_ptr_Function_v2uint = OpTypePointer Function %v2uint
-         %19 = OpConstantNull %v2uint
+         %22 = OpConstantNull %v2uint
-         %20 = OpTypeFunction %v4float
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v2uint = OpTypePointer StorageBuffer %v2uint
+         %27 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
- %abs_7f28e6 = OpFunction %void None %9
+ %abs_7f28e6 = OpFunction %void None %14
-         %12 = OpLabel
+         %17 = OpLabel
-        %res = OpVariable %_ptr_Function_v2uint Function %19
+        %res = OpVariable %_ptr_Function_v2uint Function %22
-               OpStore %res %16
+               OpStore %res %19
+         %25 = OpAccessChain %_ptr_StorageBuffer_v2uint %prevent_dce %uint_0
+         %26 = OpLoad %v2uint %res
+               OpStore %25 %26
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %20
+%vertex_main_inner = OpFunction %v4float None %27
-         %22 = OpLabel
+         %29 = OpLabel
-         %23 = OpFunctionCall %void %abs_7f28e6
+         %30 = OpFunctionCall %void %abs_7f28e6
                OpReturnValue %5
                OpFunctionEnd
-%vertex_main = OpFunction %void None %9
+%vertex_main = OpFunction %void None %14
-         %25 = OpLabel
+         %32 = OpLabel
-         %26 = OpFunctionCall %v4float %vertex_main_inner
+         %33 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %26
+               OpStore %value %33
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
-%fragment_main = OpFunction %void None %9
+%fragment_main = OpFunction %void None %14
-         %29 = OpLabel
+         %36 = OpLabel
-         %30 = OpFunctionCall %void %abs_7f28e6
+         %37 = OpFunctionCall %void %abs_7f28e6
                OpReturn
                OpFunctionEnd
-%compute_main = OpFunction %void None %9
+%compute_main = OpFunction %void None %14
-         %32 = OpLabel
+         %39 = OpLabel
-         %33 = OpFunctionCall %void %abs_7f28e6
+         %40 = OpFunctionCall %void %abs_7f28e6
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/abs/7f28e6.wgsl.expected.wgsl

@@ -1,7 +1,10 @@
 fn abs_7f28e6() {
   var res : vec2<u32> = abs(vec2<u32>(1u));
+  prevent_dce = res;
 }
 
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec2<u32>;
+
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   abs_7f28e6();

test/tint/builtins/gen/literal/abs/7faa9e.wgsl

@@ -24,7 +24,9 @@
 // fn abs(vec<2, i32>) -> vec<2, i32>
 fn abs_7faa9e() {
   var res: vec2<i32> = abs(vec2<i32>(1i));
+  prevent_dce = res;
 }
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec2<i32>;
 
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {

test/tint/builtins/gen/literal/abs/7faa9e.wgsl.expected.dxc.hlsl

@@ -1,5 +1,8 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
 void abs_7faa9e() {
   int2 res = (1).xx;
+  prevent_dce.Store2(0u, asuint(res));
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/7faa9e.wgsl.expected.fxc.hlsl

@@ -1,5 +1,8 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
 void abs_7faa9e() {
   int2 res = (1).xx;
+  prevent_dce.Store2(0u, asuint(res));
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/7faa9e.wgsl.expected.glsl

@@ -1,7 +1,12 @@
 #version 310 es
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  ivec2 inner;
+} prevent_dce;
+
 void abs_7faa9e() {
   ivec2 res = ivec2(1);
+  prevent_dce.inner = res;
 }
 
 vec4 vertex_main() {
@@ -20,8 +25,13 @@ void main() {
 #version 310 es
 precision mediump float;
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  ivec2 inner;
+} prevent_dce;
+
 void abs_7faa9e() {
   ivec2 res = ivec2(1);
+  prevent_dce.inner = res;
 }
 
 void fragment_main() {
@@ -34,8 +44,13 @@ void main() {
 }
 #version 310 es
 
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+  ivec2 inner;
+} prevent_dce;
+
 void abs_7faa9e() {
   ivec2 res = ivec2(1);
+  prevent_dce.inner = res;
 }
 
 void compute_main() {

test/tint/builtins/gen/literal/abs/7faa9e.wgsl.expected.msl

@@ -1,33 +1,34 @@
 #include <metal_stdlib>
 
 using namespace metal;
-void abs_7faa9e() {
+void abs_7faa9e(device int2* const tint_symbol_1) {
   int2 res = int2(1);
+  *(tint_symbol_1) = res;
 }
 
 struct tint_symbol {
   float4 value [[position]];
 };
 
-float4 vertex_main_inner() {
+float4 vertex_main_inner(device int2* const tint_symbol_2) {
-  abs_7faa9e();
+  abs_7faa9e(tint_symbol_2);
   return float4(0.0f);
 }
 
-vertex tint_symbol vertex_main() {
+vertex tint_symbol vertex_main(device int2* tint_symbol_3 [[buffer(0)]]) {
-  float4 const inner_result = vertex_main_inner();
+  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() {
+fragment void fragment_main(device int2* tint_symbol_4 [[buffer(0)]]) {
-  abs_7faa9e();
+  abs_7faa9e(tint_symbol_4);
   return;
 }
 
-kernel void compute_main() {
+kernel void compute_main(device int2* tint_symbol_5 [[buffer(0)]]) {
-  abs_7faa9e();
+  abs_7faa9e(tint_symbol_5);
   return;
 }
 

test/tint/builtins/gen/literal/abs/7faa9e.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 %abs_7faa9e "abs_7faa9e"
                OpName %res "res"
                OpName %vertex_main_inner "vertex_main_inner"
@@ -20,6 +23,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
@@ -28,41 +35,50 @@
 %_ptr_Output_float = OpTypePointer Output %float
           %8 = OpConstantNull %float
 %vertex_point_size = OpVariable %_ptr_Output_float Output %8
-       %void = OpTypeVoid
-          %9 = OpTypeFunction %void
         %int = OpTypeInt 32 1
       %v2int = OpTypeVector %int 2
+%prevent_dce_block = OpTypeStruct %v2int
+%_ptr_StorageBuffer_prevent_dce_block = OpTypePointer StorageBuffer %prevent_dce_block
+%prevent_dce = OpVariable %_ptr_StorageBuffer_prevent_dce_block StorageBuffer
+       %void = OpTypeVoid
+         %14 = OpTypeFunction %void
       %int_1 = OpConstant %int 1
-         %16 = OpConstantComposite %v2int %int_1 %int_1
+         %19 = OpConstantComposite %v2int %int_1 %int_1
 %_ptr_Function_v2int = OpTypePointer Function %v2int
-         %19 = OpConstantNull %v2int
+         %22 = OpConstantNull %v2int
-         %20 = OpTypeFunction %v4float
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v2int = OpTypePointer StorageBuffer %v2int
+         %28 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
- %abs_7faa9e = OpFunction %void None %9
+ %abs_7faa9e = OpFunction %void None %14
-         %12 = OpLabel
+         %17 = OpLabel
-        %res = OpVariable %_ptr_Function_v2int Function %19
+        %res = OpVariable %_ptr_Function_v2int Function %22
-               OpStore %res %16
+               OpStore %res %19
+         %26 = OpAccessChain %_ptr_StorageBuffer_v2int %prevent_dce %uint_0
+         %27 = OpLoad %v2int %res
+               OpStore %26 %27
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %20
+%vertex_main_inner = OpFunction %v4float None %28
-         %22 = OpLabel
+         %30 = OpLabel
-         %23 = OpFunctionCall %void %abs_7faa9e
+         %31 = OpFunctionCall %void %abs_7faa9e
                OpReturnValue %5
                OpFunctionEnd
-%vertex_main = OpFunction %void None %9
+%vertex_main = OpFunction %void None %14
-         %25 = OpLabel
+         %33 = OpLabel
-         %26 = OpFunctionCall %v4float %vertex_main_inner
+         %34 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %26
+               OpStore %value %34
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
-%fragment_main = OpFunction %void None %9
+%fragment_main = OpFunction %void None %14
-         %29 = OpLabel
+         %37 = OpLabel
-         %30 = OpFunctionCall %void %abs_7faa9e
+         %38 = OpFunctionCall %void %abs_7faa9e
                OpReturn
                OpFunctionEnd
-%compute_main = OpFunction %void None %9
+%compute_main = OpFunction %void None %14
-         %32 = OpLabel
+         %40 = OpLabel
-         %33 = OpFunctionCall %void %abs_7faa9e
+         %41 = OpFunctionCall %void %abs_7faa9e
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/abs/7faa9e.wgsl.expected.wgsl

@@ -1,7 +1,10 @@
 fn abs_7faa9e() {
   var res : vec2<i32> = abs(vec2<i32>(1i));
+  prevent_dce = res;
 }
 
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec2<i32>;
+
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   abs_7faa9e();

test/tint/builtins/gen/literal/abs/82ff9d.wgsl

@@ -25,7 +25,6 @@
 fn abs_82ff9d() {
   var res = abs(vec2(1.));
 }
-
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   abs_82ff9d();

test/tint/builtins/gen/literal/abs/8ca9b1.wgsl

@@ -25,7 +25,6 @@
 fn abs_8ca9b1() {
   var res = abs(vec4(1));
 }
-
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   abs_8ca9b1();

test/tint/builtins/gen/literal/abs/9c80a6.wgsl

@@ -24,7 +24,9 @@
 // fn abs(vec<4, i32>) -> vec<4, i32>
 fn abs_9c80a6() {
   var res: vec4<i32> = abs(vec4<i32>(1i));
+  prevent_dce = res;
 }
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<i32>;
 
 @vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {

test/tint/builtins/gen/literal/abs/9c80a6.wgsl.expected.dxc.hlsl

@@ -1,5 +1,8 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
 void abs_9c80a6() {
   int4 res = (1).xxxx;
+  prevent_dce.Store4(0u, asuint(res));
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/9c80a6.wgsl.expected.fxc.hlsl

@@ -1,5 +1,8 @@
+RWByteAddressBuffer prevent_dce : register(u0, space2);
+
 void abs_9c80a6() {
   int4 res = (1).xxxx;
+  prevent_dce.Store4(0u, asuint(res));
 }
 
 struct tint_symbol {