Changes between Initial Version and Version 1 of Ticket #19864


Ignore:
Timestamp:
Feb 19, 2026, 10:30:43 AM (9 days ago)
Author:
Eric Pettersen
Comment:

Legend:

Unmodified
Added
Removed
Modified

  • Ticket #19864

    • 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: various OpenMM/OpenCL errors

  • Ticket #19864 – Description

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