encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

tint/writer: Handle unnested, uniform matCx2 matrices 

Shuffle the transform orders to ensure that these are embedded in a structure before running the Std140 transform.

Add more end-to-end tests for these.

As pointed out in tint:1673, arrays of matrices are not correctly decomposed by the Std140 transform.
This will be addressed by a later change.

Bug: tint:1673
Change-Id: I47c93e458ff48578922d576819792e8ed3a5723c
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/102541
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Zhaoming Jiang <zhaoming.jiang@intel.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: David Neto <dneto@google.com>
This commit is contained in:
Ben Clayton 2022-09-17 00:38:19 +00:00 committed by Dawn LUCI CQ
parent 2bfd9f8c83
commit 84b43d61fa
406 changed files with 9658 additions and 2134 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

9
src/tint/writer/glsl/generator_impl.cc
View File

@ -221,10 +221,8 @@ SanitizedResult Sanitize(const Program* in,
manager.Add<transform::CanonicalizeEntryPointIO>(); manager.Add<transform::CanonicalizeEntryPointIO>();
manager.Add<transform::ExpandCompoundAssignment>(); manager.Add<transform::ExpandCompoundAssignment>();
manager.Add<transform::PromoteSideEffectsToDecl>(); manager.Add<transform::PromoteSideEffectsToDecl>();
manager.Add<transform::Std140>(); // Must come after PromoteSideEffectsToDecl
manager.Add<transform::PadStructs>(); manager.Add<transform::PadStructs>();
manager.Add<transform::UnwindDiscardFunctions>(); manager.Add<transform::UnwindDiscardFunctions>();
manager.Add<transform::SimplifyPointers>();
manager.Add<transform::RemovePhonies>(); manager.Add<transform::RemovePhonies>();
@ -245,6 +243,13 @@ SanitizedResult Sanitize(const Program* in,
manager.Add<transform::PromoteInitializersToLet>(); manager.Add<transform::PromoteInitializersToLet>();
manager.Add<transform::AddEmptyEntryPoint>(); manager.Add<transform::AddEmptyEntryPoint>();
manager.Add<transform::AddBlockAttribute>(); manager.Add<transform::AddBlockAttribute>();
// Std140 must come after PromoteSideEffectsToDecl and AddBlockAttribute
// Std140 must come before SimplifyPointers.
manager.Add<transform::Std140>();
manager.Add<transform::SimplifyPointers>();
data.Add<transform::CanonicalizeEntryPointIO::Config>( data.Add<transform::CanonicalizeEntryPointIO::Config>(
transform::CanonicalizeEntryPointIO::ShaderStyle::kGlsl); transform::CanonicalizeEntryPointIO::ShaderStyle::kGlsl);

12
src/tint/writer/spirv/generator_impl.cc
View File

@ -77,19 +77,25 @@ SanitizedResult Sanitize(const Program* in, const Options& options) {
manager.Add<transform::RemoveUnreachableStatements>(); manager.Add<transform::RemoveUnreachableStatements>();
manager.Add<transform::ExpandCompoundAssignment>(); manager.Add<transform::ExpandCompoundAssignment>();
manager.Add<transform::PromoteSideEffectsToDecl>(); manager.Add<transform::PromoteSideEffectsToDecl>();
manager.Add<transform::Std140>(); // Must come after PromoteSideEffectsToDecl
manager.Add<transform::UnwindDiscardFunctions>(); manager.Add<transform::UnwindDiscardFunctions>();
manager.Add<transform::SimplifyPointers>(); // Required for arrayLength() manager.Add<transform::SimplifyPointers>(); // Required for arrayLength()
manager.Add<transform::RemovePhonies>(); manager.Add<transform::RemovePhonies>();
manager.Add<transform::VectorizeScalarMatrixConstructors>(); manager.Add<transform::VectorizeScalarMatrixConstructors>();
manager.Add<transform::VectorizeMatrixConversions>(); manager.Add<transform::VectorizeMatrixConversions>();
manager.Add<transform::ForLoopToLoop>(); // Must come after manager.Add<transform::WhileToLoop>(); // ZeroInitWorkgroupMemory
manager.Add<transform::WhileToLoop>(); // ZeroInitWorkgroupMemory
manager.Add<transform::CanonicalizeEntryPointIO>(); manager.Add<transform::CanonicalizeEntryPointIO>();
manager.Add<transform::AddEmptyEntryPoint>(); manager.Add<transform::AddEmptyEntryPoint>();
manager.Add<transform::AddBlockAttribute>(); manager.Add<transform::AddBlockAttribute>();
// Std140 must come after PromoteSideEffectsToDecl, AddBlockAttribute
manager.Add<transform::Std140>();
// VarForDynamicIndex must come after Std140
manager.Add<transform::VarForDynamicIndex>(); manager.Add<transform::VarForDynamicIndex>();
// ForLoopToLoop must come after Std140, ZeroInitWorkgroupMemory
manager.Add<transform::ForLoopToLoop>();
data.Add<transform::CanonicalizeEntryPointIO::Config>( data.Add<transform::CanonicalizeEntryPointIO::Config>(
transform::CanonicalizeEntryPointIO::Config( transform::CanonicalizeEntryPointIO::Config(
transform::CanonicalizeEntryPointIO::ShaderStyle::kSpirv, 0xFFFFFFFF, transform::CanonicalizeEntryPointIO::ShaderStyle::kSpirv, 0xFFFFFFFF,

6
test/tint/buffer/uniform/dynamic_index/read.wgsl.expected.glsl
View File

@ -11,6 +11,8 @@ struct Inner {
ivec2 h; ivec2 h;
mat2x3 i; mat2x3 i;
mat3x2 j; mat3x2 j;
uint pad;
uint pad_1;
ivec4 k[4]; ivec4 k[4];
}; };
@ -32,10 +34,6 @@ struct Inner_std140 {
ivec4 k[4]; ivec4 k[4];
}; };
struct S {
Inner arr[8];
};
layout(binding = 0, std140) uniform S_std140_ubo { layout(binding = 0, std140) uniform S_std140_ubo {
Inner_std140 arr[8]; Inner_std140 arr[8];
} s; } s;

78
test/tint/buffer/uniform/dynamic_index/read.wgsl.expected.spvasm
View File

@ -1,7 +1,7 @@
; SPIR-V ; SPIR-V
; Version: 1.3 ; Version: 1.3
; Generator: Google Tint Compiler; 0 ; Generator: Google Tint Compiler; 0
; Bound: 86 ; Bound: 92
; Schema: 0 ; Schema: 0
OpCapability Shader OpCapability Shader
OpMemoryModel Logical GLSL450 OpMemoryModel Logical GLSL450
@ -76,12 +76,13 @@
%mat3v2float = OpTypeMatrix %v2float 3 %mat3v2float = OpTypeMatrix %v2float 3
%21 = OpTypeFunction %mat3v2float %uint %21 = OpTypeFunction %mat3v2float %uint
%uint_0 = OpConstant %uint 0 %uint_0 = OpConstant %uint 0
%_ptr_Uniform_Inner_std140 = OpTypePointer Uniform %Inner_std140
%uint_9 = OpConstant %uint 9 %uint_9 = OpConstant %uint 9
%_ptr_Uniform_v2float = OpTypePointer Uniform %v2float %_ptr_Uniform_v2float = OpTypePointer Uniform %v2float
%uint_10 = OpConstant %uint 10 %uint_10 = OpConstant %uint 10
%uint_11 = OpConstant %uint 11 %uint_11 = OpConstant %uint 11
%void = OpTypeVoid %void = OpTypeVoid
%38 = OpTypeFunction %void %uint %44 = OpTypeFunction %void %uint
%_ptr_Uniform_v3int = OpTypePointer Uniform %v3int %_ptr_Uniform_v3int = OpTypePointer Uniform %v3int
%uint_1 = OpConstant %uint 1 %uint_1 = OpConstant %uint 1
%_ptr_Uniform_int = OpTypePointer Uniform %int %_ptr_Uniform_int = OpTypePointer Uniform %int
@ -98,48 +99,49 @@
%_ptr_Uniform_mat2v3float = OpTypePointer Uniform %mat2v3float %_ptr_Uniform_mat2v3float = OpTypePointer Uniform %mat2v3float
%uint_12 = OpConstant %uint 12 %uint_12 = OpConstant %uint 12
%_ptr_Uniform__arr_v4int_uint_4 = OpTypePointer Uniform %_arr_v4int_uint_4 %_ptr_Uniform__arr_v4int_uint_4 = OpTypePointer Uniform %_arr_v4int_uint_4
%81 = OpTypeFunction %void %87 = OpTypeFunction %void
%load_s_arr_p0_j = OpFunction %mat3v2float None %21 %load_s_arr_p0_j = OpFunction %mat3v2float None %21
%p0 = OpFunctionParameter %uint %p0 = OpFunctionParameter %uint
%25 = OpLabel %25 = OpLabel
%29 = OpAccessChain %_ptr_Uniform_v2float %s %uint_0 %p0 %uint_9 %29 = OpAccessChain %_ptr_Uniform_Inner_std140 %s %uint_0 %p0
%30 = OpLoad %v2float %29 %33 = OpAccessChain %_ptr_Uniform_v2float %29 %uint_9
%32 = OpAccessChain %_ptr_Uniform_v2float %s %uint_0 %p0 %uint_10 %34 = OpLoad %v2float %33
%33 = OpLoad %v2float %32 %37 = OpAccessChain %_ptr_Uniform_v2float %29 %uint_10
%35 = OpAccessChain %_ptr_Uniform_v2float %s %uint_0 %p0 %uint_11 %38 = OpLoad %v2float %37
%36 = OpLoad %v2float %35 %41 = OpAccessChain %_ptr_Uniform_v2float %29 %uint_11
%37 = OpCompositeConstruct %mat3v2float %30 %33 %36 %42 = OpLoad %v2float %41
OpReturnValue %37 %43 = OpCompositeConstruct %mat3v2float %34 %38 %42
OpReturnValue %43
OpFunctionEnd OpFunctionEnd
%main_inner = OpFunction %void None %38 %main_inner = OpFunction %void None %44
%idx = OpFunctionParameter %uint %idx = OpFunctionParameter %uint
%42 = OpLabel %48 = OpLabel
%44 = OpAccessChain %_ptr_Uniform_v3int %s %uint_0 %idx %uint_0 %50 = OpAccessChain %_ptr_Uniform_v3int %s %uint_0 %idx %uint_0
%45 = OpLoad %v3int %44 %51 = OpLoad %v3int %50
%48 = OpAccessChain %_ptr_Uniform_int %s %uint_0 %idx %uint_1 %54 = OpAccessChain %_ptr_Uniform_int %s %uint_0 %idx %uint_1
%49 = OpLoad %int %48 %55 = OpLoad %int %54
%52 = OpAccessChain %_ptr_Uniform_v3uint %s %uint_0 %idx %uint_2 %58 = OpAccessChain %_ptr_Uniform_v3uint %s %uint_0 %idx %uint_2
%53 = OpLoad %v3uint %52 %59 = OpLoad %v3uint %58
%56 = OpAccessChain %_ptr_Uniform_uint %s %uint_0 %idx %uint_3 %62 = OpAccessChain %_ptr_Uniform_uint %s %uint_0 %idx %uint_3
%57 = OpLoad %uint %56 %63 = OpLoad %uint %62
%59 = OpAccessChain %_ptr_Uniform_v3float %s %uint_0 %idx %uint_4 %65 = OpAccessChain %_ptr_Uniform_v3float %s %uint_0 %idx %uint_4
%60 = OpLoad %v3float %59 %66 = OpLoad %v3float %65
%63 = OpAccessChain %_ptr_Uniform_float %s %uint_0 %idx %uint_5 %69 = OpAccessChain %_ptr_Uniform_float %s %uint_0 %idx %uint_5
%64 = OpLoad %float %63 %70 = OpLoad %float %69
%67 = OpAccessChain %_ptr_Uniform_v2int %s %uint_0 %idx %uint_6 %73 = OpAccessChain %_ptr_Uniform_v2int %s %uint_0 %idx %uint_6
%68 = OpLoad %v2int %67 %74 = OpLoad %v2int %73
%70 = OpAccessChain %_ptr_Uniform_v2int %s %uint_0 %idx %uint_7 %76 = OpAccessChain %_ptr_Uniform_v2int %s %uint_0 %idx %uint_7
%71 = OpLoad %v2int %70 %77 = OpLoad %v2int %76
%73 = OpAccessChain %_ptr_Uniform_mat2v3float %s %uint_0 %idx %uint_8 %79 = OpAccessChain %_ptr_Uniform_mat2v3float %s %uint_0 %idx %uint_8
%74 = OpLoad %mat2v3float %73 %80 = OpLoad %mat2v3float %79
%75 = OpFunctionCall %mat3v2float %load_s_arr_p0_j %idx %81 = OpFunctionCall %mat3v2float %load_s_arr_p0_j %idx
%79 = OpAccessChain %_ptr_Uniform__arr_v4int_uint_4 %s %uint_0 %idx %uint_12 %85 = OpAccessChain %_ptr_Uniform__arr_v4int_uint_4 %s %uint_0 %idx %uint_12
%80 = OpLoad %_arr_v4int_uint_4 %79 %86 = OpLoad %_arr_v4int_uint_4 %85
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd
%main = OpFunction %void None %81 %main = OpFunction %void None %87
%83 = OpLabel %89 = OpLabel
%85 = OpLoad %uint %idx_1 %91 = OpLoad %uint %idx_1
%84 = OpFunctionCall %void %main_inner %85 %90 = OpFunctionCall %void %main_inner %91
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

15
test/tint/buffer/uniform/static_index/read.wgsl.expected.glsl
View File

@ -7,21 +7,6 @@ struct Inner {
uint pad_2; uint pad_2;
}; };
struct S {
ivec3 a;
int b;
uvec3 c;
uint d;
vec3 e;
float f;
ivec2 g;
ivec2 h;
mat2x3 i;
mat3x2 j;
Inner k;
Inner l[4];
};
layout(binding = 0, std140) uniform S_std140_ubo { layout(binding = 0, std140) uniform S_std140_ubo {
ivec3 a; ivec3 a;
int b; int b;

68
test/tint/buffer/uniform/static_index/read.wgsl.expected.spvasm
View File

@ -1,7 +1,7 @@
; SPIR-V ; SPIR-V
; Version: 1.3 ; Version: 1.3
; Generator: Google Tint Compiler; 0 ; Generator: Google Tint Compiler; 0
; Bound: 78 ; Bound: 82
; Schema: 0 ; Schema: 0
OpCapability Shader OpCapability Shader
OpMemoryModel Logical GLSL450 OpMemoryModel Logical GLSL450
@ -71,7 +71,7 @@
%uint_10 = OpConstant %uint 10 %uint_10 = OpConstant %uint 10
%uint_11 = OpConstant %uint 11 %uint_11 = OpConstant %uint 11
%void = OpTypeVoid %void = OpTypeVoid
%31 = OpTypeFunction %void %35 = OpTypeFunction %void
%uint_0 = OpConstant %uint 0 %uint_0 = OpConstant %uint 0
%_ptr_Uniform_v3int = OpTypePointer Uniform %v3int %_ptr_Uniform_v3int = OpTypePointer Uniform %v3int
%uint_1 = OpConstant %uint 1 %uint_1 = OpConstant %uint 1
@ -94,39 +94,39 @@
%_ptr_Uniform__arr_Inner_uint_4 = OpTypePointer Uniform %_arr_Inner_uint_4 %_ptr_Uniform__arr_Inner_uint_4 = OpTypePointer Uniform %_arr_Inner_uint_4
%load_s_j = OpFunction %mat3v2float None %16 %load_s_j = OpFunction %mat3v2float None %16
%19 = OpLabel %19 = OpLabel
%22 = OpAccessChain %_ptr_Uniform_v2float %s %uint_9 %24 = OpAccessChain %_ptr_Uniform_v2float %s %uint_9
%23 = OpLoad %v2float %22 %25 = OpLoad %v2float %24
%25 = OpAccessChain %_ptr_Uniform_v2float %s %uint_10 %28 = OpAccessChain %_ptr_Uniform_v2float %s %uint_10
%26 = OpLoad %v2float %25
%28 = OpAccessChain %_ptr_Uniform_v2float %s %uint_11
%29 = OpLoad %v2float %28 %29 = OpLoad %v2float %28
%30 = OpCompositeConstruct %mat3v2float %23 %26 %29 %32 = OpAccessChain %_ptr_Uniform_v2float %s %uint_11
OpReturnValue %30 %33 = OpLoad %v2float %32
%34 = OpCompositeConstruct %mat3v2float %25 %29 %33
OpReturnValue %34
OpFunctionEnd OpFunctionEnd
%main = OpFunction %void None %31 %main = OpFunction %void None %35
%34 = OpLabel %38 = OpLabel
%37 = OpAccessChain %_ptr_Uniform_v3int %s %uint_0 %41 = OpAccessChain %_ptr_Uniform_v3int %s %uint_0
%38 = OpLoad %v3int %37 %42 = OpLoad %v3int %41
%41 = OpAccessChain %_ptr_Uniform_int %s %uint_1 %45 = OpAccessChain %_ptr_Uniform_int %s %uint_1
%42 = OpLoad %int %41 %46 = OpLoad %int %45
%45 = OpAccessChain %_ptr_Uniform_v3uint %s %uint_2 %49 = OpAccessChain %_ptr_Uniform_v3uint %s %uint_2
%46 = OpLoad %v3uint %45 %50 = OpLoad %v3uint %49
%49 = OpAccessChain %_ptr_Uniform_uint %s %uint_3 %53 = OpAccessChain %_ptr_Uniform_uint %s %uint_3
%50 = OpLoad %uint %49 %54 = OpLoad %uint %53
%52 = OpAccessChain %_ptr_Uniform_v3float %s %uint_4 %56 = OpAccessChain %_ptr_Uniform_v3float %s %uint_4
%53 = OpLoad %v3float %52 %57 = OpLoad %v3float %56
%56 = OpAccessChain %_ptr_Uniform_float %s %uint_5 %60 = OpAccessChain %_ptr_Uniform_float %s %uint_5
%57 = OpLoad %float %56 %61 = OpLoad %float %60
%60 = OpAccessChain %_ptr_Uniform_v2int %s %uint_6 %64 = OpAccessChain %_ptr_Uniform_v2int %s %uint_6
%61 = OpLoad %v2int %60 %65 = OpLoad %v2int %64
%63 = OpAccessChain %_ptr_Uniform_v2int %s %uint_7 %67 = OpAccessChain %_ptr_Uniform_v2int %s %uint_7
%64 = OpLoad %v2int %63 %68 = OpLoad %v2int %67
%67 = OpAccessChain %_ptr_Uniform_mat2v3float %s %uint_8 %71 = OpAccessChain %_ptr_Uniform_mat2v3float %s %uint_8
%68 = OpLoad %mat2v3float %67 %72 = OpLoad %mat2v3float %71
%69 = OpFunctionCall %mat3v2float %load_s_j %73 = OpFunctionCall %mat3v2float %load_s_j
%72 = OpAccessChain %_ptr_Uniform_Inner %s %uint_12 %76 = OpAccessChain %_ptr_Uniform_Inner %s %uint_12
%73 = OpLoad %Inner %72 %77 = OpLoad %Inner %76
%76 = OpAccessChain %_ptr_Uniform__arr_Inner_uint_4 %s %uint_13 %80 = OpAccessChain %_ptr_Uniform__arr_Inner_uint_4 %s %uint_13
%77 = OpLoad %_arr_Inner_uint_4 %76 %81 = OpLoad %_arr_Inner_uint_4 %80
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

15
test/tint/buffer/uniform/std140/array/mat2x2/dynamic_index_via_ptr.wgsl Normal file
View File

@ -0,0 +1,15 @@
@group(0) @binding(0) var<uniform> a : array<mat2x2<f32>, 4>;
var<private> counter = 0;
fn i() -> i32 { counter++; return counter; }
@compute @workgroup_size(1)
fn f() {
let p_a = &a;
let p_a_i = &((*p_a)[i()]);
let p_a_i_i = &((*p_a_i)[i()]);
let l_a : array<mat2x2<f32>, 4> = *p_a;
let l_a_i : mat2x2<f32> = *p_a_i;
let l_a_i_i : vec2<f32> = *p_a_i_i;
}

40
test/tint/buffer/uniform/std140/array/mat2x2/dynamic_index_via_ptr.wgsl.expected.dxc.hlsl Normal file
View File

@ -0,0 +1,40 @@
cbuffer cbuffer_a : register(b0, space0) {
uint4 a[4];
};
static int counter = 0;
int i() {
counter = (counter + 1);
return counter;
}
float2x2 tint_symbol_1(uint4 buffer[4], uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
uint4 ubo_load = buffer[scalar_offset / 4];
const uint scalar_offset_1 = ((offset + 8u)) / 4;
uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
return float2x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)));
}
typedef float2x2 tint_symbol_ret[4];
tint_symbol_ret tint_symbol(uint4 buffer[4], uint offset) {
float2x2 arr[4] = (float2x2[4])0;
{
[loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
arr[i_1] = tint_symbol_1(buffer, (offset + (i_1 * 16u)));
}
}
return arr;
}
[numthreads(1, 1, 1)]
void f() {
const int p_a_i_save = i();
const int p_a_i_i_save = i();
const float2x2 l_a[4] = tint_symbol(a, 0u);
const float2x2 l_a_i = tint_symbol_1(a, (16u * uint(p_a_i_save)));
const uint scalar_offset_2 = (((16u * uint(p_a_i_save)) + (8u * uint(p_a_i_i_save)))) / 4;
uint4 ubo_load_2 = a[scalar_offset_2 / 4];
const float2 l_a_i_i = asfloat(((scalar_offset_2 & 2) ? ubo_load_2.zw : ubo_load_2.xy));
return;
}

40
test/tint/buffer/uniform/std140/array/mat2x2/dynamic_index_via_ptr.wgsl.expected.fxc.hlsl Normal file
View File

@ -0,0 +1,40 @@
cbuffer cbuffer_a : register(b0, space0) {
uint4 a[4];
};
static int counter = 0;
int i() {
counter = (counter + 1);
return counter;
}
float2x2 tint_symbol_1(uint4 buffer[4], uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
uint4 ubo_load = buffer[scalar_offset / 4];
const uint scalar_offset_1 = ((offset + 8u)) / 4;
uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
return float2x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)));
}
typedef float2x2 tint_symbol_ret[4];
tint_symbol_ret tint_symbol(uint4 buffer[4], uint offset) {
float2x2 arr[4] = (float2x2[4])0;
{
[loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
arr[i_1] = tint_symbol_1(buffer, (offset + (i_1 * 16u)));
}
}
return arr;
}
[numthreads(1, 1, 1)]
void f() {
const int p_a_i_save = i();
const int p_a_i_i_save = i();
const float2x2 l_a[4] = tint_symbol(a, 0u);
const float2x2 l_a_i = tint_symbol_1(a, (16u * uint(p_a_i_save)));
const uint scalar_offset_2 = (((16u * uint(p_a_i_save)) + (8u * uint(p_a_i_i_save)))) / 4;
uint4 ubo_load_2 = a[scalar_offset_2 / 4];
const float2 l_a_i_i = asfloat(((scalar_offset_2 & 2) ? ubo_load_2.zw : ubo_load_2.xy));
return;
}

27
test/tint/buffer/uniform/std140/array/mat2x2/dynamic_index_via_ptr.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,27 @@
#version 310 es
layout(binding = 0, std140) uniform a_block_ubo {
mat2 inner[4];
} a;
int counter = 0;
int i() {
counter = (counter + 1);
return counter;
}
void f() {
int tint_symbol = i();
int p_a_i_save = tint_symbol;
int tint_symbol_1 = i();
int p_a_i_i_save = tint_symbol_1;
mat2 l_a[4] = a.inner;
mat2 l_a_i = a.inner[p_a_i_save];
vec2 l_a_i_i = a.inner[p_a_i_save][p_a_i_i_save];
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
f();
return;
}

33
test/tint/buffer/uniform/std140/array/mat2x2/dynamic_index_via_ptr.wgsl.expected.msl Normal file
View File

@ -0,0 +1,33 @@
#include <metal_stdlib>
using namespace metal;
template<typename T, size_t N>
struct tint_array {
const constant T& operator[](size_t i) const constant { return elements[i]; }
device T& operator[](size_t i) device { return elements[i]; }
const device T& operator[](size_t i) const device { return elements[i]; }
thread T& operator[](size_t i) thread { return elements[i]; }
const thread T& operator[](size_t i) const thread { return elements[i]; }
threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
T elements[N];
};
int i() {
thread int tint_symbol_2 = 0;
tint_symbol_2 = as_type<int>((as_type<uint>(tint_symbol_2) + as_type<uint>(1)));
return tint_symbol_2;
}
kernel void f(const constant tint_array<float2x2, 4>* tint_symbol_3 [[buffer(0)]]) {
int const tint_symbol = i();
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const p_a_i_i_save = tint_symbol_1;
tint_array<float2x2, 4> const l_a = *(tint_symbol_3);
float2x2 const l_a_i = (*(tint_symbol_3))[p_a_i_save];
float2 const l_a_i_i = (*(tint_symbol_3))[p_a_i_save][p_a_i_i_save];
return;
}

64
test/tint/buffer/uniform/std140/array/mat2x2/dynamic_index_via_ptr.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,64 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 37
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %a_block "a_block"
OpMemberName %a_block 0 "inner"
OpName %a "a"
OpName %counter "counter"
OpName %i "i"
OpName %f "f"
OpDecorate %a_block Block
OpMemberDecorate %a_block 0 Offset 0
OpMemberDecorate %a_block 0 ColMajor
OpMemberDecorate %a_block 0 MatrixStride 8
OpDecorate %_arr_mat2v2float_uint_4 ArrayStride 16
OpDecorate %a NonWritable
OpDecorate %a DescriptorSet 0
OpDecorate %a Binding 0
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%mat2v2float = OpTypeMatrix %v2float 2
%uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4
%_arr_mat2v2float_uint_4 = OpTypeArray %mat2v2float %uint_4
%a_block = OpTypeStruct %_arr_mat2v2float_uint_4
%_ptr_Uniform_a_block = OpTypePointer Uniform %a_block
%a = OpVariable %_ptr_Uniform_a_block Uniform
%int = OpTypeInt 32 1
%11 = OpConstantNull %int
%_ptr_Private_int = OpTypePointer Private %int
%counter = OpVariable %_ptr_Private_int Private %11
%14 = OpTypeFunction %int
%int_1 = OpConstant %int 1
%void = OpTypeVoid
%21 = OpTypeFunction %void
%uint_0 = OpConstant %uint 0
%_ptr_Uniform__arr_mat2v2float_uint_4 = OpTypePointer Uniform %_arr_mat2v2float_uint_4
%_ptr_Uniform_mat2v2float = OpTypePointer Uniform %mat2v2float
%_ptr_Uniform_v2float = OpTypePointer Uniform %v2float
%i = OpFunction %int None %14
%16 = OpLabel
%17 = OpLoad %int %counter
%19 = OpIAdd %int %17 %int_1
OpStore %counter %19
%20 = OpLoad %int %counter
OpReturnValue %20
OpFunctionEnd
%f = OpFunction %void None %21
%24 = OpLabel
%25 = OpFunctionCall %int %i
%26 = OpFunctionCall %int %i
%29 = OpAccessChain %_ptr_Uniform__arr_mat2v2float_uint_4 %a %uint_0
%30 = OpLoad %_arr_mat2v2float_uint_4 %29
%32 = OpAccessChain %_ptr_Uniform_mat2v2float %a %uint_0 %25
%33 = OpLoad %mat2v2float %32
%35 = OpAccessChain %_ptr_Uniform_v2float %a %uint_0 %25 %26
%36 = OpLoad %v2float %35
OpReturn
OpFunctionEnd

18
test/tint/buffer/uniform/std140/array/mat2x2/dynamic_index_via_ptr.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,18 @@
@group(0) @binding(0) var<uniform> a : array<mat2x2<f32>, 4>;
var<private> counter = 0;
fn i() -> i32 {
counter++;
return counter;
}
@compute @workgroup_size(1)
fn f() {
let p_a = &(a);
let p_a_i = &((*(p_a))[i()]);
let p_a_i_i = &((*(p_a_i))[i()]);
let l_a : array<mat2x2<f32>, 4> = *(p_a);
let l_a_i : mat2x2<f32> = *(p_a_i);
let l_a_i_i : vec2<f32> = *(p_a_i_i);
}

12
test/tint/buffer/uniform/std140/array/mat2x2/static_index_via_ptr.wgsl Normal file
View File

@ -0,0 +1,12 @@
@group(0) @binding(0) var<uniform> a : array<mat2x2<f32>, 4>;
@compute @workgroup_size(1)
fn f() {
let p_a = &a;
let p_a_2 = &((*p_a)[2]);
let p_a_2_1 = &((*p_a_2)[1]);
let l_a : array<mat2x2<f32>, 4> = *p_a;
let l_a_i : mat2x2<f32> = *p_a_2;
let l_a_i_i : vec2<f32> = *p_a_2_1;
}

30
test/tint/buffer/uniform/std140/array/mat2x2/static_index_via_ptr.wgsl.expected.dxc.hlsl Normal file
View File

@ -0,0 +1,30 @@
cbuffer cbuffer_a : register(b0, space0) {
uint4 a[4];
};
float2x2 tint_symbol_1(uint4 buffer[4], uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
uint4 ubo_load = buffer[scalar_offset / 4];
const uint scalar_offset_1 = ((offset + 8u)) / 4;
uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
return float2x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)));
}
typedef float2x2 tint_symbol_ret[4];
tint_symbol_ret tint_symbol(uint4 buffer[4], uint offset) {
float2x2 arr[4] = (float2x2[4])0;
{
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
arr[i] = tint_symbol_1(buffer, (offset + (i * 16u)));
}
}
return arr;
}
[numthreads(1, 1, 1)]
void f() {
const float2x2 l_a[4] = tint_symbol(a, 0u);
const float2x2 l_a_i = tint_symbol_1(a, 32u);
const float2 l_a_i_i = asfloat(a[2].zw);
return;
}

30
test/tint/buffer/uniform/std140/array/mat2x2/static_index_via_ptr.wgsl.expected.fxc.hlsl Normal file
View File

@ -0,0 +1,30 @@
cbuffer cbuffer_a : register(b0, space0) {
uint4 a[4];
};
float2x2 tint_symbol_1(uint4 buffer[4], uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
uint4 ubo_load = buffer[scalar_offset / 4];
const uint scalar_offset_1 = ((offset + 8u)) / 4;
uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
return float2x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)));
}
typedef float2x2 tint_symbol_ret[4];
tint_symbol_ret tint_symbol(uint4 buffer[4], uint offset) {
float2x2 arr[4] = (float2x2[4])0;
{
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
arr[i] = tint_symbol_1(buffer, (offset + (i * 16u)));
}
}
return arr;
}
[numthreads(1, 1, 1)]
void f() {
const float2x2 l_a[4] = tint_symbol(a, 0u);
const float2x2 l_a_i = tint_symbol_1(a, 32u);
const float2 l_a_i_i = asfloat(a[2].zw);
return;
}

17
test/tint/buffer/uniform/std140/array/mat2x2/static_index_via_ptr.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,17 @@
#version 310 es
layout(binding = 0, std140) uniform a_block_ubo {
mat2 inner[4];
} a;
void f() {
mat2 l_a[4] = a.inner;
mat2 l_a_i = a.inner[2];
vec2 l_a_i_i = a.inner[2][1];
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
f();
return;
}

23
test/tint/buffer/uniform/std140/array/mat2x2/static_index_via_ptr.wgsl.expected.msl Normal file
View File

@ -0,0 +1,23 @@
#include <metal_stdlib>
using namespace metal;
template<typename T, size_t N>
struct tint_array {
const constant T& operator[](size_t i) const constant { return elements[i]; }
device T& operator[](size_t i) device { return elements[i]; }
const device T& operator[](size_t i) const device { return elements[i]; }
thread T& operator[](size_t i) thread { return elements[i]; }
const thread T& operator[](size_t i) const thread { return elements[i]; }
threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
T elements[N];
};
kernel void f(const constant tint_array<float2x2, 4>* tint_symbol [[buffer(0)]]) {
tint_array<float2x2, 4> const l_a = *(tint_symbol);
float2x2 const l_a_i = (*(tint_symbol))[2];
float2 const l_a_i_i = (*(tint_symbol))[2][1];
return;
}

49
test/tint/buffer/uniform/std140/array/mat2x2/static_index_via_ptr.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,49 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 27
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %a_block "a_block"
OpMemberName %a_block 0 "inner"
OpName %a "a"
OpName %f "f"
OpDecorate %a_block Block
OpMemberDecorate %a_block 0 Offset 0
OpMemberDecorate %a_block 0 ColMajor
OpMemberDecorate %a_block 0 MatrixStride 8
OpDecorate %_arr_mat2v2float_uint_4 ArrayStride 16
OpDecorate %a NonWritable
OpDecorate %a DescriptorSet 0
OpDecorate %a Binding 0
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%mat2v2float = OpTypeMatrix %v2float 2
%uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4
%_arr_mat2v2float_uint_4 = OpTypeArray %mat2v2float %uint_4
%a_block = OpTypeStruct %_arr_mat2v2float_uint_4
%_ptr_Uniform_a_block = OpTypePointer Uniform %a_block
%a = OpVariable %_ptr_Uniform_a_block Uniform
%void = OpTypeVoid
%10 = OpTypeFunction %void
%uint_0 = OpConstant %uint 0
%_ptr_Uniform__arr_mat2v2float_uint_4 = OpTypePointer Uniform %_arr_mat2v2float_uint_4
%int = OpTypeInt 32 1
%int_2 = OpConstant %int 2
%_ptr_Uniform_mat2v2float = OpTypePointer Uniform %mat2v2float
%int_1 = OpConstant %int 1
%_ptr_Uniform_v2float = OpTypePointer Uniform %v2float
%f = OpFunction %void None %10
%13 = OpLabel
%16 = OpAccessChain %_ptr_Uniform__arr_mat2v2float_uint_4 %a %uint_0
%17 = OpLoad %_arr_mat2v2float_uint_4 %16
%21 = OpAccessChain %_ptr_Uniform_mat2v2float %a %uint_0 %int_2
%22 = OpLoad %mat2v2float %21
%25 = OpAccessChain %_ptr_Uniform_v2float %a %uint_0 %int_2 %int_1
%26 = OpLoad %v2float %25
OpReturn
OpFunctionEnd

11
test/tint/buffer/uniform/std140/array/mat2x2/static_index_via_ptr.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,11 @@
@group(0) @binding(0) var<uniform> a : array<mat2x2<f32>, 4>;
@compute @workgroup_size(1)
fn f() {
let p_a = &(a);
let p_a_2 = &((*(p_a))[2]);
let p_a_2_1 = &((*(p_a_2))[1]);
let l_a : array<mat2x2<f32>, 4> = *(p_a);
let l_a_i : mat2x2<f32> = *(p_a_2);
let l_a_i_i : vec2<f32> = *(p_a_2_1);
}

8
test/tint/buffer/uniform/std140/array/mat2x2/to_builtin.wgsl Normal file
View File

@ -0,0 +1,8 @@
@group(0) @binding(0) var<uniform> u : array<mat2x2<f32>, 4>;
@compute @workgroup_size(1)
fn f() {
let t = transpose(u[2]);
let l = length(u[0][1].yx);
let a = abs(u[0][1].yx.x);
}

19
test/tint/buffer/uniform/std140/array/mat2x2/to_builtin.wgsl.expected.dxc.hlsl Normal file
View File

@ -0,0 +1,19 @@
cbuffer cbuffer_u : register(b0, space0) {
uint4 u[4];
};
float2x2 tint_symbol(uint4 buffer[4], uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
uint4 ubo_load = buffer[scalar_offset / 4];
const uint scalar_offset_1 = ((offset + 8u)) / 4;
uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
return float2x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)));
}
[numthreads(1, 1, 1)]
void f() {
const float2x2 t = transpose(tint_symbol(u, 32u));
const float l = length(asfloat(u[0].zw).yx);
const float a = abs(asfloat(u[0].zw).yx.x);
return;
}

19
test/tint/buffer/uniform/std140/array/mat2x2/to_builtin.wgsl.expected.fxc.hlsl Normal file
View File

@ -0,0 +1,19 @@
cbuffer cbuffer_u : register(b0, space0) {
uint4 u[4];
};
float2x2 tint_symbol(uint4 buffer[4], uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
uint4 ubo_load = buffer[scalar_offset / 4];
const uint scalar_offset_1 = ((offset + 8u)) / 4;
uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
return float2x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)));
}
[numthreads(1, 1, 1)]
void f() {
const float2x2 t = transpose(tint_symbol(u, 32u));
const float l = length(asfloat(u[0].zw).yx);
const float a = abs(asfloat(u[0].zw).yx.x);
return;
}

17
test/tint/buffer/uniform/std140/array/mat2x2/to_builtin.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,17 @@
#version 310 es
layout(binding = 0, std140) uniform u_block_ubo {
mat2 inner[4];
} u;
void f() {
mat2 t = transpose(u.inner[2]);
float l = length(u.inner[0][1].yx);
float a = abs(u.inner[0][1].yx.x);
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
f();
return;
}

23
test/tint/buffer/uniform/std140/array/mat2x2/to_builtin.wgsl.expected.msl Normal file
View File

@ -0,0 +1,23 @@
#include <metal_stdlib>
using namespace metal;
template<typename T, size_t N>
struct tint_array {
const constant T& operator[](size_t i) const constant { return elements[i]; }
device T& operator[](size_t i) device { return elements[i]; }
const device T& operator[](size_t i) const device { return elements[i]; }
thread T& operator[](size_t i) thread { return elements[i]; }
const thread T& operator[](size_t i) const thread { return elements[i]; }
threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
T elements[N];
};
kernel void f(const constant tint_array<float2x2, 4>* tint_symbol [[buffer(0)]]) {
float2x2 const t = transpose((*(tint_symbol))[2]);
float const l = length(float2((*(tint_symbol))[0][1]).yx);
float const a = fabs(float2((*(tint_symbol))[0][1]).yx[0]);
return;
}

56
test/tint/buffer/uniform/std140/array/mat2x2/to_builtin.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,56 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 34
; Schema: 0
OpCapability Shader
%22 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %u_block "u_block"
OpMemberName %u_block 0 "inner"
OpName %u "u"
OpName %f "f"
OpDecorate %u_block Block
OpMemberDecorate %u_block 0 Offset 0
OpMemberDecorate %u_block 0 ColMajor
OpMemberDecorate %u_block 0 MatrixStride 8
OpDecorate %_arr_mat2v2float_uint_4 ArrayStride 16
OpDecorate %u NonWritable
OpDecorate %u DescriptorSet 0
OpDecorate %u Binding 0
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%mat2v2float = OpTypeMatrix %v2float 2
%uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4
%_arr_mat2v2float_uint_4 = OpTypeArray %mat2v2float %uint_4
%u_block = OpTypeStruct %_arr_mat2v2float_uint_4
%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
%u = OpVariable %_ptr_Uniform_u_block Uniform
%void = OpTypeVoid
%10 = OpTypeFunction %void
%uint_0 = OpConstant %uint 0
%int = OpTypeInt 32 1
%int_2 = OpConstant %int 2
%_ptr_Uniform_mat2v2float = OpTypePointer Uniform %mat2v2float
%23 = OpConstantNull %int
%int_1 = OpConstant %int 1
%_ptr_Uniform_v2float = OpTypePointer Uniform %v2float
%f = OpFunction %void None %10
%13 = OpLabel
%19 = OpAccessChain %_ptr_Uniform_mat2v2float %u %uint_0 %int_2
%20 = OpLoad %mat2v2float %19
%14 = OpTranspose %mat2v2float %20
%26 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %23 %int_1
%27 = OpLoad %v2float %26
%28 = OpVectorShuffle %v2float %27 %27 1 0
%21 = OpExtInst %float %22 Length %28
%30 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %23 %int_1
%31 = OpLoad %v2float %30
%32 = OpVectorShuffle %v2float %31 %31 1 0
%33 = OpCompositeExtract %float %32 0
%29 = OpExtInst %float %22 FAbs %33
OpReturn
OpFunctionEnd

8
test/tint/buffer/uniform/std140/array/mat2x2/to_builtin.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,8 @@
@group(0) @binding(0) var<uniform> u : array<mat2x2<f32>, 4>;
@compute @workgroup_size(1)
fn f() {
let t = transpose(u[2]);
let l = length(u[0][1].yx);
let a = abs(u[0][1].yx.x);
}

14
test/tint/buffer/uniform/std140/array/mat2x2/to_fn.wgsl Normal file
View File

@ -0,0 +1,14 @@
@group(0) @binding(0) var<uniform> u : array<mat2x2<f32>, 4>;
fn a(a : array<mat2x2<f32>, 4>) {}
fn b(m : mat2x2<f32>) {}
fn c(v : vec2<f32>) {}
fn d(f : f32) {}
@compute @workgroup_size(1)
fn f() {
a(u);
b(u[1]);
c(u[1][0].yx);
d(u[1][0].yx.x);
}

43
test/tint/buffer/uniform/std140/array/mat2x2/to_fn.wgsl.expected.dxc.hlsl Normal file
View File

@ -0,0 +1,43 @@
cbuffer cbuffer_u : register(b0, space0) {
uint4 u[4];
};
void a(float2x2 a_1[4]) {
}
void b(float2x2 m) {
}
void c(float2 v) {
}
void d(float f_1) {
}
float2x2 tint_symbol_1(uint4 buffer[4], uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
uint4 ubo_load = buffer[scalar_offset / 4];
const uint scalar_offset_1 = ((offset + 8u)) / 4;
uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
return float2x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)));
}
typedef float2x2 tint_symbol_ret[4];
tint_symbol_ret tint_symbol(uint4 buffer[4], uint offset) {
float2x2 arr[4] = (float2x2[4])0;
{
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
arr[i] = tint_symbol_1(buffer, (offset + (i * 16u)));
}
}
return arr;
}
[numthreads(1, 1, 1)]
void f() {
a(tint_symbol(u, 0u));
b(tint_symbol_1(u, 16u));
c(asfloat(u[1].xy).yx);
d(asfloat(u[1].xy).yx.x);
return;
}

43
test/tint/buffer/uniform/std140/array/mat2x2/to_fn.wgsl.expected.fxc.hlsl Normal file
View File

@ -0,0 +1,43 @@
cbuffer cbuffer_u : register(b0, space0) {
uint4 u[4];
};
void a(float2x2 a_1[4]) {
}
void b(float2x2 m) {
}
void c(float2 v) {
}
void d(float f_1) {
}
float2x2 tint_symbol_1(uint4 buffer[4], uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
uint4 ubo_load = buffer[scalar_offset / 4];
const uint scalar_offset_1 = ((offset + 8u)) / 4;
uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
return float2x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)));
}
typedef float2x2 tint_symbol_ret[4];
tint_symbol_ret tint_symbol(uint4 buffer[4], uint offset) {
float2x2 arr[4] = (float2x2[4])0;
{
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
arr[i] = tint_symbol_1(buffer, (offset + (i * 16u)));
}
}
return arr;
}
[numthreads(1, 1, 1)]
void f() {
a(tint_symbol(u, 0u));
b(tint_symbol_1(u, 16u));
c(asfloat(u[1].xy).yx);
d(asfloat(u[1].xy).yx.x);
return;
}

30
test/tint/buffer/uniform/std140/array/mat2x2/to_fn.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,30 @@
#version 310 es
layout(binding = 0, std140) uniform u_block_ubo {
mat2 inner[4];
} u;
void a(mat2 a_1[4]) {
}
void b(mat2 m) {
}
void c(vec2 v) {
}
void d(float f_1) {
}
void f() {
a(u.inner);
b(u.inner[1]);
c(u.inner[1][0].yx);
d(u.inner[1][0].yx.x);
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
f();
return;
}

36
test/tint/buffer/uniform/std140/array/mat2x2/to_fn.wgsl.expected.msl Normal file
View File

@ -0,0 +1,36 @@
#include <metal_stdlib>
using namespace metal;
template<typename T, size_t N>
struct tint_array {
const constant T& operator[](size_t i) const constant { return elements[i]; }
device T& operator[](size_t i) device { return elements[i]; }
const device T& operator[](size_t i) const device { return elements[i]; }
thread T& operator[](size_t i) thread { return elements[i]; }
const thread T& operator[](size_t i) const thread { return elements[i]; }
threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
T elements[N];
};
void a(tint_array<float2x2, 4> a_1) {
}
void b(float2x2 m) {
}
void c(float2 v) {
}
void d(float f_1) {
}
kernel void f(const constant tint_array<float2x2, 4>* tint_symbol [[buffer(0)]]) {
a(*(tint_symbol));
b((*(tint_symbol))[1]);
c(float2((*(tint_symbol))[1][0]).yx);
d(float2((*(tint_symbol))[1][0]).yx[0]);
return;
}

90
test/tint/buffer/uniform/std140/array/mat2x2/to_fn.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,90 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 52
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %u_block "u_block"
OpMemberName %u_block 0 "inner"
OpName %u "u"
OpName %a "a"
OpName %a_1 "a_1"
OpName %b "b"
OpName %m "m"
OpName %c "c"
OpName %v "v"
OpName %d "d"
OpName %f_1 "f_1"
OpName %f "f"
OpDecorate %u_block Block
OpMemberDecorate %u_block 0 Offset 0
OpMemberDecorate %u_block 0 ColMajor
OpMemberDecorate %u_block 0 MatrixStride 8
OpDecorate %_arr_mat2v2float_uint_4 ArrayStride 16
OpDecorate %u NonWritable
OpDecorate %u DescriptorSet 0
OpDecorate %u Binding 0
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%mat2v2float = OpTypeMatrix %v2float 2
%uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4
%_arr_mat2v2float_uint_4 = OpTypeArray %mat2v2float %uint_4
%u_block = OpTypeStruct %_arr_mat2v2float_uint_4
%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
%u = OpVariable %_ptr_Uniform_u_block Uniform
%void = OpTypeVoid
%10 = OpTypeFunction %void %_arr_mat2v2float_uint_4
%15 = OpTypeFunction %void %mat2v2float
%19 = OpTypeFunction %void %v2float
%23 = OpTypeFunction %void %float
%27 = OpTypeFunction %void
%uint_0 = OpConstant %uint 0
%_ptr_Uniform__arr_mat2v2float_uint_4 = OpTypePointer Uniform %_arr_mat2v2float_uint_4
%int = OpTypeInt 32 1
%int_1 = OpConstant %int 1
%_ptr_Uniform_mat2v2float = OpTypePointer Uniform %mat2v2float
%42 = OpConstantNull %int
%_ptr_Uniform_v2float = OpTypePointer Uniform %v2float
%a = OpFunction %void None %10
%a_1 = OpFunctionParameter %_arr_mat2v2float_uint_4
%14 = OpLabel
OpReturn
OpFunctionEnd
%b = OpFunction %void None %15
%m = OpFunctionParameter %mat2v2float
%18 = OpLabel
OpReturn
OpFunctionEnd
%c = OpFunction %void None %19
%v = OpFunctionParameter %v2float
%22 = OpLabel
OpReturn
OpFunctionEnd
%d = OpFunction %void None %23
%f_1 = OpFunctionParameter %float
%26 = OpLabel
OpReturn
OpFunctionEnd
%f = OpFunction %void None %27
%29 = OpLabel
%33 = OpAccessChain %_ptr_Uniform__arr_mat2v2float_uint_4 %u %uint_0
%34 = OpLoad %_arr_mat2v2float_uint_4 %33
%30 = OpFunctionCall %void %a %34
%39 = OpAccessChain %_ptr_Uniform_mat2v2float %u %uint_0 %int_1
%40 = OpLoad %mat2v2float %39
%35 = OpFunctionCall %void %b %40
%44 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %int_1 %42
%45 = OpLoad %v2float %44
%46 = OpVectorShuffle %v2float %45 %45 1 0
%41 = OpFunctionCall %void %c %46
%48 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %int_1 %42
%49 = OpLoad %v2float %48
%50 = OpVectorShuffle %v2float %49 %49 1 0
%51 = OpCompositeExtract %float %50 0
%47 = OpFunctionCall %void %d %51
OpReturn
OpFunctionEnd

21
test/tint/buffer/uniform/std140/array/mat2x2/to_fn.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,21 @@
@group(0) @binding(0) var<uniform> u : array<mat2x2<f32>, 4>;
fn a(a : array<mat2x2<f32>, 4>) {
}
fn b(m : mat2x2<f32>) {
}
fn c(v : vec2<f32>) {
}
fn d(f : f32) {
}
@compute @workgroup_size(1)
fn f() {
a(u);
b(u[1]);
c(u[1][0].yx);
d(u[1][0].yx.x);
}

16
test/tint/buffer/uniform/std140/array/mat2x2/to_private.wgsl Normal file
View File

@ -0,0 +1,16 @@
struct S {
before : i32,
m : mat2x2<f32>,
after : i32,
}
@group(0) @binding(0) var<uniform> u : array<mat2x2<f32>, 4>;
var<private> p : array<mat2x2<f32>, 4>;
@compute @workgroup_size(1)
fn f() {
p = u;
p[1] = u[2];
p[1][0] = u[0][1].yx;
p[1][0].x = u[0][1].x;
}

32
test/tint/buffer/uniform/std140/array/mat2x2/to_private.wgsl.expected.dxc.hlsl Normal file
View File

@ -0,0 +1,32 @@
cbuffer cbuffer_u : register(b0, space0) {
uint4 u[4];
};
static float2x2 p[4] = (float2x2[4])0;
float2x2 tint_symbol_1(uint4 buffer[4], uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
uint4 ubo_load = buffer[scalar_offset / 4];
const uint scalar_offset_1 = ((offset + 8u)) / 4;
uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
return float2x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)));
}
typedef float2x2 tint_symbol_ret[4];
tint_symbol_ret tint_symbol(uint4 buffer[4], uint offset) {
float2x2 arr[4] = (float2x2[4])0;
{
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
arr[i] = tint_symbol_1(buffer, (offset + (i * 16u)));
}
}
return arr;
}
[numthreads(1, 1, 1)]
void f() {
p = tint_symbol(u, 0u);
p[1] = tint_symbol_1(u, 32u);
p[1][0] = asfloat(u[0].zw).yx;
p[1][0].x = asfloat(u[0].z);
return;
}

32
test/tint/buffer/uniform/std140/array/mat2x2/to_private.wgsl.expected.fxc.hlsl Normal file
View File

@ -0,0 +1,32 @@
cbuffer cbuffer_u : register(b0, space0) {
uint4 u[4];
};
static float2x2 p[4] = (float2x2[4])0;
float2x2 tint_symbol_1(uint4 buffer[4], uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
uint4 ubo_load = buffer[scalar_offset / 4];
const uint scalar_offset_1 = ((offset + 8u)) / 4;
uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
return float2x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)));
}
typedef float2x2 tint_symbol_ret[4];
tint_symbol_ret tint_symbol(uint4 buffer[4], uint offset) {
float2x2 arr[4] = (float2x2[4])0;
{
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
arr[i] = tint_symbol_1(buffer, (offset + (i * 16u)));
}
}
return arr;
}
[numthreads(1, 1, 1)]
void f() {
p = tint_symbol(u, 0u);
p[1] = tint_symbol_1(u, 32u);
p[1][0] = asfloat(u[0].zw).yx;
p[1][0].x = asfloat(u[0].z);
return;
}

25
test/tint/buffer/uniform/std140/array/mat2x2/to_private.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,25 @@
#version 310 es
struct S {
int before;
mat2 m;
int after;
};
layout(binding = 0, std140) uniform u_block_ubo {
mat2 inner[4];
} u;
mat2 p[4] = mat2[4](mat2(0.0f, 0.0f, 0.0f, 0.0f), mat2(0.0f, 0.0f, 0.0f, 0.0f), mat2(0.0f, 0.0f, 0.0f, 0.0f), mat2(0.0f, 0.0f, 0.0f, 0.0f));
void f() {
p = u.inner;
p[1] = u.inner[2];
p[1][0] = u.inner[0][1].yx;
p[1][0].x = u.inner[0][1].x;
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
f();
return;
}

31
test/tint/buffer/uniform/std140/array/mat2x2/to_private.wgsl.expected.msl Normal file
View File

@ -0,0 +1,31 @@
#include <metal_stdlib>
using namespace metal;
template<typename T, size_t N>
struct tint_array {
const constant T& operator[](size_t i) const constant { return elements[i]; }
device T& operator[](size_t i) device { return elements[i]; }
const device T& operator[](size_t i) const device { return elements[i]; }
thread T& operator[](size_t i) thread { return elements[i]; }
const thread T& operator[](size_t i) const thread { return elements[i]; }
threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
T elements[N];
};
struct S {
int before;
float2x2 m;
int after;
};
kernel void f(const constant tint_array<float2x2, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<float2x2, 4> tint_symbol = {};
tint_symbol = *(tint_symbol_1);
tint_symbol[1] = (*(tint_symbol_1))[2];
tint_symbol[1][0] = float2((*(tint_symbol_1))[0][1]).yx;
tint_symbol[1][0][0] = (*(tint_symbol_1))[0][1][0];
return;
}

68
test/tint/buffer/uniform/std140/array/mat2x2/to_private.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,68 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 41
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %u_block "u_block"
OpMemberName %u_block 0 "inner"
OpName %u "u"
OpName %p "p"
OpName %f "f"
OpDecorate %u_block Block
OpMemberDecorate %u_block 0 Offset 0
OpMemberDecorate %u_block 0 ColMajor
OpMemberDecorate %u_block 0 MatrixStride 8
OpDecorate %_arr_mat2v2float_uint_4 ArrayStride 16
OpDecorate %u NonWritable
OpDecorate %u DescriptorSet 0
OpDecorate %u Binding 0
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%mat2v2float = OpTypeMatrix %v2float 2
%uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4
%_arr_mat2v2float_uint_4 = OpTypeArray %mat2v2float %uint_4
%u_block = OpTypeStruct %_arr_mat2v2float_uint_4
%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
%u = OpVariable %_ptr_Uniform_u_block Uniform
%_ptr_Private__arr_mat2v2float_uint_4 = OpTypePointer Private %_arr_mat2v2float_uint_4
%12 = OpConstantNull %_arr_mat2v2float_uint_4
%p = OpVariable %_ptr_Private__arr_mat2v2float_uint_4 Private %12
%void = OpTypeVoid
%13 = OpTypeFunction %void
%uint_0 = OpConstant %uint 0
%_ptr_Uniform__arr_mat2v2float_uint_4 = OpTypePointer Uniform %_arr_mat2v2float_uint_4
%int = OpTypeInt 32 1
%int_1 = OpConstant %int 1
%_ptr_Private_mat2v2float = OpTypePointer Private %mat2v2float
%int_2 = OpConstant %int 2
%_ptr_Uniform_mat2v2float = OpTypePointer Uniform %mat2v2float
%29 = OpConstantNull %int
%_ptr_Private_v2float = OpTypePointer Private %v2float
%_ptr_Uniform_v2float = OpTypePointer Uniform %v2float
%_ptr_Private_float = OpTypePointer Private %float
%_ptr_Uniform_float = OpTypePointer Uniform %float
%f = OpFunction %void None %13
%16 = OpLabel
%19 = OpAccessChain %_ptr_Uniform__arr_mat2v2float_uint_4 %u %uint_0
%20 = OpLoad %_arr_mat2v2float_uint_4 %19
OpStore %p %20
%24 = OpAccessChain %_ptr_Private_mat2v2float %p %int_1
%27 = OpAccessChain %_ptr_Uniform_mat2v2float %u %uint_0 %int_2
%28 = OpLoad %mat2v2float %27
OpStore %24 %28
%31 = OpAccessChain %_ptr_Private_v2float %p %int_1 %29
%33 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %29 %int_1
%34 = OpLoad %v2float %33
%35 = OpVectorShuffle %v2float %34 %34 1 0
OpStore %31 %35
%37 = OpAccessChain %_ptr_Private_float %p %int_1 %29 %uint_0
%39 = OpAccessChain %_ptr_Uniform_float %u %uint_0 %29 %int_1 %uint_0
%40 = OpLoad %float %39
OpStore %37 %40
OpReturn
OpFunctionEnd

17
test/tint/buffer/uniform/std140/array/mat2x2/to_private.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,17 @@
struct S {
before : i32,
m : mat2x2<f32>,
after : i32,
}
@group(0) @binding(0) var<uniform> u : array<mat2x2<f32>, 4>;
var<private> p : array<mat2x2<f32>, 4>;
@compute @workgroup_size(1)
fn f() {
p = u;
p[1] = u[2];
p[1][0] = u[0][1].yx;
p[1][0].x = u[0][1].x;
}

16
test/tint/buffer/uniform/std140/array/mat2x2/to_storage.wgsl Normal file
View File

@ -0,0 +1,16 @@
struct S {
before : i32,
m : mat2x2<f32>,
after : i32,
}
@group(0) @binding(0) var<uniform> u : array<mat2x2<f32>, 4>;
@group(0) @binding(1) var<storage, read_write> s : array<mat2x2<f32>, 4>;
@compute @workgroup_size(1)
fn f() {
s = u;
s[1] = u[2];
s[1][0] = u[0][1].yx;
s[1][0].x = u[0][1].x;
}

46
test/tint/buffer/uniform/std140/array/mat2x2/to_storage.wgsl.expected.dxc.hlsl Normal file
View File

@ -0,0 +1,46 @@
cbuffer cbuffer_u : register(b0, space0) {
uint4 u[4];
};
RWByteAddressBuffer s : register(u1, space0);
void tint_symbol_1(RWByteAddressBuffer buffer, uint offset, float2x2 value) {
buffer.Store2((offset + 0u), asuint(value[0u]));
buffer.Store2((offset + 8u), asuint(value[1u]));
}
void tint_symbol(RWByteAddressBuffer buffer, uint offset, float2x2 value[4]) {
float2x2 array[4] = value;
{
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
tint_symbol_1(buffer, (offset + (i * 16u)), array[i]);
}
}
}
float2x2 tint_symbol_4(uint4 buffer[4], uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
uint4 ubo_load = buffer[scalar_offset / 4];
const uint scalar_offset_1 = ((offset + 8u)) / 4;
uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
return float2x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)));
}
typedef float2x2 tint_symbol_3_ret[4];
tint_symbol_3_ret tint_symbol_3(uint4 buffer[4], uint offset) {
float2x2 arr[4] = (float2x2[4])0;
{
[loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
arr[i_1] = tint_symbol_4(buffer, (offset + (i_1 * 16u)));
}
}
return arr;
}
[numthreads(1, 1, 1)]
void f() {
tint_symbol(s, 0u, tint_symbol_3(u, 0u));
tint_symbol_1(s, 16u, tint_symbol_4(u, 32u));
s.Store2(16u, asuint(asfloat(u[0].zw).yx));
s.Store(16u, asuint(asfloat(u[0].z)));
return;
}

46
test/tint/buffer/uniform/std140/array/mat2x2/to_storage.wgsl.expected.fxc.hlsl Normal file
View File

@ -0,0 +1,46 @@
cbuffer cbuffer_u : register(b0, space0) {
uint4 u[4];
};
RWByteAddressBuffer s : register(u1, space0);
void tint_symbol_1(RWByteAddressBuffer buffer, uint offset, float2x2 value) {
buffer.Store2((offset + 0u), asuint(value[0u]));
buffer.Store2((offset + 8u), asuint(value[1u]));
}
void tint_symbol(RWByteAddressBuffer buffer, uint offset, float2x2 value[4]) {
float2x2 array[4] = value;
{
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
tint_symbol_1(buffer, (offset + (i * 16u)), array[i]);
}
}
}
float2x2 tint_symbol_4(uint4 buffer[4], uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
uint4 ubo_load = buffer[scalar_offset / 4];
const uint scalar_offset_1 = ((offset + 8u)) / 4;
uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
return float2x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)));
}
typedef float2x2 tint_symbol_3_ret[4];
tint_symbol_3_ret tint_symbol_3(uint4 buffer[4], uint offset) {
float2x2 arr[4] = (float2x2[4])0;
{
[loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
arr[i_1] = tint_symbol_4(buffer, (offset + (i_1 * 16u)));
}
}
return arr;
}
[numthreads(1, 1, 1)]
void f() {
tint_symbol(s, 0u, tint_symbol_3(u, 0u));
tint_symbol_1(s, 16u, tint_symbol_4(u, 32u));
s.Store2(16u, asuint(asfloat(u[0].zw).yx));
s.Store(16u, asuint(asfloat(u[0].z)));
return;
}

28
test/tint/buffer/uniform/std140/array/mat2x2/to_storage.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,28 @@
#version 310 es
struct S {
int before;
mat2 m;
int after;
};
layout(binding = 0, std140) uniform u_block_ubo {
mat2 inner[4];
} u;
layout(binding = 1, std430) buffer u_block_ssbo {
mat2 inner[4];
} s;
void f() {
s.inner = u.inner;
s.inner[1] = u.inner[2];
s.inner[1][0] = u.inner[0][1].yx;
s.inner[1][0].x = u.inner[0][1].x;
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
f();
return;
}

30
test/tint/buffer/uniform/std140/array/mat2x2/to_storage.wgsl.expected.msl Normal file
View File

@ -0,0 +1,30 @@
#include <metal_stdlib>
using namespace metal;
template<typename T, size_t N>
struct tint_array {
const constant T& operator[](size_t i) const constant { return elements[i]; }
device T& operator[](size_t i) device { return elements[i]; }
const device T& operator[](size_t i) const device { return elements[i]; }
thread T& operator[](size_t i) thread { return elements[i]; }
const thread T& operator[](size_t i) const thread { return elements[i]; }
threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
T elements[N];
};
struct S {
int before;
float2x2 m;
int after;
};
kernel void f(device tint_array<float2x2, 4>* tint_symbol [[buffer(1)]], const constant tint_array<float2x2, 4>* tint_symbol_1 [[buffer(0)]]) {
*(tint_symbol) = *(tint_symbol_1);
(*(tint_symbol))[1] = (*(tint_symbol_1))[2];
(*(tint_symbol))[1][0] = float2((*(tint_symbol_1))[0][1]).yx;
(*(tint_symbol))[1][0][0] = (*(tint_symbol_1))[0][1][0];
return;
}

71
test/tint/buffer/uniform/std140/array/mat2x2/to_storage.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,71 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 42
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %u_block "u_block"
OpMemberName %u_block 0 "inner"
OpName %u "u"
OpName %s "s"
OpName %f "f"
OpDecorate %u_block Block
OpMemberDecorate %u_block 0 Offset 0
OpMemberDecorate %u_block 0 ColMajor
OpMemberDecorate %u_block 0 MatrixStride 8
OpDecorate %_arr_mat2v2float_uint_4 ArrayStride 16
OpDecorate %u NonWritable
OpDecorate %u DescriptorSet 0
OpDecorate %u Binding 0
OpDecorate %s DescriptorSet 0
OpDecorate %s Binding 1
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%mat2v2float = OpTypeMatrix %v2float 2
%uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4
%_arr_mat2v2float_uint_4 = OpTypeArray %mat2v2float %uint_4
%u_block = OpTypeStruct %_arr_mat2v2float_uint_4
%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
%u = OpVariable %_ptr_Uniform_u_block Uniform
%_ptr_StorageBuffer_u_block = OpTypePointer StorageBuffer %u_block
%s = OpVariable %_ptr_StorageBuffer_u_block StorageBuffer
%void = OpTypeVoid
%12 = OpTypeFunction %void
%uint_0 = OpConstant %uint 0
%_ptr_StorageBuffer__arr_mat2v2float_uint_4 = OpTypePointer StorageBuffer %_arr_mat2v2float_uint_4
%_ptr_Uniform__arr_mat2v2float_uint_4 = OpTypePointer Uniform %_arr_mat2v2float_uint_4
%int = OpTypeInt 32 1
%int_1 = OpConstant %int 1
%_ptr_StorageBuffer_mat2v2float = OpTypePointer StorageBuffer %mat2v2float
%int_2 = OpConstant %int 2
%_ptr_Uniform_mat2v2float = OpTypePointer Uniform %mat2v2float
%30 = OpConstantNull %int
%_ptr_StorageBuffer_v2float = OpTypePointer StorageBuffer %v2float
%_ptr_Uniform_v2float = OpTypePointer Uniform %v2float
%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
%_ptr_Uniform_float = OpTypePointer Uniform %float
%f = OpFunction %void None %12
%15 = OpLabel
%18 = OpAccessChain %_ptr_StorageBuffer__arr_mat2v2float_uint_4 %s %uint_0
%20 = OpAccessChain %_ptr_Uniform__arr_mat2v2float_uint_4 %u %uint_0
%21 = OpLoad %_arr_mat2v2float_uint_4 %20
OpStore %18 %21
%25 = OpAccessChain %_ptr_StorageBuffer_mat2v2float %s %uint_0 %int_1
%28 = OpAccessChain %_ptr_Uniform_mat2v2float %u %uint_0 %int_2
%29 = OpLoad %mat2v2float %28
OpStore %25 %29
%32 = OpAccessChain %_ptr_StorageBuffer_v2float %s %uint_0 %int_1 %30
%34 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %30 %int_1
%35 = OpLoad %v2float %34
%36 = OpVectorShuffle %v2float %35 %35 1 0
OpStore %32 %36
%38 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_0 %int_1 %30 %uint_0
%40 = OpAccessChain %_ptr_Uniform_float %u %uint_0 %30 %int_1 %uint_0
%41 = OpLoad %float %40
OpStore %38 %41
OpReturn
OpFunctionEnd

17
test/tint/buffer/uniform/std140/array/mat2x2/to_storage.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,17 @@
struct S {
before : i32,
m : mat2x2<f32>,
after : i32,
}
@group(0) @binding(0) var<uniform> u : array<mat2x2<f32>, 4>;
@group(0) @binding(1) var<storage, read_write> s : array<mat2x2<f32>, 4>;
@compute @workgroup_size(1)
fn f() {
s = u;
s[1] = u[2];
s[1][0] = u[0][1].yx;
s[1][0].x = u[0][1].x;
}

16
test/tint/buffer/uniform/std140/array/mat2x2/to_workgroup.wgsl Normal file
View File

@ -0,0 +1,16 @@
struct S {
before : i32,
m : mat2x2<f32>,
after : i32,
}
@group(0) @binding(0) var<uniform> u : array<mat2x2<f32>, 4>;
var<workgroup> w : array<mat2x2<f32>, 4>;
@compute @workgroup_size(1)
fn f() {
w = u;
w[1] = u[2];
w[1][0] = u[0][1].yx;
w[1][0].x = u[0][1].x;
}

47
test/tint/buffer/uniform/std140/array/mat2x2/to_workgroup.wgsl.expected.dxc.hlsl Normal file
View File

@ -0,0 +1,47 @@
cbuffer cbuffer_u : register(b0, space0) {
uint4 u[4];
};
groupshared float2x2 w[4];
struct tint_symbol_1 {
uint local_invocation_index : SV_GroupIndex;
};
float2x2 tint_symbol_3(uint4 buffer[4], uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
uint4 ubo_load = buffer[scalar_offset / 4];
const uint scalar_offset_1 = ((offset + 8u)) / 4;
uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
return float2x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)));
}
typedef float2x2 tint_symbol_2_ret[4];
tint_symbol_2_ret tint_symbol_2(uint4 buffer[4], uint offset) {
float2x2 arr[4] = (float2x2[4])0;
{
[loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
arr[i_1] = tint_symbol_3(buffer, (offset + (i_1 * 16u)));
}
}
return arr;
}
void f_inner(uint local_invocation_index) {
{
[loop] for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
const uint i = idx;
w[i] = float2x2((0.0f).xx, (0.0f).xx);
}
}
GroupMemoryBarrierWithGroupSync();
w = tint_symbol_2(u, 0u);
w[1] = tint_symbol_3(u, 32u);
w[1][0] = asfloat(u[0].zw).yx;
w[1][0].x = asfloat(u[0].z);
}
[numthreads(1, 1, 1)]
void f(tint_symbol_1 tint_symbol) {
f_inner(tint_symbol.local_invocation_index);
return;
}

47
test/tint/buffer/uniform/std140/array/mat2x2/to_workgroup.wgsl.expected.fxc.hlsl Normal file
View File

@ -0,0 +1,47 @@
cbuffer cbuffer_u : register(b0, space0) {
uint4 u[4];
};
groupshared float2x2 w[4];
struct tint_symbol_1 {
uint local_invocation_index : SV_GroupIndex;
};
float2x2 tint_symbol_3(uint4 buffer[4], uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
uint4 ubo_load = buffer[scalar_offset / 4];
const uint scalar_offset_1 = ((offset + 8u)) / 4;
uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
return float2x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)));
}
typedef float2x2 tint_symbol_2_ret[4];
tint_symbol_2_ret tint_symbol_2(uint4 buffer[4], uint offset) {
float2x2 arr[4] = (float2x2[4])0;
{
[loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
arr[i_1] = tint_symbol_3(buffer, (offset + (i_1 * 16u)));
}
}
return arr;
}
void f_inner(uint local_invocation_index) {
{
[loop] for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
const uint i = idx;
w[i] = float2x2((0.0f).xx, (0.0f).xx);
}
}
GroupMemoryBarrierWithGroupSync();
w = tint_symbol_2(u, 0u);
w[1] = tint_symbol_3(u, 32u);
w[1][0] = asfloat(u[0].zw).yx;
w[1][0].x = asfloat(u[0].z);
}
[numthreads(1, 1, 1)]
void f(tint_symbol_1 tint_symbol) {
f_inner(tint_symbol.local_invocation_index);
return;
}

32
test/tint/buffer/uniform/std140/array/mat2x2/to_workgroup.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,32 @@
#version 310 es
struct S {
int before;
mat2 m;
int after;
};
layout(binding = 0, std140) uniform u_block_ubo {
mat2 inner[4];
} u;
shared mat2 w[4];
void f(uint local_invocation_index) {
{
for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint i = idx;
w[i] = mat2(vec2(0.0f), vec2(0.0f));
}
}
barrier();
w = u.inner;
w[1] = u.inner[2];
w[1][0] = u.inner[0][1].yx;
w[1][0].x = u.inner[0][1].x;
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
f(gl_LocalInvocationIndex);
return;
}

44
test/tint/buffer/uniform/std140/array/mat2x2/to_workgroup.wgsl.expected.msl Normal file
View File

@ -0,0 +1,44 @@
#include <metal_stdlib>
using namespace metal;
template<typename T, size_t N>
struct tint_array {
const constant T& operator[](size_t i) const constant { return elements[i]; }
device T& operator[](size_t i) device { return elements[i]; }
const device T& operator[](size_t i) const device { return elements[i]; }
thread T& operator[](size_t i) thread { return elements[i]; }
const thread T& operator[](size_t i) const thread { return elements[i]; }
threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
T elements[N];
};
struct tint_symbol_5 {
tint_array<float2x2, 4> w;
};
struct S {
int before;
float2x2 m;
int after;
};
void f_inner(uint local_invocation_index, threadgroup tint_array<float2x2, 4>* const tint_symbol, const constant tint_array<float2x2, 4>* const tint_symbol_1) {
for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
(*(tint_symbol))[i] = float2x2(float2(0.0f), float2(0.0f));
}
threadgroup_barrier(mem_flags::mem_threadgroup);
*(tint_symbol) = *(tint_symbol_1);
(*(tint_symbol))[1] = (*(tint_symbol_1))[2];
(*(tint_symbol))[1][0] = float2((*(tint_symbol_1))[0][1]).yx;
(*(tint_symbol))[1][0][0] = (*(tint_symbol_1))[0][1][0];
}
kernel void f(const constant tint_array<float2x2, 4>* tint_symbol_4 [[buffer(0)]], threadgroup tint_symbol_5* tint_symbol_3 [[threadgroup(0)]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
threadgroup tint_array<float2x2, 4>* const tint_symbol_2 = &((*(tint_symbol_3)).w);
f_inner(local_invocation_index, tint_symbol_2, tint_symbol_4);
return;
}

115
test/tint/buffer/uniform/std140/array/mat2x2/to_workgroup.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,115 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 70
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f" %local_invocation_index_1
OpExecutionMode %f LocalSize 1 1 1
OpName %local_invocation_index_1 "local_invocation_index_1"
OpName %u_block "u_block"
OpMemberName %u_block 0 "inner"
OpName %u "u"
OpName %w "w"
OpName %f_inner "f_inner"
OpName %local_invocation_index "local_invocation_index"
OpName %idx "idx"
OpName %f "f"
OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
OpDecorate %u_block Block
OpMemberDecorate %u_block 0 Offset 0
OpMemberDecorate %u_block 0 ColMajor
OpMemberDecorate %u_block 0 MatrixStride 8
OpDecorate %_arr_mat2v2float_uint_4 ArrayStride 16
OpDecorate %u NonWritable
OpDecorate %u DescriptorSet 0
OpDecorate %u Binding 0
%uint = OpTypeInt 32 0
%_ptr_Input_uint = OpTypePointer Input %uint
%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%mat2v2float = OpTypeMatrix %v2float 2
%uint_4 = OpConstant %uint 4
%_arr_mat2v2float_uint_4 = OpTypeArray %mat2v2float %uint_4
%u_block = OpTypeStruct %_arr_mat2v2float_uint_4
%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
%u = OpVariable %_ptr_Uniform_u_block Uniform
%_ptr_Workgroup__arr_mat2v2float_uint_4 = OpTypePointer Workgroup %_arr_mat2v2float_uint_4
%w = OpVariable %_ptr_Workgroup__arr_mat2v2float_uint_4 Workgroup
%void = OpTypeVoid
%14 = OpTypeFunction %void %uint
%_ptr_Function_uint = OpTypePointer Function %uint
%21 = OpConstantNull %uint
%bool = OpTypeBool
%_ptr_Workgroup_mat2v2float = OpTypePointer Workgroup %mat2v2float
%35 = OpConstantNull %mat2v2float
%uint_1 = OpConstant %uint 1
%uint_2 = OpConstant %uint 2
%uint_264 = OpConstant %uint 264
%uint_0 = OpConstant %uint 0
%_ptr_Uniform__arr_mat2v2float_uint_4 = OpTypePointer Uniform %_arr_mat2v2float_uint_4
%int = OpTypeInt 32 1
%int_1 = OpConstant %int 1
%int_2 = OpConstant %int 2
%_ptr_Uniform_mat2v2float = OpTypePointer Uniform %mat2v2float
%53 = OpConstantNull %int
%_ptr_Workgroup_v2float = OpTypePointer Workgroup %v2float
%_ptr_Uniform_v2float = OpTypePointer Uniform %v2float
%_ptr_Workgroup_float = OpTypePointer Workgroup %float
%_ptr_Uniform_float = OpTypePointer Uniform %float
%65 = OpTypeFunction %void
%f_inner = OpFunction %void None %14
%local_invocation_index = OpFunctionParameter %uint
%18 = OpLabel
%idx = OpVariable %_ptr_Function_uint Function %21
OpStore %idx %local_invocation_index
OpBranch %22
%22 = OpLabel
OpLoopMerge %23 %24 None
OpBranch %25
%25 = OpLabel
%27 = OpLoad %uint %idx
%28 = OpULessThan %bool %27 %uint_4
%26 = OpLogicalNot %bool %28
OpSelectionMerge %30 None
OpBranchConditional %26 %31 %30
%31 = OpLabel
OpBranch %23
%30 = OpLabel
%32 = OpLoad %uint %idx
%34 = OpAccessChain %_ptr_Workgroup_mat2v2float %w %32
OpStore %34 %35
OpBranch %24
%24 = OpLabel
%36 = OpLoad %uint %idx
%38 = OpIAdd %uint %36 %uint_1
OpStore %idx %38
OpBranch %22
%23 = OpLabel
OpControlBarrier %uint_2 %uint_2 %uint_264
%44 = OpAccessChain %_ptr_Uniform__arr_mat2v2float_uint_4 %u %uint_0
%45 = OpLoad %_arr_mat2v2float_uint_4 %44
OpStore %w %45
%48 = OpAccessChain %_ptr_Workgroup_mat2v2float %w %int_1
%51 = OpAccessChain %_ptr_Uniform_mat2v2float %u %uint_0 %int_2
%52 = OpLoad %mat2v2float %51
OpStore %48 %52
%55 = OpAccessChain %_ptr_Workgroup_v2float %w %int_1 %53
%57 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %53 %int_1
%58 = OpLoad %v2float %57
%59 = OpVectorShuffle %v2float %58 %58 1 0
OpStore %55 %59
%61 = OpAccessChain %_ptr_Workgroup_float %w %int_1 %53 %uint_0
%63 = OpAccessChain %_ptr_Uniform_float %u %uint_0 %53 %int_1 %uint_0
%64 = OpLoad %float %63
OpStore %61 %64
OpReturn
OpFunctionEnd
%f = OpFunction %void None %65
%67 = OpLabel
%69 = OpLoad %uint %local_invocation_index_1
%68 = OpFunctionCall %void %f_inner %69
OpReturn
OpFunctionEnd

17
test/tint/buffer/uniform/std140/array/mat2x2/to_workgroup.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,17 @@
struct S {
before : i32,
m : mat2x2<f32>,
after : i32,
}
@group(0) @binding(0) var<uniform> u : array<mat2x2<f32>, 4>;
var<workgroup> w : array<mat2x2<f32>, 4>;
@compute @workgroup_size(1)
fn f() {
w = u;
w[1] = u[2];
w[1][0] = u[0][1].yx;
w[1][0].x = u[0][1].x;
}

15
test/tint/buffer/uniform/std140/array/mat4x2/dynamic_index_via_ptr.wgsl Normal file
View File

@ -0,0 +1,15 @@
@group(0) @binding(0) var<uniform> a : array<mat4x2<f32>, 4>;
var<private> counter = 0;
fn i() -> i32 { counter++; return counter; }
@compute @workgroup_size(1)
fn f() {
let p_a = &a;
let p_a_i = &((*p_a)[i()]);
let p_a_i_i = &((*p_a_i)[i()]);
let l_a : array<mat4x2<f32>, 4> = *p_a;
let l_a_i : mat4x2<f32> = *p_a_i;
let l_a_i_i : vec2<f32> = *p_a_i_i;
}

44
test/tint/buffer/uniform/std140/array/mat4x2/dynamic_index_via_ptr.wgsl.expected.dxc.hlsl Normal file
View File

@ -0,0 +1,44 @@
cbuffer cbuffer_a : register(b0, space0) {
uint4 a[8];
};
static int counter = 0;
int i() {
counter = (counter + 1);
return counter;
}
float4x2 tint_symbol_1(uint4 buffer[8], uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
uint4 ubo_load = buffer[scalar_offset / 4];
const uint scalar_offset_1 = ((offset + 8u)) / 4;
uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
const uint scalar_offset_2 = ((offset + 16u)) / 4;
uint4 ubo_load_2 = buffer[scalar_offset_2 / 4];
const uint scalar_offset_3 = ((offset + 24u)) / 4;
uint4 ubo_load_3 = buffer[scalar_offset_3 / 4];
return float4x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)), asfloat(((scalar_offset_2 & 2) ? ubo_load_2.zw : ubo_load_2.xy)), asfloat(((scalar_offset_3 & 2) ? ubo_load_3.zw : ubo_load_3.xy)));
}
typedef float4x2 tint_symbol_ret[4];
tint_symbol_ret tint_symbol(uint4 buffer[8], uint offset) {
float4x2 arr[4] = (float4x2[4])0;
{
[loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
arr[i_1] = tint_symbol_1(buffer, (offset + (i_1 * 32u)));
}
}
return arr;
}
[numthreads(1, 1, 1)]
void f() {
const int p_a_i_save = i();
const int p_a_i_i_save = i();
const float4x2 l_a[4] = tint_symbol(a, 0u);
const float4x2 l_a_i = tint_symbol_1(a, (32u * uint(p_a_i_save)));
const uint scalar_offset_4 = (((32u * uint(p_a_i_save)) + (8u * uint(p_a_i_i_save)))) / 4;
uint4 ubo_load_4 = a[scalar_offset_4 / 4];
const float2 l_a_i_i = asfloat(((scalar_offset_4 & 2) ? ubo_load_4.zw : ubo_load_4.xy));
return;
}

44
test/tint/buffer/uniform/std140/array/mat4x2/dynamic_index_via_ptr.wgsl.expected.fxc.hlsl Normal file
View File

@ -0,0 +1,44 @@
cbuffer cbuffer_a : register(b0, space0) {
uint4 a[8];
};
static int counter = 0;
int i() {
counter = (counter + 1);
return counter;
}
float4x2 tint_symbol_1(uint4 buffer[8], uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
uint4 ubo_load = buffer[scalar_offset / 4];
const uint scalar_offset_1 = ((offset + 8u)) / 4;
uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
const uint scalar_offset_2 = ((offset + 16u)) / 4;
uint4 ubo_load_2 = buffer[scalar_offset_2 / 4];
const uint scalar_offset_3 = ((offset + 24u)) / 4;
uint4 ubo_load_3 = buffer[scalar_offset_3 / 4];
return float4x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)), asfloat(((scalar_offset_2 & 2) ? ubo_load_2.zw : ubo_load_2.xy)), asfloat(((scalar_offset_3 & 2) ? ubo_load_3.zw : ubo_load_3.xy)));
}
typedef float4x2 tint_symbol_ret[4];
tint_symbol_ret tint_symbol(uint4 buffer[8], uint offset) {
float4x2 arr[4] = (float4x2[4])0;
{
[loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
arr[i_1] = tint_symbol_1(buffer, (offset + (i_1 * 32u)));
}
}
return arr;
}
[numthreads(1, 1, 1)]
void f() {
const int p_a_i_save = i();
const int p_a_i_i_save = i();
const float4x2 l_a[4] = tint_symbol(a, 0u);
const float4x2 l_a_i = tint_symbol_1(a, (32u * uint(p_a_i_save)));
const uint scalar_offset_4 = (((32u * uint(p_a_i_save)) + (8u * uint(p_a_i_i_save)))) / 4;
uint4 ubo_load_4 = a[scalar_offset_4 / 4];
const float2 l_a_i_i = asfloat(((scalar_offset_4 & 2) ? ubo_load_4.zw : ubo_load_4.xy));
return;
}

27
test/tint/buffer/uniform/std140/array/mat4x2/dynamic_index_via_ptr.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,27 @@
#version 310 es
layout(binding = 0, std140) uniform a_block_ubo {
mat4x2 inner[4];
} a;
int counter = 0;
int i() {
counter = (counter + 1);
return counter;
}
void f() {
int tint_symbol = i();
int p_a_i_save = tint_symbol;
int tint_symbol_1 = i();
int p_a_i_i_save = tint_symbol_1;
mat4x2 l_a[4] = a.inner;
mat4x2 l_a_i = a.inner[p_a_i_save];
vec2 l_a_i_i = a.inner[p_a_i_save][p_a_i_i_save];
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
f();
return;
}

33
test/tint/buffer/uniform/std140/array/mat4x2/dynamic_index_via_ptr.wgsl.expected.msl Normal file
View File

@ -0,0 +1,33 @@
#include <metal_stdlib>
using namespace metal;
template<typename T, size_t N>
struct tint_array {
const constant T& operator[](size_t i) const constant { return elements[i]; }
device T& operator[](size_t i) device { return elements[i]; }
const device T& operator[](size_t i) const device { return elements[i]; }
thread T& operator[](size_t i) thread { return elements[i]; }
const thread T& operator[](size_t i) const thread { return elements[i]; }
threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
T elements[N];
};
int i() {
thread int tint_symbol_2 = 0;
tint_symbol_2 = as_type<int>((as_type<uint>(tint_symbol_2) + as_type<uint>(1)));
return tint_symbol_2;
}
kernel void f(const constant tint_array<float4x2, 4>* tint_symbol_3 [[buffer(0)]]) {
int const tint_symbol = i();
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const p_a_i_i_save = tint_symbol_1;
tint_array<float4x2, 4> const l_a = *(tint_symbol_3);
float4x2 const l_a_i = (*(tint_symbol_3))[p_a_i_save];
float2 const l_a_i_i = (*(tint_symbol_3))[p_a_i_save][p_a_i_i_save];
return;
}

64
test/tint/buffer/uniform/std140/array/mat4x2/dynamic_index_via_ptr.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,64 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 37
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %a_block "a_block"
OpMemberName %a_block 0 "inner"
OpName %a "a"
OpName %counter "counter"
OpName %i "i"
OpName %f "f"
OpDecorate %a_block Block
OpMemberDecorate %a_block 0 Offset 0
OpMemberDecorate %a_block 0 ColMajor
OpMemberDecorate %a_block 0 MatrixStride 8
OpDecorate %_arr_mat4v2float_uint_4 ArrayStride 32
OpDecorate %a NonWritable
OpDecorate %a DescriptorSet 0
OpDecorate %a Binding 0
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%mat4v2float = OpTypeMatrix %v2float 4
%uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4
%_arr_mat4v2float_uint_4 = OpTypeArray %mat4v2float %uint_4
%a_block = OpTypeStruct %_arr_mat4v2float_uint_4
%_ptr_Uniform_a_block = OpTypePointer Uniform %a_block
%a = OpVariable %_ptr_Uniform_a_block Uniform
%int = OpTypeInt 32 1
%11 = OpConstantNull %int
%_ptr_Private_int = OpTypePointer Private %int
%counter = OpVariable %_ptr_Private_int Private %11
%14 = OpTypeFunction %int
%int_1 = OpConstant %int 1
%void = OpTypeVoid
%21 = OpTypeFunction %void
%uint_0 = OpConstant %uint 0
%_ptr_Uniform__arr_mat4v2float_uint_4 = OpTypePointer Uniform %_arr_mat4v2float_uint_4
%_ptr_Uniform_mat4v2float = OpTypePointer Uniform %mat4v2float
%_ptr_Uniform_v2float = OpTypePointer Uniform %v2float
%i = OpFunction %int None %14
%16 = OpLabel
%17 = OpLoad %int %counter
%19 = OpIAdd %int %17 %int_1
OpStore %counter %19
%20 = OpLoad %int %counter
OpReturnValue %20
OpFunctionEnd
%f = OpFunction %void None %21
%24 = OpLabel
%25 = OpFunctionCall %int %i
%26 = OpFunctionCall %int %i
%29 = OpAccessChain %_ptr_Uniform__arr_mat4v2float_uint_4 %a %uint_0
%30 = OpLoad %_arr_mat4v2float_uint_4 %29
%32 = OpAccessChain %_ptr_Uniform_mat4v2float %a %uint_0 %25
%33 = OpLoad %mat4v2float %32
%35 = OpAccessChain %_ptr_Uniform_v2float %a %uint_0 %25 %26
%36 = OpLoad %v2float %35
OpReturn
OpFunctionEnd

18
test/tint/buffer/uniform/std140/array/mat4x2/dynamic_index_via_ptr.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,18 @@
@group(0) @binding(0) var<uniform> a : array<mat4x2<f32>, 4>;
var<private> counter = 0;
fn i() -> i32 {
counter++;
return counter;
}
@compute @workgroup_size(1)
fn f() {
let p_a = &(a);
let p_a_i = &((*(p_a))[i()]);
let p_a_i_i = &((*(p_a_i))[i()]);
let l_a : array<mat4x2<f32>, 4> = *(p_a);
let l_a_i : mat4x2<f32> = *(p_a_i);
let l_a_i_i : vec2<f32> = *(p_a_i_i);
}

12
test/tint/buffer/uniform/std140/array/mat4x2/static_index_via_ptr.wgsl Normal file
View File

@ -0,0 +1,12 @@
@group(0) @binding(0) var<uniform> a : array<mat4x2<f32>, 4>;
@compute @workgroup_size(1)
fn f() {
let p_a = &a;
let p_a_2 = &((*p_a)[2]);
let p_a_2_1 = &((*p_a_2)[1]);
let l_a : array<mat4x2<f32>, 4> = *p_a;
let l_a_i : mat4x2<f32> = *p_a_2;
let l_a_i_i : vec2<f32> = *p_a_2_1;
}

34
test/tint/buffer/uniform/std140/array/mat4x2/static_index_via_ptr.wgsl.expected.dxc.hlsl Normal file
View File

@ -0,0 +1,34 @@
cbuffer cbuffer_a : register(b0, space0) {
uint4 a[8];
};
float4x2 tint_symbol_1(uint4 buffer[8], uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
uint4 ubo_load = buffer[scalar_offset / 4];
const uint scalar_offset_1 = ((offset + 8u)) / 4;
uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
const uint scalar_offset_2 = ((offset + 16u)) / 4;
uint4 ubo_load_2 = buffer[scalar_offset_2 / 4];
const uint scalar_offset_3 = ((offset + 24u)) / 4;
uint4 ubo_load_3 = buffer[scalar_offset_3 / 4];
return float4x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)), asfloat(((scalar_offset_2 & 2) ? ubo_load_2.zw : ubo_load_2.xy)), asfloat(((scalar_offset_3 & 2) ? ubo_load_3.zw : ubo_load_3.xy)));
}
typedef float4x2 tint_symbol_ret[4];
tint_symbol_ret tint_symbol(uint4 buffer[8], uint offset) {
float4x2 arr[4] = (float4x2[4])0;
{
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
arr[i] = tint_symbol_1(buffer, (offset + (i * 32u)));
}
}
return arr;
}
[numthreads(1, 1, 1)]
void f() {
const float4x2 l_a[4] = tint_symbol(a, 0u);
const float4x2 l_a_i = tint_symbol_1(a, 64u);
const float2 l_a_i_i = asfloat(a[4].zw);
return;
}

34
test/tint/buffer/uniform/std140/array/mat4x2/static_index_via_ptr.wgsl.expected.fxc.hlsl Normal file
View File

@ -0,0 +1,34 @@
cbuffer cbuffer_a : register(b0, space0) {
uint4 a[8];
};
float4x2 tint_symbol_1(uint4 buffer[8], uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
uint4 ubo_load = buffer[scalar_offset / 4];
const uint scalar_offset_1 = ((offset + 8u)) / 4;
uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
const uint scalar_offset_2 = ((offset + 16u)) / 4;
uint4 ubo_load_2 = buffer[scalar_offset_2 / 4];
const uint scalar_offset_3 = ((offset + 24u)) / 4;
uint4 ubo_load_3 = buffer[scalar_offset_3 / 4];
return float4x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)), asfloat(((scalar_offset_2 & 2) ? ubo_load_2.zw : ubo_load_2.xy)), asfloat(((scalar_offset_3 & 2) ? ubo_load_3.zw : ubo_load_3.xy)));
}
typedef float4x2 tint_symbol_ret[4];
tint_symbol_ret tint_symbol(uint4 buffer[8], uint offset) {
float4x2 arr[4] = (float4x2[4])0;
{
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
arr[i] = tint_symbol_1(buffer, (offset + (i * 32u)));
}
}
return arr;
}
[numthreads(1, 1, 1)]
void f() {
const float4x2 l_a[4] = tint_symbol(a, 0u);
const float4x2 l_a_i = tint_symbol_1(a, 64u);
const float2 l_a_i_i = asfloat(a[4].zw);
return;
}

17
test/tint/buffer/uniform/std140/array/mat4x2/static_index_via_ptr.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,17 @@
#version 310 es
layout(binding = 0, std140) uniform a_block_ubo {
mat4x2 inner[4];
} a;
void f() {
mat4x2 l_a[4] = a.inner;
mat4x2 l_a_i = a.inner[2];
vec2 l_a_i_i = a.inner[2][1];
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
f();
return;
}

23
test/tint/buffer/uniform/std140/array/mat4x2/static_index_via_ptr.wgsl.expected.msl Normal file
View File

@ -0,0 +1,23 @@
#include <metal_stdlib>
using namespace metal;
template<typename T, size_t N>
struct tint_array {
const constant T& operator[](size_t i) const constant { return elements[i]; }
device T& operator[](size_t i) device { return elements[i]; }
const device T& operator[](size_t i) const device { return elements[i]; }
thread T& operator[](size_t i) thread { return elements[i]; }
const thread T& operator[](size_t i) const thread { return elements[i]; }
threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
T elements[N];
};
kernel void f(const constant tint_array<float4x2, 4>* tint_symbol [[buffer(0)]]) {
tint_array<float4x2, 4> const l_a = *(tint_symbol);
float4x2 const l_a_i = (*(tint_symbol))[2];
float2 const l_a_i_i = (*(tint_symbol))[2][1];
return;
}

49
test/tint/buffer/uniform/std140/array/mat4x2/static_index_via_ptr.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,49 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 27
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %a_block "a_block"
OpMemberName %a_block 0 "inner"
OpName %a "a"
OpName %f "f"
OpDecorate %a_block Block
OpMemberDecorate %a_block 0 Offset 0
OpMemberDecorate %a_block 0 ColMajor
OpMemberDecorate %a_block 0 MatrixStride 8
OpDecorate %_arr_mat4v2float_uint_4 ArrayStride 32
OpDecorate %a NonWritable
OpDecorate %a DescriptorSet 0
OpDecorate %a Binding 0
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%mat4v2float = OpTypeMatrix %v2float 4
%uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4
%_arr_mat4v2float_uint_4 = OpTypeArray %mat4v2float %uint_4
%a_block = OpTypeStruct %_arr_mat4v2float_uint_4
%_ptr_Uniform_a_block = OpTypePointer Uniform %a_block
%a = OpVariable %_ptr_Uniform_a_block Uniform
%void = OpTypeVoid
%10 = OpTypeFunction %void
%uint_0 = OpConstant %uint 0
%_ptr_Uniform__arr_mat4v2float_uint_4 = OpTypePointer Uniform %_arr_mat4v2float_uint_4
%int = OpTypeInt 32 1
%int_2 = OpConstant %int 2
%_ptr_Uniform_mat4v2float = OpTypePointer Uniform %mat4v2float
%int_1 = OpConstant %int 1
%_ptr_Uniform_v2float = OpTypePointer Uniform %v2float
%f = OpFunction %void None %10
%13 = OpLabel
%16 = OpAccessChain %_ptr_Uniform__arr_mat4v2float_uint_4 %a %uint_0
%17 = OpLoad %_arr_mat4v2float_uint_4 %16
%21 = OpAccessChain %_ptr_Uniform_mat4v2float %a %uint_0 %int_2
%22 = OpLoad %mat4v2float %21
%25 = OpAccessChain %_ptr_Uniform_v2float %a %uint_0 %int_2 %int_1
%26 = OpLoad %v2float %25
OpReturn
OpFunctionEnd

11
test/tint/buffer/uniform/std140/array/mat4x2/static_index_via_ptr.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,11 @@
@group(0) @binding(0) var<uniform> a : array<mat4x2<f32>, 4>;
@compute @workgroup_size(1)
fn f() {
let p_a = &(a);
let p_a_2 = &((*(p_a))[2]);
let p_a_2_1 = &((*(p_a_2))[1]);
let l_a : array<mat4x2<f32>, 4> = *(p_a);
let l_a_i : mat4x2<f32> = *(p_a_2);
let l_a_i_i : vec2<f32> = *(p_a_2_1);
}

14
test/tint/buffer/uniform/std140/array/mat4x2/to_fn.wgsl Normal file
View File

@ -0,0 +1,14 @@
@group(0) @binding(0) var<uniform> u : array<mat4x2<f32>, 4>;
fn a(a : array<mat4x2<f32>, 4>) {}
fn b(m : mat4x2<f32>) {}
fn c(v : vec2<f32>) {}
fn d(f : f32) {}
@compute @workgroup_size(1)
fn f() {
a(u);
b(u[1]);
c(u[1][0].yx);
d(u[1][0].yx.x);
}

47
test/tint/buffer/uniform/std140/array/mat4x2/to_fn.wgsl.expected.dxc.hlsl Normal file
View File

@ -0,0 +1,47 @@
cbuffer cbuffer_u : register(b0, space0) {
uint4 u[8];
};
void a(float4x2 a_1[4]) {
}
void b(float4x2 m) {
}
void c(float2 v) {
}
void d(float f_1) {
}
float4x2 tint_symbol_1(uint4 buffer[8], uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
uint4 ubo_load = buffer[scalar_offset / 4];
const uint scalar_offset_1 = ((offset + 8u)) / 4;
uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
const uint scalar_offset_2 = ((offset + 16u)) / 4;
uint4 ubo_load_2 = buffer[scalar_offset_2 / 4];
const uint scalar_offset_3 = ((offset + 24u)) / 4;
uint4 ubo_load_3 = buffer[scalar_offset_3 / 4];
return float4x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)), asfloat(((scalar_offset_2 & 2) ? ubo_load_2.zw : ubo_load_2.xy)), asfloat(((scalar_offset_3 & 2) ? ubo_load_3.zw : ubo_load_3.xy)));
}
typedef float4x2 tint_symbol_ret[4];
tint_symbol_ret tint_symbol(uint4 buffer[8], uint offset) {
float4x2 arr[4] = (float4x2[4])0;
{
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
arr[i] = tint_symbol_1(buffer, (offset + (i * 32u)));
}
}
return arr;
}
[numthreads(1, 1, 1)]
void f() {
a(tint_symbol(u, 0u));
b(tint_symbol_1(u, 32u));
c(asfloat(u[2].xy).yx);
d(asfloat(u[2].xy).yx.x);
return;
}

47
test/tint/buffer/uniform/std140/array/mat4x2/to_fn.wgsl.expected.fxc.hlsl Normal file
View File

@ -0,0 +1,47 @@
cbuffer cbuffer_u : register(b0, space0) {
uint4 u[8];
};
void a(float4x2 a_1[4]) {
}
void b(float4x2 m) {
}
void c(float2 v) {
}
void d(float f_1) {
}
float4x2 tint_symbol_1(uint4 buffer[8], uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
uint4 ubo_load = buffer[scalar_offset / 4];
const uint scalar_offset_1 = ((offset + 8u)) / 4;
uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
const uint scalar_offset_2 = ((offset + 16u)) / 4;
uint4 ubo_load_2 = buffer[scalar_offset_2 / 4];
const uint scalar_offset_3 = ((offset + 24u)) / 4;
uint4 ubo_load_3 = buffer[scalar_offset_3 / 4];
return float4x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)), asfloat(((scalar_offset_2 & 2) ? ubo_load_2.zw : ubo_load_2.xy)), asfloat(((scalar_offset_3 & 2) ? ubo_load_3.zw : ubo_load_3.xy)));
}
typedef float4x2 tint_symbol_ret[4];
tint_symbol_ret tint_symbol(uint4 buffer[8], uint offset) {
float4x2 arr[4] = (float4x2[4])0;
{
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
arr[i] = tint_symbol_1(buffer, (offset + (i * 32u)));
}
}
return arr;
}
[numthreads(1, 1, 1)]
void f() {
a(tint_symbol(u, 0u));
b(tint_symbol_1(u, 32u));
c(asfloat(u[2].xy).yx);
d(asfloat(u[2].xy).yx.x);
return;
}

30
test/tint/buffer/uniform/std140/array/mat4x2/to_fn.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,30 @@
#version 310 es
layout(binding = 0, std140) uniform u_block_ubo {
mat4x2 inner[4];
} u;
void a(mat4x2 a_1[4]) {
}
void b(mat4x2 m) {
}
void c(vec2 v) {
}
void d(float f_1) {
}
void f() {
a(u.inner);
b(u.inner[1]);
c(u.inner[1][0].yx);
d(u.inner[1][0].yx.x);
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
f();
return;
}

36
test/tint/buffer/uniform/std140/array/mat4x2/to_fn.wgsl.expected.msl Normal file
View File

@ -0,0 +1,36 @@
#include <metal_stdlib>
using namespace metal;
template<typename T, size_t N>
struct tint_array {
const constant T& operator[](size_t i) const constant { return elements[i]; }
device T& operator[](size_t i) device { return elements[i]; }
const device T& operator[](size_t i) const device { return elements[i]; }
thread T& operator[](size_t i) thread { return elements[i]; }
const thread T& operator[](size_t i) const thread { return elements[i]; }
threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
T elements[N];
};
void a(tint_array<float4x2, 4> a_1) {
}
void b(float4x2 m) {
}
void c(float2 v) {
}
void d(float f_1) {
}
kernel void f(const constant tint_array<float4x2, 4>* tint_symbol [[buffer(0)]]) {
a(*(tint_symbol));
b((*(tint_symbol))[1]);
c(float2((*(tint_symbol))[1][0]).yx);
d(float2((*(tint_symbol))[1][0]).yx[0]);
return;
}

90
test/tint/buffer/uniform/std140/array/mat4x2/to_fn.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,90 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 52
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %u_block "u_block"
OpMemberName %u_block 0 "inner"
OpName %u "u"
OpName %a "a"
OpName %a_1 "a_1"
OpName %b "b"
OpName %m "m"
OpName %c "c"
OpName %v "v"
OpName %d "d"
OpName %f_1 "f_1"
OpName %f "f"
OpDecorate %u_block Block
OpMemberDecorate %u_block 0 Offset 0
OpMemberDecorate %u_block 0 ColMajor
OpMemberDecorate %u_block 0 MatrixStride 8
OpDecorate %_arr_mat4v2float_uint_4 ArrayStride 32
OpDecorate %u NonWritable
OpDecorate %u DescriptorSet 0
OpDecorate %u Binding 0
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%mat4v2float = OpTypeMatrix %v2float 4
%uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4
%_arr_mat4v2float_uint_4 = OpTypeArray %mat4v2float %uint_4
%u_block = OpTypeStruct %_arr_mat4v2float_uint_4
%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
%u = OpVariable %_ptr_Uniform_u_block Uniform
%void = OpTypeVoid
%10 = OpTypeFunction %void %_arr_mat4v2float_uint_4
%15 = OpTypeFunction %void %mat4v2float
%19 = OpTypeFunction %void %v2float
%23 = OpTypeFunction %void %float
%27 = OpTypeFunction %void
%uint_0 = OpConstant %uint 0
%_ptr_Uniform__arr_mat4v2float_uint_4 = OpTypePointer Uniform %_arr_mat4v2float_uint_4
%int = OpTypeInt 32 1
%int_1 = OpConstant %int 1
%_ptr_Uniform_mat4v2float = OpTypePointer Uniform %mat4v2float
%42 = OpConstantNull %int
%_ptr_Uniform_v2float = OpTypePointer Uniform %v2float
%a = OpFunction %void None %10
%a_1 = OpFunctionParameter %_arr_mat4v2float_uint_4
%14 = OpLabel
OpReturn
OpFunctionEnd
%b = OpFunction %void None %15
%m = OpFunctionParameter %mat4v2float
%18 = OpLabel
OpReturn
OpFunctionEnd
%c = OpFunction %void None %19
%v = OpFunctionParameter %v2float
%22 = OpLabel
OpReturn
OpFunctionEnd
%d = OpFunction %void None %23
%f_1 = OpFunctionParameter %float
%26 = OpLabel
OpReturn
OpFunctionEnd
%f = OpFunction %void None %27
%29 = OpLabel
%33 = OpAccessChain %_ptr_Uniform__arr_mat4v2float_uint_4 %u %uint_0
%34 = OpLoad %_arr_mat4v2float_uint_4 %33
%30 = OpFunctionCall %void %a %34
%39 = OpAccessChain %_ptr_Uniform_mat4v2float %u %uint_0 %int_1
%40 = OpLoad %mat4v2float %39
%35 = OpFunctionCall %void %b %40
%44 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %int_1 %42
%45 = OpLoad %v2float %44
%46 = OpVectorShuffle %v2float %45 %45 1 0
%41 = OpFunctionCall %void %c %46
%48 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %int_1 %42
%49 = OpLoad %v2float %48
%50 = OpVectorShuffle %v2float %49 %49 1 0
%51 = OpCompositeExtract %float %50 0
%47 = OpFunctionCall %void %d %51
OpReturn
OpFunctionEnd

21
test/tint/buffer/uniform/std140/array/mat4x2/to_fn.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,21 @@
@group(0) @binding(0) var<uniform> u : array<mat4x2<f32>, 4>;
fn a(a : array<mat4x2<f32>, 4>) {
}
fn b(m : mat4x2<f32>) {
}
fn c(v : vec2<f32>) {
}
fn d(f : f32) {
}
@compute @workgroup_size(1)
fn f() {
a(u);
b(u[1]);
c(u[1][0].yx);
d(u[1][0].yx.x);
}

16
test/tint/buffer/uniform/std140/array/mat4x2/to_private.wgsl Normal file
View File

@ -0,0 +1,16 @@
struct S {
before : i32,
m : mat4x2<f32>,
after : i32,
}
@group(0) @binding(0) var<uniform> u : array<mat4x2<f32>, 4>;
var<private> p : array<mat4x2<f32>, 4>;
@compute @workgroup_size(1)
fn f() {
p = u;
p[1] = u[2];
p[1][0] = u[0][1].yx;
p[1][0].x = u[0][1].x;
}

36
test/tint/buffer/uniform/std140/array/mat4x2/to_private.wgsl.expected.dxc.hlsl Normal file
View File

@ -0,0 +1,36 @@
cbuffer cbuffer_u : register(b0, space0) {
uint4 u[8];
};
static float4x2 p[4] = (float4x2[4])0;
float4x2 tint_symbol_1(uint4 buffer[8], uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
uint4 ubo_load = buffer[scalar_offset / 4];
const uint scalar_offset_1 = ((offset + 8u)) / 4;
uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
const uint scalar_offset_2 = ((offset + 16u)) / 4;
uint4 ubo_load_2 = buffer[scalar_offset_2 / 4];
const uint scalar_offset_3 = ((offset + 24u)) / 4;
uint4 ubo_load_3 = buffer[scalar_offset_3 / 4];
return float4x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)), asfloat(((scalar_offset_2 & 2) ? ubo_load_2.zw : ubo_load_2.xy)), asfloat(((scalar_offset_3 & 2) ? ubo_load_3.zw : ubo_load_3.xy)));
}
typedef float4x2 tint_symbol_ret[4];
tint_symbol_ret tint_symbol(uint4 buffer[8], uint offset) {
float4x2 arr[4] = (float4x2[4])0;
{
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
arr[i] = tint_symbol_1(buffer, (offset + (i * 32u)));
}
}
return arr;
}
[numthreads(1, 1, 1)]
void f() {
p = tint_symbol(u, 0u);
p[1] = tint_symbol_1(u, 64u);
p[1][0] = asfloat(u[0].zw).yx;
p[1][0].x = asfloat(u[0].z);
return;
}

36
test/tint/buffer/uniform/std140/array/mat4x2/to_private.wgsl.expected.fxc.hlsl Normal file
View File

@ -0,0 +1,36 @@
cbuffer cbuffer_u : register(b0, space0) {
uint4 u[8];
};
static float4x2 p[4] = (float4x2[4])0;
float4x2 tint_symbol_1(uint4 buffer[8], uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
uint4 ubo_load = buffer[scalar_offset / 4];
const uint scalar_offset_1 = ((offset + 8u)) / 4;
uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
const uint scalar_offset_2 = ((offset + 16u)) / 4;
uint4 ubo_load_2 = buffer[scalar_offset_2 / 4];
const uint scalar_offset_3 = ((offset + 24u)) / 4;
uint4 ubo_load_3 = buffer[scalar_offset_3 / 4];
return float4x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)), asfloat(((scalar_offset_2 & 2) ? ubo_load_2.zw : ubo_load_2.xy)), asfloat(((scalar_offset_3 & 2) ? ubo_load_3.zw : ubo_load_3.xy)));
}
typedef float4x2 tint_symbol_ret[4];
tint_symbol_ret tint_symbol(uint4 buffer[8], uint offset) {
float4x2 arr[4] = (float4x2[4])0;
{
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
arr[i] = tint_symbol_1(buffer, (offset + (i * 32u)));
}
}
return arr;
}
[numthreads(1, 1, 1)]
void f() {
p = tint_symbol(u, 0u);
p[1] = tint_symbol_1(u, 64u);
p[1][0] = asfloat(u[0].zw).yx;
p[1][0].x = asfloat(u[0].z);
return;
}

25
test/tint/buffer/uniform/std140/array/mat4x2/to_private.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,25 @@
#version 310 es
struct S {
int before;
mat4x2 m;
int after;
};
layout(binding = 0, std140) uniform u_block_ubo {
mat4x2 inner[4];
} u;
mat4x2 p[4] = mat4x2[4](mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f), mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f), mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f), mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f));
void f() {
p = u.inner;
p[1] = u.inner[2];
p[1][0] = u.inner[0][1].yx;
p[1][0].x = u.inner[0][1].x;
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
f();
return;
}

31
test/tint/buffer/uniform/std140/array/mat4x2/to_private.wgsl.expected.msl Normal file
View File

@ -0,0 +1,31 @@
#include <metal_stdlib>
using namespace metal;
template<typename T, size_t N>
struct tint_array {
const constant T& operator[](size_t i) const constant { return elements[i]; }
device T& operator[](size_t i) device { return elements[i]; }
const device T& operator[](size_t i) const device { return elements[i]; }
thread T& operator[](size_t i) thread { return elements[i]; }
const thread T& operator[](size_t i) const thread { return elements[i]; }
threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
T elements[N];
};
struct S {
int before;
float4x2 m;
int after;
};
kernel void f(const constant tint_array<float4x2, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<float4x2, 4> tint_symbol = {};
tint_symbol = *(tint_symbol_1);
tint_symbol[1] = (*(tint_symbol_1))[2];
tint_symbol[1][0] = float2((*(tint_symbol_1))[0][1]).yx;
tint_symbol[1][0][0] = (*(tint_symbol_1))[0][1][0];
return;
}

68
test/tint/buffer/uniform/std140/array/mat4x2/to_private.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,68 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 41
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %u_block "u_block"
OpMemberName %u_block 0 "inner"
OpName %u "u"
OpName %p "p"
OpName %f "f"
OpDecorate %u_block Block
OpMemberDecorate %u_block 0 Offset 0
OpMemberDecorate %u_block 0 ColMajor
OpMemberDecorate %u_block 0 MatrixStride 8
OpDecorate %_arr_mat4v2float_uint_4 ArrayStride 32
OpDecorate %u NonWritable
OpDecorate %u DescriptorSet 0
OpDecorate %u Binding 0
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%mat4v2float = OpTypeMatrix %v2float 4
%uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4
%_arr_mat4v2float_uint_4 = OpTypeArray %mat4v2float %uint_4
%u_block = OpTypeStruct %_arr_mat4v2float_uint_4
%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
%u = OpVariable %_ptr_Uniform_u_block Uniform
%_ptr_Private__arr_mat4v2float_uint_4 = OpTypePointer Private %_arr_mat4v2float_uint_4
%12 = OpConstantNull %_arr_mat4v2float_uint_4
%p = OpVariable %_ptr_Private__arr_mat4v2float_uint_4 Private %12
%void = OpTypeVoid
%13 = OpTypeFunction %void
%uint_0 = OpConstant %uint 0
%_ptr_Uniform__arr_mat4v2float_uint_4 = OpTypePointer Uniform %_arr_mat4v2float_uint_4
%int = OpTypeInt 32 1
%int_1 = OpConstant %int 1
%_ptr_Private_mat4v2float = OpTypePointer Private %mat4v2float
%int_2 = OpConstant %int 2
%_ptr_Uniform_mat4v2float = OpTypePointer Uniform %mat4v2float
%29 = OpConstantNull %int
%_ptr_Private_v2float = OpTypePointer Private %v2float
%_ptr_Uniform_v2float = OpTypePointer Uniform %v2float
%_ptr_Private_float = OpTypePointer Private %float
%_ptr_Uniform_float = OpTypePointer Uniform %float
%f = OpFunction %void None %13
%16 = OpLabel
%19 = OpAccessChain %_ptr_Uniform__arr_mat4v2float_uint_4 %u %uint_0
%20 = OpLoad %_arr_mat4v2float_uint_4 %19
OpStore %p %20
%24 = OpAccessChain %_ptr_Private_mat4v2float %p %int_1
%27 = OpAccessChain %_ptr_Uniform_mat4v2float %u %uint_0 %int_2
%28 = OpLoad %mat4v2float %27
OpStore %24 %28
%31 = OpAccessChain %_ptr_Private_v2float %p %int_1 %29
%33 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %29 %int_1
%34 = OpLoad %v2float %33
%35 = OpVectorShuffle %v2float %34 %34 1 0
OpStore %31 %35
%37 = OpAccessChain %_ptr_Private_float %p %int_1 %29 %uint_0
%39 = OpAccessChain %_ptr_Uniform_float %u %uint_0 %29 %int_1 %uint_0
%40 = OpLoad %float %39
OpStore %37 %40
OpReturn
OpFunctionEnd

17
test/tint/buffer/uniform/std140/array/mat4x2/to_private.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,17 @@
struct S {
before : i32,
m : mat4x2<f32>,
after : i32,
}
@group(0) @binding(0) var<uniform> u : array<mat4x2<f32>, 4>;
var<private> p : array<mat4x2<f32>, 4>;
@compute @workgroup_size(1)
fn f() {
p = u;
p[1] = u[2];
p[1][0] = u[0][1].yx;
p[1][0].x = u[0][1].x;
}

16
test/tint/buffer/uniform/std140/array/mat4x2/to_storage.wgsl Normal file
View File

@ -0,0 +1,16 @@
struct S {
before : i32,
m : mat4x2<f32>,
after : i32,
}
@group(0) @binding(0) var<uniform> u : array<mat4x2<f32>, 4>;
@group(0) @binding(1) var<storage, read_write> s : array<mat4x2<f32>, 4>;
@compute @workgroup_size(1)
fn f() {
s = u;
s[1] = u[2];
s[1][0] = u[0][1].yx;
s[1][0].x = u[0][1].x;
}

52
test/tint/buffer/uniform/std140/array/mat4x2/to_storage.wgsl.expected.dxc.hlsl Normal file
View File

@ -0,0 +1,52 @@
cbuffer cbuffer_u : register(b0, space0) {
uint4 u[8];
};
RWByteAddressBuffer s : register(u1, space0);
void tint_symbol_1(RWByteAddressBuffer buffer, uint offset, float4x2 value) {
buffer.Store2((offset + 0u), asuint(value[0u]));
buffer.Store2((offset + 8u), asuint(value[1u]));
buffer.Store2((offset + 16u), asuint(value[2u]));
buffer.Store2((offset + 24u), asuint(value[3u]));
}
void tint_symbol(RWByteAddressBuffer buffer, uint offset, float4x2 value[4]) {
float4x2 array[4] = value;
{
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
tint_symbol_1(buffer, (offset + (i * 32u)), array[i]);
}
}
}
float4x2 tint_symbol_4(uint4 buffer[8], uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
uint4 ubo_load = buffer[scalar_offset / 4];
const uint scalar_offset_1 = ((offset + 8u)) / 4;
uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
const uint scalar_offset_2 = ((offset + 16u)) / 4;
uint4 ubo_load_2 = buffer[scalar_offset_2 / 4];
const uint scalar_offset_3 = ((offset + 24u)) / 4;
uint4 ubo_load_3 = buffer[scalar_offset_3 / 4];
return float4x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)), asfloat(((scalar_offset_2 & 2) ? ubo_load_2.zw : ubo_load_2.xy)), asfloat(((scalar_offset_3 & 2) ? ubo_load_3.zw : ubo_load_3.xy)));
}
typedef float4x2 tint_symbol_3_ret[4];
tint_symbol_3_ret tint_symbol_3(uint4 buffer[8], uint offset) {
float4x2 arr[4] = (float4x2[4])0;
{
[loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
arr[i_1] = tint_symbol_4(buffer, (offset + (i_1 * 32u)));
}
}
return arr;
}
[numthreads(1, 1, 1)]
void f() {
tint_symbol(s, 0u, tint_symbol_3(u, 0u));
tint_symbol_1(s, 32u, tint_symbol_4(u, 64u));
s.Store2(32u, asuint(asfloat(u[0].zw).yx));
s.Store(32u, asuint(asfloat(u[0].z)));
return;
}

52
test/tint/buffer/uniform/std140/array/mat4x2/to_storage.wgsl.expected.fxc.hlsl Normal file
View File

@ -0,0 +1,52 @@
cbuffer cbuffer_u : register(b0, space0) {
uint4 u[8];
};
RWByteAddressBuffer s : register(u1, space0);
void tint_symbol_1(RWByteAddressBuffer buffer, uint offset, float4x2 value) {
buffer.Store2((offset + 0u), asuint(value[0u]));
buffer.Store2((offset + 8u), asuint(value[1u]));
buffer.Store2((offset + 16u), asuint(value[2u]));
buffer.Store2((offset + 24u), asuint(value[3u]));
}
void tint_symbol(RWByteAddressBuffer buffer, uint offset, float4x2 value[4]) {
float4x2 array[4] = value;
{
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
tint_symbol_1(buffer, (offset + (i * 32u)), array[i]);
}
}
}
float4x2 tint_symbol_4(uint4 buffer[8], uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
uint4 ubo_load = buffer[scalar_offset / 4];
const uint scalar_offset_1 = ((offset + 8u)) / 4;
uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
const uint scalar_offset_2 = ((offset + 16u)) / 4;
uint4 ubo_load_2 = buffer[scalar_offset_2 / 4];
const uint scalar_offset_3 = ((offset + 24u)) / 4;
uint4 ubo_load_3 = buffer[scalar_offset_3 / 4];
return float4x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)), asfloat(((scalar_offset_2 & 2) ? ubo_load_2.zw : ubo_load_2.xy)), asfloat(((scalar_offset_3 & 2) ? ubo_load_3.zw : ubo_load_3.xy)));
}
typedef float4x2 tint_symbol_3_ret[4];
tint_symbol_3_ret tint_symbol_3(uint4 buffer[8], uint offset) {
float4x2 arr[4] = (float4x2[4])0;
{
[loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
arr[i_1] = tint_symbol_4(buffer, (offset + (i_1 * 32u)));
}
}
return arr;
}
[numthreads(1, 1, 1)]
void f() {
tint_symbol(s, 0u, tint_symbol_3(u, 0u));
tint_symbol_1(s, 32u, tint_symbol_4(u, 64u));
s.Store2(32u, asuint(asfloat(u[0].zw).yx));
s.Store(32u, asuint(asfloat(u[0].z)));
return;
}

28
test/tint/buffer/uniform/std140/array/mat4x2/to_storage.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,28 @@
#version 310 es
struct S {
int before;
mat4x2 m;
int after;
};
layout(binding = 0, std140) uniform u_block_ubo {
mat4x2 inner[4];
} u;
layout(binding = 1, std430) buffer u_block_ssbo {
mat4x2 inner[4];
} s;
void f() {
s.inner = u.inner;
s.inner[1] = u.inner[2];
s.inner[1][0] = u.inner[0][1].yx;
s.inner[1][0].x = u.inner[0][1].x;
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
f();
return;
}

30
test/tint/buffer/uniform/std140/array/mat4x2/to_storage.wgsl.expected.msl Normal file
View File

@ -0,0 +1,30 @@
#include <metal_stdlib>
using namespace metal;
template<typename T, size_t N>
struct tint_array {
const constant T& operator[](size_t i) const constant { return elements[i]; }
device T& operator[](size_t i) device { return elements[i]; }
const device T& operator[](size_t i) const device { return elements[i]; }
thread T& operator[](size_t i) thread { return elements[i]; }
const thread T& operator[](size_t i) const thread { return elements[i]; }
threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
T elements[N];
};
struct S {
int before;
float4x2 m;
int after;
};
kernel void f(device tint_array<float4x2, 4>* tint_symbol [[buffer(1)]], const constant tint_array<float4x2, 4>* tint_symbol_1 [[buffer(0)]]) {
*(tint_symbol) = *(tint_symbol_1);
(*(tint_symbol))[1] = (*(tint_symbol_1))[2];
(*(tint_symbol))[1][0] = float2((*(tint_symbol_1))[0][1]).yx;
(*(tint_symbol))[1][0][0] = (*(tint_symbol_1))[0][1][0];
return;
}

71
test/tint/buffer/uniform/std140/array/mat4x2/to_storage.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,71 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 42
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %u_block "u_block"
OpMemberName %u_block 0 "inner"
OpName %u "u"
OpName %s "s"
OpName %f "f"
OpDecorate %u_block Block
OpMemberDecorate %u_block 0 Offset 0
OpMemberDecorate %u_block 0 ColMajor
OpMemberDecorate %u_block 0 MatrixStride 8
OpDecorate %_arr_mat4v2float_uint_4 ArrayStride 32
OpDecorate %u NonWritable
OpDecorate %u DescriptorSet 0
OpDecorate %u Binding 0
OpDecorate %s DescriptorSet 0
OpDecorate %s Binding 1
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%mat4v2float = OpTypeMatrix %v2float 4
%uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4
%_arr_mat4v2float_uint_4 = OpTypeArray %mat4v2float %uint_4
%u_block = OpTypeStruct %_arr_mat4v2float_uint_4
%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
%u = OpVariable %_ptr_Uniform_u_block Uniform
%_ptr_StorageBuffer_u_block = OpTypePointer StorageBuffer %u_block
%s = OpVariable %_ptr_StorageBuffer_u_block StorageBuffer
%void = OpTypeVoid
%12 = OpTypeFunction %void
%uint_0 = OpConstant %uint 0
%_ptr_StorageBuffer__arr_mat4v2float_uint_4 = OpTypePointer StorageBuffer %_arr_mat4v2float_uint_4
%_ptr_Uniform__arr_mat4v2float_uint_4 = OpTypePointer Uniform %_arr_mat4v2float_uint_4
%int = OpTypeInt 32 1
%int_1 = OpConstant %int 1
%_ptr_StorageBuffer_mat4v2float = OpTypePointer StorageBuffer %mat4v2float
%int_2 = OpConstant %int 2
%_ptr_Uniform_mat4v2float = OpTypePointer Uniform %mat4v2float
%30 = OpConstantNull %int
%_ptr_StorageBuffer_v2float = OpTypePointer StorageBuffer %v2float
%_ptr_Uniform_v2float = OpTypePointer Uniform %v2float
%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
%_ptr_Uniform_float = OpTypePointer Uniform %float
%f = OpFunction %void None %12
%15 = OpLabel
%18 = OpAccessChain %_ptr_StorageBuffer__arr_mat4v2float_uint_4 %s %uint_0
%20 = OpAccessChain %_ptr_Uniform__arr_mat4v2float_uint_4 %u %uint_0
%21 = OpLoad %_arr_mat4v2float_uint_4 %20
OpStore %18 %21
%25 = OpAccessChain %_ptr_StorageBuffer_mat4v2float %s %uint_0 %int_1
%28 = OpAccessChain %_ptr_Uniform_mat4v2float %u %uint_0 %int_2
%29 = OpLoad %mat4v2float %28
OpStore %25 %29
%32 = OpAccessChain %_ptr_StorageBuffer_v2float %s %uint_0 %int_1 %30
%34 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %30 %int_1
%35 = OpLoad %v2float %34
%36 = OpVectorShuffle %v2float %35 %35 1 0
OpStore %32 %36
%38 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_0 %int_1 %30 %uint_0
%40 = OpAccessChain %_ptr_Uniform_float %u %uint_0 %30 %int_1 %uint_0
%41 = OpLoad %float %40
OpStore %38 %41
OpReturn
OpFunctionEnd

17
test/tint/buffer/uniform/std140/array/mat4x2/to_storage.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,17 @@
struct S {
before : i32,
m : mat4x2<f32>,
after : i32,
}
@group(0) @binding(0) var<uniform> u : array<mat4x2<f32>, 4>;
@group(0) @binding(1) var<storage, read_write> s : array<mat4x2<f32>, 4>;
@compute @workgroup_size(1)
fn f() {
s = u;
s[1] = u[2];
s[1][0] = u[0][1].yx;
s[1][0].x = u[0][1].x;
}

16
test/tint/buffer/uniform/std140/array/mat4x2/to_workgroup.wgsl Normal file
View File

@ -0,0 +1,16 @@
struct S {
before : i32,
m : mat4x2<f32>,
after : i32,
}
@group(0) @binding(0) var<uniform> u : array<mat4x2<f32>, 4>;
var<workgroup> w : array<mat4x2<f32>, 4>;
@compute @workgroup_size(1)
fn f() {
w = u;
w[1] = u[2];
w[1][0] = u[0][1].yx;
w[1][0].x = u[0][1].x;
}

51
test/tint/buffer/uniform/std140/array/mat4x2/to_workgroup.wgsl.expected.dxc.hlsl Normal file
View File

@ -0,0 +1,51 @@
cbuffer cbuffer_u : register(b0, space0) {
uint4 u[8];
};
groupshared float4x2 w[4];
struct tint_symbol_1 {
uint local_invocation_index : SV_GroupIndex;
};
float4x2 tint_symbol_3(uint4 buffer[8], uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
uint4 ubo_load = buffer[scalar_offset / 4];
const uint scalar_offset_1 = ((offset + 8u)) / 4;
uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
const uint scalar_offset_2 = ((offset + 16u)) / 4;
uint4 ubo_load_2 = buffer[scalar_offset_2 / 4];
const uint scalar_offset_3 = ((offset + 24u)) / 4;
uint4 ubo_load_3 = buffer[scalar_offset_3 / 4];
return float4x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)), asfloat(((scalar_offset_2 & 2) ? ubo_load_2.zw : ubo_load_2.xy)), asfloat(((scalar_offset_3 & 2) ? ubo_load_3.zw : ubo_load_3.xy)));
}
typedef float4x2 tint_symbol_2_ret[4];
tint_symbol_2_ret tint_symbol_2(uint4 buffer[8], uint offset) {
float4x2 arr[4] = (float4x2[4])0;
{
[loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
arr[i_1] = tint_symbol_3(buffer, (offset + (i_1 * 32u)));
}
}
return arr;
}
void f_inner(uint local_invocation_index) {
{
[loop] for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
const uint i = idx;
w[i] = float4x2((0.0f).xx, (0.0f).xx, (0.0f).xx, (0.0f).xx);
}
}
GroupMemoryBarrierWithGroupSync();
w = tint_symbol_2(u, 0u);
w[1] = tint_symbol_3(u, 64u);
w[1][0] = asfloat(u[0].zw).yx;
w[1][0].x = asfloat(u[0].z);
}
[numthreads(1, 1, 1)]
void f(tint_symbol_1 tint_symbol) {
f_inner(tint_symbol.local_invocation_index);
return;
}

51
test/tint/buffer/uniform/std140/array/mat4x2/to_workgroup.wgsl.expected.fxc.hlsl Normal file
View File

@ -0,0 +1,51 @@
cbuffer cbuffer_u : register(b0, space0) {
uint4 u[8];
};
groupshared float4x2 w[4];
struct tint_symbol_1 {
uint local_invocation_index : SV_GroupIndex;
};
float4x2 tint_symbol_3(uint4 buffer[8], uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
uint4 ubo_load = buffer[scalar_offset / 4];
const uint scalar_offset_1 = ((offset + 8u)) / 4;
uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
const uint scalar_offset_2 = ((offset + 16u)) / 4;
uint4 ubo_load_2 = buffer[scalar_offset_2 / 4];
const uint scalar_offset_3 = ((offset + 24u)) / 4;
uint4 ubo_load_3 = buffer[scalar_offset_3 / 4];
return float4x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)), asfloat(((scalar_offset_2 & 2) ? ubo_load_2.zw : ubo_load_2.xy)), asfloat(((scalar_offset_3 & 2) ? ubo_load_3.zw : ubo_load_3.xy)));
}
typedef float4x2 tint_symbol_2_ret[4];
tint_symbol_2_ret tint_symbol_2(uint4 buffer[8], uint offset) {
float4x2 arr[4] = (float4x2[4])0;
{
[loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
arr[i_1] = tint_symbol_3(buffer, (offset + (i_1 * 32u)));
}
}
return arr;
}
void f_inner(uint local_invocation_index) {
{
[loop] for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
const uint i = idx;
w[i] = float4x2((0.0f).xx, (0.0f).xx, (0.0f).xx, (0.0f).xx);
}
}
GroupMemoryBarrierWithGroupSync();
w = tint_symbol_2(u, 0u);
w[1] = tint_symbol_3(u, 64u);
w[1][0] = asfloat(u[0].zw).yx;
w[1][0].x = asfloat(u[0].z);
}
[numthreads(1, 1, 1)]
void f(tint_symbol_1 tint_symbol) {
f_inner(tint_symbol.local_invocation_index);
return;
}

32
test/tint/buffer/uniform/std140/array/mat4x2/to_workgroup.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,32 @@
#version 310 es
struct S {
int before;
mat4x2 m;
int after;
};
layout(binding = 0, std140) uniform u_block_ubo {
mat4x2 inner[4];
} u;
shared mat4x2 w[4];
void f(uint local_invocation_index) {
{
for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint i = idx;
w[i] = mat4x2(vec2(0.0f), vec2(0.0f), vec2(0.0f), vec2(0.0f));
}
}
barrier();
w = u.inner;
w[1] = u.inner[2];
w[1][0] = u.inner[0][1].yx;
w[1][0].x = u.inner[0][1].x;
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
f(gl_LocalInvocationIndex);
return;
}

44
test/tint/buffer/uniform/std140/array/mat4x2/to_workgroup.wgsl.expected.msl Normal file
View File

@ -0,0 +1,44 @@
#include <metal_stdlib>
using namespace metal;
template<typename T, size_t N>
struct tint_array {
const constant T& operator[](size_t i) const constant { return elements[i]; }
device T& operator[](size_t i) device { return elements[i]; }
const device T& operator[](size_t i) const device { return elements[i]; }
thread T& operator[](size_t i) thread { return elements[i]; }
const thread T& operator[](size_t i) const thread { return elements[i]; }
threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
T elements[N];
};
struct tint_symbol_5 {
tint_array<float4x2, 4> w;
};
struct S {
int before;
float4x2 m;
int after;
};
void f_inner(uint local_invocation_index, threadgroup tint_array<float4x2, 4>* const tint_symbol, const constant tint_array<float4x2, 4>* const tint_symbol_1) {
for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
(*(tint_symbol))[i] = float4x2(float2(0.0f), float2(0.0f), float2(0.0f), float2(0.0f));
}
threadgroup_barrier(mem_flags::mem_threadgroup);
*(tint_symbol) = *(tint_symbol_1);
(*(tint_symbol))[1] = (*(tint_symbol_1))[2];
(*(tint_symbol))[1][0] = float2((*(tint_symbol_1))[0][1]).yx;
(*(tint_symbol))[1][0][0] = (*(tint_symbol_1))[0][1][0];
}
kernel void f(const constant tint_array<float4x2, 4>* tint_symbol_4 [[buffer(0)]], threadgroup tint_symbol_5* tint_symbol_3 [[threadgroup(0)]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
threadgroup tint_array<float4x2, 4>* const tint_symbol_2 = &((*(tint_symbol_3)).w);
f_inner(local_invocation_index, tint_symbol_2, tint_symbol_4);
return;
}

115
test/tint/buffer/uniform/std140/array/mat4x2/to_workgroup.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,115 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 70
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f" %local_invocation_index_1
OpExecutionMode %f LocalSize 1 1 1
OpName %local_invocation_index_1 "local_invocation_index_1"
OpName %u_block "u_block"
OpMemberName %u_block 0 "inner"
OpName %u "u"
OpName %w "w"
OpName %f_inner "f_inner"
OpName %local_invocation_index "local_invocation_index"
OpName %idx "idx"
OpName %f "f"
OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
OpDecorate %u_block Block
OpMemberDecorate %u_block 0 Offset 0
OpMemberDecorate %u_block 0 ColMajor
OpMemberDecorate %u_block 0 MatrixStride 8
OpDecorate %_arr_mat4v2float_uint_4 ArrayStride 32
OpDecorate %u NonWritable
OpDecorate %u DescriptorSet 0
OpDecorate %u Binding 0
%uint = OpTypeInt 32 0
%_ptr_Input_uint = OpTypePointer Input %uint
%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%mat4v2float = OpTypeMatrix %v2float 4
%uint_4 = OpConstant %uint 4
%_arr_mat4v2float_uint_4 = OpTypeArray %mat4v2float %uint_4
%u_block = OpTypeStruct %_arr_mat4v2float_uint_4
%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
%u = OpVariable %_ptr_Uniform_u_block Uniform
%_ptr_Workgroup__arr_mat4v2float_uint_4 = OpTypePointer Workgroup %_arr_mat4v2float_uint_4
%w = OpVariable %_ptr_Workgroup__arr_mat4v2float_uint_4 Workgroup
%void = OpTypeVoid
%14 = OpTypeFunction %void %uint
%_ptr_Function_uint = OpTypePointer Function %uint
%21 = OpConstantNull %uint
%bool = OpTypeBool
%_ptr_Workgroup_mat4v2float = OpTypePointer Workgroup %mat4v2float
%35 = OpConstantNull %mat4v2float
%uint_1 = OpConstant %uint 1
%uint_2 = OpConstant %uint 2
%uint_264 = OpConstant %uint 264
%uint_0 = OpConstant %uint 0
%_ptr_Uniform__arr_mat4v2float_uint_4 = OpTypePointer Uniform %_arr_mat4v2float_uint_4
%int = OpTypeInt 32 1
%int_1 = OpConstant %int 1
%int_2 = OpConstant %int 2
%_ptr_Uniform_mat4v2float = OpTypePointer Uniform %mat4v2float
%53 = OpConstantNull %int
%_ptr_Workgroup_v2float = OpTypePointer Workgroup %v2float
%_ptr_Uniform_v2float = OpTypePointer Uniform %v2float
%_ptr_Workgroup_float = OpTypePointer Workgroup %float
%_ptr_Uniform_float = OpTypePointer Uniform %float
%65 = OpTypeFunction %void
%f_inner = OpFunction %void None %14
%local_invocation_index = OpFunctionParameter %uint
%18 = OpLabel
%idx = OpVariable %_ptr_Function_uint Function %21
OpStore %idx %local_invocation_index
OpBranch %22
%22 = OpLabel
OpLoopMerge %23 %24 None
OpBranch %25
%25 = OpLabel
%27 = OpLoad %uint %idx
%28 = OpULessThan %bool %27 %uint_4
%26 = OpLogicalNot %bool %28
OpSelectionMerge %30 None
OpBranchConditional %26 %31 %30
%31 = OpLabel
OpBranch %23
%30 = OpLabel
%32 = OpLoad %uint %idx
%34 = OpAccessChain %_ptr_Workgroup_mat4v2float %w %32
OpStore %34 %35
OpBranch %24
%24 = OpLabel
%36 = OpLoad %uint %idx
%38 = OpIAdd %uint %36 %uint_1
OpStore %idx %38
OpBranch %22
%23 = OpLabel
OpControlBarrier %uint_2 %uint_2 %uint_264
%44 = OpAccessChain %_ptr_Uniform__arr_mat4v2float_uint_4 %u %uint_0
%45 = OpLoad %_arr_mat4v2float_uint_4 %44
OpStore %w %45
%48 = OpAccessChain %_ptr_Workgroup_mat4v2float %w %int_1
%51 = OpAccessChain %_ptr_Uniform_mat4v2float %u %uint_0 %int_2
%52 = OpLoad %mat4v2float %51
OpStore %48 %52
%55 = OpAccessChain %_ptr_Workgroup_v2float %w %int_1 %53
%57 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %53 %int_1
%58 = OpLoad %v2float %57
%59 = OpVectorShuffle %v2float %58 %58 1 0
OpStore %55 %59
%61 = OpAccessChain %_ptr_Workgroup_float %w %int_1 %53 %uint_0
%63 = OpAccessChain %_ptr_Uniform_float %u %uint_0 %53 %int_1 %uint_0
%64 = OpLoad %float %63
OpStore %61 %64
OpReturn
OpFunctionEnd
%f = OpFunction %void None %65
%67 = OpLabel
%69 = OpLoad %uint %local_invocation_index_1
%68 = OpFunctionCall %void %f_inner %69
OpReturn
OpFunctionEnd

17
test/tint/buffer/uniform/std140/array/mat4x2/to_workgroup.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,17 @@
struct S {
before : i32,
m : mat4x2<f32>,
after : i32,
}
@group(0) @binding(0) var<uniform> u : array<mat4x2<f32>, 4>;
var<workgroup> w : array<mat4x2<f32>, 4>;
@compute @workgroup_size(1)
fn f() {
w = u;
w[1] = u[2];
w[1][0] = u[0][1].yx;
w[1][0].x = u[0][1].x;
}

73
test/tint/buffer/uniform/std140/mat2x2/to_builtin.wgsl.expected.spvasm
View File

@ -1,73 +0,0 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 41
; Schema: 0
OpCapability Shader
%31 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %u_block "u_block"
OpMemberName %u_block 0 "inner"
OpName %S_std140 "S_std140"
OpMemberName %S_std140 0 "before"
OpMemberName %S_std140 1 "m_0"
OpMemberName %S_std140 2 "m_1"
OpMemberName %S_std140 3 "after"
OpName %u "u"
OpName %load_u_2_m "load_u_2_m"
OpName %f "f"
OpDecorate %u_block Block
OpMemberDecorate %u_block 0 Offset 0
OpMemberDecorate %S_std140 0 Offset 0
OpMemberDecorate %S_std140 1 Offset 8
OpMemberDecorate %S_std140 2 Offset 16
OpMemberDecorate %S_std140 3 Offset 24
OpDecorate %_arr_S_std140_uint_4 ArrayStride 32
OpDecorate %u NonWritable
OpDecorate %u DescriptorSet 0
OpDecorate %u Binding 0
%int = OpTypeInt 32 1
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%S_std140 = OpTypeStruct %int %v2float %v2float %int
%uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4
%_arr_S_std140_uint_4 = OpTypeArray %S_std140 %uint_4
%u_block = OpTypeStruct %_arr_S_std140_uint_4
%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
%u = OpVariable %_ptr_Uniform_u_block Uniform
%mat2v2float = OpTypeMatrix %v2float 2
%11 = OpTypeFunction %mat2v2float
%uint_0 = OpConstant %uint 0
%uint_2 = OpConstant %uint 2
%uint_1 = OpConstant %uint 1
%_ptr_Uniform_v2float = OpTypePointer Uniform %v2float
%void = OpTypeVoid
%24 = OpTypeFunction %void
%32 = OpConstantNull %uint
%load_u_2_m = OpFunction %mat2v2float None %11
%14 = OpLabel
%19 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %uint_2 %uint_1
%20 = OpLoad %v2float %19
%21 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %uint_2 %uint_2
%22 = OpLoad %v2float %21
%23 = OpCompositeConstruct %mat2v2float %20 %22
OpReturnValue %23
OpFunctionEnd
%f = OpFunction %void None %24
%27 = OpLabel
%29 = OpFunctionCall %mat2v2float %load_u_2_m
%28 = OpTranspose %mat2v2float %29
%33 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %32 %uint_2
%34 = OpLoad %v2float %33
%35 = OpVectorShuffle %v2float %34 %34 1 0
%30 = OpExtInst %float %31 Length %35
%37 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %32 %uint_2
%38 = OpLoad %v2float %37
%39 = OpVectorShuffle %v2float %38 %38 1 0
%40 = OpCompositeExtract %float %39 0
%36 = OpExtInst %float %31 FAbs %40
OpReturn
OpFunctionEnd

53
test/tint/buffer/uniform/std140/mat2x2/to_private.wgsl.expected.glsl
View File

@ -1,53 +0,0 @@
#version 310 es
struct S {
int before;
mat2 m;
int after;
};
struct S_std140 {
int before;
uint pad;
vec2 m_0;
vec2 m_1;
int after;
uint pad_1;
};
layout(binding = 0, std140) uniform u_block_ubo {
S_std140 inner[4];
} u;
S p[4] = S[4](S(0, mat2(0.0f, 0.0f, 0.0f, 0.0f), 0), S(0, mat2(0.0f, 0.0f, 0.0f, 0.0f), 0), S(0, mat2(0.0f, 0.0f, 0.0f, 0.0f), 0), S(0, mat2(0.0f, 0.0f, 0.0f, 0.0f), 0));
S conv_S(S_std140 val) {
S tint_symbol = S(val.before, mat2(val.m_0, val.m_1), val.after);
return tint_symbol;
}
S[4] conv_arr_4_S(S_std140 val[4]) {
S arr[4] = S[4](S(0, mat2(0.0f, 0.0f, 0.0f, 0.0f), 0), S(0, mat2(0.0f, 0.0f, 0.0f, 0.0f), 0), S(0, mat2(0.0f, 0.0f, 0.0f, 0.0f), 0), S(0, mat2(0.0f, 0.0f, 0.0f, 0.0f), 0));
{
for(uint i = 0u; (i < 4u); i = (i + 1u)) {
arr[i] = conv_S(val[i]);
}
}
return arr;
}
mat2 load_u_2_m() {
return mat2(u.inner[2u].m_0, u.inner[2u].m_1);
}
void f() {
p = conv_arr_4_S(u.inner);
p[1] = conv_S(u.inner[2u]);
p[3].m = load_u_2_m();
p[1].m[0] = u.inner[0u].m_1.yx;
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
f();
return;
}

144
test/tint/buffer/uniform/std140/mat3x2/dynamic_index_via_ptr.wgsl.expected.glsl
View File

@ -1,144 +0,0 @@
#version 310 es
struct Inner {
mat3x2 m;
};
struct Inner_std140 {
vec2 m_0;
vec2 m_1;
vec2 m_2;
uint pad;
uint pad_1;
uint pad_2;
uint pad_3;
uint pad_4;
uint pad_5;
uint pad_6;
uint pad_7;
uint pad_8;
uint pad_9;
};
struct Outer {
Inner a[4];
};
struct Outer_std140 {
Inner_std140 a[4];
};
layout(binding = 0, std140) uniform a_block_ubo {
Outer_std140 inner[4];
} a;
int counter = 0;
int i() {
counter = (counter + 1);
return counter;
}
Inner conv_Inner(Inner_std140 val) {
Inner tint_symbol_4 = Inner(mat3x2(val.m_0, val.m_1, val.m_2));
return tint_symbol_4;
}
Inner[4] conv_arr_4_Inner(Inner_std140 val[4]) {
Inner arr[4] = Inner[4](Inner(mat3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)));
{
for(uint i = 0u; (i < 4u); i = (i + 1u)) {
arr[i] = conv_Inner(val[i]);
}
}
return arr;
}
Outer conv_Outer(Outer_std140 val) {
Outer tint_symbol_5 = Outer(conv_arr_4_Inner(val.a));
return tint_symbol_5;
}
Outer[4] conv_arr_4_Outer(Outer_std140 val[4]) {
Outer arr[4] = Outer[4](Outer(Inner[4](Inner(mat3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)))), Outer(Inner[4](Inner(mat3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)))), Outer(Inner[4](Inner(mat3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)))), Outer(Inner[4](Inner(mat3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)))));
{
for(uint i = 0u; (i < 4u); i = (i + 1u)) {
arr[i] = conv_Outer(val[i]);
}
}
return arr;
}
mat3x2 load_a_p0_a_p1_m(uint p0, uint p1) {
uint s_save = p0;
uint s_save_1 = p1;
return mat3x2(a.inner[s_save].a[s_save_1].m_0, a.inner[s_save].a[s_save_1].m_1, a.inner[s_save].a[s_save_1].m_2);
}
vec2 load_a_p0_a_p1_m_p2(uint p0, uint p1, uint p2) {
switch(p2) {
case 0u: {
return a.inner[p0].a[p1].m_0;
break;
}
case 1u: {
return a.inner[p0].a[p1].m_1;
break;
}
case 2u: {
return a.inner[p0].a[p1].m_2;
break;
}
default: {
return vec2(0.0f);
break;
}
}
}
float load_a_p0_a_p1_m_p2_p3(uint p0, uint p1, uint p2, uint p3) {
switch(p2) {
case 0u: {
return a.inner[p0].a[p1].m_0[p3];
break;
}
case 1u: {
return a.inner[p0].a[p1].m_1[p3];
break;
}
case 2u: {
return a.inner[p0].a[p1].m_2[p3];
break;
}
default: {
return 0.0f;
break;
}
}
}
void f() {
int I = 1;
Outer p_a[4] = conv_arr_4_Outer(a.inner);
int tint_symbol = i();
Outer p_a_i = conv_Outer(a.inner[tint_symbol]);
Inner p_a_i_a[4] = conv_arr_4_Inner(a.inner[tint_symbol].a);
int tint_symbol_1 = i();
Inner p_a_i_a_i = conv_Inner(a.inner[tint_symbol].a[tint_symbol_1]);
mat3x2 p_a_i_a_i_m = load_a_p0_a_p1_m(uint(tint_symbol), uint(tint_symbol_1));
int tint_symbol_2 = i();
vec2 p_a_i_a_i_m_i = load_a_p0_a_p1_m_p2(uint(tint_symbol), uint(tint_symbol_1), uint(tint_symbol_2));
Outer l_a[4] = conv_arr_4_Outer(a.inner);
Outer l_a_i = conv_Outer(a.inner[tint_symbol]);
Inner l_a_i_a[4] = conv_arr_4_Inner(a.inner[tint_symbol].a);
Inner l_a_i_a_i = conv_Inner(a.inner[tint_symbol].a[tint_symbol_1]);
mat3x2 l_a_i_a_i_m = load_a_p0_a_p1_m(uint(tint_symbol), uint(tint_symbol_1));
vec2 l_a_i_a_i_m_i = load_a_p0_a_p1_m_p2(uint(tint_symbol), uint(tint_symbol_1), uint(tint_symbol_2));
int tint_symbol_3 = i();
float l_a_i_a_i_m_i_i = load_a_p0_a_p1_m_p2_p3(uint(tint_symbol), uint(tint_symbol_1), uint(tint_symbol_2), uint(tint_symbol_3));
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
f();
return;
}

Some files were not shown because too many files have changed in this diff Show More