Update buffer/storage e2e tests.
Update the buffer/storage end-to-end tests to make sure that the result
is stored back into a storage buffer.
Change-Id: I9ee835b3ea05400dd02e29e5bf4a77509ec18034
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/164640
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: dan sinclair <dsinclair@chromium.org>
diff --git a/test/tint/buffer/storage/dynamic_index/read.wgsl b/test/tint/buffer/storage/dynamic_index/read.wgsl
index e0995f7..17e3319 100644
--- a/test/tint/buffer/storage/dynamic_index/read.wgsl
+++ b/test/tint/buffer/storage/dynamic_index/read.wgsl
@@ -28,6 +28,7 @@
};
@binding(0) @group(0) var<storage, read> sb : S;
+@group(0) @binding(1) var<storage, read_write> s: i32;
@compute @workgroup_size(1)
fn main(@builtin(local_invocation_index) idx : u32) {
@@ -53,4 +54,15 @@
let mat4x3_f32 : mat4x3<f32> = sb.arr[idx].mat4x3_f32;
let mat4x4_f32 : mat4x4<f32> = sb.arr[idx].mat4x4_f32;
let arr2_vec3_f32 : array<vec3<f32>, 2> = sb.arr[idx].arr2_vec3_f32;
+
+ s = i32(scalar_f32) + scalar_i32 + i32(scalar_u32) +
+ i32(vec2_f32.x) + vec2_i32.x + i32(vec2_u32.x) +
+ i32(vec3_f32.y) + vec3_i32.y + i32(vec3_u32.y) +
+ i32(vec4_f32.z) + vec4_i32.z + i32(vec4_u32.z) +
+ i32(mat2x2_f32[0].x) + i32(mat2x3_f32[0].x) +
+ i32(mat2x4_f32[0].x) + i32(mat3x2_f32[0].x) +
+ i32(mat3x3_f32[0].x) + i32(mat3x4_f32[0].x) +
+ i32(mat4x2_f32[0].x) + i32(mat4x3_f32[0].x) +
+ i32(mat4x4_f32[0].x) + i32(arr2_vec3_f32[0].x);
+
}
diff --git a/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.dxc.hlsl b/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.dxc.hlsl
index 990bb19..2bf86f8 100644
--- a/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.dxc.hlsl
+++ b/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.dxc.hlsl
@@ -1,4 +1,9 @@
+int tint_ftoi(float v) {
+ return ((v < 2147483520.0f) ? ((v < -2147483648.0f) ? -2147483648 : int(v)) : 2147483647);
+}
+
ByteAddressBuffer sb : register(t0);
+RWByteAddressBuffer s : register(u1);
struct tint_symbol_1 {
uint idx : SV_GroupIndex;
@@ -74,6 +79,7 @@
const float4x3 mat4x3_f32 = sb_load_19(((544u * idx) + 384u));
const float4x4 mat4x4_f32 = sb_load_20(((544u * idx) + 448u));
const float3 arr2_vec3_f32[2] = sb_load_21(((544u * idx) + 512u));
+ s.Store(0u, asuint((((((((((((((((((((((tint_ftoi(scalar_f32) + scalar_i32) + int(scalar_u32)) + tint_ftoi(vec2_f32.x)) + vec2_i32.x) + int(vec2_u32.x)) + tint_ftoi(vec3_f32.y)) + vec3_i32.y) + int(vec3_u32.y)) + tint_ftoi(vec4_f32.z)) + vec4_i32.z) + int(vec4_u32.z)) + tint_ftoi(mat2x2_f32[0].x)) + tint_ftoi(mat2x3_f32[0].x)) + tint_ftoi(mat2x4_f32[0].x)) + tint_ftoi(mat3x2_f32[0].x)) + tint_ftoi(mat3x3_f32[0].x)) + tint_ftoi(mat3x4_f32[0].x)) + tint_ftoi(mat4x2_f32[0].x)) + tint_ftoi(mat4x3_f32[0].x)) + tint_ftoi(mat4x4_f32[0].x)) + tint_ftoi(arr2_vec3_f32[0].x))));
}
[numthreads(1, 1, 1)]
diff --git a/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.fxc.hlsl b/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.fxc.hlsl
index 990bb19..2bf86f8 100644
--- a/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.fxc.hlsl
+++ b/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.fxc.hlsl
@@ -1,4 +1,9 @@
+int tint_ftoi(float v) {
+ return ((v < 2147483520.0f) ? ((v < -2147483648.0f) ? -2147483648 : int(v)) : 2147483647);
+}
+
ByteAddressBuffer sb : register(t0);
+RWByteAddressBuffer s : register(u1);
struct tint_symbol_1 {
uint idx : SV_GroupIndex;
@@ -74,6 +79,7 @@
const float4x3 mat4x3_f32 = sb_load_19(((544u * idx) + 384u));
const float4x4 mat4x4_f32 = sb_load_20(((544u * idx) + 448u));
const float3 arr2_vec3_f32[2] = sb_load_21(((544u * idx) + 512u));
+ s.Store(0u, asuint((((((((((((((((((((((tint_ftoi(scalar_f32) + scalar_i32) + int(scalar_u32)) + tint_ftoi(vec2_f32.x)) + vec2_i32.x) + int(vec2_u32.x)) + tint_ftoi(vec3_f32.y)) + vec3_i32.y) + int(vec3_u32.y)) + tint_ftoi(vec4_f32.z)) + vec4_i32.z) + int(vec4_u32.z)) + tint_ftoi(mat2x2_f32[0].x)) + tint_ftoi(mat2x3_f32[0].x)) + tint_ftoi(mat2x4_f32[0].x)) + tint_ftoi(mat3x2_f32[0].x)) + tint_ftoi(mat3x3_f32[0].x)) + tint_ftoi(mat3x4_f32[0].x)) + tint_ftoi(mat4x2_f32[0].x)) + tint_ftoi(mat4x3_f32[0].x)) + tint_ftoi(mat4x4_f32[0].x)) + tint_ftoi(arr2_vec3_f32[0].x))));
}
[numthreads(1, 1, 1)]
diff --git a/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.glsl b/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.glsl
index b497053..78a14d5 100644
--- a/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.glsl
+++ b/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.glsl
@@ -1,5 +1,9 @@
#version 310 es
+int tint_ftoi(float v) {
+ return ((v < 2147483520.0f) ? ((v < -2147483648.0f) ? (-2147483647 - 1) : int(v)) : 2147483647);
+}
+
struct Inner {
float scalar_f32;
int scalar_i32;
@@ -37,6 +41,10 @@
Inner arr[];
} sb;
+layout(binding = 1, std430) buffer s_block_ssbo {
+ int inner;
+} s;
+
void tint_symbol(uint idx) {
float scalar_f32 = sb.arr[idx].scalar_f32;
int scalar_i32 = sb.arr[idx].scalar_i32;
@@ -60,6 +68,7 @@
mat4x3 mat4x3_f32 = sb.arr[idx].mat4x3_f32;
mat4 mat4x4_f32 = sb.arr[idx].mat4x4_f32;
vec3 arr2_vec3_f32[2] = sb.arr[idx].arr2_vec3_f32;
+ s.inner = (((((((((((((((((((((tint_ftoi(scalar_f32) + scalar_i32) + int(scalar_u32)) + tint_ftoi(vec2_f32.x)) + vec2_i32.x) + int(vec2_u32.x)) + tint_ftoi(vec3_f32.y)) + vec3_i32.y) + int(vec3_u32.y)) + tint_ftoi(vec4_f32.z)) + vec4_i32.z) + int(vec4_u32.z)) + tint_ftoi(mat2x2_f32[0].x)) + tint_ftoi(mat2x3_f32[0].x)) + tint_ftoi(mat2x4_f32[0].x)) + tint_ftoi(mat3x2_f32[0].x)) + tint_ftoi(mat3x3_f32[0].x)) + tint_ftoi(mat3x4_f32[0].x)) + tint_ftoi(mat4x2_f32[0].x)) + tint_ftoi(mat4x3_f32[0].x)) + tint_ftoi(mat4x4_f32[0].x)) + tint_ftoi(arr2_vec3_f32[0].x));
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
diff --git a/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.msl b/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.msl
index 48bc0ff..641edaa 100644
--- a/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.msl
+++ b/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.msl
@@ -86,6 +86,10 @@
return result;
}
+int tint_ftoi(float v) {
+ return select(2147483647, select(int(v), (-2147483647 - 1), (v < -2147483648.0f)), (v < 2147483520.0f));
+}
+
struct Inner {
float scalar_f32;
int scalar_i32;
@@ -115,7 +119,7 @@
tint_array<Inner, 1> arr;
};
-void tint_symbol_inner(uint idx, const device S_tint_packed_vec3* const tint_symbol_1) {
+void tint_symbol_inner(uint idx, const device S_tint_packed_vec3* const tint_symbol_1, device int* const tint_symbol_2) {
float const scalar_f32 = (*(tint_symbol_1)).arr[idx].scalar_f32;
int const scalar_i32 = (*(tint_symbol_1)).arr[idx].scalar_i32;
uint const scalar_u32 = (*(tint_symbol_1)).arr[idx].scalar_u32;
@@ -138,10 +142,11 @@
float4x3 const mat4x3_f32 = tint_unpack_vec3_in_composite_2((*(tint_symbol_1)).arr[idx].mat4x3_f32);
float4x4 const mat4x4_f32 = (*(tint_symbol_1)).arr[idx].mat4x4_f32;
tint_array<float3, 2> const arr2_vec3_f32 = tint_unpack_vec3_in_composite_3((*(tint_symbol_1)).arr[idx].arr2_vec3_f32);
+ *(tint_symbol_2) = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(tint_ftoi(scalar_f32)) + as_type<uint>(scalar_i32)))) + as_type<uint>(int(scalar_u32))))) + as_type<uint>(tint_ftoi(vec2_f32[0]))))) + as_type<uint>(vec2_i32[0])))) + as_type<uint>(int(vec2_u32[0]))))) + as_type<uint>(tint_ftoi(vec3_f32[1]))))) + as_type<uint>(vec3_i32[1])))) + as_type<uint>(int(vec3_u32[1]))))) + as_type<uint>(tint_ftoi(vec4_f32[2]))))) + as_type<uint>(vec4_i32[2])))) + as_type<uint>(int(vec4_u32[2]))))) + as_type<uint>(tint_ftoi(mat2x2_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat2x3_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat2x4_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat3x2_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat3x3_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat3x4_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat4x2_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat4x3_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat4x4_f32[0][0]))))) + as_type<uint>(tint_ftoi(arr2_vec3_f32[0][0]))));
}
-kernel void tint_symbol(const device S_tint_packed_vec3* tint_symbol_2 [[buffer(0)]], uint idx [[thread_index_in_threadgroup]]) {
- tint_symbol_inner(idx, tint_symbol_2);
+kernel void tint_symbol(const device S_tint_packed_vec3* tint_symbol_3 [[buffer(1)]], device int* tint_symbol_4 [[buffer(0)]], uint idx [[thread_index_in_threadgroup]]) {
+ tint_symbol_inner(idx, tint_symbol_3, tint_symbol_4);
return;
}
diff --git a/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.spvasm b/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.spvasm
index ea7441a..d512bec 100644
--- a/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.spvasm
+++ b/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 128
+; Bound: 215
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -34,6 +34,11 @@
OpMemberName %Inner 20 "mat4x4_f32"
OpMemberName %Inner 21 "arr2_vec3_f32"
OpName %sb "sb"
+ OpName %s_block "s_block"
+ OpMemberName %s_block 0 "inner"
+ OpName %s "s"
+ OpName %tint_ftoi "tint_ftoi"
+ OpName %v "v"
OpName %main_inner "main_inner"
OpName %idx "idx"
OpName %main "main"
@@ -85,6 +90,10 @@
OpDecorate %sb NonWritable
OpDecorate %sb Binding 0
OpDecorate %sb DescriptorSet 0
+ OpDecorate %s_block Block
+ OpMemberDecorate %s_block 0 Offset 0
+ OpDecorate %s DescriptorSet 0
+ OpDecorate %s Binding 1
%uint = OpTypeInt 32 0
%_ptr_Input_uint = OpTypePointer Input %uint
%idx_1 = OpVariable %_ptr_Input_uint Input
@@ -115,8 +124,17 @@
%S = OpTypeStruct %_runtimearr_Inner
%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
%sb = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+ %s_block = OpTypeStruct %int
+%_ptr_StorageBuffer_s_block = OpTypePointer StorageBuffer %s_block
+ %s = OpVariable %_ptr_StorageBuffer_s_block StorageBuffer
+ %34 = OpTypeFunction %int %float
+%float_2_14748352e_09 = OpConstant %float 2.14748352e+09
+ %bool = OpTypeBool
+%float_n2_14748365e_09 = OpConstant %float -2.14748365e+09
+%int_n2147483648 = OpConstant %int -2147483648
+%int_2147483647 = OpConstant %int 2147483647
%void = OpTypeVoid
- %31 = OpTypeFunction %void %uint
+ %48 = OpTypeFunction %void %uint
%uint_0 = OpConstant %uint 0
%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
%uint_1 = OpConstant %uint 1
@@ -160,59 +178,140 @@
%_ptr_StorageBuffer_mat4v4float = OpTypePointer StorageBuffer %mat4v4float
%uint_21 = OpConstant %uint 21
%_ptr_StorageBuffer__arr_v3float_uint_2 = OpTypePointer StorageBuffer %_arr_v3float_uint_2
- %123 = OpTypeFunction %void
- %main_inner = OpFunction %void None %31
+ %170 = OpConstantNull %int
+ %210 = OpTypeFunction %void
+ %tint_ftoi = OpFunction %int None %34
+ %v = OpFunctionParameter %float
+ %37 = OpLabel
+ %40 = OpFOrdLessThan %bool %v %float_2_14748352e_09
+ %44 = OpFOrdLessThan %bool %v %float_n2_14748365e_09
+ %46 = OpConvertFToS %int %v
+ %42 = OpSelect %int %44 %int_n2147483648 %46
+ %38 = OpSelect %int %40 %42 %int_2147483647
+ OpReturnValue %38
+ OpFunctionEnd
+ %main_inner = OpFunction %void None %48
%idx = OpFunctionParameter %uint
- %35 = OpLabel
- %38 = OpAccessChain %_ptr_StorageBuffer_float %sb %uint_0 %idx %uint_0
- %39 = OpLoad %float %38
- %42 = OpAccessChain %_ptr_StorageBuffer_int %sb %uint_0 %idx %uint_1
- %43 = OpLoad %int %42
- %45 = OpAccessChain %_ptr_StorageBuffer_uint %sb %uint_0 %idx %uint_2
- %46 = OpLoad %uint %45
- %49 = OpAccessChain %_ptr_StorageBuffer_v2float %sb %uint_0 %idx %uint_3
- %50 = OpLoad %v2float %49
- %53 = OpAccessChain %_ptr_StorageBuffer_v2int %sb %uint_0 %idx %uint_4
- %54 = OpLoad %v2int %53
- %57 = OpAccessChain %_ptr_StorageBuffer_v2uint %sb %uint_0 %idx %uint_5
- %58 = OpLoad %v2uint %57
- %61 = OpAccessChain %_ptr_StorageBuffer_v3float %sb %uint_0 %idx %uint_6
- %62 = OpLoad %v3float %61
- %65 = OpAccessChain %_ptr_StorageBuffer_v3int %sb %uint_0 %idx %uint_7
- %66 = OpLoad %v3int %65
- %69 = OpAccessChain %_ptr_StorageBuffer_v3uint %sb %uint_0 %idx %uint_8
- %70 = OpLoad %v3uint %69
- %73 = OpAccessChain %_ptr_StorageBuffer_v4float %sb %uint_0 %idx %uint_9
- %74 = OpLoad %v4float %73
- %77 = OpAccessChain %_ptr_StorageBuffer_v4int %sb %uint_0 %idx %uint_10
- %78 = OpLoad %v4int %77
- %81 = OpAccessChain %_ptr_StorageBuffer_v4uint %sb %uint_0 %idx %uint_11
- %82 = OpLoad %v4uint %81
- %85 = OpAccessChain %_ptr_StorageBuffer_mat2v2float %sb %uint_0 %idx %uint_12
- %86 = OpLoad %mat2v2float %85
- %89 = OpAccessChain %_ptr_StorageBuffer_mat2v3float %sb %uint_0 %idx %uint_13
- %90 = OpLoad %mat2v3float %89
- %93 = OpAccessChain %_ptr_StorageBuffer_mat2v4float %sb %uint_0 %idx %uint_14
- %94 = OpLoad %mat2v4float %93
- %97 = OpAccessChain %_ptr_StorageBuffer_mat3v2float %sb %uint_0 %idx %uint_15
- %98 = OpLoad %mat3v2float %97
- %101 = OpAccessChain %_ptr_StorageBuffer_mat3v3float %sb %uint_0 %idx %uint_16
- %102 = OpLoad %mat3v3float %101
- %105 = OpAccessChain %_ptr_StorageBuffer_mat3v4float %sb %uint_0 %idx %uint_17
- %106 = OpLoad %mat3v4float %105
- %109 = OpAccessChain %_ptr_StorageBuffer_mat4v2float %sb %uint_0 %idx %uint_18
- %110 = OpLoad %mat4v2float %109
- %113 = OpAccessChain %_ptr_StorageBuffer_mat4v3float %sb %uint_0 %idx %uint_19
- %114 = OpLoad %mat4v3float %113
- %117 = OpAccessChain %_ptr_StorageBuffer_mat4v4float %sb %uint_0 %idx %uint_20
- %118 = OpLoad %mat4v4float %117
- %121 = OpAccessChain %_ptr_StorageBuffer__arr_v3float_uint_2 %sb %uint_0 %idx %uint_21
- %122 = OpLoad %_arr_v3float_uint_2 %121
+ %52 = OpLabel
+ %55 = OpAccessChain %_ptr_StorageBuffer_float %sb %uint_0 %idx %uint_0
+ %56 = OpLoad %float %55
+ %59 = OpAccessChain %_ptr_StorageBuffer_int %sb %uint_0 %idx %uint_1
+ %60 = OpLoad %int %59
+ %62 = OpAccessChain %_ptr_StorageBuffer_uint %sb %uint_0 %idx %uint_2
+ %63 = OpLoad %uint %62
+ %66 = OpAccessChain %_ptr_StorageBuffer_v2float %sb %uint_0 %idx %uint_3
+ %67 = OpLoad %v2float %66
+ %70 = OpAccessChain %_ptr_StorageBuffer_v2int %sb %uint_0 %idx %uint_4
+ %71 = OpLoad %v2int %70
+ %74 = OpAccessChain %_ptr_StorageBuffer_v2uint %sb %uint_0 %idx %uint_5
+ %75 = OpLoad %v2uint %74
+ %78 = OpAccessChain %_ptr_StorageBuffer_v3float %sb %uint_0 %idx %uint_6
+ %79 = OpLoad %v3float %78
+ %82 = OpAccessChain %_ptr_StorageBuffer_v3int %sb %uint_0 %idx %uint_7
+ %83 = OpLoad %v3int %82
+ %86 = OpAccessChain %_ptr_StorageBuffer_v3uint %sb %uint_0 %idx %uint_8
+ %87 = OpLoad %v3uint %86
+ %90 = OpAccessChain %_ptr_StorageBuffer_v4float %sb %uint_0 %idx %uint_9
+ %91 = OpLoad %v4float %90
+ %94 = OpAccessChain %_ptr_StorageBuffer_v4int %sb %uint_0 %idx %uint_10
+ %95 = OpLoad %v4int %94
+ %98 = OpAccessChain %_ptr_StorageBuffer_v4uint %sb %uint_0 %idx %uint_11
+ %99 = OpLoad %v4uint %98
+ %102 = OpAccessChain %_ptr_StorageBuffer_mat2v2float %sb %uint_0 %idx %uint_12
+ %103 = OpLoad %mat2v2float %102
+ %106 = OpAccessChain %_ptr_StorageBuffer_mat2v3float %sb %uint_0 %idx %uint_13
+ %107 = OpLoad %mat2v3float %106
+ %110 = OpAccessChain %_ptr_StorageBuffer_mat2v4float %sb %uint_0 %idx %uint_14
+ %111 = OpLoad %mat2v4float %110
+ %114 = OpAccessChain %_ptr_StorageBuffer_mat3v2float %sb %uint_0 %idx %uint_15
+ %115 = OpLoad %mat3v2float %114
+ %118 = OpAccessChain %_ptr_StorageBuffer_mat3v3float %sb %uint_0 %idx %uint_16
+ %119 = OpLoad %mat3v3float %118
+ %122 = OpAccessChain %_ptr_StorageBuffer_mat3v4float %sb %uint_0 %idx %uint_17
+ %123 = OpLoad %mat3v4float %122
+ %126 = OpAccessChain %_ptr_StorageBuffer_mat4v2float %sb %uint_0 %idx %uint_18
+ %127 = OpLoad %mat4v2float %126
+ %130 = OpAccessChain %_ptr_StorageBuffer_mat4v3float %sb %uint_0 %idx %uint_19
+ %131 = OpLoad %mat4v3float %130
+ %134 = OpAccessChain %_ptr_StorageBuffer_mat4v4float %sb %uint_0 %idx %uint_20
+ %135 = OpLoad %mat4v4float %134
+ %138 = OpAccessChain %_ptr_StorageBuffer__arr_v3float_uint_2 %sb %uint_0 %idx %uint_21
+ %139 = OpLoad %_arr_v3float_uint_2 %138
+ %140 = OpAccessChain %_ptr_StorageBuffer_int %s %uint_0
+ %141 = OpFunctionCall %int %tint_ftoi %56
+ %142 = OpIAdd %int %141 %60
+ %143 = OpBitcast %int %63
+ %144 = OpIAdd %int %142 %143
+ %146 = OpCompositeExtract %float %67 0
+ %145 = OpFunctionCall %int %tint_ftoi %146
+ %147 = OpIAdd %int %144 %145
+ %148 = OpCompositeExtract %int %71 0
+ %149 = OpIAdd %int %147 %148
+ %151 = OpCompositeExtract %uint %75 0
+ %150 = OpBitcast %int %151
+ %152 = OpIAdd %int %149 %150
+ %154 = OpCompositeExtract %float %79 1
+ %153 = OpFunctionCall %int %tint_ftoi %154
+ %155 = OpIAdd %int %152 %153
+ %156 = OpCompositeExtract %int %83 1
+ %157 = OpIAdd %int %155 %156
+ %159 = OpCompositeExtract %uint %87 1
+ %158 = OpBitcast %int %159
+ %160 = OpIAdd %int %157 %158
+ %162 = OpCompositeExtract %float %91 2
+ %161 = OpFunctionCall %int %tint_ftoi %162
+ %163 = OpIAdd %int %160 %161
+ %164 = OpCompositeExtract %int %95 2
+ %165 = OpIAdd %int %163 %164
+ %167 = OpCompositeExtract %uint %99 2
+ %166 = OpBitcast %int %167
+ %168 = OpIAdd %int %165 %166
+ %171 = OpCompositeExtract %v2float %103 0
+ %172 = OpCompositeExtract %float %171 0
+ %169 = OpFunctionCall %int %tint_ftoi %172
+ %173 = OpIAdd %int %168 %169
+ %175 = OpCompositeExtract %v3float %107 0
+ %176 = OpCompositeExtract %float %175 0
+ %174 = OpFunctionCall %int %tint_ftoi %176
+ %177 = OpIAdd %int %173 %174
+ %179 = OpCompositeExtract %v4float %111 0
+ %180 = OpCompositeExtract %float %179 0
+ %178 = OpFunctionCall %int %tint_ftoi %180
+ %181 = OpIAdd %int %177 %178
+ %183 = OpCompositeExtract %v2float %115 0
+ %184 = OpCompositeExtract %float %183 0
+ %182 = OpFunctionCall %int %tint_ftoi %184
+ %185 = OpIAdd %int %181 %182
+ %187 = OpCompositeExtract %v3float %119 0
+ %188 = OpCompositeExtract %float %187 0
+ %186 = OpFunctionCall %int %tint_ftoi %188
+ %189 = OpIAdd %int %185 %186
+ %191 = OpCompositeExtract %v4float %123 0
+ %192 = OpCompositeExtract %float %191 0
+ %190 = OpFunctionCall %int %tint_ftoi %192
+ %193 = OpIAdd %int %189 %190
+ %195 = OpCompositeExtract %v2float %127 0
+ %196 = OpCompositeExtract %float %195 0
+ %194 = OpFunctionCall %int %tint_ftoi %196
+ %197 = OpIAdd %int %193 %194
+ %199 = OpCompositeExtract %v3float %131 0
+ %200 = OpCompositeExtract %float %199 0
+ %198 = OpFunctionCall %int %tint_ftoi %200
+ %201 = OpIAdd %int %197 %198
+ %203 = OpCompositeExtract %v4float %135 0
+ %204 = OpCompositeExtract %float %203 0
+ %202 = OpFunctionCall %int %tint_ftoi %204
+ %205 = OpIAdd %int %201 %202
+ %207 = OpCompositeExtract %v3float %139 0
+ %208 = OpCompositeExtract %float %207 0
+ %206 = OpFunctionCall %int %tint_ftoi %208
+ %209 = OpIAdd %int %205 %206
+ OpStore %140 %209
OpReturn
OpFunctionEnd
- %main = OpFunction %void None %123
- %125 = OpLabel
- %127 = OpLoad %uint %idx_1
- %126 = OpFunctionCall %void %main_inner %127
+ %main = OpFunction %void None %210
+ %212 = OpLabel
+ %214 = OpLoad %uint %idx_1
+ %213 = OpFunctionCall %void %main_inner %214
OpReturn
OpFunctionEnd
diff --git a/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.wgsl b/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.wgsl
index 6e7b7b7..6ab9b13 100644
--- a/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.wgsl
+++ b/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.wgsl
@@ -29,6 +29,8 @@
@binding(0) @group(0) var<storage, read> sb : S;
+@group(0) @binding(1) var<storage, read_write> s : i32;
+
@compute @workgroup_size(1)
fn main(@builtin(local_invocation_index) idx : u32) {
let scalar_f32 : f32 = sb.arr[idx].scalar_f32;
@@ -53,4 +55,5 @@
let mat4x3_f32 : mat4x3<f32> = sb.arr[idx].mat4x3_f32;
let mat4x4_f32 : mat4x4<f32> = sb.arr[idx].mat4x4_f32;
let arr2_vec3_f32 : array<vec3<f32>, 2> = sb.arr[idx].arr2_vec3_f32;
+ s = (((((((((((((((((((((i32(scalar_f32) + scalar_i32) + i32(scalar_u32)) + i32(vec2_f32.x)) + vec2_i32.x) + i32(vec2_u32.x)) + i32(vec3_f32.y)) + vec3_i32.y) + i32(vec3_u32.y)) + i32(vec4_f32.z)) + vec4_i32.z) + i32(vec4_u32.z)) + i32(mat2x2_f32[0].x)) + i32(mat2x3_f32[0].x)) + i32(mat2x4_f32[0].x)) + i32(mat3x2_f32[0].x)) + i32(mat3x3_f32[0].x)) + i32(mat3x4_f32[0].x)) + i32(mat4x2_f32[0].x)) + i32(mat4x3_f32[0].x)) + i32(mat4x4_f32[0].x)) + i32(arr2_vec3_f32[0].x));
}
diff --git a/test/tint/buffer/storage/dynamic_index/read_f16.wgsl b/test/tint/buffer/storage/dynamic_index/read_f16.wgsl
index 9351d4b..37cd346 100644
--- a/test/tint/buffer/storage/dynamic_index/read_f16.wgsl
+++ b/test/tint/buffer/storage/dynamic_index/read_f16.wgsl
@@ -44,6 +44,7 @@
};
@binding(0) @group(0) var<storage, read> sb : S;
+@group(0) @binding(1) var<storage, read_write> s: i32;
@compute @workgroup_size(1)
fn main(@builtin(local_invocation_index) idx : u32) {
@@ -83,4 +84,20 @@
let mat4x4_f16 : mat4x4<f16> = sb.arr[idx].mat4x4_f16;
let arr2_vec3_f32 : array<vec3<f32>, 2> = sb.arr[idx].arr2_vec3_f32;
let arr2_mat4x2_f16 : array<mat4x2<f16>, 2> = sb.arr[idx].arr2_mat4x2_f16;
+
+ s = i32(scalar_f32) + scalar_i32 + i32(scalar_u32) + i32(scalar_f16) +
+ i32(vec2_f32.x) + vec2_i32.x + i32(vec2_u32.x) + i32(vec2_f16.x) +
+ i32(vec3_f32.y) + vec3_i32.y + i32(vec3_u32.y) + i32(vec3_f16.y) +
+ i32(vec4_f32.z) + vec4_i32.z + i32(vec4_u32.z) + i32(vec4_f16.z) +
+ i32(mat2x2_f32[0].x) + i32(mat2x3_f32[0].x) +
+ i32(mat2x4_f32[0].x) + i32(mat3x2_f32[0].x) +
+ i32(mat3x3_f32[0].x) + i32(mat3x4_f32[0].x) +
+ i32(mat4x2_f32[0].x) + i32(mat4x3_f32[0].x) +
+ i32(mat4x4_f32[0].x) +
+ i32(mat2x2_f16[0].x) + i32(mat2x3_f16[0].x) +
+ i32(mat2x4_f16[0].x) + i32(mat3x2_f16[0].x) +
+ i32(mat3x3_f16[0].x) + i32(mat3x4_f16[0].x) +
+ i32(mat4x2_f16[0].x) + i32(mat4x3_f16[0].x) +
+ i32(mat4x4_f16[0].x) +
+ i32(arr2_mat4x2_f16[0][0].x) + i32(arr2_vec3_f32[0].x);
}
diff --git a/test/tint/buffer/storage/dynamic_index/read_f16.wgsl.expected.dxc.hlsl b/test/tint/buffer/storage/dynamic_index/read_f16.wgsl.expected.dxc.hlsl
index bab24e4..2bf9779 100644
--- a/test/tint/buffer/storage/dynamic_index/read_f16.wgsl.expected.dxc.hlsl
+++ b/test/tint/buffer/storage/dynamic_index/read_f16.wgsl.expected.dxc.hlsl
@@ -1,4 +1,9 @@
+int tint_ftoi(float v) {
+ return ((v < 2147483520.0f) ? ((v < -2147483648.0f) ? -2147483648 : int(v)) : 2147483647);
+}
+
ByteAddressBuffer sb : register(t0);
+RWByteAddressBuffer s : register(u1);
struct tint_symbol_1 {
uint idx : SV_GroupIndex;
@@ -135,6 +140,7 @@
const matrix<float16_t, 4, 4> mat4x4_f16 = sb_load_33(((800u * idx) + 696u));
const float3 arr2_vec3_f32[2] = sb_load_34(((800u * idx) + 736u));
matrix<float16_t, 4, 2> arr2_mat4x2_f16[2] = sb_load_35(((800u * idx) + 768u));
+ s.Store(0u, asuint((((((((((((((((((((((((((((((((((((tint_ftoi(scalar_f32) + scalar_i32) + int(scalar_u32)) + int(scalar_f16)) + tint_ftoi(vec2_f32.x)) + vec2_i32.x) + int(vec2_u32.x)) + int(vec2_f16.x)) + tint_ftoi(vec3_f32.y)) + vec3_i32.y) + int(vec3_u32.y)) + int(vec3_f16.y)) + tint_ftoi(vec4_f32.z)) + vec4_i32.z) + int(vec4_u32.z)) + int(vec4_f16.z)) + tint_ftoi(mat2x2_f32[0].x)) + tint_ftoi(mat2x3_f32[0].x)) + tint_ftoi(mat2x4_f32[0].x)) + tint_ftoi(mat3x2_f32[0].x)) + tint_ftoi(mat3x3_f32[0].x)) + tint_ftoi(mat3x4_f32[0].x)) + tint_ftoi(mat4x2_f32[0].x)) + tint_ftoi(mat4x3_f32[0].x)) + tint_ftoi(mat4x4_f32[0].x)) + int(mat2x2_f16[0].x)) + int(mat2x3_f16[0].x)) + int(mat2x4_f16[0].x)) + int(mat3x2_f16[0].x)) + int(mat3x3_f16[0].x)) + int(mat3x4_f16[0].x)) + int(mat4x2_f16[0].x)) + int(mat4x3_f16[0].x)) + int(mat4x4_f16[0].x)) + int(arr2_mat4x2_f16[0][0].x)) + tint_ftoi(arr2_vec3_f32[0].x))));
}
[numthreads(1, 1, 1)]
diff --git a/test/tint/buffer/storage/dynamic_index/read_f16.wgsl.expected.glsl b/test/tint/buffer/storage/dynamic_index/read_f16.wgsl.expected.glsl
index d150509..2dedec5 100644
--- a/test/tint/buffer/storage/dynamic_index/read_f16.wgsl.expected.glsl
+++ b/test/tint/buffer/storage/dynamic_index/read_f16.wgsl.expected.glsl
@@ -1,6 +1,10 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
+int tint_ftoi(float v) {
+ return ((v < 2147483520.0f) ? ((v < -2147483648.0f) ? (-2147483647 - 1) : int(v)) : 2147483647);
+}
+
struct Inner {
float scalar_f32;
int scalar_i32;
@@ -57,6 +61,10 @@
Inner arr[];
} sb;
+layout(binding = 1, std430) buffer s_block_ssbo {
+ int inner;
+} s;
+
void tint_symbol(uint idx) {
float scalar_f32 = sb.arr[idx].scalar_f32;
int scalar_i32 = sb.arr[idx].scalar_i32;
@@ -94,6 +102,7 @@
f16mat4 mat4x4_f16 = sb.arr[idx].mat4x4_f16;
vec3 arr2_vec3_f32[2] = sb.arr[idx].arr2_vec3_f32;
f16mat4x2 arr2_mat4x2_f16[2] = sb.arr[idx].arr2_mat4x2_f16;
+ s.inner = (((((((((((((((((((((((((((((((((((tint_ftoi(scalar_f32) + scalar_i32) + int(scalar_u32)) + int(scalar_f16)) + tint_ftoi(vec2_f32.x)) + vec2_i32.x) + int(vec2_u32.x)) + int(vec2_f16.x)) + tint_ftoi(vec3_f32.y)) + vec3_i32.y) + int(vec3_u32.y)) + int(vec3_f16.y)) + tint_ftoi(vec4_f32.z)) + vec4_i32.z) + int(vec4_u32.z)) + int(vec4_f16.z)) + tint_ftoi(mat2x2_f32[0].x)) + tint_ftoi(mat2x3_f32[0].x)) + tint_ftoi(mat2x4_f32[0].x)) + tint_ftoi(mat3x2_f32[0].x)) + tint_ftoi(mat3x3_f32[0].x)) + tint_ftoi(mat3x4_f32[0].x)) + tint_ftoi(mat4x2_f32[0].x)) + tint_ftoi(mat4x3_f32[0].x)) + tint_ftoi(mat4x4_f32[0].x)) + int(mat2x2_f16[0].x)) + int(mat2x3_f16[0].x)) + int(mat2x4_f16[0].x)) + int(mat3x2_f16[0].x)) + int(mat3x3_f16[0].x)) + int(mat3x4_f16[0].x)) + int(mat4x2_f16[0].x)) + int(mat4x3_f16[0].x)) + int(mat4x4_f16[0].x)) + int(arr2_mat4x2_f16[0][0].x)) + tint_ftoi(arr2_vec3_f32[0].x));
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
diff --git a/test/tint/buffer/storage/dynamic_index/read_f16.wgsl.expected.msl b/test/tint/buffer/storage/dynamic_index/read_f16.wgsl.expected.msl
index 0acb114..714f1de 100644
--- a/test/tint/buffer/storage/dynamic_index/read_f16.wgsl.expected.msl
+++ b/test/tint/buffer/storage/dynamic_index/read_f16.wgsl.expected.msl
@@ -133,6 +133,10 @@
return result;
}
+int tint_ftoi(float v) {
+ return select(2147483647, select(int(v), (-2147483647 - 1), (v < -2147483648.0f)), (v < 2147483520.0f));
+}
+
struct Inner {
float scalar_f32;
int scalar_i32;
@@ -176,7 +180,7 @@
tint_array<Inner, 1> arr;
};
-void tint_symbol_inner(uint idx, const device S_tint_packed_vec3* const tint_symbol_1) {
+void tint_symbol_inner(uint idx, const device S_tint_packed_vec3* const tint_symbol_1, device int* const tint_symbol_2) {
float const scalar_f32 = (*(tint_symbol_1)).arr[idx].scalar_f32;
int const scalar_i32 = (*(tint_symbol_1)).arr[idx].scalar_i32;
uint const scalar_u32 = (*(tint_symbol_1)).arr[idx].scalar_u32;
@@ -213,10 +217,11 @@
half4x4 const mat4x4_f16 = (*(tint_symbol_1)).arr[idx].mat4x4_f16;
tint_array<float3, 2> const arr2_vec3_f32 = tint_unpack_vec3_in_composite_6((*(tint_symbol_1)).arr[idx].arr2_vec3_f32);
tint_array<half4x2, 2> const arr2_mat4x2_f16 = (*(tint_symbol_1)).arr[idx].arr2_mat4x2_f16;
+ *(tint_symbol_2) = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(tint_ftoi(scalar_f32)) + as_type<uint>(scalar_i32)))) + as_type<uint>(int(scalar_u32))))) + as_type<uint>(int(scalar_f16))))) + as_type<uint>(tint_ftoi(vec2_f32[0]))))) + as_type<uint>(vec2_i32[0])))) + as_type<uint>(int(vec2_u32[0]))))) + as_type<uint>(int(vec2_f16[0]))))) + as_type<uint>(tint_ftoi(vec3_f32[1]))))) + as_type<uint>(vec3_i32[1])))) + as_type<uint>(int(vec3_u32[1]))))) + as_type<uint>(int(vec3_f16[1]))))) + as_type<uint>(tint_ftoi(vec4_f32[2]))))) + as_type<uint>(vec4_i32[2])))) + as_type<uint>(int(vec4_u32[2]))))) + as_type<uint>(int(vec4_f16[2]))))) + as_type<uint>(tint_ftoi(mat2x2_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat2x3_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat2x4_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat3x2_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat3x3_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat3x4_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat4x2_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat4x3_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat4x4_f32[0][0]))))) + as_type<uint>(int(mat2x2_f16[0][0]))))) + as_type<uint>(int(mat2x3_f16[0][0]))))) + as_type<uint>(int(mat2x4_f16[0][0]))))) + as_type<uint>(int(mat3x2_f16[0][0]))))) + as_type<uint>(int(mat3x3_f16[0][0]))))) + as_type<uint>(int(mat3x4_f16[0][0]))))) + as_type<uint>(int(mat4x2_f16[0][0]))))) + as_type<uint>(int(mat4x3_f16[0][0]))))) + as_type<uint>(int(mat4x4_f16[0][0]))))) + as_type<uint>(int(arr2_mat4x2_f16[0][0][0]))))) + as_type<uint>(tint_ftoi(arr2_vec3_f32[0][0]))));
}
-kernel void tint_symbol(const device S_tint_packed_vec3* tint_symbol_2 [[buffer(0)]], uint idx [[thread_index_in_threadgroup]]) {
- tint_symbol_inner(idx, tint_symbol_2);
+kernel void tint_symbol(const device S_tint_packed_vec3* tint_symbol_3 [[buffer(1)]], device int* tint_symbol_4 [[buffer(0)]], uint idx [[thread_index_in_threadgroup]]) {
+ tint_symbol_inner(idx, tint_symbol_3, tint_symbol_4);
return;
}
diff --git a/test/tint/buffer/storage/dynamic_index/read_f16.wgsl.expected.spvasm b/test/tint/buffer/storage/dynamic_index/read_f16.wgsl.expected.spvasm
index d5f4b75..4774b03 100644
--- a/test/tint/buffer/storage/dynamic_index/read_f16.wgsl.expected.spvasm
+++ b/test/tint/buffer/storage/dynamic_index/read_f16.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 198
+; Bound: 337
; Schema: 0
OpCapability Shader
OpCapability Float16
@@ -52,6 +52,11 @@
OpMemberName %Inner 34 "arr2_vec3_f32"
OpMemberName %Inner 35 "arr2_mat4x2_f16"
OpName %sb "sb"
+ OpName %s_block "s_block"
+ OpMemberName %s_block 0 "inner"
+ OpName %s "s"
+ OpName %tint_ftoi "tint_ftoi"
+ OpName %v "v"
OpName %main_inner "main_inner"
OpName %idx "idx"
OpName %main "main"
@@ -138,6 +143,10 @@
OpDecorate %sb NonWritable
OpDecorate %sb Binding 0
OpDecorate %sb DescriptorSet 0
+ OpDecorate %s_block Block
+ OpMemberDecorate %s_block 0 Offset 0
+ OpDecorate %s DescriptorSet 0
+ OpDecorate %s Binding 1
%uint = OpTypeInt 32 0
%_ptr_Input_uint = OpTypePointer Input %uint
%idx_1 = OpVariable %_ptr_Input_uint Input
@@ -182,8 +191,17 @@
%S = OpTypeStruct %_runtimearr_Inner
%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
%sb = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+ %s_block = OpTypeStruct %int
+%_ptr_StorageBuffer_s_block = OpTypePointer StorageBuffer %s_block
+ %s = OpVariable %_ptr_StorageBuffer_s_block StorageBuffer
+ %48 = OpTypeFunction %int %float
+%float_2_14748352e_09 = OpConstant %float 2.14748352e+09
+ %bool = OpTypeBool
+%float_n2_14748365e_09 = OpConstant %float -2.14748365e+09
+%int_n2147483648 = OpConstant %int -2147483648
+%int_2147483647 = OpConstant %int 2147483647
%void = OpTypeVoid
- %45 = OpTypeFunction %void %uint
+ %62 = OpTypeFunction %void %uint
%uint_0 = OpConstant %uint 0
%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
%uint_1 = OpConstant %uint 1
@@ -255,87 +273,220 @@
%_ptr_StorageBuffer__arr_v3float_uint_2 = OpTypePointer StorageBuffer %_arr_v3float_uint_2
%uint_35 = OpConstant %uint 35
%_ptr_StorageBuffer__arr_mat4v2half_uint_2 = OpTypePointer StorageBuffer %_arr_mat4v2half_uint_2
- %193 = OpTypeFunction %void
- %main_inner = OpFunction %void None %45
+ %251 = OpConstantNull %int
+ %332 = OpTypeFunction %void
+ %tint_ftoi = OpFunction %int None %48
+ %v = OpFunctionParameter %float
+ %51 = OpLabel
+ %54 = OpFOrdLessThan %bool %v %float_2_14748352e_09
+ %58 = OpFOrdLessThan %bool %v %float_n2_14748365e_09
+ %60 = OpConvertFToS %int %v
+ %56 = OpSelect %int %58 %int_n2147483648 %60
+ %52 = OpSelect %int %54 %56 %int_2147483647
+ OpReturnValue %52
+ OpFunctionEnd
+ %main_inner = OpFunction %void None %62
%idx = OpFunctionParameter %uint
- %49 = OpLabel
- %52 = OpAccessChain %_ptr_StorageBuffer_float %sb %uint_0 %idx %uint_0
- %53 = OpLoad %float %52
- %56 = OpAccessChain %_ptr_StorageBuffer_int %sb %uint_0 %idx %uint_1
- %57 = OpLoad %int %56
- %59 = OpAccessChain %_ptr_StorageBuffer_uint %sb %uint_0 %idx %uint_2
- %60 = OpLoad %uint %59
- %63 = OpAccessChain %_ptr_StorageBuffer_half %sb %uint_0 %idx %uint_3
- %64 = OpLoad %half %63
- %67 = OpAccessChain %_ptr_StorageBuffer_v2float %sb %uint_0 %idx %uint_4
- %68 = OpLoad %v2float %67
- %71 = OpAccessChain %_ptr_StorageBuffer_v2int %sb %uint_0 %idx %uint_5
- %72 = OpLoad %v2int %71
- %75 = OpAccessChain %_ptr_StorageBuffer_v2uint %sb %uint_0 %idx %uint_6
- %76 = OpLoad %v2uint %75
- %79 = OpAccessChain %_ptr_StorageBuffer_v2half %sb %uint_0 %idx %uint_7
- %80 = OpLoad %v2half %79
- %83 = OpAccessChain %_ptr_StorageBuffer_v3float %sb %uint_0 %idx %uint_8
- %84 = OpLoad %v3float %83
- %87 = OpAccessChain %_ptr_StorageBuffer_v3int %sb %uint_0 %idx %uint_9
- %88 = OpLoad %v3int %87
- %91 = OpAccessChain %_ptr_StorageBuffer_v3uint %sb %uint_0 %idx %uint_10
- %92 = OpLoad %v3uint %91
- %95 = OpAccessChain %_ptr_StorageBuffer_v3half %sb %uint_0 %idx %uint_11
- %96 = OpLoad %v3half %95
- %99 = OpAccessChain %_ptr_StorageBuffer_v4float %sb %uint_0 %idx %uint_12
- %100 = OpLoad %v4float %99
- %103 = OpAccessChain %_ptr_StorageBuffer_v4int %sb %uint_0 %idx %uint_13
- %104 = OpLoad %v4int %103
- %107 = OpAccessChain %_ptr_StorageBuffer_v4uint %sb %uint_0 %idx %uint_14
- %108 = OpLoad %v4uint %107
- %111 = OpAccessChain %_ptr_StorageBuffer_v4half %sb %uint_0 %idx %uint_15
- %112 = OpLoad %v4half %111
- %115 = OpAccessChain %_ptr_StorageBuffer_mat2v2float %sb %uint_0 %idx %uint_16
- %116 = OpLoad %mat2v2float %115
- %119 = OpAccessChain %_ptr_StorageBuffer_mat2v3float %sb %uint_0 %idx %uint_17
- %120 = OpLoad %mat2v3float %119
- %123 = OpAccessChain %_ptr_StorageBuffer_mat2v4float %sb %uint_0 %idx %uint_18
- %124 = OpLoad %mat2v4float %123
- %127 = OpAccessChain %_ptr_StorageBuffer_mat3v2float %sb %uint_0 %idx %uint_19
- %128 = OpLoad %mat3v2float %127
- %131 = OpAccessChain %_ptr_StorageBuffer_mat3v3float %sb %uint_0 %idx %uint_20
- %132 = OpLoad %mat3v3float %131
- %135 = OpAccessChain %_ptr_StorageBuffer_mat3v4float %sb %uint_0 %idx %uint_21
- %136 = OpLoad %mat3v4float %135
- %139 = OpAccessChain %_ptr_StorageBuffer_mat4v2float %sb %uint_0 %idx %uint_22
- %140 = OpLoad %mat4v2float %139
- %143 = OpAccessChain %_ptr_StorageBuffer_mat4v3float %sb %uint_0 %idx %uint_23
- %144 = OpLoad %mat4v3float %143
- %147 = OpAccessChain %_ptr_StorageBuffer_mat4v4float %sb %uint_0 %idx %uint_24
- %148 = OpLoad %mat4v4float %147
- %151 = OpAccessChain %_ptr_StorageBuffer_mat2v2half %sb %uint_0 %idx %uint_25
- %152 = OpLoad %mat2v2half %151
- %155 = OpAccessChain %_ptr_StorageBuffer_mat2v3half %sb %uint_0 %idx %uint_26
- %156 = OpLoad %mat2v3half %155
- %159 = OpAccessChain %_ptr_StorageBuffer_mat2v4half %sb %uint_0 %idx %uint_27
- %160 = OpLoad %mat2v4half %159
- %163 = OpAccessChain %_ptr_StorageBuffer_mat3v2half %sb %uint_0 %idx %uint_28
- %164 = OpLoad %mat3v2half %163
- %167 = OpAccessChain %_ptr_StorageBuffer_mat3v3half %sb %uint_0 %idx %uint_29
- %168 = OpLoad %mat3v3half %167
- %171 = OpAccessChain %_ptr_StorageBuffer_mat3v4half %sb %uint_0 %idx %uint_30
- %172 = OpLoad %mat3v4half %171
- %175 = OpAccessChain %_ptr_StorageBuffer_mat4v2half %sb %uint_0 %idx %uint_31
- %176 = OpLoad %mat4v2half %175
- %179 = OpAccessChain %_ptr_StorageBuffer_mat4v3half %sb %uint_0 %idx %uint_32
- %180 = OpLoad %mat4v3half %179
- %183 = OpAccessChain %_ptr_StorageBuffer_mat4v4half %sb %uint_0 %idx %uint_33
- %184 = OpLoad %mat4v4half %183
- %187 = OpAccessChain %_ptr_StorageBuffer__arr_v3float_uint_2 %sb %uint_0 %idx %uint_34
- %188 = OpLoad %_arr_v3float_uint_2 %187
- %191 = OpAccessChain %_ptr_StorageBuffer__arr_mat4v2half_uint_2 %sb %uint_0 %idx %uint_35
- %192 = OpLoad %_arr_mat4v2half_uint_2 %191
+ %66 = OpLabel
+ %69 = OpAccessChain %_ptr_StorageBuffer_float %sb %uint_0 %idx %uint_0
+ %70 = OpLoad %float %69
+ %73 = OpAccessChain %_ptr_StorageBuffer_int %sb %uint_0 %idx %uint_1
+ %74 = OpLoad %int %73
+ %76 = OpAccessChain %_ptr_StorageBuffer_uint %sb %uint_0 %idx %uint_2
+ %77 = OpLoad %uint %76
+ %80 = OpAccessChain %_ptr_StorageBuffer_half %sb %uint_0 %idx %uint_3
+ %81 = OpLoad %half %80
+ %84 = OpAccessChain %_ptr_StorageBuffer_v2float %sb %uint_0 %idx %uint_4
+ %85 = OpLoad %v2float %84
+ %88 = OpAccessChain %_ptr_StorageBuffer_v2int %sb %uint_0 %idx %uint_5
+ %89 = OpLoad %v2int %88
+ %92 = OpAccessChain %_ptr_StorageBuffer_v2uint %sb %uint_0 %idx %uint_6
+ %93 = OpLoad %v2uint %92
+ %96 = OpAccessChain %_ptr_StorageBuffer_v2half %sb %uint_0 %idx %uint_7
+ %97 = OpLoad %v2half %96
+ %100 = OpAccessChain %_ptr_StorageBuffer_v3float %sb %uint_0 %idx %uint_8
+ %101 = OpLoad %v3float %100
+ %104 = OpAccessChain %_ptr_StorageBuffer_v3int %sb %uint_0 %idx %uint_9
+ %105 = OpLoad %v3int %104
+ %108 = OpAccessChain %_ptr_StorageBuffer_v3uint %sb %uint_0 %idx %uint_10
+ %109 = OpLoad %v3uint %108
+ %112 = OpAccessChain %_ptr_StorageBuffer_v3half %sb %uint_0 %idx %uint_11
+ %113 = OpLoad %v3half %112
+ %116 = OpAccessChain %_ptr_StorageBuffer_v4float %sb %uint_0 %idx %uint_12
+ %117 = OpLoad %v4float %116
+ %120 = OpAccessChain %_ptr_StorageBuffer_v4int %sb %uint_0 %idx %uint_13
+ %121 = OpLoad %v4int %120
+ %124 = OpAccessChain %_ptr_StorageBuffer_v4uint %sb %uint_0 %idx %uint_14
+ %125 = OpLoad %v4uint %124
+ %128 = OpAccessChain %_ptr_StorageBuffer_v4half %sb %uint_0 %idx %uint_15
+ %129 = OpLoad %v4half %128
+ %132 = OpAccessChain %_ptr_StorageBuffer_mat2v2float %sb %uint_0 %idx %uint_16
+ %133 = OpLoad %mat2v2float %132
+ %136 = OpAccessChain %_ptr_StorageBuffer_mat2v3float %sb %uint_0 %idx %uint_17
+ %137 = OpLoad %mat2v3float %136
+ %140 = OpAccessChain %_ptr_StorageBuffer_mat2v4float %sb %uint_0 %idx %uint_18
+ %141 = OpLoad %mat2v4float %140
+ %144 = OpAccessChain %_ptr_StorageBuffer_mat3v2float %sb %uint_0 %idx %uint_19
+ %145 = OpLoad %mat3v2float %144
+ %148 = OpAccessChain %_ptr_StorageBuffer_mat3v3float %sb %uint_0 %idx %uint_20
+ %149 = OpLoad %mat3v3float %148
+ %152 = OpAccessChain %_ptr_StorageBuffer_mat3v4float %sb %uint_0 %idx %uint_21
+ %153 = OpLoad %mat3v4float %152
+ %156 = OpAccessChain %_ptr_StorageBuffer_mat4v2float %sb %uint_0 %idx %uint_22
+ %157 = OpLoad %mat4v2float %156
+ %160 = OpAccessChain %_ptr_StorageBuffer_mat4v3float %sb %uint_0 %idx %uint_23
+ %161 = OpLoad %mat4v3float %160
+ %164 = OpAccessChain %_ptr_StorageBuffer_mat4v4float %sb %uint_0 %idx %uint_24
+ %165 = OpLoad %mat4v4float %164
+ %168 = OpAccessChain %_ptr_StorageBuffer_mat2v2half %sb %uint_0 %idx %uint_25
+ %169 = OpLoad %mat2v2half %168
+ %172 = OpAccessChain %_ptr_StorageBuffer_mat2v3half %sb %uint_0 %idx %uint_26
+ %173 = OpLoad %mat2v3half %172
+ %176 = OpAccessChain %_ptr_StorageBuffer_mat2v4half %sb %uint_0 %idx %uint_27
+ %177 = OpLoad %mat2v4half %176
+ %180 = OpAccessChain %_ptr_StorageBuffer_mat3v2half %sb %uint_0 %idx %uint_28
+ %181 = OpLoad %mat3v2half %180
+ %184 = OpAccessChain %_ptr_StorageBuffer_mat3v3half %sb %uint_0 %idx %uint_29
+ %185 = OpLoad %mat3v3half %184
+ %188 = OpAccessChain %_ptr_StorageBuffer_mat3v4half %sb %uint_0 %idx %uint_30
+ %189 = OpLoad %mat3v4half %188
+ %192 = OpAccessChain %_ptr_StorageBuffer_mat4v2half %sb %uint_0 %idx %uint_31
+ %193 = OpLoad %mat4v2half %192
+ %196 = OpAccessChain %_ptr_StorageBuffer_mat4v3half %sb %uint_0 %idx %uint_32
+ %197 = OpLoad %mat4v3half %196
+ %200 = OpAccessChain %_ptr_StorageBuffer_mat4v4half %sb %uint_0 %idx %uint_33
+ %201 = OpLoad %mat4v4half %200
+ %204 = OpAccessChain %_ptr_StorageBuffer__arr_v3float_uint_2 %sb %uint_0 %idx %uint_34
+ %205 = OpLoad %_arr_v3float_uint_2 %204
+ %208 = OpAccessChain %_ptr_StorageBuffer__arr_mat4v2half_uint_2 %sb %uint_0 %idx %uint_35
+ %209 = OpLoad %_arr_mat4v2half_uint_2 %208
+ %210 = OpAccessChain %_ptr_StorageBuffer_int %s %uint_0
+ %211 = OpFunctionCall %int %tint_ftoi %70
+ %212 = OpIAdd %int %211 %74
+ %213 = OpBitcast %int %77
+ %214 = OpIAdd %int %212 %213
+ %215 = OpConvertFToS %int %81
+ %216 = OpIAdd %int %214 %215
+ %218 = OpCompositeExtract %float %85 0
+ %217 = OpFunctionCall %int %tint_ftoi %218
+ %219 = OpIAdd %int %216 %217
+ %220 = OpCompositeExtract %int %89 0
+ %221 = OpIAdd %int %219 %220
+ %223 = OpCompositeExtract %uint %93 0
+ %222 = OpBitcast %int %223
+ %224 = OpIAdd %int %221 %222
+ %226 = OpCompositeExtract %half %97 0
+ %225 = OpConvertFToS %int %226
+ %227 = OpIAdd %int %224 %225
+ %229 = OpCompositeExtract %float %101 1
+ %228 = OpFunctionCall %int %tint_ftoi %229
+ %230 = OpIAdd %int %227 %228
+ %231 = OpCompositeExtract %int %105 1
+ %232 = OpIAdd %int %230 %231
+ %234 = OpCompositeExtract %uint %109 1
+ %233 = OpBitcast %int %234
+ %235 = OpIAdd %int %232 %233
+ %237 = OpCompositeExtract %half %113 1
+ %236 = OpConvertFToS %int %237
+ %238 = OpIAdd %int %235 %236
+ %240 = OpCompositeExtract %float %117 2
+ %239 = OpFunctionCall %int %tint_ftoi %240
+ %241 = OpIAdd %int %238 %239
+ %242 = OpCompositeExtract %int %121 2
+ %243 = OpIAdd %int %241 %242
+ %245 = OpCompositeExtract %uint %125 2
+ %244 = OpBitcast %int %245
+ %246 = OpIAdd %int %243 %244
+ %248 = OpCompositeExtract %half %129 2
+ %247 = OpConvertFToS %int %248
+ %249 = OpIAdd %int %246 %247
+ %252 = OpCompositeExtract %v2float %133 0
+ %253 = OpCompositeExtract %float %252 0
+ %250 = OpFunctionCall %int %tint_ftoi %253
+ %254 = OpIAdd %int %249 %250
+ %256 = OpCompositeExtract %v3float %137 0
+ %257 = OpCompositeExtract %float %256 0
+ %255 = OpFunctionCall %int %tint_ftoi %257
+ %258 = OpIAdd %int %254 %255
+ %260 = OpCompositeExtract %v4float %141 0
+ %261 = OpCompositeExtract %float %260 0
+ %259 = OpFunctionCall %int %tint_ftoi %261
+ %262 = OpIAdd %int %258 %259
+ %264 = OpCompositeExtract %v2float %145 0
+ %265 = OpCompositeExtract %float %264 0
+ %263 = OpFunctionCall %int %tint_ftoi %265
+ %266 = OpIAdd %int %262 %263
+ %268 = OpCompositeExtract %v3float %149 0
+ %269 = OpCompositeExtract %float %268 0
+ %267 = OpFunctionCall %int %tint_ftoi %269
+ %270 = OpIAdd %int %266 %267
+ %272 = OpCompositeExtract %v4float %153 0
+ %273 = OpCompositeExtract %float %272 0
+ %271 = OpFunctionCall %int %tint_ftoi %273
+ %274 = OpIAdd %int %270 %271
+ %276 = OpCompositeExtract %v2float %157 0
+ %277 = OpCompositeExtract %float %276 0
+ %275 = OpFunctionCall %int %tint_ftoi %277
+ %278 = OpIAdd %int %274 %275
+ %280 = OpCompositeExtract %v3float %161 0
+ %281 = OpCompositeExtract %float %280 0
+ %279 = OpFunctionCall %int %tint_ftoi %281
+ %282 = OpIAdd %int %278 %279
+ %284 = OpCompositeExtract %v4float %165 0
+ %285 = OpCompositeExtract %float %284 0
+ %283 = OpFunctionCall %int %tint_ftoi %285
+ %286 = OpIAdd %int %282 %283
+ %288 = OpCompositeExtract %v2half %169 0
+ %289 = OpCompositeExtract %half %288 0
+ %287 = OpConvertFToS %int %289
+ %290 = OpIAdd %int %286 %287
+ %292 = OpCompositeExtract %v3half %173 0
+ %293 = OpCompositeExtract %half %292 0
+ %291 = OpConvertFToS %int %293
+ %294 = OpIAdd %int %290 %291
+ %296 = OpCompositeExtract %v4half %177 0
+ %297 = OpCompositeExtract %half %296 0
+ %295 = OpConvertFToS %int %297
+ %298 = OpIAdd %int %294 %295
+ %300 = OpCompositeExtract %v2half %181 0
+ %301 = OpCompositeExtract %half %300 0
+ %299 = OpConvertFToS %int %301
+ %302 = OpIAdd %int %298 %299
+ %304 = OpCompositeExtract %v3half %185 0
+ %305 = OpCompositeExtract %half %304 0
+ %303 = OpConvertFToS %int %305
+ %306 = OpIAdd %int %302 %303
+ %308 = OpCompositeExtract %v4half %189 0
+ %309 = OpCompositeExtract %half %308 0
+ %307 = OpConvertFToS %int %309
+ %310 = OpIAdd %int %306 %307
+ %312 = OpCompositeExtract %v2half %193 0
+ %313 = OpCompositeExtract %half %312 0
+ %311 = OpConvertFToS %int %313
+ %314 = OpIAdd %int %310 %311
+ %316 = OpCompositeExtract %v3half %197 0
+ %317 = OpCompositeExtract %half %316 0
+ %315 = OpConvertFToS %int %317
+ %318 = OpIAdd %int %314 %315
+ %320 = OpCompositeExtract %v4half %201 0
+ %321 = OpCompositeExtract %half %320 0
+ %319 = OpConvertFToS %int %321
+ %322 = OpIAdd %int %318 %319
+ %324 = OpCompositeExtract %mat4v2half %209 0
+ %325 = OpCompositeExtract %v2half %324 0
+ %326 = OpCompositeExtract %half %325 0
+ %323 = OpConvertFToS %int %326
+ %327 = OpIAdd %int %322 %323
+ %329 = OpCompositeExtract %v3float %205 0
+ %330 = OpCompositeExtract %float %329 0
+ %328 = OpFunctionCall %int %tint_ftoi %330
+ %331 = OpIAdd %int %327 %328
+ OpStore %210 %331
OpReturn
OpFunctionEnd
- %main = OpFunction %void None %193
- %195 = OpLabel
- %197 = OpLoad %uint %idx_1
- %196 = OpFunctionCall %void %main_inner %197
+ %main = OpFunction %void None %332
+ %334 = OpLabel
+ %336 = OpLoad %uint %idx_1
+ %335 = OpFunctionCall %void %main_inner %336
OpReturn
OpFunctionEnd
diff --git a/test/tint/buffer/storage/dynamic_index/read_f16.wgsl.expected.wgsl b/test/tint/buffer/storage/dynamic_index/read_f16.wgsl.expected.wgsl
index 5a5e16b..25d4883 100644
--- a/test/tint/buffer/storage/dynamic_index/read_f16.wgsl.expected.wgsl
+++ b/test/tint/buffer/storage/dynamic_index/read_f16.wgsl.expected.wgsl
@@ -45,6 +45,8 @@
@binding(0) @group(0) var<storage, read> sb : S;
+@group(0) @binding(1) var<storage, read_write> s : i32;
+
@compute @workgroup_size(1)
fn main(@builtin(local_invocation_index) idx : u32) {
let scalar_f32 : f32 = sb.arr[idx].scalar_f32;
@@ -83,4 +85,5 @@
let mat4x4_f16 : mat4x4<f16> = sb.arr[idx].mat4x4_f16;
let arr2_vec3_f32 : array<vec3<f32>, 2> = sb.arr[idx].arr2_vec3_f32;
let arr2_mat4x2_f16 : array<mat4x2<f16>, 2> = sb.arr[idx].arr2_mat4x2_f16;
+ s = (((((((((((((((((((((((((((((((((((i32(scalar_f32) + scalar_i32) + i32(scalar_u32)) + i32(scalar_f16)) + i32(vec2_f32.x)) + vec2_i32.x) + i32(vec2_u32.x)) + i32(vec2_f16.x)) + i32(vec3_f32.y)) + vec3_i32.y) + i32(vec3_u32.y)) + i32(vec3_f16.y)) + i32(vec4_f32.z)) + vec4_i32.z) + i32(vec4_u32.z)) + i32(vec4_f16.z)) + i32(mat2x2_f32[0].x)) + i32(mat2x3_f32[0].x)) + i32(mat2x4_f32[0].x)) + i32(mat3x2_f32[0].x)) + i32(mat3x3_f32[0].x)) + i32(mat3x4_f32[0].x)) + i32(mat4x2_f32[0].x)) + i32(mat4x3_f32[0].x)) + i32(mat4x4_f32[0].x)) + i32(mat2x2_f16[0].x)) + i32(mat2x3_f16[0].x)) + i32(mat2x4_f16[0].x)) + i32(mat3x2_f16[0].x)) + i32(mat3x3_f16[0].x)) + i32(mat3x4_f16[0].x)) + i32(mat4x2_f16[0].x)) + i32(mat4x3_f16[0].x)) + i32(mat4x4_f16[0].x)) + i32(arr2_mat4x2_f16[0][0].x)) + i32(arr2_vec3_f32[0].x));
}
diff --git a/test/tint/buffer/storage/static_index/read.wgsl b/test/tint/buffer/storage/static_index/read.wgsl
index 80b9d10..bd62562 100644
--- a/test/tint/buffer/storage/static_index/read.wgsl
+++ b/test/tint/buffer/storage/static_index/read.wgsl
@@ -31,6 +31,7 @@
};
@binding(0) @group(0) var<storage, read> sb : S;
+@group(0) @binding(1) var<storage, read_write> s: i32;
@compute @workgroup_size(1)
fn main() {
@@ -58,4 +59,15 @@
let arr2_vec3_f32 = sb.arr2_vec3_f32;
let struct_inner = sb.struct_inner;
let array_struct_inner = sb.array_struct_inner;
+
+ s = i32(scalar_f32) + scalar_i32 + i32(scalar_u32) +
+ i32(vec2_f32.x) + vec2_i32.x + i32(vec2_u32.x) +
+ i32(vec3_f32.y) + vec3_i32.y + i32(vec3_u32.y) +
+ i32(vec4_f32.z) + vec4_i32.z + i32(vec4_u32.z) +
+ i32(mat2x2_f32[0].x) + i32(mat2x3_f32[0].x) +
+ i32(mat2x4_f32[0].x) + i32(mat3x2_f32[0].x) +
+ i32(mat3x3_f32[0].x) + i32(mat3x4_f32[0].x) +
+ i32(mat4x2_f32[0].x) + i32(mat4x3_f32[0].x) +
+ i32(mat4x4_f32[0].x) + i32(arr2_vec3_f32[0].x) +
+ struct_inner.scalar_i32 + array_struct_inner[0].scalar_i32;
}
diff --git a/test/tint/buffer/storage/static_index/read.wgsl.expected.dxc.hlsl b/test/tint/buffer/storage/static_index/read.wgsl.expected.dxc.hlsl
index 0b684a5..aa94d2f 100644
--- a/test/tint/buffer/storage/static_index/read.wgsl.expected.dxc.hlsl
+++ b/test/tint/buffer/storage/static_index/read.wgsl.expected.dxc.hlsl
@@ -1,9 +1,14 @@
+int tint_ftoi(float v) {
+ return ((v < 2147483520.0f) ? ((v < -2147483648.0f) ? -2147483648 : int(v)) : 2147483647);
+}
+
struct Inner {
int scalar_i32;
float scalar_f32;
};
ByteAddressBuffer sb : register(t0);
+RWByteAddressBuffer s : register(u1);
float2x2 sb_load_12(uint offset) {
return float2x2(asfloat(sb.Load2((offset + 0u))), asfloat(sb.Load2((offset + 8u))));
@@ -94,5 +99,6 @@
const float3 arr2_vec3_f32[2] = sb_load_21(512u);
const Inner struct_inner = sb_load_22(544u);
const Inner array_struct_inner[4] = sb_load_23(552u);
+ s.Store(0u, asuint((((((((((((((((((((((((tint_ftoi(scalar_f32) + scalar_i32) + int(scalar_u32)) + tint_ftoi(vec2_f32.x)) + vec2_i32.x) + int(vec2_u32.x)) + tint_ftoi(vec3_f32.y)) + vec3_i32.y) + int(vec3_u32.y)) + tint_ftoi(vec4_f32.z)) + vec4_i32.z) + int(vec4_u32.z)) + tint_ftoi(mat2x2_f32[0].x)) + tint_ftoi(mat2x3_f32[0].x)) + tint_ftoi(mat2x4_f32[0].x)) + tint_ftoi(mat3x2_f32[0].x)) + tint_ftoi(mat3x3_f32[0].x)) + tint_ftoi(mat3x4_f32[0].x)) + tint_ftoi(mat4x2_f32[0].x)) + tint_ftoi(mat4x3_f32[0].x)) + tint_ftoi(mat4x4_f32[0].x)) + tint_ftoi(arr2_vec3_f32[0].x)) + struct_inner.scalar_i32) + array_struct_inner[0].scalar_i32)));
return;
}
diff --git a/test/tint/buffer/storage/static_index/read.wgsl.expected.fxc.hlsl b/test/tint/buffer/storage/static_index/read.wgsl.expected.fxc.hlsl
index 0b684a5..aa94d2f 100644
--- a/test/tint/buffer/storage/static_index/read.wgsl.expected.fxc.hlsl
+++ b/test/tint/buffer/storage/static_index/read.wgsl.expected.fxc.hlsl
@@ -1,9 +1,14 @@
+int tint_ftoi(float v) {
+ return ((v < 2147483520.0f) ? ((v < -2147483648.0f) ? -2147483648 : int(v)) : 2147483647);
+}
+
struct Inner {
int scalar_i32;
float scalar_f32;
};
ByteAddressBuffer sb : register(t0);
+RWByteAddressBuffer s : register(u1);
float2x2 sb_load_12(uint offset) {
return float2x2(asfloat(sb.Load2((offset + 0u))), asfloat(sb.Load2((offset + 8u))));
@@ -94,5 +99,6 @@
const float3 arr2_vec3_f32[2] = sb_load_21(512u);
const Inner struct_inner = sb_load_22(544u);
const Inner array_struct_inner[4] = sb_load_23(552u);
+ s.Store(0u, asuint((((((((((((((((((((((((tint_ftoi(scalar_f32) + scalar_i32) + int(scalar_u32)) + tint_ftoi(vec2_f32.x)) + vec2_i32.x) + int(vec2_u32.x)) + tint_ftoi(vec3_f32.y)) + vec3_i32.y) + int(vec3_u32.y)) + tint_ftoi(vec4_f32.z)) + vec4_i32.z) + int(vec4_u32.z)) + tint_ftoi(mat2x2_f32[0].x)) + tint_ftoi(mat2x3_f32[0].x)) + tint_ftoi(mat2x4_f32[0].x)) + tint_ftoi(mat3x2_f32[0].x)) + tint_ftoi(mat3x3_f32[0].x)) + tint_ftoi(mat3x4_f32[0].x)) + tint_ftoi(mat4x2_f32[0].x)) + tint_ftoi(mat4x3_f32[0].x)) + tint_ftoi(mat4x4_f32[0].x)) + tint_ftoi(arr2_vec3_f32[0].x)) + struct_inner.scalar_i32) + array_struct_inner[0].scalar_i32)));
return;
}
diff --git a/test/tint/buffer/storage/static_index/read.wgsl.expected.glsl b/test/tint/buffer/storage/static_index/read.wgsl.expected.glsl
index 59613b5..185f228 100644
--- a/test/tint/buffer/storage/static_index/read.wgsl.expected.glsl
+++ b/test/tint/buffer/storage/static_index/read.wgsl.expected.glsl
@@ -1,5 +1,9 @@
#version 310 es
+int tint_ftoi(float v) {
+ return ((v < 2147483520.0f) ? ((v < -2147483648.0f) ? (-2147483647 - 1) : int(v)) : 2147483647);
+}
+
struct Inner {
int scalar_i32;
float scalar_f32;
@@ -46,6 +50,10 @@
S inner;
} sb;
+layout(binding = 1, std430) buffer s_block_ssbo {
+ int inner;
+} s;
+
void tint_symbol() {
float scalar_f32 = sb.inner.scalar_f32;
int scalar_i32 = sb.inner.scalar_i32;
@@ -71,6 +79,7 @@
vec3 arr2_vec3_f32[2] = sb.inner.arr2_vec3_f32;
Inner struct_inner = sb.inner.struct_inner;
Inner array_struct_inner[4] = sb.inner.array_struct_inner;
+ s.inner = (((((((((((((((((((((((tint_ftoi(scalar_f32) + scalar_i32) + int(scalar_u32)) + tint_ftoi(vec2_f32.x)) + vec2_i32.x) + int(vec2_u32.x)) + tint_ftoi(vec3_f32.y)) + vec3_i32.y) + int(vec3_u32.y)) + tint_ftoi(vec4_f32.z)) + vec4_i32.z) + int(vec4_u32.z)) + tint_ftoi(mat2x2_f32[0].x)) + tint_ftoi(mat2x3_f32[0].x)) + tint_ftoi(mat2x4_f32[0].x)) + tint_ftoi(mat3x2_f32[0].x)) + tint_ftoi(mat3x3_f32[0].x)) + tint_ftoi(mat3x4_f32[0].x)) + tint_ftoi(mat4x2_f32[0].x)) + tint_ftoi(mat4x3_f32[0].x)) + tint_ftoi(mat4x4_f32[0].x)) + tint_ftoi(arr2_vec3_f32[0].x)) + struct_inner.scalar_i32) + array_struct_inner[0].scalar_i32);
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
diff --git a/test/tint/buffer/storage/static_index/read.wgsl.expected.msl b/test/tint/buffer/storage/static_index/read.wgsl.expected.msl
index df35c27..6293a49 100644
--- a/test/tint/buffer/storage/static_index/read.wgsl.expected.msl
+++ b/test/tint/buffer/storage/static_index/read.wgsl.expected.msl
@@ -90,6 +90,10 @@
return result;
}
+int tint_ftoi(float v) {
+ return select(2147483647, select(int(v), (-2147483647 - 1), (v < -2147483648.0f)), (v < 2147483520.0f));
+}
+
struct S {
float scalar_f32;
int scalar_i32;
@@ -117,7 +121,7 @@
tint_array<Inner, 4> array_struct_inner;
};
-kernel void tint_symbol(const device S_tint_packed_vec3* tint_symbol_1 [[buffer(0)]]) {
+kernel void tint_symbol(const device S_tint_packed_vec3* tint_symbol_1 [[buffer(1)]], device int* tint_symbol_2 [[buffer(0)]]) {
float const scalar_f32 = (*(tint_symbol_1)).scalar_f32;
int const scalar_i32 = (*(tint_symbol_1)).scalar_i32;
uint const scalar_u32 = (*(tint_symbol_1)).scalar_u32;
@@ -142,6 +146,7 @@
tint_array<float3, 2> const arr2_vec3_f32 = tint_unpack_vec3_in_composite_3((*(tint_symbol_1)).arr2_vec3_f32);
Inner const struct_inner = (*(tint_symbol_1)).struct_inner;
tint_array<Inner, 4> const array_struct_inner = (*(tint_symbol_1)).array_struct_inner;
+ *(tint_symbol_2) = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(tint_ftoi(scalar_f32)) + as_type<uint>(scalar_i32)))) + as_type<uint>(int(scalar_u32))))) + as_type<uint>(tint_ftoi(vec2_f32[0]))))) + as_type<uint>(vec2_i32[0])))) + as_type<uint>(int(vec2_u32[0]))))) + as_type<uint>(tint_ftoi(vec3_f32[1]))))) + as_type<uint>(vec3_i32[1])))) + as_type<uint>(int(vec3_u32[1]))))) + as_type<uint>(tint_ftoi(vec4_f32[2]))))) + as_type<uint>(vec4_i32[2])))) + as_type<uint>(int(vec4_u32[2]))))) + as_type<uint>(tint_ftoi(mat2x2_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat2x3_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat2x4_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat3x2_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat3x3_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat3x4_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat4x2_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat4x3_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat4x4_f32[0][0]))))) + as_type<uint>(tint_ftoi(arr2_vec3_f32[0][0]))))) + as_type<uint>(struct_inner.scalar_i32)))) + as_type<uint>(array_struct_inner[0].scalar_i32)));
return;
}
diff --git a/test/tint/buffer/storage/static_index/read.wgsl.expected.spvasm b/test/tint/buffer/storage/static_index/read.wgsl.expected.spvasm
index c137494..0886466 100644
--- a/test/tint/buffer/storage/static_index/read.wgsl.expected.spvasm
+++ b/test/tint/buffer/storage/static_index/read.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 129
+; Bound: 221
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -38,6 +38,11 @@
OpMemberName %Inner 1 "scalar_f32"
OpMemberName %S 23 "array_struct_inner"
OpName %sb "sb"
+ OpName %s_block "s_block"
+ OpMemberName %s_block 0 "inner"
+ OpName %s "s"
+ OpName %tint_ftoi "tint_ftoi"
+ OpName %v "v"
OpName %main "main"
OpDecorate %sb_block Block
OpMemberDecorate %sb_block 0 Offset 0
@@ -90,6 +95,10 @@
OpDecorate %sb NonWritable
OpDecorate %sb Binding 0
OpDecorate %sb DescriptorSet 0
+ OpDecorate %s_block Block
+ OpMemberDecorate %s_block 0 Offset 0
+ OpDecorate %s DescriptorSet 0
+ OpDecorate %s Binding 1
%float = OpTypeFloat 32
%int = OpTypeInt 32 1
%uint = OpTypeInt 32 0
@@ -120,8 +129,17 @@
%sb_block = OpTypeStruct %S
%_ptr_StorageBuffer_sb_block = OpTypePointer StorageBuffer %sb_block
%sb = OpVariable %_ptr_StorageBuffer_sb_block StorageBuffer
+ %s_block = OpTypeStruct %int
+%_ptr_StorageBuffer_s_block = OpTypePointer StorageBuffer %s_block
+ %s = OpVariable %_ptr_StorageBuffer_s_block StorageBuffer
+ %34 = OpTypeFunction %int %float
+%float_2_14748352e_09 = OpConstant %float 2.14748352e+09
+ %bool = OpTypeBool
+%float_n2_14748365e_09 = OpConstant %float -2.14748365e+09
+%int_n2147483648 = OpConstant %int -2147483648
+%int_2147483647 = OpConstant %int 2147483647
%void = OpTypeVoid
- %31 = OpTypeFunction %void
+ %48 = OpTypeFunction %void
%uint_0 = OpConstant %uint 0
%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
%uint_1 = OpConstant %uint 1
@@ -168,55 +186,141 @@
%_ptr_StorageBuffer_Inner = OpTypePointer StorageBuffer %Inner
%uint_23 = OpConstant %uint 23
%_ptr_StorageBuffer__arr_Inner_uint_4 = OpTypePointer StorageBuffer %_arr_Inner_uint_4
- %main = OpFunction %void None %31
- %34 = OpLabel
- %37 = OpAccessChain %_ptr_StorageBuffer_float %sb %uint_0 %uint_0
- %38 = OpLoad %float %37
- %41 = OpAccessChain %_ptr_StorageBuffer_int %sb %uint_0 %uint_1
- %42 = OpLoad %int %41
- %44 = OpAccessChain %_ptr_StorageBuffer_uint %sb %uint_0 %uint_2
- %45 = OpLoad %uint %44
- %48 = OpAccessChain %_ptr_StorageBuffer_v2float %sb %uint_0 %uint_3
- %49 = OpLoad %v2float %48
- %51 = OpAccessChain %_ptr_StorageBuffer_v2int %sb %uint_0 %uint_4
- %52 = OpLoad %v2int %51
- %55 = OpAccessChain %_ptr_StorageBuffer_v2uint %sb %uint_0 %uint_5
- %56 = OpLoad %v2uint %55
- %59 = OpAccessChain %_ptr_StorageBuffer_v3float %sb %uint_0 %uint_6
- %60 = OpLoad %v3float %59
- %63 = OpAccessChain %_ptr_StorageBuffer_v3int %sb %uint_0 %uint_7
- %64 = OpLoad %v3int %63
- %67 = OpAccessChain %_ptr_StorageBuffer_v3uint %sb %uint_0 %uint_8
- %68 = OpLoad %v3uint %67
- %71 = OpAccessChain %_ptr_StorageBuffer_v4float %sb %uint_0 %uint_9
- %72 = OpLoad %v4float %71
- %75 = OpAccessChain %_ptr_StorageBuffer_v4int %sb %uint_0 %uint_10
- %76 = OpLoad %v4int %75
- %79 = OpAccessChain %_ptr_StorageBuffer_v4uint %sb %uint_0 %uint_11
- %80 = OpLoad %v4uint %79
- %83 = OpAccessChain %_ptr_StorageBuffer_mat2v2float %sb %uint_0 %uint_12
- %84 = OpLoad %mat2v2float %83
- %87 = OpAccessChain %_ptr_StorageBuffer_mat2v3float %sb %uint_0 %uint_13
- %88 = OpLoad %mat2v3float %87
- %91 = OpAccessChain %_ptr_StorageBuffer_mat2v4float %sb %uint_0 %uint_14
- %92 = OpLoad %mat2v4float %91
- %95 = OpAccessChain %_ptr_StorageBuffer_mat3v2float %sb %uint_0 %uint_15
- %96 = OpLoad %mat3v2float %95
- %99 = OpAccessChain %_ptr_StorageBuffer_mat3v3float %sb %uint_0 %uint_16
- %100 = OpLoad %mat3v3float %99
- %103 = OpAccessChain %_ptr_StorageBuffer_mat3v4float %sb %uint_0 %uint_17
- %104 = OpLoad %mat3v4float %103
- %107 = OpAccessChain %_ptr_StorageBuffer_mat4v2float %sb %uint_0 %uint_18
- %108 = OpLoad %mat4v2float %107
- %111 = OpAccessChain %_ptr_StorageBuffer_mat4v3float %sb %uint_0 %uint_19
- %112 = OpLoad %mat4v3float %111
- %115 = OpAccessChain %_ptr_StorageBuffer_mat4v4float %sb %uint_0 %uint_20
- %116 = OpLoad %mat4v4float %115
- %119 = OpAccessChain %_ptr_StorageBuffer__arr_v3float_uint_2 %sb %uint_0 %uint_21
- %120 = OpLoad %_arr_v3float_uint_2 %119
- %123 = OpAccessChain %_ptr_StorageBuffer_Inner %sb %uint_0 %uint_22
- %124 = OpLoad %Inner %123
- %127 = OpAccessChain %_ptr_StorageBuffer__arr_Inner_uint_4 %sb %uint_0 %uint_23
- %128 = OpLoad %_arr_Inner_uint_4 %127
+ %176 = OpConstantNull %int
+ %tint_ftoi = OpFunction %int None %34
+ %v = OpFunctionParameter %float
+ %37 = OpLabel
+ %40 = OpFOrdLessThan %bool %v %float_2_14748352e_09
+ %44 = OpFOrdLessThan %bool %v %float_n2_14748365e_09
+ %46 = OpConvertFToS %int %v
+ %42 = OpSelect %int %44 %int_n2147483648 %46
+ %38 = OpSelect %int %40 %42 %int_2147483647
+ OpReturnValue %38
+ OpFunctionEnd
+ %main = OpFunction %void None %48
+ %51 = OpLabel
+ %54 = OpAccessChain %_ptr_StorageBuffer_float %sb %uint_0 %uint_0
+ %55 = OpLoad %float %54
+ %58 = OpAccessChain %_ptr_StorageBuffer_int %sb %uint_0 %uint_1
+ %59 = OpLoad %int %58
+ %61 = OpAccessChain %_ptr_StorageBuffer_uint %sb %uint_0 %uint_2
+ %62 = OpLoad %uint %61
+ %65 = OpAccessChain %_ptr_StorageBuffer_v2float %sb %uint_0 %uint_3
+ %66 = OpLoad %v2float %65
+ %68 = OpAccessChain %_ptr_StorageBuffer_v2int %sb %uint_0 %uint_4
+ %69 = OpLoad %v2int %68
+ %72 = OpAccessChain %_ptr_StorageBuffer_v2uint %sb %uint_0 %uint_5
+ %73 = OpLoad %v2uint %72
+ %76 = OpAccessChain %_ptr_StorageBuffer_v3float %sb %uint_0 %uint_6
+ %77 = OpLoad %v3float %76
+ %80 = OpAccessChain %_ptr_StorageBuffer_v3int %sb %uint_0 %uint_7
+ %81 = OpLoad %v3int %80
+ %84 = OpAccessChain %_ptr_StorageBuffer_v3uint %sb %uint_0 %uint_8
+ %85 = OpLoad %v3uint %84
+ %88 = OpAccessChain %_ptr_StorageBuffer_v4float %sb %uint_0 %uint_9
+ %89 = OpLoad %v4float %88
+ %92 = OpAccessChain %_ptr_StorageBuffer_v4int %sb %uint_0 %uint_10
+ %93 = OpLoad %v4int %92
+ %96 = OpAccessChain %_ptr_StorageBuffer_v4uint %sb %uint_0 %uint_11
+ %97 = OpLoad %v4uint %96
+ %100 = OpAccessChain %_ptr_StorageBuffer_mat2v2float %sb %uint_0 %uint_12
+ %101 = OpLoad %mat2v2float %100
+ %104 = OpAccessChain %_ptr_StorageBuffer_mat2v3float %sb %uint_0 %uint_13
+ %105 = OpLoad %mat2v3float %104
+ %108 = OpAccessChain %_ptr_StorageBuffer_mat2v4float %sb %uint_0 %uint_14
+ %109 = OpLoad %mat2v4float %108
+ %112 = OpAccessChain %_ptr_StorageBuffer_mat3v2float %sb %uint_0 %uint_15
+ %113 = OpLoad %mat3v2float %112
+ %116 = OpAccessChain %_ptr_StorageBuffer_mat3v3float %sb %uint_0 %uint_16
+ %117 = OpLoad %mat3v3float %116
+ %120 = OpAccessChain %_ptr_StorageBuffer_mat3v4float %sb %uint_0 %uint_17
+ %121 = OpLoad %mat3v4float %120
+ %124 = OpAccessChain %_ptr_StorageBuffer_mat4v2float %sb %uint_0 %uint_18
+ %125 = OpLoad %mat4v2float %124
+ %128 = OpAccessChain %_ptr_StorageBuffer_mat4v3float %sb %uint_0 %uint_19
+ %129 = OpLoad %mat4v3float %128
+ %132 = OpAccessChain %_ptr_StorageBuffer_mat4v4float %sb %uint_0 %uint_20
+ %133 = OpLoad %mat4v4float %132
+ %136 = OpAccessChain %_ptr_StorageBuffer__arr_v3float_uint_2 %sb %uint_0 %uint_21
+ %137 = OpLoad %_arr_v3float_uint_2 %136
+ %140 = OpAccessChain %_ptr_StorageBuffer_Inner %sb %uint_0 %uint_22
+ %141 = OpLoad %Inner %140
+ %144 = OpAccessChain %_ptr_StorageBuffer__arr_Inner_uint_4 %sb %uint_0 %uint_23
+ %145 = OpLoad %_arr_Inner_uint_4 %144
+ %146 = OpAccessChain %_ptr_StorageBuffer_int %s %uint_0
+ %147 = OpFunctionCall %int %tint_ftoi %55
+ %148 = OpIAdd %int %147 %59
+ %149 = OpBitcast %int %62
+ %150 = OpIAdd %int %148 %149
+ %152 = OpCompositeExtract %float %66 0
+ %151 = OpFunctionCall %int %tint_ftoi %152
+ %153 = OpIAdd %int %150 %151
+ %154 = OpCompositeExtract %int %69 0
+ %155 = OpIAdd %int %153 %154
+ %157 = OpCompositeExtract %uint %73 0
+ %156 = OpBitcast %int %157
+ %158 = OpIAdd %int %155 %156
+ %160 = OpCompositeExtract %float %77 1
+ %159 = OpFunctionCall %int %tint_ftoi %160
+ %161 = OpIAdd %int %158 %159
+ %162 = OpCompositeExtract %int %81 1
+ %163 = OpIAdd %int %161 %162
+ %165 = OpCompositeExtract %uint %85 1
+ %164 = OpBitcast %int %165
+ %166 = OpIAdd %int %163 %164
+ %168 = OpCompositeExtract %float %89 2
+ %167 = OpFunctionCall %int %tint_ftoi %168
+ %169 = OpIAdd %int %166 %167
+ %170 = OpCompositeExtract %int %93 2
+ %171 = OpIAdd %int %169 %170
+ %173 = OpCompositeExtract %uint %97 2
+ %172 = OpBitcast %int %173
+ %174 = OpIAdd %int %171 %172
+ %177 = OpCompositeExtract %v2float %101 0
+ %178 = OpCompositeExtract %float %177 0
+ %175 = OpFunctionCall %int %tint_ftoi %178
+ %179 = OpIAdd %int %174 %175
+ %181 = OpCompositeExtract %v3float %105 0
+ %182 = OpCompositeExtract %float %181 0
+ %180 = OpFunctionCall %int %tint_ftoi %182
+ %183 = OpIAdd %int %179 %180
+ %185 = OpCompositeExtract %v4float %109 0
+ %186 = OpCompositeExtract %float %185 0
+ %184 = OpFunctionCall %int %tint_ftoi %186
+ %187 = OpIAdd %int %183 %184
+ %189 = OpCompositeExtract %v2float %113 0
+ %190 = OpCompositeExtract %float %189 0
+ %188 = OpFunctionCall %int %tint_ftoi %190
+ %191 = OpIAdd %int %187 %188
+ %193 = OpCompositeExtract %v3float %117 0
+ %194 = OpCompositeExtract %float %193 0
+ %192 = OpFunctionCall %int %tint_ftoi %194
+ %195 = OpIAdd %int %191 %192
+ %197 = OpCompositeExtract %v4float %121 0
+ %198 = OpCompositeExtract %float %197 0
+ %196 = OpFunctionCall %int %tint_ftoi %198
+ %199 = OpIAdd %int %195 %196
+ %201 = OpCompositeExtract %v2float %125 0
+ %202 = OpCompositeExtract %float %201 0
+ %200 = OpFunctionCall %int %tint_ftoi %202
+ %203 = OpIAdd %int %199 %200
+ %205 = OpCompositeExtract %v3float %129 0
+ %206 = OpCompositeExtract %float %205 0
+ %204 = OpFunctionCall %int %tint_ftoi %206
+ %207 = OpIAdd %int %203 %204
+ %209 = OpCompositeExtract %v4float %133 0
+ %210 = OpCompositeExtract %float %209 0
+ %208 = OpFunctionCall %int %tint_ftoi %210
+ %211 = OpIAdd %int %207 %208
+ %213 = OpCompositeExtract %v3float %137 0
+ %214 = OpCompositeExtract %float %213 0
+ %212 = OpFunctionCall %int %tint_ftoi %214
+ %215 = OpIAdd %int %211 %212
+ %216 = OpCompositeExtract %int %141 0
+ %217 = OpIAdd %int %215 %216
+ %218 = OpCompositeExtract %Inner %145 0
+ %219 = OpCompositeExtract %int %218 0
+ %220 = OpIAdd %int %217 %219
+ OpStore %146 %220
OpReturn
OpFunctionEnd
diff --git a/test/tint/buffer/storage/static_index/read.wgsl.expected.wgsl b/test/tint/buffer/storage/static_index/read.wgsl.expected.wgsl
index a978937..c125b1a 100644
--- a/test/tint/buffer/storage/static_index/read.wgsl.expected.wgsl
+++ b/test/tint/buffer/storage/static_index/read.wgsl.expected.wgsl
@@ -32,6 +32,8 @@
@binding(0) @group(0) var<storage, read> sb : S;
+@group(0) @binding(1) var<storage, read_write> s : i32;
+
@compute @workgroup_size(1)
fn main() {
let scalar_f32 = sb.scalar_f32;
@@ -58,4 +60,5 @@
let arr2_vec3_f32 = sb.arr2_vec3_f32;
let struct_inner = sb.struct_inner;
let array_struct_inner = sb.array_struct_inner;
+ s = (((((((((((((((((((((((i32(scalar_f32) + scalar_i32) + i32(scalar_u32)) + i32(vec2_f32.x)) + vec2_i32.x) + i32(vec2_u32.x)) + i32(vec3_f32.y)) + vec3_i32.y) + i32(vec3_u32.y)) + i32(vec4_f32.z)) + vec4_i32.z) + i32(vec4_u32.z)) + i32(mat2x2_f32[0].x)) + i32(mat2x3_f32[0].x)) + i32(mat2x4_f32[0].x)) + i32(mat3x2_f32[0].x)) + i32(mat3x3_f32[0].x)) + i32(mat3x4_f32[0].x)) + i32(mat4x2_f32[0].x)) + i32(mat4x3_f32[0].x)) + i32(mat4x4_f32[0].x)) + i32(arr2_vec3_f32[0].x)) + struct_inner.scalar_i32) + array_struct_inner[0].scalar_i32);
}
diff --git a/test/tint/buffer/storage/static_index/read_f16.wgsl b/test/tint/buffer/storage/static_index/read_f16.wgsl
index 5e8d2d1..0e84707 100644
--- a/test/tint/buffer/storage/static_index/read_f16.wgsl
+++ b/test/tint/buffer/storage/static_index/read_f16.wgsl
@@ -48,6 +48,7 @@
};
@binding(0) @group(0) var<storage, read> sb : S;
+@group(0) @binding(1) var<storage, read_write> s: i32;
@compute @workgroup_size(1)
fn main() {
@@ -89,4 +90,21 @@
let arr2_mat4x2_f16 = sb.arr2_mat4x2_f16;
let struct_inner = sb.struct_inner;
let array_struct_inner = sb.array_struct_inner;
+
+ s = i32(scalar_f32) + scalar_i32 + i32(scalar_u32) + i32(scalar_f16) +
+ i32(vec2_f32.x) + vec2_i32.x + i32(vec2_u32.x) + i32(vec2_f16.x) +
+ i32(vec3_f32.y) + vec3_i32.y + i32(vec3_u32.y) + i32(vec3_f16.y) +
+ i32(vec4_f32.z) + vec4_i32.z + i32(vec4_u32.z) + i32(vec4_f16.z) +
+ i32(mat2x2_f32[0].x) + i32(mat2x3_f32[0].x) +
+ i32(mat2x4_f32[0].x) + i32(mat3x2_f32[0].x) +
+ i32(mat3x3_f32[0].x) + i32(mat3x4_f32[0].x) +
+ i32(mat4x2_f32[0].x) + i32(mat4x3_f32[0].x) +
+ i32(mat4x4_f32[0].x) +
+ i32(mat2x2_f16[0].x) + i32(mat2x3_f16[0].x) +
+ i32(mat2x4_f16[0].x) + i32(mat3x2_f16[0].x) +
+ i32(mat3x3_f16[0].x) + i32(mat3x4_f16[0].x) +
+ i32(mat4x2_f16[0].x) + i32(mat4x3_f16[0].x) +
+ i32(mat4x4_f16[0].x) +
+ i32(arr2_vec3_f32[0].x) + i32(arr2_mat4x2_f16[0][0].x) +
+ struct_inner.scalar_i32 + array_struct_inner[0].scalar_i32;
}
diff --git a/test/tint/buffer/storage/static_index/read_f16.wgsl.expected.dxc.hlsl b/test/tint/buffer/storage/static_index/read_f16.wgsl.expected.dxc.hlsl
index 361d496..ed43aad 100644
--- a/test/tint/buffer/storage/static_index/read_f16.wgsl.expected.dxc.hlsl
+++ b/test/tint/buffer/storage/static_index/read_f16.wgsl.expected.dxc.hlsl
@@ -1,3 +1,7 @@
+int tint_ftoi(float v) {
+ return ((v < 2147483520.0f) ? ((v < -2147483648.0f) ? -2147483648 : int(v)) : 2147483647);
+}
+
struct Inner {
int scalar_i32;
float scalar_f32;
@@ -5,6 +9,7 @@
};
ByteAddressBuffer sb : register(t0);
+RWByteAddressBuffer s : register(u1);
float2x2 sb_load_16(uint offset) {
return float2x2(asfloat(sb.Load2((offset + 0u))), asfloat(sb.Load2((offset + 8u))));
@@ -156,5 +161,6 @@
matrix<float16_t, 4, 2> arr2_mat4x2_f16[2] = sb_load_35(768u);
const Inner struct_inner = sb_load_36(800u);
const Inner array_struct_inner[4] = sb_load_37(812u);
+ s.Store(0u, asuint((((((((((((((((((((((((((((((((((((((tint_ftoi(scalar_f32) + scalar_i32) + int(scalar_u32)) + int(scalar_f16)) + tint_ftoi(vec2_f32.x)) + vec2_i32.x) + int(vec2_u32.x)) + int(vec2_f16.x)) + tint_ftoi(vec3_f32.y)) + vec3_i32.y) + int(vec3_u32.y)) + int(vec3_f16.y)) + tint_ftoi(vec4_f32.z)) + vec4_i32.z) + int(vec4_u32.z)) + int(vec4_f16.z)) + tint_ftoi(mat2x2_f32[0].x)) + tint_ftoi(mat2x3_f32[0].x)) + tint_ftoi(mat2x4_f32[0].x)) + tint_ftoi(mat3x2_f32[0].x)) + tint_ftoi(mat3x3_f32[0].x)) + tint_ftoi(mat3x4_f32[0].x)) + tint_ftoi(mat4x2_f32[0].x)) + tint_ftoi(mat4x3_f32[0].x)) + tint_ftoi(mat4x4_f32[0].x)) + int(mat2x2_f16[0].x)) + int(mat2x3_f16[0].x)) + int(mat2x4_f16[0].x)) + int(mat3x2_f16[0].x)) + int(mat3x3_f16[0].x)) + int(mat3x4_f16[0].x)) + int(mat4x2_f16[0].x)) + int(mat4x3_f16[0].x)) + int(mat4x4_f16[0].x)) + tint_ftoi(arr2_vec3_f32[0].x)) + int(arr2_mat4x2_f16[0][0].x)) + struct_inner.scalar_i32) + array_struct_inner[0].scalar_i32)));
return;
}
diff --git a/test/tint/buffer/storage/static_index/read_f16.wgsl.expected.glsl b/test/tint/buffer/storage/static_index/read_f16.wgsl.expected.glsl
index 49d5d53..d6e11f6 100644
--- a/test/tint/buffer/storage/static_index/read_f16.wgsl.expected.glsl
+++ b/test/tint/buffer/storage/static_index/read_f16.wgsl.expected.glsl
@@ -1,6 +1,10 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
+int tint_ftoi(float v) {
+ return ((v < 2147483520.0f) ? ((v < -2147483648.0f) ? (-2147483647 - 1) : int(v)) : 2147483647);
+}
+
struct Inner {
int scalar_i32;
float scalar_f32;
@@ -66,6 +70,10 @@
S inner;
} sb;
+layout(binding = 1, std430) buffer s_block_ssbo {
+ int inner;
+} s;
+
void tint_symbol() {
float scalar_f32 = sb.inner.scalar_f32;
int scalar_i32 = sb.inner.scalar_i32;
@@ -105,6 +113,7 @@
f16mat4x2 arr2_mat4x2_f16[2] = sb.inner.arr2_mat4x2_f16;
Inner struct_inner = sb.inner.struct_inner;
Inner array_struct_inner[4] = sb.inner.array_struct_inner;
+ s.inner = (((((((((((((((((((((((((((((((((((((tint_ftoi(scalar_f32) + scalar_i32) + int(scalar_u32)) + int(scalar_f16)) + tint_ftoi(vec2_f32.x)) + vec2_i32.x) + int(vec2_u32.x)) + int(vec2_f16.x)) + tint_ftoi(vec3_f32.y)) + vec3_i32.y) + int(vec3_u32.y)) + int(vec3_f16.y)) + tint_ftoi(vec4_f32.z)) + vec4_i32.z) + int(vec4_u32.z)) + int(vec4_f16.z)) + tint_ftoi(mat2x2_f32[0].x)) + tint_ftoi(mat2x3_f32[0].x)) + tint_ftoi(mat2x4_f32[0].x)) + tint_ftoi(mat3x2_f32[0].x)) + tint_ftoi(mat3x3_f32[0].x)) + tint_ftoi(mat3x4_f32[0].x)) + tint_ftoi(mat4x2_f32[0].x)) + tint_ftoi(mat4x3_f32[0].x)) + tint_ftoi(mat4x4_f32[0].x)) + int(mat2x2_f16[0].x)) + int(mat2x3_f16[0].x)) + int(mat2x4_f16[0].x)) + int(mat3x2_f16[0].x)) + int(mat3x3_f16[0].x)) + int(mat3x4_f16[0].x)) + int(mat4x2_f16[0].x)) + int(mat4x3_f16[0].x)) + int(mat4x4_f16[0].x)) + tint_ftoi(arr2_vec3_f32[0].x)) + int(arr2_mat4x2_f16[0][0].x)) + struct_inner.scalar_i32) + array_struct_inner[0].scalar_i32);
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
diff --git a/test/tint/buffer/storage/static_index/read_f16.wgsl.expected.msl b/test/tint/buffer/storage/static_index/read_f16.wgsl.expected.msl
index 18267f3..66ae4dd 100644
--- a/test/tint/buffer/storage/static_index/read_f16.wgsl.expected.msl
+++ b/test/tint/buffer/storage/static_index/read_f16.wgsl.expected.msl
@@ -139,6 +139,10 @@
return result;
}
+int tint_ftoi(float v) {
+ return select(2147483647, select(int(v), (-2147483647 - 1), (v < -2147483648.0f)), (v < 2147483520.0f));
+}
+
struct S {
float scalar_f32;
int scalar_i32;
@@ -180,7 +184,7 @@
tint_array<Inner, 4> array_struct_inner;
};
-kernel void tint_symbol(const device S_tint_packed_vec3* tint_symbol_1 [[buffer(0)]]) {
+kernel void tint_symbol(const device S_tint_packed_vec3* tint_symbol_1 [[buffer(1)]], device int* tint_symbol_2 [[buffer(0)]]) {
float const scalar_f32 = (*(tint_symbol_1)).scalar_f32;
int const scalar_i32 = (*(tint_symbol_1)).scalar_i32;
uint const scalar_u32 = (*(tint_symbol_1)).scalar_u32;
@@ -219,6 +223,7 @@
tint_array<half4x2, 2> const arr2_mat4x2_f16 = (*(tint_symbol_1)).arr2_mat4x2_f16;
Inner const struct_inner = (*(tint_symbol_1)).struct_inner;
tint_array<Inner, 4> const array_struct_inner = (*(tint_symbol_1)).array_struct_inner;
+ *(tint_symbol_2) = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(tint_ftoi(scalar_f32)) + as_type<uint>(scalar_i32)))) + as_type<uint>(int(scalar_u32))))) + as_type<uint>(int(scalar_f16))))) + as_type<uint>(tint_ftoi(vec2_f32[0]))))) + as_type<uint>(vec2_i32[0])))) + as_type<uint>(int(vec2_u32[0]))))) + as_type<uint>(int(vec2_f16[0]))))) + as_type<uint>(tint_ftoi(vec3_f32[1]))))) + as_type<uint>(vec3_i32[1])))) + as_type<uint>(int(vec3_u32[1]))))) + as_type<uint>(int(vec3_f16[1]))))) + as_type<uint>(tint_ftoi(vec4_f32[2]))))) + as_type<uint>(vec4_i32[2])))) + as_type<uint>(int(vec4_u32[2]))))) + as_type<uint>(int(vec4_f16[2]))))) + as_type<uint>(tint_ftoi(mat2x2_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat2x3_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat2x4_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat3x2_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat3x3_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat3x4_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat4x2_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat4x3_f32[0][0]))))) + as_type<uint>(tint_ftoi(mat4x4_f32[0][0]))))) + as_type<uint>(int(mat2x2_f16[0][0]))))) + as_type<uint>(int(mat2x3_f16[0][0]))))) + as_type<uint>(int(mat2x4_f16[0][0]))))) + as_type<uint>(int(mat3x2_f16[0][0]))))) + as_type<uint>(int(mat3x3_f16[0][0]))))) + as_type<uint>(int(mat3x4_f16[0][0]))))) + as_type<uint>(int(mat4x2_f16[0][0]))))) + as_type<uint>(int(mat4x3_f16[0][0]))))) + as_type<uint>(int(mat4x4_f16[0][0]))))) + as_type<uint>(tint_ftoi(arr2_vec3_f32[0][0]))))) + as_type<uint>(int(arr2_mat4x2_f16[0][0][0]))))) + as_type<uint>(struct_inner.scalar_i32)))) + as_type<uint>(array_struct_inner[0].scalar_i32)));
return;
}
diff --git a/test/tint/buffer/storage/static_index/read_f16.wgsl.expected.spvasm b/test/tint/buffer/storage/static_index/read_f16.wgsl.expected.spvasm
index fcaa0d9..31e2cbd 100644
--- a/test/tint/buffer/storage/static_index/read_f16.wgsl.expected.spvasm
+++ b/test/tint/buffer/storage/static_index/read_f16.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 199
+; Bound: 343
; Schema: 0
OpCapability Shader
OpCapability Float16
@@ -57,6 +57,11 @@
OpMemberName %Inner 2 "scalar_f16"
OpMemberName %S 37 "array_struct_inner"
OpName %sb "sb"
+ OpName %s_block "s_block"
+ OpMemberName %s_block 0 "inner"
+ OpName %s "s"
+ OpName %tint_ftoi "tint_ftoi"
+ OpName %v "v"
OpName %main "main"
OpDecorate %sb_block Block
OpMemberDecorate %sb_block 0 Offset 0
@@ -145,6 +150,10 @@
OpDecorate %sb NonWritable
OpDecorate %sb Binding 0
OpDecorate %sb DescriptorSet 0
+ OpDecorate %s_block Block
+ OpMemberDecorate %s_block 0 Offset 0
+ OpDecorate %s DescriptorSet 0
+ OpDecorate %s Binding 1
%float = OpTypeFloat 32
%int = OpTypeInt 32 1
%uint = OpTypeInt 32 0
@@ -189,8 +198,17 @@
%sb_block = OpTypeStruct %S
%_ptr_StorageBuffer_sb_block = OpTypePointer StorageBuffer %sb_block
%sb = OpVariable %_ptr_StorageBuffer_sb_block StorageBuffer
+ %s_block = OpTypeStruct %int
+%_ptr_StorageBuffer_s_block = OpTypePointer StorageBuffer %s_block
+ %s = OpVariable %_ptr_StorageBuffer_s_block StorageBuffer
+ %48 = OpTypeFunction %int %float
+%float_2_14748352e_09 = OpConstant %float 2.14748352e+09
+ %bool = OpTypeBool
+%float_n2_14748365e_09 = OpConstant %float -2.14748365e+09
+%int_n2147483648 = OpConstant %int -2147483648
+%int_2147483647 = OpConstant %int 2147483647
%void = OpTypeVoid
- %45 = OpTypeFunction %void
+ %62 = OpTypeFunction %void
%uint_0 = OpConstant %uint 0
%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
%uint_1 = OpConstant %uint 1
@@ -265,83 +283,221 @@
%_ptr_StorageBuffer_Inner = OpTypePointer StorageBuffer %Inner
%uint_37 = OpConstant %uint 37
%_ptr_StorageBuffer__arr_Inner_uint_4 = OpTypePointer StorageBuffer %_arr_Inner_uint_4
- %main = OpFunction %void None %45
- %48 = OpLabel
- %51 = OpAccessChain %_ptr_StorageBuffer_float %sb %uint_0 %uint_0
- %52 = OpLoad %float %51
- %55 = OpAccessChain %_ptr_StorageBuffer_int %sb %uint_0 %uint_1
- %56 = OpLoad %int %55
- %58 = OpAccessChain %_ptr_StorageBuffer_uint %sb %uint_0 %uint_2
- %59 = OpLoad %uint %58
- %62 = OpAccessChain %_ptr_StorageBuffer_half %sb %uint_0 %uint_3
- %63 = OpLoad %half %62
- %65 = OpAccessChain %_ptr_StorageBuffer_v2float %sb %uint_0 %uint_4
- %66 = OpLoad %v2float %65
- %69 = OpAccessChain %_ptr_StorageBuffer_v2int %sb %uint_0 %uint_5
- %70 = OpLoad %v2int %69
- %73 = OpAccessChain %_ptr_StorageBuffer_v2uint %sb %uint_0 %uint_6
- %74 = OpLoad %v2uint %73
- %77 = OpAccessChain %_ptr_StorageBuffer_v2half %sb %uint_0 %uint_7
- %78 = OpLoad %v2half %77
- %81 = OpAccessChain %_ptr_StorageBuffer_v3float %sb %uint_0 %uint_8
- %82 = OpLoad %v3float %81
- %85 = OpAccessChain %_ptr_StorageBuffer_v3int %sb %uint_0 %uint_9
- %86 = OpLoad %v3int %85
- %89 = OpAccessChain %_ptr_StorageBuffer_v3uint %sb %uint_0 %uint_10
- %90 = OpLoad %v3uint %89
- %93 = OpAccessChain %_ptr_StorageBuffer_v3half %sb %uint_0 %uint_11
- %94 = OpLoad %v3half %93
- %97 = OpAccessChain %_ptr_StorageBuffer_v4float %sb %uint_0 %uint_12
- %98 = OpLoad %v4float %97
- %101 = OpAccessChain %_ptr_StorageBuffer_v4int %sb %uint_0 %uint_13
- %102 = OpLoad %v4int %101
- %105 = OpAccessChain %_ptr_StorageBuffer_v4uint %sb %uint_0 %uint_14
- %106 = OpLoad %v4uint %105
- %109 = OpAccessChain %_ptr_StorageBuffer_v4half %sb %uint_0 %uint_15
- %110 = OpLoad %v4half %109
- %113 = OpAccessChain %_ptr_StorageBuffer_mat2v2float %sb %uint_0 %uint_16
- %114 = OpLoad %mat2v2float %113
- %117 = OpAccessChain %_ptr_StorageBuffer_mat2v3float %sb %uint_0 %uint_17
- %118 = OpLoad %mat2v3float %117
- %121 = OpAccessChain %_ptr_StorageBuffer_mat2v4float %sb %uint_0 %uint_18
- %122 = OpLoad %mat2v4float %121
- %125 = OpAccessChain %_ptr_StorageBuffer_mat3v2float %sb %uint_0 %uint_19
- %126 = OpLoad %mat3v2float %125
- %129 = OpAccessChain %_ptr_StorageBuffer_mat3v3float %sb %uint_0 %uint_20
- %130 = OpLoad %mat3v3float %129
- %133 = OpAccessChain %_ptr_StorageBuffer_mat3v4float %sb %uint_0 %uint_21
- %134 = OpLoad %mat3v4float %133
- %137 = OpAccessChain %_ptr_StorageBuffer_mat4v2float %sb %uint_0 %uint_22
- %138 = OpLoad %mat4v2float %137
- %141 = OpAccessChain %_ptr_StorageBuffer_mat4v3float %sb %uint_0 %uint_23
- %142 = OpLoad %mat4v3float %141
- %145 = OpAccessChain %_ptr_StorageBuffer_mat4v4float %sb %uint_0 %uint_24
- %146 = OpLoad %mat4v4float %145
- %149 = OpAccessChain %_ptr_StorageBuffer_mat2v2half %sb %uint_0 %uint_25
- %150 = OpLoad %mat2v2half %149
- %153 = OpAccessChain %_ptr_StorageBuffer_mat2v3half %sb %uint_0 %uint_26
- %154 = OpLoad %mat2v3half %153
- %157 = OpAccessChain %_ptr_StorageBuffer_mat2v4half %sb %uint_0 %uint_27
- %158 = OpLoad %mat2v4half %157
- %161 = OpAccessChain %_ptr_StorageBuffer_mat3v2half %sb %uint_0 %uint_28
- %162 = OpLoad %mat3v2half %161
- %165 = OpAccessChain %_ptr_StorageBuffer_mat3v3half %sb %uint_0 %uint_29
- %166 = OpLoad %mat3v3half %165
- %169 = OpAccessChain %_ptr_StorageBuffer_mat3v4half %sb %uint_0 %uint_30
- %170 = OpLoad %mat3v4half %169
- %173 = OpAccessChain %_ptr_StorageBuffer_mat4v2half %sb %uint_0 %uint_31
- %174 = OpLoad %mat4v2half %173
- %177 = OpAccessChain %_ptr_StorageBuffer_mat4v3half %sb %uint_0 %uint_32
- %178 = OpLoad %mat4v3half %177
- %181 = OpAccessChain %_ptr_StorageBuffer_mat4v4half %sb %uint_0 %uint_33
- %182 = OpLoad %mat4v4half %181
- %185 = OpAccessChain %_ptr_StorageBuffer__arr_v3float_uint_2 %sb %uint_0 %uint_34
- %186 = OpLoad %_arr_v3float_uint_2 %185
- %189 = OpAccessChain %_ptr_StorageBuffer__arr_mat4v2half_uint_2 %sb %uint_0 %uint_35
- %190 = OpLoad %_arr_mat4v2half_uint_2 %189
- %193 = OpAccessChain %_ptr_StorageBuffer_Inner %sb %uint_0 %uint_36
- %194 = OpLoad %Inner %193
- %197 = OpAccessChain %_ptr_StorageBuffer__arr_Inner_uint_4 %sb %uint_0 %uint_37
- %198 = OpLoad %_arr_Inner_uint_4 %197
+ %257 = OpConstantNull %int
+ %tint_ftoi = OpFunction %int None %48
+ %v = OpFunctionParameter %float
+ %51 = OpLabel
+ %54 = OpFOrdLessThan %bool %v %float_2_14748352e_09
+ %58 = OpFOrdLessThan %bool %v %float_n2_14748365e_09
+ %60 = OpConvertFToS %int %v
+ %56 = OpSelect %int %58 %int_n2147483648 %60
+ %52 = OpSelect %int %54 %56 %int_2147483647
+ OpReturnValue %52
+ OpFunctionEnd
+ %main = OpFunction %void None %62
+ %65 = OpLabel
+ %68 = OpAccessChain %_ptr_StorageBuffer_float %sb %uint_0 %uint_0
+ %69 = OpLoad %float %68
+ %72 = OpAccessChain %_ptr_StorageBuffer_int %sb %uint_0 %uint_1
+ %73 = OpLoad %int %72
+ %75 = OpAccessChain %_ptr_StorageBuffer_uint %sb %uint_0 %uint_2
+ %76 = OpLoad %uint %75
+ %79 = OpAccessChain %_ptr_StorageBuffer_half %sb %uint_0 %uint_3
+ %80 = OpLoad %half %79
+ %82 = OpAccessChain %_ptr_StorageBuffer_v2float %sb %uint_0 %uint_4
+ %83 = OpLoad %v2float %82
+ %86 = OpAccessChain %_ptr_StorageBuffer_v2int %sb %uint_0 %uint_5
+ %87 = OpLoad %v2int %86
+ %90 = OpAccessChain %_ptr_StorageBuffer_v2uint %sb %uint_0 %uint_6
+ %91 = OpLoad %v2uint %90
+ %94 = OpAccessChain %_ptr_StorageBuffer_v2half %sb %uint_0 %uint_7
+ %95 = OpLoad %v2half %94
+ %98 = OpAccessChain %_ptr_StorageBuffer_v3float %sb %uint_0 %uint_8
+ %99 = OpLoad %v3float %98
+ %102 = OpAccessChain %_ptr_StorageBuffer_v3int %sb %uint_0 %uint_9
+ %103 = OpLoad %v3int %102
+ %106 = OpAccessChain %_ptr_StorageBuffer_v3uint %sb %uint_0 %uint_10
+ %107 = OpLoad %v3uint %106
+ %110 = OpAccessChain %_ptr_StorageBuffer_v3half %sb %uint_0 %uint_11
+ %111 = OpLoad %v3half %110
+ %114 = OpAccessChain %_ptr_StorageBuffer_v4float %sb %uint_0 %uint_12
+ %115 = OpLoad %v4float %114
+ %118 = OpAccessChain %_ptr_StorageBuffer_v4int %sb %uint_0 %uint_13
+ %119 = OpLoad %v4int %118
+ %122 = OpAccessChain %_ptr_StorageBuffer_v4uint %sb %uint_0 %uint_14
+ %123 = OpLoad %v4uint %122
+ %126 = OpAccessChain %_ptr_StorageBuffer_v4half %sb %uint_0 %uint_15
+ %127 = OpLoad %v4half %126
+ %130 = OpAccessChain %_ptr_StorageBuffer_mat2v2float %sb %uint_0 %uint_16
+ %131 = OpLoad %mat2v2float %130
+ %134 = OpAccessChain %_ptr_StorageBuffer_mat2v3float %sb %uint_0 %uint_17
+ %135 = OpLoad %mat2v3float %134
+ %138 = OpAccessChain %_ptr_StorageBuffer_mat2v4float %sb %uint_0 %uint_18
+ %139 = OpLoad %mat2v4float %138
+ %142 = OpAccessChain %_ptr_StorageBuffer_mat3v2float %sb %uint_0 %uint_19
+ %143 = OpLoad %mat3v2float %142
+ %146 = OpAccessChain %_ptr_StorageBuffer_mat3v3float %sb %uint_0 %uint_20
+ %147 = OpLoad %mat3v3float %146
+ %150 = OpAccessChain %_ptr_StorageBuffer_mat3v4float %sb %uint_0 %uint_21
+ %151 = OpLoad %mat3v4float %150
+ %154 = OpAccessChain %_ptr_StorageBuffer_mat4v2float %sb %uint_0 %uint_22
+ %155 = OpLoad %mat4v2float %154
+ %158 = OpAccessChain %_ptr_StorageBuffer_mat4v3float %sb %uint_0 %uint_23
+ %159 = OpLoad %mat4v3float %158
+ %162 = OpAccessChain %_ptr_StorageBuffer_mat4v4float %sb %uint_0 %uint_24
+ %163 = OpLoad %mat4v4float %162
+ %166 = OpAccessChain %_ptr_StorageBuffer_mat2v2half %sb %uint_0 %uint_25
+ %167 = OpLoad %mat2v2half %166
+ %170 = OpAccessChain %_ptr_StorageBuffer_mat2v3half %sb %uint_0 %uint_26
+ %171 = OpLoad %mat2v3half %170
+ %174 = OpAccessChain %_ptr_StorageBuffer_mat2v4half %sb %uint_0 %uint_27
+ %175 = OpLoad %mat2v4half %174
+ %178 = OpAccessChain %_ptr_StorageBuffer_mat3v2half %sb %uint_0 %uint_28
+ %179 = OpLoad %mat3v2half %178
+ %182 = OpAccessChain %_ptr_StorageBuffer_mat3v3half %sb %uint_0 %uint_29
+ %183 = OpLoad %mat3v3half %182
+ %186 = OpAccessChain %_ptr_StorageBuffer_mat3v4half %sb %uint_0 %uint_30
+ %187 = OpLoad %mat3v4half %186
+ %190 = OpAccessChain %_ptr_StorageBuffer_mat4v2half %sb %uint_0 %uint_31
+ %191 = OpLoad %mat4v2half %190
+ %194 = OpAccessChain %_ptr_StorageBuffer_mat4v3half %sb %uint_0 %uint_32
+ %195 = OpLoad %mat4v3half %194
+ %198 = OpAccessChain %_ptr_StorageBuffer_mat4v4half %sb %uint_0 %uint_33
+ %199 = OpLoad %mat4v4half %198
+ %202 = OpAccessChain %_ptr_StorageBuffer__arr_v3float_uint_2 %sb %uint_0 %uint_34
+ %203 = OpLoad %_arr_v3float_uint_2 %202
+ %206 = OpAccessChain %_ptr_StorageBuffer__arr_mat4v2half_uint_2 %sb %uint_0 %uint_35
+ %207 = OpLoad %_arr_mat4v2half_uint_2 %206
+ %210 = OpAccessChain %_ptr_StorageBuffer_Inner %sb %uint_0 %uint_36
+ %211 = OpLoad %Inner %210
+ %214 = OpAccessChain %_ptr_StorageBuffer__arr_Inner_uint_4 %sb %uint_0 %uint_37
+ %215 = OpLoad %_arr_Inner_uint_4 %214
+ %216 = OpAccessChain %_ptr_StorageBuffer_int %s %uint_0
+ %217 = OpFunctionCall %int %tint_ftoi %69
+ %218 = OpIAdd %int %217 %73
+ %219 = OpBitcast %int %76
+ %220 = OpIAdd %int %218 %219
+ %221 = OpConvertFToS %int %80
+ %222 = OpIAdd %int %220 %221
+ %224 = OpCompositeExtract %float %83 0
+ %223 = OpFunctionCall %int %tint_ftoi %224
+ %225 = OpIAdd %int %222 %223
+ %226 = OpCompositeExtract %int %87 0
+ %227 = OpIAdd %int %225 %226
+ %229 = OpCompositeExtract %uint %91 0
+ %228 = OpBitcast %int %229
+ %230 = OpIAdd %int %227 %228
+ %232 = OpCompositeExtract %half %95 0
+ %231 = OpConvertFToS %int %232
+ %233 = OpIAdd %int %230 %231
+ %235 = OpCompositeExtract %float %99 1
+ %234 = OpFunctionCall %int %tint_ftoi %235
+ %236 = OpIAdd %int %233 %234
+ %237 = OpCompositeExtract %int %103 1
+ %238 = OpIAdd %int %236 %237
+ %240 = OpCompositeExtract %uint %107 1
+ %239 = OpBitcast %int %240
+ %241 = OpIAdd %int %238 %239
+ %243 = OpCompositeExtract %half %111 1
+ %242 = OpConvertFToS %int %243
+ %244 = OpIAdd %int %241 %242
+ %246 = OpCompositeExtract %float %115 2
+ %245 = OpFunctionCall %int %tint_ftoi %246
+ %247 = OpIAdd %int %244 %245
+ %248 = OpCompositeExtract %int %119 2
+ %249 = OpIAdd %int %247 %248
+ %251 = OpCompositeExtract %uint %123 2
+ %250 = OpBitcast %int %251
+ %252 = OpIAdd %int %249 %250
+ %254 = OpCompositeExtract %half %127 2
+ %253 = OpConvertFToS %int %254
+ %255 = OpIAdd %int %252 %253
+ %258 = OpCompositeExtract %v2float %131 0
+ %259 = OpCompositeExtract %float %258 0
+ %256 = OpFunctionCall %int %tint_ftoi %259
+ %260 = OpIAdd %int %255 %256
+ %262 = OpCompositeExtract %v3float %135 0
+ %263 = OpCompositeExtract %float %262 0
+ %261 = OpFunctionCall %int %tint_ftoi %263
+ %264 = OpIAdd %int %260 %261
+ %266 = OpCompositeExtract %v4float %139 0
+ %267 = OpCompositeExtract %float %266 0
+ %265 = OpFunctionCall %int %tint_ftoi %267
+ %268 = OpIAdd %int %264 %265
+ %270 = OpCompositeExtract %v2float %143 0
+ %271 = OpCompositeExtract %float %270 0
+ %269 = OpFunctionCall %int %tint_ftoi %271
+ %272 = OpIAdd %int %268 %269
+ %274 = OpCompositeExtract %v3float %147 0
+ %275 = OpCompositeExtract %float %274 0
+ %273 = OpFunctionCall %int %tint_ftoi %275
+ %276 = OpIAdd %int %272 %273
+ %278 = OpCompositeExtract %v4float %151 0
+ %279 = OpCompositeExtract %float %278 0
+ %277 = OpFunctionCall %int %tint_ftoi %279
+ %280 = OpIAdd %int %276 %277
+ %282 = OpCompositeExtract %v2float %155 0
+ %283 = OpCompositeExtract %float %282 0
+ %281 = OpFunctionCall %int %tint_ftoi %283
+ %284 = OpIAdd %int %280 %281
+ %286 = OpCompositeExtract %v3float %159 0
+ %287 = OpCompositeExtract %float %286 0
+ %285 = OpFunctionCall %int %tint_ftoi %287
+ %288 = OpIAdd %int %284 %285
+ %290 = OpCompositeExtract %v4float %163 0
+ %291 = OpCompositeExtract %float %290 0
+ %289 = OpFunctionCall %int %tint_ftoi %291
+ %292 = OpIAdd %int %288 %289
+ %294 = OpCompositeExtract %v2half %167 0
+ %295 = OpCompositeExtract %half %294 0
+ %293 = OpConvertFToS %int %295
+ %296 = OpIAdd %int %292 %293
+ %298 = OpCompositeExtract %v3half %171 0
+ %299 = OpCompositeExtract %half %298 0
+ %297 = OpConvertFToS %int %299
+ %300 = OpIAdd %int %296 %297
+ %302 = OpCompositeExtract %v4half %175 0
+ %303 = OpCompositeExtract %half %302 0
+ %301 = OpConvertFToS %int %303
+ %304 = OpIAdd %int %300 %301
+ %306 = OpCompositeExtract %v2half %179 0
+ %307 = OpCompositeExtract %half %306 0
+ %305 = OpConvertFToS %int %307
+ %308 = OpIAdd %int %304 %305
+ %310 = OpCompositeExtract %v3half %183 0
+ %311 = OpCompositeExtract %half %310 0
+ %309 = OpConvertFToS %int %311
+ %312 = OpIAdd %int %308 %309
+ %314 = OpCompositeExtract %v4half %187 0
+ %315 = OpCompositeExtract %half %314 0
+ %313 = OpConvertFToS %int %315
+ %316 = OpIAdd %int %312 %313
+ %318 = OpCompositeExtract %v2half %191 0
+ %319 = OpCompositeExtract %half %318 0
+ %317 = OpConvertFToS %int %319
+ %320 = OpIAdd %int %316 %317
+ %322 = OpCompositeExtract %v3half %195 0
+ %323 = OpCompositeExtract %half %322 0
+ %321 = OpConvertFToS %int %323
+ %324 = OpIAdd %int %320 %321
+ %326 = OpCompositeExtract %v4half %199 0
+ %327 = OpCompositeExtract %half %326 0
+ %325 = OpConvertFToS %int %327
+ %328 = OpIAdd %int %324 %325
+ %330 = OpCompositeExtract %v3float %203 0
+ %331 = OpCompositeExtract %float %330 0
+ %329 = OpFunctionCall %int %tint_ftoi %331
+ %332 = OpIAdd %int %328 %329
+ %334 = OpCompositeExtract %mat4v2half %207 0
+ %335 = OpCompositeExtract %v2half %334 0
+ %336 = OpCompositeExtract %half %335 0
+ %333 = OpConvertFToS %int %336
+ %337 = OpIAdd %int %332 %333
+ %338 = OpCompositeExtract %int %211 0
+ %339 = OpIAdd %int %337 %338
+ %340 = OpCompositeExtract %Inner %215 0
+ %341 = OpCompositeExtract %int %340 0
+ %342 = OpIAdd %int %339 %341
+ OpStore %216 %342
OpReturn
OpFunctionEnd
diff --git a/test/tint/buffer/storage/static_index/read_f16.wgsl.expected.wgsl b/test/tint/buffer/storage/static_index/read_f16.wgsl.expected.wgsl
index cb7c9e1..8a3dfef 100644
--- a/test/tint/buffer/storage/static_index/read_f16.wgsl.expected.wgsl
+++ b/test/tint/buffer/storage/static_index/read_f16.wgsl.expected.wgsl
@@ -49,6 +49,8 @@
@binding(0) @group(0) var<storage, read> sb : S;
+@group(0) @binding(1) var<storage, read_write> s : i32;
+
@compute @workgroup_size(1)
fn main() {
let scalar_f32 = sb.scalar_f32;
@@ -89,4 +91,5 @@
let arr2_mat4x2_f16 = sb.arr2_mat4x2_f16;
let struct_inner = sb.struct_inner;
let array_struct_inner = sb.array_struct_inner;
+ s = (((((((((((((((((((((((((((((((((((((i32(scalar_f32) + scalar_i32) + i32(scalar_u32)) + i32(scalar_f16)) + i32(vec2_f32.x)) + vec2_i32.x) + i32(vec2_u32.x)) + i32(vec2_f16.x)) + i32(vec3_f32.y)) + vec3_i32.y) + i32(vec3_u32.y)) + i32(vec3_f16.y)) + i32(vec4_f32.z)) + vec4_i32.z) + i32(vec4_u32.z)) + i32(vec4_f16.z)) + i32(mat2x2_f32[0].x)) + i32(mat2x3_f32[0].x)) + i32(mat2x4_f32[0].x)) + i32(mat3x2_f32[0].x)) + i32(mat3x3_f32[0].x)) + i32(mat3x4_f32[0].x)) + i32(mat4x2_f32[0].x)) + i32(mat4x3_f32[0].x)) + i32(mat4x4_f32[0].x)) + i32(mat2x2_f16[0].x)) + i32(mat2x3_f16[0].x)) + i32(mat2x4_f16[0].x)) + i32(mat3x2_f16[0].x)) + i32(mat3x3_f16[0].x)) + i32(mat3x4_f16[0].x)) + i32(mat4x2_f16[0].x)) + i32(mat4x3_f16[0].x)) + i32(mat4x4_f16[0].x)) + i32(arr2_vec3_f32[0].x)) + i32(arr2_mat4x2_f16[0][0].x)) + struct_inner.scalar_i32) + array_struct_inner[0].scalar_i32);
}
