• 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    float3x4 col_major0;
11    float3x4 col_major1;
12};
13
14struct SSBORow
15{
16    float2x3 row_major0;
17    float2x3 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    float3x2 loaded = float3x2(v_29.col_major0[0].xy, v_29.col_major0[1].xy, v_29.col_major0[2].xy);
26    (device float2&)v_29.col_major1[0] = loaded[0];
27    (device float2&)v_29.col_major1[1] = loaded[1];
28    (device float2&)v_29.col_major1[2] = loaded[2];
29}
30
31static inline __attribute__((always_inline))
32void load_store_to_variable_row_major(device SSBORow& v_41)
33{
34    float3x2 loaded = transpose(v_41.row_major0);
35    v_41.row_major0 = transpose(loaded);
36}
37
38static inline __attribute__((always_inline))
39void copy_col_major_to_col_major(device SSBOCol& v_29)
40{
41    (device float2&)v_29.col_major0[0] = float3x2(v_29.col_major1[0].xy, v_29.col_major1[1].xy, v_29.col_major1[2].xy)[0];
42    (device float2&)v_29.col_major0[1] = float3x2(v_29.col_major1[0].xy, v_29.col_major1[1].xy, v_29.col_major1[2].xy)[1];
43    (device float2&)v_29.col_major0[2] = float3x2(v_29.col_major1[0].xy, v_29.col_major1[1].xy, v_29.col_major1[2].xy)[2];
44}
45
46static inline __attribute__((always_inline))
47void copy_col_major_to_row_major(device SSBOCol& v_29, device SSBORow& v_41)
48{
49    v_41.row_major0 = transpose(float3x2(v_29.col_major0[0].xy, v_29.col_major0[1].xy, v_29.col_major0[2].xy));
50}
51
52static inline __attribute__((always_inline))
53void copy_row_major_to_col_major(device SSBOCol& v_29, device SSBORow& v_41)
54{
55    (device float2&)v_29.col_major0[0] = float2(v_41.row_major0[0][0], v_41.row_major0[1][0]);
56    (device float2&)v_29.col_major0[1] = float2(v_41.row_major0[0][1], v_41.row_major0[1][1]);
57    (device float2&)v_29.col_major0[2] = float2(v_41.row_major0[0][2], v_41.row_major0[1][2]);
58}
59
60static inline __attribute__((always_inline))
61void copy_row_major_to_row_major(device SSBORow& v_41)
62{
63    v_41.row_major0 = v_41.row_major1;
64}
65
66static inline __attribute__((always_inline))
67void copy_columns(device SSBOCol& v_29, device SSBORow& v_41)
68{
69    (device float2&)v_29.col_major0[1] = float2(v_41.row_major0[0][1], v_41.row_major0[1][1]);
70    v_41.row_major0[0][1] = v_29.col_major0[1].x;
71    v_41.row_major0[1][1] = v_29.col_major0[1].y;
72}
73
74static inline __attribute__((always_inline))
75void copy_elements(device SSBOCol& v_29, device SSBORow& v_41)
76{
77    ((device float*)&v_29.col_major0[0])[1u] = ((device float*)&v_41.row_major0[1u])[0];
78    ((device float*)&v_41.row_major0[1u])[0] = ((device float*)&v_29.col_major0[0])[1u];
79}
80
81kernel void main0(device SSBOCol& v_29 [[buffer(0)]], device SSBORow& v_41 [[buffer(1)]])
82{
83    load_store_to_variable_col_major(v_29);
84    load_store_to_variable_row_major(v_41);
85    copy_col_major_to_col_major(v_29);
86    copy_col_major_to_row_major(v_29, v_41);
87    copy_row_major_to_col_major(v_29, v_41);
88    copy_row_major_to_row_major(v_41);
89    copy_columns(v_29, v_41);
90    copy_elements(v_29, v_41);
91}
92
93