1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
|
#pragma clang diagnostic ignored "-Wmissing-prototypes"
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct SSBOCol
{
float2x4 col_major0;
float2x4 col_major1;
};
struct SSBORow
{
float2x4 row_major0;
float2x4 row_major1;
};
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
static inline __attribute__((always_inline))
void load_store_to_variable_col_major(device SSBOCol& v_29)
{
float2x2 loaded = float2x2(v_29.col_major0[0].xy, v_29.col_major0[1].xy);
v_29.col_major1[0].xy = loaded[0];
v_29.col_major1[1].xy = loaded[1];
}
static inline __attribute__((always_inline))
void load_store_to_variable_row_major(device SSBORow& v_41)
{
float2x2 loaded = transpose(float2x2(v_41.row_major0[0].xy, v_41.row_major0[1].xy));
v_41.row_major0[0].xy = float2(loaded[0][0], loaded[1][0]);
v_41.row_major0[1].xy = float2(loaded[0][1], loaded[1][1]);
}
static inline __attribute__((always_inline))
void copy_col_major_to_col_major(device SSBOCol& v_29)
{
v_29.col_major0[0].xy = float2x2(v_29.col_major1[0].xy, v_29.col_major1[1].xy)[0];
v_29.col_major0[1].xy = float2x2(v_29.col_major1[0].xy, v_29.col_major1[1].xy)[1];
}
static inline __attribute__((always_inline))
void copy_col_major_to_row_major(device SSBOCol& v_29, device SSBORow& v_41)
{
v_41.row_major0[0].xy = float2(float2x2(v_29.col_major0[0].xy, v_29.col_major0[1].xy)[0][0], float2x2(v_29.col_major0[0].xy, v_29.col_major0[1].xy)[1][0]);
v_41.row_major0[1].xy = float2(float2x2(v_29.col_major0[0].xy, v_29.col_major0[1].xy)[0][1], float2x2(v_29.col_major0[0].xy, v_29.col_major0[1].xy)[1][1]);
}
static inline __attribute__((always_inline))
void copy_row_major_to_col_major(device SSBOCol& v_29, device SSBORow& v_41)
{
v_29.col_major0[0].xy = float2(float2x2(v_41.row_major0[0].xy, v_41.row_major0[1].xy)[0][0], float2x2(v_41.row_major0[0].xy, v_41.row_major0[1].xy)[1][0]);
v_29.col_major0[1].xy = float2(float2x2(v_41.row_major0[0].xy, v_41.row_major0[1].xy)[0][1], float2x2(v_41.row_major0[0].xy, v_41.row_major0[1].xy)[1][1]);
}
static inline __attribute__((always_inline))
void copy_row_major_to_row_major(device SSBORow& v_41)
{
v_41.row_major0[0].xy = float2x2(v_41.row_major1[0].xy, v_41.row_major1[1].xy)[0];
v_41.row_major0[1].xy = float2x2(v_41.row_major1[0].xy, v_41.row_major1[1].xy)[1];
}
static inline __attribute__((always_inline))
void copy_columns(device SSBOCol& v_29, device SSBORow& v_41)
{
v_29.col_major0[1].xy = float2(v_41.row_major0[0][1], v_41.row_major0[1][1]);
v_41.row_major0[0][1] = v_29.col_major0[1].x;
v_41.row_major0[1][1] = v_29.col_major0[1].y;
}
static inline __attribute__((always_inline))
void copy_elements(device SSBOCol& v_29, device SSBORow& v_41)
{
v_29.col_major0[0].y = v_41.row_major0[1u][0];
v_41.row_major0[1u][0] = v_29.col_major0[0].y;
}
kernel void main0(device SSBOCol& v_29 [[buffer(0)]], device SSBORow& v_41 [[buffer(1)]])
{
load_store_to_variable_col_major(v_29);
load_store_to_variable_row_major(v_41);
copy_col_major_to_col_major(v_29);
copy_col_major_to_row_major(v_29, v_41);
copy_row_major_to_col_major(v_29, v_41);
copy_row_major_to_row_major(v_41);
copy_columns(v_29, v_41);
copy_elements(v_29, v_41);
}
|