writer/hlsl: Simplify UBO accesses for static indexing
Use the new semantic constant value information to significantly reduce the complex indexing logic emitted for UBO accesses.
This will dramatically reduce the number of `for` loops that are decayed to `while` loops.
Change-Id: I1b0adb5edde2b4ed39c6beafc2e28106b86e0edd
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57701
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index 0f5fad6..372e9de 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -624,11 +624,27 @@
ast::CallExpression* expr,
const transform::DecomposeMemoryAccess::Intrinsic* intrinsic) {
const auto& params = expr->params();
+ auto* offset_arg = builder_.Sem().Get(params[1]);
- std::string scalar_offset = UniqueIdentifier("scalar_offset");
- {
+ uint32_t scalar_offset_value = 0;
+ std::string scalar_offset_expr;
+
+ // If true, use scalar_offset_value, otherwise use scalar_offset_expr
+ bool scalar_offset_constant = false;
+
+ if (auto val = offset_arg->ConstantValue()) {
+ TINT_ASSERT(Writer, val.Type()->Is<sem::U32>());
+ scalar_offset_value = val.Elements()[0].u32;
+ scalar_offset_value /= 4; // bytes -> scalar index
+ scalar_offset_constant = true;
+ }
+
+ if (!scalar_offset_constant) {
+ // UBO offset not compile-time known.
+ // Calculate the scalar offset into a temporary.
+ scalar_offset_expr = UniqueIdentifier("scalar_offset");
auto pre = line();
- pre << "const uint " << scalar_offset << " = (";
+ pre << "const uint " << scalar_offset_expr << " = (";
if (!EmitExpression(pre, params[1])) { // offset
return false;
}
@@ -649,32 +665,37 @@
if (!EmitExpression(out, params[0])) { // buffer
return false;
}
- out << "[" << scalar_offset << " / 4][" << scalar_offset << " % 4]";
+ if (scalar_offset_constant) {
+ char swizzle[] = {'x', 'y', 'z', 'w'};
+ out << "[" << (scalar_offset_value / 4) << "]."
+ << swizzle[scalar_offset_value & 3];
+ } else {
+ out << "[" << scalar_offset_expr << " / 4][" << scalar_offset_expr
+ << " % 4]";
+ }
return true;
};
// Has a minimum alignment of 8 bytes, so is either .xy or .zw
auto load_vec2 = [&] {
- std::string ubo_load = UniqueIdentifier("ubo_load");
-
- {
- auto pre = line();
- pre << "uint4 " << ubo_load << " = ";
- if (!EmitExpression(pre, params[0])) { // buffer
+ if (scalar_offset_constant) {
+ if (!EmitExpression(out, params[0])) { // buffer
return false;
}
- pre << "[" << scalar_offset << " / 4];";
+ out << "[" << (scalar_offset_value / 4) << "]";
+ out << ((scalar_offset_value & 2) == 0 ? ".xy" : ".zw");
+ } else {
+ std::string ubo_load = UniqueIdentifier("ubo_load");
+ {
+ auto pre = line();
+ pre << "uint4 " << ubo_load << " = ";
+ if (!EmitExpression(pre, params[0])) { // buffer
+ return false;
+ }
+ pre << "[" << scalar_offset_expr << " / 4];";
+ }
+ out << "((" << scalar_offset_expr << " & 2) ? " << ubo_load
+ << ".zw : " << ubo_load << ".xy)";
}
-
- out << "((" << scalar_offset << " & 2) ? " << ubo_load
- << ".zw : " << ubo_load << ".xy)";
- return true;
- };
- // vec3 has a minimum alignment of 16 bytes, so is just a .xyz swizzle
- auto load_vec3 = [&] {
- if (!EmitExpression(out, params[0])) { // buffer
- return false;
- }
- out << "[" << scalar_offset << " / 4].xyz";
return true;
};
// vec4 has a minimum alignment of 16 bytes, easiest case
@@ -682,7 +703,19 @@
if (!EmitExpression(out, params[0])) { // buffer
return false;
}
- out << "[" << scalar_offset << " / 4]";
+ if (scalar_offset_constant) {
+ out << "[" << (scalar_offset_value / 4) << "]";
+ } else {
+ out << "[" << scalar_offset_expr << " / 4]";
+ }
+ return true;
+ };
+ // vec3 has a minimum alignment of 16 bytes, so is just a .xyz swizzle
+ auto load_vec3 = [&] {
+ if (!load_vec4()) {
+ return false;
+ }
+ out << ".xyz";
return true;
};
switch (intrinsic->type) {
diff --git a/src/writer/hlsl/generator_impl_function_test.cc b/src/writer/hlsl/generator_impl_function_test.cc
index 57de37c..59ea1e7 100644
--- a/src/writer/hlsl/generator_impl_function_test.cc
+++ b/src/writer/hlsl/generator_impl_function_test.cc
@@ -335,7 +335,7 @@
Stage(ast::PipelineStage::kFragment),
});
- GeneratorImpl& gen = Build();
+ GeneratorImpl& gen = SanitizeAndBuild();
ASSERT_TRUE(gen.Generate()) << gen.error();
EXPECT_EQ(gen.result(), R"(cbuffer cbuffer_ubo : register(b0, space1) {
@@ -343,7 +343,7 @@
};
float sub_func(float param) {
- return ubo.coord.x;
+ return asfloat(ubo[0].x);
}
void frag_main() {
diff --git a/test/buffer/storage/dynamic_index/read.wgsl b/test/buffer/storage/dynamic_index/read.wgsl
new file mode 100644
index 0000000..8295c4d
--- /dev/null
+++ b/test/buffer/storage/dynamic_index/read.wgsl
@@ -0,0 +1,31 @@
+struct Inner {
+ a : vec3<i32>;
+ b : i32;
+ c : vec3<u32>;
+ d : u32;
+ e : vec3<f32>;
+ f : f32;
+ g : mat2x3<f32>;
+ h : mat3x2<f32>;
+ i : [[stride(16)]] array<vec4<i32>, 4>;
+};
+
+[[block]]
+struct S {
+ arr : array<Inner>;
+};
+
+[[binding(0), group(0)]] var<storage, read> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main([[builtin(local_invocation_index)]] idx : u32) {
+ let a = s.arr[idx].a;
+ let b = s.arr[idx].b;
+ let c = s.arr[idx].c;
+ let d = s.arr[idx].d;
+ let e = s.arr[idx].e;
+ let f = s.arr[idx].f;
+ let g = s.arr[idx].g;
+ let h = s.arr[idx].h;
+ let i = s.arr[idx].i;
+}
diff --git a/test/buffer/storage/dynamic_index/read.wgsl.expected.hlsl b/test/buffer/storage/dynamic_index/read.wgsl.expected.hlsl
new file mode 100644
index 0000000..89fc6b9
--- /dev/null
+++ b/test/buffer/storage/dynamic_index/read.wgsl.expected.hlsl
@@ -0,0 +1,34 @@
+float2x3 tint_symbol_8(ByteAddressBuffer buffer, uint offset) {
+ return float2x3(asfloat(buffer.Load3((offset + 0u))), asfloat(buffer.Load3((offset + 16u))));
+}
+
+float3x2 tint_symbol_10(ByteAddressBuffer buffer, uint offset) {
+ return float3x2(asfloat(buffer.Load2((offset + 0u))), asfloat(buffer.Load2((offset + 8u))), asfloat(buffer.Load2((offset + 16u))));
+}
+
+typedef int4 tint_symbol_12_ret[4];
+tint_symbol_12_ret tint_symbol_12(ByteAddressBuffer buffer, uint offset) {
+ const int4 tint_symbol_13[4] = {asint(buffer.Load4((offset + 0u))), asint(buffer.Load4((offset + 16u))), asint(buffer.Load4((offset + 32u))), asint(buffer.Load4((offset + 48u)))};
+ return tint_symbol_13;
+}
+
+ByteAddressBuffer s : register(t0, space0);
+
+struct tint_symbol_1 {
+ uint idx : SV_GroupIndex;
+};
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+ const uint idx = tint_symbol.idx;
+ const int3 a = asint(s.Load3((176u * idx)));
+ const int b = asint(s.Load(((176u * idx) + 12u)));
+ const uint3 c = s.Load3(((176u * idx) + 16u));
+ const uint d = s.Load(((176u * idx) + 28u));
+ const float3 e = asfloat(s.Load3(((176u * idx) + 32u)));
+ const float f = asfloat(s.Load(((176u * idx) + 44u)));
+ const float2x3 g = tint_symbol_8(s, ((176u * idx) + 48u));
+ const float3x2 h = tint_symbol_10(s, ((176u * idx) + 80u));
+ const int4 i[4] = tint_symbol_12(s, ((176u * idx) + 112u));
+ return;
+}
diff --git a/test/buffer/storage/dynamic_index/read.wgsl.expected.msl b/test/buffer/storage/dynamic_index/read.wgsl.expected.msl
new file mode 100644
index 0000000..b9aa14f
--- /dev/null
+++ b/test/buffer/storage/dynamic_index/read.wgsl.expected.msl
@@ -0,0 +1,35 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_array_wrapper {
+ /* 0x0000 */ int4 arr[4];
+};
+struct Inner {
+ /* 0x0000 */ packed_int3 a;
+ /* 0x000c */ int b;
+ /* 0x0010 */ packed_uint3 c;
+ /* 0x001c */ uint d;
+ /* 0x0020 */ packed_float3 e;
+ /* 0x002c */ float f;
+ /* 0x0030 */ float2x3 g;
+ /* 0x0050 */ float3x2 h;
+ /* 0x0068 */ int8_t tint_pad[8];
+ /* 0x0070 */ tint_array_wrapper i;
+};
+struct S {
+ /* 0x0000 */ Inner arr[1];
+};
+
+kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]], const device S& s [[buffer(0)]]) {
+ int3 const a = s.arr[idx].a;
+ int const b = s.arr[idx].b;
+ uint3 const c = s.arr[idx].c;
+ uint const d = s.arr[idx].d;
+ float3 const e = s.arr[idx].e;
+ float const f = s.arr[idx].f;
+ float2x3 const g = s.arr[idx].g;
+ float3x2 const h = s.arr[idx].h;
+ tint_array_wrapper const i = s.arr[idx].i;
+ return;
+}
+
diff --git a/test/buffer/storage/dynamic_index/read.wgsl.expected.spvasm b/test/buffer/storage/dynamic_index/read.wgsl.expected.spvasm
new file mode 100644
index 0000000..e5df892
--- /dev/null
+++ b/test/buffer/storage/dynamic_index/read.wgsl.expected.spvasm
@@ -0,0 +1,114 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 68
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main" %tint_symbol
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %S "S"
+ OpMemberName %S 0 "arr"
+ OpName %Inner "Inner"
+ OpMemberName %Inner 0 "a"
+ OpMemberName %Inner 1 "b"
+ OpMemberName %Inner 2 "c"
+ OpMemberName %Inner 3 "d"
+ OpMemberName %Inner 4 "e"
+ OpMemberName %Inner 5 "f"
+ OpMemberName %Inner 6 "g"
+ OpMemberName %Inner 7 "h"
+ OpMemberName %Inner 8 "i"
+ OpName %s "s"
+ OpName %tint_symbol "tint_symbol"
+ OpName %main "main"
+ OpDecorate %S Block
+ OpMemberDecorate %S 0 Offset 0
+ OpMemberDecorate %Inner 0 Offset 0
+ OpMemberDecorate %Inner 1 Offset 12
+ OpMemberDecorate %Inner 2 Offset 16
+ OpMemberDecorate %Inner 3 Offset 28
+ OpMemberDecorate %Inner 4 Offset 32
+ OpMemberDecorate %Inner 5 Offset 44
+ OpMemberDecorate %Inner 6 Offset 48
+ OpMemberDecorate %Inner 6 ColMajor
+ OpMemberDecorate %Inner 6 MatrixStride 16
+ OpMemberDecorate %Inner 7 Offset 80
+ OpMemberDecorate %Inner 7 ColMajor
+ OpMemberDecorate %Inner 7 MatrixStride 8
+ OpMemberDecorate %Inner 8 Offset 112
+ OpDecorate %_arr_v4int_uint_4 ArrayStride 16
+ OpDecorate %_runtimearr_Inner ArrayStride 176
+ OpDecorate %s NonWritable
+ OpDecorate %s Binding 0
+ OpDecorate %s DescriptorSet 0
+ OpDecorate %tint_symbol BuiltIn LocalInvocationIndex
+ %int = OpTypeInt 32 1
+ %v3int = OpTypeVector %int 3
+ %uint = OpTypeInt 32 0
+ %v3uint = OpTypeVector %uint 3
+ %float = OpTypeFloat 32
+ %v3float = OpTypeVector %float 3
+%mat2v3float = OpTypeMatrix %v3float 2
+ %v2float = OpTypeVector %float 2
+%mat3v2float = OpTypeMatrix %v2float 3
+ %v4int = OpTypeVector %int 4
+ %uint_4 = OpConstant %uint 4
+%_arr_v4int_uint_4 = OpTypeArray %v4int %uint_4
+ %Inner = OpTypeStruct %v3int %int %v3uint %uint %v3float %float %mat2v3float %mat3v2float %_arr_v4int_uint_4
+%_runtimearr_Inner = OpTypeRuntimeArray %Inner
+ %S = OpTypeStruct %_runtimearr_Inner
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+ %s = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+%_ptr_Input_uint = OpTypePointer Input %uint
+%tint_symbol = OpVariable %_ptr_Input_uint Input
+ %void = OpTypeVoid
+ %20 = OpTypeFunction %void
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v3int = OpTypePointer StorageBuffer %v3int
+ %uint_1 = OpConstant %uint 1
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+ %uint_2 = OpConstant %uint 2
+%_ptr_StorageBuffer_v3uint = OpTypePointer StorageBuffer %v3uint
+ %uint_3 = OpConstant %uint 3
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+%_ptr_StorageBuffer_v3float = OpTypePointer StorageBuffer %v3float
+ %uint_5 = OpConstant %uint 5
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+ %uint_6 = OpConstant %uint 6
+%_ptr_StorageBuffer_mat2v3float = OpTypePointer StorageBuffer %mat2v3float
+ %uint_7 = OpConstant %uint 7
+%_ptr_StorageBuffer_mat3v2float = OpTypePointer StorageBuffer %mat3v2float
+ %uint_8 = OpConstant %uint 8
+%_ptr_StorageBuffer__arr_v4int_uint_4 = OpTypePointer StorageBuffer %_arr_v4int_uint_4
+ %main = OpFunction %void None %20
+ %23 = OpLabel
+ %25 = OpLoad %uint %tint_symbol
+ %27 = OpAccessChain %_ptr_StorageBuffer_v3int %s %uint_0 %25 %uint_0
+ %28 = OpLoad %v3int %27
+ %29 = OpLoad %uint %tint_symbol
+ %32 = OpAccessChain %_ptr_StorageBuffer_int %s %uint_0 %29 %uint_1
+ %33 = OpLoad %int %32
+ %34 = OpLoad %uint %tint_symbol
+ %37 = OpAccessChain %_ptr_StorageBuffer_v3uint %s %uint_0 %34 %uint_2
+ %38 = OpLoad %v3uint %37
+ %39 = OpLoad %uint %tint_symbol
+ %42 = OpAccessChain %_ptr_StorageBuffer_uint %s %uint_0 %39 %uint_3
+ %43 = OpLoad %uint %42
+ %44 = OpLoad %uint %tint_symbol
+ %46 = OpAccessChain %_ptr_StorageBuffer_v3float %s %uint_0 %44 %uint_4
+ %47 = OpLoad %v3float %46
+ %48 = OpLoad %uint %tint_symbol
+ %51 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_0 %48 %uint_5
+ %52 = OpLoad %float %51
+ %53 = OpLoad %uint %tint_symbol
+ %56 = OpAccessChain %_ptr_StorageBuffer_mat2v3float %s %uint_0 %53 %uint_6
+ %57 = OpLoad %mat2v3float %56
+ %58 = OpLoad %uint %tint_symbol
+ %61 = OpAccessChain %_ptr_StorageBuffer_mat3v2float %s %uint_0 %58 %uint_7
+ %62 = OpLoad %mat3v2float %61
+ %63 = OpLoad %uint %tint_symbol
+ %66 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %s %uint_0 %63 %uint_8
+ %67 = OpLoad %_arr_v4int_uint_4 %66
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/storage/dynamic_index/read.wgsl.expected.wgsl b/test/buffer/storage/dynamic_index/read.wgsl.expected.wgsl
new file mode 100644
index 0000000..3af1fa8
--- /dev/null
+++ b/test/buffer/storage/dynamic_index/read.wgsl.expected.wgsl
@@ -0,0 +1,31 @@
+struct Inner {
+ a : vec3<i32>;
+ b : i32;
+ c : vec3<u32>;
+ d : u32;
+ e : vec3<f32>;
+ f : f32;
+ g : mat2x3<f32>;
+ h : mat3x2<f32>;
+ i : [[stride(16)]] array<vec4<i32>, 4>;
+};
+
+[[block]]
+struct S {
+ arr : array<Inner>;
+};
+
+[[binding(0), group(0)]] var<storage, read> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main([[builtin(local_invocation_index)]] idx : u32) {
+ let a = s.arr[idx].a;
+ let b = s.arr[idx].b;
+ let c = s.arr[idx].c;
+ let d = s.arr[idx].d;
+ let e = s.arr[idx].e;
+ let f = s.arr[idx].f;
+ let g = s.arr[idx].g;
+ let h = s.arr[idx].h;
+ let i = s.arr[idx].i;
+}
diff --git a/test/buffer/storage/dynamic_index/write.wgsl b/test/buffer/storage/dynamic_index/write.wgsl
new file mode 100644
index 0000000..0e4b480
--- /dev/null
+++ b/test/buffer/storage/dynamic_index/write.wgsl
@@ -0,0 +1,31 @@
+struct Inner {
+ a : vec3<i32>;
+ b : i32;
+ c : vec3<u32>;
+ d : u32;
+ e : vec3<f32>;
+ f : f32;
+ g : mat2x3<f32>;
+ h : mat3x2<f32>;
+ i : [[stride(16)]] array<vec4<i32>, 4>;
+};
+
+[[block]]
+struct S {
+ arr : array<Inner>;
+};
+
+[[binding(0), group(0)]] var<storage, read_write> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main([[builtin(local_invocation_index)]] idx : u32) {
+ s.arr[idx].a = vec3<i32>();
+ s.arr[idx].b = i32();
+ s.arr[idx].c = vec3<u32>();
+ s.arr[idx].d = u32();
+ s.arr[idx].e = vec3<f32>();
+ s.arr[idx].f = f32();
+ s.arr[idx].g = mat2x3<f32>();
+ s.arr[idx].h = mat3x2<f32>();
+ s.arr[idx].i = [[stride(16)]] array<vec4<i32>, 4>();
+}
diff --git a/test/buffer/storage/dynamic_index/write.wgsl.expected.hlsl b/test/buffer/storage/dynamic_index/write.wgsl.expected.hlsl
new file mode 100644
index 0000000..9d03558
--- /dev/null
+++ b/test/buffer/storage/dynamic_index/write.wgsl.expected.hlsl
@@ -0,0 +1,39 @@
+void tint_symbol_8(RWByteAddressBuffer buffer, uint offset, float2x3 value) {
+ buffer.Store3((offset + 0u), asuint(value[0u]));
+ buffer.Store3((offset + 16u), asuint(value[1u]));
+}
+
+void tint_symbol_10(RWByteAddressBuffer buffer, uint offset, float3x2 value) {
+ buffer.Store2((offset + 0u), asuint(value[0u]));
+ buffer.Store2((offset + 8u), asuint(value[1u]));
+ buffer.Store2((offset + 16u), asuint(value[2u]));
+}
+
+void tint_symbol_12(RWByteAddressBuffer buffer, uint offset, int4 value[4]) {
+ buffer.Store4((offset + 0u), asuint(value[0u]));
+ buffer.Store4((offset + 16u), asuint(value[1u]));
+ buffer.Store4((offset + 32u), asuint(value[2u]));
+ buffer.Store4((offset + 48u), asuint(value[3u]));
+}
+
+RWByteAddressBuffer s : register(u0, space0);
+
+struct tint_symbol_1 {
+ uint idx : SV_GroupIndex;
+};
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+ const uint idx = tint_symbol.idx;
+ s.Store3((176u * idx), asuint(int3(0, 0, 0)));
+ s.Store(((176u * idx) + 12u), asuint(0));
+ s.Store3(((176u * idx) + 16u), asuint(uint3(0u, 0u, 0u)));
+ s.Store(((176u * idx) + 28u), asuint(0u));
+ s.Store3(((176u * idx) + 32u), asuint(float3(0.0f, 0.0f, 0.0f)));
+ s.Store(((176u * idx) + 44u), asuint(0.0f));
+ tint_symbol_8(s, ((176u * idx) + 48u), float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f));
+ tint_symbol_10(s, ((176u * idx) + 80u), float3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f));
+ const int4 tint_symbol_13[4] = (int4[4])0;
+ tint_symbol_12(s, ((176u * idx) + 112u), tint_symbol_13);
+ return;
+}
diff --git a/test/buffer/storage/dynamic_index/write.wgsl.expected.msl b/test/buffer/storage/dynamic_index/write.wgsl.expected.msl
new file mode 100644
index 0000000..be0d4db
--- /dev/null
+++ b/test/buffer/storage/dynamic_index/write.wgsl.expected.msl
@@ -0,0 +1,36 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_array_wrapper {
+ /* 0x0000 */ int4 arr[4];
+};
+struct Inner {
+ /* 0x0000 */ packed_int3 a;
+ /* 0x000c */ int b;
+ /* 0x0010 */ packed_uint3 c;
+ /* 0x001c */ uint d;
+ /* 0x0020 */ packed_float3 e;
+ /* 0x002c */ float f;
+ /* 0x0030 */ float2x3 g;
+ /* 0x0050 */ float3x2 h;
+ /* 0x0068 */ int8_t tint_pad[8];
+ /* 0x0070 */ tint_array_wrapper i;
+};
+struct S {
+ /* 0x0000 */ Inner arr[1];
+};
+
+kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]], device S& s [[buffer(0)]]) {
+ s.arr[idx].a = int3();
+ s.arr[idx].b = int();
+ s.arr[idx].c = uint3();
+ s.arr[idx].d = uint();
+ s.arr[idx].e = float3();
+ s.arr[idx].f = float();
+ s.arr[idx].g = float2x3();
+ s.arr[idx].h = float3x2();
+ tint_array_wrapper const tint_symbol_2 = {.arr={}};
+ s.arr[idx].i = tint_symbol_2;
+ return;
+}
+
diff --git a/test/buffer/storage/dynamic_index/write.wgsl.expected.spvasm b/test/buffer/storage/dynamic_index/write.wgsl.expected.spvasm
new file mode 100644
index 0000000..ca9c2b3
--- /dev/null
+++ b/test/buffer/storage/dynamic_index/write.wgsl.expected.spvasm
@@ -0,0 +1,122 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 68
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main" %tint_symbol
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %S "S"
+ OpMemberName %S 0 "arr"
+ OpName %Inner "Inner"
+ OpMemberName %Inner 0 "a"
+ OpMemberName %Inner 1 "b"
+ OpMemberName %Inner 2 "c"
+ OpMemberName %Inner 3 "d"
+ OpMemberName %Inner 4 "e"
+ OpMemberName %Inner 5 "f"
+ OpMemberName %Inner 6 "g"
+ OpMemberName %Inner 7 "h"
+ OpMemberName %Inner 8 "i"
+ OpName %s "s"
+ OpName %tint_symbol "tint_symbol"
+ OpName %main "main"
+ OpDecorate %S Block
+ OpMemberDecorate %S 0 Offset 0
+ OpMemberDecorate %Inner 0 Offset 0
+ OpMemberDecorate %Inner 1 Offset 12
+ OpMemberDecorate %Inner 2 Offset 16
+ OpMemberDecorate %Inner 3 Offset 28
+ OpMemberDecorate %Inner 4 Offset 32
+ OpMemberDecorate %Inner 5 Offset 44
+ OpMemberDecorate %Inner 6 Offset 48
+ OpMemberDecorate %Inner 6 ColMajor
+ OpMemberDecorate %Inner 6 MatrixStride 16
+ OpMemberDecorate %Inner 7 Offset 80
+ OpMemberDecorate %Inner 7 ColMajor
+ OpMemberDecorate %Inner 7 MatrixStride 8
+ OpMemberDecorate %Inner 8 Offset 112
+ OpDecorate %_arr_v4int_uint_4 ArrayStride 16
+ OpDecorate %_runtimearr_Inner ArrayStride 176
+ OpDecorate %s Binding 0
+ OpDecorate %s DescriptorSet 0
+ OpDecorate %tint_symbol BuiltIn LocalInvocationIndex
+ %int = OpTypeInt 32 1
+ %v3int = OpTypeVector %int 3
+ %uint = OpTypeInt 32 0
+ %v3uint = OpTypeVector %uint 3
+ %float = OpTypeFloat 32
+ %v3float = OpTypeVector %float 3
+%mat2v3float = OpTypeMatrix %v3float 2
+ %v2float = OpTypeVector %float 2
+%mat3v2float = OpTypeMatrix %v2float 3
+ %v4int = OpTypeVector %int 4
+ %uint_4 = OpConstant %uint 4
+%_arr_v4int_uint_4 = OpTypeArray %v4int %uint_4
+ %Inner = OpTypeStruct %v3int %int %v3uint %uint %v3float %float %mat2v3float %mat3v2float %_arr_v4int_uint_4
+%_runtimearr_Inner = OpTypeRuntimeArray %Inner
+ %S = OpTypeStruct %_runtimearr_Inner
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+ %s = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+%_ptr_Input_uint = OpTypePointer Input %uint
+%tint_symbol = OpVariable %_ptr_Input_uint Input
+ %void = OpTypeVoid
+ %20 = OpTypeFunction %void
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v3int = OpTypePointer StorageBuffer %v3int
+ %28 = OpConstantNull %v3int
+ %uint_1 = OpConstant %uint 1
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+ %33 = OpConstantNull %int
+ %uint_2 = OpConstant %uint 2
+%_ptr_StorageBuffer_v3uint = OpTypePointer StorageBuffer %v3uint
+ %38 = OpConstantNull %v3uint
+ %uint_3 = OpConstant %uint 3
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+ %43 = OpConstantNull %uint
+%_ptr_StorageBuffer_v3float = OpTypePointer StorageBuffer %v3float
+ %47 = OpConstantNull %v3float
+ %uint_5 = OpConstant %uint 5
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+ %52 = OpConstantNull %float
+ %uint_6 = OpConstant %uint 6
+%_ptr_StorageBuffer_mat2v3float = OpTypePointer StorageBuffer %mat2v3float
+ %57 = OpConstantNull %mat2v3float
+ %uint_7 = OpConstant %uint 7
+%_ptr_StorageBuffer_mat3v2float = OpTypePointer StorageBuffer %mat3v2float
+ %62 = OpConstantNull %mat3v2float
+ %uint_8 = OpConstant %uint 8
+%_ptr_StorageBuffer__arr_v4int_uint_4 = OpTypePointer StorageBuffer %_arr_v4int_uint_4
+ %67 = OpConstantNull %_arr_v4int_uint_4
+ %main = OpFunction %void None %20
+ %23 = OpLabel
+ %25 = OpLoad %uint %tint_symbol
+ %27 = OpAccessChain %_ptr_StorageBuffer_v3int %s %uint_0 %25 %uint_0
+ OpStore %27 %28
+ %29 = OpLoad %uint %tint_symbol
+ %32 = OpAccessChain %_ptr_StorageBuffer_int %s %uint_0 %29 %uint_1
+ OpStore %32 %33
+ %34 = OpLoad %uint %tint_symbol
+ %37 = OpAccessChain %_ptr_StorageBuffer_v3uint %s %uint_0 %34 %uint_2
+ OpStore %37 %38
+ %39 = OpLoad %uint %tint_symbol
+ %42 = OpAccessChain %_ptr_StorageBuffer_uint %s %uint_0 %39 %uint_3
+ OpStore %42 %43
+ %44 = OpLoad %uint %tint_symbol
+ %46 = OpAccessChain %_ptr_StorageBuffer_v3float %s %uint_0 %44 %uint_4
+ OpStore %46 %47
+ %48 = OpLoad %uint %tint_symbol
+ %51 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_0 %48 %uint_5
+ OpStore %51 %52
+ %53 = OpLoad %uint %tint_symbol
+ %56 = OpAccessChain %_ptr_StorageBuffer_mat2v3float %s %uint_0 %53 %uint_6
+ OpStore %56 %57
+ %58 = OpLoad %uint %tint_symbol
+ %61 = OpAccessChain %_ptr_StorageBuffer_mat3v2float %s %uint_0 %58 %uint_7
+ OpStore %61 %62
+ %63 = OpLoad %uint %tint_symbol
+ %66 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %s %uint_0 %63 %uint_8
+ OpStore %66 %67
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/storage/dynamic_index/write.wgsl.expected.wgsl b/test/buffer/storage/dynamic_index/write.wgsl.expected.wgsl
new file mode 100644
index 0000000..d6aea47
--- /dev/null
+++ b/test/buffer/storage/dynamic_index/write.wgsl.expected.wgsl
@@ -0,0 +1,31 @@
+struct Inner {
+ a : vec3<i32>;
+ b : i32;
+ c : vec3<u32>;
+ d : u32;
+ e : vec3<f32>;
+ f : f32;
+ g : mat2x3<f32>;
+ h : mat3x2<f32>;
+ i : [[stride(16)]] array<vec4<i32>, 4>;
+};
+
+[[block]]
+struct S {
+ arr : array<Inner>;
+};
+
+[[binding(0), group(0)]] var<storage, read_write> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main([[builtin(local_invocation_index)]] idx : u32) {
+ s.arr[idx].a = vec3<i32>();
+ s.arr[idx].b = i32();
+ s.arr[idx].c = vec3<u32>();
+ s.arr[idx].d = u32();
+ s.arr[idx].e = vec3<f32>();
+ s.arr[idx].f = f32();
+ s.arr[idx].g = mat2x3<f32>();
+ s.arr[idx].h = mat3x2<f32>();
+ s.arr[idx].i = [[stride(16)]] array<vec4<i32>, 4>();
+}
diff --git a/test/buffer/storage/static_index/read.wgsl b/test/buffer/storage/static_index/read.wgsl
new file mode 100644
index 0000000..096492d
--- /dev/null
+++ b/test/buffer/storage/static_index/read.wgsl
@@ -0,0 +1,33 @@
+struct Inner {
+ x : i32;
+};
+
+[[block]]
+struct S {
+ a : vec3<i32>;
+ b : i32;
+ c : vec3<u32>;
+ d : u32;
+ e : vec3<f32>;
+ f : f32;
+ g : mat2x3<f32>;
+ h : mat3x2<f32>;
+ i : Inner;
+ j : [[stride(16)]] array<Inner, 4>;
+};
+
+[[binding(0), group(0)]] var<storage, read> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let a = s.a;
+ let b = s.b;
+ let c = s.c;
+ let d = s.d;
+ let e = s.e;
+ let f = s.f;
+ let g = s.g;
+ let h = s.h;
+ let i = s.i;
+ let j = s.j;
+}
diff --git a/test/buffer/storage/static_index/read.wgsl.expected.hlsl b/test/buffer/storage/static_index/read.wgsl.expected.hlsl
new file mode 100644
index 0000000..95ed272
--- /dev/null
+++ b/test/buffer/storage/static_index/read.wgsl.expected.hlsl
@@ -0,0 +1,42 @@
+struct Inner {
+ int x;
+};
+struct tint_padded_array_element {
+ Inner el;
+};
+
+float2x3 tint_symbol_6(ByteAddressBuffer buffer, uint offset) {
+ return float2x3(asfloat(buffer.Load3((offset + 0u))), asfloat(buffer.Load3((offset + 16u))));
+}
+
+float3x2 tint_symbol_8(ByteAddressBuffer buffer, uint offset) {
+ return float3x2(asfloat(buffer.Load2((offset + 0u))), asfloat(buffer.Load2((offset + 8u))), asfloat(buffer.Load2((offset + 16u))));
+}
+
+Inner tint_symbol_9(ByteAddressBuffer buffer, uint offset) {
+ const Inner tint_symbol_11 = {asint(buffer.Load((offset + 0u)))};
+ return tint_symbol_11;
+}
+
+typedef tint_padded_array_element tint_symbol_10_ret[4];
+tint_symbol_10_ret tint_symbol_10(ByteAddressBuffer buffer, uint offset) {
+ const tint_padded_array_element tint_symbol_12[4] = {{tint_symbol_9(buffer, (offset + 0u))}, {tint_symbol_9(buffer, (offset + 16u))}, {tint_symbol_9(buffer, (offset + 32u))}, {tint_symbol_9(buffer, (offset + 48u))}};
+ return tint_symbol_12;
+}
+
+ByteAddressBuffer s : register(t0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+ const int3 a = asint(s.Load3(0u));
+ const int b = asint(s.Load(12u));
+ const uint3 c = s.Load3(16u);
+ const uint d = s.Load(28u);
+ const float3 e = asfloat(s.Load3(32u));
+ const float f = asfloat(s.Load(44u));
+ const float2x3 g = tint_symbol_6(s, 48u);
+ const float3x2 h = tint_symbol_8(s, 80u);
+ const Inner i = tint_symbol_9(s, 104u);
+ const tint_padded_array_element j[4] = tint_symbol_10(s, 108u);
+ return;
+}
diff --git a/test/buffer/storage/static_index/read.wgsl.expected.msl b/test/buffer/storage/static_index/read.wgsl.expected.msl
new file mode 100644
index 0000000..316ed31
--- /dev/null
+++ b/test/buffer/storage/static_index/read.wgsl.expected.msl
@@ -0,0 +1,41 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Inner {
+ /* 0x0000 */ int x;
+};
+struct tint_padded_array_element {
+ /* 0x0000 */ Inner el;
+ /* 0x0004 */ int8_t tint_pad[12];
+};
+struct tint_array_wrapper {
+ /* 0x0000 */ tint_padded_array_element arr[4];
+};
+struct S {
+ /* 0x0000 */ packed_int3 a;
+ /* 0x000c */ int b;
+ /* 0x0010 */ packed_uint3 c;
+ /* 0x001c */ uint d;
+ /* 0x0020 */ packed_float3 e;
+ /* 0x002c */ float f;
+ /* 0x0030 */ float2x3 g;
+ /* 0x0050 */ float3x2 h;
+ /* 0x0068 */ Inner i;
+ /* 0x006c */ tint_array_wrapper j;
+ /* 0x00ac */ int8_t tint_pad_1[4];
+};
+
+kernel void tint_symbol(const device S& s [[buffer(0)]]) {
+ int3 const a = s.a;
+ int const b = s.b;
+ uint3 const c = s.c;
+ uint const d = s.d;
+ float3 const e = s.e;
+ float const f = s.f;
+ float2x3 const g = s.g;
+ float3x2 const h = s.h;
+ Inner const i = s.i;
+ tint_array_wrapper const j = s.j;
+ return;
+}
+
diff --git a/test/buffer/storage/static_index/read.wgsl.expected.spvasm b/test/buffer/storage/static_index/read.wgsl.expected.spvasm
new file mode 100644
index 0000000..98bf962
--- /dev/null
+++ b/test/buffer/storage/static_index/read.wgsl.expected.spvasm
@@ -0,0 +1,104 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 59
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %S "S"
+ OpMemberName %S 0 "a"
+ OpMemberName %S 1 "b"
+ OpMemberName %S 2 "c"
+ OpMemberName %S 3 "d"
+ OpMemberName %S 4 "e"
+ OpMemberName %S 5 "f"
+ OpMemberName %S 6 "g"
+ OpMemberName %S 7 "h"
+ OpMemberName %S 8 "i"
+ OpName %Inner "Inner"
+ OpMemberName %Inner 0 "x"
+ OpMemberName %S 9 "j"
+ OpName %s "s"
+ OpName %main "main"
+ OpDecorate %S Block
+ OpMemberDecorate %S 0 Offset 0
+ OpMemberDecorate %S 1 Offset 12
+ OpMemberDecorate %S 2 Offset 16
+ OpMemberDecorate %S 3 Offset 28
+ OpMemberDecorate %S 4 Offset 32
+ OpMemberDecorate %S 5 Offset 44
+ OpMemberDecorate %S 6 Offset 48
+ OpMemberDecorate %S 6 ColMajor
+ OpMemberDecorate %S 6 MatrixStride 16
+ OpMemberDecorate %S 7 Offset 80
+ OpMemberDecorate %S 7 ColMajor
+ OpMemberDecorate %S 7 MatrixStride 8
+ OpMemberDecorate %S 8 Offset 104
+ OpMemberDecorate %Inner 0 Offset 0
+ OpMemberDecorate %S 9 Offset 108
+ OpDecorate %_arr_Inner_uint_4 ArrayStride 16
+ OpDecorate %s NonWritable
+ OpDecorate %s Binding 0
+ OpDecorate %s DescriptorSet 0
+ %int = OpTypeInt 32 1
+ %v3int = OpTypeVector %int 3
+ %uint = OpTypeInt 32 0
+ %v3uint = OpTypeVector %uint 3
+ %float = OpTypeFloat 32
+ %v3float = OpTypeVector %float 3
+%mat2v3float = OpTypeMatrix %v3float 2
+ %v2float = OpTypeVector %float 2
+%mat3v2float = OpTypeMatrix %v2float 3
+ %Inner = OpTypeStruct %int
+ %uint_4 = OpConstant %uint 4
+%_arr_Inner_uint_4 = OpTypeArray %Inner %uint_4
+ %S = OpTypeStruct %v3int %int %v3uint %uint %v3float %float %mat2v3float %mat3v2float %Inner %_arr_Inner_uint_4
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+ %s = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+ %void = OpTypeVoid
+ %16 = OpTypeFunction %void
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v3int = OpTypePointer StorageBuffer %v3int
+ %uint_1 = OpConstant %uint 1
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+ %uint_2 = OpConstant %uint 2
+%_ptr_StorageBuffer_v3uint = OpTypePointer StorageBuffer %v3uint
+ %uint_3 = OpConstant %uint 3
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+%_ptr_StorageBuffer_v3float = OpTypePointer StorageBuffer %v3float
+ %uint_5 = OpConstant %uint 5
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+ %uint_6 = OpConstant %uint 6
+%_ptr_StorageBuffer_mat2v3float = OpTypePointer StorageBuffer %mat2v3float
+ %uint_7 = OpConstant %uint 7
+%_ptr_StorageBuffer_mat3v2float = OpTypePointer StorageBuffer %mat3v2float
+ %uint_8 = OpConstant %uint 8
+%_ptr_StorageBuffer_Inner = OpTypePointer StorageBuffer %Inner
+ %uint_9 = OpConstant %uint 9
+%_ptr_StorageBuffer__arr_Inner_uint_4 = OpTypePointer StorageBuffer %_arr_Inner_uint_4
+ %main = OpFunction %void None %16
+ %19 = OpLabel
+ %22 = OpAccessChain %_ptr_StorageBuffer_v3int %s %uint_0
+ %23 = OpLoad %v3int %22
+ %26 = OpAccessChain %_ptr_StorageBuffer_int %s %uint_1
+ %27 = OpLoad %int %26
+ %30 = OpAccessChain %_ptr_StorageBuffer_v3uint %s %uint_2
+ %31 = OpLoad %v3uint %30
+ %34 = OpAccessChain %_ptr_StorageBuffer_uint %s %uint_3
+ %35 = OpLoad %uint %34
+ %37 = OpAccessChain %_ptr_StorageBuffer_v3float %s %uint_4
+ %38 = OpLoad %v3float %37
+ %41 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_5
+ %42 = OpLoad %float %41
+ %45 = OpAccessChain %_ptr_StorageBuffer_mat2v3float %s %uint_6
+ %46 = OpLoad %mat2v3float %45
+ %49 = OpAccessChain %_ptr_StorageBuffer_mat3v2float %s %uint_7
+ %50 = OpLoad %mat3v2float %49
+ %53 = OpAccessChain %_ptr_StorageBuffer_Inner %s %uint_8
+ %54 = OpLoad %Inner %53
+ %57 = OpAccessChain %_ptr_StorageBuffer__arr_Inner_uint_4 %s %uint_9
+ %58 = OpLoad %_arr_Inner_uint_4 %57
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/storage/static_index/read.wgsl.expected.wgsl b/test/buffer/storage/static_index/read.wgsl.expected.wgsl
new file mode 100644
index 0000000..5dcc611
--- /dev/null
+++ b/test/buffer/storage/static_index/read.wgsl.expected.wgsl
@@ -0,0 +1,33 @@
+struct Inner {
+ x : i32;
+};
+
+[[block]]
+struct S {
+ a : vec3<i32>;
+ b : i32;
+ c : vec3<u32>;
+ d : u32;
+ e : vec3<f32>;
+ f : f32;
+ g : mat2x3<f32>;
+ h : mat3x2<f32>;
+ i : Inner;
+ j : [[stride(16)]] array<Inner, 4>;
+};
+
+[[binding(0), group(0)]] var<storage, read> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let a = s.a;
+ let b = s.b;
+ let c = s.c;
+ let d = s.d;
+ let e = s.e;
+ let f = s.f;
+ let g = s.g;
+ let h = s.h;
+ let i = s.i;
+ let j = s.j;
+}
diff --git a/test/buffer/storage/static_index/write.wgsl b/test/buffer/storage/static_index/write.wgsl
new file mode 100644
index 0000000..6b6affb
--- /dev/null
+++ b/test/buffer/storage/static_index/write.wgsl
@@ -0,0 +1,33 @@
+struct Inner {
+ x : i32;
+};
+
+[[block]]
+struct S {
+ a : vec3<i32>;
+ b : i32;
+ c : vec3<u32>;
+ d : u32;
+ e : vec3<f32>;
+ f : f32;
+ g : mat2x3<f32>;
+ h : mat3x2<f32>;
+ i : Inner;
+ j : [[stride(16)]] array<Inner, 4>;
+};
+
+[[binding(0), group(0)]] var<storage, write> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ s.a = vec3<i32>();
+ s.b = i32();
+ s.c = vec3<u32>();
+ s.d = u32();
+ s.e = vec3<f32>();
+ s.f = f32();
+ s.g = mat2x3<f32>();
+ s.h = mat3x2<f32>();
+ s.i = Inner();
+ s.j = [[stride(16)]] array<Inner, 4>();
+}
diff --git a/test/buffer/storage/static_index/write.wgsl.expected.hlsl b/test/buffer/storage/static_index/write.wgsl.expected.hlsl
new file mode 100644
index 0000000..52c36c6
--- /dev/null
+++ b/test/buffer/storage/static_index/write.wgsl.expected.hlsl
@@ -0,0 +1,47 @@
+struct Inner {
+ int x;
+};
+struct tint_padded_array_element {
+ Inner el;
+};
+
+void tint_symbol_6(RWByteAddressBuffer buffer, uint offset, float2x3 value) {
+ buffer.Store3((offset + 0u), asuint(value[0u]));
+ buffer.Store3((offset + 16u), asuint(value[1u]));
+}
+
+void tint_symbol_8(RWByteAddressBuffer buffer, uint offset, float3x2 value) {
+ buffer.Store2((offset + 0u), asuint(value[0u]));
+ buffer.Store2((offset + 8u), asuint(value[1u]));
+ buffer.Store2((offset + 16u), asuint(value[2u]));
+}
+
+void tint_symbol_9(RWByteAddressBuffer buffer, uint offset, Inner value) {
+ buffer.Store((offset + 0u), asuint(value.x));
+}
+
+void tint_symbol_10(RWByteAddressBuffer buffer, uint offset, tint_padded_array_element value[4]) {
+ tint_symbol_9(buffer, (offset + 0u), value[0u].el);
+ tint_symbol_9(buffer, (offset + 16u), value[1u].el);
+ tint_symbol_9(buffer, (offset + 32u), value[2u].el);
+ tint_symbol_9(buffer, (offset + 48u), value[3u].el);
+}
+
+RWByteAddressBuffer s : register(u0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+ s.Store3(0u, asuint(int3(0, 0, 0)));
+ s.Store(12u, asuint(0));
+ s.Store3(16u, asuint(uint3(0u, 0u, 0u)));
+ s.Store(28u, asuint(0u));
+ s.Store3(32u, asuint(float3(0.0f, 0.0f, 0.0f)));
+ s.Store(44u, asuint(0.0f));
+ tint_symbol_6(s, 48u, float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f));
+ tint_symbol_8(s, 80u, float3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f));
+ const Inner tint_symbol_11 = (Inner)0;
+ tint_symbol_9(s, 104u, tint_symbol_11);
+ const tint_padded_array_element tint_symbol_12[4] = (tint_padded_array_element[4])0;
+ tint_symbol_10(s, 108u, tint_symbol_12);
+ return;
+}
diff --git a/test/buffer/storage/static_index/write.wgsl.expected.msl b/test/buffer/storage/static_index/write.wgsl.expected.msl
new file mode 100644
index 0000000..eab1f4f
--- /dev/null
+++ b/test/buffer/storage/static_index/write.wgsl.expected.msl
@@ -0,0 +1,43 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Inner {
+ /* 0x0000 */ int x;
+};
+struct tint_padded_array_element {
+ /* 0x0000 */ Inner el;
+ /* 0x0004 */ int8_t tint_pad[12];
+};
+struct tint_array_wrapper {
+ /* 0x0000 */ tint_padded_array_element arr[4];
+};
+struct S {
+ /* 0x0000 */ packed_int3 a;
+ /* 0x000c */ int b;
+ /* 0x0010 */ packed_uint3 c;
+ /* 0x001c */ uint d;
+ /* 0x0020 */ packed_float3 e;
+ /* 0x002c */ float f;
+ /* 0x0030 */ float2x3 g;
+ /* 0x0050 */ float3x2 h;
+ /* 0x0068 */ Inner i;
+ /* 0x006c */ tint_array_wrapper j;
+ /* 0x00ac */ int8_t tint_pad_1[4];
+};
+
+kernel void tint_symbol(device S& s [[buffer(0)]]) {
+ s.a = int3();
+ s.b = int();
+ s.c = uint3();
+ s.d = uint();
+ s.e = float3();
+ s.f = float();
+ s.g = float2x3();
+ s.h = float3x2();
+ Inner const tint_symbol_1 = {};
+ s.i = tint_symbol_1;
+ tint_array_wrapper const tint_symbol_2 = {.arr={}};
+ s.j = tint_symbol_2;
+ return;
+}
+
diff --git a/test/buffer/storage/static_index/write.wgsl.expected.spvasm b/test/buffer/storage/static_index/write.wgsl.expected.spvasm
new file mode 100644
index 0000000..faa44c5
--- /dev/null
+++ b/test/buffer/storage/static_index/write.wgsl.expected.spvasm
@@ -0,0 +1,114 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 59
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %S "S"
+ OpMemberName %S 0 "a"
+ OpMemberName %S 1 "b"
+ OpMemberName %S 2 "c"
+ OpMemberName %S 3 "d"
+ OpMemberName %S 4 "e"
+ OpMemberName %S 5 "f"
+ OpMemberName %S 6 "g"
+ OpMemberName %S 7 "h"
+ OpMemberName %S 8 "i"
+ OpName %Inner "Inner"
+ OpMemberName %Inner 0 "x"
+ OpMemberName %S 9 "j"
+ OpName %s "s"
+ OpName %main "main"
+ OpDecorate %S Block
+ OpMemberDecorate %S 0 Offset 0
+ OpMemberDecorate %S 1 Offset 12
+ OpMemberDecorate %S 2 Offset 16
+ OpMemberDecorate %S 3 Offset 28
+ OpMemberDecorate %S 4 Offset 32
+ OpMemberDecorate %S 5 Offset 44
+ OpMemberDecorate %S 6 Offset 48
+ OpMemberDecorate %S 6 ColMajor
+ OpMemberDecorate %S 6 MatrixStride 16
+ OpMemberDecorate %S 7 Offset 80
+ OpMemberDecorate %S 7 ColMajor
+ OpMemberDecorate %S 7 MatrixStride 8
+ OpMemberDecorate %S 8 Offset 104
+ OpMemberDecorate %Inner 0 Offset 0
+ OpMemberDecorate %S 9 Offset 108
+ OpDecorate %_arr_Inner_uint_4 ArrayStride 16
+ OpDecorate %s NonReadable
+ OpDecorate %s Binding 0
+ OpDecorate %s DescriptorSet 0
+ %int = OpTypeInt 32 1
+ %v3int = OpTypeVector %int 3
+ %uint = OpTypeInt 32 0
+ %v3uint = OpTypeVector %uint 3
+ %float = OpTypeFloat 32
+ %v3float = OpTypeVector %float 3
+%mat2v3float = OpTypeMatrix %v3float 2
+ %v2float = OpTypeVector %float 2
+%mat3v2float = OpTypeMatrix %v2float 3
+ %Inner = OpTypeStruct %int
+ %uint_4 = OpConstant %uint 4
+%_arr_Inner_uint_4 = OpTypeArray %Inner %uint_4
+ %S = OpTypeStruct %v3int %int %v3uint %uint %v3float %float %mat2v3float %mat3v2float %Inner %_arr_Inner_uint_4
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+ %s = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+ %void = OpTypeVoid
+ %16 = OpTypeFunction %void
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v3int = OpTypePointer StorageBuffer %v3int
+ %23 = OpConstantNull %v3int
+ %uint_1 = OpConstant %uint 1
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+ %27 = OpConstantNull %int
+ %uint_2 = OpConstant %uint 2
+%_ptr_StorageBuffer_v3uint = OpTypePointer StorageBuffer %v3uint
+ %31 = OpConstantNull %v3uint
+ %uint_3 = OpConstant %uint 3
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+ %35 = OpConstantNull %uint
+%_ptr_StorageBuffer_v3float = OpTypePointer StorageBuffer %v3float
+ %38 = OpConstantNull %v3float
+ %uint_5 = OpConstant %uint 5
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+ %42 = OpConstantNull %float
+ %uint_6 = OpConstant %uint 6
+%_ptr_StorageBuffer_mat2v3float = OpTypePointer StorageBuffer %mat2v3float
+ %46 = OpConstantNull %mat2v3float
+ %uint_7 = OpConstant %uint 7
+%_ptr_StorageBuffer_mat3v2float = OpTypePointer StorageBuffer %mat3v2float
+ %50 = OpConstantNull %mat3v2float
+ %uint_8 = OpConstant %uint 8
+%_ptr_StorageBuffer_Inner = OpTypePointer StorageBuffer %Inner
+ %54 = OpConstantNull %Inner
+ %uint_9 = OpConstant %uint 9
+%_ptr_StorageBuffer__arr_Inner_uint_4 = OpTypePointer StorageBuffer %_arr_Inner_uint_4
+ %58 = OpConstantNull %_arr_Inner_uint_4
+ %main = OpFunction %void None %16
+ %19 = OpLabel
+ %22 = OpAccessChain %_ptr_StorageBuffer_v3int %s %uint_0
+ OpStore %22 %23
+ %26 = OpAccessChain %_ptr_StorageBuffer_int %s %uint_1
+ OpStore %26 %27
+ %30 = OpAccessChain %_ptr_StorageBuffer_v3uint %s %uint_2
+ OpStore %30 %31
+ %34 = OpAccessChain %_ptr_StorageBuffer_uint %s %uint_3
+ OpStore %34 %35
+ %37 = OpAccessChain %_ptr_StorageBuffer_v3float %s %uint_4
+ OpStore %37 %38
+ %41 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_5
+ OpStore %41 %42
+ %45 = OpAccessChain %_ptr_StorageBuffer_mat2v3float %s %uint_6
+ OpStore %45 %46
+ %49 = OpAccessChain %_ptr_StorageBuffer_mat3v2float %s %uint_7
+ OpStore %49 %50
+ %53 = OpAccessChain %_ptr_StorageBuffer_Inner %s %uint_8
+ OpStore %53 %54
+ %57 = OpAccessChain %_ptr_StorageBuffer__arr_Inner_uint_4 %s %uint_9
+ OpStore %57 %58
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/storage/static_index/write.wgsl.expected.wgsl b/test/buffer/storage/static_index/write.wgsl.expected.wgsl
new file mode 100644
index 0000000..16725de
--- /dev/null
+++ b/test/buffer/storage/static_index/write.wgsl.expected.wgsl
@@ -0,0 +1,33 @@
+struct Inner {
+ x : i32;
+};
+
+[[block]]
+struct S {
+ a : vec3<i32>;
+ b : i32;
+ c : vec3<u32>;
+ d : u32;
+ e : vec3<f32>;
+ f : f32;
+ g : mat2x3<f32>;
+ h : mat3x2<f32>;
+ i : Inner;
+ j : [[stride(16)]] array<Inner, 4>;
+};
+
+[[binding(0), group(0)]] var<storage, write> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ s.a = vec3<i32>();
+ s.b = i32();
+ s.c = vec3<u32>();
+ s.d = u32();
+ s.e = vec3<f32>();
+ s.f = f32();
+ s.g = mat2x3<f32>();
+ s.h = mat3x2<f32>();
+ s.i = Inner();
+ s.j = [[stride(16)]] array<Inner, 4>();
+}
diff --git a/test/buffer/uniform/dynamic_index/read.wgsl b/test/buffer/uniform/dynamic_index/read.wgsl
new file mode 100644
index 0000000..12cc0c0
--- /dev/null
+++ b/test/buffer/uniform/dynamic_index/read.wgsl
@@ -0,0 +1,33 @@
+struct Inner {
+ a : vec3<i32>;
+ b : i32;
+ c : vec3<u32>;
+ d : u32;
+ e : vec3<f32>;
+ f : f32;
+ g : vec2<i32>;
+ h : vec2<i32>;
+ i : mat2x3<f32>;
+ [[align(16)]] j : mat3x2<f32>;
+ [[align(16)]] k : [[stride(16)]] array<vec4<i32>, 4>;
+};
+
+[[block]]
+struct S {
+ arr : array<Inner, 8>;
+};
+
+[[binding(0), group(0)]] var<uniform> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main([[builtin(local_invocation_index)]] idx : u32) {
+ let a = s.arr[idx].a;
+ let b = s.arr[idx].b;
+ let c = s.arr[idx].c;
+ let d = s.arr[idx].d;
+ let e = s.arr[idx].e;
+ let f = s.arr[idx].f;
+ let g = s.arr[idx].g;
+ let h = s.arr[idx].h;
+ let i = s.arr[idx].i;
+}
diff --git a/test/buffer/uniform/dynamic_index/read.wgsl.expected.hlsl b/test/buffer/uniform/dynamic_index/read.wgsl.expected.hlsl
new file mode 100644
index 0000000..378e877
--- /dev/null
+++ b/test/buffer/uniform/dynamic_index/read.wgsl.expected.hlsl
@@ -0,0 +1,38 @@
+float2x3 tint_symbol_9(uint4 buffer[96], uint offset) {
+ const uint scalar_offset = ((offset + 0u)) / 4;
+ const uint scalar_offset_1 = ((offset + 16u)) / 4;
+ return float2x3(asfloat(buffer[scalar_offset / 4].xyz), asfloat(buffer[scalar_offset_1 / 4].xyz));
+}
+
+cbuffer cbuffer_s : register(b0, space0) {
+ uint4 s[96];
+};
+
+struct tint_symbol_1 {
+ uint idx : SV_GroupIndex;
+};
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+ const uint idx = tint_symbol.idx;
+ const uint scalar_offset_2 = ((192u * idx)) / 4;
+ const int3 a = asint(s[scalar_offset_2 / 4].xyz);
+ const uint scalar_offset_3 = (((192u * idx) + 12u)) / 4;
+ const int b = asint(s[scalar_offset_3 / 4][scalar_offset_3 % 4]);
+ const uint scalar_offset_4 = (((192u * idx) + 16u)) / 4;
+ const uint3 c = s[scalar_offset_4 / 4].xyz;
+ const uint scalar_offset_5 = (((192u * idx) + 28u)) / 4;
+ const uint d = s[scalar_offset_5 / 4][scalar_offset_5 % 4];
+ const uint scalar_offset_6 = (((192u * idx) + 32u)) / 4;
+ const float3 e = asfloat(s[scalar_offset_6 / 4].xyz);
+ const uint scalar_offset_7 = (((192u * idx) + 44u)) / 4;
+ const float f = asfloat(s[scalar_offset_7 / 4][scalar_offset_7 % 4]);
+ const uint scalar_offset_8 = (((192u * idx) + 48u)) / 4;
+ uint4 ubo_load = s[scalar_offset_8 / 4];
+ const int2 g = asint(((scalar_offset_8 & 2) ? ubo_load.zw : ubo_load.xy));
+ const uint scalar_offset_9 = (((192u * idx) + 56u)) / 4;
+ uint4 ubo_load_1 = s[scalar_offset_9 / 4];
+ const int2 h = asint(((scalar_offset_9 & 2) ? ubo_load_1.zw : ubo_load_1.xy));
+ const float2x3 i = tint_symbol_9(s, ((192u * idx) + 64u));
+ return;
+}
diff --git a/test/buffer/uniform/dynamic_index/read.wgsl.expected.msl b/test/buffer/uniform/dynamic_index/read.wgsl.expected.msl
new file mode 100644
index 0000000..08c4e7d
--- /dev/null
+++ b/test/buffer/uniform/dynamic_index/read.wgsl.expected.msl
@@ -0,0 +1,40 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_array_wrapper {
+ /* 0x0000 */ int4 arr[4];
+};
+struct Inner {
+ /* 0x0000 */ packed_int3 a;
+ /* 0x000c */ int b;
+ /* 0x0010 */ packed_uint3 c;
+ /* 0x001c */ uint d;
+ /* 0x0020 */ packed_float3 e;
+ /* 0x002c */ float f;
+ /* 0x0030 */ packed_int2 g;
+ /* 0x0038 */ packed_int2 h;
+ /* 0x0040 */ float2x3 i;
+ /* 0x0060 */ float3x2 j;
+ /* 0x0078 */ int8_t tint_pad[8];
+ /* 0x0080 */ tint_array_wrapper k;
+};
+struct tint_array_wrapper_1 {
+ /* 0x0000 */ Inner arr[8];
+};
+struct S {
+ /* 0x0000 */ tint_array_wrapper_1 arr;
+};
+
+kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]], constant S& s [[buffer(0)]]) {
+ int3 const a = s.arr.arr[idx].a;
+ int const b = s.arr.arr[idx].b;
+ uint3 const c = s.arr.arr[idx].c;
+ uint const d = s.arr.arr[idx].d;
+ float3 const e = s.arr.arr[idx].e;
+ float const f = s.arr.arr[idx].f;
+ int2 const g = s.arr.arr[idx].g;
+ int2 const h = s.arr.arr[idx].h;
+ float2x3 const i = s.arr.arr[idx].i;
+ return;
+}
+
diff --git a/test/buffer/uniform/dynamic_index/read.wgsl.expected.spvasm b/test/buffer/uniform/dynamic_index/read.wgsl.expected.spvasm
new file mode 100644
index 0000000..4ea6207
--- /dev/null
+++ b/test/buffer/uniform/dynamic_index/read.wgsl.expected.spvasm
@@ -0,0 +1,118 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 68
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main" %tint_symbol
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %S "S"
+ OpMemberName %S 0 "arr"
+ OpName %Inner "Inner"
+ OpMemberName %Inner 0 "a"
+ OpMemberName %Inner 1 "b"
+ OpMemberName %Inner 2 "c"
+ OpMemberName %Inner 3 "d"
+ OpMemberName %Inner 4 "e"
+ OpMemberName %Inner 5 "f"
+ OpMemberName %Inner 6 "g"
+ OpMemberName %Inner 7 "h"
+ OpMemberName %Inner 8 "i"
+ OpMemberName %Inner 9 "j"
+ OpMemberName %Inner 10 "k"
+ OpName %s "s"
+ OpName %tint_symbol "tint_symbol"
+ OpName %main "main"
+ OpDecorate %S Block
+ OpMemberDecorate %S 0 Offset 0
+ OpMemberDecorate %Inner 0 Offset 0
+ OpMemberDecorate %Inner 1 Offset 12
+ OpMemberDecorate %Inner 2 Offset 16
+ OpMemberDecorate %Inner 3 Offset 28
+ OpMemberDecorate %Inner 4 Offset 32
+ OpMemberDecorate %Inner 5 Offset 44
+ OpMemberDecorate %Inner 6 Offset 48
+ OpMemberDecorate %Inner 7 Offset 56
+ OpMemberDecorate %Inner 8 Offset 64
+ OpMemberDecorate %Inner 8 ColMajor
+ OpMemberDecorate %Inner 8 MatrixStride 16
+ OpMemberDecorate %Inner 9 Offset 96
+ OpMemberDecorate %Inner 9 ColMajor
+ OpMemberDecorate %Inner 9 MatrixStride 8
+ OpMemberDecorate %Inner 10 Offset 128
+ OpDecorate %_arr_v4int_uint_4 ArrayStride 16
+ OpDecorate %_arr_Inner_uint_8 ArrayStride 192
+ OpDecorate %s NonWritable
+ OpDecorate %s Binding 0
+ OpDecorate %s DescriptorSet 0
+ OpDecorate %tint_symbol BuiltIn LocalInvocationIndex
+ %int = OpTypeInt 32 1
+ %v3int = OpTypeVector %int 3
+ %uint = OpTypeInt 32 0
+ %v3uint = OpTypeVector %uint 3
+ %float = OpTypeFloat 32
+ %v3float = OpTypeVector %float 3
+ %v2int = OpTypeVector %int 2
+%mat2v3float = OpTypeMatrix %v3float 2
+ %v2float = OpTypeVector %float 2
+%mat3v2float = OpTypeMatrix %v2float 3
+ %v4int = OpTypeVector %int 4
+ %uint_4 = OpConstant %uint 4
+%_arr_v4int_uint_4 = OpTypeArray %v4int %uint_4
+ %Inner = OpTypeStruct %v3int %int %v3uint %uint %v3float %float %v2int %v2int %mat2v3float %mat3v2float %_arr_v4int_uint_4
+ %uint_8 = OpConstant %uint 8
+%_arr_Inner_uint_8 = OpTypeArray %Inner %uint_8
+ %S = OpTypeStruct %_arr_Inner_uint_8
+%_ptr_Uniform_S = OpTypePointer Uniform %S
+ %s = OpVariable %_ptr_Uniform_S Uniform
+%_ptr_Input_uint = OpTypePointer Input %uint
+%tint_symbol = OpVariable %_ptr_Input_uint Input
+ %void = OpTypeVoid
+ %22 = OpTypeFunction %void
+ %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_v3int = OpTypePointer Uniform %v3int
+ %uint_1 = OpConstant %uint 1
+%_ptr_Uniform_int = OpTypePointer Uniform %int
+ %uint_2 = OpConstant %uint 2
+%_ptr_Uniform_v3uint = OpTypePointer Uniform %v3uint
+ %uint_3 = OpConstant %uint 3
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+%_ptr_Uniform_v3float = OpTypePointer Uniform %v3float
+ %uint_5 = OpConstant %uint 5
+%_ptr_Uniform_float = OpTypePointer Uniform %float
+ %uint_6 = OpConstant %uint 6
+%_ptr_Uniform_v2int = OpTypePointer Uniform %v2int
+ %uint_7 = OpConstant %uint 7
+%_ptr_Uniform_mat2v3float = OpTypePointer Uniform %mat2v3float
+ %main = OpFunction %void None %22
+ %25 = OpLabel
+ %27 = OpLoad %uint %tint_symbol
+ %29 = OpAccessChain %_ptr_Uniform_v3int %s %uint_0 %27 %uint_0
+ %30 = OpLoad %v3int %29
+ %31 = OpLoad %uint %tint_symbol
+ %34 = OpAccessChain %_ptr_Uniform_int %s %uint_0 %31 %uint_1
+ %35 = OpLoad %int %34
+ %36 = OpLoad %uint %tint_symbol
+ %39 = OpAccessChain %_ptr_Uniform_v3uint %s %uint_0 %36 %uint_2
+ %40 = OpLoad %v3uint %39
+ %41 = OpLoad %uint %tint_symbol
+ %44 = OpAccessChain %_ptr_Uniform_uint %s %uint_0 %41 %uint_3
+ %45 = OpLoad %uint %44
+ %46 = OpLoad %uint %tint_symbol
+ %48 = OpAccessChain %_ptr_Uniform_v3float %s %uint_0 %46 %uint_4
+ %49 = OpLoad %v3float %48
+ %50 = OpLoad %uint %tint_symbol
+ %53 = OpAccessChain %_ptr_Uniform_float %s %uint_0 %50 %uint_5
+ %54 = OpLoad %float %53
+ %55 = OpLoad %uint %tint_symbol
+ %58 = OpAccessChain %_ptr_Uniform_v2int %s %uint_0 %55 %uint_6
+ %59 = OpLoad %v2int %58
+ %60 = OpLoad %uint %tint_symbol
+ %62 = OpAccessChain %_ptr_Uniform_v2int %s %uint_0 %60 %uint_7
+ %63 = OpLoad %v2int %62
+ %64 = OpLoad %uint %tint_symbol
+ %66 = OpAccessChain %_ptr_Uniform_mat2v3float %s %uint_0 %64 %uint_8
+ %67 = OpLoad %mat2v3float %66
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/uniform/dynamic_index/read.wgsl.expected.wgsl b/test/buffer/uniform/dynamic_index/read.wgsl.expected.wgsl
new file mode 100644
index 0000000..7377f9e
--- /dev/null
+++ b/test/buffer/uniform/dynamic_index/read.wgsl.expected.wgsl
@@ -0,0 +1,35 @@
+struct Inner {
+ a : vec3<i32>;
+ b : i32;
+ c : vec3<u32>;
+ d : u32;
+ e : vec3<f32>;
+ f : f32;
+ g : vec2<i32>;
+ h : vec2<i32>;
+ i : mat2x3<f32>;
+ [[align(16)]]
+ j : mat3x2<f32>;
+ [[align(16)]]
+ k : [[stride(16)]] array<vec4<i32>, 4>;
+};
+
+[[block]]
+struct S {
+ arr : array<Inner, 8>;
+};
+
+[[binding(0), group(0)]] var<uniform> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main([[builtin(local_invocation_index)]] idx : u32) {
+ let a = s.arr[idx].a;
+ let b = s.arr[idx].b;
+ let c = s.arr[idx].c;
+ let d = s.arr[idx].d;
+ let e = s.arr[idx].e;
+ let f = s.arr[idx].f;
+ let g = s.arr[idx].g;
+ let h = s.arr[idx].h;
+ let i = s.arr[idx].i;
+}
diff --git a/test/buffer/uniform/static_index/read.wgsl b/test/buffer/uniform/static_index/read.wgsl
new file mode 100644
index 0000000..c3ad9ad
--- /dev/null
+++ b/test/buffer/uniform/static_index/read.wgsl
@@ -0,0 +1,37 @@
+struct Inner {
+ x : i32;
+};
+
+[[block]]
+struct S {
+ a : vec3<i32>;
+ b : i32;
+ c : vec3<u32>;
+ d : u32;
+ e : vec3<f32>;
+ f : f32;
+ g : vec2<i32>;
+ h : vec2<i32>;
+ i : mat2x3<f32>;
+ j : mat3x2<f32>;
+ [[align(16)]] k : Inner;
+ [[align(16)]] l : [[stride(16)]] array<Inner, 4>;
+};
+
+[[binding(0), group(0)]] var<uniform> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let a = s.a;
+ let b = s.b;
+ let c = s.c;
+ let d = s.d;
+ let e = s.e;
+ let f = s.f;
+ let g = s.g;
+ let h = s.h;
+ let i = s.i;
+ let j = s.j;
+ let k = s.k;
+ let l = s.l;
+}
diff --git a/test/buffer/uniform/static_index/read.wgsl.expected.hlsl b/test/buffer/uniform/static_index/read.wgsl.expected.hlsl
new file mode 100644
index 0000000..7acef2b
--- /dev/null
+++ b/test/buffer/uniform/static_index/read.wgsl.expected.hlsl
@@ -0,0 +1,55 @@
+struct Inner {
+ int x;
+};
+struct tint_padded_array_element {
+ Inner el;
+};
+
+float2x3 tint_symbol_7(uint4 buffer[13], uint offset) {
+ const uint scalar_offset = ((offset + 0u)) / 4;
+ const uint scalar_offset_1 = ((offset + 16u)) / 4;
+ return float2x3(asfloat(buffer[scalar_offset / 4].xyz), asfloat(buffer[scalar_offset_1 / 4].xyz));
+}
+
+float3x2 tint_symbol_9(uint4 buffer[13], uint offset) {
+ const uint scalar_offset_2 = ((offset + 0u)) / 4;
+ uint4 ubo_load = buffer[scalar_offset_2 / 4];
+ const uint scalar_offset_3 = ((offset + 8u)) / 4;
+ uint4 ubo_load_1 = buffer[scalar_offset_3 / 4];
+ const uint scalar_offset_4 = ((offset + 16u)) / 4;
+ uint4 ubo_load_2 = buffer[scalar_offset_4 / 4];
+ return float3x2(asfloat(((scalar_offset_2 & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_3 & 2) ? ubo_load_1.zw : ubo_load_1.xy)), asfloat(((scalar_offset_4 & 2) ? ubo_load_2.zw : ubo_load_2.xy)));
+}
+
+Inner tint_symbol_10(uint4 buffer[13], uint offset) {
+ const uint scalar_offset_5 = ((offset + 0u)) / 4;
+ const Inner tint_symbol_12 = {asint(buffer[scalar_offset_5 / 4][scalar_offset_5 % 4])};
+ return tint_symbol_12;
+}
+
+typedef tint_padded_array_element tint_symbol_11_ret[4];
+tint_symbol_11_ret tint_symbol_11(uint4 buffer[13], uint offset) {
+ const tint_padded_array_element tint_symbol_13[4] = {{tint_symbol_10(buffer, (offset + 0u))}, {tint_symbol_10(buffer, (offset + 16u))}, {tint_symbol_10(buffer, (offset + 32u))}, {tint_symbol_10(buffer, (offset + 48u))}};
+ return tint_symbol_13;
+}
+
+cbuffer cbuffer_s : register(b0, space0) {
+ uint4 s[13];
+};
+
+[numthreads(1, 1, 1)]
+void main() {
+ const int3 a = asint(s[0].xyz);
+ const int b = asint(s[0].w);
+ const uint3 c = s[1].xyz;
+ const uint d = s[1].w;
+ const float3 e = asfloat(s[2].xyz);
+ const float f = asfloat(s[2].w);
+ const int2 g = asint(s[3].xy);
+ const int2 h = asint(s[3].zw);
+ const float2x3 i = tint_symbol_7(s, 64u);
+ const float3x2 j = tint_symbol_9(s, 96u);
+ const Inner k = tint_symbol_10(s, 128u);
+ const tint_padded_array_element l[4] = tint_symbol_11(s, 144u);
+ return;
+}
diff --git a/test/buffer/uniform/static_index/read.wgsl.expected.msl b/test/buffer/uniform/static_index/read.wgsl.expected.msl
new file mode 100644
index 0000000..bc3ce89
--- /dev/null
+++ b/test/buffer/uniform/static_index/read.wgsl.expected.msl
@@ -0,0 +1,46 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Inner {
+ /* 0x0000 */ int x;
+};
+struct tint_padded_array_element {
+ /* 0x0000 */ Inner el;
+ /* 0x0004 */ int8_t tint_pad[12];
+};
+struct tint_array_wrapper {
+ /* 0x0000 */ tint_padded_array_element arr[4];
+};
+struct S {
+ /* 0x0000 */ packed_int3 a;
+ /* 0x000c */ int b;
+ /* 0x0010 */ packed_uint3 c;
+ /* 0x001c */ uint d;
+ /* 0x0020 */ packed_float3 e;
+ /* 0x002c */ float f;
+ /* 0x0030 */ packed_int2 g;
+ /* 0x0038 */ packed_int2 h;
+ /* 0x0040 */ float2x3 i;
+ /* 0x0060 */ float3x2 j;
+ /* 0x0078 */ int8_t tint_pad_1[8];
+ /* 0x0080 */ Inner k;
+ /* 0x0084 */ int8_t tint_pad_2[12];
+ /* 0x0090 */ tint_array_wrapper l;
+};
+
+kernel void tint_symbol(constant S& s [[buffer(0)]]) {
+ int3 const a = s.a;
+ int const b = s.b;
+ uint3 const c = s.c;
+ uint const d = s.d;
+ float3 const e = s.e;
+ float const f = s.f;
+ int2 const g = s.g;
+ int2 const h = s.h;
+ float2x3 const i = s.i;
+ float3x2 const j = s.j;
+ Inner const k = s.k;
+ tint_array_wrapper const l = s.l;
+ return;
+}
+
diff --git a/test/buffer/uniform/static_index/read.wgsl.expected.spvasm b/test/buffer/uniform/static_index/read.wgsl.expected.spvasm
new file mode 100644
index 0000000..25854c6
--- /dev/null
+++ b/test/buffer/uniform/static_index/read.wgsl.expected.spvasm
@@ -0,0 +1,116 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 67
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %S "S"
+ OpMemberName %S 0 "a"
+ OpMemberName %S 1 "b"
+ OpMemberName %S 2 "c"
+ OpMemberName %S 3 "d"
+ OpMemberName %S 4 "e"
+ OpMemberName %S 5 "f"
+ OpMemberName %S 6 "g"
+ OpMemberName %S 7 "h"
+ OpMemberName %S 8 "i"
+ OpMemberName %S 9 "j"
+ OpMemberName %S 10 "k"
+ OpName %Inner "Inner"
+ OpMemberName %Inner 0 "x"
+ OpMemberName %S 11 "l"
+ OpName %s "s"
+ OpName %main "main"
+ OpDecorate %S Block
+ OpMemberDecorate %S 0 Offset 0
+ OpMemberDecorate %S 1 Offset 12
+ OpMemberDecorate %S 2 Offset 16
+ OpMemberDecorate %S 3 Offset 28
+ OpMemberDecorate %S 4 Offset 32
+ OpMemberDecorate %S 5 Offset 44
+ OpMemberDecorate %S 6 Offset 48
+ OpMemberDecorate %S 7 Offset 56
+ OpMemberDecorate %S 8 Offset 64
+ OpMemberDecorate %S 8 ColMajor
+ OpMemberDecorate %S 8 MatrixStride 16
+ OpMemberDecorate %S 9 Offset 96
+ OpMemberDecorate %S 9 ColMajor
+ OpMemberDecorate %S 9 MatrixStride 8
+ OpMemberDecorate %S 10 Offset 128
+ OpMemberDecorate %Inner 0 Offset 0
+ OpMemberDecorate %S 11 Offset 144
+ OpDecorate %_arr_Inner_uint_4 ArrayStride 16
+ OpDecorate %s NonWritable
+ OpDecorate %s Binding 0
+ OpDecorate %s DescriptorSet 0
+ %int = OpTypeInt 32 1
+ %v3int = OpTypeVector %int 3
+ %uint = OpTypeInt 32 0
+ %v3uint = OpTypeVector %uint 3
+ %float = OpTypeFloat 32
+ %v3float = OpTypeVector %float 3
+ %v2int = OpTypeVector %int 2
+%mat2v3float = OpTypeMatrix %v3float 2
+ %v2float = OpTypeVector %float 2
+%mat3v2float = OpTypeMatrix %v2float 3
+ %Inner = OpTypeStruct %int
+ %uint_4 = OpConstant %uint 4
+%_arr_Inner_uint_4 = OpTypeArray %Inner %uint_4
+ %S = OpTypeStruct %v3int %int %v3uint %uint %v3float %float %v2int %v2int %mat2v3float %mat3v2float %Inner %_arr_Inner_uint_4
+%_ptr_Uniform_S = OpTypePointer Uniform %S
+ %s = OpVariable %_ptr_Uniform_S Uniform
+ %void = OpTypeVoid
+ %17 = OpTypeFunction %void
+ %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_v3int = OpTypePointer Uniform %v3int
+ %uint_1 = OpConstant %uint 1
+%_ptr_Uniform_int = OpTypePointer Uniform %int
+ %uint_2 = OpConstant %uint 2
+%_ptr_Uniform_v3uint = OpTypePointer Uniform %v3uint
+ %uint_3 = OpConstant %uint 3
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+%_ptr_Uniform_v3float = OpTypePointer Uniform %v3float
+ %uint_5 = OpConstant %uint 5
+%_ptr_Uniform_float = OpTypePointer Uniform %float
+ %uint_6 = OpConstant %uint 6
+%_ptr_Uniform_v2int = OpTypePointer Uniform %v2int
+ %uint_7 = OpConstant %uint 7
+ %uint_8 = OpConstant %uint 8
+%_ptr_Uniform_mat2v3float = OpTypePointer Uniform %mat2v3float
+ %uint_9 = OpConstant %uint 9
+%_ptr_Uniform_mat3v2float = OpTypePointer Uniform %mat3v2float
+ %uint_10 = OpConstant %uint 10
+%_ptr_Uniform_Inner = OpTypePointer Uniform %Inner
+ %uint_11 = OpConstant %uint 11
+%_ptr_Uniform__arr_Inner_uint_4 = OpTypePointer Uniform %_arr_Inner_uint_4
+ %main = OpFunction %void None %17
+ %20 = OpLabel
+ %23 = OpAccessChain %_ptr_Uniform_v3int %s %uint_0
+ %24 = OpLoad %v3int %23
+ %27 = OpAccessChain %_ptr_Uniform_int %s %uint_1
+ %28 = OpLoad %int %27
+ %31 = OpAccessChain %_ptr_Uniform_v3uint %s %uint_2
+ %32 = OpLoad %v3uint %31
+ %35 = OpAccessChain %_ptr_Uniform_uint %s %uint_3
+ %36 = OpLoad %uint %35
+ %38 = OpAccessChain %_ptr_Uniform_v3float %s %uint_4
+ %39 = OpLoad %v3float %38
+ %42 = OpAccessChain %_ptr_Uniform_float %s %uint_5
+ %43 = OpLoad %float %42
+ %46 = OpAccessChain %_ptr_Uniform_v2int %s %uint_6
+ %47 = OpLoad %v2int %46
+ %49 = OpAccessChain %_ptr_Uniform_v2int %s %uint_7
+ %50 = OpLoad %v2int %49
+ %53 = OpAccessChain %_ptr_Uniform_mat2v3float %s %uint_8
+ %54 = OpLoad %mat2v3float %53
+ %57 = OpAccessChain %_ptr_Uniform_mat3v2float %s %uint_9
+ %58 = OpLoad %mat3v2float %57
+ %61 = OpAccessChain %_ptr_Uniform_Inner %s %uint_10
+ %62 = OpLoad %Inner %61
+ %65 = OpAccessChain %_ptr_Uniform__arr_Inner_uint_4 %s %uint_11
+ %66 = OpLoad %_arr_Inner_uint_4 %65
+ OpReturn
+ OpFunctionEnd
diff --git a/test/buffer/uniform/static_index/read.wgsl.expected.wgsl b/test/buffer/uniform/static_index/read.wgsl.expected.wgsl
new file mode 100644
index 0000000..d856b8c
--- /dev/null
+++ b/test/buffer/uniform/static_index/read.wgsl.expected.wgsl
@@ -0,0 +1,39 @@
+struct Inner {
+ x : i32;
+};
+
+[[block]]
+struct S {
+ a : vec3<i32>;
+ b : i32;
+ c : vec3<u32>;
+ d : u32;
+ e : vec3<f32>;
+ f : f32;
+ g : vec2<i32>;
+ h : vec2<i32>;
+ i : mat2x3<f32>;
+ j : mat3x2<f32>;
+ [[align(16)]]
+ k : Inner;
+ [[align(16)]]
+ l : [[stride(16)]] array<Inner, 4>;
+};
+
+[[binding(0), group(0)]] var<uniform> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let a = s.a;
+ let b = s.b;
+ let c = s.c;
+ let d = s.d;
+ let e = s.e;
+ let f = s.f;
+ let g = s.g;
+ let h = s.h;
+ let i = s.i;
+ let j = s.j;
+ let k = s.k;
+ let l = s.l;
+}
diff --git a/test/bug/dawn/947.wgsl.expected.hlsl b/test/bug/dawn/947.wgsl.expected.hlsl
index a65f3b1..56a4250 100644
--- a/test/bug/dawn/947.wgsl.expected.hlsl
+++ b/test/bug/dawn/947.wgsl.expected.hlsl
@@ -19,20 +19,11 @@
float2 texcoord[3] = {float2(-0.5f, 0.0f), float2(1.5f, 0.0f), float2(0.5f, 2.0f)};
VertexOutputs output = (VertexOutputs)0;
output.position = float4(((texcoord[VertexIndex] * 2.0f) - float2(1.0f, 1.0f)), 0.0f, 1.0f);
- const uint scalar_offset = (4u) / 4;
- bool flipY = (asfloat(uniforms[scalar_offset / 4][scalar_offset % 4]) < 0.0f);
+ bool flipY = (asfloat(uniforms[0].y) < 0.0f);
if (flipY) {
- const uint scalar_offset_1 = (0u) / 4;
- uint4 ubo_load = uniforms[scalar_offset_1 / 4];
- const uint scalar_offset_2 = (8u) / 4;
- uint4 ubo_load_1 = uniforms[scalar_offset_2 / 4];
- output.texcoords = ((((texcoord[VertexIndex] * asfloat(((scalar_offset_1 & 2) ? ubo_load.zw : ubo_load.xy))) + asfloat(((scalar_offset_2 & 2) ? ubo_load_1.zw : ubo_load_1.xy))) * float2(1.0f, -1.0f)) + float2(0.0f, 1.0f));
+ output.texcoords = ((((texcoord[VertexIndex] * asfloat(uniforms[0].xy)) + asfloat(uniforms[0].zw)) * float2(1.0f, -1.0f)) + float2(0.0f, 1.0f));
} else {
- const uint scalar_offset_3 = (0u) / 4;
- uint4 ubo_load_2 = uniforms[scalar_offset_3 / 4];
- const uint scalar_offset_4 = (8u) / 4;
- uint4 ubo_load_3 = uniforms[scalar_offset_4 / 4];
- output.texcoords = ((((texcoord[VertexIndex] * float2(1.0f, -1.0f)) + float2(0.0f, 1.0f)) * asfloat(((scalar_offset_3 & 2) ? ubo_load_2.zw : ubo_load_2.xy))) + asfloat(((scalar_offset_4 & 2) ? ubo_load_3.zw : ubo_load_3.xy)));
+ output.texcoords = ((((texcoord[VertexIndex] * float2(1.0f, -1.0f)) + float2(0.0f, 1.0f)) * asfloat(uniforms[0].xy)) + asfloat(uniforms[0].zw));
}
const tint_symbol_2 tint_symbol_8 = {output.texcoords, output.position};
return tint_symbol_8;
diff --git a/test/bug/tint/534.wgsl.expected.hlsl b/test/bug/tint/534.wgsl.expected.hlsl
index e7a735c..da61ba4 100644
--- a/test/bug/tint/534.wgsl.expected.hlsl
+++ b/test/bug/tint/534.wgsl.expected.hlsl
@@ -30,8 +30,7 @@
int2 size = tint_tmp;
int2 dstTexCoord = int2(GlobalInvocationID.xy);
int2 srcTexCoord = dstTexCoord;
- const uint scalar_offset = (0u) / 4;
- if ((uniforms[scalar_offset / 4][scalar_offset % 4] == 1u)) {
+ if ((uniforms[0].x == 1u)) {
srcTexCoord.y = ((size.y - dstTexCoord.y) - 1);
}
float4 srcColor = src.Load(int3(srcTexCoord, 0));
@@ -40,17 +39,13 @@
uint4 srcColorBits = uint4(0u, 0u, 0u, 0u);
uint4 dstColorBits = uint4(dstColor);
{
- uint i = 0u;
- while (true) {
- const uint scalar_offset_1 = (12u) / 4;
- if (!((i < uniforms[scalar_offset_1 / 4][scalar_offset_1 % 4]))) { break; }
+ for(uint i = 0u; (i < uniforms[0].w); i = (i + 1u)) {
Set_uint4(srcColorBits, i, ConvertToFp16FloatValue(srcColor[i]));
bool tint_tmp_1 = success;
if (tint_tmp_1) {
tint_tmp_1 = (srcColorBits[i] == dstColorBits[i]);
}
success = (tint_tmp_1);
- i = (i + 1u);
}
}
uint outputIndex = ((GlobalInvocationID.y * uint(size.x)) + GlobalInvocationID.x);
diff --git a/test/bug/tint/744.wgsl.expected.hlsl b/test/bug/tint/744.wgsl.expected.hlsl
index 47f47a7..c7bab90 100644
--- a/test/bug/tint/744.wgsl.expected.hlsl
+++ b/test/bug/tint/744.wgsl.expected.hlsl
@@ -13,10 +13,8 @@
void main(tint_symbol_1 tint_symbol) {
const uint3 global_id = tint_symbol.global_id;
const uint2 resultCell = uint2(global_id.y, global_id.x);
- const uint scalar_offset = (4u) / 4;
- const uint dimInner = uniforms[scalar_offset / 4][scalar_offset % 4];
- const uint scalar_offset_1 = (20u) / 4;
- const uint dimOutter = uniforms[scalar_offset_1 / 4][scalar_offset_1 % 4];
+ const uint dimInner = uniforms[0].y;
+ const uint dimOutter = uniforms[1].y;
uint result = 0u;
{
for(uint i = 0u; (i < dimInner); i = (i + 1u)) {
diff --git a/test/bug/tint/749.spvasm.expected.hlsl b/test/bug/tint/749.spvasm.expected.hlsl
index 5d8b15b..c0db54c 100644
--- a/test/bug/tint/749.spvasm.expected.hlsl
+++ b/test/bug/tint/749.spvasm.expected.hlsl
@@ -861,9 +861,7 @@
const float2 x_762 = uv;
uv = float2(0.0f, 0.0f);
uv = x_762;
- const uint scalar_offset = (0u) / 4;
- uint4 ubo_load = x_188[scalar_offset / 4];
- const float2 x_191 = asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy));
+ const float2 x_191 = asfloat(x_188[0].xy);
const QuicksortObject x_763 = obj;
const int tint_symbol_51[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
const QuicksortObject tint_symbol_52 = {tint_symbol_51};
diff --git a/test/bug/tint/913.wgsl.expected.hlsl b/test/bug/tint/913.wgsl.expected.hlsl
index 18dca3f..2e66d2a 100644
--- a/test/bug/tint/913.wgsl.expected.hlsl
+++ b/test/bug/tint/913.wgsl.expected.hlsl
@@ -25,23 +25,17 @@
const uint2 dstTexCoord = uint2(GlobalInvocationID.xy);
const float4 nonCoveredColor = float4(0.0f, 1.0f, 0.0f, 1.0f);
bool success = true;
- const uint scalar_offset = (16u) / 4;
- bool tint_tmp_4 = (dstTexCoord.x < uniforms[scalar_offset / 4][scalar_offset % 4]);
+ bool tint_tmp_4 = (dstTexCoord.x < uniforms[1].x);
if (!tint_tmp_4) {
- const uint scalar_offset_1 = (20u) / 4;
- tint_tmp_4 = (dstTexCoord.y < uniforms[scalar_offset_1 / 4][scalar_offset_1 % 4]);
+ tint_tmp_4 = (dstTexCoord.y < uniforms[1].y);
}
bool tint_tmp_3 = (tint_tmp_4);
if (!tint_tmp_3) {
- const uint scalar_offset_2 = (16u) / 4;
- const uint scalar_offset_3 = (24u) / 4;
- tint_tmp_3 = (dstTexCoord.x >= (uniforms[scalar_offset_2 / 4][scalar_offset_2 % 4] + uniforms[scalar_offset_3 / 4][scalar_offset_3 % 4]));
+ tint_tmp_3 = (dstTexCoord.x >= (uniforms[1].x + uniforms[1].z));
}
bool tint_tmp_2 = (tint_tmp_3);
if (!tint_tmp_2) {
- const uint scalar_offset_4 = (20u) / 4;
- const uint scalar_offset_5 = (28u) / 4;
- tint_tmp_2 = (dstTexCoord.y >= (uniforms[scalar_offset_4 / 4][scalar_offset_4 % 4] + uniforms[scalar_offset_5 / 4][scalar_offset_5 % 4]));
+ tint_tmp_2 = (dstTexCoord.y >= (uniforms[1].y + uniforms[1].w));
}
if ((tint_tmp_2)) {
bool tint_tmp_5 = success;
@@ -50,19 +44,13 @@
}
success = (tint_tmp_5);
} else {
- const uint scalar_offset_6 = (16u) / 4;
- uint4 ubo_load = uniforms[scalar_offset_6 / 4];
- const uint scalar_offset_7 = (8u) / 4;
- uint4 ubo_load_1 = uniforms[scalar_offset_7 / 4];
- uint2 srcTexCoord = ((dstTexCoord - ((scalar_offset_6 & 2) ? ubo_load.zw : ubo_load.xy)) + ((scalar_offset_7 & 2) ? ubo_load_1.zw : ubo_load_1.xy));
- const uint scalar_offset_8 = (0u) / 4;
- if ((uniforms[scalar_offset_8 / 4][scalar_offset_8 % 4] == 1u)) {
+ uint2 srcTexCoord = ((dstTexCoord - uniforms[1].xy) + uniforms[0].zw);
+ if ((uniforms[0].x == 1u)) {
srcTexCoord.y = ((uint(srcSize.y) - srcTexCoord.y) - 1u);
}
const float4 srcColor = src.Load(int3(srcTexCoord, 0));
const float4 dstColor = tint_symbol.Load(int3(dstTexCoord, 0));
- const uint scalar_offset_9 = (4u) / 4;
- if ((uniforms[scalar_offset_9 / 4][scalar_offset_9 % 4] == 2u)) {
+ if ((uniforms[0].y == 2u)) {
bool tint_tmp_7 = success;
if (tint_tmp_7) {
tint_tmp_7 = aboutEqual(dstColor.r, srcColor.r);
diff --git a/test/bug/tint/914.wgsl.expected.hlsl b/test/bug/tint/914.wgsl.expected.hlsl
index 5e0ebb4..00be3a9 100644
--- a/test/bug/tint/914.wgsl.expected.hlsl
+++ b/test/bug/tint/914.wgsl.expected.hlsl
@@ -6,45 +6,36 @@
};
float mm_readA(uint row, uint col) {
- const uint scalar_offset = (0u) / 4;
- bool tint_tmp = (row < uniforms[scalar_offset / 4][scalar_offset % 4]);
+ bool tint_tmp = (row < uniforms[0].x);
if (tint_tmp) {
- const uint scalar_offset_1 = (4u) / 4;
- tint_tmp = (col < uniforms[scalar_offset_1 / 4][scalar_offset_1 % 4]);
+ tint_tmp = (col < uniforms[0].y);
}
if ((tint_tmp)) {
- const uint scalar_offset_2 = (4u) / 4;
- const float result = asfloat(firstMatrix.Load((4u * ((row * uniforms[scalar_offset_2 / 4][scalar_offset_2 % 4]) + col))));
+ const float result = asfloat(firstMatrix.Load((4u * ((row * uniforms[0].y) + col))));
return result;
}
return 0.0f;
}
float mm_readB(uint row, uint col) {
- const uint scalar_offset_3 = (4u) / 4;
- bool tint_tmp_1 = (row < uniforms[scalar_offset_3 / 4][scalar_offset_3 % 4]);
+ bool tint_tmp_1 = (row < uniforms[0].y);
if (tint_tmp_1) {
- const uint scalar_offset_4 = (8u) / 4;
- tint_tmp_1 = (col < uniforms[scalar_offset_4 / 4][scalar_offset_4 % 4]);
+ tint_tmp_1 = (col < uniforms[0].z);
}
if ((tint_tmp_1)) {
- const uint scalar_offset_5 = (8u) / 4;
- const float result = asfloat(secondMatrix.Load((4u * ((row * uniforms[scalar_offset_5 / 4][scalar_offset_5 % 4]) + col))));
+ const float result = asfloat(secondMatrix.Load((4u * ((row * uniforms[0].z) + col))));
return result;
}
return 0.0f;
}
void mm_write(uint row, uint col, float value) {
- const uint scalar_offset_6 = (0u) / 4;
- bool tint_tmp_2 = (row < uniforms[scalar_offset_6 / 4][scalar_offset_6 % 4]);
+ bool tint_tmp_2 = (row < uniforms[0].x);
if (tint_tmp_2) {
- const uint scalar_offset_7 = (8u) / 4;
- tint_tmp_2 = (col < uniforms[scalar_offset_7 / 4][scalar_offset_7 % 4]);
+ tint_tmp_2 = (col < uniforms[0].z);
}
if ((tint_tmp_2)) {
- const uint scalar_offset_8 = (8u) / 4;
- const uint index = (col + (row * uniforms[scalar_offset_8 / 4][scalar_offset_8 % 4]));
+ const uint index = (col + (row * uniforms[0].z));
resultMatrix.Store((4u * index), asuint(value));
}
}
@@ -93,8 +84,7 @@
const uint tileCol = (local_id.x * ColPerThread);
const uint globalRow = (global_id.y * RowPerThread);
const uint globalCol = (global_id.x * ColPerThread);
- const uint scalar_offset_9 = (4u) / 4;
- const uint numTiles = (((uniforms[scalar_offset_9 / 4][scalar_offset_9 % 4] - 1u) / TileInner) + 1u);
+ const uint numTiles = (((uniforms[0].y - 1u) / TileInner) + 1u);
float acc[16] = (float[16])0;
float ACached = 0.0f;
float BCached[4] = (float[4])0;
diff --git a/test/bug/tint/922.wgsl.expected.hlsl b/test/bug/tint/922.wgsl.expected.hlsl
index 964eb97..60eb716 100644
--- a/test/bug/tint/922.wgsl.expected.hlsl
+++ b/test/bug/tint/922.wgsl.expected.hlsl
@@ -196,8 +196,7 @@
const float4 _e49 = Mul(_e35, _e48);
gl_Position = _e49;
v_Color = a_Color1;
- const uint scalar_offset_9 = (32u) / 4;
- const float4 _e52 = asfloat(global1[scalar_offset_9 / 4]);
+ const float4 _e52 = asfloat(global1[2]);
if ((_e52.x == 2.0f)) {
{
const float3 _e59 = a_Normal1;
diff --git a/test/bug/tint/942.wgsl.expected.hlsl b/test/bug/tint/942.wgsl.expected.hlsl
index 695e8ae..ba42782 100644
--- a/test/bug/tint/942.wgsl.expected.hlsl
+++ b/test/bug/tint/942.wgsl.expected.hlsl
@@ -33,20 +33,17 @@
}
}
GroupMemoryBarrierWithGroupSync();
- const uint scalar_offset = (0u) / 4;
- const uint filterOffset = ((params[scalar_offset / 4][scalar_offset % 4] - 1u) / 2u);
+ const uint filterOffset = ((params[0].x - 1u) / 2u);
int3 tint_tmp;
inputTex.GetDimensions(0, tint_tmp.x, tint_tmp.y, tint_tmp.z);
const int2 dims = tint_tmp.xy;
- const uint scalar_offset_1 = (4u) / 4;
- const int2 baseIndex = (int2(((WorkGroupID.xy * uint2(params[scalar_offset_1 / 4][scalar_offset_1 % 4], 4u)) + (LocalInvocationID.xy * uint2(4u, 1u)))) - int2(int(filterOffset), 0));
+ const int2 baseIndex = (int2(((WorkGroupID.xy * uint2(params[0].y, 4u)) + (LocalInvocationID.xy * uint2(4u, 1u)))) - int2(int(filterOffset), 0));
{
for(uint r = 0u; (r < 4u); r = (r + 1u)) {
{
for(uint c = 0u; (c < 4u); c = (c + 1u)) {
int2 loadIndex = (baseIndex + int2(int(c), int(r)));
- const uint scalar_offset_2 = (0u) / 4;
- if ((flip[scalar_offset_2 / 4][scalar_offset_2 % 4] != 0u)) {
+ if ((flip[0].x != 0u)) {
loadIndex = loadIndex.yx;
}
tile[r][((4u * LocalInvocationID.x) + c)] = inputTex.SampleLevel(samp, ((float2(loadIndex) + float2(0.25f, 0.25f)) / float2(dims)), 0.0f).rgb;
@@ -60,8 +57,7 @@
{
for(uint c = 0u; (c < 4u); c = (c + 1u)) {
int2 writeIndex = (baseIndex + int2(int(c), int(r)));
- const uint scalar_offset_3 = (0u) / 4;
- if ((flip[scalar_offset_3 / 4][scalar_offset_3 % 4] != 0u)) {
+ if ((flip[0].x != 0u)) {
writeIndex = writeIndex.yx;
}
const uint center = ((4u * LocalInvocationID.x) + c);
@@ -76,14 +72,9 @@
if ((tint_tmp_1)) {
float3 acc = float3(0.0f, 0.0f, 0.0f);
{
- uint f = 0u;
- while (true) {
- const uint scalar_offset_4 = (0u) / 4;
- if (!((f < params[scalar_offset_4 / 4][scalar_offset_4 % 4]))) { break; }
+ for(uint f = 0u; (f < params[0].x); f = (f + 1u)) {
uint i = ((center + f) - filterOffset);
- const uint scalar_offset_5 = (0u) / 4;
- acc = (acc + ((1.0f / float(params[scalar_offset_5 / 4][scalar_offset_5 % 4])) * tile[r][i]));
- f = (f + 1u);
+ acc = (acc + ((1.0f / float(params[0].x)) * tile[r][i]));
}
}
outputTex[writeIndex] = float4(acc, 1.0f);
diff --git a/test/bug/tint/943.spvasm.expected.hlsl b/test/bug/tint/943.spvasm.expected.hlsl
index a317e46..0087ca2 100644
--- a/test/bug/tint/943.spvasm.expected.hlsl
+++ b/test/bug/tint/943.spvasm.expected.hlsl
@@ -33,10 +33,8 @@
int2 param_10 = int2(0, 0);
int2 param_11 = int2(0, 0);
float x_430 = 0.0f;
- const uint scalar_offset = (20u) / 4;
- const int x_417 = asint(x_48[scalar_offset / 4][scalar_offset % 4]);
- const uint scalar_offset_1 = (24u) / 4;
- const int x_419 = asint(x_48[scalar_offset_1 / 4][scalar_offset_1 % 4]);
+ const int x_417 = asint(x_48[1].y);
+ const int x_419 = asint(x_48[1].z);
batchASize = (x_417 * x_419);
const int x_421 = row;
const int x_422 = col;
@@ -64,10 +62,8 @@
int2 param_12 = int2(0, 0);
int2 param_13 = int2(0, 0);
float x_468 = 0.0f;
- const uint scalar_offset_2 = (36u) / 4;
- const int x_455 = asint(x_48[scalar_offset_2 / 4][scalar_offset_2 % 4]);
- const uint scalar_offset_3 = (40u) / 4;
- const int x_457 = asint(x_48[scalar_offset_3 / 4][scalar_offset_3 % 4]);
+ const int x_455 = asint(x_48[2].y);
+ const int x_457 = asint(x_48[2].z);
batchBSize = (x_455 * x_457);
const int x_459 = row_1;
const int x_460 = col_1;
@@ -92,10 +88,8 @@
int getOutputFlatIndex_vi3_(inout int3 coords) {
const int3 x_99 = coords;
- const uint scalar_offset_4 = (64u) / 4;
- const int x_105 = asint(x_48[scalar_offset_4 / 4][scalar_offset_4 % 4]);
- const uint scalar_offset_5 = (68u) / 4;
- const int x_107 = asint(x_48[scalar_offset_5 / 4][scalar_offset_5 % 4]);
+ const int x_105 = asint(x_48[4].x);
+ const int x_107 = asint(x_48[4].y);
return int(dot(float3(x_99), float3(int3(x_105, x_107, 1))));
}
@@ -330,14 +324,11 @@
int param_18 = 0;
int param_19 = 0;
int param_20 = 0;
- const uint scalar_offset_6 = (20u) / 4;
- const int x_67 = asint(x_48[scalar_offset_6 / 4][scalar_offset_6 % 4]);
+ const int x_67 = asint(x_48[1].y);
dimAOuter_1 = x_67;
- const uint scalar_offset_7 = (24u) / 4;
- const int x_71 = asint(x_48[scalar_offset_7 / 4][scalar_offset_7 % 4]);
+ const int x_71 = asint(x_48[1].z);
dimInner_1 = x_71;
- const uint scalar_offset_8 = (40u) / 4;
- const int x_75 = asint(x_48[scalar_offset_8 / 4][scalar_offset_8 % 4]);
+ const int x_75 = asint(x_48[2].z);
dimBOuter_1 = x_75;
const uint x_505 = gl_GlobalInvocationID.z;
batch = asint(x_505);
diff --git a/test/bug/tint/948.wgsl.expected.hlsl b/test/bug/tint/948.wgsl.expected.hlsl
index d6f3a02..694cd53 100644
--- a/test/bug/tint/948.wgsl.expected.hlsl
+++ b/test/bug/tint/948.wgsl.expected.hlsl
@@ -22,8 +22,7 @@
float4x4 getFrameData_f1_(inout float frameID) {
float fX = 0.0f;
const float x_15 = frameID;
- const uint scalar_offset = (108u) / 4;
- const float x_25 = asfloat(x_20[scalar_offset / 4][scalar_offset % 4]);
+ const float x_25 = asfloat(x_20[6].w);
fX = (x_15 / x_25);
const float4 x_40 = frameMapTexture.SampleBias(frameMapSampler, float2(fX, 0.0f), 0.0f);
const float4 x_47 = frameMapTexture.SampleBias(frameMapSampler, float2(fX, 0.25f), 0.0f);
@@ -55,16 +54,11 @@
const float x_91 = tileUV.y;
tileUV.y = (1.0f - x_91);
tileID = floor(tUV);
- const uint scalar_offset_1 = (96u) / 4;
- uint4 ubo_load = x_20[scalar_offset_1 / 4];
- const float2 x_101 = asfloat(((scalar_offset_1 & 2) ? ubo_load.zw : ubo_load.xy));
+ const float2 x_101 = asfloat(x_20[6].xy);
sheetUnits = (float2(1.0f, 1.0f) / x_101);
- const uint scalar_offset_2 = (108u) / 4;
- const float x_106 = asfloat(x_20[scalar_offset_2 / 4][scalar_offset_2 % 4]);
+ const float x_106 = asfloat(x_20[6].w);
spriteUnits = (1.0f / x_106);
- const uint scalar_offset_3 = (88u) / 4;
- uint4 ubo_load_1 = x_20[scalar_offset_3 / 4];
- const float2 x_111 = asfloat(((scalar_offset_3 & 2) ? ubo_load_1.zw : ubo_load_1.xy));
+ const float2 x_111 = asfloat(x_20[5].zw);
stageUnits = (float2(1.0f, 1.0f) / x_111);
i = 0;
{
@@ -72,18 +66,14 @@
switch(i) {
case 1: {
const float2 x_150 = tileID;
- const uint scalar_offset_4 = (88u) / 4;
- uint4 ubo_load_2 = x_20[scalar_offset_4 / 4];
- const float2 x_154 = asfloat(((scalar_offset_4 & 2) ? ubo_load_2.zw : ubo_load_2.xy));
+ const float2 x_154 = asfloat(x_20[5].zw);
const float4 x_156 = tileMapsTexture1.SampleBias(tileMapsSampler, ((x_150 + float2(0.5f, 0.5f)) / x_154), 0.0f);
frameID_1 = x_156.x;
break;
}
case 0: {
const float2 x_136 = tileID;
- const uint scalar_offset_5 = (88u) / 4;
- uint4 ubo_load_3 = x_20[scalar_offset_5 / 4];
- const float2 x_140 = asfloat(((scalar_offset_5 & 2) ? ubo_load_3.zw : ubo_load_3.xy));
+ const float2 x_140 = asfloat(x_20[5].zw);
const float4 x_142 = tileMapsTexture0.SampleBias(tileMapsSampler, ((x_136 + float2(0.5f, 0.5f)) / x_140), 0.0f);
frameID_1 = x_142.x;
break;
@@ -93,14 +83,12 @@
}
}
const float x_166 = frameID_1;
- const uint scalar_offset_6 = (108u) / 4;
- const float x_169 = asfloat(x_20[scalar_offset_6 / 4][scalar_offset_6 % 4]);
+ const float x_169 = asfloat(x_20[6].w);
const float4 x_172 = animationMapTexture.SampleBias(animationMapSampler, float2(((x_166 + 0.5f) / x_169), 0.0f), 0.0f);
animationData = x_172;
const float x_174 = animationData.y;
if ((x_174 > 0.0f)) {
- const uint scalar_offset_7 = (0u) / 4;
- const float x_181 = asfloat(x_20[scalar_offset_7 / 4][scalar_offset_7 % 4]);
+ const float x_181 = asfloat(x_20[0].x);
const float x_184 = animationData.z;
mt = ((x_181 * x_184) % 1.0f);
f = 0.0f;
@@ -113,8 +101,7 @@
break;
}
const float x_208 = frameID_1;
- const uint scalar_offset_8 = (108u) / 4;
- const float x_211 = asfloat(x_20[scalar_offset_8 / 4][scalar_offset_8 % 4]);
+ const float x_211 = asfloat(x_20[6].w);
const float4 x_217 = animationMapTexture.SampleBias(animationMapSampler, float2(((x_208 + 0.5f) / x_211), (0.125f * f)), 0.0f);
animationData = x_217;
}
@@ -124,9 +111,7 @@
const float4x4 x_225 = getFrameData_f1_(param);
frameData = x_225;
const float4 x_228 = frameData[0];
- const uint scalar_offset_9 = (96u) / 4;
- uint4 ubo_load_4 = x_20[scalar_offset_9 / 4];
- const float2 x_231 = asfloat(((scalar_offset_9 & 2) ? ubo_load_4.zw : ubo_load_4.xy));
+ const float2 x_231 = asfloat(x_20[6].xy);
frameSize = (float2(x_228.w, x_228.z) / x_231);
const float4 x_235 = frameData[0];
offset_1 = (float2(x_235.x, x_235.y) * sheetUnits);
@@ -156,8 +141,7 @@
}
}
}
- const uint scalar_offset_10 = (112u) / 4;
- const float3 x_310 = asfloat(x_20[scalar_offset_10 / 4].xyz);
+ const float3 x_310 = asfloat(x_20[7].xyz);
const float4 x_311 = color;
const float3 x_313 = (float3(x_311.x, x_311.y, x_311.z) * x_310);
color = float4(x_313.x, x_313.y, x_313.z, color.w);
diff --git a/test/bug/tint/949.wgsl.expected.hlsl b/test/bug/tint/949.wgsl.expected.hlsl
index 2bb720a..5aca5bf 100644
--- a/test/bug/tint/949.wgsl.expected.hlsl
+++ b/test/bug/tint/949.wgsl.expected.hlsl
@@ -200,17 +200,14 @@
const float4 x_262 = TextureSamplerTexture.Sample(TextureSamplerSampler, vMainuv);
tempTextureRead = x_262;
const float4 x_264 = tempTextureRead;
- const uint scalar_offset = (160u) / 4;
- const float x_273 = asfloat(x_269[scalar_offset / 4][scalar_offset % 4]);
+ const float x_273 = asfloat(x_269[10].x);
rgb = (float3(x_264.x, x_264.y, x_264.z) * x_273);
- const uint scalar_offset_1 = (144u) / 4;
- const float3 x_279 = asfloat(x_269[scalar_offset_1 / 4].xyz);
+ const float3 x_279 = asfloat(x_269[9].xyz);
const float4 x_282 = v_output1;
output5 = normalize((x_279 - float3(x_282.x, x_282.y, x_282.z)));
output4 = float4(0.0f, 0.0f, 0.0f, 0.0f);
uvOffset = float2(0.0f, 0.0f);
- const uint scalar_offset_2 = (128u) / 4;
- const float x_292 = asfloat(x_269[scalar_offset_2 / 4][scalar_offset_2 % 4]);
+ const float x_292 = asfloat(x_269[8].x);
normalScale = (1.0f / x_292);
if (gl_FrontFacing) {
x_299 = v_uv;
@@ -223,9 +220,7 @@
const float4 x_317 = v_output1;
param_4 = float3(x_317.x, x_317.y, x_317.z);
param_5 = TBNUV;
- const uint scalar_offset_3 = (168u) / 4;
- uint4 ubo_load = x_269[scalar_offset_3 / 4];
- const float2 x_324 = asfloat(((scalar_offset_3 & 2) ? ubo_load.zw : ubo_load.xy));
+ const float2 x_324 = asfloat(x_269[10].zw);
param_6 = x_324;
const float3x3 x_325 = cotangent_frame_vf3_vf3_vf2_vf2_(param_3, param_4, param_5, param_6);
TBN = x_325;
@@ -234,8 +229,7 @@
invTBN = x_329;
const float3 x_334 = mul(-(output5), invTBN);
parallaxLimit = (length(float2(x_334.x, x_334.y)) / mul(-(output5), invTBN).z);
- const uint scalar_offset_4 = (156u) / 4;
- const float x_345 = asfloat(x_269[scalar_offset_4 / 4][scalar_offset_4 % 4]);
+ const float x_345 = asfloat(x_269[9].w);
parallaxLimit = (parallaxLimit * x_345);
const float3 x_352 = mul(-(output5), invTBN);
vOffsetDir = normalize(float2(x_352.x, x_352.y));
@@ -270,8 +264,7 @@
parallaxOcclusion_0 = vCurrOffset;
uvOffset = parallaxOcclusion_0;
const float4 x_452 = TextureSamplerTexture.Sample(TextureSamplerSampler, (v_uv + uvOffset));
- const uint scalar_offset_5 = (128u) / 4;
- const float x_454 = asfloat(x_269[scalar_offset_5 / 4][scalar_offset_5 % 4]);
+ const float x_454 = asfloat(x_269[8].x);
param_8 = TBN;
param_9 = float3(x_452.x, x_452.y, x_452.z);
param_10 = (1.0f / x_454);
@@ -282,8 +275,7 @@
tempTextureRead1 = x_475;
const float4 x_477 = tempTextureRead1;
rgb1 = float3(x_477.x, x_477.y, x_477.z);
- const uint scalar_offset_6 = (144u) / 4;
- const float3 x_481 = asfloat(x_269[scalar_offset_6 / 4].xyz);
+ const float3 x_481 = asfloat(x_269[9].xyz);
const float4 x_482 = v_output1;
viewDirectionW_1 = normalize((x_481 - float3(x_482.x, x_482.y, x_482.z)));
shadow = 1.0f;
@@ -294,17 +286,13 @@
normalW = float3(x_494.x, x_494.y, x_494.z);
param_11 = viewDirectionW_1;
param_12 = normalW;
- const uint scalar_offset_7 = (0u) / 4;
- const float4 x_507 = asfloat(light0[scalar_offset_7 / 4]);
+ const float4 x_507 = asfloat(light0[0]);
param_13 = x_507;
- const uint scalar_offset_8 = (16u) / 4;
- const float4 x_510 = asfloat(light0[scalar_offset_8 / 4]);
+ const float4 x_510 = asfloat(light0[1]);
param_14 = float3(x_510.x, x_510.y, x_510.z);
- const uint scalar_offset_9 = (32u) / 4;
- const float4 x_514 = asfloat(light0[scalar_offset_9 / 4]);
+ const float4 x_514 = asfloat(light0[2]);
param_15 = float3(x_514.x, x_514.y, x_514.z);
- const uint scalar_offset_10 = (48u) / 4;
- const float3 x_518 = asfloat(light0[scalar_offset_10 / 4].xyz);
+ const float3 x_518 = asfloat(light0[3].xyz);
param_16 = x_518;
param_17 = glossiness_1;
const lightingInfo x_521 = computeHemisphericLighting_vf3_vf3_vf4_vf3_vf3_vf3_f1_(param_11, param_12, param_13, param_14, param_15, param_16, param_17);
diff --git a/test/bug/tint/951.spvasm.expected.hlsl b/test/bug/tint/951.spvasm.expected.hlsl
index 24ddc28..86a56b5 100644
--- a/test/bug/tint/951.spvasm.expected.hlsl
+++ b/test/bug/tint/951.spvasm.expected.hlsl
@@ -36,8 +36,7 @@
const uint x_61 = gl_GlobalInvocationID.x;
index = asint(x_61);
const int x_63 = index;
- const uint scalar_offset = (16u) / 4;
- const int x_70 = asint(x_24[scalar_offset / 4][scalar_offset % 4]);
+ const int x_70 = asint(x_24[1].x);
if ((x_63 < x_70)) {
const float x_75 = getAAtOutCoords_();
a_1 = x_75;
diff --git a/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.hlsl b/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.hlsl
index 45ec648..634b49e 100644
--- a/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.hlsl
+++ b/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.hlsl
@@ -4,7 +4,6 @@
[numthreads(1, 1, 1)]
void main() {
- const uint scalar_offset = (0u) / 4;
- const int use = (asint(v[scalar_offset / 4][scalar_offset % 4]) + 1);
+ const int use = (asint(v[0].x) + 1);
return;
}
diff --git a/test/samples/compute_boids.wgsl.expected.hlsl b/test/samples/compute_boids.wgsl.expected.hlsl
index 4643faa..3a88307 100644
--- a/test/samples/compute_boids.wgsl.expected.hlsl
+++ b/test/samples/compute_boids.wgsl.expected.hlsl
@@ -59,17 +59,14 @@
}
pos = asfloat(particlesA.Load2((16u * i))).xy;
vel = asfloat(particlesA.Load2(((16u * i) + 8u))).xy;
- const uint scalar_offset = (4u) / 4;
- if ((distance(pos, vPos) < asfloat(params[scalar_offset / 4][scalar_offset % 4]))) {
+ if ((distance(pos, vPos) < asfloat(params[0].y))) {
cMass = (cMass + pos);
cMassCount = (cMassCount + 1);
}
- const uint scalar_offset_1 = (8u) / 4;
- if ((distance(pos, vPos) < asfloat(params[scalar_offset_1 / 4][scalar_offset_1 % 4]))) {
+ if ((distance(pos, vPos) < asfloat(params[0].z))) {
colVel = (colVel - (pos - vPos));
}
- const uint scalar_offset_2 = (12u) / 4;
- if ((distance(pos, vPos) < asfloat(params[scalar_offset_2 / 4][scalar_offset_2 % 4]))) {
+ if ((distance(pos, vPos) < asfloat(params[0].w))) {
cVel = (cVel + vel);
cVelCount = (cVelCount + 1);
}
@@ -81,13 +78,9 @@
if ((cVelCount > 0)) {
cVel = (cVel / float2(float(cVelCount), float(cVelCount)));
}
- const uint scalar_offset_3 = (16u) / 4;
- const uint scalar_offset_4 = (20u) / 4;
- const uint scalar_offset_5 = (24u) / 4;
- vVel = (((vVel + (cMass * asfloat(params[scalar_offset_3 / 4][scalar_offset_3 % 4]))) + (colVel * asfloat(params[scalar_offset_4 / 4][scalar_offset_4 % 4]))) + (cVel * asfloat(params[scalar_offset_5 / 4][scalar_offset_5 % 4])));
+ vVel = (((vVel + (cMass * asfloat(params[1].x))) + (colVel * asfloat(params[1].y))) + (cVel * asfloat(params[1].z)));
vVel = (normalize(vVel) * clamp(length(vVel), 0.0f, 0.100000001f));
- const uint scalar_offset_6 = (0u) / 4;
- vPos = (vPos + (vVel * asfloat(params[scalar_offset_6 / 4][scalar_offset_6 % 4])));
+ vPos = (vPos + (vVel * asfloat(params[0].x)));
if ((vPos.x < -1.0f)) {
vPos.x = 1.0f;
}
