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