1*c8dee2aaSAndroid Build Coastguard Worker#include <metal_stdlib> 2*c8dee2aaSAndroid Build Coastguard Worker#include <simd/simd.h> 3*c8dee2aaSAndroid Build Coastguard Worker#ifdef __clang__ 4*c8dee2aaSAndroid Build Coastguard Worker#pragma clang diagnostic ignored "-Wall" 5*c8dee2aaSAndroid Build Coastguard Worker#endif 6*c8dee2aaSAndroid Build Coastguard Workerusing namespace metal; 7*c8dee2aaSAndroid Build Coastguard Workerstruct Inputs { 8*c8dee2aaSAndroid Build Coastguard Worker uint3 sk_GlobalInvocationID; 9*c8dee2aaSAndroid Build Coastguard Worker}; 10*c8dee2aaSAndroid Build Coastguard Workerstruct Globals { 11*c8dee2aaSAndroid Build Coastguard Worker texture2d<half, access::read> texIn; 12*c8dee2aaSAndroid Build Coastguard Worker texture2d<half, access::write> texOut; 13*c8dee2aaSAndroid Build Coastguard Worker}; 14*c8dee2aaSAndroid Build Coastguard Workerkernel void computeMain(uint3 sk_GlobalInvocationID [[thread_position_in_grid]], texture2d<half, access::read> texIn [[texture(0)]], texture2d<half, access::write> texOut [[texture(1)]]) { 15*c8dee2aaSAndroid Build Coastguard Worker Globals _globals{texIn, texOut}; 16*c8dee2aaSAndroid Build Coastguard Worker (void)_globals; 17*c8dee2aaSAndroid Build Coastguard Worker Inputs _in = { sk_GlobalInvocationID }; 18*c8dee2aaSAndroid Build Coastguard Worker if (_in.sk_GlobalInvocationID.x < _globals.texIn.get_width() && _in.sk_GlobalInvocationID.y < _globals.texIn.get_height()) { 19*c8dee2aaSAndroid Build Coastguard Worker half4 _0_color = _globals.texIn.read(_in.sk_GlobalInvocationID.xy); 20*c8dee2aaSAndroid Build Coastguard Worker _0_color.xyz = half3(dot(_0_color.xyz, half3(0.22h, 0.67h, 0.11h))); 21*c8dee2aaSAndroid Build Coastguard Worker half4 gray = _0_color; 22*c8dee2aaSAndroid Build Coastguard Worker _globals.texOut.write(gray, _in.sk_GlobalInvocationID.xy); 23*c8dee2aaSAndroid Build Coastguard Worker } 24*c8dee2aaSAndroid Build Coastguard Worker return; 25*c8dee2aaSAndroid Build Coastguard Worker} 26