blob: 051427d9d67a44ba7f254301b21e6b079aff91b7 [file] [log] [blame]
Natalie Chouinarde27b1662024-08-16 16:24:35 +00001#include <metal_stdlib>
2
3using namespace metal;
4struct tint_private_vars_struct {
5 uint tint_msl_thread_index_in_quadgroup;
6};
Natalie Chouinarde31fa9d2024-08-14 20:52:07 +00007
8
Natalie Chouinarde27b1662024-08-16 16:24:35 +00009int4 tint_msl_quadSwapDiagonal(int4 e, thread tint_private_vars_struct* const tint_private_vars) {
10 return quad_shuffle(e,((*(tint_private_vars)).tint_msl_thread_index_in_quadgroup ^ 3u));
11}
Natalie Chouinarde31fa9d2024-08-14 20:52:07 +000012
Natalie Chouinarde27b1662024-08-16 16:24:35 +000013int4 quadSwapDiagonal_a665b1(thread tint_private_vars_struct* const tint_private_vars) {
14 int4 res = tint_msl_quadSwapDiagonal(int4(1), tint_private_vars);
Natalie Chouinarde31fa9d2024-08-14 20:52:07 +000015 return res;
16}
17
Natalie Chouinarde27b1662024-08-16 16:24:35 +000018fragment void fragment_main(device int4* tint_symbol [[buffer(0)]], uint tint_thread_index_in_quadgroup [[thread_index_in_quadgroup]]) {
19 thread tint_private_vars_struct tint_private_vars = {};
20 {
21 tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup;
22 }
23 *(tint_symbol) = quadSwapDiagonal_a665b1(&(tint_private_vars));
24 return;
Natalie Chouinarde31fa9d2024-08-14 20:52:07 +000025}
26
Natalie Chouinarde27b1662024-08-16 16:24:35 +000027kernel void compute_main(device int4* tint_symbol_1 [[buffer(0)]], uint tint_thread_index_in_quadgroup_1 [[thread_index_in_quadgroup]]) {
28 thread tint_private_vars_struct tint_private_vars = {};
29 {
30 tint_private_vars.tint_msl_thread_index_in_quadgroup = tint_thread_index_in_quadgroup_1;
31 }
32 *(tint_symbol_1) = quadSwapDiagonal_a665b1(&(tint_private_vars));
33 return;
Natalie Chouinarde31fa9d2024-08-14 20:52:07 +000034}
35