Changes between Initial Version and Version 1 of Ticket #20455


Ignore:
Timestamp:
Jun 5, 2026, 9:07:36 AM (2 days ago)
Author:
Eric Pettersen
Comment:

Legend:

Unmodified
Added
Removed
Modified

  • Ticket #20455

    • Property Component Unassigned → Third Party
    • Property Owner set to Tristan Croll
    • Property Platform → all
    • Property Project → ChimeraX
    • Property Status new → assigned
    • Property Summary ChimeraX bug report submission → isolde sim start: OpenMMException: Error compiling kernel

  • Ticket #20455 – Description

    initial v1  
    19741974simd_sum 
    19751975^ 
    1976 program_source:317:1: error: implicit declaration of function 'simd_sum' is
    1977 invalid in OpenCL 
    1978 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    1979 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    1980 ^ 
    1981 program_source:305:26: note: expanded from macro '\ 
    1982 DECLARE_I_REDUCTION_UNIFORM' 
    1983 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    1984 ^ 
    1985 <scratch space>:17:1: note: expanded from here 
    1986 simd_sum 
    1987 ^ 
    1988 program_source:317:1: error: implicit declaration of function 'simd_sum' is
    1989 invalid in OpenCL 
    1990 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    1991 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    1992 ^ 
    1993 program_source:308:26: note: expanded from macro '\ 
    1994 DECLARE_F_REDUCTION_UNIFORM' 
    1995 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    1996 ^ 
    1997 <scratch space>:19:1: note: expanded from here 
    1998 simd_sum 
    1999 ^ 
    2000 program_source:318:1: error: implicit declaration of function
    2001 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    2002 DECLARE_REDUCTION_UNIFORM(prefix_inclusive_sum, scan_inclusive_add) 
    2003 ^ 
    2004 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2005 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2006 ^ 
    2007 program_source:305:26: note: expanded from macro '\ 
    2008 DECLARE_I_REDUCTION_UNIFORM' 
    2009 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2010 ^ 
    2011 <scratch space>:21:1: note: expanded from here 
    2012 simd_prefix_inclusive_sum 
    2013 ^ 
    2014 program_source:318:1: error: implicit declaration of function
    2015 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    2016 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2017 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2018 ^ 
    2019 program_source:305:26: note: expanded from macro '\ 
    2020 DECLARE_I_REDUCTION_UNIFORM' 
    2021 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2022 ^ 
    2023 <scratch space>:21:1: note: expanded from here 
    2024 simd_prefix_inclusive_sum 
    2025 ^ 
    2026 program_source:318:1: error: implicit declaration of function
    2027 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    2028 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2029 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2030 ^ 
    2031 program_source:308:26: note: expanded from macro '\ 
    2032 DECLARE_F_REDUCTION_UNIFORM' 
    2033 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2034 ^ 
    2035 <scratch space>:23:1: note: expanded from here 
    2036 simd_prefix_inclusive_sum 
    2037 ^ 
    2038 program_source:319:1: error: implicit declaration of function
    2039 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    2040 DECLARE_REDUCTION_UNIFORM(prefix_exclusive_sum, scan_exclusive_add) 
    2041 ^ 
    2042 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2043 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2044 ^ 
    2045 program_source:305:26: note: expanded from macro '\ 
    2046 DECLARE_I_REDUCTION_UNIFORM' 
    2047 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2048 ^ 
    2049 <scratch space>:25:1: note: expanded from here 
    2050 simd_prefix_exclusive_sum 
    2051 ^ 
    2052 program_source:319:1: error: implicit declaration of function
    2053 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    2054 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2055 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2056 ^ 
    2057 program_source:305:26: note: expanded from macro '\ 
    2058 DECLARE_I_REDUCTION_UNIFORM' 
    2059 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2060 ^ 
    2061 <scratch space>:25:1: note: expanded from here 
    2062 simd_prefix_exclusive_sum 
    2063 ^ 
    2064 program_source:319:1: error: implicit declaration of function
    2065 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    2066 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2067 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2068 ^ 
    2069 program_source:308:26: note: expanded from macro '\ 
    2070 DECLARE_F_REDUCTION_UNIFORM' 
    2071 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2072 ^ 
    2073 <scratch space>:27:1: note: expanded from here 
    2074 simd_prefix_exclusive_sum 
    2075 ^ 
    2076 program_source:320:1: error: implicit declaration of function 'simd_min' is
    2077 invalid in OpenCL 
    2078 DECLARE_REDUCTION_UNIFORM(min, reduce_min) 
    2079 ^ 
    2080 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2081 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2082 ^ 
    2083 program_source:305:26: note: expanded from macro '\ 
    2084 DECLARE_I_REDUCTION_UNIFORM' 
    2085 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2086 ^ 
    2087 <scratch space>:29:1: note: expanded from here 
    2088 simd_min 
    2089 ^ 
    2090 program_source:320:1: error: implicit declaration of function 'simd_min' is
    2091 invalid in OpenCL 
    2092 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2093 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2094 ^ 
    2095 program_source:305:26: note: expanded from macro '\ 
    2096 DECLARE_I_REDUCTION_UNIFORM' 
    2097 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2098 ^ 
    2099 <scratch space>:29:1: note: expanded from here 
    2100 simd_min 
    2101 ^ 
    2102 program_source:320:1: error: implicit declaration of function 'simd_min' is
    2103 invalid in OpenCL 
    2104 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2105 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2106 ^ 
    2107 program_source:308:26: note: expanded from macro '\ 
    2108 DECLARE_F_REDUCTION_UNIFORM' 
    2109 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2110 ^ 
    2111 <scratch space>:31:1: note: expanded from here 
    2112 simd_min 
    2113 ^ 
    2114 program_source:321:1: error: implicit declaration of function 'simd_max' is
    2115 invalid in OpenCL 
    2116 DECLARE_REDUCTION_UNIFORM(max, reduce_max) 
    2117 ^ 
    2118 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2119 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2120 ^ 
    2121 program_source:305:26: note: expanded from macro '\ 
    2122 DECLARE_I_REDUCTION_UNIFORM' 
    2123 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2124 ^ 
    2125 <scratch space>:33:1: note: expanded from here 
    2126 simd_max 
    2127 ^ 
    2128 program_source:321:1: error: implicit declaration of function 'simd_max' is
    2129 invalid in OpenCL 
    2130 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2131 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2132 ^ 
    2133 program_source:305:26: note: expanded from macro '\ 
    2134 DECLARE_I_REDUCTION_UNIFORM' 
    2135 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2136 ^ 
    2137 <scratch space>:33:1: note: expanded from here 
    2138 simd_max 
    2139 ^ 
    2140 program_source:321:1: error: implicit declaration of function 'simd_max' is
    2141 invalid in OpenCL 
    2142 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2143 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2144 ^ 
    2145 program_source:308:26: note: expanded from macro '\ 
    2146 DECLARE_F_REDUCTION_UNIFORM' 
    2147 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2148 ^ 
    2149 <scratch space>:35:1: note: expanded from here 
    2150 simd_max 
    2151 ^ 
    2152 program_source:323:1: error: implicit declaration of function 'simd_shuffle'
    2153 is invalid in OpenCL 
    2154 DECLARE_SHUFFLE_UNIFORM(shuffle, shuffle) 
    2155 ^ 
    2156 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2157 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2158 ^ 
    2159 <scratch space>:37:1: note: expanded from here 
    2160 simd_shuffle 
    2161 ^ 
    2162 program_source:323:1: error: implicit declaration of function 'simd_shuffle'
    2163 is invalid in OpenCL 
    2164 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2165 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2166 ^ 
    2167 <scratch space>:37:1: note: expanded from here 
    2168 simd_shuffle 
    2169 ^ 
    2170 program_source:323:1: error: implicit declaration of function 'simd_shuffle'
    2171 is invalid in OpenCL 
    2172 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2173 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2174 ^ 
    2175 <scratch space>:37:1: note: expanded from here 
    2176 simd_shuffle 
    2177 ^ 
    2178 program_source:324:1: error: implicit declaration of function
    2179 'simd_shuffle_xor' is invalid in OpenCL 
    2180 DECLARE_SHUFFLE_UNIFORM(shuffle_xor, shuffle_xor) 
    2181 ^ 
    2182 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2183 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2184 ^ 
    2185 <scratch space>:39:1: note: expanded from here 
    2186 simd_shuffle_xor 
    2187 ^ 
    2188 program_source:324:1: error: implicit declaration of function
    2189 'simd_shuffle_xor' is invalid in OpenCL 
    2190 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2191 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2192 ^ 
    2193 <scratch space>:39:1: note: expanded from here 
    2194 simd_shuffle_xor 
    2195 ^ 
    2196 program_source:324:1: error: implicit declaration of function
    2197 'simd_shuffle_xor' is invalid in OpenCL 
    2198 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2199 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2200 ^ 
    2201 <scratch space>:39:1: note: expanded from here 
    2202 simd_shuffle_xor 
    2203 ^ 
    2204 program_source:325:1: error: implicit declaration of function
    2205 'simd_shuffle_up' is invalid in OpenCL 
    2206 DECLARE_SHUFFLE_UNIFORM(shuffle_up, shuffle_up) 
    2207 ^ 
    2208 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2209 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2210 ^ 
    2211 <scratch space>:41:1: note: expanded from here 
    2212 simd_shuffle_up 
    2213 ^ 
    2214 program_source:325:1: error: implicit declaration of function
    2215 'simd_shuffle_up' is invalid in OpenCL 
    2216 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2217 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2218 ^ 
    2219 <scratch space>:41:1: note: expanded from here 
    2220 simd_shuffle_up 
    2221 ^ 
    2222 program_source:325:1: error: implicit declaration of function
    2223 'simd_shuffle_up' is invalid in OpenCL 
    2224 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2225 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2226 ^ 
    2227 <scratch space>:41:1: note: expanded from here 
    2228 simd_shuffle_up 
    2229 ^ 
    2230 program_source:326:1: error: implicit declaration of function
    2231 'simd_shuffle_down' is invalid in OpenCL 
    2232 DECLARE_SHUFFLE_UNIFORM(shuffle_down, shuffle_down) 
    2233 ^ 
    2234 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2235 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2236 ^ 
    2237 <scratch space>:43:1: note: expanded from here 
    2238 simd_shuffle_down 
    2239 ^ 
    2240 program_source:326:1: error: implicit declaration of function
    2241 'simd_shuffle_down' is invalid in OpenCL 
    2242 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2243 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2244 ^ 
    2245 <scratch space>:43:1: note: expanded from here 
    2246 simd_shuffle_down 
    2247 ^ 
    2248 program_source:326:1: error: implicit declaration of function
    2249 'simd_shuffle_down' is invalid in OpenCL 
    2250 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2251 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2252 ^ 
    2253 <scratch space>:43:1: note: expanded from here 
    2254 simd_shuffle_down 
    2255 ^ 
    2256 program_source:327:1: error: implicit declaration of function
    2257 'simd_shuffle_rotate_down' is invalid in OpenCL 
    2258 DECLARE_SHUFFLE_UNIFORM(shuffle_rotate_down, rotate) 
    2259 ^ 
    2260 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2261 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2262 ^ 
    2263 <scratch space>:45:1: note: expanded from here 
    2264 simd_shuffle_rotate_down 
    2265 ^ 
    2266 program_source:327:1: error: implicit declaration of function
    2267 'simd_shuffle_rotate_down' is invalid in OpenCL 
    2268 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2269 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2270 ^ 
    2271 <scratch space>:45:1: note: expanded from here 
    2272 simd_shuffle_rotate_down 
    2273 ^ 
    2274 program_source:327:1: error: implicit declaration of function
    2275 'simd_shuffle_rotate_down' is invalid in OpenCL 
    2276 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2277 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2278 ^ 
    2279 <scratch space>:45:1: note: expanded from here 
    2280 simd_shuffle_rotate_down 
    2281 ^ 
    2282 program_source:329:1: error: implicit declaration of function 'simd_broadcast'
    2283 is invalid in OpenCL 
    2284 DECLARE_SHUFFLE_UNIFORM(broadcast, broadcast) 
    2285 ^ 
    2286 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2287 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2288 ^ 
    2289 <scratch space>:47:1: note: expanded from here 
    2290 simd_broadcast 
    2291 ^ 
    2292 program_source:329:1: error: implicit declaration of function 'simd_broadcast'
    2293 is invalid in OpenCL 
    2294 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2295 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2296 ^ 
    2297 <scratch space>:47:1: note: expanded from here 
    2298 simd_broadcast 
    2299 ^ 
    2300 program_source:329:1: error: implicit declaration of function 'simd_broadcast'
    2301 is invalid in OpenCL 
    2302 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2303 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2304 ^ 
    2305 <scratch space>:47:1: note: expanded from here 
    2306 simd_broadcast 
    2307 ^ 
    2308 program_source:330:1: error: implicit declaration of function
    2309 'simd_broadcast_first' is invalid in OpenCL 
    2310 DECLARE_REDUCTION_UNIFORM(broadcast_first, broadcast_first) 
    2311 ^ 
    2312 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2313 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2314 ^ 
    2315 program_source:305:26: note: expanded from macro '\ 
    2316 DECLARE_I_REDUCTION_UNIFORM' 
    2317 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2318 ^ 
    2319 <scratch space>:49:1: note: expanded from here 
    2320 simd_broadcast_first 
    2321 ^ 
    2322 program_source:330:1: error: implicit declaration of function
    2323 'simd_broadcast_first' is invalid in OpenCL 
    2324 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2325 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2326 ^ 
    2327 program_source:305:26: note: expanded from macro '\ 
    2328 DECLARE_I_REDUCTION_UNIFORM' 
    2329 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2330 ^ 
    2331 <scratch space>:49:1: note: expanded from here 
    2332 simd_broadcast_first 
    2333 ^ 
    2334 program_source:330:1: error: implicit declaration of function
    2335 'simd_broadcast_first' is invalid in OpenCL 
    2336 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2337 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2338 ^ 
    2339 program_source:308:26: note: expanded from macro '\ 
    2340 DECLARE_F_REDUCTION_UNIFORM' 
    2341 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2342 ^ 
    2343 <scratch space>:51:1: note: expanded from here 
    2344 simd_broadcast_first 
    2345 ^ 
    2346 program_source:345:1: error: implicit declaration of function 'simd_sum' is
    2347 invalid in OpenCL 
    2348 DECLARE_REDUCTION_NON_UNIFORM(sum, reduce_add) 
    2349 ^ 
    2350 program_source:338:60: note: expanded from macro
    2351 'DECLARE_REDUCTION_NON_UNIFORM' 
    2352 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2353 ^ 
    2354 program_source:333:26: note: expanded from macro '\ 
    2355 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2356 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2357 ^ 
    2358 <scratch space>:53:1: note: expanded from here 
    2359 simd_sum 
    2360 ^ 
    2361 program_source:345:1: error: implicit declaration of function 'simd_sum' is
    2362 invalid in OpenCL 
    2363 program_source:338:60: note: expanded from macro
    2364 'DECLARE_REDUCTION_NON_UNIFORM' 
    2365 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2366 ^ 
    2367 program_source:333:26: note: expanded from macro '\ 
    2368 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2369 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2370 ^ 
    2371 <scratch space>:53:1: note: expanded from here 
    2372 simd_sum 
    2373 ^ 
    2374 program_source:345:1: error: implicit declaration of function 'simd_sum' is
    2375 invalid in OpenCL 
    2376 program_source:339:54: note: expanded from macro
    2377 'DECLARE_REDUCTION_NON_UNIFORM' 
    2378 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2379 ^ 
    2380 program_source:336:26: note: expanded from macro '\ 
    2381 DECLARE_F_REDUCTION_NON_UNIFORM' 
    2382 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2383 ^ 
    2384 <scratch space>:55:1: note: expanded from here 
    2385 simd_sum 
    2386 ^ 
    2387 program_source:346:1: error: implicit declaration of function
    2388 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    2389 DECLARE_REDUCTION_NON_UNIFORM(prefix_inclusive_sum, scan_inclusive_add) 
    2390 ^ 
    2391 program_source:338:60: note: expanded from macro
    2392 'DECLARE_REDUCTION_NON_UNIFORM' 
    2393 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2394 ^ 
    2395 program_source:333:26: note: expanded from macro '\ 
    2396 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2397 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2398 ^ 
    2399 <scratch space>:57:1: note: expanded from here 
    2400 simd_prefix_inclusive_sum 
    2401 ^ 
    2402 program_source:346:1: error: implicit declaration of function
    2403 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    2404 program_source:338:60: note: expanded from macro
    2405 'DECLARE_REDUCTION_NON_UNIFORM' 
    2406 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2407 ^ 
    2408 program_source:333:26: note: expanded from macro '\ 
    2409 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2410 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2411 ^ 
    2412 <scratch space>:57:1: note: expanded from here 
    2413 simd_prefix_inclusive_sum 
    2414 ^ 
    2415 program_source:346:1: error: implicit declaration of function
    2416 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    2417 program_source:339:54: note: expanded from macro
    2418 'DECLARE_REDUCTION_NON_UNIFORM' 
    2419 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2420 ^ 
    2421 program_source:336:26: note: expanded from macro '\ 
    2422 DECLARE_F_REDUCTION_NON_UNIFORM' 
    2423 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2424 ^ 
    2425 <scratch space>:59:1: note: expanded from here 
    2426 simd_prefix_inclusive_sum 
    2427 ^ 
    2428 program_source:347:1: error: implicit declaration of function
    2429 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    2430 DECLARE_REDUCTION_NON_UNIFORM(prefix_exclusive_sum, scan_exclusive_add) 
    2431 ^ 
    2432 program_source:338:60: note: expanded from macro
    2433 'DECLARE_REDUCTION_NON_UNIFORM' 
    2434 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2435 ^ 
    2436 program_source:333:26: note: expanded from macro '\ 
    2437 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2438 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2439 ^ 
    2440 <scratch space>:61:1: note: expanded from here 
    2441 simd_prefix_exclusive_sum 
    2442 ^ 
    2443 program_source:347:1: error: implicit declaration of function
    2444 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    2445 program_source:338:60: note: expanded from macro
    2446 'DECLARE_REDUCTION_NON_UNIFORM' 
    2447 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2448 ^ 
    2449 program_source:333:26: note: expanded from macro '\ 
    2450 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2451 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2452 ^ 
    2453 <scratch space>:61:1: note: expanded from here 
    2454 simd_prefix_exclusive_sum 
    2455 ^ 
    2456 program_source:347:1: error: implicit declaration of function
    2457 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    2458 program_source:339:54: note: expanded from macro
    2459 'DECLARE_REDUCTION_NON_UNIFORM' 
    2460 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2461 ^ 
    2462 program_source:336:26: note: expanded from macro '\ 
    2463 DECLARE_F_REDUCTION_NON_UNIFORM' 
    2464 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2465 ^ 
    2466 <scratch space>:63:1: note: expanded from here 
    2467 simd_prefix_exclusive_sum 
    2468 ^ 
    2469 program_source:348:1: error: implicit declaration of function 'simd_min' is
    2470 invalid in OpenCL 
    2471 DECLARE_REDUCTION_NON_UNIFORM(min, reduce_min) 
    2472 ^ 
    2473 program_source:338:60: note: expanded from macro
    2474 'DECLARE_REDUCTION_NON_UNIFORM' 
    2475 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2476 ^ 
    2477 program_source:333:26: note: expanded from macro '\ 
    2478 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2479 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2480 ^ 
    2481 <scratch space>:65:1: note: expanded from here 
    2482 simd_min 
    2483 ^ 
    2484 program_source:348:1: error: implicit declaration of function 'simd_min' is
    2485 invalid in OpenCL 
    2486 program_source:338:60: note: expanded from macro
    2487 'DECLARE_REDUCTION_NON_UNIFORM' 
    2488 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2489 ^ 
    2490 program_source:333:26: note: expanded from macro '\ 
    2491 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2492 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2493 ^ 
    2494 <scratch space>:65:1: note: expanded from here 
    2495 simd_min 
    2496 ^ 
    2497 program_source:348:1: error: implicit declaration of function 'simd_min' is
    2498 invalid in OpenCL 
    2499 program_source:339:54: note: expanded from macro
    2500 'DECLARE_REDUCTION_NON_UNIFORM' 
    2501 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2502 ^ 
    2503 program_source:336:26: note: expanded from macro '\ 
    2504 DECLARE_F_REDUCTION_NON_UNIFORM' 
    2505 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2506 ^ 
    2507 <scratch space>:67:1: note: expanded from here 
    2508 simd_min 
    2509 ^ 
    2510 program_source:349:1: error: implicit declaration of function 'simd_max' is
    2511 invalid in OpenCL 
    2512 DECLARE_REDUCTION_NON_UNIFORM(max, reduce_max) 
    2513 ^ 
    2514 program_source:338:60: note: expanded from macro
    2515 'DECLARE_REDUCTION_NON_UNIFORM' 
    2516 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2517 ^ 
    2518 program_source:333:26: note: expanded from macro '\ 
    2519 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2520 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2521 ^ 
    2522 <scratch space>:69:1: note: expanded from here 
    2523 simd_max 
    2524 ^ 
    2525 program_source:349:1: error: implicit declaration of function 'simd_max' is
    2526 invalid in OpenCL 
    2527 program_source:338:60: note: expanded from macro
    2528 'DECLARE_REDUCTION_NON_UNIFORM' 
    2529 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2530 ^ 
    2531 program_source:333:26: note: expanded from macro '\ 
    2532 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2533 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2534 ^ 
    2535 <scratch space>:69:1: note: expanded from here 
    2536 simd_max 
    2537 ^ 
    2538 program_source:349:1: error: implicit declaration of function 'simd_max' is
    2539 invalid in OpenCL 
    2540 program_source:339:54: note: expanded from macro
    2541 'DECLARE_REDUCTION_NON_UNIFORM' 
    2542 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2543 ^ 
    2544 program_source:336:26: note: expanded from macro '\ 
    2545 DECLARE_F_REDUCTION_NON_UNIFORM' 
    2546 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2547 ^ 
    2548 <scratch space>:71:1: note: expanded from here 
    2549 simd_max 
    2550 ^ 
    2551 program_source:351:1: error: implicit declaration of function 'simd_product'
    2552 is invalid in OpenCL 
    2553 DECLARE_REDUCTION_NON_UNIFORM(product, reduce_mul) 
    2554 ^ 
    2555 program_source:338:60: note: expanded from macro
    2556 'DECLARE_REDUCTION_NON_UNIFORM' 
    2557 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2558 ^ 
    2559 program_source:333:26: note: expanded from macro '\ 
    2560 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2561 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2562 ^ 
    2563 <scratch space>:73:1: note: expanded from here 
    2564 simd_product 
    2565 ^ 
    2566 program_source:351:1: error: implicit declaration of function 'simd_product'
    2567 is invalid in OpenCL 
    2568 program_source:338:60: note: expanded from macro
    2569 'DECLARE_REDUCTION_NON_UNIFORM' 
    2570 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2571 ^ 
    2572 program_source:333:26: note: expanded from macro '\ 
    2573 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2574 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2575 ^ 
    2576 <scratch space>:73:1: note: expanded from here 
    2577 simd_product 
    2578 ^ 
    2579 program_source:351:1: error: implicit declaration of function 'simd_product'
    2580 is invalid in OpenCL 
    2581 program_source:339:54: note: expanded from macro
    2582 'DECLARE_REDUCTION_NON_UNIFORM' 
    2583 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2584 ^ 
    2585 program_source:336:26: note: expanded from macro '\ 
    2586 DECLARE_F_REDUCTION_NON_UNIFORM' 
    2587 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2588 ^ 
    2589 <scratch space>:75:1: note: expanded from here 
    2590 simd_product 
    2591 ^ 
    2592 program_source:352:1: error: implicit declaration of function
    2593 'simd_prefix_inclusive_product' is invalid in OpenCL 
    2594 DECLARE_REDUCTION_NON_UNIFORM(prefix_inclusive_product, scan_inclusive_mul) 
    2595 ^ 
    2596 program_source:338:60: note: expanded from macro
    2597 'DECLARE_REDUCTION_NON_UNIFORM' 
    2598 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2599 ^ 
    2600 program_source:333:26: note: expanded from macro '\ 
    2601 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2602 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2603 ^ 
    2604 <scratch space>:77:1: note: expanded from here 
    2605 simd_prefix_inclusive_product 
    2606 ^ 
    2607 program_source:352:1: error: implicit declaration of function
    2608 'simd_prefix_inclusive_product' is invalid in OpenCL 
    2609 program_source:338:60: note: expanded from macro
    2610 'DECLARE_REDUCTION_NON_UNIFORM' 
    2611 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2612 ^ 
    2613 program_source:333:26: note: expanded from macro '\ 
    2614 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2615 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2616 ^ 
    2617 <scratch space>:77:1: note: expanded from here 
    2618 simd_prefix_inclusive_product 
    2619 ^ 
    2620 program_source:352:1: error: implicit declaration of function
    2621 'simd_prefix_inclusive_product' is invalid in OpenCL 
    2622 program_source:339:54: note: expanded from macro
    2623 'DECLARE_REDUCTION_NON_UNIFORM' 
    2624 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2625 ^ 
    2626 program_source:336:26: note: expanded from macro '\ 
    2627 DECLARE_F_REDUCTION_NON_UNIFORM' 
    2628 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2629 ^ 
    2630 <scratch space>:79:1: note: expanded from here 
    2631 simd_prefix_inclusive_product 
    2632 ^ 
    2633 program_source:353:1: error: implicit declaration of function
    2634 'simd_prefix_exclusive_product' is invalid in OpenCL 
    2635 DECLARE_REDUCTION_NON_UNIFORM(prefix_exclusive_product, scan_exclusive_mul) 
    2636 ^ 
    2637 program_source:338:60: note: expanded from macro
    2638 'DECLARE_REDUCTION_NON_UNIFORM' 
    2639 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2640 ^ 
    2641 program_source:333:26: note: expanded from macro '\ 
    2642 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2643 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2644 ^ 
    2645 <scratch space>:81:1: note: expanded from here 
    2646 simd_prefix_exclusive_product 
    2647 ^ 
    2648 program_source:353:1: error: implicit declaration of function
    2649 'simd_prefix_exclusive_product' is invalid in OpenCL 
    2650 program_source:338:60: note: expanded from macro
    2651 'DECLARE_REDUCTION_NON_UNIFORM' 
    2652 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2653 ^ 
    2654 program_source:333:26: note: expanded from macro '\ 
    2655 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2656 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2657 ^ 
    2658 <scratch space>:81:1: note: expanded from here 
    2659 simd_prefix_exclusive_product 
    2660 ^ 
    2661 program_source:353:1: error: implicit declaration of function
    2662 'simd_prefix_exclusive_product' is invalid in OpenCL 
    2663 program_source:339:54: note: expanded from macro
    2664 'DECLARE_REDUCTION_NON_UNIFORM' 
    2665 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2666 ^ 
    2667 program_source:336:26: note: expanded from macro '\ 
    2668 DECLARE_F_REDUCTION_NON_UNIFORM' 
    2669 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2670 ^ 
    2671 <scratch space>:83:1: note: expanded from here 
    2672 simd_prefix_exclusive_product 
    2673 ^ 
    2674 program_source:354:1: error: implicit declaration of function 'simd_and' is
    2675 invalid in OpenCL 
    2676 DECLARE_I_REDUCTION_NON_UNIFORM(and, reduce_and) 
    2677 ^ 
    2678 program_source:333:26: note: expanded from macro
    2679 'DECLARE_I_REDUCTION_NON_UNIFORM' 
    2680 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2681 ^ 
    2682 <scratch space>:85:1: note: expanded from here 
    2683 simd_and 
    2684 ^ 
    2685 program_source:354:1: error: implicit declaration of function 'simd_and' is
    2686 invalid in OpenCL 
    2687 program_source:333:26: note: expanded from macro
    2688 'DECLARE_I_REDUCTION_NON_UNIFORM' 
    2689 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2690 ^ 
    2691 <scratch space>:85:1: note: expanded from here 
    2692 simd_and 
    2693 ^ 
    2694 program_source:355:1: error: implicit declaration of function 'simd_or' is
    2695 invalid in OpenCL 
    2696 DECLARE_I_REDUCTION_NON_UNIFORM(or, reduce_or) 
    2697 ^ 
    2698 program_source:333:26: note: expanded from macro
    2699 'DECLARE_I_REDUCTION_NON_UNIFORM' 
    2700 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2701 ^ 
    2702 <scratch space>:87:1: note: expanded from here 
    2703 simd_or 
    2704 ^ 
    2705 program_source:355:1: error: implicit declaration of function 'simd_or' is
    2706 invalid in OpenCL 
    2707 program_source:333:26: note: expanded from macro
    2708 'DECLARE_I_REDUCTION_NON_UNIFORM' 
    2709 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2710 ^ 
    2711 <scratch space>:87:1: note: expanded from here 
    2712 simd_or 
    2713 ^ 
    2714 program_source:356:1: error: implicit declaration of function 'simd_xor' is
    2715 invalid in OpenCL 
    2716 DECLARE_I_REDUCTION_NON_UNIFORM(xor, reduce_xor) 
    2717 ^ 
    2718 program_source:333:26: note: expanded from macro
    2719 'DECLARE_I_REDUCTION_NON_UNIFORM' 
    2720 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2721 ^ 
    2722 <scratch space>:89:1: note: expanded from here 
    2723 simd_xor 
    2724 ^ 
    2725 program_source:356:1: error: implicit declaration of function 'simd_xor' is
    2726 invalid in OpenCL 
    2727 program_source:333:26: note: expanded from macro
    2728 'DECLARE_I_REDUCTION_NON_UNIFORM' 
    2729 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2730 ^ 
    2731 <scratch space>:89:1: note: expanded from here 
    2732 simd_xor 
    2733 ^ 
    2734 program_source:358:1: error: implicit declaration of function 'simd_broadcast'
    2735 is invalid in OpenCL 
    2736 DECLARE_SHUFFLE_NON_UNIFORM(broadcast, broadcast) 
    2737 ^ 
    2738 program_source:343:24: note: expanded from macro 'DECLARE_SHUFFLE_NON_UNIFORM' 
    2739 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2740 ^ 
    2741 <scratch space>:91:1: note: expanded from here 
    2742 simd_broadcast 
    2743 ^ 
    2744 program_source:358:1: error: implicit declaration of function 'simd_broadcast'
    2745 is invalid in OpenCL 
    2746 program_source:343:24: note: expanded from macro 'DECLARE_SHUFFLE_NON_UNIFORM' 
    2747 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2748 ^ 
    2749 <scratch space>:91:1: note: expanded from here 
    2750 simd_broadcast 
    2751 ^ 
    2752 program_source:358:1: error: implicit declaration of function 'simd_broadcast'
    2753 is invalid in OpenCL 
    2754 program_source:343:24: note: expanded from macro 'DECLARE_SHUFFLE_NON_UNIFORM' 
    2755 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2756 ^ 
    2757 <scratch space>:91:1: note: expanded from here 
    2758 simd_broadcast 
    2759 ^ 
    2760 program_source:391:1: error: implicit declaration of function 'quad_sum' is
    2761 invalid in OpenCL 
    2762 DECLARE_REDUCTION_CLUSTERED(sum, reduce_add) 
    2763 ^ 
    2764 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    2765 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2766 ^ 
    2767 program_source:360:60: note: expanded from macro '\ 
    2768 DECLARE_I_REDUCTION_CLUSTERED' 
    2769 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2770 ^ 
    2771 program_source:245:71: note: expanded from macro '\ 
    2772 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    2773 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    2774 ^ 
    2775 program_source:196:12: note: expanded from macro '\ 
    2776 OVERLOAD_CLUSTERED_ARGS_1' 
    2777 return quad_##METAL_SUFFIX(x); \ 
    2778 ^ 
    2779 <scratch space>:94:1: note: expanded from here 
    2780 quad_sum 
    2781 ^ 
    2782 program_source:391:1: error: implicit declaration of function 'simd_sum' is
    2783 invalid in OpenCL 
    2784 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    2785 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2786 ^ 
    2787 program_source:360:60: note: expanded from macro '\ 
    2788 DECLARE_I_REDUCTION_CLUSTERED' 
    2789 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2790 ^ 
    2791 program_source:245:71: note: expanded from macro '\ 
    2792 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    2793 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    2794 ^ 
    2795 program_source:198:12: note: expanded from macro '\ 
    2796 OVERLOAD_CLUSTERED_ARGS_1' 
    2797 return simd_##METAL_SUFFIX(x); \ 
    2798 ^ 
    2799 <scratch space>:95:1: note: expanded from here 
    2800 simd_sum 
    2801 ^ 
    2802 program_source:391:1: error: implicit declaration of function 'quad_sum' is
    2803 invalid in OpenCL 
    2804 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    2805 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2806 ^ 
    2807 program_source:360:60: note: expanded from macro '\ 
    2808 DECLARE_I_REDUCTION_CLUSTERED' 
    2809 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2810 ^ 
    2811 program_source:246:59: note: expanded from macro '\ 
    2812 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    2813 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    2814 ^ 
    2815 program_source:196:12: note: expanded from macro '\ 
    2816 OVERLOAD_CLUSTERED_ARGS_1' 
    2817 return quad_##METAL_SUFFIX(x); \ 
    2818 ^ 
    2819 <scratch space>:96:1: note: expanded from here 
    2820 quad_sum 
    2821 ^ 
    2822 program_source:391:1: error: implicit declaration of function 'simd_sum' is
    2823 invalid in OpenCL 
    2824 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    2825 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2826 ^ 
    2827 program_source:360:60: note: expanded from macro '\ 
    2828 DECLARE_I_REDUCTION_CLUSTERED' 
    2829 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2830 ^ 
    2831 program_source:246:59: note: expanded from macro '\ 
    2832 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    2833 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    2834 ^ 
    2835 program_source:198:12: note: expanded from macro '\ 
    2836 OVERLOAD_CLUSTERED_ARGS_1' 
    2837 return simd_##METAL_SUFFIX(x); \ 
    2838 ^ 
    2839 <scratch space>:97:1: note: expanded from here 
    2840 simd_sum 
    2841 ^ 
    2842 program_source:391:1: error: implicit declaration of function 'quad_sum' is
    2843 invalid in OpenCL 
    2844 program_source:385:52: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    2845 DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2846 ^ 
    2847 program_source:363:60: note: expanded from macro '\ 
    2848 DECLARE_F_REDUCTION_CLUSTERED' 
    2849 #define DECLARE_F_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2850 ^ 
    2851 program_source:249:71: note: expanded from macro '\ 
    2852 BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1' 
    2853 #define BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    2854 ^ 
    2855 program_source:196:12: note: expanded from macro '\ 
    2856 OVERLOAD_CLUSTERED_ARGS_1' 
    2857 return quad_##METAL_SUFFIX(x); \ 
    2858 ^ 
    2859 <scratch space>:99:1: note: expanded from here 
    2860 quad_sum 
    2861 ^ 
    2862 program_source:391:1: error: implicit declaration of function 'simd_sum' is
    2863 invalid in OpenCL 
    2864 program_source:385:52: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    2865 DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2866 ^ 
    2867 program_source:363:60: note: expanded from macro '\ 
    2868 DECLARE_F_REDUCTION_CLUSTERED' 
    2869 #define DECLARE_F_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2870 ^ 
    2871 program_source:249:71: note: expanded from macro '\ 
    2872 BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1' 
    2873 #define BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    2874 ^ 
    2875 program_source:198:12: note: expanded from macro '\ 
    2876 OVERLOAD_CLUSTERED_ARGS_1' 
    2877 return simd_##METAL_SUFFIX(x); \ 
    2878 ^ 
    2879 <scratch space>:100:1: note: expanded from here 
    2880 simd_sum 
    2881 ^ 
    2882 program_source:392:1: error: implicit declaration of function 'quad_min' is
    2883 invalid in OpenCL 
    2884 DECLARE_REDUCTION_CLUSTERED(min, reduce_min) 
    2885 ^ 
    2886 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    2887 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2888 ^ 
    2889 program_source:360:60: note: expanded from macro '\ 
    2890 DECLARE_I_REDUCTION_CLUSTERED' 
    2891 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2892 ^ 
    2893 program_source:245:71: note: expanded from macro '\ 
    2894 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    2895 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    2896 ^ 
    2897 program_source:196:12: note: expanded from macro '\ 
    2898 OVERLOAD_CLUSTERED_ARGS_1' 
    2899 return quad_##METAL_SUFFIX(x); \ 
    2900 ^ 
    2901 <scratch space>:102:1: note: expanded from here 
    2902 quad_min 
    2903 ^ 
    2904 program_source:392:1: error: implicit declaration of function 'simd_min' is
    2905 invalid in OpenCL 
    2906 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    2907 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2908 ^ 
    2909 program_source:360:60: note: expanded from macro '\ 
    2910 DECLARE_I_REDUCTION_CLUSTERED' 
    2911 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2912 ^ 
    2913 program_source:245:71: note: expanded from macro '\ 
    2914 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    2915 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    2916 ^ 
    2917 program_source:198:12: note: expanded from macro '\ 
    2918 OVERLOAD_CLUSTERED_ARGS_1' 
    2919 return simd_##METAL_SUFFIX(x); \ 
    2920 ^ 
    2921 <scratch space>:103:1: note: expanded from here 
    2922 simd_min 
    2923 ^ 
    2924 program_source:392:1: error: implicit declaration of function 'quad_min' is
    2925 invalid in OpenCL 
    2926 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    2927 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2928 ^ 
    2929 program_source:360:60: note: expanded from macro '\ 
    2930 DECLARE_I_REDUCTION_CLUSTERED' 
    2931 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2932 ^ 
    2933 program_source:246:59: note: expanded from macro '\ 
    2934 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    2935 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    2936 ^ 
    2937 program_source:196:12: note: expanded from macro '\ 
    2938 OVERLOAD_CLUSTERED_ARGS_1' 
    2939 return quad_##METAL_SUFFIX(x); \ 
    2940 ^ 
    2941 <scratch space>:104:1: note: expanded from here 
    2942 quad_min 
    2943 ^ 
    2944 program_source:392:1: error: implicit declaration of function 'simd_min' is
    2945 invalid in OpenCL 
    2946 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    2947 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2948 ^ 
    2949 program_source:360:60: note: expanded from macro '\ 
    2950 DECLARE_I_REDUCTION_CLUSTERED' 
    2951 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2952 ^ 
    2953 program_source:246:59: note: expanded from macro '\ 
    2954 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    2955 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    2956 ^ 
    2957 program_source:198:12: note: expanded from macro '\ 
    2958 OVERLOAD_CLUSTERED_ARGS_1' 
    2959 return simd_##METAL_SUFFIX(x); \ 
    2960 ^ 
    2961 <scratch space>:105:1: note: expanded from here 
    2962 simd_min 
    2963 ^ 
    2964 program_source:392:1: error: implicit declaration of function 'quad_min' is
    2965 invalid in OpenCL 
    2966 program_source:385:52: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    2967 DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2968 ^ 
    2969 program_source:363:60: note: expanded from macro '\ 
    2970 DECLARE_F_REDUCTION_CLUSTERED' 
    2971 #define DECLARE_F_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2972 ^ 
    2973 program_source:249:71: note: expanded from macro '\ 
    2974 BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1' 
    2975 #define BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    2976 ^ 
    2977 program_source:196:12: note: expanded from macro '\ 
    2978 OVERLOAD_CLUSTERED_ARGS_1' 
    2979 return quad_##METAL_SUFFIX(x); \ 
    2980 ^ 
    2981 <scratch space>:107:1: note: expanded from here 
    2982 quad_min 
    2983 ^ 
    2984 program_source:392:1: error: implicit declaration of function 'simd_min' is
    2985 invalid in OpenCL 
    2986 program_source:385:52: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    2987 DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2988 ^ 
    2989 program_source:363:60: note: expanded from macro '\ 
    2990 DECLARE_F_REDUCTION_CLUSTERED' 
    2991 #define DECLARE_F_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2992 ^ 
    2993 program_source:249:71: note: expanded from macro '\ 
    2994 BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1' 
    2995 #define BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    2996 ^ 
    2997 program_source:198:12: note: expanded from macro '\ 
    2998 OVERLOAD_CLUSTERED_ARGS_1' 
    2999 return simd_##METAL_SUFFIX(x); \ 
    3000 ^ 
    3001 <scratch space>:108:1: note: expanded from here 
    3002 simd_min 
    3003 ^ 
    3004 program_source:393:1: error: implicit declaration of function 'quad_max' is
    3005 invalid in OpenCL 
    3006 DECLARE_REDUCTION_CLUSTERED(max, reduce_max) 
    3007 ^ 
    3008 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3009 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3010 ^ 
    3011 program_source:360:60: note: expanded from macro '\ 
    3012 DECLARE_I_REDUCTION_CLUSTERED' 
    3013 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3014 ^ 
    3015 program_source:245:71: note: expanded from macro '\ 
    3016 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3017 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3018 ^ 
    3019 program_source:196:12: note: expanded from macro '\ 
    3020 OVERLOAD_CLUSTERED_ARGS_1' 
    3021 return quad_##METAL_SUFFIX(x); \ 
    3022 ^ 
    3023 <scratch space>:110:1: note: expanded from here 
    3024 quad_max 
    3025 ^ 
    3026 program_source:393:1: error: implicit declaration of function 'simd_max' is
    3027 invalid in OpenCL 
    3028 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3029 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3030 ^ 
    3031 program_source:360:60: note: expanded from macro '\ 
    3032 DECLARE_I_REDUCTION_CLUSTERED' 
    3033 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3034 ^ 
    3035 program_source:245:71: note: expanded from macro '\ 
    3036 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3037 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3038 ^ 
    3039 program_source:198:12: note: expanded from macro '\ 
    3040 OVERLOAD_CLUSTERED_ARGS_1' 
    3041 return simd_##METAL_SUFFIX(x); \ 
    3042 ^ 
    3043 <scratch space>:111:1: note: expanded from here 
    3044 simd_max 
    3045 ^ 
    3046 program_source:393:1: error: implicit declaration of function 'quad_max' is
    3047 invalid in OpenCL 
    3048 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3049 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3050 ^ 
    3051 program_source:360:60: note: expanded from macro '\ 
    3052 DECLARE_I_REDUCTION_CLUSTERED' 
    3053 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3054 ^ 
    3055 program_source:246:59: note: expanded from macro '\ 
    3056 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3057 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3058 ^ 
    3059 program_source:196:12: note: expanded from macro '\ 
    3060 OVERLOAD_CLUSTERED_ARGS_1' 
    3061 return quad_##METAL_SUFFIX(x); \ 
    3062 ^ 
    3063 <scratch space>:112:1: note: expanded from here 
    3064 quad_max 
    3065 ^ 
    3066 program_source:393:1: error: implicit declaration of function 'simd_max' is
    3067 invalid in OpenCL 
    3068 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3069 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3070 ^ 
    3071 program_source:360:60: note: expanded from macro '\ 
    3072 DECLARE_I_REDUCTION_CLUSTERED' 
    3073 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3074 ^ 
    3075 program_source:246:59: note: expanded from macro '\ 
    3076 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3077 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3078 ^ 
    3079 program_source:198:12: note: expanded from macro '\ 
    3080 OVERLOAD_CLUSTERED_ARGS_1' 
    3081 return simd_##METAL_SUFFIX(x); \ 
    3082 ^ 
    3083 <scratch space>:113:1: note: expanded from here 
    3084 simd_max 
    3085 ^ 
    3086 program_source:393:1: error: implicit declaration of function 'quad_max' is
    3087 invalid in OpenCL 
    3088 program_source:385:52: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3089 DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3090 ^ 
    3091 program_source:363:60: note: expanded from macro '\ 
    3092 DECLARE_F_REDUCTION_CLUSTERED' 
    3093 #define DECLARE_F_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3094 ^ 
    3095 program_source:249:71: note: expanded from macro '\ 
    3096 BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1' 
    3097 #define BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3098 ^ 
    3099 program_source:196:12: note: expanded from macro '\ 
    3100 OVERLOAD_CLUSTERED_ARGS_1' 
    3101 return quad_##METAL_SUFFIX(x); \ 
    3102 ^ 
    3103 <scratch space>:115:1: note: expanded from here 
    3104 quad_max 
    3105 ^ 
    3106 program_source:393:1: error: implicit declaration of function 'simd_max' is
    3107 invalid in OpenCL 
    3108 program_source:385:52: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3109 DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3110 ^ 
    3111 program_source:363:60: note: expanded from macro '\ 
    3112 DECLARE_F_REDUCTION_CLUSTERED' 
    3113 #define DECLARE_F_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3114 ^ 
    3115 program_source:249:71: note: expanded from macro '\ 
    3116 BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1' 
    3117 #define BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3118 ^ 
    3119 program_source:198:12: note: expanded from macro '\ 
    3120 OVERLOAD_CLUSTERED_ARGS_1' 
    3121 return simd_##METAL_SUFFIX(x); \ 
    3122 ^ 
    3123 <scratch space>:116:1: note: expanded from here 
    3124 simd_max 
    3125 ^ 
    3126 program_source:395:1: error: implicit declaration of function 'quad_product'
    3127 is invalid in OpenCL 
    3128 DECLARE_REDUCTION_CLUSTERED(product, reduce_mul) 
    3129 ^ 
    3130 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3131 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3132 ^ 
    3133 program_source:360:60: note: expanded from macro '\ 
    3134 DECLARE_I_REDUCTION_CLUSTERED' 
    3135 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3136 ^ 
    3137 program_source:245:71: note: expanded from macro '\ 
    3138 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3139 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3140 ^ 
    3141 program_source:196:12: note: expanded from macro '\ 
    3142 OVERLOAD_CLUSTERED_ARGS_1' 
    3143 return quad_##METAL_SUFFIX(x); \ 
    3144 ^ 
    3145 <scratch space>:118:1: note: expanded from here 
    3146 quad_product 
    3147 ^ 
    3148 program_source:395:1: error: implicit declaration of function 'simd_product'
    3149 is invalid in OpenCL 
    3150 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3151 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3152 ^ 
    3153 program_source:360:60: note: expanded from macro '\ 
    3154 DECLARE_I_REDUCTION_CLUSTERED' 
    3155 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3156 ^ 
    3157 program_source:245:71: note: expanded from macro '\ 
    3158 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3159 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3160 ^ 
    3161 program_source:198:12: note: expanded from macro '\ 
    3162 OVERLOAD_CLUSTERED_ARGS_1' 
    3163 return simd_##METAL_SUFFIX(x); \ 
    3164 ^ 
    3165 <scratch space>:119:1: note: expanded from here 
    3166 simd_product 
    3167 ^ 
    3168 program_source:395:1: error: implicit declaration of function 'quad_product'
    3169 is invalid in OpenCL 
    3170 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3171 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3172 ^ 
    3173 program_source:360:60: note: expanded from macro '\ 
    3174 DECLARE_I_REDUCTION_CLUSTERED' 
    3175 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3176 ^ 
    3177 program_source:246:59: note: expanded from macro '\ 
    3178 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3179 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3180 ^ 
    3181 program_source:196:12: note: expanded from macro '\ 
    3182 OVERLOAD_CLUSTERED_ARGS_1' 
    3183 return quad_##METAL_SUFFIX(x); \ 
    3184 ^ 
    3185 <scratch space>:120:1: note: expanded from here 
    3186 quad_product 
    3187 ^ 
    3188 program_source:395:1: error: implicit declaration of function 'simd_product'
    3189 is invalid in OpenCL 
    3190 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3191 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3192 ^ 
    3193 program_source:360:60: note: expanded from macro '\ 
    3194 DECLARE_I_REDUCTION_CLUSTERED' 
    3195 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3196 ^ 
    3197 program_source:246:59: note: expanded from macro '\ 
    3198 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3199 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3200 ^ 
    3201 program_source:198:12: note: expanded from macro '\ 
    3202 OVERLOAD_CLUSTERED_ARGS_1' 
    3203 return simd_##METAL_SUFFIX(x); \ 
    3204 ^ 
    3205 <scratch space>:121:1: note: expanded from here 
    3206 simd_product 
    3207 ^ 
    3208 program_source:395:1: error: implicit declaration of function 'quad_product'
    3209 is invalid in OpenCL 
    3210 program_source:385:52: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3211 DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3212 ^ 
    3213 program_source:363:60: note: expanded from macro '\ 
    3214 DECLARE_F_REDUCTION_CLUSTERED' 
    3215 #define DECLARE_F_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3216 ^ 
    3217 program_source:249:71: note: expanded from macro '\ 
    3218 BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1' 
    3219 #define BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3220 ^ 
    3221 program_source:196:12: note: expanded from macro '\ 
    3222 OVERLOAD_CLUSTERED_ARGS_1' 
    3223 return quad_##METAL_SUFFIX(x); \ 
    3224 ^ 
    3225 <scratch space>:123:1: note: expanded from here 
    3226 quad_product 
    3227 ^ 
    3228 program_source:395:1: error: implicit declaration of function 'simd_product'
    3229 is invalid in OpenCL 
    3230 program_source:385:52: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3231 DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3232 ^ 
    3233 program_source:363:60: note: expanded from macro '\ 
    3234 DECLARE_F_REDUCTION_CLUSTERED' 
    3235 #define DECLARE_F_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3236 ^ 
    3237 program_source:249:71: note: expanded from macro '\ 
    3238 BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1' 
    3239 #define BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3240 ^ 
    3241 program_source:198:12: note: expanded from macro '\ 
    3242 OVERLOAD_CLUSTERED_ARGS_1' 
    3243 return simd_##METAL_SUFFIX(x); \ 
    3244 ^ 
    3245 <scratch space>:124:1: note: expanded from here 
    3246 simd_product 
    3247 ^ 
    3248 program_source:396:1: error: implicit declaration of function 'quad_and' is
    3249 invalid in OpenCL 
    3250 DECLARE_I_REDUCTION_CLUSTERED(and, reduce_and) 
    3251 ^ 
    3252 program_source:360:60: note: expanded from macro
    3253 'DECLARE_I_REDUCTION_CLUSTERED' 
    3254 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3255 ^ 
    3256 program_source:245:71: note: expanded from macro '\ 
    3257 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3258 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3259 ^ 
    3260 program_source:196:12: note: expanded from macro '\ 
    3261 OVERLOAD_CLUSTERED_ARGS_1' 
    3262 return quad_##METAL_SUFFIX(x); \ 
    3263 ^ 
    3264 <scratch space>:126:1: note: expanded from here 
    3265 quad_and 
    3266 ^ 
    3267 program_source:396:1: error: implicit declaration of function 'simd_and' is
    3268 invalid in OpenCL 
    3269 program_source:360:60: note: expanded from macro
    3270 'DECLARE_I_REDUCTION_CLUSTERED' 
    3271 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3272 ^ 
    3273 program_source:245:71: note: expanded from macro '\ 
    3274 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3275 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3276 ^ 
    3277 program_source:198:12: note: expanded from macro '\ 
    3278 OVERLOAD_CLUSTERED_ARGS_1' 
    3279 return simd_##METAL_SUFFIX(x); \ 
    3280 ^ 
    3281 <scratch space>:127:1: note: expanded from here 
    3282 simd_and 
    3283 ^ 
    3284 program_source:396:1: error: implicit declaration of function 'quad_and' is
    3285 invalid in OpenCL 
    3286 program_source:360:60: note: expanded from macro
    3287 'DECLARE_I_REDUCTION_CLUSTERED' 
    3288 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3289 ^ 
    3290 program_source:246:59: note: expanded from macro '\ 
    3291 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3292 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3293 ^ 
    3294 program_source:196:12: note: expanded from macro '\ 
    3295 OVERLOAD_CLUSTERED_ARGS_1' 
    3296 return quad_##METAL_SUFFIX(x); \ 
    3297 ^ 
    3298 <scratch space>:128:1: note: expanded from here 
    3299 quad_and 
    3300 ^ 
    3301 program_source:396:1: error: implicit declaration of function 'simd_and' is
    3302 invalid in OpenCL 
    3303 program_source:360:60: note: expanded from macro
    3304 'DECLARE_I_REDUCTION_CLUSTERED' 
    3305 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3306 ^ 
    3307 program_source:246:59: note: expanded from macro '\ 
    3308 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3309 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3310 ^ 
    3311 program_source:198:12: note: expanded from macro '\ 
    3312 OVERLOAD_CLUSTERED_ARGS_1' 
    3313 return simd_##METAL_SUFFIX(x); \ 
    3314 ^ 
    3315 <scratch space>:129:1: note: expanded from here 
    3316 simd_and 
    3317 ^ 
    3318 program_source:397:1: error: implicit declaration of function 'quad_or' is
    3319 invalid in OpenCL 
    3320 DECLARE_I_REDUCTION_CLUSTERED(or, reduce_or) 
    3321 ^ 
    3322 program_source:360:60: note: expanded from macro
    3323 'DECLARE_I_REDUCTION_CLUSTERED' 
    3324 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3325 ^ 
    3326 program_source:245:71: note: expanded from macro '\ 
    3327 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3328 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3329 ^ 
    3330 program_source:196:12: note: expanded from macro '\ 
    3331 OVERLOAD_CLUSTERED_ARGS_1' 
    3332 return quad_##METAL_SUFFIX(x); \ 
    3333 ^ 
    3334 <scratch space>:131:1: note: expanded from here 
    3335 quad_or 
    3336 ^ 
    3337 program_source:397:1: error: implicit declaration of function 'simd_or' is
    3338 invalid in OpenCL 
    3339 program_source:360:60: note: expanded from macro
    3340 'DECLARE_I_REDUCTION_CLUSTERED' 
    3341 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3342 ^ 
    3343 program_source:245:71: note: expanded from macro '\ 
    3344 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3345 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3346 ^ 
    3347 program_source:198:12: note: expanded from macro '\ 
    3348 OVERLOAD_CLUSTERED_ARGS_1' 
    3349 return simd_##METAL_SUFFIX(x); \ 
    3350 ^ 
    3351 <scratch space>:132:1: note: expanded from here 
    3352 simd_or 
    3353 ^ 
    3354 program_source:397:1: error: implicit declaration of function 'quad_or' is
    3355 invalid in OpenCL 
    3356 program_source:360:60: note: expanded from macro
    3357 'DECLARE_I_REDUCTION_CLUSTERED' 
    3358 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3359 ^ 
    3360 program_source:246:59: note: expanded from macro '\ 
    3361 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3362 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3363 ^ 
    3364 program_source:196:12: note: expanded from macro '\ 
    3365 OVERLOAD_CLUSTERED_ARGS_1' 
    3366 return quad_##METAL_SUFFIX(x); \ 
    3367 ^ 
    3368 <scratch space>:133:1: note: expanded from here 
    3369 quad_or 
    3370 ^ 
    3371 program_source:397:1: error: implicit declaration of function 'simd_or' is
    3372 invalid in OpenCL 
    3373 program_source:360:60: note: expanded from macro
    3374 'DECLARE_I_REDUCTION_CLUSTERED' 
    3375 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3376 ^ 
    3377 program_source:246:59: note: expanded from macro '\ 
    3378 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3379 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3380 ^ 
    3381 program_source:198:12: note: expanded from macro '\ 
    3382 OVERLOAD_CLUSTERED_ARGS_1' 
    3383 return simd_##METAL_SUFFIX(x); \ 
    3384 ^ 
    3385 <scratch space>:134:1: note: expanded from here 
    3386 simd_or 
    3387 ^ 
    3388 program_source:398:1: error: implicit declaration of function 'quad_xor' is
    3389 invalid in OpenCL 
    3390 DECLARE_I_REDUCTION_CLUSTERED(xor, reduce_xor) 
    3391 ^ 
    3392 program_source:360:60: note: expanded from macro
    3393 'DECLARE_I_REDUCTION_CLUSTERED' 
    3394 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3395 ^ 
    3396 program_source:245:71: note: expanded from macro '\ 
    3397 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3398 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3399 ^ 
    3400 program_source:196:12: note: expanded from macro '\ 
    3401 OVERLOAD_CLUSTERED_ARGS_1' 
    3402 return quad_##METAL_SUFFIX(x); \ 
    3403 ^ 
    3404 <scratch space>:136:1: note: expanded from here 
    3405 quad_xor 
    3406 ^ 
    3407 program_source:398:1: error: implicit declaration of function 'simd_xor' is
    3408 invalid in OpenCL 
    3409 program_source:360:60: note: expanded from macro
    3410 'DECLARE_I_REDUCTION_CLUSTERED' 
    3411 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3412 ^ 
    3413 program_source:245:71: note: expanded from macro '\ 
    3414 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3415 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3416 ^ 
    3417 program_source:198:12: note: expanded from macro '\ 
    3418 OVERLOAD_CLUSTERED_ARGS_1' 
    3419 return simd_##METAL_SUFFIX(x); \ 
    3420 ^ 
    3421 <scratch space>:137:1: note: expanded from here 
    3422 simd_xor 
    3423 ^ 
    3424 program_source:398:1: error: implicit declaration of function 'quad_xor' is
    3425 invalid in OpenCL 
    3426 program_source:360:60: note: expanded from macro
    3427 'DECLARE_I_REDUCTION_CLUSTERED' 
    3428 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3429 ^ 
    3430 program_source:246:59: note: expanded from macro '\ 
    3431 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3432 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3433 ^ 
    3434 program_source:196:12: note: expanded from macro '\ 
    3435 OVERLOAD_CLUSTERED_ARGS_1' 
    3436 return quad_##METAL_SUFFIX(x); \ 
    3437 ^ 
    3438 <scratch space>:138:1: note: expanded from here 
    3439 quad_xor 
    3440 ^ 
    3441 program_source:398:1: error: implicit declaration of function 'simd_xor' is
    3442 invalid in OpenCL 
    3443 program_source:360:60: note: expanded from macro
    3444 'DECLARE_I_REDUCTION_CLUSTERED' 
    3445 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3446 ^ 
    3447 program_source:246:59: note: expanded from macro '\ 
    3448 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3449 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3450 ^ 
    3451 program_source:198:12: note: expanded from macro '\ 
    3452 OVERLOAD_CLUSTERED_ARGS_1' 
    3453 return simd_##METAL_SUFFIX(x); \ 
    3454 ^ 
    3455 <scratch space>:139:1: note: expanded from here 
    3456 simd_xor 
    3457 ^ 
    3458 program_source:399:1: error: implicit declaration of function 'simd_and' is
    3459 invalid in OpenCL 
    3460 DECLARE_B_REDUCTION_CLUSTERED(and) 
    3461 ^ 
    3462 program_source:369:8: note: expanded from macro
    3463 'DECLARE_B_REDUCTION_CLUSTERED' 
    3464 return simd_##OP(select(0, 1, predicate != 0)); \ 
    3465 ^ 
    3466 <scratch space>:141:1: note: expanded from here 
    3467 simd_and 
    3468 ^ 
    3469 program_source:399:1: error: implicit declaration of function 'quad_and' is
    3470 invalid in OpenCL 
    3471 program_source:378:12: note: expanded from macro
    3472 'DECLARE_B_REDUCTION_CLUSTERED' 
    3473 return quad_##OP(x); \ 
    3474 ^ 
    3475 <scratch space>:143:1: note: expanded from here 
    3476 quad_and 
    3477 ^ 
    3478 program_source:399:1: error: implicit declaration of function 'simd_and' is
    3479 invalid in OpenCL 
    3480 program_source:380:12: note: expanded from macro
    3481 'DECLARE_B_REDUCTION_CLUSTERED' 
    3482 return simd_##OP(x); \ 
    3483 ^ 
    3484 <scratch space>:144:1: note: expanded from here 
    3485 simd_and 
    3486 ^ 
    3487 program_source:400:1: error: implicit declaration of function 'simd_or' is
    3488 invalid in OpenCL 
    3489 DECLARE_B_REDUCTION_CLUSTERED(or) 
    3490 ^ 
    3491 program_source:369:8: note: expanded from macro
    3492 'DECLARE_B_REDUCTION_CLUSTERED' 
    3493 return simd_##OP(select(0, 1, predicate != 0)); \ 
    3494 ^ 
    3495 <scratch space>:146:1: note: expanded from here 
    3496 simd_or 
    3497 ^ 
    3498 program_source:400:1: error: implicit declaration of function 'quad_or' is
    3499 invalid in OpenCL 
    3500 program_source:378:12: note: expanded from macro
    3501 'DECLARE_B_REDUCTION_CLUSTERED' 
    3502 return quad_##OP(x); \ 
    3503 ^ 
    3504 <scratch space>:148:1: note: expanded from here 
    3505 quad_or 
    3506 ^ 
    3507 program_source:400:1: error: implicit declaration of function 'simd_or' is
    3508 invalid in OpenCL 
    3509 program_source:380:12: note: expanded from macro
    3510 'DECLARE_B_REDUCTION_CLUSTERED' 
    3511 return simd_##OP(x); \ 
    3512 ^ 
    3513 <scratch space>:149:1: note: expanded from here 
    3514 simd_or 
    3515 ^ 
    3516 program_source:401:1: error: implicit declaration of function 'simd_xor' is
    3517 invalid in OpenCL 
    3518 DECLARE_B_REDUCTION_CLUSTERED(xor) 
    3519 ^ 
    3520 program_source:369:8: note: expanded from macro
    3521 'DECLARE_B_REDUCTION_CLUSTERED' 
    3522 return simd_##OP(select(0, 1, predicate != 0)); \ 
    3523 ^ 
    3524 <scratch space>:151:1: note: expanded from here 
    3525 simd_xor 
    3526 ^ 
    3527 program_source:401:1: error: implicit declaration of function 'quad_xor' is
    3528 invalid in OpenCL 
    3529 program_source:378:12: note: expanded from macro
    3530 'DECLARE_B_REDUCTION_CLUSTERED' 
    3531 return quad_##OP(x); \ 
    3532 ^ 
    3533 <scratch space>:153:1: note: expanded from here 
    3534 quad_xor 
    3535 ^ 
    3536 program_source:401:1: error: implicit declaration of function 'simd_xor' is
    3537 invalid in OpenCL 
    3538 program_source:380:12: note: expanded from macro
    3539 'DECLARE_B_REDUCTION_CLUSTERED' 
    3540 return simd_##OP(x); \ 
    3541 ^ 
    3542 <scratch space>:154:1: note: expanded from here 
    3543 simd_xor 
    3544 ^ 
    3545 program_source:403:1: error: implicit declaration of function
    3546 'quad_shuffle_rotate_down' is invalid in OpenCL 
    3547 DECLARE_SHUFFLE_CLUSTERED(shuffle_rotate_down, rotate) 
    3548 ^ 
    3549 program_source:388:56: note: expanded from macro 'DECLARE_SHUFFLE_CLUSTERED' 
    3550 #define DECLARE_SHUFFLE_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3551 ^ 
    3552 program_source:252:69: note: expanded from macro '\ 
    3553 BRIDGE_FUNCTION_CLUSTERED_ARGS_2' 
    3554 #define BRIDGE_FUNCTION_CLUSTERED_ARGS_2(METAL_SUFFIX, OPENCL_EXPR) \ 
    3555 ^ 
    3556 program_source:208:12: note: expanded from macro '\ 
    3557 OVERLOAD_CLUSTERED_ARGS_2' 
    3558 return quad_##METAL_SUFFIX(x, ushort(delta)); \ 
    3559 ^ 
    3560 <scratch space>:156:1: note: expanded from here 
    3561 quad_shuffle_rotate_down 
    3562 ^ 
    3563 program_source:403:1: error: implicit declaration of function
    3564 'simd_shuffle_rotate_down' is invalid in OpenCL 
    3565 program_source:388:56: note: expanded from macro 'DECLARE_SHUFFLE_CLUSTERED' 
    3566 #define DECLARE_SHUFFLE_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3567 ^ 
    3568 program_source:252:69: note: expanded from macro '\ 
    3569 BRIDGE_FUNCTION_CLUSTERED_ARGS_2' 
    3570 #define BRIDGE_FUNCTION_CLUSTERED_ARGS_2(METAL_SUFFIX, OPENCL_EXPR) \ 
    3571 ^ 
    3572 program_source:210:12: note: expanded from macro '\ 
    3573 OVERLOAD_CLUSTERED_ARGS_2' 
    3574 return simd_##METAL_SUFFIX(x, ushort(delta)); \ 
    3575 ^ 
    3576 <scratch space>:157:1: note: expanded from here 
    3577 simd_shuffle_rotate_down 
    3578 ^ 
    3579 program_source:403:1: error: implicit declaration of function
    3580 'quad_shuffle_rotate_down' is invalid in OpenCL 
    3581 program_source:388:56: note: expanded from macro 'DECLARE_SHUFFLE_CLUSTERED' 
    3582 #define DECLARE_SHUFFLE_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3583 ^ 
    3584 program_source:253:59: note: expanded from macro '\ 
    3585 BRIDGE_FUNCTION_CLUSTERED_ARGS_2' 
    3586 OVERLOAD_CLUSTERED_ARGS_2(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3587 ^ 
    3588 program_source:208:12: note: expanded from macro '\ 
    3589 OVERLOAD_CLUSTERED_ARGS_2' 
    3590 return quad_##METAL_SUFFIX(x, ushort(delta)); \ 
    3591 ^ 
    3592 <scratch space>:158:1: note: expanded from here 
    3593 quad_shuffle_rotate_down 
    3594 ^ 
    3595 program_source:403:1: error: implicit declaration of function
    3596 'simd_shuffle_rotate_down' is invalid in OpenCL 
    3597 program_source:388:56: note: expanded from macro 'DECLARE_SHUFFLE_CLUSTERED' 
    3598 #define DECLARE_SHUFFLE_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3599 ^ 
    3600 program_source:253:59: note: expanded from macro '\ 
    3601 BRIDGE_FUNCTION_CLUSTERED_ARGS_2' 
    3602 OVERLOAD_CLUSTERED_ARGS_2(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3603 ^ 
    3604 program_source:210:12: note: expanded from macro '\ 
    3605 OVERLOAD_CLUSTERED_ARGS_2' 
    3606 return simd_##METAL_SUFFIX(x, ushort(delta)); \ 
    3607 ^ 
    3608 <scratch space>:159:1: note: expanded from here 
    3609 simd_shuffle_rotate_down 
    3610 ^ 
    3611 program_source:403:1: error: implicit declaration of function
    3612 'quad_shuffle_rotate_down' is invalid in OpenCL 
    3613 program_source:388:56: note: expanded from macro 'DECLARE_SHUFFLE_CLUSTERED' 
    3614 #define DECLARE_SHUFFLE_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3615 ^ 
    3616 program_source:254:60: note: expanded from macro '\ 
    3617 BRIDGE_FUNCTION_CLUSTERED_ARGS_2' 
    3618 OVERLOAD_CLUSTERED_ARGS_2(METAL_SUFFIX, OPENCL_EXPR, uint) \ 
    3619 ^ 
    3620 program_source:208:12: note: expanded from macro '\ 
    3621 OVERLOAD_CLUSTERED_ARGS_2' 
    3622 return quad_##METAL_SUFFIX(x, ushort(delta)); \ 
    3623 ^ 
    3624 <scratch space>:160:1: note: expanded from here 
    3625 quad_shuffle_rotate_down 
    3626 ^ 
    3627 program_source:403:1: error: implicit declaration of function
    3628 'simd_shuffle_rotate_down' is invalid in OpenCL 
    3629 program_source:388:56: note: expanded from macro 'DECLARE_SHUFFLE_CLUSTERED' 
    3630 #define DECLARE_SHUFFLE_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3631 ^ 
    3632 program_source:254:60: note: expanded from macro '\ 
    3633 BRIDGE_FUNCTION_CLUSTERED_ARGS_2' 
    3634 OVERLOAD_CLUSTERED_ARGS_2(METAL_SUFFIX, OPENCL_EXPR, uint) \ 
    3635 ^ 
    3636 program_source:210:12: note: expanded from macro '\ 
    3637 OVERLOAD_CLUSTERED_ARGS_2' 
    3638 return simd_##METAL_SUFFIX(x, ushort(delta)); \ 
    3639 ^ 
    3640 <scratch space>:161:1: note: expanded from here 
    3641 simd_shuffle_rotate_down 
    3642 ^ 
    3643 program_source:409:1: error: illegal string literal in 'asm' 
    3644 EXPOSE_BALLOT(simd_is_first, , bool, ) 
    3645 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3646 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    3647 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    3648 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3649 program_source:410:1: error: illegal string literal in 'asm' 
    3650 EXPOSE_BALLOT(simd_all, bool expr, bool, ) 
    3651 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3652 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    3653 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    3654 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3655 program_source:411:1: error: illegal string literal in 'asm' 
    3656 EXPOSE_BALLOT(simd_any, bool expr, bool, ) 
    3657 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3658 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    3659 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    3660 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3661 program_source:412:1: error: illegal string literal in 'asm' 
    3662 EXPOSE_BALLOT(simd_ballot, bool expr, ulong, .i64) 
    3663 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3664 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    3665 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    3666 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3667 program_source:413:1: error: illegal string literal in 'asm' 
    3668 EXPOSE_BALLOT(simd_active_threads_mask, , ulong, .i64) 
    3669 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3670 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    3671 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    3672 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3673 program_source:415:1: error: illegal string literal in 'asm' 
    3674 EXPOSE_BALLOT(quad_is_first, , bool, ) 
    3675 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3676 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    3677 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    3678 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3679 program_source:416:1: error: illegal string literal in 'asm' 
    3680 EXPOSE_BALLOT(quad_all, bool expr, bool, ) 
    3681 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3682 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    3683 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    3684 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3685 program_source:417:1: error: illegal string literal in 'asm' 
    3686 EXPOSE_BALLOT(quad_any, bool expr, bool, ) 
    3687 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3688 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    3689 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    3690 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3691 program_source:418:1: error: illegal string literal in 'asm' 
    3692 EXPOSE_BALLOT(quad_ballot, bool expr, ushort, ) 
    3693 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3694 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    3695 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    3696 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3697 program_source:419:1: error: illegal string literal in 'asm' 
    3698 EXPOSE_BALLOT(quad_active_threads_mask, , ushort, ) 
    3699 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3700 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    3701 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    3702 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3703 program_source:422:23: error: implicit declaration of function 'simd_is_first'
    3704 is invalid in OpenCL 
    3705 return select(0, 1, simd_is_first()); 
    3706 ^ 
    3707 program_source:426:23: error: implicit declaration of function 'simd_all' is
    3708 invalid in OpenCL 
    3709 return select(0, 1, simd_all(predicate != 0)); 
    3710 ^ 
    3711 program_source:430:23: error: implicit declaration of function 'simd_any' is
    3712 invalid in OpenCL 
    3713 return select(0, 1, simd_any(predicate != 0)); 
    3714 ^ 
    3715 program_source:434:23: error: implicit declaration of function 'simd_all' is
    3716 invalid in OpenCL 
    3717 return select(0, 1, simd_all(predicate != 0)); 
    3718 ^ 
    3719 program_source:438:23: error: implicit declaration of function 'simd_any' is
    3720 invalid in OpenCL 
    3721 return select(0, 1, simd_any(predicate != 0)); 
    3722 ^ 
    3723 program_source:443:14: error: implicit declaration of function 'simd_ballot'
    3724 is invalid in OpenCL 
    3725 output.x = simd_ballot(predicate != 0); 
    3726 ^ 
    3727 program_source:480:9: error: illegal string literal in 'asm' 
    3728 __asm("air.simdgroup.barrier"); 
    3729 ^~~~~~~~~~~~~~~~~~~~~~~ 
    3730 program_source:487:3: error: implicit declaration of function
    3731 '__metal_simdgroup_barrier' is invalid in OpenCL 
    3732 __metal_simdgroup_barrier(flags, 1); 
    3733 ^ 
    3734 program_source:487:3: note: did you mean 'simdgroup_barrier'? 
    3735 program_source:483:6: note: 'simdgroup_barrier' declared here 
    3736 void simdgroup_barrier(int flags) { 
    3737 ^ 
    3738 program_source:681:9: error: implicit declaration of function 'simd_is_first'
    3739 is invalid in OpenCL 
    3740 if (simd_is_first()) { 
    3741 ^ 
    3742 program_source:673:45: warning: comparison of integers of different signs:
    3743 '__private unsigned int' and '__private int' [-Wsign-compare] 
    3744 for (unsigned int index = thread; index < bufferSize; index +=
    3745 get_local_size(0)) 
    3746 ~~~~~ ^ ~~~~~~~~~~ 
    3747 program_source:685:45: warning: comparison of integers of different signs:
    3748 'unsigned int' and '__private int' [-Wsign-compare] 
    3749 if (thread%(32*2) == 0 && thread+32 < workGroupSize) 
    3750 ~~~~~~~~~ ^ ~~~~~~~~~~~~~ 
    3751 program_source:689:45: warning: comparison of integers of different signs:
    3752 'unsigned int' and '__private int' [-Wsign-compare] 
    3753 if (thread%(64*2) == 0 && thread+64 < workGroupSize) 
    3754 ~~~~~~~~~ ^ ~~~~~~~~~~~~~ 
    3755 program_source:693:47: warning: comparison of integers of different signs:
    3756 'unsigned int' and '__private int' [-Wsign-compare] 
    3757 if (thread%(128*2) == 0 && thread+128 < workGroupSize) 
    3758 ~~~~~~~~~~ ^ ~~~~~~~~~~~~~ 
    3759  
    3760  
    3761 openmm.OpenMMException: Error compiling kernel: program_source:279:1: error:
    3762 illegal string literal in 'asm' 
    3763 DECLARE_REDUCTION_BASE(sum) 
    3764 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3765 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    3766 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    3767 ^ 
    3768 program_source:263:44: note: expanded from macro '\ 
    3769 DECLARE_I_REDUCTION_BASE' 
    3770 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    3771 ^ 
    3772 program_source:217:40: note: expanded from macro '\ 
    3773 EXPOSE_FUNCTION_I_ARGS_1' 
    3774 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    3775 ^ 
    3776 program_source:164:9: note: expanded from macro '\ 
    3777 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    3778 __asm("air." #EXPR "." #AIR_TYPE); \ 
    3779 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3780 program_source:279:1: error: illegal string literal in 'asm' 
    3781 DECLARE_REDUCTION_BASE(sum) 
    3782 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3783 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    3784 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    3785 ^ 
    3786 program_source:263:44: note: expanded from macro '\ 
    3787 DECLARE_I_REDUCTION_BASE' 
    3788 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    3789 ^ 
    3790 program_source:218:51: note: expanded from macro '\ 
    3791 EXPOSE_FUNCTION_I_ARGS_1' 
    3792 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    3793 ^ 
    3794 program_source:164:9: note: expanded from macro '\ 
    3795 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    3796 __asm("air." #EXPR "." #AIR_TYPE); \ 
    3797 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3798 program_source:279:1: error: illegal string literal in 'asm' 
    3799 DECLARE_REDUCTION_BASE(sum) 
    3800 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3801 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    3802 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    3803 ^ 
    3804 program_source:264:43: note: expanded from macro '\ 
    3805 DECLARE_I_REDUCTION_BASE' 
    3806 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    3807 ^ 
    3808 program_source:217:40: note: expanded from macro '\ 
    3809 EXPOSE_FUNCTION_I_ARGS_1' 
    3810 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    3811 ^ 
    3812 program_source:164:9: note: expanded from macro '\ 
    3813 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    3814 __asm("air." #EXPR "." #AIR_TYPE); \ 
    3815 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3816 program_source:279:1: error: illegal string literal in 'asm' 
    3817 DECLARE_REDUCTION_BASE(sum) 
    3818 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3819 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    3820 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    3821 ^ 
    3822 program_source:264:43: note: expanded from macro '\ 
    3823 DECLARE_I_REDUCTION_BASE' 
    3824 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    3825 ^ 
    3826 program_source:218:51: note: expanded from macro '\ 
    3827 EXPOSE_FUNCTION_I_ARGS_1' 
    3828 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    3829 ^ 
    3830 program_source:164:9: note: expanded from macro '\ 
    3831 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    3832 __asm("air." #EXPR "." #AIR_TYPE); \ 
    3833 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3834 program_source:279:1: error: illegal string literal in 'asm' 
    3835 DECLARE_REDUCTION_BASE(sum) 
    3836 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3837 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    3838 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    3839 ^ 
    3840 program_source:267:44: note: expanded from macro '\ 
    3841 DECLARE_F_REDUCTION_BASE' 
    3842 #define DECLARE_F_REDUCTION_BASE(METAL_OP) \ 
    3843 ^ 
    3844 program_source:221:40: note: expanded from macro '\ 
    3845 EXPOSE_FUNCTION_F_ARGS_1' 
    3846 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    3847 ^ 
    3848 program_source:164:9: note: expanded from macro '\ 
    3849 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    3850 __asm("air." #EXPR "." #AIR_TYPE); \ 
    3851 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3852 program_source:279:1: error: illegal string literal in 'asm' 
    3853 DECLARE_REDUCTION_BASE(sum) 
    3854 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3855 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    3856 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    3857 ^ 
    3858 program_source:268:43: note: expanded from macro '\ 
    3859 DECLARE_F_REDUCTION_BASE' 
    3860 EXPOSE_FUNCTION_F_ARGS_1(quad_##METAL_OP) \ 
    3861 ^ 
    3862 program_source:221:40: note: expanded from macro '\ 
    3863 EXPOSE_FUNCTION_F_ARGS_1' 
    3864 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    3865 ^ 
    3866 program_source:164:9: note: expanded from macro '\ 
    3867 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    3868 __asm("air." #EXPR "." #AIR_TYPE); \ 
    3869 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3870 program_source:280:1: error: illegal string literal in 'asm' 
    3871 DECLARE_REDUCTION_BASE(prefix_inclusive_sum) 
    3872 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3873 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    3874 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    3875 ^ 
    3876 program_source:263:44: note: expanded from macro '\ 
    3877 DECLARE_I_REDUCTION_BASE' 
    3878 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    3879 ^ 
    3880 program_source:217:40: note: expanded from macro '\ 
    3881 EXPOSE_FUNCTION_I_ARGS_1' 
    3882 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    3883 ^ 
    3884 program_source:164:9: note: expanded from macro '\ 
    3885 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    3886 __asm("air." #EXPR "." #AIR_TYPE); \ 
    3887 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3888 program_source:280:1: error: illegal string literal in 'asm' 
    3889 DECLARE_REDUCTION_BASE(prefix_inclusive_sum) 
    3890 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3891 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    3892 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    3893 ^ 
    3894 program_source:263:44: note: expanded from macro '\ 
    3895 DECLARE_I_REDUCTION_BASE' 
    3896 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    3897 ^ 
    3898 program_source:218:51: note: expanded from macro '\ 
    3899 EXPOSE_FUNCTION_I_ARGS_1' 
    3900 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    3901 ^ 
    3902 program_source:164:9: note: expanded from macro '\ 
    3903 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    3904 __asm("air." #EXPR "." #AIR_TYPE); \ 
    3905 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3906 program_source:280:1: error: illegal string literal in 'asm' 
    3907 DECLARE_REDUCTION_BASE(prefix_inclusive_sum) 
    3908 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3909 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    3910 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    3911 ^ 
    3912 program_source:264:43: note: expanded from macro '\ 
    3913 DECLARE_I_REDUCTION_BASE' 
    3914 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    3915 ^ 
    3916 program_source:217:40: note: expanded from macro '\ 
    3917 EXPOSE_FUNCTION_I_ARGS_1' 
    3918 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    3919 ^ 
    3920 program_source:164:9: note: expanded from macro '\ 
    3921 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    3922 __asm("air." #EXPR "." #AIR_TYPE); \ 
    3923 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3924 program_source:280:1: error: illegal string literal in 'asm' 
    3925 DECLARE_REDUCTION_BASE(prefix_inclusive_sum) 
    3926 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3927 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    3928 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    3929 ^ 
    3930 program_source:264:43: note: expanded from macro '\ 
    3931 DECLARE_I_REDUCTION_BASE' 
    3932 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    3933 ^ 
    3934 program_source:218:51: note: expanded from macro '\ 
    3935 EXPOSE_FUNCTION_I_ARGS_1' 
    3936 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    3937 ^ 
    3938 program_source:164:9: note: expanded from macro '\ 
    3939 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    3940 __asm("air." #EXPR "." #AIR_TYPE); \ 
    3941 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3942 program_source:280:1: error: illegal string literal in 'asm' 
    3943 DECLARE_REDUCTION_BASE(prefix_inclusive_sum) 
    3944 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3945 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    3946 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    3947 ^ 
    3948 program_source:267:44: note: expanded from macro '\ 
    3949 DECLARE_F_REDUCTION_BASE' 
    3950 #define DECLARE_F_REDUCTION_BASE(METAL_OP) \ 
    3951 ^ 
    3952 program_source:221:40: note: expanded from macro '\ 
    3953 EXPOSE_FUNCTION_F_ARGS_1' 
    3954 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    3955 ^ 
    3956 program_source:164:9: note: expanded from macro '\ 
    3957 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    3958 __asm("air." #EXPR "." #AIR_TYPE); \ 
    3959 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3960 program_source:280:1: error: illegal string literal in 'asm' 
    3961 DECLARE_REDUCTION_BASE(prefix_inclusive_sum) 
    3962 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3963 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    3964 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    3965 ^ 
    3966 program_source:268:43: note: expanded from macro '\ 
    3967 DECLARE_F_REDUCTION_BASE' 
    3968 EXPOSE_FUNCTION_F_ARGS_1(quad_##METAL_OP) \ 
    3969 ^ 
    3970 program_source:221:40: note: expanded from macro '\ 
    3971 EXPOSE_FUNCTION_F_ARGS_1' 
    3972 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    3973 ^ 
    3974 program_source:164:9: note: expanded from macro '\ 
    3975 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    3976 __asm("air." #EXPR "." #AIR_TYPE); \ 
    3977 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3978 program_source:281:1: error: illegal string literal in 'asm' 
    3979 DECLARE_REDUCTION_BASE(prefix_exclusive_sum) 
    3980 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3981 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    3982 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    3983 ^ 
    3984 program_source:263:44: note: expanded from macro '\ 
    3985 DECLARE_I_REDUCTION_BASE' 
    3986 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    3987 ^ 
    3988 program_source:217:40: note: expanded from macro '\ 
    3989 EXPOSE_FUNCTION_I_ARGS_1' 
    3990 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    3991 ^ 
    3992 program_source:164:9: note: expanded from macro '\ 
    3993 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    3994 __asm("air." #EXPR "." #AIR_TYPE); \ 
    3995 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3996 program_source:281:1: error: illegal string literal in 'asm' 
    3997 DECLARE_REDUCTION_BASE(prefix_exclusive_sum) 
    3998 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3999 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4000 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4001 ^ 
    4002 program_source:263:44: note: expanded from macro '\ 
    4003 DECLARE_I_REDUCTION_BASE' 
    4004 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4005 ^ 
    4006 program_source:218:51: note: expanded from macro '\ 
    4007 EXPOSE_FUNCTION_I_ARGS_1' 
    4008 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4009 ^ 
    4010 program_source:164:9: note: expanded from macro '\ 
    4011 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4012 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4013 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4014 program_source:281:1: error: illegal string literal in 'asm' 
    4015 DECLARE_REDUCTION_BASE(prefix_exclusive_sum) 
    4016 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4017 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4018 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4019 ^ 
    4020 program_source:264:43: note: expanded from macro '\ 
    4021 DECLARE_I_REDUCTION_BASE' 
    4022 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4023 ^ 
    4024 program_source:217:40: note: expanded from macro '\ 
    4025 EXPOSE_FUNCTION_I_ARGS_1' 
    4026 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4027 ^ 
    4028 program_source:164:9: note: expanded from macro '\ 
    4029 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4030 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4031 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4032 program_source:281:1: error: illegal string literal in 'asm' 
    4033 DECLARE_REDUCTION_BASE(prefix_exclusive_sum) 
    4034 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4035 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4036 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4037 ^ 
    4038 program_source:264:43: note: expanded from macro '\ 
    4039 DECLARE_I_REDUCTION_BASE' 
    4040 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4041 ^ 
    4042 program_source:218:51: note: expanded from macro '\ 
    4043 EXPOSE_FUNCTION_I_ARGS_1' 
    4044 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4045 ^ 
    4046 program_source:164:9: note: expanded from macro '\ 
    4047 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4048 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4049 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4050 program_source:281:1: error: illegal string literal in 'asm' 
    4051 DECLARE_REDUCTION_BASE(prefix_exclusive_sum) 
    4052 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4053 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4054 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4055 ^ 
    4056 program_source:267:44: note: expanded from macro '\ 
    4057 DECLARE_F_REDUCTION_BASE' 
    4058 #define DECLARE_F_REDUCTION_BASE(METAL_OP) \ 
    4059 ^ 
    4060 program_source:221:40: note: expanded from macro '\ 
    4061 EXPOSE_FUNCTION_F_ARGS_1' 
    4062 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4063 ^ 
    4064 program_source:164:9: note: expanded from macro '\ 
    4065 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4066 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4067 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4068 program_source:281:1: error: illegal string literal in 'asm' 
    4069 DECLARE_REDUCTION_BASE(prefix_exclusive_sum) 
    4070 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4071 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4072 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4073 ^ 
    4074 program_source:268:43: note: expanded from macro '\ 
    4075 DECLARE_F_REDUCTION_BASE' 
    4076 EXPOSE_FUNCTION_F_ARGS_1(quad_##METAL_OP) \ 
    4077 ^ 
    4078 program_source:221:40: note: expanded from macro '\ 
    4079 EXPOSE_FUNCTION_F_ARGS_1' 
    4080 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4081 ^ 
    4082 program_source:164:9: note: expanded from macro '\ 
    4083 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4084 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4085 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4086 program_source:282:1: error: illegal string literal in 'asm' 
    4087 DECLARE_REDUCTION_BASE(min) 
    4088 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4089 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4090 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4091 ^ 
    4092 program_source:263:44: note: expanded from macro '\ 
    4093 DECLARE_I_REDUCTION_BASE' 
    4094 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4095 ^ 
    4096 program_source:217:40: note: expanded from macro '\ 
    4097 EXPOSE_FUNCTION_I_ARGS_1' 
    4098 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4099 ^ 
    4100 program_source:164:9: note: expanded from macro '\ 
    4101 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4102 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4103 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4104 program_source:282:1: error: illegal string literal in 'asm' 
    4105 DECLARE_REDUCTION_BASE(min) 
    4106 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4107 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4108 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4109 ^ 
    4110 program_source:263:44: note: expanded from macro '\ 
    4111 DECLARE_I_REDUCTION_BASE' 
    4112 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4113 ^ 
    4114 program_source:218:51: note: expanded from macro '\ 
    4115 EXPOSE_FUNCTION_I_ARGS_1' 
    4116 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4117 ^ 
    4118 program_source:164:9: note: expanded from macro '\ 
    4119 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4120 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4121 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4122 program_source:282:1: error: illegal string literal in 'asm' 
    4123 DECLARE_REDUCTION_BASE(min) 
    4124 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4125 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4126 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4127 ^ 
    4128 program_source:264:43: note: expanded from macro '\ 
    4129 DECLARE_I_REDUCTION_BASE' 
    4130 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4131 ^ 
    4132 program_source:217:40: note: expanded from macro '\ 
    4133 EXPOSE_FUNCTION_I_ARGS_1' 
    4134 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4135 ^ 
    4136 program_source:164:9: note: expanded from macro '\ 
    4137 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4138 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4139 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4140 program_source:282:1: error: illegal string literal in 'asm' 
    4141 DECLARE_REDUCTION_BASE(min) 
    4142 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4143 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4144 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4145 ^ 
    4146 program_source:264:43: note: expanded from macro '\ 
    4147 DECLARE_I_REDUCTION_BASE' 
    4148 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4149 ^ 
    4150 program_source:218:51: note: expanded from macro '\ 
    4151 EXPOSE_FUNCTION_I_ARGS_1' 
    4152 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4153 ^ 
    4154 program_source:164:9: note: expanded from macro '\ 
    4155 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4156 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4157 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4158 program_source:282:1: error: illegal string literal in 'asm' 
    4159 DECLARE_REDUCTION_BASE(min) 
    4160 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4161 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4162 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4163 ^ 
    4164 program_source:267:44: note: expanded from macro '\ 
    4165 DECLARE_F_REDUCTION_BASE' 
    4166 #define DECLARE_F_REDUCTION_BASE(METAL_OP) \ 
    4167 ^ 
    4168 program_source:221:40: note: expanded from macro '\ 
    4169 EXPOSE_FUNCTION_F_ARGS_1' 
    4170 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4171 ^ 
    4172 program_source:164:9: note: expanded from macro '\ 
    4173 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4174 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4175 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4176 program_source:282:1: error: illegal string literal in 'asm' 
    4177 DECLARE_REDUCTION_BASE(min) 
    4178 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4179 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4180 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4181 ^ 
    4182 program_source:268:43: note: expanded from macro '\ 
    4183 DECLARE_F_REDUCTION_BASE' 
    4184 EXPOSE_FUNCTION_F_ARGS_1(quad_##METAL_OP) \ 
    4185 ^ 
    4186 program_source:221:40: note: expanded from macro '\ 
    4187 EXPOSE_FUNCTION_F_ARGS_1' 
    4188 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4189 ^ 
    4190 program_source:164:9: note: expanded from macro '\ 
    4191 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4192 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4193 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4194 program_source:283:1: error: illegal string literal in 'asm' 
    4195 DECLARE_REDUCTION_BASE(max) 
    4196 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4197 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4198 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4199 ^ 
    4200 program_source:263:44: note: expanded from macro '\ 
    4201 DECLARE_I_REDUCTION_BASE' 
    4202 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4203 ^ 
    4204 program_source:217:40: note: expanded from macro '\ 
    4205 EXPOSE_FUNCTION_I_ARGS_1' 
    4206 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4207 ^ 
    4208 program_source:164:9: note: expanded from macro '\ 
    4209 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4210 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4211 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4212 program_source:283:1: error: illegal string literal in 'asm' 
    4213 DECLARE_REDUCTION_BASE(max) 
    4214 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4215 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4216 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4217 ^ 
    4218 program_source:263:44: note: expanded from macro '\ 
    4219 DECLARE_I_REDUCTION_BASE' 
    4220 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4221 ^ 
    4222 program_source:218:51: note: expanded from macro '\ 
    4223 EXPOSE_FUNCTION_I_ARGS_1' 
    4224 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4225 ^ 
    4226 program_source:164:9: note: expanded from macro '\ 
    4227 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4228 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4229 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4230 program_source:283:1: error: illegal string literal in 'asm' 
    4231 DECLARE_REDUCTION_BASE(max) 
    4232 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4233 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4234 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4235 ^ 
    4236 program_source:264:43: note: expanded from macro '\ 
    4237 DECLARE_I_REDUCTION_BASE' 
    4238 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4239 ^ 
    4240 program_source:217:40: note: expanded from macro '\ 
    4241 EXPOSE_FUNCTION_I_ARGS_1' 
    4242 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4243 ^ 
    4244 program_source:164:9: note: expanded from macro '\ 
    4245 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4246 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4247 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4248 program_source:283:1: error: illegal string literal in 'asm' 
    4249 DECLARE_REDUCTION_BASE(max) 
    4250 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4251 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4252 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4253 ^ 
    4254 program_source:264:43: note: expanded from macro '\ 
    4255 DECLARE_I_REDUCTION_BASE' 
    4256 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4257 ^ 
    4258 program_source:218:51: note: expanded from macro '\ 
    4259 EXPOSE_FUNCTION_I_ARGS_1' 
    4260 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4261 ^ 
    4262 program_source:164:9: note: expanded from macro '\ 
    4263 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4264 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4265 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4266 program_source:283:1: error: illegal string literal in 'asm' 
    4267 DECLARE_REDUCTION_BASE(max) 
    4268 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4269 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4270 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4271 ^ 
    4272 program_source:267:44: note: expanded from macro '\ 
    4273 DECLARE_F_REDUCTION_BASE' 
    4274 #define DECLARE_F_REDUCTION_BASE(METAL_OP) \ 
    4275 ^ 
    4276 program_source:221:40: note: expanded from macro '\ 
    4277 EXPOSE_FUNCTION_F_ARGS_1' 
    4278 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4279 ^ 
    4280 program_source:164:9: note: expanded from macro '\ 
    4281 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4282 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4283 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4284 program_source:283:1: error: illegal string literal in 'asm' 
    4285 DECLARE_REDUCTION_BASE(max) 
    4286 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4287 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4288 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4289 ^ 
    4290 program_source:268:43: note: expanded from macro '\ 
    4291 DECLARE_F_REDUCTION_BASE' 
    4292 EXPOSE_FUNCTION_F_ARGS_1(quad_##METAL_OP) \ 
    4293 ^ 
    4294 program_source:221:40: note: expanded from macro '\ 
    4295 EXPOSE_FUNCTION_F_ARGS_1' 
    4296 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4297 ^ 
    4298 program_source:164:9: note: expanded from macro '\ 
    4299 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4300 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4301 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4302 program_source:285:1: error: illegal string literal in 'asm' 
    4303 DECLARE_REDUCTION_BASE(product) 
    4304 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4305 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4306 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4307 ^ 
    4308 program_source:263:44: note: expanded from macro '\ 
    4309 DECLARE_I_REDUCTION_BASE' 
    4310 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4311 ^ 
    4312 program_source:217:40: note: expanded from macro '\ 
    4313 EXPOSE_FUNCTION_I_ARGS_1' 
    4314 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4315 ^ 
    4316 program_source:164:9: note: expanded from macro '\ 
    4317 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4318 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4319 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4320 program_source:285:1: error: illegal string literal in 'asm' 
    4321 DECLARE_REDUCTION_BASE(product) 
    4322 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4323 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4324 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4325 ^ 
    4326 program_source:263:44: note: expanded from macro '\ 
    4327 DECLARE_I_REDUCTION_BASE' 
    4328 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4329 ^ 
    4330 program_source:218:51: note: expanded from macro '\ 
    4331 EXPOSE_FUNCTION_I_ARGS_1' 
    4332 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4333 ^ 
    4334 program_source:164:9: note: expanded from macro '\ 
    4335 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4336 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4337 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4338 program_source:285:1: error: illegal string literal in 'asm' 
    4339 DECLARE_REDUCTION_BASE(product) 
    4340 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4341 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4342 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4343 ^ 
    4344 program_source:264:43: note: expanded from macro '\ 
    4345 DECLARE_I_REDUCTION_BASE' 
    4346 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4347 ^ 
    4348 program_source:217:40: note: expanded from macro '\ 
    4349 EXPOSE_FUNCTION_I_ARGS_1' 
    4350 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4351 ^ 
    4352 program_source:164:9: note: expanded from macro '\ 
    4353 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4354 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4355 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4356 program_source:285:1: error: illegal string literal in 'asm' 
    4357 DECLARE_REDUCTION_BASE(product) 
    4358 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4359 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4360 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4361 ^ 
    4362 program_source:264:43: note: expanded from macro '\ 
    4363 DECLARE_I_REDUCTION_BASE' 
    4364 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4365 ^ 
    4366 program_source:218:51: note: expanded from macro '\ 
    4367 EXPOSE_FUNCTION_I_ARGS_1' 
    4368 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4369 ^ 
    4370 program_source:164:9: note: expanded from macro '\ 
    4371 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4372 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4373 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4374 program_source:285:1: error: illegal string literal in 'asm' 
    4375 DECLARE_REDUCTION_BASE(product) 
    4376 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4377 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4378 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4379 ^ 
    4380 program_source:267:44: note: expanded from macro '\ 
    4381 DECLARE_F_REDUCTION_BASE' 
    4382 #define DECLARE_F_REDUCTION_BASE(METAL_OP) \ 
    4383 ^ 
    4384 program_source:221:40: note: expanded from macro '\ 
    4385 EXPOSE_FUNCTION_F_ARGS_1' 
    4386 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4387 ^ 
    4388 program_source:164:9: note: expanded from macro '\ 
    4389 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4390 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4391 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4392 program_source:285:1: error: illegal string literal in 'asm' 
    4393 DECLARE_REDUCTION_BASE(product) 
    4394 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4395 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4396 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4397 ^ 
    4398 program_source:268:43: note: expanded from macro '\ 
    4399 DECLARE_F_REDUCTION_BASE' 
    4400 EXPOSE_FUNCTION_F_ARGS_1(quad_##METAL_OP) \ 
    4401 ^ 
    4402 program_source:221:40: note: expanded from macro '\ 
    4403 EXPOSE_FUNCTION_F_ARGS_1' 
    4404 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4405 ^ 
    4406 program_source:164:9: note: expanded from macro '\ 
    4407 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4408 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4409 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4410 program_source:286:1: error: illegal string literal in 'asm' 
    4411 DECLARE_REDUCTION_BASE(prefix_inclusive_product) 
    4412 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4413 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4414 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4415 ^ 
    4416 program_source:263:44: note: expanded from macro '\ 
    4417 DECLARE_I_REDUCTION_BASE' 
    4418 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4419 ^ 
    4420 program_source:217:40: note: expanded from macro '\ 
    4421 EXPOSE_FUNCTION_I_ARGS_1' 
    4422 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4423 ^ 
    4424 program_source:164:9: note: expanded from macro '\ 
    4425 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4426 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4427 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4428 program_source:286:1: error: illegal string literal in 'asm' 
    4429 DECLARE_REDUCTION_BASE(prefix_inclusive_product) 
    4430 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4431 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4432 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4433 ^ 
    4434 program_source:263:44: note: expanded from macro '\ 
    4435 DECLARE_I_REDUCTION_BASE' 
    4436 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4437 ^ 
    4438 program_source:218:51: note: expanded from macro '\ 
    4439 EXPOSE_FUNCTION_I_ARGS_1' 
    4440 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4441 ^ 
    4442 program_source:164:9: note: expanded from macro '\ 
    4443 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4444 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4445 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4446 program_source:286:1: error: illegal string literal in 'asm' 
    4447 DECLARE_REDUCTION_BASE(prefix_inclusive_product) 
    4448 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4449 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4450 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4451 ^ 
    4452 program_source:264:43: note: expanded from macro '\ 
    4453 DECLARE_I_REDUCTION_BASE' 
    4454 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4455 ^ 
    4456 program_source:217:40: note: expanded from macro '\ 
    4457 EXPOSE_FUNCTION_I_ARGS_1' 
    4458 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4459 ^ 
    4460 program_source:164:9: note: expanded from macro '\ 
    4461 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4462 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4463 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4464 program_source:286:1: error: illegal string literal in 'asm' 
    4465 DECLARE_REDUCTION_BASE(prefix_inclusive_product) 
    4466 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4467 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4468 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4469 ^ 
    4470 program_source:264:43: note: expanded from macro '\ 
    4471 DECLARE_I_REDUCTION_BASE' 
    4472 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4473 ^ 
    4474 program_source:218:51: note: expanded from macro '\ 
    4475 EXPOSE_FUNCTION_I_ARGS_1' 
    4476 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4477 ^ 
    4478 program_source:164:9: note: expanded from macro '\ 
    4479 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4480 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4481 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4482 program_source:286:1: error: illegal string literal in 'asm' 
    4483 DECLARE_REDUCTION_BASE(prefix_inclusive_product) 
    4484 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4485 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4486 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4487 ^ 
    4488 program_source:267:44: note: expanded from macro '\ 
    4489 DECLARE_F_REDUCTION_BASE' 
    4490 #define DECLARE_F_REDUCTION_BASE(METAL_OP) \ 
    4491 ^ 
    4492 program_source:221:40: note: expanded from macro '\ 
    4493 EXPOSE_FUNCTION_F_ARGS_1' 
    4494 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4495 ^ 
    4496 program_source:164:9: note: expanded from macro '\ 
    4497 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4498 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4499 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4500 program_source:286:1: error: illegal string literal in 'asm' 
    4501 DECLARE_REDUCTION_BASE(prefix_inclusive_product) 
    4502 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4503 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4504 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4505 ^ 
    4506 program_source:268:43: note: expanded from macro '\ 
    4507 DECLARE_F_REDUCTION_BASE' 
    4508 EXPOSE_FUNCTION_F_ARGS_1(quad_##METAL_OP) \ 
    4509 ^ 
    4510 program_source:221:40: note: expanded from macro '\ 
    4511 EXPOSE_FUNCTION_F_ARGS_1' 
    4512 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4513 ^ 
    4514 program_source:164:9: note: expanded from macro '\ 
    4515 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4516 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4517 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4518 program_source:287:1: error: illegal string literal in 'asm' 
    4519 DECLARE_REDUCTION_BASE(prefix_exclusive_product) 
    4520 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4521 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4522 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4523 ^ 
    4524 program_source:263:44: note: expanded from macro '\ 
    4525 DECLARE_I_REDUCTION_BASE' 
    4526 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4527 ^ 
    4528 program_source:217:40: note: expanded from macro '\ 
    4529 EXPOSE_FUNCTION_I_ARGS_1' 
    4530 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4531 ^ 
    4532 program_source:164:9: note: expanded from macro '\ 
    4533 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4534 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4535 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4536 program_source:287:1: error: illegal string literal in 'asm' 
    4537 DECLARE_REDUCTION_BASE(prefix_exclusive_product) 
    4538 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4539 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4540 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4541 ^ 
    4542 program_source:263:44: note: expanded from macro '\ 
    4543 DECLARE_I_REDUCTION_BASE' 
    4544 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4545 ^ 
    4546 program_source:218:51: note: expanded from macro '\ 
    4547 EXPOSE_FUNCTION_I_ARGS_1' 
    4548 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4549 ^ 
    4550 program_source:164:9: note: expanded from macro '\ 
    4551 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4552 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4553 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4554 program_source:287:1: error: illegal string literal in 'asm' 
    4555 DECLARE_REDUCTION_BASE(prefix_exclusive_product) 
    4556 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4557 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4558 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4559 ^ 
    4560 program_source:264:43: note: expanded from macro '\ 
    4561 DECLARE_I_REDUCTION_BASE' 
    4562 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4563 ^ 
    4564 program_source:217:40: note: expanded from macro '\ 
    4565 EXPOSE_FUNCTION_I_ARGS_1' 
    4566 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4567 ^ 
    4568 program_source:164:9: note: expanded from macro '\ 
    4569 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4570 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4571 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4572 program_source:287:1: error: illegal string literal in 'asm' 
    4573 DECLARE_REDUCTION_BASE(prefix_exclusive_product) 
    4574 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4575 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4576 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4577 ^ 
    4578 program_source:264:43: note: expanded from macro '\ 
    4579 DECLARE_I_REDUCTION_BASE' 
    4580 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4581 ^ 
    4582 program_source:218:51: note: expanded from macro '\ 
    4583 EXPOSE_FUNCTION_I_ARGS_1' 
    4584 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4585 ^ 
    4586 program_source:164:9: note: expanded from macro '\ 
    4587 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4588 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4589 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4590 program_source:287:1: error: illegal string literal in 'asm' 
    4591 DECLARE_REDUCTION_BASE(prefix_exclusive_product) 
    4592 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4593 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4594 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4595 ^ 
    4596 program_source:267:44: note: expanded from macro '\ 
    4597 DECLARE_F_REDUCTION_BASE' 
    4598 #define DECLARE_F_REDUCTION_BASE(METAL_OP) \ 
    4599 ^ 
    4600 program_source:221:40: note: expanded from macro '\ 
    4601 EXPOSE_FUNCTION_F_ARGS_1' 
    4602 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4603 ^ 
    4604 program_source:164:9: note: expanded from macro '\ 
    4605 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4606 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4607 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4608 program_source:287:1: error: illegal string literal in 'asm' 
    4609 DECLARE_REDUCTION_BASE(prefix_exclusive_product) 
    4610 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4611 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4612 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4613 ^ 
    4614 program_source:268:43: note: expanded from macro '\ 
    4615 DECLARE_F_REDUCTION_BASE' 
    4616 EXPOSE_FUNCTION_F_ARGS_1(quad_##METAL_OP) \ 
    4617 ^ 
    4618 program_source:221:40: note: expanded from macro '\ 
    4619 EXPOSE_FUNCTION_F_ARGS_1' 
    4620 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4621 ^ 
    4622 program_source:164:9: note: expanded from macro '\ 
    4623 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4624 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4625 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4626 program_source:288:1: error: illegal string literal in 'asm' 
    4627 DECLARE_I_REDUCTION_BASE(and) 
    4628 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4629 program_source:263:44: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    4630 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4631 ^ 
    4632 program_source:217:40: note: expanded from macro '\ 
    4633 EXPOSE_FUNCTION_I_ARGS_1' 
    4634 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4635 ^ 
    4636 program_source:164:9: note: expanded from macro '\ 
    4637 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4638 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4639 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4640 program_source:288:1: error: illegal string literal in 'asm' 
    4641 DECLARE_I_REDUCTION_BASE(and) 
    4642 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4643 program_source:263:44: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    4644 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4645 ^ 
    4646 program_source:218:51: note: expanded from macro '\ 
    4647 EXPOSE_FUNCTION_I_ARGS_1' 
    4648 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4649 ^ 
    4650 program_source:164:9: note: expanded from macro '\ 
    4651 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4652 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4653 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4654 program_source:288:1: error: illegal string literal in 'asm' 
    4655 DECLARE_I_REDUCTION_BASE(and) 
    4656 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4657 program_source:264:43: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    4658 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4659 ^ 
    4660 program_source:217:40: note: expanded from macro '\ 
    4661 EXPOSE_FUNCTION_I_ARGS_1' 
    4662 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4663 ^ 
    4664 program_source:164:9: note: expanded from macro '\ 
    4665 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4666 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4667 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4668 program_source:288:1: error: illegal string literal in 'asm' 
    4669 DECLARE_I_REDUCTION_BASE(and) 
    4670 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4671 program_source:264:43: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    4672 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4673 ^ 
    4674 program_source:218:51: note: expanded from macro '\ 
    4675 EXPOSE_FUNCTION_I_ARGS_1' 
    4676 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4677 ^ 
    4678 program_source:164:9: note: expanded from macro '\ 
    4679 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4680 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4681 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4682 program_source:289:1: error: illegal string literal in 'asm' 
    4683 DECLARE_I_REDUCTION_BASE(or) 
    4684 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4685 program_source:263:44: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    4686 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4687 ^ 
    4688 program_source:217:40: note: expanded from macro '\ 
    4689 EXPOSE_FUNCTION_I_ARGS_1' 
    4690 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4691 ^ 
    4692 program_source:164:9: note: expanded from macro '\ 
    4693 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4694 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4695 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4696 program_source:289:1: error: illegal string literal in 'asm' 
    4697 DECLARE_I_REDUCTION_BASE(or) 
    4698 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4699 program_source:263:44: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    4700 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4701 ^ 
    4702 program_source:218:51: note: expanded from macro '\ 
    4703 EXPOSE_FUNCTION_I_ARGS_1' 
    4704 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4705 ^ 
    4706 program_source:164:9: note: expanded from macro '\ 
    4707 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4708 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4709 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4710 program_source:289:1: error: illegal string literal in 'asm' 
    4711 DECLARE_I_REDUCTION_BASE(or) 
    4712 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4713 program_source:264:43: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    4714 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4715 ^ 
    4716 program_source:217:40: note: expanded from macro '\ 
    4717 EXPOSE_FUNCTION_I_ARGS_1' 
    4718 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4719 ^ 
    4720 program_source:164:9: note: expanded from macro '\ 
    4721 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4722 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4723 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4724 program_source:289:1: error: illegal string literal in 'asm' 
    4725 DECLARE_I_REDUCTION_BASE(or) 
    4726 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4727 program_source:264:43: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    4728 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4729 ^ 
    4730 program_source:218:51: note: expanded from macro '\ 
    4731 EXPOSE_FUNCTION_I_ARGS_1' 
    4732 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4733 ^ 
    4734 program_source:164:9: note: expanded from macro '\ 
    4735 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4736 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4737 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4738 program_source:290:1: error: illegal string literal in 'asm' 
    4739 DECLARE_I_REDUCTION_BASE(xor) 
    4740 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4741 program_source:263:44: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    4742 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4743 ^ 
    4744 program_source:217:40: note: expanded from macro '\ 
    4745 EXPOSE_FUNCTION_I_ARGS_1' 
    4746 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4747 ^ 
    4748 program_source:164:9: note: expanded from macro '\ 
    4749 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4750 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4751 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4752 program_source:290:1: error: illegal string literal in 'asm' 
    4753 DECLARE_I_REDUCTION_BASE(xor) 
    4754 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4755 program_source:263:44: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    4756 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4757 ^ 
    4758 program_source:218:51: note: expanded from macro '\ 
    4759 EXPOSE_FUNCTION_I_ARGS_1' 
    4760 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4761 ^ 
    4762 program_source:164:9: note: expanded from macro '\ 
    4763 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4764 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4765 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4766 program_source:290:1: error: illegal string literal in 'asm' 
    4767 DECLARE_I_REDUCTION_BASE(xor) 
    4768 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4769 program_source:264:43: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    4770 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4771 ^ 
    4772 program_source:217:40: note: expanded from macro '\ 
    4773 EXPOSE_FUNCTION_I_ARGS_1' 
    4774 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4775 ^ 
    4776 program_source:164:9: note: expanded from macro '\ 
    4777 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4778 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4779 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4780 program_source:290:1: error: illegal string literal in 'asm' 
    4781 DECLARE_I_REDUCTION_BASE(xor) 
    4782 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4783 program_source:264:43: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    4784 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4785 ^ 
    4786 program_source:218:51: note: expanded from macro '\ 
    4787 EXPOSE_FUNCTION_I_ARGS_1' 
    4788 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4789 ^ 
    4790 program_source:164:9: note: expanded from macro '\ 
    4791 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4792 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4793 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4794 program_source:292:1: error: illegal string literal in 'asm' 
    4795 DECLARE_SHUFFLE_BASE(broadcast) 
    4796 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4797 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    4798 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    4799 ^ 
    4800 program_source:224:38: note: expanded from macro '\ 
    4801 EXPOSE_FUNCTION_ARGS_2' 
    4802 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    4803 ^ 
    4804 program_source:168:9: note: expanded from macro '\ 
    4805 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    4806 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4807 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4808 program_source:292:1: error: illegal string literal in 'asm' 
    4809 DECLARE_SHUFFLE_BASE(broadcast) 
    4810 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4811 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    4812 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    4813 ^ 
    4814 program_source:225:51: note: expanded from macro '\ 
    4815 EXPOSE_FUNCTION_ARGS_2' 
    4816 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    4817 ^ 
    4818 program_source:168:9: note: expanded from macro '\ 
    4819 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    4820 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4821 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4822 program_source:292:1: error: illegal string literal in 'asm' 
    4823 DECLARE_SHUFFLE_BASE(broadcast) 
    4824 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4825 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    4826 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    4827 ^ 
    4828 program_source:226:52: note: expanded from macro '\ 
    4829 EXPOSE_FUNCTION_ARGS_2' 
    4830 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    4831 ^ 
    4832 program_source:168:9: note: expanded from macro '\ 
    4833 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    4834 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4835 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4836 program_source:292:1: error: illegal string literal in 'asm' 
    4837 DECLARE_SHUFFLE_BASE(broadcast) 
    4838 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4839 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    4840 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    4841 ^ 
    4842 program_source:224:38: note: expanded from macro '\ 
    4843 EXPOSE_FUNCTION_ARGS_2' 
    4844 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    4845 ^ 
    4846 program_source:168:9: note: expanded from macro '\ 
    4847 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    4848 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4849 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4850 program_source:292:1: error: illegal string literal in 'asm' 
    4851 DECLARE_SHUFFLE_BASE(broadcast) 
    4852 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4853 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    4854 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    4855 ^ 
    4856 program_source:225:51: note: expanded from macro '\ 
    4857 EXPOSE_FUNCTION_ARGS_2' 
    4858 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    4859 ^ 
    4860 program_source:168:9: note: expanded from macro '\ 
    4861 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    4862 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4863 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4864 program_source:292:1: error: illegal string literal in 'asm' 
    4865 DECLARE_SHUFFLE_BASE(broadcast) 
    4866 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4867 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    4868 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    4869 ^ 
    4870 program_source:226:52: note: expanded from macro '\ 
    4871 EXPOSE_FUNCTION_ARGS_2' 
    4872 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    4873 ^ 
    4874 program_source:168:9: note: expanded from macro '\ 
    4875 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    4876 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4877 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4878 program_source:293:1: error: illegal string literal in 'asm' 
    4879 DECLARE_REDUCTION_BASE(broadcast_first) 
    4880 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4881 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4882 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4883 ^ 
    4884 program_source:263:44: note: expanded from macro '\ 
    4885 DECLARE_I_REDUCTION_BASE' 
    4886 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4887 ^ 
    4888 program_source:217:40: note: expanded from macro '\ 
    4889 EXPOSE_FUNCTION_I_ARGS_1' 
    4890 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4891 ^ 
    4892 program_source:164:9: note: expanded from macro '\ 
    4893 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4894 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4895 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4896 program_source:293:1: error: illegal string literal in 'asm' 
    4897 DECLARE_REDUCTION_BASE(broadcast_first) 
    4898 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4899 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4900 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4901 ^ 
    4902 program_source:263:44: note: expanded from macro '\ 
    4903 DECLARE_I_REDUCTION_BASE' 
    4904 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4905 ^ 
    4906 program_source:218:51: note: expanded from macro '\ 
    4907 EXPOSE_FUNCTION_I_ARGS_1' 
    4908 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4909 ^ 
    4910 program_source:164:9: note: expanded from macro '\ 
    4911 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4912 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4913 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4914 program_source:293:1: error: illegal string literal in 'asm' 
    4915 DECLARE_REDUCTION_BASE(broadcast_first) 
    4916 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4917 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4918 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4919 ^ 
    4920 program_source:264:43: note: expanded from macro '\ 
    4921 DECLARE_I_REDUCTION_BASE' 
    4922 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4923 ^ 
    4924 program_source:217:40: note: expanded from macro '\ 
    4925 EXPOSE_FUNCTION_I_ARGS_1' 
    4926 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4927 ^ 
    4928 program_source:164:9: note: expanded from macro '\ 
    4929 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4930 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4931 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4932 program_source:293:1: error: illegal string literal in 'asm' 
    4933 DECLARE_REDUCTION_BASE(broadcast_first) 
    4934 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4935 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4936 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4937 ^ 
    4938 program_source:264:43: note: expanded from macro '\ 
    4939 DECLARE_I_REDUCTION_BASE' 
    4940 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4941 ^ 
    4942 program_source:218:51: note: expanded from macro '\ 
    4943 EXPOSE_FUNCTION_I_ARGS_1' 
    4944 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4945 ^ 
    4946 program_source:164:9: note: expanded from macro '\ 
    4947 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4948 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4949 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4950 program_source:293:1: error: illegal string literal in 'asm' 
    4951 DECLARE_REDUCTION_BASE(broadcast_first) 
    4952 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4953 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4954 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4955 ^ 
    4956 program_source:267:44: note: expanded from macro '\ 
    4957 DECLARE_F_REDUCTION_BASE' 
    4958 #define DECLARE_F_REDUCTION_BASE(METAL_OP) \ 
    4959 ^ 
    4960 program_source:221:40: note: expanded from macro '\ 
    4961 EXPOSE_FUNCTION_F_ARGS_1' 
    4962 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4963 ^ 
    4964 program_source:164:9: note: expanded from macro '\ 
    4965 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4966 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4967 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4968 program_source:293:1: error: illegal string literal in 'asm' 
    4969 DECLARE_REDUCTION_BASE(broadcast_first) 
    4970 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4971 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4972 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4973 ^ 
    4974 program_source:268:43: note: expanded from macro '\ 
    4975 DECLARE_F_REDUCTION_BASE' 
    4976 EXPOSE_FUNCTION_F_ARGS_1(quad_##METAL_OP) \ 
    4977 ^ 
    4978 program_source:221:40: note: expanded from macro '\ 
    4979 EXPOSE_FUNCTION_F_ARGS_1' 
    4980 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4981 ^ 
    4982 program_source:164:9: note: expanded from macro '\ 
    4983 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4984 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4985 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4986 program_source:295:1: error: illegal string literal in 'asm' 
    4987 DECLARE_SHUFFLE_BASE(shuffle) 
    4988 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4989 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    4990 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    4991 ^ 
    4992 program_source:224:38: note: expanded from macro '\ 
    4993 EXPOSE_FUNCTION_ARGS_2' 
    4994 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    4995 ^ 
    4996 program_source:168:9: note: expanded from macro '\ 
    4997 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    4998 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4999 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5000 program_source:295:1: error: illegal string literal in 'asm' 
    5001 DECLARE_SHUFFLE_BASE(shuffle) 
    5002 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5003 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5004 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5005 ^ 
    5006 program_source:225:51: note: expanded from macro '\ 
    5007 EXPOSE_FUNCTION_ARGS_2' 
    5008 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    5009 ^ 
    5010 program_source:168:9: note: expanded from macro '\ 
    5011 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5012 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5013 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5014 program_source:295:1: error: illegal string literal in 'asm' 
    5015 DECLARE_SHUFFLE_BASE(shuffle) 
    5016 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5017 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5018 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5019 ^ 
    5020 program_source:226:52: note: expanded from macro '\ 
    5021 EXPOSE_FUNCTION_ARGS_2' 
    5022 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    5023 ^ 
    5024 program_source:168:9: note: expanded from macro '\ 
    5025 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5026 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5027 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5028 program_source:295:1: error: illegal string literal in 'asm' 
    5029 DECLARE_SHUFFLE_BASE(shuffle) 
    5030 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5031 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5032 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5033 ^ 
    5034 program_source:224:38: note: expanded from macro '\ 
    5035 EXPOSE_FUNCTION_ARGS_2' 
    5036 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    5037 ^ 
    5038 program_source:168:9: note: expanded from macro '\ 
    5039 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5040 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5041 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5042 program_source:295:1: error: illegal string literal in 'asm' 
    5043 DECLARE_SHUFFLE_BASE(shuffle) 
    5044 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5045 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5046 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5047 ^ 
    5048 program_source:225:51: note: expanded from macro '\ 
    5049 EXPOSE_FUNCTION_ARGS_2' 
    5050 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    5051 ^ 
    5052 program_source:168:9: note: expanded from macro '\ 
    5053 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5054 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5055 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5056 program_source:295:1: error: illegal string literal in 'asm' 
    5057 DECLARE_SHUFFLE_BASE(shuffle) 
    5058 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5059 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5060 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5061 ^ 
    5062 program_source:226:52: note: expanded from macro '\ 
    5063 EXPOSE_FUNCTION_ARGS_2' 
    5064 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    5065 ^ 
    5066 program_source:168:9: note: expanded from macro '\ 
    5067 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5068 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5069 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5070 program_source:296:1: error: illegal string literal in 'asm' 
    5071 DECLARE_SHUFFLE_BASE(shuffle_xor) 
    5072 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5073 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5074 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5075 ^ 
    5076 program_source:224:38: note: expanded from macro '\ 
    5077 EXPOSE_FUNCTION_ARGS_2' 
    5078 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    5079 ^ 
    5080 program_source:168:9: note: expanded from macro '\ 
    5081 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5082 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5083 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5084 program_source:296:1: error: illegal string literal in 'asm' 
    5085 DECLARE_SHUFFLE_BASE(shuffle_xor) 
    5086 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5087 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5088 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5089 ^ 
    5090 program_source:225:51: note: expanded from macro '\ 
    5091 EXPOSE_FUNCTION_ARGS_2' 
    5092 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    5093 ^ 
    5094 program_source:168:9: note: expanded from macro '\ 
    5095 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5096 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5097 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5098 program_source:296:1: error: illegal string literal in 'asm' 
    5099 DECLARE_SHUFFLE_BASE(shuffle_xor) 
    5100 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5101 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5102 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5103 ^ 
    5104 program_source:226:52: note: expanded from macro '\ 
    5105 EXPOSE_FUNCTION_ARGS_2' 
    5106 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    5107 ^ 
    5108 program_source:168:9: note: expanded from macro '\ 
    5109 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5110 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5111 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5112 program_source:296:1: error: illegal string literal in 'asm' 
    5113 DECLARE_SHUFFLE_BASE(shuffle_xor) 
    5114 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5115 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5116 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5117 ^ 
    5118 program_source:224:38: note: expanded from macro '\ 
    5119 EXPOSE_FUNCTION_ARGS_2' 
    5120 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    5121 ^ 
    5122 program_source:168:9: note: expanded from macro '\ 
    5123 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5124 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5125 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5126 program_source:296:1: error: illegal string literal in 'asm' 
    5127 DECLARE_SHUFFLE_BASE(shuffle_xor) 
    5128 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5129 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5130 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5131 ^ 
    5132 program_source:225:51: note: expanded from macro '\ 
    5133 EXPOSE_FUNCTION_ARGS_2' 
    5134 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    5135 ^ 
    5136 program_source:168:9: note: expanded from macro '\ 
    5137 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5138 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5139 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5140 program_source:296:1: error: illegal string literal in 'asm' 
    5141 DECLARE_SHUFFLE_BASE(shuffle_xor) 
    5142 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5143 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5144 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5145 ^ 
    5146 program_source:226:52: note: expanded from macro '\ 
    5147 EXPOSE_FUNCTION_ARGS_2' 
    5148 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    5149 ^ 
    5150 program_source:168:9: note: expanded from macro '\ 
    5151 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5152 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5153 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5154 program_source:297:1: error: illegal string literal in 'asm' 
    5155 DECLARE_SHUFFLE_BASE(shuffle_up) 
    5156 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5157 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5158 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5159 ^ 
    5160 program_source:224:38: note: expanded from macro '\ 
    5161 EXPOSE_FUNCTION_ARGS_2' 
    5162 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    5163 ^ 
    5164 program_source:168:9: note: expanded from macro '\ 
    5165 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5166 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5167 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5168 program_source:297:1: error: illegal string literal in 'asm' 
    5169 DECLARE_SHUFFLE_BASE(shuffle_up) 
    5170 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5171 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5172 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5173 ^ 
    5174 program_source:225:51: note: expanded from macro '\ 
    5175 EXPOSE_FUNCTION_ARGS_2' 
    5176 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    5177 ^ 
    5178 program_source:168:9: note: expanded from macro '\ 
    5179 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5180 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5181 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5182 program_source:297:1: error: illegal string literal in 'asm' 
    5183 DECLARE_SHUFFLE_BASE(shuffle_up) 
    5184 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5185 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5186 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5187 ^ 
    5188 program_source:226:52: note: expanded from macro '\ 
    5189 EXPOSE_FUNCTION_ARGS_2' 
    5190 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    5191 ^ 
    5192 program_source:168:9: note: expanded from macro '\ 
    5193 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5194 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5195 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5196 program_source:297:1: error: illegal string literal in 'asm' 
    5197 DECLARE_SHUFFLE_BASE(shuffle_up) 
    5198 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5199 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5200 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5201 ^ 
    5202 program_source:224:38: note: expanded from macro '\ 
    5203 EXPOSE_FUNCTION_ARGS_2' 
    5204 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    5205 ^ 
    5206 program_source:168:9: note: expanded from macro '\ 
    5207 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5208 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5209 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5210 program_source:297:1: error: illegal string literal in 'asm' 
    5211 DECLARE_SHUFFLE_BASE(shuffle_up) 
    5212 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5213 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5214 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5215 ^ 
    5216 program_source:225:51: note: expanded from macro '\ 
    5217 EXPOSE_FUNCTION_ARGS_2' 
    5218 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    5219 ^ 
    5220 program_source:168:9: note: expanded from macro '\ 
    5221 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5222 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5223 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5224 program_source:297:1: error: illegal string literal in 'asm' 
    5225 DECLARE_SHUFFLE_BASE(shuffle_up) 
    5226 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5227 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5228 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5229 ^ 
    5230 program_source:226:52: note: expanded from macro '\ 
    5231 EXPOSE_FUNCTION_ARGS_2' 
    5232 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    5233 ^ 
    5234 program_source:168:9: note: expanded from macro '\ 
    5235 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5236 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5237 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5238 program_source:298:1: error: illegal string literal in 'asm' 
    5239 DECLARE_SHUFFLE_BASE(shuffle_down) 
    5240 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5241 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5242 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5243 ^ 
    5244 program_source:224:38: note: expanded from macro '\ 
    5245 EXPOSE_FUNCTION_ARGS_2' 
    5246 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    5247 ^ 
    5248 program_source:168:9: note: expanded from macro '\ 
    5249 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5250 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5251 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5252 program_source:298:1: error: illegal string literal in 'asm' 
    5253 DECLARE_SHUFFLE_BASE(shuffle_down) 
    5254 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5255 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5256 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5257 ^ 
    5258 program_source:225:51: note: expanded from macro '\ 
    5259 EXPOSE_FUNCTION_ARGS_2' 
    5260 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    5261 ^ 
    5262 program_source:168:9: note: expanded from macro '\ 
    5263 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5264 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5265 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5266 program_source:298:1: error: illegal string literal in 'asm' 
    5267 DECLARE_SHUFFLE_BASE(shuffle_down) 
    5268 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5269 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5270 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5271 ^ 
    5272 program_source:226:52: note: expanded from macro '\ 
    5273 EXPOSE_FUNCTION_ARGS_2' 
    5274 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    5275 ^ 
    5276 program_source:168:9: note: expanded from macro '\ 
    5277 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5278 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5279 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5280 program_source:298:1: error: illegal string literal in 'asm' 
    5281 DECLARE_SHUFFLE_BASE(shuffle_down) 
    5282 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5283 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5284 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5285 ^ 
    5286 program_source:224:38: note: expanded from macro '\ 
    5287 EXPOSE_FUNCTION_ARGS_2' 
    5288 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    5289 ^ 
    5290 program_source:168:9: note: expanded from macro '\ 
    5291 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5292 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5293 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5294 program_source:298:1: error: illegal string literal in 'asm' 
    5295 DECLARE_SHUFFLE_BASE(shuffle_down) 
    5296 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5297 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5298 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5299 ^ 
    5300 program_source:225:51: note: expanded from macro '\ 
    5301 EXPOSE_FUNCTION_ARGS_2' 
    5302 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    5303 ^ 
    5304 program_source:168:9: note: expanded from macro '\ 
    5305 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5306 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5307 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5308 program_source:298:1: error: illegal string literal in 'asm' 
    5309 DECLARE_SHUFFLE_BASE(shuffle_down) 
    5310 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5311 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5312 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5313 ^ 
    5314 program_source:226:52: note: expanded from macro '\ 
    5315 EXPOSE_FUNCTION_ARGS_2' 
    5316 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    5317 ^ 
    5318 program_source:168:9: note: expanded from macro '\ 
    5319 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5320 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5321 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5322 program_source:299:1: error: illegal string literal in 'asm' 
    5323 DECLARE_SHUFFLE_BASE(shuffle_rotate_up) 
    5324 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5325 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5326 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5327 ^ 
    5328 program_source:224:38: note: expanded from macro '\ 
    5329 EXPOSE_FUNCTION_ARGS_2' 
    5330 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    5331 ^ 
    5332 program_source:168:9: note: expanded from macro '\ 
    5333 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5334 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5335 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5336 program_source:299:1: error: illegal string literal in 'asm' 
    5337 DECLARE_SHUFFLE_BASE(shuffle_rotate_up) 
    5338 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5339 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5340 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5341 ^ 
    5342 program_source:225:51: note: expanded from macro '\ 
    5343 EXPOSE_FUNCTION_ARGS_2' 
    5344 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    5345 ^ 
    5346 program_source:168:9: note: expanded from macro '\ 
    5347 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5348 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5349 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5350 program_source:299:1: error: illegal string literal in 'asm' 
    5351 DECLARE_SHUFFLE_BASE(shuffle_rotate_up) 
    5352 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5353 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5354 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5355 ^ 
    5356 program_source:226:52: note: expanded from macro '\ 
    5357 EXPOSE_FUNCTION_ARGS_2' 
    5358 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    5359 ^ 
    5360 program_source:168:9: note: expanded from macro '\ 
    5361 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5362 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5363 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5364 program_source:299:1: error: illegal string literal in 'asm' 
    5365 DECLARE_SHUFFLE_BASE(shuffle_rotate_up) 
    5366 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5367 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5368 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5369 ^ 
    5370 program_source:224:38: note: expanded from macro '\ 
    5371 EXPOSE_FUNCTION_ARGS_2' 
    5372 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    5373 ^ 
    5374 program_source:168:9: note: expanded from macro '\ 
    5375 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5376 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5377 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5378 program_source:299:1: error: illegal string literal in 'asm' 
    5379 DECLARE_SHUFFLE_BASE(shuffle_rotate_up) 
    5380 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5381 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5382 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5383 ^ 
    5384 program_source:225:51: note: expanded from macro '\ 
    5385 EXPOSE_FUNCTION_ARGS_2' 
    5386 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    5387 ^ 
    5388 program_source:168:9: note: expanded from macro '\ 
    5389 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5390 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5391 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5392 program_source:299:1: error: illegal string literal in 'asm' 
    5393 DECLARE_SHUFFLE_BASE(shuffle_rotate_up) 
    5394 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5395 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5396 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5397 ^ 
    5398 program_source:226:52: note: expanded from macro '\ 
    5399 EXPOSE_FUNCTION_ARGS_2' 
    5400 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    5401 ^ 
    5402 program_source:168:9: note: expanded from macro '\ 
    5403 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5404 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5405 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5406 program_source:300:1: error: illegal string literal in 'asm' 
    5407 DECLARE_SHUFFLE_BASE(shuffle_rotate_down) 
    5408 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5409 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5410 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5411 ^ 
    5412 program_source:224:38: note: expanded from macro '\ 
    5413 EXPOSE_FUNCTION_ARGS_2' 
    5414 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    5415 ^ 
    5416 program_source:168:9: note: expanded from macro '\ 
    5417 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5418 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5419 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5420 program_source:300:1: error: illegal string literal in 'asm' 
    5421 DECLARE_SHUFFLE_BASE(shuffle_rotate_down) 
    5422 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5423 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5424 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5425 ^ 
    5426 program_source:225:51: note: expanded from macro '\ 
    5427 EXPOSE_FUNCTION_ARGS_2' 
    5428 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    5429 ^ 
    5430 program_source:168:9: note: expanded from macro '\ 
    5431 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5432 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5433 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5434 program_source:300:1: error: illegal string literal in 'asm' 
    5435 DECLARE_SHUFFLE_BASE(shuffle_rotate_down) 
    5436 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5437 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5438 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5439 ^ 
    5440 program_source:226:52: note: expanded from macro '\ 
    5441 EXPOSE_FUNCTION_ARGS_2' 
    5442 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    5443 ^ 
    5444 program_source:168:9: note: expanded from macro '\ 
    5445 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5446 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5447 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5448 program_source:300:1: error: illegal string literal in 'asm' 
    5449 DECLARE_SHUFFLE_BASE(shuffle_rotate_down) 
    5450 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5451 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5452 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5453 ^ 
    5454 program_source:224:38: note: expanded from macro '\ 
    5455 EXPOSE_FUNCTION_ARGS_2' 
    5456 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    5457 ^ 
    5458 program_source:168:9: note: expanded from macro '\ 
    5459 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5460 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5461 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5462 program_source:300:1: error: illegal string literal in 'asm' 
    5463 DECLARE_SHUFFLE_BASE(shuffle_rotate_down) 
    5464 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5465 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5466 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5467 ^ 
    5468 program_source:225:51: note: expanded from macro '\ 
    5469 EXPOSE_FUNCTION_ARGS_2' 
    5470 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    5471 ^ 
    5472 program_source:168:9: note: expanded from macro '\ 
    5473 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5474 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5475 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5476 program_source:300:1: error: illegal string literal in 'asm' 
    5477 DECLARE_SHUFFLE_BASE(shuffle_rotate_down) 
    5478 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5479 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5480 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5481 ^ 
    5482 program_source:226:52: note: expanded from macro '\ 
    5483 EXPOSE_FUNCTION_ARGS_2' 
    5484 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    5485 ^ 
    5486 program_source:168:9: note: expanded from macro '\ 
    5487 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5488 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5489 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5490 program_source:302:1: error: illegal string literal in 'asm' 
    5491 EXPOSE_FUNCTION_I_ARGS_3(ctz) 
    5492 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5493 program_source:229:40: note: expanded from macro 'EXPOSE_FUNCTION_I_ARGS_3' 
    5494 #define EXPOSE_FUNCTION_I_ARGS_3(EXPR) \ 
    5495 ^ 
    5496 program_source:172:9: note: expanded from macro '\ 
    5497 EXPOSE_FUNCTION_OVERLOAD_ARGS_3' 
    5498 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5499 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5500 program_source:302:1: error: implicit declaration of function '___metal_ctz'
    5501 is invalid in OpenCL 
    5502 program_source:229:40: note: expanded from macro 'EXPOSE_FUNCTION_I_ARGS_3' 
    5503 #define EXPOSE_FUNCTION_I_ARGS_3(EXPR) \ 
    5504 ^ 
    5505 program_source:175:10: note: expanded from macro '\ 
    5506 EXPOSE_FUNCTION_OVERLOAD_ARGS_3' 
    5507 return ___metal_##EXPR(x, false); \ 
    5508 ^ 
    5509 :12:1: note: expanded from here 
    5510 ___metal_ctz 
    5511 ^ 
    5512 program_source:302:1: error: illegal string literal in 'asm' 
    5513 EXPOSE_FUNCTION_I_ARGS_3(ctz) 
    5514 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5515 program_source:230:49: note: expanded from macro 'EXPOSE_FUNCTION_I_ARGS_3' 
    5516 EXPOSE_FUNCTION_OVERLOAD_ARGS_3(EXPR, int, i32) \ 
    5517 ^ 
    5518 program_source:172:9: note: expanded from macro '\ 
    5519 EXPOSE_FUNCTION_OVERLOAD_ARGS_3' 
    5520 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5521 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5522 program_source:302:1: error: implicit declaration of function '___metal_ctz'
    5523 is invalid in OpenCL 
    5524 program_source:230:49: note: expanded from macro 'EXPOSE_FUNCTION_I_ARGS_3' 
    5525 EXPOSE_FUNCTION_OVERLOAD_ARGS_3(EXPR, int, i32) \ 
    5526 ^ 
    5527 program_source:175:10: note: expanded from macro '\ 
    5528 EXPOSE_FUNCTION_OVERLOAD_ARGS_3' 
    5529 return ___metal_##EXPR(x, false); \ 
    5530 ^ 
    5531 :16:1: note: expanded from here 
    5532 ___metal_ctz 
    5533 ^ 
    5534 program_source:317:1: error: implicit declaration of function 'simd_sum' is
    5535 invalid in OpenCL 
    5536 DECLARE_REDUCTION_UNIFORM(sum, reduce_add) 
    5537 ^ 
    5538 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5539 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5540 ^ 
    5541 program_source:305:26: note: expanded from macro '\ 
    5542 DECLARE_I_REDUCTION_UNIFORM' 
    5543 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5544 ^ 
    5545 :17:1: note: expanded from here 
    5546 simd_sum 
    5547 ^ 
    5548 program_source:317:1: error: implicit declaration of function 'simd_sum' is
    5549 invalid in OpenCL 
    5550 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5551 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5552 ^ 
    5553 program_source:305:26: note: expanded from macro '\ 
    5554 DECLARE_I_REDUCTION_UNIFORM' 
    5555 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5556 ^ 
    5557 :17:1: note: expanded from here 
    5558 simd_sum 
    5559 ^ 
    5560 program_source:317:1: error: implicit declaration of function 'simd_sum' is
    5561 invalid in OpenCL 
    5562 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5563 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5564 ^ 
    5565 program_source:308:26: note: expanded from macro '\ 
    5566 DECLARE_F_REDUCTION_UNIFORM' 
    5567 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5568 ^ 
    5569 :19:1: note: expanded from here 
    5570 simd_sum 
    5571 ^ 
    5572 program_source:318:1: error: implicit declaration of function
    5573 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    5574 DECLARE_REDUCTION_UNIFORM(prefix_inclusive_sum, scan_inclusive_add) 
    5575 ^ 
    5576 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5577 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5578 ^ 
    5579 program_source:305:26: note: expanded from macro '\ 
    5580 DECLARE_I_REDUCTION_UNIFORM' 
    5581 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5582 ^ 
    5583 :21:1: note: expanded from here 
    5584 simd_prefix_inclusive_sum 
    5585 ^ 
    5586 program_source:318:1: error: implicit declaration of function
    5587 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    5588 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5589 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5590 ^ 
    5591 program_source:305:26: note: expanded from macro '\ 
    5592 DECLARE_I_REDUCTION_UNIFORM' 
    5593 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5594 ^ 
    5595 :21:1: note: expanded from here 
    5596 simd_prefix_inclusive_sum 
    5597 ^ 
    5598 program_source:318:1: error: implicit declaration of function
    5599 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    5600 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5601 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5602 ^ 
    5603 program_source:308:26: note: expanded from macro '\ 
    5604 DECLARE_F_REDUCTION_UNIFORM' 
    5605 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5606 ^ 
    5607 :23:1: note: expanded from here 
    5608 simd_prefix_inclusive_sum 
    5609 ^ 
    5610 program_source:319:1: error: implicit declaration of function
    5611 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    5612 DECLARE_REDUCTION_UNIFORM(prefix_exclusive_sum, scan_exclusive_add) 
    5613 ^ 
    5614 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5615 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5616 ^ 
    5617 program_source:305:26: note: expanded from macro '\ 
    5618 DECLARE_I_REDUCTION_UNIFORM' 
    5619 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5620 ^ 
    5621 :25:1: note: expanded from here 
    5622 simd_prefix_exclusive_sum 
    5623 ^ 
    5624 program_source:319:1: error: implicit declaration of function
    5625 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    5626 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5627 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5628 ^ 
    5629 program_source:305:26: note: expanded from macro '\ 
    5630 DECLARE_I_REDUCTION_UNIFORM' 
    5631 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5632 ^ 
    5633 :25:1: note: expanded from here 
    5634 simd_prefix_exclusive_sum 
    5635 ^ 
    5636 program_source:319:1: error: implicit declaration of function
    5637 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    5638 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5639 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5640 ^ 
    5641 program_source:308:26: note: expanded from macro '\ 
    5642 DECLARE_F_REDUCTION_UNIFORM' 
    5643 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5644 ^ 
    5645 :27:1: note: expanded from here 
    5646 simd_prefix_exclusive_sum 
    5647 ^ 
    5648 program_source:320:1: error: implicit declaration of function 'simd_min' is
    5649 invalid in OpenCL 
    5650 DECLARE_REDUCTION_UNIFORM(min, reduce_min) 
    5651 ^ 
    5652 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5653 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5654 ^ 
    5655 program_source:305:26: note: expanded from macro '\ 
    5656 DECLARE_I_REDUCTION_UNIFORM' 
    5657 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5658 ^ 
    5659 :29:1: note: expanded from here 
    5660 simd_min 
    5661 ^ 
    5662 program_source:320:1: error: implicit declaration of function 'simd_min' is
    5663 invalid in OpenCL 
    5664 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5665 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5666 ^ 
    5667 program_source:305:26: note: expanded from macro '\ 
    5668 DECLARE_I_REDUCTION_UNIFORM' 
    5669 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5670 ^ 
    5671 :29:1: note: expanded from here 
    5672 simd_min 
    5673 ^ 
    5674 program_source:320:1: error: implicit declaration of function 'simd_min' is
    5675 invalid in OpenCL 
    5676 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5677 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5678 ^ 
    5679 program_source:308:26: note: expanded from macro '\ 
    5680 DECLARE_F_REDUCTION_UNIFORM' 
    5681 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5682 ^ 
    5683 :31:1: note: expanded from here 
    5684 simd_min 
    5685 ^ 
    5686 program_source:321:1: error: implicit declaration of function 'simd_max' is
    5687 invalid in OpenCL 
    5688 DECLARE_REDUCTION_UNIFORM(max, reduce_max) 
    5689 ^ 
    5690 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5691 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5692 ^ 
    5693 program_source:305:26: note: expanded from macro '\ 
    5694 DECLARE_I_REDUCTION_UNIFORM' 
    5695 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5696 ^ 
    5697 :33:1: note: expanded from here 
    5698 simd_max 
    5699 ^ 
    5700 program_source:321:1: error: implicit declaration of function 'simd_max' is
    5701 invalid in OpenCL 
    5702 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5703 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5704 ^ 
    5705 program_source:305:26: note: expanded from macro '\ 
    5706 DECLARE_I_REDUCTION_UNIFORM' 
    5707 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5708 ^ 
    5709 :33:1: note: expanded from here 
    5710 simd_max 
    5711 ^ 
    5712 program_source:321:1: error: implicit declaration of function 'simd_max' is
    5713 invalid in OpenCL 
    5714 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5715 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5716 ^ 
    5717 program_source:308:26: note: expanded from macro '\ 
    5718 DECLARE_F_REDUCTION_UNIFORM' 
    5719 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5720 ^ 
    5721 :35:1: note: expanded from here 
    5722 simd_max 
    5723 ^ 
    5724 program_source:323:1: error: implicit declaration of function 'simd_shuffle'
    5725 is invalid in OpenCL 
    5726 DECLARE_SHUFFLE_UNIFORM(shuffle, shuffle) 
    5727 ^ 
    5728 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5729 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5730 ^ 
    5731 :37:1: note: expanded from here 
    5732 simd_shuffle 
    5733 ^ 
    5734 program_source:323:1: error: implicit declaration of function 'simd_shuffle'
    5735 is invalid in OpenCL 
    5736 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5737 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5738 ^ 
    5739 :37:1: note: expanded from here 
    5740 simd_shuffle 
    5741 ^ 
    5742 program_source:323:1: error: implicit declaration of function 'simd_shuffle'
    5743 is invalid in OpenCL 
    5744 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5745 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5746 ^ 
    5747 :37:1: note: expanded from here 
    5748 simd_shuffle 
    5749 ^ 
    5750 program_source:324:1: error: implicit declaration of function
    5751 'simd_shuffle_xor' is invalid in OpenCL 
    5752 DECLARE_SHUFFLE_UNIFORM(shuffle_xor, shuffle_xor) 
    5753 ^ 
    5754 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5755 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5756 ^ 
    5757 :39:1: note: expanded from here 
    5758 simd_shuffle_xor 
    5759 ^ 
    5760 program_source:324:1: error: implicit declaration of function
    5761 'simd_shuffle_xor' is invalid in OpenCL 
    5762 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5763 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5764 ^ 
    5765 :39:1: note: expanded from here 
    5766 simd_shuffle_xor 
    5767 ^ 
    5768 program_source:324:1: error: implicit declaration of function
    5769 'simd_shuffle_xor' is invalid in OpenCL 
    5770 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5771 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5772 ^ 
    5773 :39:1: note: expanded from here 
    5774 simd_shuffle_xor 
    5775 ^ 
    5776 program_source:325:1: error: implicit declaration of function
    5777 'simd_shuffle_up' is invalid in OpenCL 
    5778 DECLARE_SHUFFLE_UNIFORM(shuffle_up, shuffle_up) 
    5779 ^ 
    5780 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5781 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5782 ^ 
    5783 :41:1: note: expanded from here 
    5784 simd_shuffle_up 
    5785 ^ 
    5786 program_source:325:1: error: implicit declaration of function
    5787 'simd_shuffle_up' is invalid in OpenCL 
    5788 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5789 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5790 ^ 
    5791 :41:1: note: expanded from here 
    5792 simd_shuffle_up 
    5793 ^ 
    5794 program_source:325:1: error: implicit declaration of function
    5795 'simd_shuffle_up' is invalid in OpenCL 
    5796 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5797 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5798 ^ 
    5799 :41:1: note: expanded from here 
    5800 simd_shuffle_up 
    5801 ^ 
    5802 program_source:326:1: error: implicit declaration of function
    5803 'simd_shuffle_down' is invalid in OpenCL 
    5804 DECLARE_SHUFFLE_UNIFORM(shuffle_down, shuffle_down) 
    5805 ^ 
    5806 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5807 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5808 ^ 
    5809 :43:1: note: expanded from here 
    5810 simd_shuffle_down 
    5811 ^ 
    5812 program_source:326:1: error: implicit declaration of function
    5813 'simd_shuffle_down' is invalid in OpenCL 
    5814 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5815 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5816 ^ 
    5817 :43:1: note: expanded from here 
    5818 simd_shuffle_down 
    5819 ^ 
    5820 program_source:326:1: error: implicit declaration of function
    5821 'simd_shuffle_down' is invalid in OpenCL 
    5822 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5823 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5824 ^ 
    5825 :43:1: note: expanded from here 
    5826 simd_shuffle_down 
    5827 ^ 
    5828 program_source:327:1: error: implicit declaration of function
    5829 'simd_shuffle_rotate_down' is invalid in OpenCL 
    5830 DECLARE_SHUFFLE_UNIFORM(shuffle_rotate_down, rotate) 
    5831 ^ 
    5832 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5833 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5834 ^ 
    5835 :45:1: note: expanded from here 
    5836 simd_shuffle_rotate_down 
    5837 ^ 
    5838 program_source:327:1: error: implicit declaration of function
    5839 'simd_shuffle_rotate_down' is invalid in OpenCL 
    5840 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5841 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5842 ^ 
    5843 :45:1: note: expanded from here 
    5844 simd_shuffle_rotate_down 
    5845 ^ 
    5846 program_source:327:1: error: implicit declaration of function
    5847 'simd_shuffle_rotate_down' is invalid in OpenCL 
    5848 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5849 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5850 ^ 
    5851 :45:1: note: expanded from here 
    5852 simd_shuffle_rotate_down 
    5853 ^ 
    5854 program_source:329:1: error: implicit declaration of function 'simd_broadcast'
    5855 is invalid in OpenCL 
    5856 DECLARE_SHUFFLE_UNIFORM(broadcast, broadcast) 
    5857 ^ 
    5858 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5859 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5860 ^ 
    5861 :47:1: note: expanded from here 
    5862 simd_broadcast 
    5863 ^ 
    5864 program_source:329:1: error: implicit declaration of function 'simd_broadcast'
    5865 is invalid in OpenCL 
    5866 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5867 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5868 ^ 
    5869 :47:1: note: expanded from here 
    5870 simd_broadcast 
    5871 ^ 
    5872 program_source:329:1: error: implicit declaration of function 'simd_broadcast'
    5873 is invalid in OpenCL 
    5874 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5875 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5876 ^ 
    5877 :47:1: note: expanded from here 
    5878 simd_broadcast 
    5879 ^ 
    5880 program_source:330:1: error: implicit declaration of function
    5881 'simd_broadcast_first' is invalid in OpenCL 
    5882 DECLARE_REDUCTION_UNIFORM(broadcast_first, broadcast_first) 
    5883 ^ 
    5884 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5885 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5886 ^ 
    5887 program_source:305:26: note: expanded from macro '\ 
    5888 DECLARE_I_REDUCTION_UNIFORM' 
    5889 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5890 ^ 
    5891 :49:1: note: expanded from here 
    5892 simd_broadcast_first 
    5893 ^ 
    5894 program_source:330:1: error: implicit declaration of function
    5895 'simd_broadcast_first' is invalid in OpenCL 
    5896 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5897 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5898 ^ 
    5899 program_source:305:26: note: expanded from macro '\ 
    5900 DECLARE_I_REDUCTION_UNIFORM' 
    5901 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5902 ^ 
    5903 :49:1: note: expanded from here 
    5904 simd_broadcast_first 
    5905 ^ 
    5906 program_source:330:1: error: implicit declaration of function
    5907 'simd_broadcast_first' is invalid in OpenCL 
    5908 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5909 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5910 ^ 
    5911 program_source:308:26: note: expanded from macro '\ 
    5912 DECLARE_F_REDUCTION_UNIFORM' 
    5913 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5914 ^ 
    5915 :51:1: note: expanded from here 
    5916 simd_broadcast_first 
    5917 ^ 
    5918 program_source:345:1: error: implicit declaration of function 'simd_sum' is
    5919 invalid in OpenCL 
    5920 DECLARE_REDUCTION_NON_UNIFORM(sum, reduce_add) 
    5921 ^ 
    5922 program_source:338:60: note: expanded from macro
    5923 'DECLARE_REDUCTION_NON_UNIFORM' 
    5924 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5925 ^ 
    5926 program_source:333:26: note: expanded from macro '\ 
    5927 DECLARE_I_REDUCTION_NON_UNIFORM' 
    5928 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    5929 ^ 
    5930 :53:1: note: expanded from here 
    5931 simd_sum 
    5932 ^ 
    5933 program_source:345:1: error: implicit declaration of function 'simd_sum' is
    5934 invalid in OpenCL 
    5935 program_source:338:60: note: expanded from macro
    5936 'DECLARE_REDUCTION_NON_UNIFORM' 
    5937 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5938 ^ 
    5939 program_source:333:26: note: expanded from macro '\ 
    5940 DECLARE_I_REDUCTION_NON_UNIFORM' 
    5941 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    5942 ^ 
    5943 :53:1: note: expanded from here 
    5944 simd_sum 
    5945 ^ 
    5946 program_source:345:1: error: implicit declaration of function 'simd_sum' is
    5947 invalid in OpenCL 
    5948 program_source:339:54: note: expanded from macro
    5949 'DECLARE_REDUCTION_NON_UNIFORM' 
    5950 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5951 ^ 
    5952 program_source:336:26: note: expanded from macro '\ 
    5953 DECLARE_F_REDUCTION_NON_UNIFORM' 
    5954 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    5955 ^ 
    5956 :55:1: note: expanded from here 
    5957 simd_sum 
    5958 ^ 
     1976[deleted to fit within ticket limits]
     1977
    59591978program_source:346:1: error: implicit declaration of function
    59601979'simd_prefix_inclusive_sum' is invalid in OpenCL