Natalie Chouinard | e7de00b | 2024-08-06 20:54:16 +0000 | [diff] [blame] | 1 | #include <metal_stdlib> |
| 2 | using namespace metal; |
| 3 | |
| 4 | struct tint_module_vars_struct { |
James Price | 61c2d1e | 2024-08-27 21:26:49 +0000 | [diff] [blame] | 5 | device packed_half3* prevent_dce; |
Natalie Chouinard | e7de00b | 2024-08-06 20:54:16 +0000 | [diff] [blame] | 6 | }; |
| 7 | |
| 8 | half3 subgroupExclusiveAdd_e58e23() { |
| 9 | half3 arg_0 = half3(1.0h); |
| 10 | half3 res = simd_prefix_exclusive_sum(arg_0); |
| 11 | return res; |
| 12 | } |
| 13 | |
James Price | 61c2d1e | 2024-08-27 21:26:49 +0000 | [diff] [blame] | 14 | fragment void fragment_main(device packed_half3* prevent_dce [[buffer(0)]]) { |
Natalie Chouinard | e9e450e | 2024-08-08 23:16:40 +0000 | [diff] [blame] | 15 | tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce}; |
James Price | 61c2d1e | 2024-08-27 21:26:49 +0000 | [diff] [blame] | 16 | (*tint_module_vars.prevent_dce) = packed_half3(subgroupExclusiveAdd_e58e23()); |
Natalie Chouinard | e9e450e | 2024-08-08 23:16:40 +0000 | [diff] [blame] | 17 | } |
| 18 | |
James Price | 61c2d1e | 2024-08-27 21:26:49 +0000 | [diff] [blame] | 19 | kernel void compute_main(device packed_half3* prevent_dce [[buffer(0)]]) { |
Natalie Chouinard | e7de00b | 2024-08-06 20:54:16 +0000 | [diff] [blame] | 20 | tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce}; |
James Price | 61c2d1e | 2024-08-27 21:26:49 +0000 | [diff] [blame] | 21 | (*tint_module_vars.prevent_dce) = packed_half3(subgroupExclusiveAdd_e58e23()); |
Natalie Chouinard | e7de00b | 2024-08-06 20:54:16 +0000 | [diff] [blame] | 22 | } |