Changes between Initial Version and Version 1 of Ticket #19764


Ignore:
Timestamp:
Jan 22, 2026, 2:20:59 PM (11 hours ago)
Author:
Eric Pettersen
Comment:

Legend:

Unmodified
Added
Removed
Modified

  • Ticket #19764

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

  • Ticket #19764 – Description

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