mirror of
https://github.com/encounter/dawn-cmake.git
synced 2025-12-16 00:17:03 +00:00
tint: Add and use new Std140 transform
This transform breaks up matNx2<f32> matrices used in uniform buffers into column vectors, which fixes std140 layout rules. Used by the SPIR-V and GLSL backends. Re-enable tests that were disabled for these cases. Bug: tint:1632 Change-Id: I596d016582b4189a0b413d762b3e7eabd3504b22 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/100907 Kokoro: Kokoro <noreply+kokoro@google.com> Reviewed-by: Dan Sinclair <dsinclair@chromium.org> Commit-Queue: Ben Clayton <bclayton@chromium.org>
This commit is contained in:
Ben Clayton
committed by
Dawn LUCI CQ
Dawn LUCI CQ
parent
644d23b3f8
commit
657e61d43d
@@ -0,0 +1,33 @@
|
||||
struct Inner {
|
||||
m : mat4x2<f32>,
|
||||
}
|
||||
|
||||
struct Outer {
|
||||
a : array<Inner, 4>,
|
||||
}
|
||||
|
||||
@group(0) @binding(0) var<uniform> a : array<Outer, 4>;
|
||||
|
||||
var<private> counter = 0;
|
||||
fn i() -> i32 { counter++; return counter; }
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn f() {
|
||||
let I = 1;
|
||||
|
||||
let p_a = &a;
|
||||
let p_a_i = &((*p_a)[i()]);
|
||||
let p_a_i_a = &((*p_a_i).a);
|
||||
let p_a_i_a_i = &((*p_a_i_a)[i()]);
|
||||
let p_a_i_a_i_m = &((*p_a_i_a_i).m);
|
||||
let p_a_i_a_i_m_i = &((*p_a_i_a_i_m)[i()]);
|
||||
|
||||
|
||||
let l_a : array<Outer, 4> = *p_a;
|
||||
let l_a_i : Outer = *p_a_i;
|
||||
let l_a_i_a : array<Inner, 4> = *p_a_i_a;
|
||||
let l_a_i_a_i : Inner = *p_a_i_a_i;
|
||||
let l_a_i_a_i_m : mat4x2<f32> = *p_a_i_a_i_m;
|
||||
let l_a_i_a_i_m_i : vec2<f32> = *p_a_i_a_i_m_i;
|
||||
let l_a_i_a_i_m_i_i : f32 = (*p_a_i_a_i_m_i)[i()];
|
||||
}
|
||||
@@ -0,0 +1,83 @@
|
||||
struct Inner {
|
||||
float4x2 m;
|
||||
};
|
||||
struct Outer {
|
||||
Inner a[4];
|
||||
};
|
||||
|
||||
cbuffer cbuffer_a : register(b0, space0) {
|
||||
uint4 a[32];
|
||||
};
|
||||
static int counter = 0;
|
||||
|
||||
int i() {
|
||||
counter = (counter + 1);
|
||||
return counter;
|
||||
}
|
||||
|
||||
float4x2 tint_symbol_8(uint4 buffer[32], 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)));
|
||||
}
|
||||
|
||||
Inner tint_symbol_7(uint4 buffer[32], uint offset) {
|
||||
const Inner tint_symbol_11 = {tint_symbol_8(buffer, (offset + 0u))};
|
||||
return tint_symbol_11;
|
||||
}
|
||||
|
||||
typedef Inner tint_symbol_6_ret[4];
|
||||
tint_symbol_6_ret tint_symbol_6(uint4 buffer[32], uint offset) {
|
||||
Inner arr[4] = (Inner[4])0;
|
||||
{
|
||||
[loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
|
||||
arr[i_1] = tint_symbol_7(buffer, (offset + (i_1 * 32u)));
|
||||
}
|
||||
}
|
||||
return arr;
|
||||
}
|
||||
|
||||
Outer tint_symbol_5(uint4 buffer[32], uint offset) {
|
||||
const Outer tint_symbol_12 = {tint_symbol_6(buffer, (offset + 0u))};
|
||||
return tint_symbol_12;
|
||||
}
|
||||
|
||||
typedef Outer tint_symbol_4_ret[4];
|
||||
tint_symbol_4_ret tint_symbol_4(uint4 buffer[32], uint offset) {
|
||||
Outer arr_1[4] = (Outer[4])0;
|
||||
{
|
||||
[loop] for(uint i_2 = 0u; (i_2 < 4u); i_2 = (i_2 + 1u)) {
|
||||
arr_1[i_2] = tint_symbol_5(buffer, (offset + (i_2 * 128u)));
|
||||
}
|
||||
}
|
||||
return arr_1;
|
||||
}
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void f() {
|
||||
const int I = 1;
|
||||
const int p_a_i_save = i();
|
||||
const int p_a_i_a_i_save = i();
|
||||
const int p_a_i_a_i_m_i_save = i();
|
||||
const Outer l_a[4] = tint_symbol_4(a, 0u);
|
||||
const Outer l_a_i = tint_symbol_5(a, (128u * uint(p_a_i_save)));
|
||||
const Inner l_a_i_a[4] = tint_symbol_6(a, (128u * uint(p_a_i_save)));
|
||||
const Inner l_a_i_a_i = tint_symbol_7(a, ((128u * uint(p_a_i_save)) + (32u * uint(p_a_i_a_i_save))));
|
||||
const float4x2 l_a_i_a_i_m = tint_symbol_8(a, ((128u * uint(p_a_i_save)) + (32u * uint(p_a_i_a_i_save))));
|
||||
const uint scalar_offset_4 = ((((128u * uint(p_a_i_save)) + (32u * uint(p_a_i_a_i_save))) + (8u * uint(p_a_i_a_i_m_i_save)))) / 4;
|
||||
uint4 ubo_load_4 = a[scalar_offset_4 / 4];
|
||||
const float2 l_a_i_a_i_m_i = asfloat(((scalar_offset_4 & 2) ? ubo_load_4.zw : ubo_load_4.xy));
|
||||
const int tint_symbol = p_a_i_save;
|
||||
const int tint_symbol_1 = p_a_i_a_i_save;
|
||||
const int tint_symbol_2 = p_a_i_a_i_m_i_save;
|
||||
const int tint_symbol_3 = i();
|
||||
const uint scalar_offset_5 = (((((128u * uint(tint_symbol)) + (32u * uint(tint_symbol_1))) + (8u * uint(tint_symbol_2))) + (4u * uint(tint_symbol_3)))) / 4;
|
||||
const float l_a_i_a_i_m_i_i = asfloat(a[scalar_offset_5 / 4][scalar_offset_5 % 4]);
|
||||
return;
|
||||
}
|
||||
@@ -0,0 +1,83 @@
|
||||
struct Inner {
|
||||
float4x2 m;
|
||||
};
|
||||
struct Outer {
|
||||
Inner a[4];
|
||||
};
|
||||
|
||||
cbuffer cbuffer_a : register(b0, space0) {
|
||||
uint4 a[32];
|
||||
};
|
||||
static int counter = 0;
|
||||
|
||||
int i() {
|
||||
counter = (counter + 1);
|
||||
return counter;
|
||||
}
|
||||
|
||||
float4x2 tint_symbol_8(uint4 buffer[32], 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)));
|
||||
}
|
||||
|
||||
Inner tint_symbol_7(uint4 buffer[32], uint offset) {
|
||||
const Inner tint_symbol_11 = {tint_symbol_8(buffer, (offset + 0u))};
|
||||
return tint_symbol_11;
|
||||
}
|
||||
|
||||
typedef Inner tint_symbol_6_ret[4];
|
||||
tint_symbol_6_ret tint_symbol_6(uint4 buffer[32], uint offset) {
|
||||
Inner arr[4] = (Inner[4])0;
|
||||
{
|
||||
[loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
|
||||
arr[i_1] = tint_symbol_7(buffer, (offset + (i_1 * 32u)));
|
||||
}
|
||||
}
|
||||
return arr;
|
||||
}
|
||||
|
||||
Outer tint_symbol_5(uint4 buffer[32], uint offset) {
|
||||
const Outer tint_symbol_12 = {tint_symbol_6(buffer, (offset + 0u))};
|
||||
return tint_symbol_12;
|
||||
}
|
||||
|
||||
typedef Outer tint_symbol_4_ret[4];
|
||||
tint_symbol_4_ret tint_symbol_4(uint4 buffer[32], uint offset) {
|
||||
Outer arr_1[4] = (Outer[4])0;
|
||||
{
|
||||
[loop] for(uint i_2 = 0u; (i_2 < 4u); i_2 = (i_2 + 1u)) {
|
||||
arr_1[i_2] = tint_symbol_5(buffer, (offset + (i_2 * 128u)));
|
||||
}
|
||||
}
|
||||
return arr_1;
|
||||
}
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void f() {
|
||||
const int I = 1;
|
||||
const int p_a_i_save = i();
|
||||
const int p_a_i_a_i_save = i();
|
||||
const int p_a_i_a_i_m_i_save = i();
|
||||
const Outer l_a[4] = tint_symbol_4(a, 0u);
|
||||
const Outer l_a_i = tint_symbol_5(a, (128u * uint(p_a_i_save)));
|
||||
const Inner l_a_i_a[4] = tint_symbol_6(a, (128u * uint(p_a_i_save)));
|
||||
const Inner l_a_i_a_i = tint_symbol_7(a, ((128u * uint(p_a_i_save)) + (32u * uint(p_a_i_a_i_save))));
|
||||
const float4x2 l_a_i_a_i_m = tint_symbol_8(a, ((128u * uint(p_a_i_save)) + (32u * uint(p_a_i_a_i_save))));
|
||||
const uint scalar_offset_4 = ((((128u * uint(p_a_i_save)) + (32u * uint(p_a_i_a_i_save))) + (8u * uint(p_a_i_a_i_m_i_save)))) / 4;
|
||||
uint4 ubo_load_4 = a[scalar_offset_4 / 4];
|
||||
const float2 l_a_i_a_i_m_i = asfloat(((scalar_offset_4 & 2) ? ubo_load_4.zw : ubo_load_4.xy));
|
||||
const int tint_symbol = p_a_i_save;
|
||||
const int tint_symbol_1 = p_a_i_a_i_save;
|
||||
const int tint_symbol_2 = p_a_i_a_i_m_i_save;
|
||||
const int tint_symbol_3 = i();
|
||||
const uint scalar_offset_5 = (((((128u * uint(tint_symbol)) + (32u * uint(tint_symbol_1))) + (8u * uint(tint_symbol_2))) + (4u * uint(tint_symbol_3)))) / 4;
|
||||
const float l_a_i_a_i_m_i_i = asfloat(a[scalar_offset_5 / 4][scalar_offset_5 % 4]);
|
||||
return;
|
||||
}
|
||||
@@ -0,0 +1,147 @@
|
||||
#version 310 es
|
||||
|
||||
struct Inner {
|
||||
mat4x2 m;
|
||||
};
|
||||
|
||||
struct Inner_std140 {
|
||||
vec2 m_0;
|
||||
vec2 m_1;
|
||||
vec2 m_2;
|
||||
vec2 m_3;
|
||||
};
|
||||
|
||||
struct Outer {
|
||||
Inner a[4];
|
||||
};
|
||||
|
||||
struct Outer_std140 {
|
||||
Inner_std140 a[4];
|
||||
};
|
||||
|
||||
struct a_block {
|
||||
Outer_std140 inner[4];
|
||||
};
|
||||
|
||||
layout(binding = 0) uniform a_block_1 {
|
||||
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(mat4x2(val.m_0, val.m_1, val.m_2, val.m_3));
|
||||
return tint_symbol_4;
|
||||
}
|
||||
|
||||
Inner[4] conv_arr_4_Inner(Inner_std140 val[4]) {
|
||||
Inner arr[4] = Inner[4](Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 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(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)))), Outer(Inner[4](Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)))), Outer(Inner[4](Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)))), Outer(Inner[4](Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 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;
|
||||
}
|
||||
|
||||
mat4x2 load_a_p0_a_p1_m(uint p0, uint p1) {
|
||||
uint s_save = p0;
|
||||
uint s_save_1 = p1;
|
||||
return mat4x2(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, a.inner[s_save].a[s_save_1].m_3);
|
||||
}
|
||||
|
||||
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;
|
||||
}
|
||||
case 3u: {
|
||||
return a.inner[p0].a[p1].m_3;
|
||||
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;
|
||||
}
|
||||
case 3u: {
|
||||
return a.inner[p0].a[p1].m_3[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]);
|
||||
mat4x2 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]);
|
||||
mat4x2 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;
|
||||
}
|
||||
@@ -0,0 +1,49 @@
|
||||
#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 Inner {
|
||||
/* 0x0000 */ float4x2 m;
|
||||
};
|
||||
|
||||
struct Outer {
|
||||
/* 0x0000 */ tint_array<Inner, 4> a;
|
||||
};
|
||||
|
||||
int i() {
|
||||
thread int tint_symbol_4 = 0;
|
||||
tint_symbol_4 = as_type<int>((as_type<uint>(tint_symbol_4) + as_type<uint>(1)));
|
||||
return tint_symbol_4;
|
||||
}
|
||||
|
||||
kernel void f(const constant tint_array<Outer, 4>* tint_symbol_5 [[buffer(0)]]) {
|
||||
int const I = 1;
|
||||
int const tint_symbol = i();
|
||||
int const p_a_i_save = tint_symbol;
|
||||
int const tint_symbol_1 = i();
|
||||
int const p_a_i_a_i_save = tint_symbol_1;
|
||||
int const tint_symbol_2 = i();
|
||||
int const p_a_i_a_i_m_i_save = tint_symbol_2;
|
||||
tint_array<Outer, 4> const l_a = *(tint_symbol_5);
|
||||
Outer const l_a_i = (*(tint_symbol_5))[p_a_i_save];
|
||||
tint_array<Inner, 4> const l_a_i_a = (*(tint_symbol_5))[p_a_i_save].a;
|
||||
Inner const l_a_i_a_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save];
|
||||
float4x2 const l_a_i_a_i_m = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m;
|
||||
float2 const l_a_i_a_i_m_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save];
|
||||
int const tint_symbol_3 = i();
|
||||
float const l_a_i_a_i_m_i_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save][tint_symbol_3];
|
||||
return;
|
||||
}
|
||||
|
||||
@@ -0,0 +1,352 @@
|
||||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 228
|
||||
; 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 %Outer_std140 "Outer_std140"
|
||||
OpMemberName %Outer_std140 0 "a"
|
||||
OpName %Inner_std140 "Inner_std140"
|
||||
OpMemberName %Inner_std140 0 "m_0"
|
||||
OpMemberName %Inner_std140 1 "m_1"
|
||||
OpMemberName %Inner_std140 2 "m_2"
|
||||
OpMemberName %Inner_std140 3 "m_3"
|
||||
OpName %a "a"
|
||||
OpName %counter "counter"
|
||||
OpName %i "i"
|
||||
OpName %Inner "Inner"
|
||||
OpMemberName %Inner 0 "m"
|
||||
OpName %conv_Inner "conv_Inner"
|
||||
OpName %val "val"
|
||||
OpName %conv_arr_4_Inner "conv_arr_4_Inner"
|
||||
OpName %val_0 "val"
|
||||
OpName %arr "arr"
|
||||
OpName %i_0 "i"
|
||||
OpName %var_for_index "var_for_index"
|
||||
OpName %Outer "Outer"
|
||||
OpMemberName %Outer 0 "a"
|
||||
OpName %conv_Outer "conv_Outer"
|
||||
OpName %val_1 "val"
|
||||
OpName %conv_arr_4_Outer "conv_arr_4_Outer"
|
||||
OpName %val_2 "val"
|
||||
OpName %arr_0 "arr"
|
||||
OpName %i_1 "i"
|
||||
OpName %var_for_index_1 "var_for_index_1"
|
||||
OpName %load_a_p0_a_p1_m "load_a_p0_a_p1_m"
|
||||
OpName %p0 "p0"
|
||||
OpName %p1 "p1"
|
||||
OpName %load_a_p0_a_p1_m_p2 "load_a_p0_a_p1_m_p2"
|
||||
OpName %p0_0 "p0"
|
||||
OpName %p1_0 "p1"
|
||||
OpName %p2 "p2"
|
||||
OpName %load_a_p0_a_p1_m_p2_p3 "load_a_p0_a_p1_m_p2_p3"
|
||||
OpName %p0_1 "p0"
|
||||
OpName %p1_1 "p1"
|
||||
OpName %p2_0 "p2"
|
||||
OpName %p3 "p3"
|
||||
OpName %f "f"
|
||||
OpDecorate %a_block Block
|
||||
OpMemberDecorate %a_block 0 Offset 0
|
||||
OpMemberDecorate %Outer_std140 0 Offset 0
|
||||
OpMemberDecorate %Inner_std140 0 Offset 0
|
||||
OpMemberDecorate %Inner_std140 1 Offset 8
|
||||
OpMemberDecorate %Inner_std140 2 Offset 16
|
||||
OpMemberDecorate %Inner_std140 3 Offset 24
|
||||
OpDecorate %_arr_Inner_std140_uint_4 ArrayStride 32
|
||||
OpDecorate %_arr_Outer_std140_uint_4 ArrayStride 128
|
||||
OpDecorate %a NonWritable
|
||||
OpDecorate %a DescriptorSet 0
|
||||
OpDecorate %a Binding 0
|
||||
OpMemberDecorate %Inner 0 Offset 0
|
||||
OpMemberDecorate %Inner 0 ColMajor
|
||||
OpMemberDecorate %Inner 0 MatrixStride 8
|
||||
OpDecorate %_arr_Inner_uint_4 ArrayStride 32
|
||||
OpMemberDecorate %Outer 0 Offset 0
|
||||
OpDecorate %_arr_Outer_uint_4 ArrayStride 128
|
||||
%float = OpTypeFloat 32
|
||||
%v2float = OpTypeVector %float 2
|
||||
%Inner_std140 = OpTypeStruct %v2float %v2float %v2float %v2float
|
||||
%uint = OpTypeInt 32 0
|
||||
%uint_4 = OpConstant %uint 4
|
||||
%_arr_Inner_std140_uint_4 = OpTypeArray %Inner_std140 %uint_4
|
||||
%Outer_std140 = OpTypeStruct %_arr_Inner_std140_uint_4
|
||||
%_arr_Outer_std140_uint_4 = OpTypeArray %Outer_std140 %uint_4
|
||||
%a_block = OpTypeStruct %_arr_Outer_std140_uint_4
|
||||
%_ptr_Uniform_a_block = OpTypePointer Uniform %a_block
|
||||
%a = OpVariable %_ptr_Uniform_a_block Uniform
|
||||
%int = OpTypeInt 32 1
|
||||
%13 = OpConstantNull %int
|
||||
%_ptr_Private_int = OpTypePointer Private %int
|
||||
%counter = OpVariable %_ptr_Private_int Private %13
|
||||
%16 = OpTypeFunction %int
|
||||
%int_1 = OpConstant %int 1
|
||||
%mat4v2float = OpTypeMatrix %v2float 4
|
||||
%Inner = OpTypeStruct %mat4v2float
|
||||
%23 = OpTypeFunction %Inner %Inner_std140
|
||||
%_arr_Inner_uint_4 = OpTypeArray %Inner %uint_4
|
||||
%35 = OpTypeFunction %_arr_Inner_uint_4 %_arr_Inner_std140_uint_4
|
||||
%_ptr_Function__arr_Inner_uint_4 = OpTypePointer Function %_arr_Inner_uint_4
|
||||
%42 = OpConstantNull %_arr_Inner_uint_4
|
||||
%_ptr_Function_uint = OpTypePointer Function %uint
|
||||
%45 = OpConstantNull %uint
|
||||
%bool = OpTypeBool
|
||||
%_ptr_Function__arr_Inner_std140_uint_4 = OpTypePointer Function %_arr_Inner_std140_uint_4
|
||||
%58 = OpConstantNull %_arr_Inner_std140_uint_4
|
||||
%_ptr_Function_Inner = OpTypePointer Function %Inner
|
||||
%_ptr_Function_Inner_std140 = OpTypePointer Function %Inner_std140
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%Outer = OpTypeStruct %_arr_Inner_uint_4
|
||||
%71 = OpTypeFunction %Outer %Outer_std140
|
||||
%_arr_Outer_uint_4 = OpTypeArray %Outer %uint_4
|
||||
%79 = OpTypeFunction %_arr_Outer_uint_4 %_arr_Outer_std140_uint_4
|
||||
%_ptr_Function__arr_Outer_uint_4 = OpTypePointer Function %_arr_Outer_uint_4
|
||||
%86 = OpConstantNull %_arr_Outer_uint_4
|
||||
%_ptr_Function__arr_Outer_std140_uint_4 = OpTypePointer Function %_arr_Outer_std140_uint_4
|
||||
%99 = OpConstantNull %_arr_Outer_std140_uint_4
|
||||
%_ptr_Function_Outer = OpTypePointer Function %Outer
|
||||
%_ptr_Function_Outer_std140 = OpTypePointer Function %Outer_std140
|
||||
%111 = OpTypeFunction %mat4v2float %uint %uint
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Uniform_v2float = OpTypePointer Uniform %v2float
|
||||
%uint_2 = OpConstant %uint 2
|
||||
%uint_3 = OpConstant %uint 3
|
||||
%129 = OpTypeFunction %v2float %uint %uint %uint
|
||||
%149 = OpConstantNull %v2float
|
||||
%150 = OpTypeFunction %float %uint %uint %uint %uint
|
||||
%_ptr_Uniform_float = OpTypePointer Uniform %float
|
||||
%172 = OpConstantNull %float
|
||||
%void = OpTypeVoid
|
||||
%173 = OpTypeFunction %void
|
||||
%_ptr_Uniform__arr_Outer_std140_uint_4 = OpTypePointer Uniform %_arr_Outer_std140_uint_4
|
||||
%_ptr_Uniform_Outer_std140 = OpTypePointer Uniform %Outer_std140
|
||||
%_ptr_Uniform__arr_Inner_std140_uint_4 = OpTypePointer Uniform %_arr_Inner_std140_uint_4
|
||||
%_ptr_Uniform_Inner_std140 = OpTypePointer Uniform %Inner_std140
|
||||
%i = OpFunction %int None %16
|
||||
%18 = OpLabel
|
||||
%19 = OpLoad %int %counter
|
||||
%21 = OpIAdd %int %19 %int_1
|
||||
OpStore %counter %21
|
||||
%22 = OpLoad %int %counter
|
||||
OpReturnValue %22
|
||||
OpFunctionEnd
|
||||
%conv_Inner = OpFunction %Inner None %23
|
||||
%val = OpFunctionParameter %Inner_std140
|
||||
%28 = OpLabel
|
||||
%29 = OpCompositeExtract %v2float %val 0
|
||||
%30 = OpCompositeExtract %v2float %val 1
|
||||
%31 = OpCompositeExtract %v2float %val 2
|
||||
%32 = OpCompositeExtract %v2float %val 3
|
||||
%33 = OpCompositeConstruct %mat4v2float %29 %30 %31 %32
|
||||
%34 = OpCompositeConstruct %Inner %33
|
||||
OpReturnValue %34
|
||||
OpFunctionEnd
|
||||
%conv_arr_4_Inner = OpFunction %_arr_Inner_uint_4 None %35
|
||||
%val_0 = OpFunctionParameter %_arr_Inner_std140_uint_4
|
||||
%39 = OpLabel
|
||||
%arr = OpVariable %_ptr_Function__arr_Inner_uint_4 Function %42
|
||||
%i_0 = OpVariable %_ptr_Function_uint Function %45
|
||||
%var_for_index = OpVariable %_ptr_Function__arr_Inner_std140_uint_4 Function %58
|
||||
OpBranch %46
|
||||
%46 = OpLabel
|
||||
OpLoopMerge %47 %48 None
|
||||
OpBranch %49
|
||||
%49 = OpLabel
|
||||
%51 = OpLoad %uint %i_0
|
||||
%52 = OpULessThan %bool %51 %uint_4
|
||||
%50 = OpLogicalNot %bool %52
|
||||
OpSelectionMerge %54 None
|
||||
OpBranchConditional %50 %55 %54
|
||||
%55 = OpLabel
|
||||
OpBranch %47
|
||||
%54 = OpLabel
|
||||
OpStore %var_for_index %val_0
|
||||
%59 = OpLoad %uint %i_0
|
||||
%61 = OpAccessChain %_ptr_Function_Inner %arr %59
|
||||
%63 = OpLoad %uint %i_0
|
||||
%65 = OpAccessChain %_ptr_Function_Inner_std140 %var_for_index %63
|
||||
%66 = OpLoad %Inner_std140 %65
|
||||
%62 = OpFunctionCall %Inner %conv_Inner %66
|
||||
OpStore %61 %62
|
||||
OpBranch %48
|
||||
%48 = OpLabel
|
||||
%67 = OpLoad %uint %i_0
|
||||
%69 = OpIAdd %uint %67 %uint_1
|
||||
OpStore %i_0 %69
|
||||
OpBranch %46
|
||||
%47 = OpLabel
|
||||
%70 = OpLoad %_arr_Inner_uint_4 %arr
|
||||
OpReturnValue %70
|
||||
OpFunctionEnd
|
||||
%conv_Outer = OpFunction %Outer None %71
|
||||
%val_1 = OpFunctionParameter %Outer_std140
|
||||
%75 = OpLabel
|
||||
%77 = OpCompositeExtract %_arr_Inner_std140_uint_4 %val_1 0
|
||||
%76 = OpFunctionCall %_arr_Inner_uint_4 %conv_arr_4_Inner %77
|
||||
%78 = OpCompositeConstruct %Outer %76
|
||||
OpReturnValue %78
|
||||
OpFunctionEnd
|
||||
%conv_arr_4_Outer = OpFunction %_arr_Outer_uint_4 None %79
|
||||
%val_2 = OpFunctionParameter %_arr_Outer_std140_uint_4
|
||||
%83 = OpLabel
|
||||
%arr_0 = OpVariable %_ptr_Function__arr_Outer_uint_4 Function %86
|
||||
%i_1 = OpVariable %_ptr_Function_uint Function %45
|
||||
%var_for_index_1 = OpVariable %_ptr_Function__arr_Outer_std140_uint_4 Function %99
|
||||
OpBranch %88
|
||||
%88 = OpLabel
|
||||
OpLoopMerge %89 %90 None
|
||||
OpBranch %91
|
||||
%91 = OpLabel
|
||||
%93 = OpLoad %uint %i_1
|
||||
%94 = OpULessThan %bool %93 %uint_4
|
||||
%92 = OpLogicalNot %bool %94
|
||||
OpSelectionMerge %95 None
|
||||
OpBranchConditional %92 %96 %95
|
||||
%96 = OpLabel
|
||||
OpBranch %89
|
||||
%95 = OpLabel
|
||||
OpStore %var_for_index_1 %val_2
|
||||
%100 = OpLoad %uint %i_1
|
||||
%102 = OpAccessChain %_ptr_Function_Outer %arr_0 %100
|
||||
%104 = OpLoad %uint %i_1
|
||||
%106 = OpAccessChain %_ptr_Function_Outer_std140 %var_for_index_1 %104
|
||||
%107 = OpLoad %Outer_std140 %106
|
||||
%103 = OpFunctionCall %Outer %conv_Outer %107
|
||||
OpStore %102 %103
|
||||
OpBranch %90
|
||||
%90 = OpLabel
|
||||
%108 = OpLoad %uint %i_1
|
||||
%109 = OpIAdd %uint %108 %uint_1
|
||||
OpStore %i_1 %109
|
||||
OpBranch %88
|
||||
%89 = OpLabel
|
||||
%110 = OpLoad %_arr_Outer_uint_4 %arr_0
|
||||
OpReturnValue %110
|
||||
OpFunctionEnd
|
||||
%load_a_p0_a_p1_m = OpFunction %mat4v2float None %111
|
||||
%p0 = OpFunctionParameter %uint
|
||||
%p1 = OpFunctionParameter %uint
|
||||
%115 = OpLabel
|
||||
%118 = OpAccessChain %_ptr_Uniform_v2float %a %uint_0 %p0 %uint_0 %p1 %uint_0
|
||||
%119 = OpLoad %v2float %118
|
||||
%120 = OpAccessChain %_ptr_Uniform_v2float %a %uint_0 %p0 %uint_0 %p1 %uint_1
|
||||
%121 = OpLoad %v2float %120
|
||||
%123 = OpAccessChain %_ptr_Uniform_v2float %a %uint_0 %p0 %uint_0 %p1 %uint_2
|
||||
%124 = OpLoad %v2float %123
|
||||
%126 = OpAccessChain %_ptr_Uniform_v2float %a %uint_0 %p0 %uint_0 %p1 %uint_3
|
||||
%127 = OpLoad %v2float %126
|
||||
%128 = OpCompositeConstruct %mat4v2float %119 %121 %124 %127
|
||||
OpReturnValue %128
|
||||
OpFunctionEnd
|
||||
%load_a_p0_a_p1_m_p2 = OpFunction %v2float None %129
|
||||
%p0_0 = OpFunctionParameter %uint
|
||||
%p1_0 = OpFunctionParameter %uint
|
||||
%p2 = OpFunctionParameter %uint
|
||||
%134 = OpLabel
|
||||
OpSelectionMerge %135 None
|
||||
OpSwitch %p2 %136 0 %137 1 %138 2 %139 3 %140
|
||||
%137 = OpLabel
|
||||
%141 = OpAccessChain %_ptr_Uniform_v2float %a %uint_0 %p0_0 %uint_0 %p1_0 %uint_0
|
||||
%142 = OpLoad %v2float %141
|
||||
OpReturnValue %142
|
||||
%138 = OpLabel
|
||||
%143 = OpAccessChain %_ptr_Uniform_v2float %a %uint_0 %p0_0 %uint_0 %p1_0 %uint_1
|
||||
%144 = OpLoad %v2float %143
|
||||
OpReturnValue %144
|
||||
%139 = OpLabel
|
||||
%145 = OpAccessChain %_ptr_Uniform_v2float %a %uint_0 %p0_0 %uint_0 %p1_0 %uint_2
|
||||
%146 = OpLoad %v2float %145
|
||||
OpReturnValue %146
|
||||
%140 = OpLabel
|
||||
%147 = OpAccessChain %_ptr_Uniform_v2float %a %uint_0 %p0_0 %uint_0 %p1_0 %uint_3
|
||||
%148 = OpLoad %v2float %147
|
||||
OpReturnValue %148
|
||||
%136 = OpLabel
|
||||
OpReturnValue %149
|
||||
%135 = OpLabel
|
||||
OpReturnValue %149
|
||||
OpFunctionEnd
|
||||
%load_a_p0_a_p1_m_p2_p3 = OpFunction %float None %150
|
||||
%p0_1 = OpFunctionParameter %uint
|
||||
%p1_1 = OpFunctionParameter %uint
|
||||
%p2_0 = OpFunctionParameter %uint
|
||||
%p3 = OpFunctionParameter %uint
|
||||
%156 = OpLabel
|
||||
OpSelectionMerge %157 None
|
||||
OpSwitch %p2_0 %158 0 %159 1 %160 2 %161 3 %162
|
||||
%159 = OpLabel
|
||||
%164 = OpAccessChain %_ptr_Uniform_float %a %uint_0 %p0_1 %uint_0 %p1_1 %uint_0 %p3
|
||||
%165 = OpLoad %float %164
|
||||
OpReturnValue %165
|
||||
%160 = OpLabel
|
||||
%166 = OpAccessChain %_ptr_Uniform_float %a %uint_0 %p0_1 %uint_0 %p1_1 %uint_1 %p3
|
||||
%167 = OpLoad %float %166
|
||||
OpReturnValue %167
|
||||
%161 = OpLabel
|
||||
%168 = OpAccessChain %_ptr_Uniform_float %a %uint_0 %p0_1 %uint_0 %p1_1 %uint_2 %p3
|
||||
%169 = OpLoad %float %168
|
||||
OpReturnValue %169
|
||||
%162 = OpLabel
|
||||
%170 = OpAccessChain %_ptr_Uniform_float %a %uint_0 %p0_1 %uint_0 %p1_1 %uint_3 %p3
|
||||
%171 = OpLoad %float %170
|
||||
OpReturnValue %171
|
||||
%158 = OpLabel
|
||||
OpReturnValue %172
|
||||
%157 = OpLabel
|
||||
OpReturnValue %172
|
||||
OpFunctionEnd
|
||||
%f = OpFunction %void None %173
|
||||
%176 = OpLabel
|
||||
%179 = OpAccessChain %_ptr_Uniform__arr_Outer_std140_uint_4 %a %uint_0
|
||||
%180 = OpLoad %_arr_Outer_std140_uint_4 %179
|
||||
%177 = OpFunctionCall %_arr_Outer_uint_4 %conv_arr_4_Outer %180
|
||||
%181 = OpFunctionCall %int %i
|
||||
%184 = OpAccessChain %_ptr_Uniform_Outer_std140 %a %uint_0 %181
|
||||
%185 = OpLoad %Outer_std140 %184
|
||||
%182 = OpFunctionCall %Outer %conv_Outer %185
|
||||
%188 = OpAccessChain %_ptr_Uniform__arr_Inner_std140_uint_4 %a %uint_0 %181 %uint_0
|
||||
%189 = OpLoad %_arr_Inner_std140_uint_4 %188
|
||||
%186 = OpFunctionCall %_arr_Inner_uint_4 %conv_arr_4_Inner %189
|
||||
%190 = OpFunctionCall %int %i
|
||||
%193 = OpAccessChain %_ptr_Uniform_Inner_std140 %a %uint_0 %181 %uint_0 %190
|
||||
%194 = OpLoad %Inner_std140 %193
|
||||
%191 = OpFunctionCall %Inner %conv_Inner %194
|
||||
%196 = OpBitcast %uint %181
|
||||
%197 = OpBitcast %uint %190
|
||||
%195 = OpFunctionCall %mat4v2float %load_a_p0_a_p1_m %196 %197
|
||||
%198 = OpFunctionCall %int %i
|
||||
%200 = OpBitcast %uint %181
|
||||
%201 = OpBitcast %uint %190
|
||||
%202 = OpBitcast %uint %198
|
||||
%199 = OpFunctionCall %v2float %load_a_p0_a_p1_m_p2 %200 %201 %202
|
||||
%204 = OpAccessChain %_ptr_Uniform__arr_Outer_std140_uint_4 %a %uint_0
|
||||
%205 = OpLoad %_arr_Outer_std140_uint_4 %204
|
||||
%203 = OpFunctionCall %_arr_Outer_uint_4 %conv_arr_4_Outer %205
|
||||
%207 = OpAccessChain %_ptr_Uniform_Outer_std140 %a %uint_0 %181
|
||||
%208 = OpLoad %Outer_std140 %207
|
||||
%206 = OpFunctionCall %Outer %conv_Outer %208
|
||||
%210 = OpAccessChain %_ptr_Uniform__arr_Inner_std140_uint_4 %a %uint_0 %181 %uint_0
|
||||
%211 = OpLoad %_arr_Inner_std140_uint_4 %210
|
||||
%209 = OpFunctionCall %_arr_Inner_uint_4 %conv_arr_4_Inner %211
|
||||
%213 = OpAccessChain %_ptr_Uniform_Inner_std140 %a %uint_0 %181 %uint_0 %190
|
||||
%214 = OpLoad %Inner_std140 %213
|
||||
%212 = OpFunctionCall %Inner %conv_Inner %214
|
||||
%216 = OpBitcast %uint %181
|
||||
%217 = OpBitcast %uint %190
|
||||
%215 = OpFunctionCall %mat4v2float %load_a_p0_a_p1_m %216 %217
|
||||
%219 = OpBitcast %uint %181
|
||||
%220 = OpBitcast %uint %190
|
||||
%221 = OpBitcast %uint %198
|
||||
%218 = OpFunctionCall %v2float %load_a_p0_a_p1_m_p2 %219 %220 %221
|
||||
%222 = OpFunctionCall %int %i
|
||||
%224 = OpBitcast %uint %181
|
||||
%225 = OpBitcast %uint %190
|
||||
%226 = OpBitcast %uint %198
|
||||
%227 = OpBitcast %uint %222
|
||||
%223 = OpFunctionCall %float %load_a_p0_a_p1_m_p2_p3 %224 %225 %226 %227
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@@ -0,0 +1,34 @@
|
||||
struct Inner {
|
||||
m : mat4x2<f32>,
|
||||
}
|
||||
|
||||
struct Outer {
|
||||
a : array<Inner, 4>,
|
||||
}
|
||||
|
||||
@group(0) @binding(0) var<uniform> a : array<Outer, 4>;
|
||||
|
||||
var<private> counter = 0;
|
||||
|
||||
fn i() -> i32 {
|
||||
counter++;
|
||||
return counter;
|
||||
}
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn f() {
|
||||
let I = 1;
|
||||
let p_a = &(a);
|
||||
let p_a_i = &((*(p_a))[i()]);
|
||||
let p_a_i_a = &((*(p_a_i)).a);
|
||||
let p_a_i_a_i = &((*(p_a_i_a))[i()]);
|
||||
let p_a_i_a_i_m = &((*(p_a_i_a_i)).m);
|
||||
let p_a_i_a_i_m_i = &((*(p_a_i_a_i_m))[i()]);
|
||||
let l_a : array<Outer, 4> = *(p_a);
|
||||
let l_a_i : Outer = *(p_a_i);
|
||||
let l_a_i_a : array<Inner, 4> = *(p_a_i_a);
|
||||
let l_a_i_a_i : Inner = *(p_a_i_a_i);
|
||||
let l_a_i_a_i_m : mat4x2<f32> = *(p_a_i_a_i_m);
|
||||
let l_a_i_a_i_m_i : vec2<f32> = *(p_a_i_a_i_m_i);
|
||||
let l_a_i_a_i_m_i_i : f32 = (*(p_a_i_a_i_m_i))[i()];
|
||||
}
|
||||
@@ -0,0 +1,30 @@
|
||||
struct Inner {
|
||||
m : mat4x2<f32>,
|
||||
}
|
||||
|
||||
struct Outer {
|
||||
a : array<Inner, 4>,
|
||||
}
|
||||
|
||||
@group(0) @binding(0) var<uniform> a : array<Outer, 4>;
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn f() {
|
||||
let I = 1;
|
||||
|
||||
let p_a = &a;
|
||||
let p_a_3 = &((*p_a)[3]);
|
||||
let p_a_3_a = &((*p_a_3).a);
|
||||
let p_a_3_a_2 = &((*p_a_3_a)[2]);
|
||||
let p_a_3_a_2_m = &((*p_a_3_a_2).m);
|
||||
let p_a_3_a_2_m_1 = &((*p_a_3_a_2_m)[1]);
|
||||
|
||||
|
||||
let l_a : array<Outer, 4> = *p_a;
|
||||
let l_a_3 : Outer = *p_a_3;
|
||||
let l_a_3_a : array<Inner, 4> = *p_a_3_a;
|
||||
let l_a_3_a_2 : Inner = *p_a_3_a_2;
|
||||
let l_a_3_a_2_m : mat4x2<f32> = *p_a_3_a_2_m;
|
||||
let l_a_3_a_2_m_1 : vec2<f32> = *p_a_3_a_2_m_1;
|
||||
let l_a_3_a_2_m_1_0 : f32 = (*p_a_3_a_2_m_1)[0];
|
||||
}
|
||||
@@ -0,0 +1,67 @@
|
||||
struct Inner {
|
||||
float4x2 m;
|
||||
};
|
||||
struct Outer {
|
||||
Inner a[4];
|
||||
};
|
||||
|
||||
cbuffer cbuffer_a : register(b0, space0) {
|
||||
uint4 a[32];
|
||||
};
|
||||
|
||||
float4x2 tint_symbol_4(uint4 buffer[32], 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)));
|
||||
}
|
||||
|
||||
Inner tint_symbol_3(uint4 buffer[32], uint offset) {
|
||||
const Inner tint_symbol_7 = {tint_symbol_4(buffer, (offset + 0u))};
|
||||
return tint_symbol_7;
|
||||
}
|
||||
|
||||
typedef Inner tint_symbol_2_ret[4];
|
||||
tint_symbol_2_ret tint_symbol_2(uint4 buffer[32], uint offset) {
|
||||
Inner arr[4] = (Inner[4])0;
|
||||
{
|
||||
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
|
||||
arr[i] = tint_symbol_3(buffer, (offset + (i * 32u)));
|
||||
}
|
||||
}
|
||||
return arr;
|
||||
}
|
||||
|
||||
Outer tint_symbol_1(uint4 buffer[32], uint offset) {
|
||||
const Outer tint_symbol_8 = {tint_symbol_2(buffer, (offset + 0u))};
|
||||
return tint_symbol_8;
|
||||
}
|
||||
|
||||
typedef Outer tint_symbol_ret[4];
|
||||
tint_symbol_ret tint_symbol(uint4 buffer[32], uint offset) {
|
||||
Outer arr_1[4] = (Outer[4])0;
|
||||
{
|
||||
[loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
|
||||
arr_1[i_1] = tint_symbol_1(buffer, (offset + (i_1 * 128u)));
|
||||
}
|
||||
}
|
||||
return arr_1;
|
||||
}
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void f() {
|
||||
const int I = 1;
|
||||
const Outer l_a[4] = tint_symbol(a, 0u);
|
||||
const Outer l_a_3 = tint_symbol_1(a, 384u);
|
||||
const Inner l_a_3_a[4] = tint_symbol_2(a, 384u);
|
||||
const Inner l_a_3_a_2 = tint_symbol_3(a, 448u);
|
||||
const float4x2 l_a_3_a_2_m = tint_symbol_4(a, 448u);
|
||||
const float2 l_a_3_a_2_m_1 = asfloat(a[28].zw);
|
||||
const float l_a_3_a_2_m_1_0 = asfloat(a[28].z);
|
||||
return;
|
||||
}
|
||||
@@ -0,0 +1,67 @@
|
||||
struct Inner {
|
||||
float4x2 m;
|
||||
};
|
||||
struct Outer {
|
||||
Inner a[4];
|
||||
};
|
||||
|
||||
cbuffer cbuffer_a : register(b0, space0) {
|
||||
uint4 a[32];
|
||||
};
|
||||
|
||||
float4x2 tint_symbol_4(uint4 buffer[32], 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)));
|
||||
}
|
||||
|
||||
Inner tint_symbol_3(uint4 buffer[32], uint offset) {
|
||||
const Inner tint_symbol_7 = {tint_symbol_4(buffer, (offset + 0u))};
|
||||
return tint_symbol_7;
|
||||
}
|
||||
|
||||
typedef Inner tint_symbol_2_ret[4];
|
||||
tint_symbol_2_ret tint_symbol_2(uint4 buffer[32], uint offset) {
|
||||
Inner arr[4] = (Inner[4])0;
|
||||
{
|
||||
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
|
||||
arr[i] = tint_symbol_3(buffer, (offset + (i * 32u)));
|
||||
}
|
||||
}
|
||||
return arr;
|
||||
}
|
||||
|
||||
Outer tint_symbol_1(uint4 buffer[32], uint offset) {
|
||||
const Outer tint_symbol_8 = {tint_symbol_2(buffer, (offset + 0u))};
|
||||
return tint_symbol_8;
|
||||
}
|
||||
|
||||
typedef Outer tint_symbol_ret[4];
|
||||
tint_symbol_ret tint_symbol(uint4 buffer[32], uint offset) {
|
||||
Outer arr_1[4] = (Outer[4])0;
|
||||
{
|
||||
[loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
|
||||
arr_1[i_1] = tint_symbol_1(buffer, (offset + (i_1 * 128u)));
|
||||
}
|
||||
}
|
||||
return arr_1;
|
||||
}
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void f() {
|
||||
const int I = 1;
|
||||
const Outer l_a[4] = tint_symbol(a, 0u);
|
||||
const Outer l_a_3 = tint_symbol_1(a, 384u);
|
||||
const Inner l_a_3_a[4] = tint_symbol_2(a, 384u);
|
||||
const Inner l_a_3_a_2 = tint_symbol_3(a, 448u);
|
||||
const float4x2 l_a_3_a_2_m = tint_symbol_4(a, 448u);
|
||||
const float2 l_a_3_a_2_m_1 = asfloat(a[28].zw);
|
||||
const float l_a_3_a_2_m_1_0 = asfloat(a[28].z);
|
||||
return;
|
||||
}
|
||||
@@ -0,0 +1,85 @@
|
||||
#version 310 es
|
||||
|
||||
struct Inner {
|
||||
mat4x2 m;
|
||||
};
|
||||
|
||||
struct Inner_std140 {
|
||||
vec2 m_0;
|
||||
vec2 m_1;
|
||||
vec2 m_2;
|
||||
vec2 m_3;
|
||||
};
|
||||
|
||||
struct Outer {
|
||||
Inner a[4];
|
||||
};
|
||||
|
||||
struct Outer_std140 {
|
||||
Inner_std140 a[4];
|
||||
};
|
||||
|
||||
struct a_block {
|
||||
Outer_std140 inner[4];
|
||||
};
|
||||
|
||||
layout(binding = 0) uniform a_block_1 {
|
||||
Outer_std140 inner[4];
|
||||
} a;
|
||||
|
||||
Inner conv_Inner(Inner_std140 val) {
|
||||
Inner tint_symbol = Inner(mat4x2(val.m_0, val.m_1, val.m_2, val.m_3));
|
||||
return tint_symbol;
|
||||
}
|
||||
|
||||
Inner[4] conv_arr_4_Inner(Inner_std140 val[4]) {
|
||||
Inner arr[4] = Inner[4](Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 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_1 = Outer(conv_arr_4_Inner(val.a));
|
||||
return tint_symbol_1;
|
||||
}
|
||||
|
||||
Outer[4] conv_arr_4_Outer(Outer_std140 val[4]) {
|
||||
Outer arr[4] = Outer[4](Outer(Inner[4](Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)))), Outer(Inner[4](Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)))), Outer(Inner[4](Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)))), Outer(Inner[4](Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)), Inner(mat4x2(0.0f, 0.0f, 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;
|
||||
}
|
||||
|
||||
mat4x2 load_a_3_a_2_m() {
|
||||
return mat4x2(a.inner[3u].a[2u].m_0, a.inner[3u].a[2u].m_1, a.inner[3u].a[2u].m_2, a.inner[3u].a[2u].m_3);
|
||||
}
|
||||
|
||||
void f() {
|
||||
int I = 1;
|
||||
Outer p_a[4] = conv_arr_4_Outer(a.inner);
|
||||
Outer p_a_3 = conv_Outer(a.inner[3u]);
|
||||
Inner p_a_3_a[4] = conv_arr_4_Inner(a.inner[3u].a);
|
||||
Inner p_a_3_a_2 = conv_Inner(a.inner[3u].a[2u]);
|
||||
mat4x2 p_a_3_a_2_m = load_a_3_a_2_m();
|
||||
vec2 p_a_3_a_2_m_1 = a.inner[3u].a[2u].m_1;
|
||||
Outer l_a[4] = conv_arr_4_Outer(a.inner);
|
||||
Outer l_a_3 = conv_Outer(a.inner[3u]);
|
||||
Inner l_a_3_a[4] = conv_arr_4_Inner(a.inner[3u].a);
|
||||
Inner l_a_3_a_2 = conv_Inner(a.inner[3u].a[2u]);
|
||||
mat4x2 l_a_3_a_2_m = load_a_3_a_2_m();
|
||||
vec2 l_a_3_a_2_m_1 = a.inner[3u].a[2u].m_1;
|
||||
float l_a_3_a_2_m_1_0 = a.inner[3u].a[2u].m_1[0u];
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
void main() {
|
||||
f();
|
||||
return;
|
||||
}
|
||||
@@ -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];
|
||||
};
|
||||
|
||||
struct Inner {
|
||||
/* 0x0000 */ float4x2 m;
|
||||
};
|
||||
|
||||
struct Outer {
|
||||
/* 0x0000 */ tint_array<Inner, 4> a;
|
||||
};
|
||||
|
||||
kernel void f(const constant tint_array<Outer, 4>* tint_symbol [[buffer(0)]]) {
|
||||
int const I = 1;
|
||||
tint_array<Outer, 4> const l_a = *(tint_symbol);
|
||||
Outer const l_a_3 = (*(tint_symbol))[3];
|
||||
tint_array<Inner, 4> const l_a_3_a = (*(tint_symbol))[3].a;
|
||||
Inner const l_a_3_a_2 = (*(tint_symbol))[3].a[2];
|
||||
float4x2 const l_a_3_a_2_m = (*(tint_symbol))[3].a[2].m;
|
||||
float2 const l_a_3_a_2_m_1 = (*(tint_symbol))[3].a[2].m[1];
|
||||
float const l_a_3_a_2_m_1_0 = (*(tint_symbol))[3].a[2].m[1][0];
|
||||
return;
|
||||
}
|
||||
|
||||
@@ -0,0 +1,249 @@
|
||||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 159
|
||||
; 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 %Outer_std140 "Outer_std140"
|
||||
OpMemberName %Outer_std140 0 "a"
|
||||
OpName %Inner_std140 "Inner_std140"
|
||||
OpMemberName %Inner_std140 0 "m_0"
|
||||
OpMemberName %Inner_std140 1 "m_1"
|
||||
OpMemberName %Inner_std140 2 "m_2"
|
||||
OpMemberName %Inner_std140 3 "m_3"
|
||||
OpName %a "a"
|
||||
OpName %Inner "Inner"
|
||||
OpMemberName %Inner 0 "m"
|
||||
OpName %conv_Inner "conv_Inner"
|
||||
OpName %val "val"
|
||||
OpName %conv_arr_4_Inner "conv_arr_4_Inner"
|
||||
OpName %val_0 "val"
|
||||
OpName %arr "arr"
|
||||
OpName %i "i"
|
||||
OpName %var_for_index "var_for_index"
|
||||
OpName %Outer "Outer"
|
||||
OpMemberName %Outer 0 "a"
|
||||
OpName %conv_Outer "conv_Outer"
|
||||
OpName %val_1 "val"
|
||||
OpName %conv_arr_4_Outer "conv_arr_4_Outer"
|
||||
OpName %val_2 "val"
|
||||
OpName %arr_0 "arr"
|
||||
OpName %i_0 "i"
|
||||
OpName %var_for_index_1 "var_for_index_1"
|
||||
OpName %load_a_3_a_2_m "load_a_3_a_2_m"
|
||||
OpName %f "f"
|
||||
OpDecorate %a_block Block
|
||||
OpMemberDecorate %a_block 0 Offset 0
|
||||
OpMemberDecorate %Outer_std140 0 Offset 0
|
||||
OpMemberDecorate %Inner_std140 0 Offset 0
|
||||
OpMemberDecorate %Inner_std140 1 Offset 8
|
||||
OpMemberDecorate %Inner_std140 2 Offset 16
|
||||
OpMemberDecorate %Inner_std140 3 Offset 24
|
||||
OpDecorate %_arr_Inner_std140_uint_4 ArrayStride 32
|
||||
OpDecorate %_arr_Outer_std140_uint_4 ArrayStride 128
|
||||
OpDecorate %a NonWritable
|
||||
OpDecorate %a DescriptorSet 0
|
||||
OpDecorate %a Binding 0
|
||||
OpMemberDecorate %Inner 0 Offset 0
|
||||
OpMemberDecorate %Inner 0 ColMajor
|
||||
OpMemberDecorate %Inner 0 MatrixStride 8
|
||||
OpDecorate %_arr_Inner_uint_4 ArrayStride 32
|
||||
OpMemberDecorate %Outer 0 Offset 0
|
||||
OpDecorate %_arr_Outer_uint_4 ArrayStride 128
|
||||
%float = OpTypeFloat 32
|
||||
%v2float = OpTypeVector %float 2
|
||||
%Inner_std140 = OpTypeStruct %v2float %v2float %v2float %v2float
|
||||
%uint = OpTypeInt 32 0
|
||||
%uint_4 = OpConstant %uint 4
|
||||
%_arr_Inner_std140_uint_4 = OpTypeArray %Inner_std140 %uint_4
|
||||
%Outer_std140 = OpTypeStruct %_arr_Inner_std140_uint_4
|
||||
%_arr_Outer_std140_uint_4 = OpTypeArray %Outer_std140 %uint_4
|
||||
%a_block = OpTypeStruct %_arr_Outer_std140_uint_4
|
||||
%_ptr_Uniform_a_block = OpTypePointer Uniform %a_block
|
||||
%a = OpVariable %_ptr_Uniform_a_block Uniform
|
||||
%mat4v2float = OpTypeMatrix %v2float 4
|
||||
%Inner = OpTypeStruct %mat4v2float
|
||||
%12 = OpTypeFunction %Inner %Inner_std140
|
||||
%_arr_Inner_uint_4 = OpTypeArray %Inner %uint_4
|
||||
%24 = OpTypeFunction %_arr_Inner_uint_4 %_arr_Inner_std140_uint_4
|
||||
%_ptr_Function__arr_Inner_uint_4 = OpTypePointer Function %_arr_Inner_uint_4
|
||||
%31 = OpConstantNull %_arr_Inner_uint_4
|
||||
%_ptr_Function_uint = OpTypePointer Function %uint
|
||||
%34 = OpConstantNull %uint
|
||||
%bool = OpTypeBool
|
||||
%_ptr_Function__arr_Inner_std140_uint_4 = OpTypePointer Function %_arr_Inner_std140_uint_4
|
||||
%47 = OpConstantNull %_arr_Inner_std140_uint_4
|
||||
%_ptr_Function_Inner = OpTypePointer Function %Inner
|
||||
%_ptr_Function_Inner_std140 = OpTypePointer Function %Inner_std140
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%Outer = OpTypeStruct %_arr_Inner_uint_4
|
||||
%60 = OpTypeFunction %Outer %Outer_std140
|
||||
%_arr_Outer_uint_4 = OpTypeArray %Outer %uint_4
|
||||
%68 = OpTypeFunction %_arr_Outer_uint_4 %_arr_Outer_std140_uint_4
|
||||
%_ptr_Function__arr_Outer_uint_4 = OpTypePointer Function %_arr_Outer_uint_4
|
||||
%75 = OpConstantNull %_arr_Outer_uint_4
|
||||
%_ptr_Function__arr_Outer_std140_uint_4 = OpTypePointer Function %_arr_Outer_std140_uint_4
|
||||
%88 = OpConstantNull %_arr_Outer_std140_uint_4
|
||||
%_ptr_Function_Outer = OpTypePointer Function %Outer
|
||||
%_ptr_Function_Outer_std140 = OpTypePointer Function %Outer_std140
|
||||
%100 = OpTypeFunction %mat4v2float
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%uint_3 = OpConstant %uint 3
|
||||
%uint_2 = OpConstant %uint 2
|
||||
%_ptr_Uniform_v2float = OpTypePointer Uniform %v2float
|
||||
%void = OpTypeVoid
|
||||
%116 = OpTypeFunction %void
|
||||
%int = OpTypeInt 32 1
|
||||
%int_1 = OpConstant %int 1
|
||||
%_ptr_Uniform__arr_Outer_std140_uint_4 = OpTypePointer Uniform %_arr_Outer_std140_uint_4
|
||||
%_ptr_Uniform_Outer_std140 = OpTypePointer Uniform %Outer_std140
|
||||
%_ptr_Uniform__arr_Inner_std140_uint_4 = OpTypePointer Uniform %_arr_Inner_std140_uint_4
|
||||
%_ptr_Uniform_Inner_std140 = OpTypePointer Uniform %Inner_std140
|
||||
%_ptr_Uniform_float = OpTypePointer Uniform %float
|
||||
%conv_Inner = OpFunction %Inner None %12
|
||||
%val = OpFunctionParameter %Inner_std140
|
||||
%17 = OpLabel
|
||||
%18 = OpCompositeExtract %v2float %val 0
|
||||
%19 = OpCompositeExtract %v2float %val 1
|
||||
%20 = OpCompositeExtract %v2float %val 2
|
||||
%21 = OpCompositeExtract %v2float %val 3
|
||||
%22 = OpCompositeConstruct %mat4v2float %18 %19 %20 %21
|
||||
%23 = OpCompositeConstruct %Inner %22
|
||||
OpReturnValue %23
|
||||
OpFunctionEnd
|
||||
%conv_arr_4_Inner = OpFunction %_arr_Inner_uint_4 None %24
|
||||
%val_0 = OpFunctionParameter %_arr_Inner_std140_uint_4
|
||||
%28 = OpLabel
|
||||
%arr = OpVariable %_ptr_Function__arr_Inner_uint_4 Function %31
|
||||
%i = OpVariable %_ptr_Function_uint Function %34
|
||||
%var_for_index = OpVariable %_ptr_Function__arr_Inner_std140_uint_4 Function %47
|
||||
OpBranch %35
|
||||
%35 = OpLabel
|
||||
OpLoopMerge %36 %37 None
|
||||
OpBranch %38
|
||||
%38 = OpLabel
|
||||
%40 = OpLoad %uint %i
|
||||
%41 = OpULessThan %bool %40 %uint_4
|
||||
%39 = OpLogicalNot %bool %41
|
||||
OpSelectionMerge %43 None
|
||||
OpBranchConditional %39 %44 %43
|
||||
%44 = OpLabel
|
||||
OpBranch %36
|
||||
%43 = OpLabel
|
||||
OpStore %var_for_index %val_0
|
||||
%48 = OpLoad %uint %i
|
||||
%50 = OpAccessChain %_ptr_Function_Inner %arr %48
|
||||
%52 = OpLoad %uint %i
|
||||
%54 = OpAccessChain %_ptr_Function_Inner_std140 %var_for_index %52
|
||||
%55 = OpLoad %Inner_std140 %54
|
||||
%51 = OpFunctionCall %Inner %conv_Inner %55
|
||||
OpStore %50 %51
|
||||
OpBranch %37
|
||||
%37 = OpLabel
|
||||
%56 = OpLoad %uint %i
|
||||
%58 = OpIAdd %uint %56 %uint_1
|
||||
OpStore %i %58
|
||||
OpBranch %35
|
||||
%36 = OpLabel
|
||||
%59 = OpLoad %_arr_Inner_uint_4 %arr
|
||||
OpReturnValue %59
|
||||
OpFunctionEnd
|
||||
%conv_Outer = OpFunction %Outer None %60
|
||||
%val_1 = OpFunctionParameter %Outer_std140
|
||||
%64 = OpLabel
|
||||
%66 = OpCompositeExtract %_arr_Inner_std140_uint_4 %val_1 0
|
||||
%65 = OpFunctionCall %_arr_Inner_uint_4 %conv_arr_4_Inner %66
|
||||
%67 = OpCompositeConstruct %Outer %65
|
||||
OpReturnValue %67
|
||||
OpFunctionEnd
|
||||
%conv_arr_4_Outer = OpFunction %_arr_Outer_uint_4 None %68
|
||||
%val_2 = OpFunctionParameter %_arr_Outer_std140_uint_4
|
||||
%72 = OpLabel
|
||||
%arr_0 = OpVariable %_ptr_Function__arr_Outer_uint_4 Function %75
|
||||
%i_0 = OpVariable %_ptr_Function_uint Function %34
|
||||
%var_for_index_1 = OpVariable %_ptr_Function__arr_Outer_std140_uint_4 Function %88
|
||||
OpBranch %77
|
||||
%77 = OpLabel
|
||||
OpLoopMerge %78 %79 None
|
||||
OpBranch %80
|
||||
%80 = OpLabel
|
||||
%82 = OpLoad %uint %i_0
|
||||
%83 = OpULessThan %bool %82 %uint_4
|
||||
%81 = OpLogicalNot %bool %83
|
||||
OpSelectionMerge %84 None
|
||||
OpBranchConditional %81 %85 %84
|
||||
%85 = OpLabel
|
||||
OpBranch %78
|
||||
%84 = OpLabel
|
||||
OpStore %var_for_index_1 %val_2
|
||||
%89 = OpLoad %uint %i_0
|
||||
%91 = OpAccessChain %_ptr_Function_Outer %arr_0 %89
|
||||
%93 = OpLoad %uint %i_0
|
||||
%95 = OpAccessChain %_ptr_Function_Outer_std140 %var_for_index_1 %93
|
||||
%96 = OpLoad %Outer_std140 %95
|
||||
%92 = OpFunctionCall %Outer %conv_Outer %96
|
||||
OpStore %91 %92
|
||||
OpBranch %79
|
||||
%79 = OpLabel
|
||||
%97 = OpLoad %uint %i_0
|
||||
%98 = OpIAdd %uint %97 %uint_1
|
||||
OpStore %i_0 %98
|
||||
OpBranch %77
|
||||
%78 = OpLabel
|
||||
%99 = OpLoad %_arr_Outer_uint_4 %arr_0
|
||||
OpReturnValue %99
|
||||
OpFunctionEnd
|
||||
%load_a_3_a_2_m = OpFunction %mat4v2float None %100
|
||||
%102 = OpLabel
|
||||
%107 = OpAccessChain %_ptr_Uniform_v2float %a %uint_0 %uint_3 %uint_0 %uint_2 %uint_0
|
||||
%108 = OpLoad %v2float %107
|
||||
%109 = OpAccessChain %_ptr_Uniform_v2float %a %uint_0 %uint_3 %uint_0 %uint_2 %uint_1
|
||||
%110 = OpLoad %v2float %109
|
||||
%111 = OpAccessChain %_ptr_Uniform_v2float %a %uint_0 %uint_3 %uint_0 %uint_2 %uint_2
|
||||
%112 = OpLoad %v2float %111
|
||||
%113 = OpAccessChain %_ptr_Uniform_v2float %a %uint_0 %uint_3 %uint_0 %uint_2 %uint_3
|
||||
%114 = OpLoad %v2float %113
|
||||
%115 = OpCompositeConstruct %mat4v2float %108 %110 %112 %114
|
||||
OpReturnValue %115
|
||||
OpFunctionEnd
|
||||
%f = OpFunction %void None %116
|
||||
%119 = OpLabel
|
||||
%124 = OpAccessChain %_ptr_Uniform__arr_Outer_std140_uint_4 %a %uint_0
|
||||
%125 = OpLoad %_arr_Outer_std140_uint_4 %124
|
||||
%122 = OpFunctionCall %_arr_Outer_uint_4 %conv_arr_4_Outer %125
|
||||
%128 = OpAccessChain %_ptr_Uniform_Outer_std140 %a %uint_0 %uint_3
|
||||
%129 = OpLoad %Outer_std140 %128
|
||||
%126 = OpFunctionCall %Outer %conv_Outer %129
|
||||
%132 = OpAccessChain %_ptr_Uniform__arr_Inner_std140_uint_4 %a %uint_0 %uint_3 %uint_0
|
||||
%133 = OpLoad %_arr_Inner_std140_uint_4 %132
|
||||
%130 = OpFunctionCall %_arr_Inner_uint_4 %conv_arr_4_Inner %133
|
||||
%136 = OpAccessChain %_ptr_Uniform_Inner_std140 %a %uint_0 %uint_3 %uint_0 %uint_2
|
||||
%137 = OpLoad %Inner_std140 %136
|
||||
%134 = OpFunctionCall %Inner %conv_Inner %137
|
||||
%138 = OpFunctionCall %mat4v2float %load_a_3_a_2_m
|
||||
%139 = OpAccessChain %_ptr_Uniform_v2float %a %uint_0 %uint_3 %uint_0 %uint_2 %uint_1
|
||||
%140 = OpLoad %v2float %139
|
||||
%142 = OpAccessChain %_ptr_Uniform__arr_Outer_std140_uint_4 %a %uint_0
|
||||
%143 = OpLoad %_arr_Outer_std140_uint_4 %142
|
||||
%141 = OpFunctionCall %_arr_Outer_uint_4 %conv_arr_4_Outer %143
|
||||
%145 = OpAccessChain %_ptr_Uniform_Outer_std140 %a %uint_0 %uint_3
|
||||
%146 = OpLoad %Outer_std140 %145
|
||||
%144 = OpFunctionCall %Outer %conv_Outer %146
|
||||
%148 = OpAccessChain %_ptr_Uniform__arr_Inner_std140_uint_4 %a %uint_0 %uint_3 %uint_0
|
||||
%149 = OpLoad %_arr_Inner_std140_uint_4 %148
|
||||
%147 = OpFunctionCall %_arr_Inner_uint_4 %conv_arr_4_Inner %149
|
||||
%151 = OpAccessChain %_ptr_Uniform_Inner_std140 %a %uint_0 %uint_3 %uint_0 %uint_2
|
||||
%152 = OpLoad %Inner_std140 %151
|
||||
%150 = OpFunctionCall %Inner %conv_Inner %152
|
||||
%153 = OpFunctionCall %mat4v2float %load_a_3_a_2_m
|
||||
%154 = OpAccessChain %_ptr_Uniform_v2float %a %uint_0 %uint_3 %uint_0 %uint_2 %uint_1
|
||||
%155 = OpLoad %v2float %154
|
||||
%157 = OpAccessChain %_ptr_Uniform_float %a %uint_0 %uint_3 %uint_0 %uint_2 %uint_1 %34
|
||||
%158 = OpLoad %float %157
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@@ -0,0 +1,27 @@
|
||||
struct Inner {
|
||||
m : mat4x2<f32>,
|
||||
}
|
||||
|
||||
struct Outer {
|
||||
a : array<Inner, 4>,
|
||||
}
|
||||
|
||||
@group(0) @binding(0) var<uniform> a : array<Outer, 4>;
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn f() {
|
||||
let I = 1;
|
||||
let p_a = &(a);
|
||||
let p_a_3 = &((*(p_a))[3]);
|
||||
let p_a_3_a = &((*(p_a_3)).a);
|
||||
let p_a_3_a_2 = &((*(p_a_3_a))[2]);
|
||||
let p_a_3_a_2_m = &((*(p_a_3_a_2)).m);
|
||||
let p_a_3_a_2_m_1 = &((*(p_a_3_a_2_m))[1]);
|
||||
let l_a : array<Outer, 4> = *(p_a);
|
||||
let l_a_3 : Outer = *(p_a_3);
|
||||
let l_a_3_a : array<Inner, 4> = *(p_a_3_a);
|
||||
let l_a_3_a_2 : Inner = *(p_a_3_a_2);
|
||||
let l_a_3_a_2_m : mat4x2<f32> = *(p_a_3_a_2_m);
|
||||
let l_a_3_a_2_m_1 : vec2<f32> = *(p_a_3_a_2_m_1);
|
||||
let l_a_3_a_2_m_1_0 : f32 = (*(p_a_3_a_2_m_1))[0];
|
||||
}
|
||||
14
test/tint/buffer/uniform/std140/mat4x2/to_builtin.wgsl
Normal file
14
test/tint/buffer/uniform/std140/mat4x2/to_builtin.wgsl
Normal file
@@ -0,0 +1,14 @@
|
||||
struct S {
|
||||
before : i32,
|
||||
m : mat4x2<f32>,
|
||||
after : i32,
|
||||
}
|
||||
|
||||
@group(0) @binding(0) var<uniform> u : array<S, 4>;
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn f() {
|
||||
let t = transpose(u[2].m);
|
||||
let l = length(u[0].m[1].yx);
|
||||
let a = abs(u[0].m[1].yx.x);
|
||||
}
|
||||
@@ -0,0 +1,23 @@
|
||||
cbuffer cbuffer_u : register(b0, space0) {
|
||||
uint4 u[12];
|
||||
};
|
||||
|
||||
float4x2 tint_symbol(uint4 buffer[12], 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)));
|
||||
}
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void f() {
|
||||
const float2x4 t = transpose(tint_symbol(u, 104u));
|
||||
const float l = length(asfloat(u[1].xy).yx);
|
||||
const float a = abs(asfloat(u[1].xy).yx.x);
|
||||
return;
|
||||
}
|
||||
@@ -0,0 +1,23 @@
|
||||
cbuffer cbuffer_u : register(b0, space0) {
|
||||
uint4 u[12];
|
||||
};
|
||||
|
||||
float4x2 tint_symbol(uint4 buffer[12], 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)));
|
||||
}
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void f() {
|
||||
const float2x4 t = transpose(tint_symbol(u, 104u));
|
||||
const float l = length(asfloat(u[1].xy).yx);
|
||||
const float a = abs(asfloat(u[1].xy).yx.x);
|
||||
return;
|
||||
}
|
||||
@@ -0,0 +1,40 @@
|
||||
#version 310 es
|
||||
|
||||
struct S {
|
||||
int before;
|
||||
mat4x2 m;
|
||||
int after;
|
||||
};
|
||||
|
||||
struct S_std140 {
|
||||
int before;
|
||||
vec2 m_0;
|
||||
vec2 m_1;
|
||||
vec2 m_2;
|
||||
vec2 m_3;
|
||||
int after;
|
||||
};
|
||||
|
||||
struct u_block {
|
||||
S_std140 inner[4];
|
||||
};
|
||||
|
||||
layout(binding = 0) uniform u_block_1 {
|
||||
S_std140 inner[4];
|
||||
} u;
|
||||
|
||||
mat4x2 load_u_2_m() {
|
||||
return mat4x2(u.inner[2u].m_0, u.inner[2u].m_1, u.inner[2u].m_2, u.inner[2u].m_3);
|
||||
}
|
||||
|
||||
void f() {
|
||||
mat2x4 t = transpose(load_u_2_m());
|
||||
float l = length(u.inner[0u].m_1.yx);
|
||||
float a = abs(u.inner[0u].m_1.yx[0u]);
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
void main() {
|
||||
f();
|
||||
return;
|
||||
}
|
||||
@@ -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 {
|
||||
/* 0x0000 */ int before;
|
||||
/* 0x0004 */ tint_array<int8_t, 4> tint_pad;
|
||||
/* 0x0008 */ float4x2 m;
|
||||
/* 0x0028 */ int after;
|
||||
/* 0x002c */ tint_array<int8_t, 4> tint_pad_1;
|
||||
};
|
||||
|
||||
kernel void f(const constant tint_array<S, 4>* tint_symbol [[buffer(0)]]) {
|
||||
float2x4 const t = transpose((*(tint_symbol))[2].m);
|
||||
float const l = length(float2((*(tint_symbol))[0].m[1]).yx);
|
||||
float const a = fabs(float2((*(tint_symbol))[0].m[1]).yx[0]);
|
||||
return;
|
||||
}
|
||||
|
||||
@@ -0,0 +1,84 @@
|
||||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 48
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
%38 = 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 "m_2"
|
||||
OpMemberName %S_std140 4 "m_3"
|
||||
OpMemberName %S_std140 5 "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
|
||||
OpMemberDecorate %S_std140 4 Offset 32
|
||||
OpMemberDecorate %S_std140 5 Offset 40
|
||||
OpDecorate %_arr_S_std140_uint_4 ArrayStride 48
|
||||
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 %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
|
||||
%mat4v2float = OpTypeMatrix %v2float 4
|
||||
%11 = OpTypeFunction %mat4v2float
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%uint_2 = OpConstant %uint 2
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%_ptr_Uniform_v2float = OpTypePointer Uniform %v2float
|
||||
%uint_3 = OpConstant %uint 3
|
||||
%void = OpTypeVoid
|
||||
%29 = OpTypeFunction %void
|
||||
%v4float = OpTypeVector %float 4
|
||||
%mat2v4float = OpTypeMatrix %v4float 2
|
||||
%39 = OpConstantNull %uint
|
||||
%load_u_2_m = OpFunction %mat4v2float 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
|
||||
%24 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %uint_2 %uint_3
|
||||
%25 = OpLoad %v2float %24
|
||||
%26 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %uint_2 %uint_4
|
||||
%27 = OpLoad %v2float %26
|
||||
%28 = OpCompositeConstruct %mat4v2float %20 %22 %25 %27
|
||||
OpReturnValue %28
|
||||
OpFunctionEnd
|
||||
%f = OpFunction %void None %29
|
||||
%32 = OpLabel
|
||||
%36 = OpFunctionCall %mat4v2float %load_u_2_m
|
||||
%33 = OpTranspose %mat2v4float %36
|
||||
%40 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %39 %uint_2
|
||||
%41 = OpLoad %v2float %40
|
||||
%42 = OpVectorShuffle %v2float %41 %41 1 0
|
||||
%37 = OpExtInst %float %38 Length %42
|
||||
%44 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %39 %uint_2
|
||||
%45 = OpLoad %v2float %44
|
||||
%46 = OpVectorShuffle %v2float %45 %45 1 0
|
||||
%47 = OpCompositeExtract %float %46 0
|
||||
%43 = OpExtInst %float %38 FAbs %47
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@@ -0,0 +1,14 @@
|
||||
struct S {
|
||||
before : i32,
|
||||
m : mat4x2<f32>,
|
||||
after : i32,
|
||||
}
|
||||
|
||||
@group(0) @binding(0) var<uniform> u : array<S, 4>;
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn f() {
|
||||
let t = transpose(u[2].m);
|
||||
let l = length(u[0].m[1].yx);
|
||||
let a = abs(u[0].m[1].yx.x);
|
||||
}
|
||||
22
test/tint/buffer/uniform/std140/mat4x2/to_fn.wgsl
Normal file
22
test/tint/buffer/uniform/std140/mat4x2/to_fn.wgsl
Normal file
@@ -0,0 +1,22 @@
|
||||
struct S {
|
||||
before : i32,
|
||||
m : mat4x2<f32>,
|
||||
after : i32,
|
||||
}
|
||||
|
||||
@group(0) @binding(0) var<uniform> u : array<S, 4>;
|
||||
|
||||
fn a(a : array<S, 4>) {}
|
||||
fn b(s : S) {}
|
||||
fn c(m : mat4x2<f32>) {}
|
||||
fn d(v : vec2<f32>) {}
|
||||
fn e(f : f32) {}
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn f() {
|
||||
a(u);
|
||||
b(u[2]);
|
||||
c(u[2].m);
|
||||
d(u[0].m[1].yx);
|
||||
e(u[0].m[1].yx.x);
|
||||
}
|
||||
@@ -0,0 +1,64 @@
|
||||
struct S {
|
||||
int before;
|
||||
float4x2 m;
|
||||
int after;
|
||||
};
|
||||
|
||||
cbuffer cbuffer_u : register(b0, space0) {
|
||||
uint4 u[12];
|
||||
};
|
||||
|
||||
void a(S a_1[4]) {
|
||||
}
|
||||
|
||||
void b(S s) {
|
||||
}
|
||||
|
||||
void c(float4x2 m) {
|
||||
}
|
||||
|
||||
void d(float2 v) {
|
||||
}
|
||||
|
||||
void e(float f_1) {
|
||||
}
|
||||
|
||||
float4x2 tint_symbol_3(uint4 buffer[12], 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)));
|
||||
}
|
||||
|
||||
S tint_symbol_1(uint4 buffer[12], uint offset) {
|
||||
const uint scalar_offset_4 = ((offset + 0u)) / 4;
|
||||
const uint scalar_offset_5 = ((offset + 40u)) / 4;
|
||||
const S tint_symbol_5 = {asint(buffer[scalar_offset_4 / 4][scalar_offset_4 % 4]), tint_symbol_3(buffer, (offset + 8u)), asint(buffer[scalar_offset_5 / 4][scalar_offset_5 % 4])};
|
||||
return tint_symbol_5;
|
||||
}
|
||||
|
||||
typedef S tint_symbol_ret[4];
|
||||
tint_symbol_ret tint_symbol(uint4 buffer[12], uint offset) {
|
||||
S arr[4] = (S[4])0;
|
||||
{
|
||||
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
|
||||
arr[i] = tint_symbol_1(buffer, (offset + (i * 48u)));
|
||||
}
|
||||
}
|
||||
return arr;
|
||||
}
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void f() {
|
||||
a(tint_symbol(u, 0u));
|
||||
b(tint_symbol_1(u, 96u));
|
||||
c(tint_symbol_3(u, 104u));
|
||||
d(asfloat(u[1].xy).yx);
|
||||
e(asfloat(u[1].xy).yx.x);
|
||||
return;
|
||||
}
|
||||
@@ -0,0 +1,64 @@
|
||||
struct S {
|
||||
int before;
|
||||
float4x2 m;
|
||||
int after;
|
||||
};
|
||||
|
||||
cbuffer cbuffer_u : register(b0, space0) {
|
||||
uint4 u[12];
|
||||
};
|
||||
|
||||
void a(S a_1[4]) {
|
||||
}
|
||||
|
||||
void b(S s) {
|
||||
}
|
||||
|
||||
void c(float4x2 m) {
|
||||
}
|
||||
|
||||
void d(float2 v) {
|
||||
}
|
||||
|
||||
void e(float f_1) {
|
||||
}
|
||||
|
||||
float4x2 tint_symbol_3(uint4 buffer[12], 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)));
|
||||
}
|
||||
|
||||
S tint_symbol_1(uint4 buffer[12], uint offset) {
|
||||
const uint scalar_offset_4 = ((offset + 0u)) / 4;
|
||||
const uint scalar_offset_5 = ((offset + 40u)) / 4;
|
||||
const S tint_symbol_5 = {asint(buffer[scalar_offset_4 / 4][scalar_offset_4 % 4]), tint_symbol_3(buffer, (offset + 8u)), asint(buffer[scalar_offset_5 / 4][scalar_offset_5 % 4])};
|
||||
return tint_symbol_5;
|
||||
}
|
||||
|
||||
typedef S tint_symbol_ret[4];
|
||||
tint_symbol_ret tint_symbol(uint4 buffer[12], uint offset) {
|
||||
S arr[4] = (S[4])0;
|
||||
{
|
||||
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
|
||||
arr[i] = tint_symbol_1(buffer, (offset + (i * 48u)));
|
||||
}
|
||||
}
|
||||
return arr;
|
||||
}
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void f() {
|
||||
a(tint_symbol(u, 0u));
|
||||
b(tint_symbol_1(u, 96u));
|
||||
c(tint_symbol_3(u, 104u));
|
||||
d(asfloat(u[1].xy).yx);
|
||||
e(asfloat(u[1].xy).yx.x);
|
||||
return;
|
||||
}
|
||||
@@ -0,0 +1,72 @@
|
||||
#version 310 es
|
||||
|
||||
struct S {
|
||||
int before;
|
||||
mat4x2 m;
|
||||
int after;
|
||||
};
|
||||
|
||||
struct S_std140 {
|
||||
int before;
|
||||
vec2 m_0;
|
||||
vec2 m_1;
|
||||
vec2 m_2;
|
||||
vec2 m_3;
|
||||
int after;
|
||||
};
|
||||
|
||||
struct u_block {
|
||||
S_std140 inner[4];
|
||||
};
|
||||
|
||||
layout(binding = 0) uniform u_block_1 {
|
||||
S_std140 inner[4];
|
||||
} u;
|
||||
|
||||
void a(S a_1[4]) {
|
||||
}
|
||||
|
||||
void b(S s) {
|
||||
}
|
||||
|
||||
void c(mat4x2 m) {
|
||||
}
|
||||
|
||||
void d(vec2 v) {
|
||||
}
|
||||
|
||||
void e(float f_1) {
|
||||
}
|
||||
|
||||
S conv_S(S_std140 val) {
|
||||
S tint_symbol = S(val.before, mat4x2(val.m_0, val.m_1, val.m_2, val.m_3), val.after);
|
||||
return tint_symbol;
|
||||
}
|
||||
|
||||
S[4] conv_arr_4_S(S_std140 val[4]) {
|
||||
S arr[4] = S[4](S(0, mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f), 0), S(0, mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f), 0), S(0, mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f), 0), S(0, mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 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;
|
||||
}
|
||||
|
||||
mat4x2 load_u_2_m() {
|
||||
return mat4x2(u.inner[2u].m_0, u.inner[2u].m_1, u.inner[2u].m_2, u.inner[2u].m_3);
|
||||
}
|
||||
|
||||
void f() {
|
||||
a(conv_arr_4_S(u.inner));
|
||||
b(conv_S(u.inner[2u]));
|
||||
c(load_u_2_m());
|
||||
d(u.inner[0u].m_1.yx);
|
||||
e(u.inner[0u].m_1.yx[0u]);
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
void main() {
|
||||
f();
|
||||
return;
|
||||
}
|
||||
@@ -0,0 +1,48 @@
|
||||
#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 {
|
||||
/* 0x0000 */ int before;
|
||||
/* 0x0004 */ tint_array<int8_t, 4> tint_pad;
|
||||
/* 0x0008 */ float4x2 m;
|
||||
/* 0x0028 */ int after;
|
||||
/* 0x002c */ tint_array<int8_t, 4> tint_pad_1;
|
||||
};
|
||||
|
||||
void a(tint_array<S, 4> a_1) {
|
||||
}
|
||||
|
||||
void b(S s) {
|
||||
}
|
||||
|
||||
void c(float4x2 m) {
|
||||
}
|
||||
|
||||
void d(float2 v) {
|
||||
}
|
||||
|
||||
void e(float f_1) {
|
||||
}
|
||||
|
||||
kernel void f(const constant tint_array<S, 4>* tint_symbol [[buffer(0)]]) {
|
||||
a(*(tint_symbol));
|
||||
b((*(tint_symbol))[2]);
|
||||
c((*(tint_symbol))[2].m);
|
||||
d(float2((*(tint_symbol))[0].m[1]).yx);
|
||||
e(float2((*(tint_symbol))[0].m[1]).yx[0]);
|
||||
return;
|
||||
}
|
||||
|
||||
@@ -0,0 +1,210 @@
|
||||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 122
|
||||
; 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 %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 "m_2"
|
||||
OpMemberName %S_std140 4 "m_3"
|
||||
OpMemberName %S_std140 5 "after"
|
||||
OpName %u "u"
|
||||
OpName %S "S"
|
||||
OpMemberName %S 0 "before"
|
||||
OpMemberName %S 1 "m"
|
||||
OpMemberName %S 2 "after"
|
||||
OpName %a "a"
|
||||
OpName %a_1 "a_1"
|
||||
OpName %b "b"
|
||||
OpName %s "s"
|
||||
OpName %c "c"
|
||||
OpName %m "m"
|
||||
OpName %d "d"
|
||||
OpName %v "v"
|
||||
OpName %e "e"
|
||||
OpName %f_1 "f_1"
|
||||
OpName %conv_S "conv_S"
|
||||
OpName %val "val"
|
||||
OpName %conv_arr_4_S "conv_arr_4_S"
|
||||
OpName %val_0 "val"
|
||||
OpName %arr "arr"
|
||||
OpName %i "i"
|
||||
OpName %var_for_index "var_for_index"
|
||||
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
|
||||
OpMemberDecorate %S_std140 4 Offset 32
|
||||
OpMemberDecorate %S_std140 5 Offset 40
|
||||
OpDecorate %_arr_S_std140_uint_4 ArrayStride 48
|
||||
OpDecorate %u NonWritable
|
||||
OpDecorate %u DescriptorSet 0
|
||||
OpDecorate %u Binding 0
|
||||
OpMemberDecorate %S 0 Offset 0
|
||||
OpMemberDecorate %S 1 Offset 8
|
||||
OpMemberDecorate %S 1 ColMajor
|
||||
OpMemberDecorate %S 1 MatrixStride 8
|
||||
OpMemberDecorate %S 2 Offset 40
|
||||
OpDecorate %_arr_S_uint_4 ArrayStride 48
|
||||
%int = OpTypeInt 32 1
|
||||
%float = OpTypeFloat 32
|
||||
%v2float = OpTypeVector %float 2
|
||||
%S_std140 = OpTypeStruct %int %v2float %v2float %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
|
||||
%void = OpTypeVoid
|
||||
%mat4v2float = OpTypeMatrix %v2float 4
|
||||
%S = OpTypeStruct %int %mat4v2float %int
|
||||
%_arr_S_uint_4 = OpTypeArray %S %uint_4
|
||||
%11 = OpTypeFunction %void %_arr_S_uint_4
|
||||
%19 = OpTypeFunction %void %S
|
||||
%23 = OpTypeFunction %void %mat4v2float
|
||||
%27 = OpTypeFunction %void %v2float
|
||||
%31 = OpTypeFunction %void %float
|
||||
%35 = OpTypeFunction %S %S_std140
|
||||
%47 = OpTypeFunction %_arr_S_uint_4 %_arr_S_std140_uint_4
|
||||
%_ptr_Function__arr_S_uint_4 = OpTypePointer Function %_arr_S_uint_4
|
||||
%53 = OpConstantNull %_arr_S_uint_4
|
||||
%_ptr_Function_uint = OpTypePointer Function %uint
|
||||
%56 = OpConstantNull %uint
|
||||
%bool = OpTypeBool
|
||||
%_ptr_Function__arr_S_std140_uint_4 = OpTypePointer Function %_arr_S_std140_uint_4
|
||||
%69 = OpConstantNull %_arr_S_std140_uint_4
|
||||
%_ptr_Function_S = OpTypePointer Function %S
|
||||
%_ptr_Function_S_std140 = OpTypePointer Function %S_std140
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%82 = OpTypeFunction %mat4v2float
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%uint_2 = OpConstant %uint 2
|
||||
%_ptr_Uniform_v2float = OpTypePointer Uniform %v2float
|
||||
%uint_3 = OpConstant %uint 3
|
||||
%98 = OpTypeFunction %void
|
||||
%_ptr_Uniform__arr_S_std140_uint_4 = OpTypePointer Uniform %_arr_S_std140_uint_4
|
||||
%_ptr_Uniform_S_std140 = OpTypePointer Uniform %S_std140
|
||||
%a = OpFunction %void None %11
|
||||
%a_1 = OpFunctionParameter %_arr_S_uint_4
|
||||
%18 = OpLabel
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%b = OpFunction %void None %19
|
||||
%s = OpFunctionParameter %S
|
||||
%22 = OpLabel
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%c = OpFunction %void None %23
|
||||
%m = OpFunctionParameter %mat4v2float
|
||||
%26 = OpLabel
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%d = OpFunction %void None %27
|
||||
%v = OpFunctionParameter %v2float
|
||||
%30 = OpLabel
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%e = OpFunction %void None %31
|
||||
%f_1 = OpFunctionParameter %float
|
||||
%34 = OpLabel
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%conv_S = OpFunction %S None %35
|
||||
%val = OpFunctionParameter %S_std140
|
||||
%38 = OpLabel
|
||||
%39 = OpCompositeExtract %int %val 0
|
||||
%40 = OpCompositeExtract %v2float %val 1
|
||||
%41 = OpCompositeExtract %v2float %val 2
|
||||
%42 = OpCompositeExtract %v2float %val 3
|
||||
%43 = OpCompositeExtract %v2float %val 4
|
||||
%44 = OpCompositeConstruct %mat4v2float %40 %41 %42 %43
|
||||
%45 = OpCompositeExtract %int %val 5
|
||||
%46 = OpCompositeConstruct %S %39 %44 %45
|
||||
OpReturnValue %46
|
||||
OpFunctionEnd
|
||||
%conv_arr_4_S = OpFunction %_arr_S_uint_4 None %47
|
||||
%val_0 = OpFunctionParameter %_arr_S_std140_uint_4
|
||||
%50 = OpLabel
|
||||
%arr = OpVariable %_ptr_Function__arr_S_uint_4 Function %53
|
||||
%i = OpVariable %_ptr_Function_uint Function %56
|
||||
%var_for_index = OpVariable %_ptr_Function__arr_S_std140_uint_4 Function %69
|
||||
OpBranch %57
|
||||
%57 = OpLabel
|
||||
OpLoopMerge %58 %59 None
|
||||
OpBranch %60
|
||||
%60 = OpLabel
|
||||
%62 = OpLoad %uint %i
|
||||
%63 = OpULessThan %bool %62 %uint_4
|
||||
%61 = OpLogicalNot %bool %63
|
||||
OpSelectionMerge %65 None
|
||||
OpBranchConditional %61 %66 %65
|
||||
%66 = OpLabel
|
||||
OpBranch %58
|
||||
%65 = OpLabel
|
||||
OpStore %var_for_index %val_0
|
||||
%70 = OpLoad %uint %i
|
||||
%72 = OpAccessChain %_ptr_Function_S %arr %70
|
||||
%74 = OpLoad %uint %i
|
||||
%76 = OpAccessChain %_ptr_Function_S_std140 %var_for_index %74
|
||||
%77 = OpLoad %S_std140 %76
|
||||
%73 = OpFunctionCall %S %conv_S %77
|
||||
OpStore %72 %73
|
||||
OpBranch %59
|
||||
%59 = OpLabel
|
||||
%78 = OpLoad %uint %i
|
||||
%80 = OpIAdd %uint %78 %uint_1
|
||||
OpStore %i %80
|
||||
OpBranch %57
|
||||
%58 = OpLabel
|
||||
%81 = OpLoad %_arr_S_uint_4 %arr
|
||||
OpReturnValue %81
|
||||
OpFunctionEnd
|
||||
%load_u_2_m = OpFunction %mat4v2float None %82
|
||||
%84 = OpLabel
|
||||
%88 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %uint_2 %uint_1
|
||||
%89 = OpLoad %v2float %88
|
||||
%90 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %uint_2 %uint_2
|
||||
%91 = OpLoad %v2float %90
|
||||
%93 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %uint_2 %uint_3
|
||||
%94 = OpLoad %v2float %93
|
||||
%95 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %uint_2 %uint_4
|
||||
%96 = OpLoad %v2float %95
|
||||
%97 = OpCompositeConstruct %mat4v2float %89 %91 %94 %96
|
||||
OpReturnValue %97
|
||||
OpFunctionEnd
|
||||
%f = OpFunction %void None %98
|
||||
%100 = OpLabel
|
||||
%104 = OpAccessChain %_ptr_Uniform__arr_S_std140_uint_4 %u %uint_0
|
||||
%105 = OpLoad %_arr_S_std140_uint_4 %104
|
||||
%102 = OpFunctionCall %_arr_S_uint_4 %conv_arr_4_S %105
|
||||
%101 = OpFunctionCall %void %a %102
|
||||
%109 = OpAccessChain %_ptr_Uniform_S_std140 %u %uint_0 %uint_2
|
||||
%110 = OpLoad %S_std140 %109
|
||||
%107 = OpFunctionCall %S %conv_S %110
|
||||
%106 = OpFunctionCall %void %b %107
|
||||
%112 = OpFunctionCall %mat4v2float %load_u_2_m
|
||||
%111 = OpFunctionCall %void %c %112
|
||||
%114 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %56 %uint_2
|
||||
%115 = OpLoad %v2float %114
|
||||
%116 = OpVectorShuffle %v2float %115 %115 1 0
|
||||
%113 = OpFunctionCall %void %d %116
|
||||
%118 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %56 %uint_2
|
||||
%119 = OpLoad %v2float %118
|
||||
%120 = OpVectorShuffle %v2float %119 %119 1 0
|
||||
%121 = OpCompositeExtract %float %120 0
|
||||
%117 = OpFunctionCall %void %e %121
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@@ -0,0 +1,31 @@
|
||||
struct S {
|
||||
before : i32,
|
||||
m : mat4x2<f32>,
|
||||
after : i32,
|
||||
}
|
||||
|
||||
@group(0) @binding(0) var<uniform> u : array<S, 4>;
|
||||
|
||||
fn a(a : array<S, 4>) {
|
||||
}
|
||||
|
||||
fn b(s : S) {
|
||||
}
|
||||
|
||||
fn c(m : mat4x2<f32>) {
|
||||
}
|
||||
|
||||
fn d(v : vec2<f32>) {
|
||||
}
|
||||
|
||||
fn e(f : f32) {
|
||||
}
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn f() {
|
||||
a(u);
|
||||
b(u[2]);
|
||||
c(u[2].m);
|
||||
d(u[0].m[1].yx);
|
||||
e(u[0].m[1].yx.x);
|
||||
}
|
||||
16
test/tint/buffer/uniform/std140/mat4x2/to_private.wgsl
Normal file
16
test/tint/buffer/uniform/std140/mat4x2/to_private.wgsl
Normal file
@@ -0,0 +1,16 @@
|
||||
struct S {
|
||||
before : i32,
|
||||
m : mat4x2<f32>,
|
||||
after : i32,
|
||||
}
|
||||
|
||||
@group(0) @binding(0) var<uniform> u : array<S, 4>;
|
||||
var<private> p : array<S, 4>;
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn f() {
|
||||
p = u;
|
||||
p[1] = u[2];
|
||||
p[3].m = u[2].m;
|
||||
p[1].m[0] = u[0].m[1].yx;
|
||||
}
|
||||
@@ -0,0 +1,49 @@
|
||||
struct S {
|
||||
int before;
|
||||
float4x2 m;
|
||||
int after;
|
||||
};
|
||||
|
||||
cbuffer cbuffer_u : register(b0, space0) {
|
||||
uint4 u[12];
|
||||
};
|
||||
static S p[4] = (S[4])0;
|
||||
|
||||
float4x2 tint_symbol_3(uint4 buffer[12], 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)));
|
||||
}
|
||||
|
||||
S tint_symbol_1(uint4 buffer[12], uint offset) {
|
||||
const uint scalar_offset_4 = ((offset + 0u)) / 4;
|
||||
const uint scalar_offset_5 = ((offset + 40u)) / 4;
|
||||
const S tint_symbol_5 = {asint(buffer[scalar_offset_4 / 4][scalar_offset_4 % 4]), tint_symbol_3(buffer, (offset + 8u)), asint(buffer[scalar_offset_5 / 4][scalar_offset_5 % 4])};
|
||||
return tint_symbol_5;
|
||||
}
|
||||
|
||||
typedef S tint_symbol_ret[4];
|
||||
tint_symbol_ret tint_symbol(uint4 buffer[12], uint offset) {
|
||||
S arr[4] = (S[4])0;
|
||||
{
|
||||
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
|
||||
arr[i] = tint_symbol_1(buffer, (offset + (i * 48u)));
|
||||
}
|
||||
}
|
||||
return arr;
|
||||
}
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void f() {
|
||||
p = tint_symbol(u, 0u);
|
||||
p[1] = tint_symbol_1(u, 96u);
|
||||
p[3].m = tint_symbol_3(u, 104u);
|
||||
p[1].m[0] = asfloat(u[1].xy).yx;
|
||||
return;
|
||||
}
|
||||
@@ -0,0 +1,49 @@
|
||||
struct S {
|
||||
int before;
|
||||
float4x2 m;
|
||||
int after;
|
||||
};
|
||||
|
||||
cbuffer cbuffer_u : register(b0, space0) {
|
||||
uint4 u[12];
|
||||
};
|
||||
static S p[4] = (S[4])0;
|
||||
|
||||
float4x2 tint_symbol_3(uint4 buffer[12], 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)));
|
||||
}
|
||||
|
||||
S tint_symbol_1(uint4 buffer[12], uint offset) {
|
||||
const uint scalar_offset_4 = ((offset + 0u)) / 4;
|
||||
const uint scalar_offset_5 = ((offset + 40u)) / 4;
|
||||
const S tint_symbol_5 = {asint(buffer[scalar_offset_4 / 4][scalar_offset_4 % 4]), tint_symbol_3(buffer, (offset + 8u)), asint(buffer[scalar_offset_5 / 4][scalar_offset_5 % 4])};
|
||||
return tint_symbol_5;
|
||||
}
|
||||
|
||||
typedef S tint_symbol_ret[4];
|
||||
tint_symbol_ret tint_symbol(uint4 buffer[12], uint offset) {
|
||||
S arr[4] = (S[4])0;
|
||||
{
|
||||
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
|
||||
arr[i] = tint_symbol_1(buffer, (offset + (i * 48u)));
|
||||
}
|
||||
}
|
||||
return arr;
|
||||
}
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void f() {
|
||||
p = tint_symbol(u, 0u);
|
||||
p[1] = tint_symbol_1(u, 96u);
|
||||
p[3].m = tint_symbol_3(u, 104u);
|
||||
p[1].m[0] = asfloat(u[1].xy).yx;
|
||||
return;
|
||||
}
|
||||
@@ -0,0 +1,57 @@
|
||||
#version 310 es
|
||||
|
||||
struct S {
|
||||
int before;
|
||||
mat4x2 m;
|
||||
int after;
|
||||
};
|
||||
|
||||
struct S_std140 {
|
||||
int before;
|
||||
vec2 m_0;
|
||||
vec2 m_1;
|
||||
vec2 m_2;
|
||||
vec2 m_3;
|
||||
int after;
|
||||
};
|
||||
|
||||
struct u_block {
|
||||
S_std140 inner[4];
|
||||
};
|
||||
|
||||
layout(binding = 0) uniform u_block_1 {
|
||||
S_std140 inner[4];
|
||||
} u;
|
||||
|
||||
S p[4] = S[4](S(0, mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f), 0), S(0, mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f), 0), S(0, mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f), 0), S(0, mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f), 0));
|
||||
S conv_S(S_std140 val) {
|
||||
S tint_symbol = S(val.before, mat4x2(val.m_0, val.m_1, val.m_2, val.m_3), val.after);
|
||||
return tint_symbol;
|
||||
}
|
||||
|
||||
S[4] conv_arr_4_S(S_std140 val[4]) {
|
||||
S arr[4] = S[4](S(0, mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f), 0), S(0, mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f), 0), S(0, mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f), 0), S(0, mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 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;
|
||||
}
|
||||
|
||||
mat4x2 load_u_2_m() {
|
||||
return mat4x2(u.inner[2u].m_0, u.inner[2u].m_1, u.inner[2u].m_2, u.inner[2u].m_3);
|
||||
}
|
||||
|
||||
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;
|
||||
}
|
||||
@@ -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];
|
||||
};
|
||||
|
||||
struct S {
|
||||
/* 0x0000 */ int before;
|
||||
/* 0x0004 */ tint_array<int8_t, 4> tint_pad;
|
||||
/* 0x0008 */ float4x2 m;
|
||||
/* 0x0028 */ int after;
|
||||
/* 0x002c */ tint_array<int8_t, 4> tint_pad_1;
|
||||
};
|
||||
|
||||
kernel void f(const constant tint_array<S, 4>* tint_symbol_1 [[buffer(0)]]) {
|
||||
thread tint_array<S, 4> tint_symbol = {};
|
||||
tint_symbol = *(tint_symbol_1);
|
||||
tint_symbol[1] = (*(tint_symbol_1))[2];
|
||||
tint_symbol[3].m = (*(tint_symbol_1))[2].m;
|
||||
tint_symbol[1].m[0] = float2((*(tint_symbol_1))[0].m[1]).yx;
|
||||
return;
|
||||
}
|
||||
|
||||
@@ -0,0 +1,177 @@
|
||||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 104
|
||||
; 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 %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 "m_2"
|
||||
OpMemberName %S_std140 4 "m_3"
|
||||
OpMemberName %S_std140 5 "after"
|
||||
OpName %u "u"
|
||||
OpName %S "S"
|
||||
OpMemberName %S 0 "before"
|
||||
OpMemberName %S 1 "m"
|
||||
OpMemberName %S 2 "after"
|
||||
OpName %p "p"
|
||||
OpName %conv_S "conv_S"
|
||||
OpName %val "val"
|
||||
OpName %conv_arr_4_S "conv_arr_4_S"
|
||||
OpName %val_0 "val"
|
||||
OpName %arr "arr"
|
||||
OpName %i "i"
|
||||
OpName %var_for_index "var_for_index"
|
||||
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
|
||||
OpMemberDecorate %S_std140 4 Offset 32
|
||||
OpMemberDecorate %S_std140 5 Offset 40
|
||||
OpDecorate %_arr_S_std140_uint_4 ArrayStride 48
|
||||
OpDecorate %u NonWritable
|
||||
OpDecorate %u DescriptorSet 0
|
||||
OpDecorate %u Binding 0
|
||||
OpMemberDecorate %S 0 Offset 0
|
||||
OpMemberDecorate %S 1 Offset 8
|
||||
OpMemberDecorate %S 1 ColMajor
|
||||
OpMemberDecorate %S 1 MatrixStride 8
|
||||
OpMemberDecorate %S 2 Offset 40
|
||||
OpDecorate %_arr_S_uint_4 ArrayStride 48
|
||||
%int = OpTypeInt 32 1
|
||||
%float = OpTypeFloat 32
|
||||
%v2float = OpTypeVector %float 2
|
||||
%S_std140 = OpTypeStruct %int %v2float %v2float %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
|
||||
%mat4v2float = OpTypeMatrix %v2float 4
|
||||
%S = OpTypeStruct %int %mat4v2float %int
|
||||
%_arr_S_uint_4 = OpTypeArray %S %uint_4
|
||||
%_ptr_Private__arr_S_uint_4 = OpTypePointer Private %_arr_S_uint_4
|
||||
%16 = OpConstantNull %_arr_S_uint_4
|
||||
%p = OpVariable %_ptr_Private__arr_S_uint_4 Private %16
|
||||
%17 = OpTypeFunction %S %S_std140
|
||||
%29 = OpTypeFunction %_arr_S_uint_4 %_arr_S_std140_uint_4
|
||||
%_ptr_Function__arr_S_uint_4 = OpTypePointer Function %_arr_S_uint_4
|
||||
%_ptr_Function_uint = OpTypePointer Function %uint
|
||||
%37 = OpConstantNull %uint
|
||||
%bool = OpTypeBool
|
||||
%_ptr_Function__arr_S_std140_uint_4 = OpTypePointer Function %_arr_S_std140_uint_4
|
||||
%50 = OpConstantNull %_arr_S_std140_uint_4
|
||||
%_ptr_Function_S = OpTypePointer Function %S
|
||||
%_ptr_Function_S_std140 = OpTypePointer Function %S_std140
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%63 = OpTypeFunction %mat4v2float
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%uint_2 = OpConstant %uint 2
|
||||
%_ptr_Uniform_v2float = OpTypePointer Uniform %v2float
|
||||
%uint_3 = OpConstant %uint 3
|
||||
%void = OpTypeVoid
|
||||
%79 = OpTypeFunction %void
|
||||
%_ptr_Uniform__arr_S_std140_uint_4 = OpTypePointer Uniform %_arr_S_std140_uint_4
|
||||
%int_1 = OpConstant %int 1
|
||||
%_ptr_Private_S = OpTypePointer Private %S
|
||||
%_ptr_Uniform_S_std140 = OpTypePointer Uniform %S_std140
|
||||
%int_3 = OpConstant %int 3
|
||||
%_ptr_Private_mat4v2float = OpTypePointer Private %mat4v2float
|
||||
%98 = OpConstantNull %int
|
||||
%_ptr_Private_v2float = OpTypePointer Private %v2float
|
||||
%conv_S = OpFunction %S None %17
|
||||
%val = OpFunctionParameter %S_std140
|
||||
%20 = OpLabel
|
||||
%21 = OpCompositeExtract %int %val 0
|
||||
%22 = OpCompositeExtract %v2float %val 1
|
||||
%23 = OpCompositeExtract %v2float %val 2
|
||||
%24 = OpCompositeExtract %v2float %val 3
|
||||
%25 = OpCompositeExtract %v2float %val 4
|
||||
%26 = OpCompositeConstruct %mat4v2float %22 %23 %24 %25
|
||||
%27 = OpCompositeExtract %int %val 5
|
||||
%28 = OpCompositeConstruct %S %21 %26 %27
|
||||
OpReturnValue %28
|
||||
OpFunctionEnd
|
||||
%conv_arr_4_S = OpFunction %_arr_S_uint_4 None %29
|
||||
%val_0 = OpFunctionParameter %_arr_S_std140_uint_4
|
||||
%32 = OpLabel
|
||||
%arr = OpVariable %_ptr_Function__arr_S_uint_4 Function %16
|
||||
%i = OpVariable %_ptr_Function_uint Function %37
|
||||
%var_for_index = OpVariable %_ptr_Function__arr_S_std140_uint_4 Function %50
|
||||
OpBranch %38
|
||||
%38 = OpLabel
|
||||
OpLoopMerge %39 %40 None
|
||||
OpBranch %41
|
||||
%41 = OpLabel
|
||||
%43 = OpLoad %uint %i
|
||||
%44 = OpULessThan %bool %43 %uint_4
|
||||
%42 = OpLogicalNot %bool %44
|
||||
OpSelectionMerge %46 None
|
||||
OpBranchConditional %42 %47 %46
|
||||
%47 = OpLabel
|
||||
OpBranch %39
|
||||
%46 = OpLabel
|
||||
OpStore %var_for_index %val_0
|
||||
%51 = OpLoad %uint %i
|
||||
%53 = OpAccessChain %_ptr_Function_S %arr %51
|
||||
%55 = OpLoad %uint %i
|
||||
%57 = OpAccessChain %_ptr_Function_S_std140 %var_for_index %55
|
||||
%58 = OpLoad %S_std140 %57
|
||||
%54 = OpFunctionCall %S %conv_S %58
|
||||
OpStore %53 %54
|
||||
OpBranch %40
|
||||
%40 = OpLabel
|
||||
%59 = OpLoad %uint %i
|
||||
%61 = OpIAdd %uint %59 %uint_1
|
||||
OpStore %i %61
|
||||
OpBranch %38
|
||||
%39 = OpLabel
|
||||
%62 = OpLoad %_arr_S_uint_4 %arr
|
||||
OpReturnValue %62
|
||||
OpFunctionEnd
|
||||
%load_u_2_m = OpFunction %mat4v2float None %63
|
||||
%65 = OpLabel
|
||||
%69 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %uint_2 %uint_1
|
||||
%70 = OpLoad %v2float %69
|
||||
%71 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %uint_2 %uint_2
|
||||
%72 = OpLoad %v2float %71
|
||||
%74 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %uint_2 %uint_3
|
||||
%75 = OpLoad %v2float %74
|
||||
%76 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %uint_2 %uint_4
|
||||
%77 = OpLoad %v2float %76
|
||||
%78 = OpCompositeConstruct %mat4v2float %70 %72 %75 %77
|
||||
OpReturnValue %78
|
||||
OpFunctionEnd
|
||||
%f = OpFunction %void None %79
|
||||
%82 = OpLabel
|
||||
%85 = OpAccessChain %_ptr_Uniform__arr_S_std140_uint_4 %u %uint_0
|
||||
%86 = OpLoad %_arr_S_std140_uint_4 %85
|
||||
%83 = OpFunctionCall %_arr_S_uint_4 %conv_arr_4_S %86
|
||||
OpStore %p %83
|
||||
%89 = OpAccessChain %_ptr_Private_S %p %int_1
|
||||
%92 = OpAccessChain %_ptr_Uniform_S_std140 %u %uint_0 %uint_2
|
||||
%93 = OpLoad %S_std140 %92
|
||||
%90 = OpFunctionCall %S %conv_S %93
|
||||
OpStore %89 %90
|
||||
%96 = OpAccessChain %_ptr_Private_mat4v2float %p %int_3 %uint_1
|
||||
%97 = OpFunctionCall %mat4v2float %load_u_2_m
|
||||
OpStore %96 %97
|
||||
%100 = OpAccessChain %_ptr_Private_v2float %p %int_1 %uint_1 %98
|
||||
%101 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %37 %uint_2
|
||||
%102 = OpLoad %v2float %101
|
||||
%103 = OpVectorShuffle %v2float %102 %102 1 0
|
||||
OpStore %100 %103
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@@ -0,0 +1,17 @@
|
||||
struct S {
|
||||
before : i32,
|
||||
m : mat4x2<f32>,
|
||||
after : i32,
|
||||
}
|
||||
|
||||
@group(0) @binding(0) var<uniform> u : array<S, 4>;
|
||||
|
||||
var<private> p : array<S, 4>;
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn f() {
|
||||
p = u;
|
||||
p[1] = u[2];
|
||||
p[3].m = u[2].m;
|
||||
p[1].m[0] = u[0].m[1].yx;
|
||||
}
|
||||
16
test/tint/buffer/uniform/std140/mat4x2/to_storage.wgsl
Normal file
16
test/tint/buffer/uniform/std140/mat4x2/to_storage.wgsl
Normal file
@@ -0,0 +1,16 @@
|
||||
struct S {
|
||||
before : i32,
|
||||
m : mat4x2<f32>,
|
||||
after : i32,
|
||||
}
|
||||
|
||||
@group(0) @binding(0) var<uniform> u : array<S, 4>;
|
||||
@group(0) @binding(1) var<storage, read_write> s : array<S, 4>;
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn f() {
|
||||
s = u;
|
||||
s[1] = u[2];
|
||||
s[3].m = u[2].m;
|
||||
s[1].m[0] = u[0].m[1].yx;
|
||||
}
|
||||
@@ -0,0 +1,71 @@
|
||||
struct S {
|
||||
int before;
|
||||
float4x2 m;
|
||||
int after;
|
||||
};
|
||||
|
||||
cbuffer cbuffer_u : register(b0, space0) {
|
||||
uint4 u[12];
|
||||
};
|
||||
RWByteAddressBuffer s : register(u1, space0);
|
||||
|
||||
void tint_symbol_3(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_1(RWByteAddressBuffer buffer, uint offset, S value) {
|
||||
buffer.Store((offset + 0u), asuint(value.before));
|
||||
tint_symbol_3(buffer, (offset + 8u), value.m);
|
||||
buffer.Store((offset + 40u), asuint(value.after));
|
||||
}
|
||||
|
||||
void tint_symbol(RWByteAddressBuffer buffer, uint offset, S value[4]) {
|
||||
S array[4] = value;
|
||||
{
|
||||
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
|
||||
tint_symbol_1(buffer, (offset + (i * 48u)), array[i]);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
float4x2 tint_symbol_8(uint4 buffer[12], 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)));
|
||||
}
|
||||
|
||||
S tint_symbol_6(uint4 buffer[12], uint offset) {
|
||||
const uint scalar_offset_4 = ((offset + 0u)) / 4;
|
||||
const uint scalar_offset_5 = ((offset + 40u)) / 4;
|
||||
const S tint_symbol_10 = {asint(buffer[scalar_offset_4 / 4][scalar_offset_4 % 4]), tint_symbol_8(buffer, (offset + 8u)), asint(buffer[scalar_offset_5 / 4][scalar_offset_5 % 4])};
|
||||
return tint_symbol_10;
|
||||
}
|
||||
|
||||
typedef S tint_symbol_5_ret[4];
|
||||
tint_symbol_5_ret tint_symbol_5(uint4 buffer[12], uint offset) {
|
||||
S arr[4] = (S[4])0;
|
||||
{
|
||||
[loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
|
||||
arr[i_1] = tint_symbol_6(buffer, (offset + (i_1 * 48u)));
|
||||
}
|
||||
}
|
||||
return arr;
|
||||
}
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void f() {
|
||||
tint_symbol(s, 0u, tint_symbol_5(u, 0u));
|
||||
tint_symbol_1(s, 48u, tint_symbol_6(u, 96u));
|
||||
tint_symbol_3(s, 152u, tint_symbol_8(u, 104u));
|
||||
s.Store2(56u, asuint(asfloat(u[1].xy).yx));
|
||||
return;
|
||||
}
|
||||
@@ -0,0 +1,71 @@
|
||||
struct S {
|
||||
int before;
|
||||
float4x2 m;
|
||||
int after;
|
||||
};
|
||||
|
||||
cbuffer cbuffer_u : register(b0, space0) {
|
||||
uint4 u[12];
|
||||
};
|
||||
RWByteAddressBuffer s : register(u1, space0);
|
||||
|
||||
void tint_symbol_3(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_1(RWByteAddressBuffer buffer, uint offset, S value) {
|
||||
buffer.Store((offset + 0u), asuint(value.before));
|
||||
tint_symbol_3(buffer, (offset + 8u), value.m);
|
||||
buffer.Store((offset + 40u), asuint(value.after));
|
||||
}
|
||||
|
||||
void tint_symbol(RWByteAddressBuffer buffer, uint offset, S value[4]) {
|
||||
S array[4] = value;
|
||||
{
|
||||
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
|
||||
tint_symbol_1(buffer, (offset + (i * 48u)), array[i]);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
float4x2 tint_symbol_8(uint4 buffer[12], 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)));
|
||||
}
|
||||
|
||||
S tint_symbol_6(uint4 buffer[12], uint offset) {
|
||||
const uint scalar_offset_4 = ((offset + 0u)) / 4;
|
||||
const uint scalar_offset_5 = ((offset + 40u)) / 4;
|
||||
const S tint_symbol_10 = {asint(buffer[scalar_offset_4 / 4][scalar_offset_4 % 4]), tint_symbol_8(buffer, (offset + 8u)), asint(buffer[scalar_offset_5 / 4][scalar_offset_5 % 4])};
|
||||
return tint_symbol_10;
|
||||
}
|
||||
|
||||
typedef S tint_symbol_5_ret[4];
|
||||
tint_symbol_5_ret tint_symbol_5(uint4 buffer[12], uint offset) {
|
||||
S arr[4] = (S[4])0;
|
||||
{
|
||||
[loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
|
||||
arr[i_1] = tint_symbol_6(buffer, (offset + (i_1 * 48u)));
|
||||
}
|
||||
}
|
||||
return arr;
|
||||
}
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void f() {
|
||||
tint_symbol(s, 0u, tint_symbol_5(u, 0u));
|
||||
tint_symbol_1(s, 48u, tint_symbol_6(u, 96u));
|
||||
tint_symbol_3(s, 152u, tint_symbol_8(u, 104u));
|
||||
s.Store2(56u, asuint(asfloat(u[1].xy).yx));
|
||||
return;
|
||||
}
|
||||
@@ -0,0 +1,63 @@
|
||||
#version 310 es
|
||||
|
||||
struct S {
|
||||
int before;
|
||||
mat4x2 m;
|
||||
int after;
|
||||
};
|
||||
|
||||
struct S_std140 {
|
||||
int before;
|
||||
vec2 m_0;
|
||||
vec2 m_1;
|
||||
vec2 m_2;
|
||||
vec2 m_3;
|
||||
int after;
|
||||
};
|
||||
|
||||
struct u_block {
|
||||
S_std140 inner[4];
|
||||
};
|
||||
|
||||
layout(binding = 0) uniform u_block_1 {
|
||||
S_std140 inner[4];
|
||||
} u;
|
||||
|
||||
struct s_block {
|
||||
S inner[4];
|
||||
};
|
||||
|
||||
layout(binding = 1, std430) buffer s_block_1 {
|
||||
S inner[4];
|
||||
} s;
|
||||
S conv_S(S_std140 val) {
|
||||
S tint_symbol = S(val.before, mat4x2(val.m_0, val.m_1, val.m_2, val.m_3), val.after);
|
||||
return tint_symbol;
|
||||
}
|
||||
|
||||
S[4] conv_arr_4_S(S_std140 val[4]) {
|
||||
S arr[4] = S[4](S(0, mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f), 0), S(0, mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f), 0), S(0, mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f), 0), S(0, mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 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;
|
||||
}
|
||||
|
||||
mat4x2 load_u_2_m() {
|
||||
return mat4x2(u.inner[2u].m_0, u.inner[2u].m_1, u.inner[2u].m_2, u.inner[2u].m_3);
|
||||
}
|
||||
|
||||
void f() {
|
||||
s.inner = conv_arr_4_S(u.inner);
|
||||
s.inner[1] = conv_S(u.inner[2u]);
|
||||
s.inner[3].m = load_u_2_m();
|
||||
s.inner[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;
|
||||
}
|
||||
@@ -0,0 +1,32 @@
|
||||
#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 {
|
||||
/* 0x0000 */ int before;
|
||||
/* 0x0004 */ tint_array<int8_t, 4> tint_pad;
|
||||
/* 0x0008 */ float4x2 m;
|
||||
/* 0x0028 */ int after;
|
||||
/* 0x002c */ tint_array<int8_t, 4> tint_pad_1;
|
||||
};
|
||||
|
||||
kernel void f(device tint_array<S, 4>* tint_symbol [[buffer(1)]], const constant tint_array<S, 4>* tint_symbol_1 [[buffer(0)]]) {
|
||||
*(tint_symbol) = *(tint_symbol_1);
|
||||
(*(tint_symbol))[1] = (*(tint_symbol_1))[2];
|
||||
(*(tint_symbol))[3].m = (*(tint_symbol_1))[2].m;
|
||||
(*(tint_symbol))[1].m[0] = float2((*(tint_symbol_1))[0].m[1]).yx;
|
||||
return;
|
||||
}
|
||||
|
||||
@@ -0,0 +1,186 @@
|
||||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 107
|
||||
; 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 %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 "m_2"
|
||||
OpMemberName %S_std140 4 "m_3"
|
||||
OpMemberName %S_std140 5 "after"
|
||||
OpName %u "u"
|
||||
OpName %s_block "s_block"
|
||||
OpMemberName %s_block 0 "inner"
|
||||
OpName %S "S"
|
||||
OpMemberName %S 0 "before"
|
||||
OpMemberName %S 1 "m"
|
||||
OpMemberName %S 2 "after"
|
||||
OpName %s "s"
|
||||
OpName %conv_S "conv_S"
|
||||
OpName %val "val"
|
||||
OpName %conv_arr_4_S "conv_arr_4_S"
|
||||
OpName %val_0 "val"
|
||||
OpName %arr "arr"
|
||||
OpName %i "i"
|
||||
OpName %var_for_index "var_for_index"
|
||||
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
|
||||
OpMemberDecorate %S_std140 4 Offset 32
|
||||
OpMemberDecorate %S_std140 5 Offset 40
|
||||
OpDecorate %_arr_S_std140_uint_4 ArrayStride 48
|
||||
OpDecorate %u NonWritable
|
||||
OpDecorate %u DescriptorSet 0
|
||||
OpDecorate %u Binding 0
|
||||
OpDecorate %s_block Block
|
||||
OpMemberDecorate %s_block 0 Offset 0
|
||||
OpMemberDecorate %S 0 Offset 0
|
||||
OpMemberDecorate %S 1 Offset 8
|
||||
OpMemberDecorate %S 1 ColMajor
|
||||
OpMemberDecorate %S 1 MatrixStride 8
|
||||
OpMemberDecorate %S 2 Offset 40
|
||||
OpDecorate %_arr_S_uint_4 ArrayStride 48
|
||||
OpDecorate %s DescriptorSet 0
|
||||
OpDecorate %s Binding 1
|
||||
%int = OpTypeInt 32 1
|
||||
%float = OpTypeFloat 32
|
||||
%v2float = OpTypeVector %float 2
|
||||
%S_std140 = OpTypeStruct %int %v2float %v2float %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
|
||||
%mat4v2float = OpTypeMatrix %v2float 4
|
||||
%S = OpTypeStruct %int %mat4v2float %int
|
||||
%_arr_S_uint_4 = OpTypeArray %S %uint_4
|
||||
%s_block = OpTypeStruct %_arr_S_uint_4
|
||||
%_ptr_StorageBuffer_s_block = OpTypePointer StorageBuffer %s_block
|
||||
%s = OpVariable %_ptr_StorageBuffer_s_block StorageBuffer
|
||||
%17 = OpTypeFunction %S %S_std140
|
||||
%29 = OpTypeFunction %_arr_S_uint_4 %_arr_S_std140_uint_4
|
||||
%_ptr_Function__arr_S_uint_4 = OpTypePointer Function %_arr_S_uint_4
|
||||
%35 = OpConstantNull %_arr_S_uint_4
|
||||
%_ptr_Function_uint = OpTypePointer Function %uint
|
||||
%38 = OpConstantNull %uint
|
||||
%bool = OpTypeBool
|
||||
%_ptr_Function__arr_S_std140_uint_4 = OpTypePointer Function %_arr_S_std140_uint_4
|
||||
%51 = OpConstantNull %_arr_S_std140_uint_4
|
||||
%_ptr_Function_S = OpTypePointer Function %S
|
||||
%_ptr_Function_S_std140 = OpTypePointer Function %S_std140
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%64 = OpTypeFunction %mat4v2float
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%uint_2 = OpConstant %uint 2
|
||||
%_ptr_Uniform_v2float = OpTypePointer Uniform %v2float
|
||||
%uint_3 = OpConstant %uint 3
|
||||
%void = OpTypeVoid
|
||||
%80 = OpTypeFunction %void
|
||||
%_ptr_StorageBuffer__arr_S_uint_4 = OpTypePointer StorageBuffer %_arr_S_uint_4
|
||||
%_ptr_Uniform__arr_S_std140_uint_4 = OpTypePointer Uniform %_arr_S_std140_uint_4
|
||||
%int_1 = OpConstant %int 1
|
||||
%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
|
||||
%_ptr_Uniform_S_std140 = OpTypePointer Uniform %S_std140
|
||||
%int_3 = OpConstant %int 3
|
||||
%_ptr_StorageBuffer_mat4v2float = OpTypePointer StorageBuffer %mat4v2float
|
||||
%101 = OpConstantNull %int
|
||||
%_ptr_StorageBuffer_v2float = OpTypePointer StorageBuffer %v2float
|
||||
%conv_S = OpFunction %S None %17
|
||||
%val = OpFunctionParameter %S_std140
|
||||
%20 = OpLabel
|
||||
%21 = OpCompositeExtract %int %val 0
|
||||
%22 = OpCompositeExtract %v2float %val 1
|
||||
%23 = OpCompositeExtract %v2float %val 2
|
||||
%24 = OpCompositeExtract %v2float %val 3
|
||||
%25 = OpCompositeExtract %v2float %val 4
|
||||
%26 = OpCompositeConstruct %mat4v2float %22 %23 %24 %25
|
||||
%27 = OpCompositeExtract %int %val 5
|
||||
%28 = OpCompositeConstruct %S %21 %26 %27
|
||||
OpReturnValue %28
|
||||
OpFunctionEnd
|
||||
%conv_arr_4_S = OpFunction %_arr_S_uint_4 None %29
|
||||
%val_0 = OpFunctionParameter %_arr_S_std140_uint_4
|
||||
%32 = OpLabel
|
||||
%arr = OpVariable %_ptr_Function__arr_S_uint_4 Function %35
|
||||
%i = OpVariable %_ptr_Function_uint Function %38
|
||||
%var_for_index = OpVariable %_ptr_Function__arr_S_std140_uint_4 Function %51
|
||||
OpBranch %39
|
||||
%39 = OpLabel
|
||||
OpLoopMerge %40 %41 None
|
||||
OpBranch %42
|
||||
%42 = OpLabel
|
||||
%44 = OpLoad %uint %i
|
||||
%45 = OpULessThan %bool %44 %uint_4
|
||||
%43 = OpLogicalNot %bool %45
|
||||
OpSelectionMerge %47 None
|
||||
OpBranchConditional %43 %48 %47
|
||||
%48 = OpLabel
|
||||
OpBranch %40
|
||||
%47 = OpLabel
|
||||
OpStore %var_for_index %val_0
|
||||
%52 = OpLoad %uint %i
|
||||
%54 = OpAccessChain %_ptr_Function_S %arr %52
|
||||
%56 = OpLoad %uint %i
|
||||
%58 = OpAccessChain %_ptr_Function_S_std140 %var_for_index %56
|
||||
%59 = OpLoad %S_std140 %58
|
||||
%55 = OpFunctionCall %S %conv_S %59
|
||||
OpStore %54 %55
|
||||
OpBranch %41
|
||||
%41 = OpLabel
|
||||
%60 = OpLoad %uint %i
|
||||
%62 = OpIAdd %uint %60 %uint_1
|
||||
OpStore %i %62
|
||||
OpBranch %39
|
||||
%40 = OpLabel
|
||||
%63 = OpLoad %_arr_S_uint_4 %arr
|
||||
OpReturnValue %63
|
||||
OpFunctionEnd
|
||||
%load_u_2_m = OpFunction %mat4v2float None %64
|
||||
%66 = OpLabel
|
||||
%70 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %uint_2 %uint_1
|
||||
%71 = OpLoad %v2float %70
|
||||
%72 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %uint_2 %uint_2
|
||||
%73 = OpLoad %v2float %72
|
||||
%75 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %uint_2 %uint_3
|
||||
%76 = OpLoad %v2float %75
|
||||
%77 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %uint_2 %uint_4
|
||||
%78 = OpLoad %v2float %77
|
||||
%79 = OpCompositeConstruct %mat4v2float %71 %73 %76 %78
|
||||
OpReturnValue %79
|
||||
OpFunctionEnd
|
||||
%f = OpFunction %void None %80
|
||||
%83 = OpLabel
|
||||
%85 = OpAccessChain %_ptr_StorageBuffer__arr_S_uint_4 %s %uint_0
|
||||
%88 = OpAccessChain %_ptr_Uniform__arr_S_std140_uint_4 %u %uint_0
|
||||
%89 = OpLoad %_arr_S_std140_uint_4 %88
|
||||
%86 = OpFunctionCall %_arr_S_uint_4 %conv_arr_4_S %89
|
||||
OpStore %85 %86
|
||||
%92 = OpAccessChain %_ptr_StorageBuffer_S %s %uint_0 %int_1
|
||||
%95 = OpAccessChain %_ptr_Uniform_S_std140 %u %uint_0 %uint_2
|
||||
%96 = OpLoad %S_std140 %95
|
||||
%93 = OpFunctionCall %S %conv_S %96
|
||||
OpStore %92 %93
|
||||
%99 = OpAccessChain %_ptr_StorageBuffer_mat4v2float %s %uint_0 %int_3 %uint_1
|
||||
%100 = OpFunctionCall %mat4v2float %load_u_2_m
|
||||
OpStore %99 %100
|
||||
%103 = OpAccessChain %_ptr_StorageBuffer_v2float %s %uint_0 %int_1 %uint_1 %101
|
||||
%104 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %38 %uint_2
|
||||
%105 = OpLoad %v2float %104
|
||||
%106 = OpVectorShuffle %v2float %105 %105 1 0
|
||||
OpStore %103 %106
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@@ -0,0 +1,17 @@
|
||||
struct S {
|
||||
before : i32,
|
||||
m : mat4x2<f32>,
|
||||
after : i32,
|
||||
}
|
||||
|
||||
@group(0) @binding(0) var<uniform> u : array<S, 4>;
|
||||
|
||||
@group(0) @binding(1) var<storage, read_write> s : array<S, 4>;
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn f() {
|
||||
s = u;
|
||||
s[1] = u[2];
|
||||
s[3].m = u[2].m;
|
||||
s[1].m[0] = u[0].m[1].yx;
|
||||
}
|
||||
16
test/tint/buffer/uniform/std140/mat4x2/to_workgroup.wgsl
Normal file
16
test/tint/buffer/uniform/std140/mat4x2/to_workgroup.wgsl
Normal file
@@ -0,0 +1,16 @@
|
||||
struct S {
|
||||
before : i32,
|
||||
m : mat4x2<f32>,
|
||||
after : i32,
|
||||
}
|
||||
|
||||
@group(0) @binding(0) var<uniform> u : array<S, 4>;
|
||||
var<workgroup> w : array<S, 4>;
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn f() {
|
||||
w = u;
|
||||
w[1] = u[2];
|
||||
w[3].m = u[2].m;
|
||||
w[1].m[0] = u[0].m[1].yx;
|
||||
}
|
||||
@@ -0,0 +1,65 @@
|
||||
struct S {
|
||||
int before;
|
||||
float4x2 m;
|
||||
int after;
|
||||
};
|
||||
|
||||
cbuffer cbuffer_u : register(b0, space0) {
|
||||
uint4 u[12];
|
||||
};
|
||||
groupshared S w[4];
|
||||
|
||||
struct tint_symbol_1 {
|
||||
uint local_invocation_index : SV_GroupIndex;
|
||||
};
|
||||
|
||||
float4x2 tint_symbol_5(uint4 buffer[12], 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)));
|
||||
}
|
||||
|
||||
S tint_symbol_3(uint4 buffer[12], uint offset) {
|
||||
const uint scalar_offset_4 = ((offset + 0u)) / 4;
|
||||
const uint scalar_offset_5 = ((offset + 40u)) / 4;
|
||||
const S tint_symbol_8 = {asint(buffer[scalar_offset_4 / 4][scalar_offset_4 % 4]), tint_symbol_5(buffer, (offset + 8u)), asint(buffer[scalar_offset_5 / 4][scalar_offset_5 % 4])};
|
||||
return tint_symbol_8;
|
||||
}
|
||||
|
||||
typedef S tint_symbol_2_ret[4];
|
||||
tint_symbol_2_ret tint_symbol_2(uint4 buffer[12], uint offset) {
|
||||
S arr[4] = (S[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 * 48u)));
|
||||
}
|
||||
}
|
||||
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;
|
||||
const S tint_symbol_7 = (S)0;
|
||||
w[i] = tint_symbol_7;
|
||||
}
|
||||
}
|
||||
GroupMemoryBarrierWithGroupSync();
|
||||
w = tint_symbol_2(u, 0u);
|
||||
w[1] = tint_symbol_3(u, 96u);
|
||||
w[3].m = tint_symbol_5(u, 104u);
|
||||
w[1].m[0] = asfloat(u[1].xy).yx;
|
||||
}
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void f(tint_symbol_1 tint_symbol) {
|
||||
f_inner(tint_symbol.local_invocation_index);
|
||||
return;
|
||||
}
|
||||
@@ -0,0 +1,65 @@
|
||||
struct S {
|
||||
int before;
|
||||
float4x2 m;
|
||||
int after;
|
||||
};
|
||||
|
||||
cbuffer cbuffer_u : register(b0, space0) {
|
||||
uint4 u[12];
|
||||
};
|
||||
groupshared S w[4];
|
||||
|
||||
struct tint_symbol_1 {
|
||||
uint local_invocation_index : SV_GroupIndex;
|
||||
};
|
||||
|
||||
float4x2 tint_symbol_5(uint4 buffer[12], 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)));
|
||||
}
|
||||
|
||||
S tint_symbol_3(uint4 buffer[12], uint offset) {
|
||||
const uint scalar_offset_4 = ((offset + 0u)) / 4;
|
||||
const uint scalar_offset_5 = ((offset + 40u)) / 4;
|
||||
const S tint_symbol_8 = {asint(buffer[scalar_offset_4 / 4][scalar_offset_4 % 4]), tint_symbol_5(buffer, (offset + 8u)), asint(buffer[scalar_offset_5 / 4][scalar_offset_5 % 4])};
|
||||
return tint_symbol_8;
|
||||
}
|
||||
|
||||
typedef S tint_symbol_2_ret[4];
|
||||
tint_symbol_2_ret tint_symbol_2(uint4 buffer[12], uint offset) {
|
||||
S arr[4] = (S[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 * 48u)));
|
||||
}
|
||||
}
|
||||
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;
|
||||
const S tint_symbol_7 = (S)0;
|
||||
w[i] = tint_symbol_7;
|
||||
}
|
||||
}
|
||||
GroupMemoryBarrierWithGroupSync();
|
||||
w = tint_symbol_2(u, 0u);
|
||||
w[1] = tint_symbol_3(u, 96u);
|
||||
w[3].m = tint_symbol_5(u, 104u);
|
||||
w[1].m[0] = asfloat(u[1].xy).yx;
|
||||
}
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void f(tint_symbol_1 tint_symbol) {
|
||||
f_inner(tint_symbol.local_invocation_index);
|
||||
return;
|
||||
}
|
||||
@@ -0,0 +1,65 @@
|
||||
#version 310 es
|
||||
|
||||
struct S {
|
||||
int before;
|
||||
mat4x2 m;
|
||||
int after;
|
||||
};
|
||||
|
||||
struct S_std140 {
|
||||
int before;
|
||||
vec2 m_0;
|
||||
vec2 m_1;
|
||||
vec2 m_2;
|
||||
vec2 m_3;
|
||||
int after;
|
||||
};
|
||||
|
||||
struct u_block {
|
||||
S_std140 inner[4];
|
||||
};
|
||||
|
||||
layout(binding = 0) uniform u_block_1 {
|
||||
S_std140 inner[4];
|
||||
} u;
|
||||
|
||||
shared S w[4];
|
||||
S conv_S(S_std140 val) {
|
||||
S tint_symbol = S(val.before, mat4x2(val.m_0, val.m_1, val.m_2, val.m_3), val.after);
|
||||
return tint_symbol;
|
||||
}
|
||||
|
||||
S[4] conv_arr_4_S(S_std140 val[4]) {
|
||||
S arr[4] = S[4](S(0, mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f), 0), S(0, mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f), 0), S(0, mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f), 0), S(0, mat4x2(0.0f, 0.0f, 0.0f, 0.0f, 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;
|
||||
}
|
||||
|
||||
mat4x2 load_u_2_m() {
|
||||
return mat4x2(u.inner[2u].m_0, u.inner[2u].m_1, u.inner[2u].m_2, u.inner[2u].m_3);
|
||||
}
|
||||
|
||||
void f(uint local_invocation_index) {
|
||||
{
|
||||
for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
|
||||
uint i = idx;
|
||||
S tint_symbol_1 = S(0, mat4x2(vec2(0.0f), vec2(0.0f), vec2(0.0f), vec2(0.0f)), 0);
|
||||
w[i] = tint_symbol_1;
|
||||
}
|
||||
}
|
||||
barrier();
|
||||
w = conv_arr_4_S(u.inner);
|
||||
w[1] = conv_S(u.inner[2u]);
|
||||
w[3].m = load_u_2_m();
|
||||
w[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(gl_LocalInvocationIndex);
|
||||
return;
|
||||
}
|
||||
@@ -0,0 +1,47 @@
|
||||
#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 {
|
||||
/* 0x0000 */ int before;
|
||||
/* 0x0004 */ tint_array<int8_t, 4> tint_pad;
|
||||
/* 0x0008 */ float4x2 m;
|
||||
/* 0x0028 */ int after;
|
||||
/* 0x002c */ tint_array<int8_t, 4> tint_pad_1;
|
||||
};
|
||||
|
||||
struct tint_symbol_6 {
|
||||
tint_array<S, 4> w;
|
||||
};
|
||||
|
||||
void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S, 4>* const tint_symbol_2) {
|
||||
for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
|
||||
uint const i = idx;
|
||||
S const tint_symbol = S{};
|
||||
(*(tint_symbol_1))[i] = tint_symbol;
|
||||
}
|
||||
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||
*(tint_symbol_1) = *(tint_symbol_2);
|
||||
(*(tint_symbol_1))[1] = (*(tint_symbol_2))[2];
|
||||
(*(tint_symbol_1))[3].m = (*(tint_symbol_2))[2].m;
|
||||
(*(tint_symbol_1))[1].m[0] = float2((*(tint_symbol_2))[0].m[1]).yx;
|
||||
}
|
||||
|
||||
kernel void f(const constant tint_array<S, 4>* tint_symbol_5 [[buffer(0)]], threadgroup tint_symbol_6* tint_symbol_4 [[threadgroup(0)]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
|
||||
threadgroup tint_array<S, 4>* const tint_symbol_3 = &((*(tint_symbol_4)).w);
|
||||
f_inner(local_invocation_index, tint_symbol_3, tint_symbol_5);
|
||||
return;
|
||||
}
|
||||
|
||||
@@ -0,0 +1,220 @@
|
||||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 129
|
||||
; 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 %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 "m_2"
|
||||
OpMemberName %S_std140 4 "m_3"
|
||||
OpMemberName %S_std140 5 "after"
|
||||
OpName %u "u"
|
||||
OpName %S "S"
|
||||
OpMemberName %S 0 "before"
|
||||
OpMemberName %S 1 "m"
|
||||
OpMemberName %S 2 "after"
|
||||
OpName %w "w"
|
||||
OpName %conv_S "conv_S"
|
||||
OpName %val "val"
|
||||
OpName %conv_arr_4_S "conv_arr_4_S"
|
||||
OpName %val_0 "val"
|
||||
OpName %arr "arr"
|
||||
OpName %i "i"
|
||||
OpName %var_for_index "var_for_index"
|
||||
OpName %load_u_2_m "load_u_2_m"
|
||||
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 %S_std140 0 Offset 0
|
||||
OpMemberDecorate %S_std140 1 Offset 8
|
||||
OpMemberDecorate %S_std140 2 Offset 16
|
||||
OpMemberDecorate %S_std140 3 Offset 24
|
||||
OpMemberDecorate %S_std140 4 Offset 32
|
||||
OpMemberDecorate %S_std140 5 Offset 40
|
||||
OpDecorate %_arr_S_std140_uint_4 ArrayStride 48
|
||||
OpDecorate %u NonWritable
|
||||
OpDecorate %u DescriptorSet 0
|
||||
OpDecorate %u Binding 0
|
||||
OpMemberDecorate %S 0 Offset 0
|
||||
OpMemberDecorate %S 1 Offset 8
|
||||
OpMemberDecorate %S 1 ColMajor
|
||||
OpMemberDecorate %S 1 MatrixStride 8
|
||||
OpMemberDecorate %S 2 Offset 40
|
||||
OpDecorate %_arr_S_uint_4 ArrayStride 48
|
||||
%uint = OpTypeInt 32 0
|
||||
%_ptr_Input_uint = OpTypePointer Input %uint
|
||||
%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
|
||||
%int = OpTypeInt 32 1
|
||||
%float = OpTypeFloat 32
|
||||
%v2float = OpTypeVector %float 2
|
||||
%S_std140 = OpTypeStruct %int %v2float %v2float %v2float %v2float %int
|
||||
%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
|
||||
%mat4v2float = OpTypeMatrix %v2float 4
|
||||
%S = OpTypeStruct %int %mat4v2float %int
|
||||
%_arr_S_uint_4 = OpTypeArray %S %uint_4
|
||||
%_ptr_Workgroup__arr_S_uint_4 = OpTypePointer Workgroup %_arr_S_uint_4
|
||||
%w = OpVariable %_ptr_Workgroup__arr_S_uint_4 Workgroup
|
||||
%18 = OpTypeFunction %S %S_std140
|
||||
%30 = OpTypeFunction %_arr_S_uint_4 %_arr_S_std140_uint_4
|
||||
%_ptr_Function__arr_S_uint_4 = OpTypePointer Function %_arr_S_uint_4
|
||||
%36 = OpConstantNull %_arr_S_uint_4
|
||||
%_ptr_Function_uint = OpTypePointer Function %uint
|
||||
%39 = OpConstantNull %uint
|
||||
%bool = OpTypeBool
|
||||
%_ptr_Function__arr_S_std140_uint_4 = OpTypePointer Function %_arr_S_std140_uint_4
|
||||
%52 = OpConstantNull %_arr_S_std140_uint_4
|
||||
%_ptr_Function_S = OpTypePointer Function %S
|
||||
%_ptr_Function_S_std140 = OpTypePointer Function %S_std140
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%65 = OpTypeFunction %mat4v2float
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%uint_2 = OpConstant %uint 2
|
||||
%_ptr_Uniform_v2float = OpTypePointer Uniform %v2float
|
||||
%uint_3 = OpConstant %uint 3
|
||||
%void = OpTypeVoid
|
||||
%81 = OpTypeFunction %void %uint
|
||||
%_ptr_Workgroup_S = OpTypePointer Workgroup %S
|
||||
%99 = OpConstantNull %S
|
||||
%uint_264 = OpConstant %uint 264
|
||||
%_ptr_Uniform__arr_S_std140_uint_4 = OpTypePointer Uniform %_arr_S_std140_uint_4
|
||||
%int_1 = OpConstant %int 1
|
||||
%_ptr_Uniform_S_std140 = OpTypePointer Uniform %S_std140
|
||||
%int_3 = OpConstant %int 3
|
||||
%_ptr_Workgroup_mat4v2float = OpTypePointer Workgroup %mat4v2float
|
||||
%118 = OpConstantNull %int
|
||||
%_ptr_Workgroup_v2float = OpTypePointer Workgroup %v2float
|
||||
%124 = OpTypeFunction %void
|
||||
%conv_S = OpFunction %S None %18
|
||||
%val = OpFunctionParameter %S_std140
|
||||
%21 = OpLabel
|
||||
%22 = OpCompositeExtract %int %val 0
|
||||
%23 = OpCompositeExtract %v2float %val 1
|
||||
%24 = OpCompositeExtract %v2float %val 2
|
||||
%25 = OpCompositeExtract %v2float %val 3
|
||||
%26 = OpCompositeExtract %v2float %val 4
|
||||
%27 = OpCompositeConstruct %mat4v2float %23 %24 %25 %26
|
||||
%28 = OpCompositeExtract %int %val 5
|
||||
%29 = OpCompositeConstruct %S %22 %27 %28
|
||||
OpReturnValue %29
|
||||
OpFunctionEnd
|
||||
%conv_arr_4_S = OpFunction %_arr_S_uint_4 None %30
|
||||
%val_0 = OpFunctionParameter %_arr_S_std140_uint_4
|
||||
%33 = OpLabel
|
||||
%arr = OpVariable %_ptr_Function__arr_S_uint_4 Function %36
|
||||
%i = OpVariable %_ptr_Function_uint Function %39
|
||||
%var_for_index = OpVariable %_ptr_Function__arr_S_std140_uint_4 Function %52
|
||||
OpBranch %40
|
||||
%40 = OpLabel
|
||||
OpLoopMerge %41 %42 None
|
||||
OpBranch %43
|
||||
%43 = OpLabel
|
||||
%45 = OpLoad %uint %i
|
||||
%46 = OpULessThan %bool %45 %uint_4
|
||||
%44 = OpLogicalNot %bool %46
|
||||
OpSelectionMerge %48 None
|
||||
OpBranchConditional %44 %49 %48
|
||||
%49 = OpLabel
|
||||
OpBranch %41
|
||||
%48 = OpLabel
|
||||
OpStore %var_for_index %val_0
|
||||
%53 = OpLoad %uint %i
|
||||
%55 = OpAccessChain %_ptr_Function_S %arr %53
|
||||
%57 = OpLoad %uint %i
|
||||
%59 = OpAccessChain %_ptr_Function_S_std140 %var_for_index %57
|
||||
%60 = OpLoad %S_std140 %59
|
||||
%56 = OpFunctionCall %S %conv_S %60
|
||||
OpStore %55 %56
|
||||
OpBranch %42
|
||||
%42 = OpLabel
|
||||
%61 = OpLoad %uint %i
|
||||
%63 = OpIAdd %uint %61 %uint_1
|
||||
OpStore %i %63
|
||||
OpBranch %40
|
||||
%41 = OpLabel
|
||||
%64 = OpLoad %_arr_S_uint_4 %arr
|
||||
OpReturnValue %64
|
||||
OpFunctionEnd
|
||||
%load_u_2_m = OpFunction %mat4v2float None %65
|
||||
%67 = OpLabel
|
||||
%71 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %uint_2 %uint_1
|
||||
%72 = OpLoad %v2float %71
|
||||
%73 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %uint_2 %uint_2
|
||||
%74 = OpLoad %v2float %73
|
||||
%76 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %uint_2 %uint_3
|
||||
%77 = OpLoad %v2float %76
|
||||
%78 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %uint_2 %uint_4
|
||||
%79 = OpLoad %v2float %78
|
||||
%80 = OpCompositeConstruct %mat4v2float %72 %74 %77 %79
|
||||
OpReturnValue %80
|
||||
OpFunctionEnd
|
||||
%f_inner = OpFunction %void None %81
|
||||
%local_invocation_index = OpFunctionParameter %uint
|
||||
%85 = OpLabel
|
||||
%idx = OpVariable %_ptr_Function_uint Function %39
|
||||
OpStore %idx %local_invocation_index
|
||||
OpBranch %87
|
||||
%87 = OpLabel
|
||||
OpLoopMerge %88 %89 None
|
||||
OpBranch %90
|
||||
%90 = OpLabel
|
||||
%92 = OpLoad %uint %idx
|
||||
%93 = OpULessThan %bool %92 %uint_4
|
||||
%91 = OpLogicalNot %bool %93
|
||||
OpSelectionMerge %94 None
|
||||
OpBranchConditional %91 %95 %94
|
||||
%95 = OpLabel
|
||||
OpBranch %88
|
||||
%94 = OpLabel
|
||||
%96 = OpLoad %uint %idx
|
||||
%98 = OpAccessChain %_ptr_Workgroup_S %w %96
|
||||
OpStore %98 %99
|
||||
OpBranch %89
|
||||
%89 = OpLabel
|
||||
%100 = OpLoad %uint %idx
|
||||
%101 = OpIAdd %uint %100 %uint_1
|
||||
OpStore %idx %101
|
||||
OpBranch %87
|
||||
%88 = OpLabel
|
||||
OpControlBarrier %uint_2 %uint_2 %uint_264
|
||||
%106 = OpAccessChain %_ptr_Uniform__arr_S_std140_uint_4 %u %uint_0
|
||||
%107 = OpLoad %_arr_S_std140_uint_4 %106
|
||||
%104 = OpFunctionCall %_arr_S_uint_4 %conv_arr_4_S %107
|
||||
OpStore %w %104
|
||||
%109 = OpAccessChain %_ptr_Workgroup_S %w %int_1
|
||||
%112 = OpAccessChain %_ptr_Uniform_S_std140 %u %uint_0 %uint_2
|
||||
%113 = OpLoad %S_std140 %112
|
||||
%110 = OpFunctionCall %S %conv_S %113
|
||||
OpStore %109 %110
|
||||
%116 = OpAccessChain %_ptr_Workgroup_mat4v2float %w %int_3 %uint_1
|
||||
%117 = OpFunctionCall %mat4v2float %load_u_2_m
|
||||
OpStore %116 %117
|
||||
%120 = OpAccessChain %_ptr_Workgroup_v2float %w %int_1 %uint_1 %118
|
||||
%121 = OpAccessChain %_ptr_Uniform_v2float %u %uint_0 %39 %uint_2
|
||||
%122 = OpLoad %v2float %121
|
||||
%123 = OpVectorShuffle %v2float %122 %122 1 0
|
||||
OpStore %120 %123
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%f = OpFunction %void None %124
|
||||
%126 = OpLabel
|
||||
%128 = OpLoad %uint %local_invocation_index_1
|
||||
%127 = OpFunctionCall %void %f_inner %128
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@@ -0,0 +1,17 @@
|
||||
struct S {
|
||||
before : i32,
|
||||
m : mat4x2<f32>,
|
||||
after : i32,
|
||||
}
|
||||
|
||||
@group(0) @binding(0) var<uniform> u : array<S, 4>;
|
||||
|
||||
var<workgroup> w : array<S, 4>;
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn f() {
|
||||
w = u;
|
||||
w[1] = u[2];
|
||||
w[3].m = u[2].m;
|
||||
w[1].m[0] = u[0].m[1].yx;
|
||||
}
|
||||
Reference in New Issue
Block a user