Changes between Initial Version and Version 1 of Ticket #20347


Ignore:
Timestamp:
May 19, 2026, 12:57:37 PM (3 weeks ago)
Author:
Eric Pettersen
Comment:

Legend:

Unmodified
Added
Removed
Modified

  • Ticket #20347

    • 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 #20347 – Description

    initial v1  
    16621662DECLARE_SHUFFLE_BASE(shuffle_xor) 
    16631663^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1664 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1665 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    1666 ^ 
    1667 program_source:225:51: note: expanded from macro '\ 
    1668 EXPOSE_FUNCTION_ARGS_2' 
    1669 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    1670 ^ 
    1671 program_source:168:9: note: expanded from macro '\ 
    1672 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1673 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1674 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1675 program_source:296:1: error: illegal string literal in 'asm' 
    1676 DECLARE_SHUFFLE_BASE(shuffle_xor) 
    1677 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1678 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1679 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    1680 ^ 
    1681 program_source:226:52: note: expanded from macro '\ 
    1682 EXPOSE_FUNCTION_ARGS_2' 
    1683 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    1684 ^ 
    1685 program_source:168:9: note: expanded from macro '\ 
    1686 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1687 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1688 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1689 program_source:296:1: error: illegal string literal in 'asm' 
    1690 DECLARE_SHUFFLE_BASE(shuffle_xor) 
    1691 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1692 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1693 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    1694 ^ 
    1695 program_source:224:38: note: expanded from macro '\ 
    1696 EXPOSE_FUNCTION_ARGS_2' 
    1697 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    1698 ^ 
    1699 program_source:168:9: note: expanded from macro '\ 
    1700 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1701 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1702 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1703 program_source:296:1: error: illegal string literal in 'asm' 
    1704 DECLARE_SHUFFLE_BASE(shuffle_xor) 
    1705 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1706 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1707 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    1708 ^ 
    1709 program_source:225:51: note: expanded from macro '\ 
    1710 EXPOSE_FUNCTION_ARGS_2' 
    1711 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    1712 ^ 
    1713 program_source:168:9: note: expanded from macro '\ 
    1714 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1715 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1716 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1717 program_source:296:1: error: illegal string literal in 'asm' 
    1718 DECLARE_SHUFFLE_BASE(shuffle_xor) 
    1719 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1720 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1721 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    1722 ^ 
    1723 program_source:226:52: note: expanded from macro '\ 
    1724 EXPOSE_FUNCTION_ARGS_2' 
    1725 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    1726 ^ 
    1727 program_source:168:9: note: expanded from macro '\ 
    1728 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1729 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1730 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1731 program_source:297:1: error: illegal string literal in 'asm' 
    1732 DECLARE_SHUFFLE_BASE(shuffle_up) 
    1733 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1734 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1735 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    1736 ^ 
    1737 program_source:224:38: note: expanded from macro '\ 
    1738 EXPOSE_FUNCTION_ARGS_2' 
    1739 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    1740 ^ 
    1741 program_source:168:9: note: expanded from macro '\ 
    1742 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1743 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1744 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1745 program_source:297:1: error: illegal string literal in 'asm' 
    1746 DECLARE_SHUFFLE_BASE(shuffle_up) 
    1747 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1748 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1749 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    1750 ^ 
    1751 program_source:225:51: note: expanded from macro '\ 
    1752 EXPOSE_FUNCTION_ARGS_2' 
    1753 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    1754 ^ 
    1755 program_source:168:9: note: expanded from macro '\ 
    1756 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1757 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1758 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1759 program_source:297:1: error: illegal string literal in 'asm' 
    1760 DECLARE_SHUFFLE_BASE(shuffle_up) 
    1761 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1762 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1763 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    1764 ^ 
    1765 program_source:226:52: note: expanded from macro '\ 
    1766 EXPOSE_FUNCTION_ARGS_2' 
    1767 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    1768 ^ 
    1769 program_source:168:9: note: expanded from macro '\ 
    1770 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1771 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1772 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1773 program_source:297:1: error: illegal string literal in 'asm' 
    1774 DECLARE_SHUFFLE_BASE(shuffle_up) 
    1775 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1776 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1777 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    1778 ^ 
    1779 program_source:224:38: note: expanded from macro '\ 
    1780 EXPOSE_FUNCTION_ARGS_2' 
    1781 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    1782 ^ 
    1783 program_source:168:9: note: expanded from macro '\ 
    1784 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1785 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1786 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1787 program_source:297:1: error: illegal string literal in 'asm' 
    1788 DECLARE_SHUFFLE_BASE(shuffle_up) 
    1789 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1790 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1791 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    1792 ^ 
    1793 program_source:225:51: note: expanded from macro '\ 
    1794 EXPOSE_FUNCTION_ARGS_2' 
    1795 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    1796 ^ 
    1797 program_source:168:9: note: expanded from macro '\ 
    1798 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1799 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1800 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1801 program_source:297:1: error: illegal string literal in 'asm' 
    1802 DECLARE_SHUFFLE_BASE(shuffle_up) 
    1803 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1804 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1805 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    1806 ^ 
    1807 program_source:226:52: note: expanded from macro '\ 
    1808 EXPOSE_FUNCTION_ARGS_2' 
    1809 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    1810 ^ 
    1811 program_source:168:9: note: expanded from macro '\ 
    1812 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1813 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1814 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1815 program_source:298:1: error: illegal string literal in 'asm' 
    1816 DECLARE_SHUFFLE_BASE(shuffle_down) 
    1817 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1818 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1819 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    1820 ^ 
    1821 program_source:224:38: note: expanded from macro '\ 
    1822 EXPOSE_FUNCTION_ARGS_2' 
    1823 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    1824 ^ 
    1825 program_source:168:9: note: expanded from macro '\ 
    1826 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1827 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1828 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1829 program_source:298:1: error: illegal string literal in 'asm' 
    1830 DECLARE_SHUFFLE_BASE(shuffle_down) 
    1831 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1832 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1833 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    1834 ^ 
    1835 program_source:225:51: note: expanded from macro '\ 
    1836 EXPOSE_FUNCTION_ARGS_2' 
    1837 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    1838 ^ 
    1839 program_source:168:9: note: expanded from macro '\ 
    1840 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1841 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1842 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1843 program_source:298:1: error: illegal string literal in 'asm' 
    1844 DECLARE_SHUFFLE_BASE(shuffle_down) 
    1845 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1846 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1847 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    1848 ^ 
    1849 program_source:226:52: note: expanded from macro '\ 
    1850 EXPOSE_FUNCTION_ARGS_2' 
    1851 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    1852 ^ 
    1853 program_source:168:9: note: expanded from macro '\ 
    1854 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1855 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1856 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1857 program_source:298:1: error: illegal string literal in 'asm' 
    1858 DECLARE_SHUFFLE_BASE(shuffle_down) 
    1859 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1860 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1861 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    1862 ^ 
    1863 program_source:224:38: note: expanded from macro '\ 
    1864 EXPOSE_FUNCTION_ARGS_2' 
    1865 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    1866 ^ 
    1867 program_source:168:9: note: expanded from macro '\ 
    1868 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1869 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1870 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1871 program_source:298:1: error: illegal string literal in 'asm' 
    1872 DECLARE_SHUFFLE_BASE(shuffle_down) 
    1873 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1874 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1875 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    1876 ^ 
    1877 program_source:225:51: note: expanded from macro '\ 
    1878 EXPOSE_FUNCTION_ARGS_2' 
    1879 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    1880 ^ 
    1881 program_source:168:9: note: expanded from macro '\ 
    1882 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1883 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1884 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1885 program_source:298:1: error: illegal string literal in 'asm' 
    1886 DECLARE_SHUFFLE_BASE(shuffle_down) 
    1887 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1888 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1889 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    1890 ^ 
    1891 program_source:226:52: note: expanded from macro '\ 
    1892 EXPOSE_FUNCTION_ARGS_2' 
    1893 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    1894 ^ 
    1895 program_source:168:9: note: expanded from macro '\ 
    1896 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1897 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1898 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1899 program_source:299:1: error: illegal string literal in 'asm' 
    1900 DECLARE_SHUFFLE_BASE(shuffle_rotate_up) 
    1901 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1902 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1903 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    1904 ^ 
    1905 program_source:224:38: note: expanded from macro '\ 
    1906 EXPOSE_FUNCTION_ARGS_2' 
    1907 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    1908 ^ 
    1909 program_source:168:9: note: expanded from macro '\ 
    1910 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1911 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1912 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1913 program_source:299:1: error: illegal string literal in 'asm' 
    1914 DECLARE_SHUFFLE_BASE(shuffle_rotate_up) 
    1915 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1916 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1917 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    1918 ^ 
    1919 program_source:225:51: note: expanded from macro '\ 
    1920 EXPOSE_FUNCTION_ARGS_2' 
    1921 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    1922 ^ 
    1923 program_source:168:9: note: expanded from macro '\ 
    1924 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1925 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1926 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1927 program_source:299:1: error: illegal string literal in 'asm' 
    1928 DECLARE_SHUFFLE_BASE(shuffle_rotate_up) 
    1929 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1930 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1931 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    1932 ^ 
    1933 program_source:226:52: note: expanded from macro '\ 
    1934 EXPOSE_FUNCTION_ARGS_2' 
    1935 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    1936 ^ 
    1937 program_source:168:9: note: expanded from macro '\ 
    1938 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1939 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1940 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1941 program_source:299:1: error: illegal string literal in 'asm' 
    1942 DECLARE_SHUFFLE_BASE(shuffle_rotate_up) 
    1943 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1944 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1945 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    1946 ^ 
    1947 program_source:224:38: note: expanded from macro '\ 
    1948 EXPOSE_FUNCTION_ARGS_2' 
    1949 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    1950 ^ 
    1951 program_source:168:9: note: expanded from macro '\ 
    1952 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1953 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1954 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1955 program_source:299:1: error: illegal string literal in 'asm' 
    1956 DECLARE_SHUFFLE_BASE(shuffle_rotate_up) 
    1957 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1958 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1959 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    1960 ^ 
    1961 program_source:225:51: note: expanded from macro '\ 
    1962 EXPOSE_FUNCTION_ARGS_2' 
    1963 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    1964 ^ 
    1965 program_source:168:9: note: expanded from macro '\ 
    1966 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1967 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1968 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1969 program_source:299:1: error: illegal string literal in 'asm' 
    1970 DECLARE_SHUFFLE_BASE(shuffle_rotate_up) 
    1971 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1972 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1973 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    1974 ^ 
    1975 program_source:226:52: note: expanded from macro '\ 
    1976 EXPOSE_FUNCTION_ARGS_2' 
    1977 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    1978 ^ 
    1979 program_source:168:9: note: expanded from macro '\ 
    1980 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1981 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1982 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1983 program_source:300:1: error: illegal string literal in 'asm' 
    1984 DECLARE_SHUFFLE_BASE(shuffle_rotate_down) 
    1985 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1986 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1987 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    1988 ^ 
    1989 program_source:224:38: note: expanded from macro '\ 
    1990 EXPOSE_FUNCTION_ARGS_2' 
    1991 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    1992 ^ 
    1993 program_source:168:9: note: expanded from macro '\ 
    1994 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1995 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1996 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1997 program_source:300:1: error: illegal string literal in 'asm' 
    1998 DECLARE_SHUFFLE_BASE(shuffle_rotate_down) 
    1999 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2000 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    2001 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    2002 ^ 
    2003 program_source:225:51: note: expanded from macro '\ 
    2004 EXPOSE_FUNCTION_ARGS_2' 
    2005 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    2006 ^ 
    2007 program_source:168:9: note: expanded from macro '\ 
    2008 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    2009 __asm("air." #EXPR "." #AIR_TYPE); \ 
    2010 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2011 program_source:300:1: error: illegal string literal in 'asm' 
    2012 DECLARE_SHUFFLE_BASE(shuffle_rotate_down) 
    2013 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2014 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    2015 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    2016 ^ 
    2017 program_source:226:52: note: expanded from macro '\ 
    2018 EXPOSE_FUNCTION_ARGS_2' 
    2019 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    2020 ^ 
    2021 program_source:168:9: note: expanded from macro '\ 
    2022 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    2023 __asm("air." #EXPR "." #AIR_TYPE); \ 
    2024 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2025 program_source:300:1: error: illegal string literal in 'asm' 
    2026 DECLARE_SHUFFLE_BASE(shuffle_rotate_down) 
    2027 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2028 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    2029 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    2030 ^ 
    2031 program_source:224:38: note: expanded from macro '\ 
    2032 EXPOSE_FUNCTION_ARGS_2' 
    2033 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    2034 ^ 
    2035 program_source:168:9: note: expanded from macro '\ 
    2036 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    2037 __asm("air." #EXPR "." #AIR_TYPE); \ 
    2038 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2039 program_source:300:1: error: illegal string literal in 'asm' 
    2040 DECLARE_SHUFFLE_BASE(shuffle_rotate_down) 
    2041 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2042 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    2043 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    2044 ^ 
    2045 program_source:225:51: note: expanded from macro '\ 
    2046 EXPOSE_FUNCTION_ARGS_2' 
    2047 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    2048 ^ 
    2049 program_source:168:9: note: expanded from macro '\ 
    2050 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    2051 __asm("air." #EXPR "." #AIR_TYPE); \ 
    2052 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2053 program_source:300:1: error: illegal string literal in 'asm' 
    2054 DECLARE_SHUFFLE_BASE(shuffle_rotate_down) 
    2055 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2056 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    2057 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    2058 ^ 
    2059 program_source:226:52: note: expanded from macro '\ 
    2060 EXPOSE_FUNCTION_ARGS_2' 
    2061 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    2062 ^ 
    2063 program_source:168:9: note: expanded from macro '\ 
    2064 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    2065 __asm("air." #EXPR "." #AIR_TYPE); \ 
    2066 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2067 program_source:302:1: error: illegal string literal in 'asm' 
    2068 EXPOSE_FUNCTION_I_ARGS_3(ctz) 
    2069 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2070 program_source:229:40: note: expanded from macro 'EXPOSE_FUNCTION_I_ARGS_3' 
    2071 #define EXPOSE_FUNCTION_I_ARGS_3(EXPR) \ 
    2072 ^ 
    2073 program_source:172:9: note: expanded from macro '\ 
    2074 EXPOSE_FUNCTION_OVERLOAD_ARGS_3' 
    2075 __asm("air." #EXPR "." #AIR_TYPE); \ 
    2076 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2077 program_source:302:1: error: implicit declaration of function '___metal_ctz'
    2078 is invalid in OpenCL 
    2079 program_source:229:40: note: expanded from macro 'EXPOSE_FUNCTION_I_ARGS_3' 
    2080 #define EXPOSE_FUNCTION_I_ARGS_3(EXPR) \ 
    2081 ^ 
    2082 program_source:175:10: note: expanded from macro '\ 
    2083 EXPOSE_FUNCTION_OVERLOAD_ARGS_3' 
    2084 return ___metal_##EXPR(x, false); \ 
    2085 ^ 
    2086 <scratch space>:12:1: note: expanded from here 
    2087 ___metal_ctz 
    2088 ^ 
    2089 program_source:302:1: error: illegal string literal in 'asm' 
    2090 EXPOSE_FUNCTION_I_ARGS_3(ctz) 
    2091 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2092 program_source:230:49: note: expanded from macro 'EXPOSE_FUNCTION_I_ARGS_3' 
    2093 EXPOSE_FUNCTION_OVERLOAD_ARGS_3(EXPR, int, i32) \ 
    2094 ^ 
    2095 program_source:172:9: note: expanded from macro '\ 
    2096 EXPOSE_FUNCTION_OVERLOAD_ARGS_3' 
    2097 __asm("air." #EXPR "." #AIR_TYPE); \ 
    2098 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2099 program_source:302:1: error: implicit declaration of function '___metal_ctz'
    2100 is invalid in OpenCL 
    2101 program_source:230:49: note: expanded from macro 'EXPOSE_FUNCTION_I_ARGS_3' 
    2102 EXPOSE_FUNCTION_OVERLOAD_ARGS_3(EXPR, int, i32) \ 
    2103 ^ 
    2104 program_source:175:10: note: expanded from macro '\ 
    2105 EXPOSE_FUNCTION_OVERLOAD_ARGS_3' 
    2106 return ___metal_##EXPR(x, false); \ 
    2107 ^ 
    2108 <scratch space>:16:1: note: expanded from here 
    2109 ___metal_ctz 
    2110 ^ 
    2111 program_source:317:1: error: implicit declaration of function 'simd_sum' is
    2112 invalid in OpenCL 
    2113 DECLARE_REDUCTION_UNIFORM(sum, reduce_add) 
    2114 ^ 
    2115 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2116 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2117 ^ 
    2118 program_source:305:26: note: expanded from macro '\ 
    2119 DECLARE_I_REDUCTION_UNIFORM' 
    2120 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2121 ^ 
    2122 <scratch space>:17:1: note: expanded from here 
    2123 simd_sum 
    2124 ^ 
    2125 program_source:317:1: error: implicit declaration of function 'simd_sum' is
    2126 invalid in OpenCL 
    2127 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2128 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2129 ^ 
    2130 program_source:305:26: note: expanded from macro '\ 
    2131 DECLARE_I_REDUCTION_UNIFORM' 
    2132 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2133 ^ 
    2134 <scratch space>:17:1: note: expanded from here 
    2135 simd_sum 
    2136 ^ 
    2137 program_source:317:1: error: implicit declaration of function 'simd_sum' is
    2138 invalid in OpenCL 
    2139 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2140 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2141 ^ 
    2142 program_source:308:26: note: expanded from macro '\ 
    2143 DECLARE_F_REDUCTION_UNIFORM' 
    2144 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2145 ^ 
    2146 <scratch space>:19:1: note: expanded from here 
    2147 simd_sum 
    2148 ^ 
    2149 program_source:318:1: error: implicit declaration of function
    2150 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    2151 DECLARE_REDUCTION_UNIFORM(prefix_inclusive_sum, scan_inclusive_add) 
    2152 ^ 
    2153 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2154 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2155 ^ 
    2156 program_source:305:26: note: expanded from macro '\ 
    2157 DECLARE_I_REDUCTION_UNIFORM' 
    2158 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2159 ^ 
    2160 <scratch space>:21:1: note: expanded from here 
    2161 simd_prefix_inclusive_sum 
    2162 ^ 
    2163 program_source:318:1: error: implicit declaration of function
    2164 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    2165 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2166 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2167 ^ 
    2168 program_source:305:26: note: expanded from macro '\ 
    2169 DECLARE_I_REDUCTION_UNIFORM' 
    2170 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2171 ^ 
    2172 <scratch space>:21:1: note: expanded from here 
    2173 simd_prefix_inclusive_sum 
    2174 ^ 
    2175 program_source:318:1: error: implicit declaration of function
    2176 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    2177 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2178 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2179 ^ 
    2180 program_source:308:26: note: expanded from macro '\ 
    2181 DECLARE_F_REDUCTION_UNIFORM' 
    2182 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2183 ^ 
    2184 <scratch space>:23:1: note: expanded from here 
    2185 simd_prefix_inclusive_sum 
    2186 ^ 
    2187 program_source:319:1: error: implicit declaration of function
    2188 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    2189 DECLARE_REDUCTION_UNIFORM(prefix_exclusive_sum, scan_exclusive_add) 
    2190 ^ 
    2191 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2192 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2193 ^ 
    2194 program_source:305:26: note: expanded from macro '\ 
    2195 DECLARE_I_REDUCTION_UNIFORM' 
    2196 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2197 ^ 
    2198 <scratch space>:25:1: note: expanded from here 
    2199 simd_prefix_exclusive_sum 
    2200 ^ 
    2201 program_source:319:1: error: implicit declaration of function
    2202 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    2203 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2204 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2205 ^ 
    2206 program_source:305:26: note: expanded from macro '\ 
    2207 DECLARE_I_REDUCTION_UNIFORM' 
    2208 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2209 ^ 
    2210 <scratch space>:25:1: note: expanded from here 
    2211 simd_prefix_exclusive_sum 
    2212 ^ 
    2213 program_source:319:1: error: implicit declaration of function
    2214 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    2215 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2216 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2217 ^ 
    2218 program_source:308:26: note: expanded from macro '\ 
    2219 DECLARE_F_REDUCTION_UNIFORM' 
    2220 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2221 ^ 
    2222 <scratch space>:27:1: note: expanded from here 
    2223 simd_prefix_exclusive_sum 
    2224 ^ 
    2225 program_source:320:1: error: implicit declaration of function 'simd_min' is
    2226 invalid in OpenCL 
    2227 DECLARE_REDUCTION_UNIFORM(min, reduce_min) 
    2228 ^ 
    2229 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2230 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2231 ^ 
    2232 program_source:305:26: note: expanded from macro '\ 
    2233 DECLARE_I_REDUCTION_UNIFORM' 
    2234 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2235 ^ 
    2236 <scratch space>:29:1: note: expanded from here 
    2237 simd_min 
    2238 ^ 
    2239 program_source:320:1: error: implicit declaration of function 'simd_min' is
    2240 invalid in OpenCL 
    2241 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2242 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2243 ^ 
    2244 program_source:305:26: note: expanded from macro '\ 
    2245 DECLARE_I_REDUCTION_UNIFORM' 
    2246 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2247 ^ 
    2248 <scratch space>:29:1: note: expanded from here 
    2249 simd_min 
    2250 ^ 
    2251 program_source:320:1: error: implicit declaration of function 'simd_min' is
    2252 invalid in OpenCL 
    2253 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2254 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2255 ^ 
    2256 program_source:308:26: note: expanded from macro '\ 
    2257 DECLARE_F_REDUCTION_UNIFORM' 
    2258 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2259 ^ 
    2260 <scratch space>:31:1: note: expanded from here 
    2261 simd_min 
    2262 ^ 
    2263 program_source:321:1: error: implicit declaration of function 'simd_max' is
    2264 invalid in OpenCL 
    2265 DECLARE_REDUCTION_UNIFORM(max, reduce_max) 
    2266 ^ 
    2267 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2268 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2269 ^ 
    2270 program_source:305:26: note: expanded from macro '\ 
    2271 DECLARE_I_REDUCTION_UNIFORM' 
    2272 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2273 ^ 
    2274 <scratch space>:33:1: note: expanded from here 
    2275 simd_max 
    2276 ^ 
    2277 program_source:321:1: error: implicit declaration of function 'simd_max' is
    2278 invalid in OpenCL 
    2279 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2280 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2281 ^ 
    2282 program_source:305:26: note: expanded from macro '\ 
    2283 DECLARE_I_REDUCTION_UNIFORM' 
    2284 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2285 ^ 
    2286 <scratch space>:33:1: note: expanded from here 
    2287 simd_max 
    2288 ^ 
    2289 program_source:321:1: error: implicit declaration of function 'simd_max' is
    2290 invalid in OpenCL 
    2291 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2292 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2293 ^ 
    2294 program_source:308:26: note: expanded from macro '\ 
    2295 DECLARE_F_REDUCTION_UNIFORM' 
    2296 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2297 ^ 
    2298 <scratch space>:35:1: note: expanded from here 
    2299 simd_max 
    2300 ^ 
    2301 program_source:323:1: error: implicit declaration of function 'simd_shuffle'
    2302 is invalid in OpenCL 
    2303 DECLARE_SHUFFLE_UNIFORM(shuffle, shuffle) 
    2304 ^ 
    2305 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2306 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2307 ^ 
    2308 <scratch space>:37:1: note: expanded from here 
    2309 simd_shuffle 
    2310 ^ 
    2311 program_source:323:1: error: implicit declaration of function 'simd_shuffle'
    2312 is invalid in OpenCL 
    2313 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2314 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2315 ^ 
    2316 <scratch space>:37:1: note: expanded from here 
    2317 simd_shuffle 
    2318 ^ 
    2319 program_source:323:1: error: implicit declaration of function 'simd_shuffle'
    2320 is invalid in OpenCL 
    2321 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2322 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2323 ^ 
    2324 <scratch space>:37:1: note: expanded from here 
    2325 simd_shuffle 
    2326 ^ 
    2327 program_source:324:1: error: implicit declaration of function
    2328 'simd_shuffle_xor' is invalid in OpenCL 
    2329 DECLARE_SHUFFLE_UNIFORM(shuffle_xor, shuffle_xor) 
    2330 ^ 
    2331 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2332 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2333 ^ 
    2334 <scratch space>:39:1: note: expanded from here 
    2335 simd_shuffle_xor 
    2336 ^ 
    2337 program_source:324:1: error: implicit declaration of function
    2338 'simd_shuffle_xor' is invalid in OpenCL 
    2339 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2340 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2341 ^ 
    2342 <scratch space>:39:1: note: expanded from here 
    2343 simd_shuffle_xor 
    2344 ^ 
    2345 program_source:324:1: error: implicit declaration of function
    2346 'simd_shuffle_xor' is invalid in OpenCL 
    2347 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2348 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2349 ^ 
    2350 <scratch space>:39:1: note: expanded from here 
    2351 simd_shuffle_xor 
    2352 ^ 
    2353 program_source:325:1: error: implicit declaration of function
    2354 'simd_shuffle_up' is invalid in OpenCL 
    2355 DECLARE_SHUFFLE_UNIFORM(shuffle_up, shuffle_up) 
    2356 ^ 
    2357 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2358 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2359 ^ 
    2360 <scratch space>:41:1: note: expanded from here 
    2361 simd_shuffle_up 
    2362 ^ 
    2363 program_source:325:1: error: implicit declaration of function
    2364 'simd_shuffle_up' is invalid in OpenCL 
    2365 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2366 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2367 ^ 
    2368 <scratch space>:41:1: note: expanded from here 
    2369 simd_shuffle_up 
    2370 ^ 
    2371 program_source:325:1: error: implicit declaration of function
    2372 'simd_shuffle_up' is invalid in OpenCL 
    2373 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2374 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2375 ^ 
    2376 <scratch space>:41:1: note: expanded from here 
    2377 simd_shuffle_up 
    2378 ^ 
    2379 program_source:326:1: error: implicit declaration of function
    2380 'simd_shuffle_down' is invalid in OpenCL 
    2381 DECLARE_SHUFFLE_UNIFORM(shuffle_down, shuffle_down) 
    2382 ^ 
    2383 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2384 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2385 ^ 
    2386 <scratch space>:43:1: note: expanded from here 
    2387 simd_shuffle_down 
    2388 ^ 
    2389 program_source:326:1: error: implicit declaration of function
    2390 'simd_shuffle_down' is invalid in OpenCL 
    2391 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2392 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2393 ^ 
    2394 <scratch space>:43:1: note: expanded from here 
    2395 simd_shuffle_down 
    2396 ^ 
    2397 program_source:326:1: error: implicit declaration of function
    2398 'simd_shuffle_down' is invalid in OpenCL 
    2399 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2400 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2401 ^ 
    2402 <scratch space>:43:1: note: expanded from here 
    2403 simd_shuffle_down 
    2404 ^ 
    2405 program_source:327:1: error: implicit declaration of function
    2406 'simd_shuffle_rotate_down' is invalid in OpenCL 
    2407 DECLARE_SHUFFLE_UNIFORM(shuffle_rotate_down, rotate) 
    2408 ^ 
    2409 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2410 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2411 ^ 
    2412 <scratch space>:45:1: note: expanded from here 
    2413 simd_shuffle_rotate_down 
    2414 ^ 
    2415 program_source:327:1: error: implicit declaration of function
    2416 'simd_shuffle_rotate_down' is invalid in OpenCL 
    2417 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2418 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2419 ^ 
    2420 <scratch space>:45:1: note: expanded from here 
    2421 simd_shuffle_rotate_down 
    2422 ^ 
    2423 program_source:327:1: error: implicit declaration of function
    2424 'simd_shuffle_rotate_down' is invalid in OpenCL 
    2425 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2426 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2427 ^ 
    2428 <scratch space>:45:1: note: expanded from here 
    2429 simd_shuffle_rotate_down 
    2430 ^ 
    2431 program_source:329:1: error: implicit declaration of function 'simd_broadcast'
    2432 is invalid in OpenCL 
    2433 DECLARE_SHUFFLE_UNIFORM(broadcast, broadcast) 
    2434 ^ 
    2435 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2436 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2437 ^ 
    2438 <scratch space>:47:1: note: expanded from here 
    2439 simd_broadcast 
    2440 ^ 
    2441 program_source:329:1: error: implicit declaration of function 'simd_broadcast'
    2442 is invalid in OpenCL 
    2443 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2444 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2445 ^ 
    2446 <scratch space>:47:1: note: expanded from here 
    2447 simd_broadcast 
    2448 ^ 
    2449 program_source:329:1: error: implicit declaration of function 'simd_broadcast'
    2450 is invalid in OpenCL 
    2451 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2452 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2453 ^ 
    2454 <scratch space>:47:1: note: expanded from here 
    2455 simd_broadcast 
    2456 ^ 
    2457 program_source:330:1: error: implicit declaration of function
    2458 'simd_broadcast_first' is invalid in OpenCL 
    2459 DECLARE_REDUCTION_UNIFORM(broadcast_first, broadcast_first) 
    2460 ^ 
    2461 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2462 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2463 ^ 
    2464 program_source:305:26: note: expanded from macro '\ 
    2465 DECLARE_I_REDUCTION_UNIFORM' 
    2466 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2467 ^ 
    2468 <scratch space>:49:1: note: expanded from here 
    2469 simd_broadcast_first 
    2470 ^ 
    2471 program_source:330:1: error: implicit declaration of function
    2472 'simd_broadcast_first' is invalid in OpenCL 
    2473 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2474 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2475 ^ 
    2476 program_source:305:26: note: expanded from macro '\ 
    2477 DECLARE_I_REDUCTION_UNIFORM' 
    2478 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2479 ^ 
    2480 <scratch space>:49:1: note: expanded from here 
    2481 simd_broadcast_first 
    2482 ^ 
    2483 program_source:330:1: error: implicit declaration of function
    2484 'simd_broadcast_first' is invalid in OpenCL 
    2485 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2486 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2487 ^ 
    2488 program_source:308:26: note: expanded from macro '\ 
    2489 DECLARE_F_REDUCTION_UNIFORM' 
    2490 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2491 ^ 
    2492 <scratch space>:51:1: note: expanded from here 
    2493 simd_broadcast_first 
    2494 ^ 
    2495 program_source:345:1: error: implicit declaration of function 'simd_sum' is
    2496 invalid in OpenCL 
    2497 DECLARE_REDUCTION_NON_UNIFORM(sum, reduce_add) 
    2498 ^ 
    2499 program_source:338:60: note: expanded from macro
    2500 'DECLARE_REDUCTION_NON_UNIFORM' 
    2501 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2502 ^ 
    2503 program_source:333:26: note: expanded from macro '\ 
    2504 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2505 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2506 ^ 
    2507 <scratch space>:53:1: note: expanded from here 
    2508 simd_sum 
    2509 ^ 
    2510 program_source:345:1: error: implicit declaration of function 'simd_sum' is
    2511 invalid in OpenCL 
    2512 program_source:338:60: note: expanded from macro
    2513 'DECLARE_REDUCTION_NON_UNIFORM' 
    2514 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2515 ^ 
    2516 program_source:333:26: note: expanded from macro '\ 
    2517 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2518 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2519 ^ 
    2520 <scratch space>:53:1: note: expanded from here 
    2521 simd_sum 
    2522 ^ 
    2523 program_source:345:1: error: implicit declaration of function 'simd_sum' is
    2524 invalid in OpenCL 
    2525 program_source:339:54: note: expanded from macro
    2526 'DECLARE_REDUCTION_NON_UNIFORM' 
    2527 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2528 ^ 
    2529 program_source:336:26: note: expanded from macro '\ 
    2530 DECLARE_F_REDUCTION_NON_UNIFORM' 
    2531 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2532 ^ 
    2533 <scratch space>:55:1: note: expanded from here 
    2534 simd_sum 
    2535 ^ 
    2536 program_source:346:1: error: implicit declaration of function
    2537 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    2538 DECLARE_REDUCTION_NON_UNIFORM(prefix_inclusive_sum, scan_inclusive_add) 
    2539 ^ 
    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>:57:1: note: expanded from here 
    2549 simd_prefix_inclusive_sum 
    2550 ^ 
    2551 program_source:346:1: error: implicit declaration of function
    2552 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    2553 program_source:338:60: note: expanded from macro
    2554 'DECLARE_REDUCTION_NON_UNIFORM' 
    2555 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2556 ^ 
    2557 program_source:333:26: note: expanded from macro '\ 
    2558 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2559 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2560 ^ 
    2561 <scratch space>:57:1: note: expanded from here 
    2562 simd_prefix_inclusive_sum 
    2563 ^ 
    2564 program_source:346:1: error: implicit declaration of function
    2565 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    2566 program_source:339:54: note: expanded from macro
    2567 'DECLARE_REDUCTION_NON_UNIFORM' 
    2568 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2569 ^ 
    2570 program_source:336:26: note: expanded from macro '\ 
    2571 DECLARE_F_REDUCTION_NON_UNIFORM' 
    2572 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2573 ^ 
    2574 <scratch space>:59:1: note: expanded from here 
    2575 simd_prefix_inclusive_sum 
    2576 ^ 
    2577 program_source:347:1: error: implicit declaration of function
    2578 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    2579 DECLARE_REDUCTION_NON_UNIFORM(prefix_exclusive_sum, scan_exclusive_add) 
    2580 ^ 
    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>:61:1: note: expanded from here 
    2590 simd_prefix_exclusive_sum 
    2591 ^ 
    2592 program_source:347:1: error: implicit declaration of function
    2593 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    2594 program_source:338:60: note: expanded from macro
    2595 'DECLARE_REDUCTION_NON_UNIFORM' 
    2596 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2597 ^ 
    2598 program_source:333:26: note: expanded from macro '\ 
    2599 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2600 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2601 ^ 
    2602 <scratch space>:61:1: note: expanded from here 
    2603 simd_prefix_exclusive_sum 
    2604 ^ 
    2605 program_source:347:1: error: implicit declaration of function
    2606 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    2607 program_source:339:54: note: expanded from macro
    2608 'DECLARE_REDUCTION_NON_UNIFORM' 
    2609 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2610 ^ 
    2611 program_source:336:26: note: expanded from macro '\ 
    2612 DECLARE_F_REDUCTION_NON_UNIFORM' 
    2613 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2614 ^ 
    2615 <scratch space>:63:1: note: expanded from here 
    2616 simd_prefix_exclusive_sum 
    2617 ^ 
    2618 program_source:348:1: error: implicit declaration of function 'simd_min' is
    2619 invalid in OpenCL 
    2620 DECLARE_REDUCTION_NON_UNIFORM(min, reduce_min) 
    2621 ^ 
    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>:65:1: note: expanded from here 
    2631 simd_min 
    2632 ^ 
    2633 program_source:348:1: error: implicit declaration of function 'simd_min' is
    2634 invalid in OpenCL 
    2635 program_source:338:60: note: expanded from macro
    2636 'DECLARE_REDUCTION_NON_UNIFORM' 
    2637 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2638 ^ 
    2639 program_source:333:26: note: expanded from macro '\ 
    2640 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2641 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2642 ^ 
    2643 <scratch space>:65:1: note: expanded from here 
    2644 simd_min 
    2645 ^ 
    2646 program_source:348:1: error: implicit declaration of function 'simd_min' is
    2647 invalid in OpenCL 
    2648 program_source:339:54: note: expanded from macro
    2649 'DECLARE_REDUCTION_NON_UNIFORM' 
    2650 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2651 ^ 
    2652 program_source:336:26: note: expanded from macro '\ 
    2653 DECLARE_F_REDUCTION_NON_UNIFORM' 
    2654 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2655 ^ 
    2656 <scratch space>:67:1: note: expanded from here 
    2657 simd_min 
    2658 ^ 
    2659 program_source:349:1: error: implicit declaration of function 'simd_max' is
    2660 invalid in OpenCL 
    2661 DECLARE_REDUCTION_NON_UNIFORM(max, reduce_max) 
    2662 ^ 
    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>:69:1: note: expanded from here 
    2672 simd_max 
    2673 ^ 
    2674 program_source:349:1: error: implicit declaration of function 'simd_max' is
    2675 invalid in OpenCL 
    2676 program_source:338:60: note: expanded from macro
    2677 'DECLARE_REDUCTION_NON_UNIFORM' 
    2678 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2679 ^ 
    2680 program_source:333:26: note: expanded from macro '\ 
    2681 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2682 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2683 ^ 
    2684 <scratch space>:69:1: note: expanded from here 
    2685 simd_max 
    2686 ^ 
    2687 program_source:349:1: error: implicit declaration of function 'simd_max' is
    2688 invalid in OpenCL 
    2689 program_source:339:54: note: expanded from macro
    2690 'DECLARE_REDUCTION_NON_UNIFORM' 
    2691 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2692 ^ 
    2693 program_source:336:26: note: expanded from macro '\ 
    2694 DECLARE_F_REDUCTION_NON_UNIFORM' 
    2695 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2696 ^ 
    2697 <scratch space>:71:1: note: expanded from here 
    2698 simd_max 
    2699 ^ 
    2700 program_source:351:1: error: implicit declaration of function 'simd_product'
    2701 is invalid in OpenCL 
    2702 DECLARE_REDUCTION_NON_UNIFORM(product, reduce_mul) 
    2703 ^ 
    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>:73:1: note: expanded from here 
    2713 simd_product 
    2714 ^ 
    2715 program_source:351:1: error: implicit declaration of function 'simd_product'
    2716 is invalid in OpenCL 
    2717 program_source:338:60: note: expanded from macro
    2718 'DECLARE_REDUCTION_NON_UNIFORM' 
    2719 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2720 ^ 
    2721 program_source:333:26: note: expanded from macro '\ 
    2722 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2723 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2724 ^ 
    2725 <scratch space>:73:1: note: expanded from here 
    2726 simd_product 
    2727 ^ 
    2728 program_source:351:1: error: implicit declaration of function 'simd_product'
    2729 is invalid in OpenCL 
    2730 program_source:339:54: note: expanded from macro
    2731 'DECLARE_REDUCTION_NON_UNIFORM' 
    2732 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2733 ^ 
    2734 program_source:336:26: note: expanded from macro '\ 
    2735 DECLARE_F_REDUCTION_NON_UNIFORM' 
    2736 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2737 ^ 
    2738 <scratch space>:75:1: note: expanded from here 
    2739 simd_product 
    2740 ^ 
    2741 program_source:352:1: error: implicit declaration of function
    2742 'simd_prefix_inclusive_product' is invalid in OpenCL 
    2743 DECLARE_REDUCTION_NON_UNIFORM(prefix_inclusive_product, scan_inclusive_mul) 
    2744 ^ 
    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>:77:1: note: expanded from here 
    2754 simd_prefix_inclusive_product 
    2755 ^ 
    2756 program_source:352:1: error: implicit declaration of function
    2757 'simd_prefix_inclusive_product' is invalid in OpenCL 
    2758 program_source:338:60: note: expanded from macro
    2759 'DECLARE_REDUCTION_NON_UNIFORM' 
    2760 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2761 ^ 
    2762 program_source:333:26: note: expanded from macro '\ 
    2763 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2764 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2765 ^ 
    2766 <scratch space>:77:1: note: expanded from here 
    2767 simd_prefix_inclusive_product 
    2768 ^ 
    2769 program_source:352:1: error: implicit declaration of function
    2770 'simd_prefix_inclusive_product' is invalid in OpenCL 
    2771 program_source:339:54: note: expanded from macro
    2772 'DECLARE_REDUCTION_NON_UNIFORM' 
    2773 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2774 ^ 
    2775 program_source:336:26: note: expanded from macro '\ 
    2776 DECLARE_F_REDUCTION_NON_UNIFORM' 
    2777 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2778 ^ 
    2779 <scratch space>:79:1: note: expanded from here 
    2780 simd_prefix_inclusive_product 
    2781 ^ 
    2782 program_source:353:1: error: implicit declaration of function
    2783 'simd_prefix_exclusive_product' is invalid in OpenCL 
    2784 DECLARE_REDUCTION_NON_UNIFORM(prefix_exclusive_product, scan_exclusive_mul) 
    2785 ^ 
    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>:81:1: note: expanded from here 
    2795 simd_prefix_exclusive_product 
    2796 ^ 
    2797 program_source:353:1: error: implicit declaration of function
    2798 'simd_prefix_exclusive_product' is invalid in OpenCL 
    2799 program_source:338:60: note: expanded from macro
    2800 'DECLARE_REDUCTION_NON_UNIFORM' 
    2801 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2802 ^ 
    2803 program_source:333:26: note: expanded from macro '\ 
    2804 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2805 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2806 ^ 
    2807 <scratch space>:81:1: note: expanded from here 
    2808 simd_prefix_exclusive_product 
    2809 ^ 
    2810 program_source:353:1: error: implicit declaration of function
    2811 'simd_prefix_exclusive_product' is invalid in OpenCL 
    2812 program_source:339:54: note: expanded from macro
    2813 'DECLARE_REDUCTION_NON_UNIFORM' 
    2814 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2815 ^ 
    2816 program_source:336:26: note: expanded from macro '\ 
    2817 DECLARE_F_REDUCTION_NON_UNIFORM' 
    2818 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2819 ^ 
    2820 <scratch space>:83:1: note: expanded from here 
    2821 simd_prefix_exclusive_product 
    2822 ^ 
    2823 program_source:354:1: error: implicit declaration of function 'simd_and' is
    2824 invalid in OpenCL 
    2825 DECLARE_I_REDUCTION_NON_UNIFORM(and, reduce_and) 
    2826 ^ 
    2827 program_source:333:26: note: expanded from macro
    2828 'DECLARE_I_REDUCTION_NON_UNIFORM' 
    2829 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2830 ^ 
    2831 <scratch space>:85:1: note: expanded from here 
    2832 simd_and 
    2833 ^ 
    2834 program_source:354:1: error: implicit declaration of function 'simd_and' is
    2835 invalid in OpenCL 
    2836 program_source:333:26: note: expanded from macro
    2837 'DECLARE_I_REDUCTION_NON_UNIFORM' 
    2838 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2839 ^ 
    2840 <scratch space>:85:1: note: expanded from here 
    2841 simd_and 
    2842 ^ 
    2843 program_source:355:1: error: implicit declaration of function 'simd_or' is
    2844 invalid in OpenCL 
    2845 DECLARE_I_REDUCTION_NON_UNIFORM(or, reduce_or) 
    2846 ^ 
    2847 program_source:333:26: note: expanded from macro
    2848 'DECLARE_I_REDUCTION_NON_UNIFORM' 
    2849 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2850 ^ 
    2851 <scratch space>:87:1: note: expanded from here 
    2852 simd_or 
    2853 ^ 
    2854 program_source:355:1: error: implicit declaration of function 'simd_or' is
    2855 invalid in OpenCL 
    2856 program_source:333:26: note: expanded from macro
    2857 'DECLARE_I_REDUCTION_NON_UNIFORM' 
    2858 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2859 ^ 
    2860 <scratch space>:87:1: note: expanded from here 
    2861 simd_or 
    2862 ^ 
    2863 program_source:356:1: error: implicit declaration of function 'simd_xor' is
    2864 invalid in OpenCL 
    2865 DECLARE_I_REDUCTION_NON_UNIFORM(xor, reduce_xor) 
    2866 ^ 
    2867 program_source:333:26: note: expanded from macro
    2868 'DECLARE_I_REDUCTION_NON_UNIFORM' 
    2869 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2870 ^ 
    2871 <scratch space>:89:1: note: expanded from here 
    2872 simd_xor 
    2873 ^ 
    2874 program_source:356:1: error: implicit declaration of function 'simd_xor' is
    2875 invalid in OpenCL 
    2876 program_source:333:26: note: expanded from macro
    2877 'DECLARE_I_REDUCTION_NON_UNIFORM' 
    2878 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2879 ^ 
    2880 <scratch space>:89:1: note: expanded from here 
    2881 simd_xor 
    2882 ^ 
    2883 program_source:358:1: error: implicit declaration of function 'simd_broadcast'
    2884 is invalid in OpenCL 
    2885 DECLARE_SHUFFLE_NON_UNIFORM(broadcast, broadcast) 
    2886 ^ 
    2887 program_source:343:24: note: expanded from macro 'DECLARE_SHUFFLE_NON_UNIFORM' 
    2888 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2889 ^ 
    2890 <scratch space>:91:1: note: expanded from here 
    2891 simd_broadcast 
    2892 ^ 
    2893 program_source:358:1: error: implicit declaration of function 'simd_broadcast'
    2894 is invalid in OpenCL 
    2895 program_source:343:24: note: expanded from macro 'DECLARE_SHUFFLE_NON_UNIFORM' 
    2896 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2897 ^ 
    2898 <scratch space>:91:1: note: expanded from here 
    2899 simd_broadcast 
    2900 ^ 
    2901 program_source:358:1: error: implicit declaration of function 'simd_broadcast'
    2902 is invalid in OpenCL 
    2903 program_source:343:24: note: expanded from macro 'DECLARE_SHUFFLE_NON_UNIFORM' 
    2904 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2905 ^ 
    2906 <scratch space>:91:1: note: expanded from here 
    2907 simd_broadcast 
    2908 ^ 
    2909 program_source:391:1: error: implicit declaration of function 'quad_sum' is
    2910 invalid in OpenCL 
    2911 DECLARE_REDUCTION_CLUSTERED(sum, reduce_add) 
    2912 ^ 
    2913 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    2914 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2915 ^ 
    2916 program_source:360:60: note: expanded from macro '\ 
    2917 DECLARE_I_REDUCTION_CLUSTERED' 
    2918 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2919 ^ 
    2920 program_source:245:71: note: expanded from macro '\ 
    2921 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    2922 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    2923 ^ 
    2924 program_source:196:12: note: expanded from macro '\ 
    2925 OVERLOAD_CLUSTERED_ARGS_1' 
    2926 return quad_##METAL_SUFFIX(x); \ 
    2927 ^ 
    2928 <scratch space>:94:1: note: expanded from here 
    2929 quad_sum 
    2930 ^ 
    2931 program_source:391:1: error: implicit declaration of function 'simd_sum' is
    2932 invalid in OpenCL 
    2933 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    2934 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2935 ^ 
    2936 program_source:360:60: note: expanded from macro '\ 
    2937 DECLARE_I_REDUCTION_CLUSTERED' 
    2938 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2939 ^ 
    2940 program_source:245:71: note: expanded from macro '\ 
    2941 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    2942 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    2943 ^ 
    2944 program_source:198:12: note: expanded from macro '\ 
    2945 OVERLOAD_CLUSTERED_ARGS_1' 
    2946 return simd_##METAL_SUFFIX(x); \ 
    2947 ^ 
    2948 <scratch space>:95:1: note: expanded from here 
    2949 simd_sum 
    2950 ^ 
    2951 program_source:391:1: error: implicit declaration of function 'quad_sum' is
    2952 invalid in OpenCL 
    2953 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    2954 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2955 ^ 
    2956 program_source:360:60: note: expanded from macro '\ 
    2957 DECLARE_I_REDUCTION_CLUSTERED' 
    2958 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2959 ^ 
    2960 program_source:246:59: note: expanded from macro '\ 
    2961 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    2962 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    2963 ^ 
    2964 program_source:196:12: note: expanded from macro '\ 
    2965 OVERLOAD_CLUSTERED_ARGS_1' 
    2966 return quad_##METAL_SUFFIX(x); \ 
    2967 ^ 
    2968 <scratch space>:96:1: note: expanded from here 
    2969 quad_sum 
    2970 ^ 
    2971 program_source:391:1: error: implicit declaration of function 'simd_sum' is
    2972 invalid in OpenCL 
    2973 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    2974 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2975 ^ 
    2976 program_source:360:60: note: expanded from macro '\ 
    2977 DECLARE_I_REDUCTION_CLUSTERED' 
    2978 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2979 ^ 
    2980 program_source:246:59: note: expanded from macro '\ 
    2981 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    2982 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    2983 ^ 
    2984 program_source:198:12: note: expanded from macro '\ 
    2985 OVERLOAD_CLUSTERED_ARGS_1' 
    2986 return simd_##METAL_SUFFIX(x); \ 
    2987 ^ 
    2988 <scratch space>:97:1: note: expanded from here 
    2989 simd_sum 
    2990 ^ 
    2991 program_source:391:1: error: implicit declaration of function 'quad_sum' is
    2992 invalid in OpenCL 
    2993 program_source:385:52: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    2994 DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2995 ^ 
    2996 program_source:363:60: note: expanded from macro '\ 
    2997 DECLARE_F_REDUCTION_CLUSTERED' 
    2998 #define DECLARE_F_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    2999 ^ 
    3000 program_source:249:71: note: expanded from macro '\ 
    3001 BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1' 
    3002 #define BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3003 ^ 
    3004 program_source:196:12: note: expanded from macro '\ 
    3005 OVERLOAD_CLUSTERED_ARGS_1' 
    3006 return quad_##METAL_SUFFIX(x); \ 
    3007 ^ 
    3008 <scratch space>:99:1: note: expanded from here 
    3009 quad_sum 
    3010 ^ 
    3011 program_source:391:1: error: implicit declaration of function 'simd_sum' is
    3012 invalid in OpenCL 
    3013 program_source:385:52: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3014 DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3015 ^ 
    3016 program_source:363:60: note: expanded from macro '\ 
    3017 DECLARE_F_REDUCTION_CLUSTERED' 
    3018 #define DECLARE_F_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3019 ^ 
    3020 program_source:249:71: note: expanded from macro '\ 
    3021 BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1' 
    3022 #define BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3023 ^ 
    3024 program_source:198:12: note: expanded from macro '\ 
    3025 OVERLOAD_CLUSTERED_ARGS_1' 
    3026 return simd_##METAL_SUFFIX(x); \ 
    3027 ^ 
    3028 <scratch space>:100:1: note: expanded from here 
    3029 simd_sum 
    3030 ^ 
    3031 program_source:392:1: error: implicit declaration of function 'quad_min' is
    3032 invalid in OpenCL 
    3033 DECLARE_REDUCTION_CLUSTERED(min, reduce_min) 
    3034 ^ 
    3035 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3036 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3037 ^ 
    3038 program_source:360:60: note: expanded from macro '\ 
    3039 DECLARE_I_REDUCTION_CLUSTERED' 
    3040 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3041 ^ 
    3042 program_source:245:71: note: expanded from macro '\ 
    3043 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3044 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3045 ^ 
    3046 program_source:196:12: note: expanded from macro '\ 
    3047 OVERLOAD_CLUSTERED_ARGS_1' 
    3048 return quad_##METAL_SUFFIX(x); \ 
    3049 ^ 
    3050 <scratch space>:102:1: note: expanded from here 
    3051 quad_min 
    3052 ^ 
    3053 program_source:392:1: error: implicit declaration of function 'simd_min' is
    3054 invalid in OpenCL 
    3055 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3056 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3057 ^ 
    3058 program_source:360:60: note: expanded from macro '\ 
    3059 DECLARE_I_REDUCTION_CLUSTERED' 
    3060 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3061 ^ 
    3062 program_source:245:71: note: expanded from macro '\ 
    3063 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3064 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3065 ^ 
    3066 program_source:198:12: note: expanded from macro '\ 
    3067 OVERLOAD_CLUSTERED_ARGS_1' 
    3068 return simd_##METAL_SUFFIX(x); \ 
    3069 ^ 
    3070 <scratch space>:103:1: note: expanded from here 
    3071 simd_min 
    3072 ^ 
    3073 program_source:392:1: error: implicit declaration of function 'quad_min' is
    3074 invalid in OpenCL 
    3075 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3076 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3077 ^ 
    3078 program_source:360:60: note: expanded from macro '\ 
    3079 DECLARE_I_REDUCTION_CLUSTERED' 
    3080 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3081 ^ 
    3082 program_source:246:59: note: expanded from macro '\ 
    3083 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3084 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3085 ^ 
    3086 program_source:196:12: note: expanded from macro '\ 
    3087 OVERLOAD_CLUSTERED_ARGS_1' 
    3088 return quad_##METAL_SUFFIX(x); \ 
    3089 ^ 
    3090 <scratch space>:104:1: note: expanded from here 
    3091 quad_min 
    3092 ^ 
    3093 program_source:392:1: error: implicit declaration of function 'simd_min' is
    3094 invalid in OpenCL 
    3095 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3096 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3097 ^ 
    3098 program_source:360:60: note: expanded from macro '\ 
    3099 DECLARE_I_REDUCTION_CLUSTERED' 
    3100 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3101 ^ 
    3102 program_source:246:59: note: expanded from macro '\ 
    3103 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3104 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3105 ^ 
    3106 program_source:198:12: note: expanded from macro '\ 
    3107 OVERLOAD_CLUSTERED_ARGS_1' 
    3108 return simd_##METAL_SUFFIX(x); \ 
    3109 ^ 
    3110 <scratch space>:105:1: note: expanded from here 
    3111 simd_min 
    3112 ^ 
    3113 program_source:392:1: error: implicit declaration of function 'quad_min' is
    3114 invalid in OpenCL 
    3115 program_source:385:52: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3116 DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3117 ^ 
    3118 program_source:363:60: note: expanded from macro '\ 
    3119 DECLARE_F_REDUCTION_CLUSTERED' 
    3120 #define DECLARE_F_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3121 ^ 
    3122 program_source:249:71: note: expanded from macro '\ 
    3123 BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1' 
    3124 #define BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3125 ^ 
    3126 program_source:196:12: note: expanded from macro '\ 
    3127 OVERLOAD_CLUSTERED_ARGS_1' 
    3128 return quad_##METAL_SUFFIX(x); \ 
    3129 ^ 
    3130 <scratch space>:107:1: note: expanded from here 
    3131 quad_min 
    3132 ^ 
    3133 program_source:392:1: error: implicit declaration of function 'simd_min' is
    3134 invalid in OpenCL 
    3135 program_source:385:52: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3136 DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3137 ^ 
    3138 program_source:363:60: note: expanded from macro '\ 
    3139 DECLARE_F_REDUCTION_CLUSTERED' 
    3140 #define DECLARE_F_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3141 ^ 
    3142 program_source:249:71: note: expanded from macro '\ 
    3143 BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1' 
    3144 #define BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3145 ^ 
    3146 program_source:198:12: note: expanded from macro '\ 
    3147 OVERLOAD_CLUSTERED_ARGS_1' 
    3148 return simd_##METAL_SUFFIX(x); \ 
    3149 ^ 
    3150 <scratch space>:108:1: note: expanded from here 
    3151 simd_min 
    3152 ^ 
    3153 program_source:393:1: error: implicit declaration of function 'quad_max' is
    3154 invalid in OpenCL 
    3155 DECLARE_REDUCTION_CLUSTERED(max, reduce_max) 
    3156 ^ 
    3157 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3158 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3159 ^ 
    3160 program_source:360:60: note: expanded from macro '\ 
    3161 DECLARE_I_REDUCTION_CLUSTERED' 
    3162 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3163 ^ 
    3164 program_source:245:71: note: expanded from macro '\ 
    3165 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3166 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3167 ^ 
    3168 program_source:196:12: note: expanded from macro '\ 
    3169 OVERLOAD_CLUSTERED_ARGS_1' 
    3170 return quad_##METAL_SUFFIX(x); \ 
    3171 ^ 
    3172 <scratch space>:110:1: note: expanded from here 
    3173 quad_max 
    3174 ^ 
    3175 program_source:393:1: error: implicit declaration of function 'simd_max' is
    3176 invalid in OpenCL 
    3177 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3178 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3179 ^ 
    3180 program_source:360:60: note: expanded from macro '\ 
    3181 DECLARE_I_REDUCTION_CLUSTERED' 
    3182 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3183 ^ 
    3184 program_source:245:71: note: expanded from macro '\ 
    3185 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3186 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3187 ^ 
    3188 program_source:198:12: note: expanded from macro '\ 
    3189 OVERLOAD_CLUSTERED_ARGS_1' 
    3190 return simd_##METAL_SUFFIX(x); \ 
    3191 ^ 
    3192 <scratch space>:111:1: note: expanded from here 
    3193 simd_max 
    3194 ^ 
    3195 program_source:393:1: error: implicit declaration of function 'quad_max' is
    3196 invalid in OpenCL 
    3197 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3198 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3199 ^ 
    3200 program_source:360:60: note: expanded from macro '\ 
    3201 DECLARE_I_REDUCTION_CLUSTERED' 
    3202 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3203 ^ 
    3204 program_source:246:59: note: expanded from macro '\ 
    3205 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3206 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3207 ^ 
    3208 program_source:196:12: note: expanded from macro '\ 
    3209 OVERLOAD_CLUSTERED_ARGS_1' 
    3210 return quad_##METAL_SUFFIX(x); \ 
    3211 ^ 
    3212 <scratch space>:112:1: note: expanded from here 
    3213 quad_max 
    3214 ^ 
    3215 program_source:393:1: error: implicit declaration of function 'simd_max' is
    3216 invalid in OpenCL 
    3217 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3218 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3219 ^ 
    3220 program_source:360:60: note: expanded from macro '\ 
    3221 DECLARE_I_REDUCTION_CLUSTERED' 
    3222 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3223 ^ 
    3224 program_source:246:59: note: expanded from macro '\ 
    3225 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3226 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3227 ^ 
    3228 program_source:198:12: note: expanded from macro '\ 
    3229 OVERLOAD_CLUSTERED_ARGS_1' 
    3230 return simd_##METAL_SUFFIX(x); \ 
    3231 ^ 
    3232 <scratch space>:113:1: note: expanded from here 
    3233 simd_max 
    3234 ^ 
    3235 program_source:393:1: error: implicit declaration of function 'quad_max' is
    3236 invalid in OpenCL 
    3237 program_source:385:52: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3238 DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3239 ^ 
    3240 program_source:363:60: note: expanded from macro '\ 
    3241 DECLARE_F_REDUCTION_CLUSTERED' 
    3242 #define DECLARE_F_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3243 ^ 
    3244 program_source:249:71: note: expanded from macro '\ 
    3245 BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1' 
    3246 #define BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3247 ^ 
    3248 program_source:196:12: note: expanded from macro '\ 
    3249 OVERLOAD_CLUSTERED_ARGS_1' 
    3250 return quad_##METAL_SUFFIX(x); \ 
    3251 ^ 
    3252 <scratch space>:115:1: note: expanded from here 
    3253 quad_max 
    3254 ^ 
    3255 program_source:393:1: error: implicit declaration of function 'simd_max' is
    3256 invalid in OpenCL 
    3257 program_source:385:52: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3258 DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3259 ^ 
    3260 program_source:363:60: note: expanded from macro '\ 
    3261 DECLARE_F_REDUCTION_CLUSTERED' 
    3262 #define DECLARE_F_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3263 ^ 
    3264 program_source:249:71: note: expanded from macro '\ 
    3265 BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1' 
    3266 #define BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3267 ^ 
    3268 program_source:198:12: note: expanded from macro '\ 
    3269 OVERLOAD_CLUSTERED_ARGS_1' 
    3270 return simd_##METAL_SUFFIX(x); \ 
    3271 ^ 
    3272 <scratch space>:116:1: note: expanded from here 
    3273 simd_max 
    3274 ^ 
    3275 program_source:395:1: error: implicit declaration of function 'quad_product'
    3276 is invalid in OpenCL 
    3277 DECLARE_REDUCTION_CLUSTERED(product, reduce_mul) 
    3278 ^ 
    3279 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3280 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3281 ^ 
    3282 program_source:360:60: note: expanded from macro '\ 
    3283 DECLARE_I_REDUCTION_CLUSTERED' 
    3284 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3285 ^ 
    3286 program_source:245:71: note: expanded from macro '\ 
    3287 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3288 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3289 ^ 
    3290 program_source:196:12: note: expanded from macro '\ 
    3291 OVERLOAD_CLUSTERED_ARGS_1' 
    3292 return quad_##METAL_SUFFIX(x); \ 
    3293 ^ 
    3294 <scratch space>:118:1: note: expanded from here 
    3295 quad_product 
    3296 ^ 
    3297 program_source:395:1: error: implicit declaration of function 'simd_product'
    3298 is invalid in OpenCL 
    3299 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3300 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3301 ^ 
    3302 program_source:360:60: note: expanded from macro '\ 
    3303 DECLARE_I_REDUCTION_CLUSTERED' 
    3304 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3305 ^ 
    3306 program_source:245:71: note: expanded from macro '\ 
    3307 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3308 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3309 ^ 
    3310 program_source:198:12: note: expanded from macro '\ 
    3311 OVERLOAD_CLUSTERED_ARGS_1' 
    3312 return simd_##METAL_SUFFIX(x); \ 
    3313 ^ 
    3314 <scratch space>:119:1: note: expanded from here 
    3315 simd_product 
    3316 ^ 
    3317 program_source:395:1: error: implicit declaration of function 'quad_product'
    3318 is invalid in OpenCL 
    3319 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3320 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3321 ^ 
    3322 program_source:360:60: note: expanded from macro '\ 
    3323 DECLARE_I_REDUCTION_CLUSTERED' 
    3324 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3325 ^ 
    3326 program_source:246:59: note: expanded from macro '\ 
    3327 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3328 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3329 ^ 
    3330 program_source:196:12: note: expanded from macro '\ 
    3331 OVERLOAD_CLUSTERED_ARGS_1' 
    3332 return quad_##METAL_SUFFIX(x); \ 
    3333 ^ 
    3334 <scratch space>:120:1: note: expanded from here 
    3335 quad_product 
    3336 ^ 
    3337 program_source:395:1: error: implicit declaration of function 'simd_product'
    3338 is invalid in OpenCL 
    3339 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3340 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3341 ^ 
    3342 program_source:360:60: note: expanded from macro '\ 
    3343 DECLARE_I_REDUCTION_CLUSTERED' 
    3344 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3345 ^ 
    3346 program_source:246:59: note: expanded from macro '\ 
    3347 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3348 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3349 ^ 
    3350 program_source:198:12: note: expanded from macro '\ 
    3351 OVERLOAD_CLUSTERED_ARGS_1' 
    3352 return simd_##METAL_SUFFIX(x); \ 
    3353 ^ 
    3354 <scratch space>:121:1: note: expanded from here 
    3355 simd_product 
    3356 ^ 
    3357 program_source:395:1: error: implicit declaration of function 'quad_product'
    3358 is invalid in OpenCL 
    3359 program_source:385:52: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3360 DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3361 ^ 
    3362 program_source:363:60: note: expanded from macro '\ 
    3363 DECLARE_F_REDUCTION_CLUSTERED' 
    3364 #define DECLARE_F_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3365 ^ 
    3366 program_source:249:71: note: expanded from macro '\ 
    3367 BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1' 
    3368 #define BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3369 ^ 
    3370 program_source:196:12: note: expanded from macro '\ 
    3371 OVERLOAD_CLUSTERED_ARGS_1' 
    3372 return quad_##METAL_SUFFIX(x); \ 
    3373 ^ 
    3374 <scratch space>:123:1: note: expanded from here 
    3375 quad_product 
    3376 ^ 
    3377 program_source:395:1: error: implicit declaration of function 'simd_product'
    3378 is invalid in OpenCL 
    3379 program_source:385:52: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3380 DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3381 ^ 
    3382 program_source:363:60: note: expanded from macro '\ 
    3383 DECLARE_F_REDUCTION_CLUSTERED' 
    3384 #define DECLARE_F_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3385 ^ 
    3386 program_source:249:71: note: expanded from macro '\ 
    3387 BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1' 
    3388 #define BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3389 ^ 
    3390 program_source:198:12: note: expanded from macro '\ 
    3391 OVERLOAD_CLUSTERED_ARGS_1' 
    3392 return simd_##METAL_SUFFIX(x); \ 
    3393 ^ 
    3394 <scratch space>:124:1: note: expanded from here 
    3395 simd_product 
    3396 ^ 
    3397 program_source:396:1: error: implicit declaration of function 'quad_and' is
    3398 invalid in OpenCL 
    3399 DECLARE_I_REDUCTION_CLUSTERED(and, reduce_and) 
    3400 ^ 
    3401 program_source:360:60: note: expanded from macro
    3402 'DECLARE_I_REDUCTION_CLUSTERED' 
    3403 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3404 ^ 
    3405 program_source:245:71: note: expanded from macro '\ 
    3406 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3407 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3408 ^ 
    3409 program_source:196:12: note: expanded from macro '\ 
    3410 OVERLOAD_CLUSTERED_ARGS_1' 
    3411 return quad_##METAL_SUFFIX(x); \ 
    3412 ^ 
    3413 <scratch space>:126:1: note: expanded from here 
    3414 quad_and 
    3415 ^ 
    3416 program_source:396:1: error: implicit declaration of function 'simd_and' is
    3417 invalid in OpenCL 
    3418 program_source:360:60: note: expanded from macro
    3419 'DECLARE_I_REDUCTION_CLUSTERED' 
    3420 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3421 ^ 
    3422 program_source:245:71: note: expanded from macro '\ 
    3423 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3424 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3425 ^ 
    3426 program_source:198:12: note: expanded from macro '\ 
    3427 OVERLOAD_CLUSTERED_ARGS_1' 
    3428 return simd_##METAL_SUFFIX(x); \ 
    3429 ^ 
    3430 <scratch space>:127:1: note: expanded from here 
    3431 simd_and 
    3432 ^ 
    3433 program_source:396:1: error: implicit declaration of function 'quad_and' is
    3434 invalid in OpenCL 
    3435 program_source:360:60: note: expanded from macro
    3436 'DECLARE_I_REDUCTION_CLUSTERED' 
    3437 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3438 ^ 
    3439 program_source:246:59: note: expanded from macro '\ 
    3440 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3441 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3442 ^ 
    3443 program_source:196:12: note: expanded from macro '\ 
    3444 OVERLOAD_CLUSTERED_ARGS_1' 
    3445 return quad_##METAL_SUFFIX(x); \ 
    3446 ^ 
    3447 <scratch space>:128:1: note: expanded from here 
    3448 quad_and 
    3449 ^ 
    3450 program_source:396:1: error: implicit declaration of function 'simd_and' is
    3451 invalid in OpenCL 
    3452 program_source:360:60: note: expanded from macro
    3453 'DECLARE_I_REDUCTION_CLUSTERED' 
    3454 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3455 ^ 
    3456 program_source:246:59: note: expanded from macro '\ 
    3457 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3458 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3459 ^ 
    3460 program_source:198:12: note: expanded from macro '\ 
    3461 OVERLOAD_CLUSTERED_ARGS_1' 
    3462 return simd_##METAL_SUFFIX(x); \ 
    3463 ^ 
    3464 <scratch space>:129:1: note: expanded from here 
    3465 simd_and 
    3466 ^ 
    3467 program_source:397:1: error: implicit declaration of function 'quad_or' is
    3468 invalid in OpenCL 
    3469 DECLARE_I_REDUCTION_CLUSTERED(or, reduce_or) 
    3470 ^ 
    3471 program_source:360:60: note: expanded from macro
    3472 'DECLARE_I_REDUCTION_CLUSTERED' 
    3473 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3474 ^ 
    3475 program_source:245:71: note: expanded from macro '\ 
    3476 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3477 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3478 ^ 
    3479 program_source:196:12: note: expanded from macro '\ 
    3480 OVERLOAD_CLUSTERED_ARGS_1' 
    3481 return quad_##METAL_SUFFIX(x); \ 
    3482 ^ 
    3483 <scratch space>:131:1: note: expanded from here 
    3484 quad_or 
    3485 ^ 
    3486 program_source:397:1: error: implicit declaration of function 'simd_or' is
    3487 invalid in OpenCL 
    3488 program_source:360:60: note: expanded from macro
    3489 'DECLARE_I_REDUCTION_CLUSTERED' 
    3490 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3491 ^ 
    3492 program_source:245:71: note: expanded from macro '\ 
    3493 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3494 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3495 ^ 
    3496 program_source:198:12: note: expanded from macro '\ 
    3497 OVERLOAD_CLUSTERED_ARGS_1' 
    3498 return simd_##METAL_SUFFIX(x); \ 
    3499 ^ 
    3500 <scratch space>:132:1: note: expanded from here 
    3501 simd_or 
    3502 ^ 
    3503 program_source:397:1: error: implicit declaration of function 'quad_or' is
    3504 invalid in OpenCL 
    3505 program_source:360:60: note: expanded from macro
    3506 'DECLARE_I_REDUCTION_CLUSTERED' 
    3507 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3508 ^ 
    3509 program_source:246:59: note: expanded from macro '\ 
    3510 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3511 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3512 ^ 
    3513 program_source:196:12: note: expanded from macro '\ 
    3514 OVERLOAD_CLUSTERED_ARGS_1' 
    3515 return quad_##METAL_SUFFIX(x); \ 
    3516 ^ 
    3517 <scratch space>:133:1: note: expanded from here 
    3518 quad_or 
    3519 ^ 
    3520 program_source:397:1: error: implicit declaration of function 'simd_or' is
    3521 invalid in OpenCL 
    3522 program_source:360:60: note: expanded from macro
    3523 'DECLARE_I_REDUCTION_CLUSTERED' 
    3524 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3525 ^ 
    3526 program_source:246:59: note: expanded from macro '\ 
    3527 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3528 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3529 ^ 
    3530 program_source:198:12: note: expanded from macro '\ 
    3531 OVERLOAD_CLUSTERED_ARGS_1' 
    3532 return simd_##METAL_SUFFIX(x); \ 
    3533 ^ 
    3534 <scratch space>:134:1: note: expanded from here 
    3535 simd_or 
    3536 ^ 
    3537 program_source:398:1: error: implicit declaration of function 'quad_xor' is
    3538 invalid in OpenCL 
    3539 DECLARE_I_REDUCTION_CLUSTERED(xor, reduce_xor) 
    3540 ^ 
    3541 program_source:360:60: note: expanded from macro
    3542 'DECLARE_I_REDUCTION_CLUSTERED' 
    3543 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3544 ^ 
    3545 program_source:245:71: note: expanded from macro '\ 
    3546 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3547 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3548 ^ 
    3549 program_source:196:12: note: expanded from macro '\ 
    3550 OVERLOAD_CLUSTERED_ARGS_1' 
    3551 return quad_##METAL_SUFFIX(x); \ 
    3552 ^ 
    3553 <scratch space>:136:1: note: expanded from here 
    3554 quad_xor 
    3555 ^ 
    3556 program_source:398:1: error: implicit declaration of function 'simd_xor' is
    3557 invalid in OpenCL 
    3558 program_source:360:60: note: expanded from macro
    3559 'DECLARE_I_REDUCTION_CLUSTERED' 
    3560 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3561 ^ 
    3562 program_source:245:71: note: expanded from macro '\ 
    3563 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3564 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3565 ^ 
    3566 program_source:198:12: note: expanded from macro '\ 
    3567 OVERLOAD_CLUSTERED_ARGS_1' 
    3568 return simd_##METAL_SUFFIX(x); \ 
    3569 ^ 
    3570 <scratch space>:137:1: note: expanded from here 
    3571 simd_xor 
    3572 ^ 
    3573 program_source:398:1: error: implicit declaration of function 'quad_xor' is
    3574 invalid in OpenCL 
    3575 program_source:360:60: note: expanded from macro
    3576 'DECLARE_I_REDUCTION_CLUSTERED' 
    3577 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3578 ^ 
    3579 program_source:246:59: note: expanded from macro '\ 
    3580 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3581 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3582 ^ 
    3583 program_source:196:12: note: expanded from macro '\ 
    3584 OVERLOAD_CLUSTERED_ARGS_1' 
    3585 return quad_##METAL_SUFFIX(x); \ 
    3586 ^ 
    3587 <scratch space>:138:1: note: expanded from here 
    3588 quad_xor 
    3589 ^ 
    3590 program_source:398:1: error: implicit declaration of function 'simd_xor' is
    3591 invalid in OpenCL 
    3592 program_source:360:60: note: expanded from macro
    3593 'DECLARE_I_REDUCTION_CLUSTERED' 
    3594 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3595 ^ 
    3596 program_source:246:59: note: expanded from macro '\ 
    3597 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3598 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3599 ^ 
    3600 program_source:198:12: note: expanded from macro '\ 
    3601 OVERLOAD_CLUSTERED_ARGS_1' 
    3602 return simd_##METAL_SUFFIX(x); \ 
    3603 ^ 
    3604 <scratch space>:139:1: note: expanded from here 
    3605 simd_xor 
    3606 ^ 
    3607 program_source:399:1: error: implicit declaration of function 'simd_and' is
    3608 invalid in OpenCL 
    3609 DECLARE_B_REDUCTION_CLUSTERED(and) 
    3610 ^ 
    3611 program_source:369:8: note: expanded from macro
    3612 'DECLARE_B_REDUCTION_CLUSTERED' 
    3613 return simd_##OP(select(0, 1, predicate != 0)); \ 
    3614 ^ 
    3615 <scratch space>:141:1: note: expanded from here 
    3616 simd_and 
    3617 ^ 
    3618 program_source:399:1: error: implicit declaration of function 'quad_and' is
    3619 invalid in OpenCL 
    3620 program_source:378:12: note: expanded from macro
    3621 'DECLARE_B_REDUCTION_CLUSTERED' 
    3622 return quad_##OP(x); \ 
    3623 ^ 
    3624 <scratch space>:143:1: note: expanded from here 
    3625 quad_and 
    3626 ^ 
    3627 program_source:399:1: error: implicit declaration of function 'simd_and' is
    3628 invalid in OpenCL 
    3629 program_source:380:12: note: expanded from macro
    3630 'DECLARE_B_REDUCTION_CLUSTERED' 
    3631 return simd_##OP(x); \ 
    3632 ^ 
    3633 <scratch space>:144:1: note: expanded from here 
    3634 simd_and 
    3635 ^ 
    3636 program_source:400:1: error: implicit declaration of function 'simd_or' is
    3637 invalid in OpenCL 
    3638 DECLARE_B_REDUCTION_CLUSTERED(or) 
    3639 ^ 
    3640 program_source:369:8: note: expanded from macro
    3641 'DECLARE_B_REDUCTION_CLUSTERED' 
    3642 return simd_##OP(select(0, 1, predicate != 0)); \ 
    3643 ^ 
    3644 <scratch space>:146:1: note: expanded from here 
    3645 simd_or 
    3646 ^ 
    3647 program_source:400:1: error: implicit declaration of function 'quad_or' is
    3648 invalid in OpenCL 
    3649 program_source:378:12: note: expanded from macro
    3650 'DECLARE_B_REDUCTION_CLUSTERED' 
    3651 return quad_##OP(x); \ 
    3652 ^ 
    3653 <scratch space>:148:1: note: expanded from here 
    3654 quad_or 
    3655 ^ 
    3656 program_source:400:1: error: implicit declaration of function 'simd_or' is
    3657 invalid in OpenCL 
    3658 program_source:380:12: note: expanded from macro
    3659 'DECLARE_B_REDUCTION_CLUSTERED' 
    3660 return simd_##OP(x); \ 
    3661 ^ 
    3662 <scratch space>:149:1: note: expanded from here 
    3663 simd_or 
    3664 ^ 
    3665 program_source:401:1: error: implicit declaration of function 'simd_xor' is
    3666 invalid in OpenCL 
    3667 DECLARE_B_REDUCTION_CLUSTERED(xor) 
    3668 ^ 
    3669 program_source:369:8: note: expanded from macro
    3670 'DECLARE_B_REDUCTION_CLUSTERED' 
    3671 return simd_##OP(select(0, 1, predicate != 0)); \ 
    3672 ^ 
    3673 <scratch space>:151:1: note: expanded from here 
    3674 simd_xor 
    3675 ^ 
    3676 program_source:401:1: error: implicit declaration of function 'quad_xor' is
    3677 invalid in OpenCL 
    3678 program_source:378:12: note: expanded from macro
    3679 'DECLARE_B_REDUCTION_CLUSTERED' 
    3680 return quad_##OP(x); \ 
    3681 ^ 
    3682 <scratch space>:153:1: note: expanded from here 
    3683 quad_xor 
    3684 ^ 
    3685 program_source:401:1: error: implicit declaration of function 'simd_xor' is
    3686 invalid in OpenCL 
    3687 program_source:380:12: note: expanded from macro
    3688 'DECLARE_B_REDUCTION_CLUSTERED' 
    3689 return simd_##OP(x); \ 
    3690 ^ 
    3691 <scratch space>:154:1: note: expanded from here 
    3692 simd_xor 
    3693 ^ 
    3694 program_source:403:1: error: implicit declaration of function
    3695 'quad_shuffle_rotate_down' is invalid in OpenCL 
    3696 DECLARE_SHUFFLE_CLUSTERED(shuffle_rotate_down, rotate) 
    3697 ^ 
    3698 program_source:388:56: note: expanded from macro 'DECLARE_SHUFFLE_CLUSTERED' 
    3699 #define DECLARE_SHUFFLE_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3700 ^ 
    3701 program_source:252:69: note: expanded from macro '\ 
    3702 BRIDGE_FUNCTION_CLUSTERED_ARGS_2' 
    3703 #define BRIDGE_FUNCTION_CLUSTERED_ARGS_2(METAL_SUFFIX, OPENCL_EXPR) \ 
    3704 ^ 
    3705 program_source:208:12: note: expanded from macro '\ 
    3706 OVERLOAD_CLUSTERED_ARGS_2' 
    3707 return quad_##METAL_SUFFIX(x, ushort(delta)); \ 
    3708 ^ 
    3709 <scratch space>:156:1: note: expanded from here 
    3710 quad_shuffle_rotate_down 
    3711 ^ 
    3712 program_source:403:1: error: implicit declaration of function
    3713 'simd_shuffle_rotate_down' is invalid in OpenCL 
    3714 program_source:388:56: note: expanded from macro 'DECLARE_SHUFFLE_CLUSTERED' 
    3715 #define DECLARE_SHUFFLE_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3716 ^ 
    3717 program_source:252:69: note: expanded from macro '\ 
    3718 BRIDGE_FUNCTION_CLUSTERED_ARGS_2' 
    3719 #define BRIDGE_FUNCTION_CLUSTERED_ARGS_2(METAL_SUFFIX, OPENCL_EXPR) \ 
    3720 ^ 
    3721 program_source:210:12: note: expanded from macro '\ 
    3722 OVERLOAD_CLUSTERED_ARGS_2' 
    3723 return simd_##METAL_SUFFIX(x, ushort(delta)); \ 
    3724 ^ 
    3725 <scratch space>:157:1: note: expanded from here 
    3726 simd_shuffle_rotate_down 
    3727 ^ 
    3728 program_source:403:1: error: implicit declaration of function
    3729 'quad_shuffle_rotate_down' is invalid in OpenCL 
    3730 program_source:388:56: note: expanded from macro 'DECLARE_SHUFFLE_CLUSTERED' 
    3731 #define DECLARE_SHUFFLE_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3732 ^ 
    3733 program_source:253:59: note: expanded from macro '\ 
    3734 BRIDGE_FUNCTION_CLUSTERED_ARGS_2' 
    3735 OVERLOAD_CLUSTERED_ARGS_2(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3736 ^ 
    3737 program_source:208:12: note: expanded from macro '\ 
    3738 OVERLOAD_CLUSTERED_ARGS_2' 
    3739 return quad_##METAL_SUFFIX(x, ushort(delta)); \ 
    3740 ^ 
    3741 <scratch space>:158:1: note: expanded from here 
    3742 quad_shuffle_rotate_down 
    3743 ^ 
    3744 program_source:403:1: error: implicit declaration of function
    3745 'simd_shuffle_rotate_down' is invalid in OpenCL 
    3746 program_source:388:56: note: expanded from macro 'DECLARE_SHUFFLE_CLUSTERED' 
    3747 #define DECLARE_SHUFFLE_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3748 ^ 
    3749 program_source:253:59: note: expanded from macro '\ 
    3750 BRIDGE_FUNCTION_CLUSTERED_ARGS_2' 
    3751 OVERLOAD_CLUSTERED_ARGS_2(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3752 ^ 
    3753 program_source:210:12: note: expanded from macro '\ 
    3754 OVERLOAD_CLUSTERED_ARGS_2' 
    3755 return simd_##METAL_SUFFIX(x, ushort(delta)); \ 
    3756 ^ 
    3757 <scratch space>:159:1: note: expanded from here 
    3758 simd_shuffle_rotate_down 
    3759 ^ 
    3760 program_source:403:1: error: implicit declaration of function
    3761 'quad_shuffle_rotate_down' is invalid in OpenCL 
    3762 program_source:388:56: note: expanded from macro 'DECLARE_SHUFFLE_CLUSTERED' 
    3763 #define DECLARE_SHUFFLE_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3764 ^ 
    3765 program_source:254:60: note: expanded from macro '\ 
    3766 BRIDGE_FUNCTION_CLUSTERED_ARGS_2' 
    3767 OVERLOAD_CLUSTERED_ARGS_2(METAL_SUFFIX, OPENCL_EXPR, uint) \ 
    3768 ^ 
    3769 program_source:208:12: note: expanded from macro '\ 
    3770 OVERLOAD_CLUSTERED_ARGS_2' 
    3771 return quad_##METAL_SUFFIX(x, ushort(delta)); \ 
    3772 ^ 
    3773 <scratch space>:160:1: note: expanded from here 
    3774 quad_shuffle_rotate_down 
    3775 ^ 
    3776 program_source:403:1: error: implicit declaration of function
    3777 'simd_shuffle_rotate_down' is invalid in OpenCL 
    3778 program_source:388:56: note: expanded from macro 'DECLARE_SHUFFLE_CLUSTERED' 
    3779 #define DECLARE_SHUFFLE_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3780 ^ 
    3781 program_source:254:60: note: expanded from macro '\ 
    3782 BRIDGE_FUNCTION_CLUSTERED_ARGS_2' 
    3783 OVERLOAD_CLUSTERED_ARGS_2(METAL_SUFFIX, OPENCL_EXPR, uint) \ 
    3784 ^ 
    3785 program_source:210:12: note: expanded from macro '\ 
    3786 OVERLOAD_CLUSTERED_ARGS_2' 
    3787 return simd_##METAL_SUFFIX(x, ushort(delta)); \ 
    3788 ^ 
    3789 <scratch space>:161:1: note: expanded from here 
    3790 simd_shuffle_rotate_down 
    3791 ^ 
    3792 program_source:409:1: error: illegal string literal in 'asm' 
    3793 EXPOSE_BALLOT(simd_is_first, , bool, ) 
    3794 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3795 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    3796 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    3797 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3798 program_source:410:1: error: illegal string literal in 'asm' 
    3799 EXPOSE_BALLOT(simd_all, bool expr, bool, ) 
    3800 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3801 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    3802 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    3803 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3804 program_source:411:1: error: illegal string literal in 'asm' 
    3805 EXPOSE_BALLOT(simd_any, bool expr, bool, ) 
    3806 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3807 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    3808 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    3809 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3810 program_source:412:1: error: illegal string literal in 'asm' 
    3811 EXPOSE_BALLOT(simd_ballot, bool expr, ulong, .i64) 
    3812 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3813 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    3814 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    3815 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3816 program_source:413:1: error: illegal string literal in 'asm' 
    3817 EXPOSE_BALLOT(simd_active_threads_mask, , ulong, .i64) 
    3818 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3819 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    3820 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    3821 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3822 program_source:415:1: error: illegal string literal in 'asm' 
    3823 EXPOSE_BALLOT(quad_is_first, , bool, ) 
    3824 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3825 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    3826 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    3827 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3828 program_source:416:1: error: illegal string literal in 'asm' 
    3829 EXPOSE_BALLOT(quad_all, bool expr, bool, ) 
    3830 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3831 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    3832 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    3833 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3834 program_source:417:1: error: illegal string literal in 'asm' 
    3835 EXPOSE_BALLOT(quad_any, bool expr, bool, ) 
    3836 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3837 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    3838 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    3839 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3840 program_source:418:1: error: illegal string literal in 'asm' 
    3841 EXPOSE_BALLOT(quad_ballot, bool expr, ushort, ) 
    3842 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3843 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    3844 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    3845 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3846 program_source:419:1: error: illegal string literal in 'asm' 
    3847 EXPOSE_BALLOT(quad_active_threads_mask, , ushort, ) 
    3848 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3849 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    3850 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    3851 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3852 program_source:422:23: error: implicit declaration of function 'simd_is_first'
    3853 is invalid in OpenCL 
    3854 return select(0, 1, simd_is_first()); 
    3855 ^ 
    3856 program_source:426:23: error: implicit declaration of function 'simd_all' is
    3857 invalid in OpenCL 
    3858 return select(0, 1, simd_all(predicate != 0)); 
    3859 ^ 
    3860 program_source:430:23: error: implicit declaration of function 'simd_any' is
    3861 invalid in OpenCL 
    3862 return select(0, 1, simd_any(predicate != 0)); 
    3863 ^ 
    3864 program_source:434:23: error: implicit declaration of function 'simd_all' is
    3865 invalid in OpenCL 
    3866 return select(0, 1, simd_all(predicate != 0)); 
    3867 ^ 
    3868 program_source:438:23: error: implicit declaration of function 'simd_any' is
    3869 invalid in OpenCL 
    3870 return select(0, 1, simd_any(predicate != 0)); 
    3871 ^ 
    3872 program_source:443:14: error: implicit declaration of function 'simd_ballot'
    3873 is invalid in OpenCL 
    3874 output.x = simd_ballot(predicate != 0); 
    3875 ^ 
    3876 program_source:480:9: error: illegal string literal in 'asm' 
    3877 __asm("air.simdgroup.barrier"); 
    3878 ^~~~~~~~~~~~~~~~~~~~~~~ 
    3879 program_source:487:3: error: implicit declaration of function
    3880 '__metal_simdgroup_barrier' is invalid in OpenCL 
    3881 __metal_simdgroup_barrier(flags, 1); 
    3882 ^ 
    3883 program_source:487:3: note: did you mean 'simdgroup_barrier'? 
    3884 program_source:483:6: note: 'simdgroup_barrier' declared here 
    3885 void simdgroup_barrier(int flags) { 
    3886 ^ 
    3887 program_source:681:9: error: implicit declaration of function 'simd_is_first'
    3888 is invalid in OpenCL 
    3889 if (simd_is_first()) { 
    3890 ^ 
    3891 program_source:673:45: warning: comparison of integers of different signs:
    3892 '__private unsigned int' and '__private int' [-Wsign-compare] 
    3893 for (unsigned int index = thread; index < bufferSize; index +=
    3894 get_local_size(0)) 
    3895 ~~~~~ ^ ~~~~~~~~~~ 
    3896 program_source:685:45: warning: comparison of integers of different signs:
    3897 'unsigned int' and '__private int' [-Wsign-compare] 
    3898 if (thread%(32*2) == 0 && thread+32 < workGroupSize) 
    3899 ~~~~~~~~~ ^ ~~~~~~~~~~~~~ 
    3900 program_source:689:45: warning: comparison of integers of different signs:
    3901 'unsigned int' and '__private int' [-Wsign-compare] 
    3902 if (thread%(64*2) == 0 && thread+64 < workGroupSize) 
    3903 ~~~~~~~~~ ^ ~~~~~~~~~~~~~ 
    3904 program_source:693:47: warning: comparison of integers of different signs:
    3905 'unsigned int' and '__private int' [-Wsign-compare] 
    3906 if (thread%(128*2) == 0 && thread+128 < workGroupSize) 
    3907 ~~~~~~~~~~ ^ ~~~~~~~~~~~~~ 
    3908  
    3909  
    3910 openmm.OpenMMException: Error compiling kernel: program_source:279:1: error:
    3911 illegal string literal in 'asm' 
    3912 DECLARE_REDUCTION_BASE(sum) 
    3913 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3914 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    3915 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    3916 ^ 
    3917 program_source:263:44: note: expanded from macro '\ 
    3918 DECLARE_I_REDUCTION_BASE' 
    3919 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    3920 ^ 
    3921 program_source:217:40: note: expanded from macro '\ 
    3922 EXPOSE_FUNCTION_I_ARGS_1' 
    3923 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    3924 ^ 
    3925 program_source:164:9: note: expanded from macro '\ 
    3926 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    3927 __asm("air." #EXPR "." #AIR_TYPE); \ 
    3928 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3929 program_source:279:1: error: illegal string literal in 'asm' 
    3930 DECLARE_REDUCTION_BASE(sum) 
    3931 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3932 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    3933 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    3934 ^ 
    3935 program_source:263:44: note: expanded from macro '\ 
    3936 DECLARE_I_REDUCTION_BASE' 
    3937 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    3938 ^ 
    3939 program_source:218:51: note: expanded from macro '\ 
    3940 EXPOSE_FUNCTION_I_ARGS_1' 
    3941 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    3942 ^ 
    3943 program_source:164:9: note: expanded from macro '\ 
    3944 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    3945 __asm("air." #EXPR "." #AIR_TYPE); \ 
    3946 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3947 program_source:279:1: error: illegal string literal in 'asm' 
    3948 DECLARE_REDUCTION_BASE(sum) 
    3949 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3950 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    3951 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    3952 ^ 
    3953 program_source:264:43: note: expanded from macro '\ 
    3954 DECLARE_I_REDUCTION_BASE' 
    3955 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    3956 ^ 
    3957 program_source:217:40: note: expanded from macro '\ 
    3958 EXPOSE_FUNCTION_I_ARGS_1' 
    3959 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    3960 ^ 
    3961 program_source:164:9: note: expanded from macro '\ 
    3962 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    3963 __asm("air." #EXPR "." #AIR_TYPE); \ 
    3964 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3965 program_source:279:1: error: illegal string literal in 'asm' 
    3966 DECLARE_REDUCTION_BASE(sum) 
    3967 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3968 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    3969 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    3970 ^ 
    3971 program_source:264:43: note: expanded from macro '\ 
    3972 DECLARE_I_REDUCTION_BASE' 
    3973 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    3974 ^ 
    3975 program_source:218:51: note: expanded from macro '\ 
    3976 EXPOSE_FUNCTION_I_ARGS_1' 
    3977 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    3978 ^ 
    3979 program_source:164:9: note: expanded from macro '\ 
    3980 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    3981 __asm("air." #EXPR "." #AIR_TYPE); \ 
    3982 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3983 program_source:279:1: error: illegal string literal in 'asm' 
    3984 DECLARE_REDUCTION_BASE(sum) 
    3985 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3986 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    3987 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    3988 ^ 
    3989 program_source:267:44: note: expanded from macro '\ 
    3990 DECLARE_F_REDUCTION_BASE' 
    3991 #define DECLARE_F_REDUCTION_BASE(METAL_OP) \ 
    3992 ^ 
    3993 program_source:221:40: note: expanded from macro '\ 
    3994 EXPOSE_FUNCTION_F_ARGS_1' 
    3995 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    3996 ^ 
    3997 program_source:164:9: note: expanded from macro '\ 
    3998 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    3999 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4000 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4001 program_source:279:1: error: illegal string literal in 'asm' 
    4002 DECLARE_REDUCTION_BASE(sum) 
    4003 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4004 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4005 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4006 ^ 
    4007 program_source:268:43: note: expanded from macro '\ 
    4008 DECLARE_F_REDUCTION_BASE' 
    4009 EXPOSE_FUNCTION_F_ARGS_1(quad_##METAL_OP) \ 
    4010 ^ 
    4011 program_source:221:40: note: expanded from macro '\ 
    4012 EXPOSE_FUNCTION_F_ARGS_1' 
    4013 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4014 ^ 
    4015 program_source:164:9: note: expanded from macro '\ 
    4016 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4017 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4018 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4019 program_source:280:1: error: illegal string literal in 'asm' 
    4020 DECLARE_REDUCTION_BASE(prefix_inclusive_sum) 
    4021 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4022 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4023 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4024 ^ 
    4025 program_source:263:44: note: expanded from macro '\ 
    4026 DECLARE_I_REDUCTION_BASE' 
    4027 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4028 ^ 
    4029 program_source:217:40: note: expanded from macro '\ 
    4030 EXPOSE_FUNCTION_I_ARGS_1' 
    4031 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4032 ^ 
    4033 program_source:164:9: note: expanded from macro '\ 
    4034 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4035 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4036 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4037 program_source:280:1: error: illegal string literal in 'asm' 
    4038 DECLARE_REDUCTION_BASE(prefix_inclusive_sum) 
    4039 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4040 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4041 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4042 ^ 
    4043 program_source:263:44: note: expanded from macro '\ 
    4044 DECLARE_I_REDUCTION_BASE' 
    4045 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4046 ^ 
    4047 program_source:218:51: note: expanded from macro '\ 
    4048 EXPOSE_FUNCTION_I_ARGS_1' 
    4049 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4050 ^ 
    4051 program_source:164:9: note: expanded from macro '\ 
    4052 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4053 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4054 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4055 program_source:280:1: error: illegal string literal in 'asm' 
    4056 DECLARE_REDUCTION_BASE(prefix_inclusive_sum) 
    4057 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4058 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4059 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4060 ^ 
    4061 program_source:264:43: note: expanded from macro '\ 
    4062 DECLARE_I_REDUCTION_BASE' 
    4063 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4064 ^ 
    4065 program_source:217:40: note: expanded from macro '\ 
    4066 EXPOSE_FUNCTION_I_ARGS_1' 
    4067 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4068 ^ 
    4069 program_source:164:9: note: expanded from macro '\ 
    4070 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4071 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4072 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4073 program_source:280:1: error: illegal string literal in 'asm' 
    4074 DECLARE_REDUCTION_BASE(prefix_inclusive_sum) 
    4075 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4076 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4077 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4078 ^ 
    4079 program_source:264:43: note: expanded from macro '\ 
    4080 DECLARE_I_REDUCTION_BASE' 
    4081 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4082 ^ 
    4083 program_source:218:51: note: expanded from macro '\ 
    4084 EXPOSE_FUNCTION_I_ARGS_1' 
    4085 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4086 ^ 
    4087 program_source:164:9: note: expanded from macro '\ 
    4088 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4089 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4090 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4091 program_source:280:1: error: illegal string literal in 'asm' 
    4092 DECLARE_REDUCTION_BASE(prefix_inclusive_sum) 
    4093 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4094 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4095 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4096 ^ 
    4097 program_source:267:44: note: expanded from macro '\ 
    4098 DECLARE_F_REDUCTION_BASE' 
    4099 #define DECLARE_F_REDUCTION_BASE(METAL_OP) \ 
    4100 ^ 
    4101 program_source:221:40: note: expanded from macro '\ 
    4102 EXPOSE_FUNCTION_F_ARGS_1' 
    4103 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4104 ^ 
    4105 program_source:164:9: note: expanded from macro '\ 
    4106 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4107 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4108 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4109 program_source:280:1: error: illegal string literal in 'asm' 
    4110 DECLARE_REDUCTION_BASE(prefix_inclusive_sum) 
    4111 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4112 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4113 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4114 ^ 
    4115 program_source:268:43: note: expanded from macro '\ 
    4116 DECLARE_F_REDUCTION_BASE' 
    4117 EXPOSE_FUNCTION_F_ARGS_1(quad_##METAL_OP) \ 
    4118 ^ 
    4119 program_source:221:40: note: expanded from macro '\ 
    4120 EXPOSE_FUNCTION_F_ARGS_1' 
    4121 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4122 ^ 
    4123 program_source:164:9: note: expanded from macro '\ 
    4124 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4125 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4126 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4127 program_source:281:1: error: illegal string literal in 'asm' 
    4128 DECLARE_REDUCTION_BASE(prefix_exclusive_sum) 
    4129 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4130 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4131 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4132 ^ 
    4133 program_source:263:44: note: expanded from macro '\ 
    4134 DECLARE_I_REDUCTION_BASE' 
    4135 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4136 ^ 
    4137 program_source:217:40: note: expanded from macro '\ 
    4138 EXPOSE_FUNCTION_I_ARGS_1' 
    4139 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4140 ^ 
    4141 program_source:164:9: note: expanded from macro '\ 
    4142 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4143 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4144 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4145 program_source:281:1: error: illegal string literal in 'asm' 
    4146 DECLARE_REDUCTION_BASE(prefix_exclusive_sum) 
    4147 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4148 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4149 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4150 ^ 
    4151 program_source:263:44: note: expanded from macro '\ 
    4152 DECLARE_I_REDUCTION_BASE' 
    4153 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4154 ^ 
    4155 program_source:218:51: note: expanded from macro '\ 
    4156 EXPOSE_FUNCTION_I_ARGS_1' 
    4157 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4158 ^ 
    4159 program_source:164:9: note: expanded from macro '\ 
    4160 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4161 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4162 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4163 program_source:281:1: error: illegal string literal in 'asm' 
    4164 DECLARE_REDUCTION_BASE(prefix_exclusive_sum) 
    4165 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4166 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4167 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4168 ^ 
    4169 program_source:264:43: note: expanded from macro '\ 
    4170 DECLARE_I_REDUCTION_BASE' 
    4171 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4172 ^ 
    4173 program_source:217:40: note: expanded from macro '\ 
    4174 EXPOSE_FUNCTION_I_ARGS_1' 
    4175 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4176 ^ 
    4177 program_source:164:9: note: expanded from macro '\ 
    4178 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4179 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4180 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4181 program_source:281:1: error: illegal string literal in 'asm' 
    4182 DECLARE_REDUCTION_BASE(prefix_exclusive_sum) 
    4183 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4184 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4185 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4186 ^ 
    4187 program_source:264:43: note: expanded from macro '\ 
    4188 DECLARE_I_REDUCTION_BASE' 
    4189 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4190 ^ 
    4191 program_source:218:51: note: expanded from macro '\ 
    4192 EXPOSE_FUNCTION_I_ARGS_1' 
    4193 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4194 ^ 
    4195 program_source:164:9: note: expanded from macro '\ 
    4196 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4197 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4198 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4199 program_source:281:1: error: illegal string literal in 'asm' 
    4200 DECLARE_REDUCTION_BASE(prefix_exclusive_sum) 
    4201 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4202 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4203 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4204 ^ 
    4205 program_source:267:44: note: expanded from macro '\ 
    4206 DECLARE_F_REDUCTION_BASE' 
    4207 #define DECLARE_F_REDUCTION_BASE(METAL_OP) \ 
    4208 ^ 
    4209 program_source:221:40: note: expanded from macro '\ 
    4210 EXPOSE_FUNCTION_F_ARGS_1' 
    4211 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4212 ^ 
    4213 program_source:164:9: note: expanded from macro '\ 
    4214 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4215 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4216 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4217 program_source:281:1: error: illegal string literal in 'asm' 
    4218 DECLARE_REDUCTION_BASE(prefix_exclusive_sum) 
    4219 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4220 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4221 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4222 ^ 
    4223 program_source:268:43: note: expanded from macro '\ 
    4224 DECLARE_F_REDUCTION_BASE' 
    4225 EXPOSE_FUNCTION_F_ARGS_1(quad_##METAL_OP) \ 
    4226 ^ 
    4227 program_source:221:40: note: expanded from macro '\ 
    4228 EXPOSE_FUNCTION_F_ARGS_1' 
    4229 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4230 ^ 
    4231 program_source:164:9: note: expanded from macro '\ 
    4232 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4233 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4234 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4235 program_source:282:1: error: illegal string literal in 'asm' 
    4236 DECLARE_REDUCTION_BASE(min) 
    4237 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4238 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4239 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4240 ^ 
    4241 program_source:263:44: note: expanded from macro '\ 
    4242 DECLARE_I_REDUCTION_BASE' 
    4243 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4244 ^ 
    4245 program_source:217:40: note: expanded from macro '\ 
    4246 EXPOSE_FUNCTION_I_ARGS_1' 
    4247 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4248 ^ 
    4249 program_source:164:9: note: expanded from macro '\ 
    4250 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4251 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4252 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4253 program_source:282:1: error: illegal string literal in 'asm' 
    4254 DECLARE_REDUCTION_BASE(min) 
    4255 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4256 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4257 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4258 ^ 
    4259 program_source:263:44: note: expanded from macro '\ 
    4260 DECLARE_I_REDUCTION_BASE' 
    4261 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4262 ^ 
    4263 program_source:218:51: note: expanded from macro '\ 
    4264 EXPOSE_FUNCTION_I_ARGS_1' 
    4265 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4266 ^ 
    4267 program_source:164:9: note: expanded from macro '\ 
    4268 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4269 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4270 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4271 program_source:282:1: error: illegal string literal in 'asm' 
    4272 DECLARE_REDUCTION_BASE(min) 
    4273 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4274 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4275 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4276 ^ 
    4277 program_source:264:43: note: expanded from macro '\ 
    4278 DECLARE_I_REDUCTION_BASE' 
    4279 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4280 ^ 
    4281 program_source:217:40: note: expanded from macro '\ 
    4282 EXPOSE_FUNCTION_I_ARGS_1' 
    4283 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4284 ^ 
    4285 program_source:164:9: note: expanded from macro '\ 
    4286 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4287 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4288 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4289 program_source:282:1: error: illegal string literal in 'asm' 
    4290 DECLARE_REDUCTION_BASE(min) 
    4291 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4292 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4293 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4294 ^ 
    4295 program_source:264:43: note: expanded from macro '\ 
    4296 DECLARE_I_REDUCTION_BASE' 
    4297 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4298 ^ 
    4299 program_source:218:51: note: expanded from macro '\ 
    4300 EXPOSE_FUNCTION_I_ARGS_1' 
    4301 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4302 ^ 
    4303 program_source:164:9: note: expanded from macro '\ 
    4304 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4305 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4306 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4307 program_source:282:1: error: illegal string literal in 'asm' 
    4308 DECLARE_REDUCTION_BASE(min) 
    4309 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4310 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4311 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4312 ^ 
    4313 program_source:267:44: note: expanded from macro '\ 
    4314 DECLARE_F_REDUCTION_BASE' 
    4315 #define DECLARE_F_REDUCTION_BASE(METAL_OP) \ 
    4316 ^ 
    4317 program_source:221:40: note: expanded from macro '\ 
    4318 EXPOSE_FUNCTION_F_ARGS_1' 
    4319 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4320 ^ 
    4321 program_source:164:9: note: expanded from macro '\ 
    4322 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4323 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4324 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4325 program_source:282:1: error: illegal string literal in 'asm' 
    4326 DECLARE_REDUCTION_BASE(min) 
    4327 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4328 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4329 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4330 ^ 
    4331 program_source:268:43: note: expanded from macro '\ 
    4332 DECLARE_F_REDUCTION_BASE' 
    4333 EXPOSE_FUNCTION_F_ARGS_1(quad_##METAL_OP) \ 
    4334 ^ 
    4335 program_source:221:40: note: expanded from macro '\ 
    4336 EXPOSE_FUNCTION_F_ARGS_1' 
    4337 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4338 ^ 
    4339 program_source:164:9: note: expanded from macro '\ 
    4340 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4341 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4342 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4343 program_source:283:1: error: illegal string literal in 'asm' 
    4344 DECLARE_REDUCTION_BASE(max) 
    4345 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4346 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4347 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4348 ^ 
    4349 program_source:263:44: note: expanded from macro '\ 
    4350 DECLARE_I_REDUCTION_BASE' 
    4351 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4352 ^ 
    4353 program_source:217:40: note: expanded from macro '\ 
    4354 EXPOSE_FUNCTION_I_ARGS_1' 
    4355 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4356 ^ 
    4357 program_source:164:9: note: expanded from macro '\ 
    4358 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4359 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4360 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4361 program_source:283:1: error: illegal string literal in 'asm' 
    4362 DECLARE_REDUCTION_BASE(max) 
    4363 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4364 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4365 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4366 ^ 
    4367 program_source:263:44: note: expanded from macro '\ 
    4368 DECLARE_I_REDUCTION_BASE' 
    4369 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4370 ^ 
    4371 program_source:218:51: note: expanded from macro '\ 
    4372 EXPOSE_FUNCTION_I_ARGS_1' 
    4373 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4374 ^ 
    4375 program_source:164:9: note: expanded from macro '\ 
    4376 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4377 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4378 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4379 program_source:283:1: error: illegal string literal in 'asm' 
    4380 DECLARE_REDUCTION_BASE(max) 
    4381 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4382 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4383 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4384 ^ 
    4385 program_source:264:43: note: expanded from macro '\ 
    4386 DECLARE_I_REDUCTION_BASE' 
    4387 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4388 ^ 
    4389 program_source:217:40: note: expanded from macro '\ 
    4390 EXPOSE_FUNCTION_I_ARGS_1' 
    4391 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4392 ^ 
    4393 program_source:164:9: note: expanded from macro '\ 
    4394 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4395 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4396 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4397 program_source:283:1: error: illegal string literal in 'asm' 
    4398 DECLARE_REDUCTION_BASE(max) 
    4399 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4400 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4401 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4402 ^ 
    4403 program_source:264:43: note: expanded from macro '\ 
    4404 DECLARE_I_REDUCTION_BASE' 
    4405 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4406 ^ 
    4407 program_source:218:51: note: expanded from macro '\ 
    4408 EXPOSE_FUNCTION_I_ARGS_1' 
    4409 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4410 ^ 
    4411 program_source:164:9: note: expanded from macro '\ 
    4412 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4413 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4414 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4415 program_source:283:1: error: illegal string literal in 'asm' 
    4416 DECLARE_REDUCTION_BASE(max) 
    4417 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4418 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4419 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4420 ^ 
    4421 program_source:267:44: note: expanded from macro '\ 
    4422 DECLARE_F_REDUCTION_BASE' 
    4423 #define DECLARE_F_REDUCTION_BASE(METAL_OP) \ 
    4424 ^ 
    4425 program_source:221:40: note: expanded from macro '\ 
    4426 EXPOSE_FUNCTION_F_ARGS_1' 
    4427 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4428 ^ 
    4429 program_source:164:9: note: expanded from macro '\ 
    4430 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4431 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4432 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4433 program_source:283:1: error: illegal string literal in 'asm' 
    4434 DECLARE_REDUCTION_BASE(max) 
    4435 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4436 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4437 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4438 ^ 
    4439 program_source:268:43: note: expanded from macro '\ 
    4440 DECLARE_F_REDUCTION_BASE' 
    4441 EXPOSE_FUNCTION_F_ARGS_1(quad_##METAL_OP) \ 
    4442 ^ 
    4443 program_source:221:40: note: expanded from macro '\ 
    4444 EXPOSE_FUNCTION_F_ARGS_1' 
    4445 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4446 ^ 
    4447 program_source:164:9: note: expanded from macro '\ 
    4448 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4449 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4450 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4451 program_source:285:1: error: illegal string literal in 'asm' 
    4452 DECLARE_REDUCTION_BASE(product) 
    4453 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4454 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4455 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4456 ^ 
    4457 program_source:263:44: note: expanded from macro '\ 
    4458 DECLARE_I_REDUCTION_BASE' 
    4459 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4460 ^ 
    4461 program_source:217:40: note: expanded from macro '\ 
    4462 EXPOSE_FUNCTION_I_ARGS_1' 
    4463 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4464 ^ 
    4465 program_source:164:9: note: expanded from macro '\ 
    4466 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4467 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4468 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4469 program_source:285:1: error: illegal string literal in 'asm' 
    4470 DECLARE_REDUCTION_BASE(product) 
    4471 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4472 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4473 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4474 ^ 
    4475 program_source:263:44: note: expanded from macro '\ 
    4476 DECLARE_I_REDUCTION_BASE' 
    4477 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4478 ^ 
    4479 program_source:218:51: note: expanded from macro '\ 
    4480 EXPOSE_FUNCTION_I_ARGS_1' 
    4481 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4482 ^ 
    4483 program_source:164:9: note: expanded from macro '\ 
    4484 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4485 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4486 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4487 program_source:285:1: error: illegal string literal in 'asm' 
    4488 DECLARE_REDUCTION_BASE(product) 
    4489 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4490 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4491 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4492 ^ 
    4493 program_source:264:43: note: expanded from macro '\ 
    4494 DECLARE_I_REDUCTION_BASE' 
    4495 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4496 ^ 
    4497 program_source:217:40: note: expanded from macro '\ 
    4498 EXPOSE_FUNCTION_I_ARGS_1' 
    4499 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4500 ^ 
    4501 program_source:164:9: note: expanded from macro '\ 
    4502 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4503 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4504 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4505 program_source:285:1: error: illegal string literal in 'asm' 
    4506 DECLARE_REDUCTION_BASE(product) 
    4507 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4508 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4509 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4510 ^ 
    4511 program_source:264:43: note: expanded from macro '\ 
    4512 DECLARE_I_REDUCTION_BASE' 
    4513 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4514 ^ 
    4515 program_source:218:51: note: expanded from macro '\ 
    4516 EXPOSE_FUNCTION_I_ARGS_1' 
    4517 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4518 ^ 
    4519 program_source:164:9: note: expanded from macro '\ 
    4520 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4521 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4522 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4523 program_source:285:1: error: illegal string literal in 'asm' 
    4524 DECLARE_REDUCTION_BASE(product) 
    4525 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4526 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4527 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4528 ^ 
    4529 program_source:267:44: note: expanded from macro '\ 
    4530 DECLARE_F_REDUCTION_BASE' 
    4531 #define DECLARE_F_REDUCTION_BASE(METAL_OP) \ 
    4532 ^ 
    4533 program_source:221:40: note: expanded from macro '\ 
    4534 EXPOSE_FUNCTION_F_ARGS_1' 
    4535 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4536 ^ 
    4537 program_source:164:9: note: expanded from macro '\ 
    4538 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4539 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4540 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4541 program_source:285:1: error: illegal string literal in 'asm' 
    4542 DECLARE_REDUCTION_BASE(product) 
    4543 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4544 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4545 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4546 ^ 
    4547 program_source:268:43: note: expanded from macro '\ 
    4548 DECLARE_F_REDUCTION_BASE' 
    4549 EXPOSE_FUNCTION_F_ARGS_1(quad_##METAL_OP) \ 
    4550 ^ 
    4551 program_source:221:40: note: expanded from macro '\ 
    4552 EXPOSE_FUNCTION_F_ARGS_1' 
    4553 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4554 ^ 
    4555 program_source:164:9: note: expanded from macro '\ 
    4556 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4557 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4558 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4559 program_source:286:1: error: illegal string literal in 'asm' 
    4560 DECLARE_REDUCTION_BASE(prefix_inclusive_product) 
    4561 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4562 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4563 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4564 ^ 
    4565 program_source:263:44: note: expanded from macro '\ 
    4566 DECLARE_I_REDUCTION_BASE' 
    4567 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4568 ^ 
    4569 program_source:217:40: note: expanded from macro '\ 
    4570 EXPOSE_FUNCTION_I_ARGS_1' 
    4571 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4572 ^ 
    4573 program_source:164:9: note: expanded from macro '\ 
    4574 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4575 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4576 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4577 program_source:286:1: error: illegal string literal in 'asm' 
    4578 DECLARE_REDUCTION_BASE(prefix_inclusive_product) 
    4579 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4580 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4581 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4582 ^ 
    4583 program_source:263:44: note: expanded from macro '\ 
    4584 DECLARE_I_REDUCTION_BASE' 
    4585 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4586 ^ 
    4587 program_source:218:51: note: expanded from macro '\ 
    4588 EXPOSE_FUNCTION_I_ARGS_1' 
    4589 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4590 ^ 
    4591 program_source:164:9: note: expanded from macro '\ 
    4592 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4593 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4594 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4595 program_source:286:1: error: illegal string literal in 'asm' 
    4596 DECLARE_REDUCTION_BASE(prefix_inclusive_product) 
    4597 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4598 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4599 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4600 ^ 
    4601 program_source:264:43: note: expanded from macro '\ 
    4602 DECLARE_I_REDUCTION_BASE' 
    4603 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4604 ^ 
    4605 program_source:217:40: note: expanded from macro '\ 
    4606 EXPOSE_FUNCTION_I_ARGS_1' 
    4607 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4608 ^ 
    4609 program_source:164:9: note: expanded from macro '\ 
    4610 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4611 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4612 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4613 program_source:286:1: error: illegal string literal in 'asm' 
    4614 DECLARE_REDUCTION_BASE(prefix_inclusive_product) 
    4615 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4616 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4617 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4618 ^ 
    4619 program_source:264:43: note: expanded from macro '\ 
    4620 DECLARE_I_REDUCTION_BASE' 
    4621 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4622 ^ 
    4623 program_source:218:51: note: expanded from macro '\ 
    4624 EXPOSE_FUNCTION_I_ARGS_1' 
    4625 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4626 ^ 
    4627 program_source:164:9: note: expanded from macro '\ 
    4628 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4629 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4630 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4631 program_source:286:1: error: illegal string literal in 'asm' 
    4632 DECLARE_REDUCTION_BASE(prefix_inclusive_product) 
    4633 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4634 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4635 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4636 ^ 
    4637 program_source:267:44: note: expanded from macro '\ 
    4638 DECLARE_F_REDUCTION_BASE' 
    4639 #define DECLARE_F_REDUCTION_BASE(METAL_OP) \ 
    4640 ^ 
    4641 program_source:221:40: note: expanded from macro '\ 
    4642 EXPOSE_FUNCTION_F_ARGS_1' 
    4643 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4644 ^ 
    4645 program_source:164:9: note: expanded from macro '\ 
    4646 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4647 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4648 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4649 program_source:286:1: error: illegal string literal in 'asm' 
    4650 DECLARE_REDUCTION_BASE(prefix_inclusive_product) 
    4651 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4652 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4653 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4654 ^ 
    4655 program_source:268:43: note: expanded from macro '\ 
    4656 DECLARE_F_REDUCTION_BASE' 
    4657 EXPOSE_FUNCTION_F_ARGS_1(quad_##METAL_OP) \ 
    4658 ^ 
    4659 program_source:221:40: note: expanded from macro '\ 
    4660 EXPOSE_FUNCTION_F_ARGS_1' 
    4661 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4662 ^ 
    4663 program_source:164:9: note: expanded from macro '\ 
    4664 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4665 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4666 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4667 program_source:287:1: error: illegal string literal in 'asm' 
    4668 DECLARE_REDUCTION_BASE(prefix_exclusive_product) 
    4669 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4670 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4671 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4672 ^ 
    4673 program_source:263:44: note: expanded from macro '\ 
    4674 DECLARE_I_REDUCTION_BASE' 
    4675 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4676 ^ 
    4677 program_source:217:40: note: expanded from macro '\ 
    4678 EXPOSE_FUNCTION_I_ARGS_1' 
    4679 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4680 ^ 
    4681 program_source:164:9: note: expanded from macro '\ 
    4682 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4683 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4684 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4685 program_source:287:1: error: illegal string literal in 'asm' 
    4686 DECLARE_REDUCTION_BASE(prefix_exclusive_product) 
    4687 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4688 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4689 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4690 ^ 
    4691 program_source:263:44: note: expanded from macro '\ 
    4692 DECLARE_I_REDUCTION_BASE' 
    4693 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4694 ^ 
    4695 program_source:218:51: note: expanded from macro '\ 
    4696 EXPOSE_FUNCTION_I_ARGS_1' 
    4697 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4698 ^ 
    4699 program_source:164:9: note: expanded from macro '\ 
    4700 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4701 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4702 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4703 program_source:287:1: error: illegal string literal in 'asm' 
    4704 DECLARE_REDUCTION_BASE(prefix_exclusive_product) 
    4705 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4706 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4707 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4708 ^ 
    4709 program_source:264:43: note: expanded from macro '\ 
    4710 DECLARE_I_REDUCTION_BASE' 
    4711 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4712 ^ 
    4713 program_source:217:40: note: expanded from macro '\ 
    4714 EXPOSE_FUNCTION_I_ARGS_1' 
    4715 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4716 ^ 
    4717 program_source:164:9: note: expanded from macro '\ 
    4718 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4719 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4720 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4721 program_source:287:1: error: illegal string literal in 'asm' 
    4722 DECLARE_REDUCTION_BASE(prefix_exclusive_product) 
    4723 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4724 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4725 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4726 ^ 
    4727 program_source:264:43: note: expanded from macro '\ 
    4728 DECLARE_I_REDUCTION_BASE' 
    4729 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4730 ^ 
    4731 program_source:218:51: note: expanded from macro '\ 
    4732 EXPOSE_FUNCTION_I_ARGS_1' 
    4733 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4734 ^ 
    4735 program_source:164:9: note: expanded from macro '\ 
    4736 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4737 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4738 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4739 program_source:287:1: error: illegal string literal in 'asm' 
    4740 DECLARE_REDUCTION_BASE(prefix_exclusive_product) 
    4741 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4742 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4743 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4744 ^ 
    4745 program_source:267:44: note: expanded from macro '\ 
    4746 DECLARE_F_REDUCTION_BASE' 
    4747 #define DECLARE_F_REDUCTION_BASE(METAL_OP) \ 
    4748 ^ 
    4749 program_source:221:40: note: expanded from macro '\ 
    4750 EXPOSE_FUNCTION_F_ARGS_1' 
    4751 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4752 ^ 
    4753 program_source:164:9: note: expanded from macro '\ 
    4754 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4755 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4756 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4757 program_source:287:1: error: illegal string literal in 'asm' 
    4758 DECLARE_REDUCTION_BASE(prefix_exclusive_product) 
    4759 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4760 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4761 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4762 ^ 
    4763 program_source:268:43: note: expanded from macro '\ 
    4764 DECLARE_F_REDUCTION_BASE' 
    4765 EXPOSE_FUNCTION_F_ARGS_1(quad_##METAL_OP) \ 
    4766 ^ 
    4767 program_source:221:40: note: expanded from macro '\ 
    4768 EXPOSE_FUNCTION_F_ARGS_1' 
    4769 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4770 ^ 
    4771 program_source:164:9: note: expanded from macro '\ 
    4772 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4773 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4774 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4775 program_source:288:1: error: illegal string literal in 'asm' 
    4776 DECLARE_I_REDUCTION_BASE(and) 
    4777 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4778 program_source:263:44: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    4779 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4780 ^ 
    4781 program_source:217:40: note: expanded from macro '\ 
    4782 EXPOSE_FUNCTION_I_ARGS_1' 
    4783 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4784 ^ 
    4785 program_source:164:9: note: expanded from macro '\ 
    4786 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4787 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4788 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4789 program_source:288:1: error: illegal string literal in 'asm' 
    4790 DECLARE_I_REDUCTION_BASE(and) 
    4791 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4792 program_source:263:44: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    4793 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4794 ^ 
    4795 program_source:218:51: note: expanded from macro '\ 
    4796 EXPOSE_FUNCTION_I_ARGS_1' 
    4797 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    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:264:43: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    4807 EXPOSE_FUNCTION_I_ARGS_1(quad_##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:264:43: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    4821 EXPOSE_FUNCTION_I_ARGS_1(quad_##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:289:1: error: illegal string literal in 'asm' 
    4832 DECLARE_I_REDUCTION_BASE(or) 
    4833 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4834 program_source:263:44: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    4835 #define DECLARE_I_REDUCTION_BASE(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:289:1: error: illegal string literal in 'asm' 
    4846 DECLARE_I_REDUCTION_BASE(or) 
    4847 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4848 program_source:263:44: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    4849 #define DECLARE_I_REDUCTION_BASE(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:264:43: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    4863 EXPOSE_FUNCTION_I_ARGS_1(quad_##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:264:43: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    4877 EXPOSE_FUNCTION_I_ARGS_1(quad_##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:290:1: error: illegal string literal in 'asm' 
    4888 DECLARE_I_REDUCTION_BASE(xor) 
    4889 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4890 program_source:263:44: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    4891 #define DECLARE_I_REDUCTION_BASE(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:290:1: error: illegal string literal in 'asm' 
    4902 DECLARE_I_REDUCTION_BASE(xor) 
    4903 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4904 program_source:263:44: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    4905 #define DECLARE_I_REDUCTION_BASE(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:264:43: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    4919 EXPOSE_FUNCTION_I_ARGS_1(quad_##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:264:43: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    4933 EXPOSE_FUNCTION_I_ARGS_1(quad_##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:292:1: error: illegal string literal in 'asm' 
    4944 DECLARE_SHUFFLE_BASE(broadcast) 
    4945 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4946 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    4947 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    4948 ^ 
    4949 program_source:224:38: note: expanded from macro '\ 
    4950 EXPOSE_FUNCTION_ARGS_2' 
    4951 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    4952 ^ 
    4953 program_source:168:9: note: expanded from macro '\ 
    4954 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    4955 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4956 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4957 program_source:292:1: error: illegal string literal in 'asm' 
    4958 DECLARE_SHUFFLE_BASE(broadcast) 
    4959 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4960 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    4961 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    4962 ^ 
    4963 program_source:225:51: note: expanded from macro '\ 
    4964 EXPOSE_FUNCTION_ARGS_2' 
    4965 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    4966 ^ 
    4967 program_source:168:9: note: expanded from macro '\ 
    4968 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    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:226:52: note: expanded from macro '\ 
    4978 EXPOSE_FUNCTION_ARGS_2' 
    4979 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    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:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    4989 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    4990 ^ 
    4991 program_source:224:38: note: expanded from macro '\ 
    4992 EXPOSE_FUNCTION_ARGS_2' 
    4993 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    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:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5003 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5004 ^ 
    5005 program_source:225:51: note: expanded from macro '\ 
    5006 EXPOSE_FUNCTION_ARGS_2' 
    5007 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.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:226:52: note: expanded from macro '\ 
    5020 EXPOSE_FUNCTION_ARGS_2' 
    5021 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    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:293:1: error: illegal string literal in 'asm' 
    5028 DECLARE_REDUCTION_BASE(broadcast_first) 
    5029 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5030 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    5031 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    5032 ^ 
    5033 program_source:263:44: note: expanded from macro '\ 
    5034 DECLARE_I_REDUCTION_BASE' 
    5035 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    5036 ^ 
    5037 program_source:217:40: note: expanded from macro '\ 
    5038 EXPOSE_FUNCTION_I_ARGS_1' 
    5039 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    5040 ^ 
    5041 program_source:164:9: note: expanded from macro '\ 
    5042 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    5043 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5044 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5045 program_source:293:1: error: illegal string literal in 'asm' 
    5046 DECLARE_REDUCTION_BASE(broadcast_first) 
    5047 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5048 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    5049 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    5050 ^ 
    5051 program_source:263:44: note: expanded from macro '\ 
    5052 DECLARE_I_REDUCTION_BASE' 
    5053 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    5054 ^ 
    5055 program_source:218:51: note: expanded from macro '\ 
    5056 EXPOSE_FUNCTION_I_ARGS_1' 
    5057 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    5058 ^ 
    5059 program_source:164:9: note: expanded from macro '\ 
    5060 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    5061 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5062 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5063 program_source:293:1: error: illegal string literal in 'asm' 
    5064 DECLARE_REDUCTION_BASE(broadcast_first) 
    5065 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5066 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    5067 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    5068 ^ 
    5069 program_source:264:43: note: expanded from macro '\ 
    5070 DECLARE_I_REDUCTION_BASE' 
    5071 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    5072 ^ 
    5073 program_source:217:40: note: expanded from macro '\ 
    5074 EXPOSE_FUNCTION_I_ARGS_1' 
    5075 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    5076 ^ 
    5077 program_source:164:9: note: expanded from macro '\ 
    5078 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    5079 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5080 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5081 program_source:293:1: error: illegal string literal in 'asm' 
    5082 DECLARE_REDUCTION_BASE(broadcast_first) 
    5083 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5084 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    5085 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    5086 ^ 
    5087 program_source:264:43: note: expanded from macro '\ 
    5088 DECLARE_I_REDUCTION_BASE' 
    5089 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    5090 ^ 
    5091 program_source:218:51: note: expanded from macro '\ 
    5092 EXPOSE_FUNCTION_I_ARGS_1' 
    5093 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    5094 ^ 
    5095 program_source:164:9: note: expanded from macro '\ 
    5096 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    5097 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5098 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5099 program_source:293:1: error: illegal string literal in 'asm' 
    5100 DECLARE_REDUCTION_BASE(broadcast_first) 
    5101 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5102 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    5103 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    5104 ^ 
    5105 program_source:267:44: note: expanded from macro '\ 
    5106 DECLARE_F_REDUCTION_BASE' 
    5107 #define DECLARE_F_REDUCTION_BASE(METAL_OP) \ 
    5108 ^ 
    5109 program_source:221:40: note: expanded from macro '\ 
    5110 EXPOSE_FUNCTION_F_ARGS_1' 
    5111 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    5112 ^ 
    5113 program_source:164:9: note: expanded from macro '\ 
    5114 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    5115 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5116 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5117 program_source:293:1: error: illegal string literal in 'asm' 
    5118 DECLARE_REDUCTION_BASE(broadcast_first) 
    5119 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5120 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    5121 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    5122 ^ 
    5123 program_source:268:43: note: expanded from macro '\ 
    5124 DECLARE_F_REDUCTION_BASE' 
    5125 EXPOSE_FUNCTION_F_ARGS_1(quad_##METAL_OP) \ 
    5126 ^ 
    5127 program_source:221:40: note: expanded from macro '\ 
    5128 EXPOSE_FUNCTION_F_ARGS_1' 
    5129 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    5130 ^ 
    5131 program_source:164:9: note: expanded from macro '\ 
    5132 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    5133 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5134 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5135 program_source:295:1: error: illegal string literal in 'asm' 
    5136 DECLARE_SHUFFLE_BASE(shuffle) 
    5137 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5138 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5139 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5140 ^ 
    5141 program_source:224:38: note: expanded from macro '\ 
    5142 EXPOSE_FUNCTION_ARGS_2' 
    5143 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    5144 ^ 
    5145 program_source:168:9: note: expanded from macro '\ 
    5146 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5147 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5148 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5149 program_source:295:1: error: illegal string literal in 'asm' 
    5150 DECLARE_SHUFFLE_BASE(shuffle) 
    5151 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5152 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5153 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5154 ^ 
    5155 program_source:225:51: note: expanded from macro '\ 
    5156 EXPOSE_FUNCTION_ARGS_2' 
    5157 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    5158 ^ 
    5159 program_source:168:9: note: expanded from macro '\ 
    5160 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    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:226:52: note: expanded from macro '\ 
    5170 EXPOSE_FUNCTION_ARGS_2' 
    5171 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    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:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5181 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5182 ^ 
    5183 program_source:224:38: note: expanded from macro '\ 
    5184 EXPOSE_FUNCTION_ARGS_2' 
    5185 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    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:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5195 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5196 ^ 
    5197 program_source:225:51: note: expanded from macro '\ 
    5198 EXPOSE_FUNCTION_ARGS_2' 
    5199 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.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:226:52: note: expanded from macro '\ 
    5212 EXPOSE_FUNCTION_ARGS_2' 
    5213 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    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:296:1: error: illegal string literal in 'asm' 
    5220 DECLARE_SHUFFLE_BASE(shuffle_xor) 
    5221 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5222 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5223 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5224 ^ 
    5225 program_source:224:38: note: expanded from macro '\ 
    5226 EXPOSE_FUNCTION_ARGS_2' 
    5227 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    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:296:1: error: illegal string literal in 'asm' 
    5234 DECLARE_SHUFFLE_BASE(shuffle_xor) 
    5235 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5236 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5237 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5238 ^ 
    5239 program_source:225:51: note: expanded from macro '\ 
    5240 EXPOSE_FUNCTION_ARGS_2' 
    5241 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.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:226:52: note: expanded from macro '\ 
    5254 EXPOSE_FUNCTION_ARGS_2' 
    5255 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    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:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5265 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5266 ^ 
    5267 program_source:224:38: note: expanded from macro '\ 
    5268 EXPOSE_FUNCTION_ARGS_2' 
    5269 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    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:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5279 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5280 ^ 
    5281 program_source:225:51: note: expanded from macro '\ 
    5282 EXPOSE_FUNCTION_ARGS_2' 
    5283 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.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:226:52: note: expanded from macro '\ 
    5296 EXPOSE_FUNCTION_ARGS_2' 
    5297 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    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:297:1: error: illegal string literal in 'asm' 
    5304 DECLARE_SHUFFLE_BASE(shuffle_up) 
    5305 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5306 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5307 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5308 ^ 
    5309 program_source:224:38: note: expanded from macro '\ 
    5310 EXPOSE_FUNCTION_ARGS_2' 
    5311 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    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:297:1: error: illegal string literal in 'asm' 
    5318 DECLARE_SHUFFLE_BASE(shuffle_up) 
    5319 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5320 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5321 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5322 ^ 
    5323 program_source:225:51: note: expanded from macro '\ 
    5324 EXPOSE_FUNCTION_ARGS_2' 
    5325 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.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:226:52: note: expanded from macro '\ 
    5338 EXPOSE_FUNCTION_ARGS_2' 
    5339 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    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:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5349 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5350 ^ 
    5351 program_source:224:38: note: expanded from macro '\ 
    5352 EXPOSE_FUNCTION_ARGS_2' 
    5353 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    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:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5363 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5364 ^ 
    5365 program_source:225:51: note: expanded from macro '\ 
    5366 EXPOSE_FUNCTION_ARGS_2' 
    5367 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.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:226:52: note: expanded from macro '\ 
    5380 EXPOSE_FUNCTION_ARGS_2' 
    5381 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    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:298:1: error: illegal string literal in 'asm' 
    5388 DECLARE_SHUFFLE_BASE(shuffle_down) 
    5389 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5390 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5391 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5392 ^ 
    5393 program_source:224:38: note: expanded from macro '\ 
    5394 EXPOSE_FUNCTION_ARGS_2' 
    5395 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    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:298:1: error: illegal string literal in 'asm' 
    5402 DECLARE_SHUFFLE_BASE(shuffle_down) 
    5403 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5404 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5405 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5406 ^ 
    5407 program_source:225:51: note: expanded from macro '\ 
    5408 EXPOSE_FUNCTION_ARGS_2' 
    5409 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.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:226:52: note: expanded from macro '\ 
    5422 EXPOSE_FUNCTION_ARGS_2' 
    5423 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    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:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5433 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5434 ^ 
    5435 program_source:224:38: note: expanded from macro '\ 
    5436 EXPOSE_FUNCTION_ARGS_2' 
    5437 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    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:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5447 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5448 ^ 
    5449 program_source:225:51: note: expanded from macro '\ 
    5450 EXPOSE_FUNCTION_ARGS_2' 
    5451 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.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:226:52: note: expanded from macro '\ 
    5464 EXPOSE_FUNCTION_ARGS_2' 
    5465 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    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:299:1: error: illegal string literal in 'asm' 
    5472 DECLARE_SHUFFLE_BASE(shuffle_rotate_up) 
    5473 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5474 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5475 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5476 ^ 
    5477 program_source:224:38: note: expanded from macro '\ 
    5478 EXPOSE_FUNCTION_ARGS_2' 
    5479 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    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:299:1: error: illegal string literal in 'asm' 
    5486 DECLARE_SHUFFLE_BASE(shuffle_rotate_up) 
    5487 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5488 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5489 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5490 ^ 
    5491 program_source:225:51: note: expanded from macro '\ 
    5492 EXPOSE_FUNCTION_ARGS_2' 
    5493 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.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:226:52: note: expanded from macro '\ 
    5506 EXPOSE_FUNCTION_ARGS_2' 
    5507 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    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:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5517 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5518 ^ 
    5519 program_source:224:38: note: expanded from macro '\ 
    5520 EXPOSE_FUNCTION_ARGS_2' 
    5521 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    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:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5531 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5532 ^ 
    5533 program_source:225:51: note: expanded from macro '\ 
    5534 EXPOSE_FUNCTION_ARGS_2' 
    5535 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.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:226:52: note: expanded from macro '\ 
    5548 EXPOSE_FUNCTION_ARGS_2' 
    5549 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    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:300:1: error: illegal string literal in 'asm' 
    5556 DECLARE_SHUFFLE_BASE(shuffle_rotate_down) 
    5557 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5558 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5559 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5560 ^ 
    5561 program_source:224:38: note: expanded from macro '\ 
    5562 EXPOSE_FUNCTION_ARGS_2' 
    5563 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    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:300:1: error: illegal string literal in 'asm' 
    5570 DECLARE_SHUFFLE_BASE(shuffle_rotate_down) 
    5571 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5572 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5573 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5574 ^ 
    5575 program_source:225:51: note: expanded from macro '\ 
    5576 EXPOSE_FUNCTION_ARGS_2' 
    5577 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.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:226:52: note: expanded from macro '\ 
    5590 EXPOSE_FUNCTION_ARGS_2' 
    5591 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    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:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5601 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5602 ^ 
    5603 program_source:224:38: note: expanded from macro '\ 
    5604 EXPOSE_FUNCTION_ARGS_2' 
    5605 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    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:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5615 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5616 ^ 
    5617 program_source:225:51: note: expanded from macro '\ 
    5618 EXPOSE_FUNCTION_ARGS_2' 
    5619 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.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:226:52: note: expanded from macro '\ 
    5632 EXPOSE_FUNCTION_ARGS_2' 
    5633 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    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:302:1: error: illegal string literal in 'asm' 
    5640 EXPOSE_FUNCTION_I_ARGS_3(ctz) 
    5641 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5642 program_source:229:40: note: expanded from macro 'EXPOSE_FUNCTION_I_ARGS_3' 
    5643 #define EXPOSE_FUNCTION_I_ARGS_3(EXPR) \ 
    5644 ^ 
    5645 program_source:172:9: note: expanded from macro '\ 
    5646 EXPOSE_FUNCTION_OVERLOAD_ARGS_3' 
    5647 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5648 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5649 program_source:302:1: error: implicit declaration of function '___metal_ctz'
    5650 is invalid in OpenCL 
    5651 program_source:229:40: note: expanded from macro 'EXPOSE_FUNCTION_I_ARGS_3' 
    5652 #define EXPOSE_FUNCTION_I_ARGS_3(EXPR) \ 
    5653 ^ 
    5654 program_source:175:10: note: expanded from macro '\ 
    5655 EXPOSE_FUNCTION_OVERLOAD_ARGS_3' 
    5656 return ___metal_##EXPR(x, false); \ 
    5657 ^ 
    5658 :12:1: note: expanded from here 
    5659 ___metal_ctz 
    5660 ^ 
    5661 program_source:302:1: error: illegal string literal in 'asm' 
    5662 EXPOSE_FUNCTION_I_ARGS_3(ctz) 
    5663 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5664 program_source:230:49: note: expanded from macro 'EXPOSE_FUNCTION_I_ARGS_3' 
    5665 EXPOSE_FUNCTION_OVERLOAD_ARGS_3(EXPR, int, i32) \ 
    5666 ^ 
    5667 program_source:172:9: note: expanded from macro '\ 
    5668 EXPOSE_FUNCTION_OVERLOAD_ARGS_3' 
    5669 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5670 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5671 program_source:302:1: error: implicit declaration of function '___metal_ctz'
    5672 is invalid in OpenCL 
    5673 program_source:230:49: note: expanded from macro 'EXPOSE_FUNCTION_I_ARGS_3' 
    5674 EXPOSE_FUNCTION_OVERLOAD_ARGS_3(EXPR, int, i32) \ 
    5675 ^ 
    5676 program_source:175:10: note: expanded from macro '\ 
    5677 EXPOSE_FUNCTION_OVERLOAD_ARGS_3' 
    5678 return ___metal_##EXPR(x, false); \ 
    5679 ^ 
    5680 :16:1: note: expanded from here 
    5681 ___metal_ctz 
    5682 ^ 
    5683 program_source:317:1: error: implicit declaration of function 'simd_sum' is
    5684 invalid in OpenCL 
    5685 DECLARE_REDUCTION_UNIFORM(sum, reduce_add) 
    5686 ^ 
    5687 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5688 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5689 ^ 
    5690 program_source:305:26: note: expanded from macro '\ 
    5691 DECLARE_I_REDUCTION_UNIFORM' 
    5692 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5693 ^ 
    5694 :17:1: note: expanded from here 
    5695 simd_sum 
    5696 ^ 
    5697 program_source:317:1: error: implicit declaration of function 'simd_sum' is
    5698 invalid in OpenCL 
    5699 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5700 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5701 ^ 
    5702 program_source:305:26: note: expanded from macro '\ 
    5703 DECLARE_I_REDUCTION_UNIFORM' 
    5704 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5705 ^ 
    5706 :17:1: note: expanded from here 
    5707 simd_sum 
    5708 ^ 
    5709 program_source:317:1: error: implicit declaration of function 'simd_sum' is
    5710 invalid in OpenCL 
    5711 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5712 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5713 ^ 
    5714 program_source:308:26: note: expanded from macro '\ 
    5715 DECLARE_F_REDUCTION_UNIFORM' 
    5716 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5717 ^ 
    5718 :19:1: note: expanded from here 
    5719 simd_sum 
    5720 ^ 
    5721 program_source:318:1: error: implicit declaration of function
    5722 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    5723 DECLARE_REDUCTION_UNIFORM(prefix_inclusive_sum, scan_inclusive_add) 
    5724 ^ 
    5725 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5726 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5727 ^ 
    5728 program_source:305:26: note: expanded from macro '\ 
    5729 DECLARE_I_REDUCTION_UNIFORM' 
    5730 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5731 ^ 
    5732 :21:1: note: expanded from here 
    5733 simd_prefix_inclusive_sum 
    5734 ^ 
    5735 program_source:318:1: error: implicit declaration of function
    5736 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    5737 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5738 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5739 ^ 
    5740 program_source:305:26: note: expanded from macro '\ 
    5741 DECLARE_I_REDUCTION_UNIFORM' 
    5742 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5743 ^ 
    5744 :21:1: note: expanded from here 
    5745 simd_prefix_inclusive_sum 
    5746 ^ 
    5747 program_source:318:1: error: implicit declaration of function
    5748 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    5749 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5750 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5751 ^ 
    5752 program_source:308:26: note: expanded from macro '\ 
    5753 DECLARE_F_REDUCTION_UNIFORM' 
    5754 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5755 ^ 
    5756 :23:1: note: expanded from here 
    5757 simd_prefix_inclusive_sum 
    5758 ^ 
    5759 program_source:319:1: error: implicit declaration of function
    5760 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    5761 DECLARE_REDUCTION_UNIFORM(prefix_exclusive_sum, scan_exclusive_add) 
    5762 ^ 
    5763 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5764 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5765 ^ 
    5766 program_source:305:26: note: expanded from macro '\ 
    5767 DECLARE_I_REDUCTION_UNIFORM' 
    5768 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5769 ^ 
    5770 :25:1: note: expanded from here 
    5771 simd_prefix_exclusive_sum 
    5772 ^ 
    5773 program_source:319:1: error: implicit declaration of function
    5774 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    5775 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5776 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5777 ^ 
    5778 program_source:305:26: note: expanded from macro '\ 
    5779 DECLARE_I_REDUCTION_UNIFORM' 
    5780 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5781 ^ 
    5782 :25:1: note: expanded from here 
    5783 simd_prefix_exclusive_sum 
    5784 ^ 
    5785 program_source:319:1: error: implicit declaration of function
    5786 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    5787 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5788 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5789 ^ 
    5790 program_source:308:26: note: expanded from macro '\ 
    5791 DECLARE_F_REDUCTION_UNIFORM' 
    5792 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5793 ^ 
    5794 :27:1: note: expanded from here 
    5795 simd_prefix_exclusive_sum 
    5796 ^ 
    5797 program_source:320:1: error: implicit declaration of function 'simd_min' is
    5798 invalid in OpenCL 
    5799 DECLARE_REDUCTION_UNIFORM(min, reduce_min) 
    5800 ^ 
    5801 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5802 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5803 ^ 
    5804 program_source:305:26: note: expanded from macro '\ 
    5805 DECLARE_I_REDUCTION_UNIFORM' 
    5806 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5807 ^ 
    5808 :29:1: note: expanded from here 
    5809 simd_min 
    5810 ^ 
    5811 program_source:320:1: error: implicit declaration of function 'simd_min' is
    5812 invalid in OpenCL 
    5813 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5814 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5815 ^ 
    5816 program_source:305:26: note: expanded from macro '\ 
    5817 DECLARE_I_REDUCTION_UNIFORM' 
    5818 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5819 ^ 
    5820 :29:1: note: expanded from here 
    5821 simd_min 
    5822 ^ 
    5823 program_source:320:1: error: implicit declaration of function 'simd_min' is
    5824 invalid in OpenCL 
    5825 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5826 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5827 ^ 
    5828 program_source:308:26: note: expanded from macro '\ 
    5829 DECLARE_F_REDUCTION_UNIFORM' 
    5830 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5831 ^ 
    5832 :31:1: note: expanded from here 
    5833 simd_min 
    5834 ^ 
    5835 program_source:321:1: error: implicit declaration of function 'simd_max' is
    5836 invalid in OpenCL 
    5837 DECLARE_REDUCTION_UNIFORM(max, reduce_max) 
    5838 ^ 
    5839 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5840 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5841 ^ 
    5842 program_source:305:26: note: expanded from macro '\ 
    5843 DECLARE_I_REDUCTION_UNIFORM' 
    5844 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5845 ^ 
    5846 :33:1: note: expanded from here 
    5847 simd_max 
    5848 ^ 
    5849 program_source:321:1: error: implicit declaration of function 'simd_max' is
    5850 invalid in OpenCL 
    5851 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5852 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5853 ^ 
    5854 program_source:305:26: note: expanded from macro '\ 
    5855 DECLARE_I_REDUCTION_UNIFORM' 
    5856 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5857 ^ 
    5858 :33:1: note: expanded from here 
    5859 simd_max 
    5860 ^ 
    5861 program_source:321:1: error: implicit declaration of function 'simd_max' is
    5862 invalid in OpenCL 
    5863 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5864 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5865 ^ 
    5866 program_source:308:26: note: expanded from macro '\ 
    5867 DECLARE_F_REDUCTION_UNIFORM' 
    5868 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5869 ^ 
    5870 :35:1: note: expanded from here 
    5871 simd_max 
    5872 ^ 
    5873 program_source:323:1: error: implicit declaration of function 'simd_shuffle'
    5874 is invalid in OpenCL 
    5875 DECLARE_SHUFFLE_UNIFORM(shuffle, shuffle) 
    5876 ^ 
    5877 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5878 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5879 ^ 
    5880 :37:1: note: expanded from here 
    5881 simd_shuffle 
    5882 ^ 
    5883 program_source:323:1: error: implicit declaration of function 'simd_shuffle'
    5884 is invalid in OpenCL 
    5885 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5886 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5887 ^ 
    5888 :37:1: note: expanded from here 
    5889 simd_shuffle 
    5890 ^ 
    5891 program_source:323:1: error: implicit declaration of function 'simd_shuffle'
    5892 is invalid in OpenCL 
    5893 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5894 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5895 ^ 
    5896 :37:1: note: expanded from here 
    5897 simd_shuffle 
    5898 ^ 
    5899 program_source:324:1: error: implicit declaration of function
    5900 'simd_shuffle_xor' is invalid in OpenCL 
    5901 DECLARE_SHUFFLE_UNIFORM(shuffle_xor, shuffle_xor) 
    5902 ^ 
    5903 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5904 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5905 ^ 
    5906 :39:1: note: expanded from here 
    5907 simd_shuffle_xor 
    5908 ^ 
    5909 program_source:324:1: error: implicit declaration of function
    5910 'simd_shuffle_xor' is invalid in OpenCL 
    5911 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5912 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5913 ^ 
    5914 :39:1: note: expanded from here 
    5915 simd_shuffle_xor 
    5916 ^ 
    5917 program_source:324:1: error: implicit declaration of function
    5918 'simd_shuffle_xor' is invalid in OpenCL 
    5919 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5920 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5921 ^ 
    5922 :39:1: note: expanded from here 
    5923 simd_shuffle_xor 
    5924 ^ 
    5925 program_source:325:1: error: implicit declaration of function
    5926 'simd_shuffle_up' is invalid in OpenCL 
    5927 DECLARE_SHUFFLE_UNIFORM(shuffle_up, shuffle_up) 
    5928 ^ 
    5929 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5930 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5931 ^ 
    5932 :41:1: note: expanded from here 
    5933 simd_shuffle_up 
    5934 ^ 
    5935 program_source:325:1: error: implicit declaration of function
    5936 'simd_shuffle_up' is invalid in OpenCL 
    5937 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5938 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5939 ^ 
    5940 :41:1: note: expanded from here 
    5941 simd_shuffle_up 
    5942 ^ 
    5943 program_source:325:1: error: implicit declaration of function
    5944 'simd_shuffle_up' is invalid in OpenCL 
    5945 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5946 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5947 ^ 
    5948 :41:1: note: expanded from here 
    5949 simd_shuffle_up 
    5950 ^ 
    5951 program_source:326:1: error: implicit declaration of function
    5952 'simd_shuffle_down' is invalid in OpenCL 
    5953 DECLARE_SHUFFLE_UNIFORM(shuffle_down, shuffle_down) 
    5954 ^ 
    5955 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5956 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5957 ^ 
    5958 :43:1: note: expanded from here 
    5959 simd_shuffle_down 
    5960 ^ 
    5961 program_source:326:1: error: implicit declaration of function
    5962 'simd_shuffle_down' is invalid in OpenCL 
    5963 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5964 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5965 ^ 
    5966 :43:1: note: expanded from here 
    5967 simd_shuffle_down 
    5968 ^ 
    5969 program_source:326:1: error: implicit declaration of function
    5970 'simd_shuffle_down' is invalid in OpenCL 
    5971 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5972 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5973 ^ 
    5974 :43:1: note: expanded from here 
    5975 simd_shuffle_down 
    5976 ^ 
    5977 program_source:327:1: error: implicit declaration of function
    5978 'simd_shuffle_rotate_down' is invalid in OpenCL 
    5979 DECLARE_SHUFFLE_UNIFORM(shuffle_rotate_down, rotate) 
    5980 ^ 
    5981 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5982 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5983 ^ 
    5984 :45:1: note: expanded from here 
    5985 simd_shuffle_rotate_down 
    5986 ^ 
    5987 program_source:327:1: error: implicit declaration of function
    5988 'simd_shuffle_rotate_down' is invalid in OpenCL 
    5989 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5990 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5991 ^ 
    5992 :45:1: note: expanded from here 
    5993 simd_shuffle_rotate_down 
    5994 ^ 
    5995 program_source:327:1: error: implicit declaration of function
    5996 'simd_shuffle_rotate_down' is invalid in OpenCL 
    5997 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    5998 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5999 ^ 
    6000 :45:1: note: expanded from here 
    6001 simd_shuffle_rotate_down 
    6002 ^ 
    6003 program_source:329:1: error: implicit declaration of function 'simd_broadcast'
    6004 is invalid in OpenCL 
    6005 DECLARE_SHUFFLE_UNIFORM(broadcast, broadcast) 
    6006 ^ 
    6007 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    6008 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6009 ^ 
    6010 :47:1: note: expanded from here 
    6011 simd_broadcast 
    6012 ^ 
    6013 program_source:329:1: error: implicit declaration of function 'simd_broadcast'
    6014 is invalid in OpenCL 
    6015 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    6016 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6017 ^ 
    6018 :47:1: note: expanded from here 
    6019 simd_broadcast 
    6020 ^ 
    6021 program_source:329:1: error: implicit declaration of function 'simd_broadcast'
    6022 is invalid in OpenCL 
    6023 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    6024 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6025 ^ 
    6026 :47:1: note: expanded from here 
    6027 simd_broadcast 
    6028 ^ 
    6029 program_source:330:1: error: implicit declaration of function
    6030 'simd_broadcast_first' is invalid in OpenCL 
    6031 DECLARE_REDUCTION_UNIFORM(broadcast_first, broadcast_first) 
    6032 ^ 
    6033 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    6034 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6035 ^ 
    6036 program_source:305:26: note: expanded from macro '\ 
    6037 DECLARE_I_REDUCTION_UNIFORM' 
    6038 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6039 ^ 
    6040 :49:1: note: expanded from here 
    6041 simd_broadcast_first 
    6042 ^ 
    6043 program_source:330:1: error: implicit declaration of function
    6044 'simd_broadcast_first' is invalid in OpenCL 
    6045 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    6046 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6047 ^ 
    6048 program_source:305:26: note: expanded from macro '\ 
    6049 DECLARE_I_REDUCTION_UNIFORM' 
    6050 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6051 ^ 
    6052 :49:1: note: expanded from here 
    6053 simd_broadcast_first 
    6054 ^ 
    6055 program_source:330:1: error: implicit declaration of function
    6056 'simd_broadcast_first' is invalid in OpenCL 
    6057 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    6058 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6059 ^ 
    6060 program_source:308:26: note: expanded from macro '\ 
    6061 DECLARE_F_REDUCTION_UNIFORM' 
    6062 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6063 ^ 
    6064 :51:1: note: expanded from here 
    6065 simd_broadcast_first 
    6066 ^ 
    6067 program_source:345:1: error: implicit declaration of function 'simd_sum' is
    6068 invalid in OpenCL 
    6069 DECLARE_REDUCTION_NON_UNIFORM(sum, reduce_add) 
    6070 ^ 
    6071 program_source:338:60: note: expanded from macro
    6072 'DECLARE_REDUCTION_NON_UNIFORM' 
    6073 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6074 ^ 
    6075 program_source:333:26: note: expanded from macro '\ 
    6076 DECLARE_I_REDUCTION_NON_UNIFORM' 
    6077 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6078 ^ 
    6079 :53:1: note: expanded from here 
    6080 simd_sum 
    6081 ^ 
    6082 program_source:345:1: error: implicit declaration of function 'simd_sum' is
    6083 invalid in OpenCL 
    6084 program_source:338:60: note: expanded from macro
    6085 'DECLARE_REDUCTION_NON_UNIFORM' 
    6086 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6087 ^ 
    6088 program_source:333:26: note: expanded from macro '\ 
    6089 DECLARE_I_REDUCTION_NON_UNIFORM' 
    6090 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6091 ^ 
    6092 :53:1: note: expanded from here 
    6093 simd_sum 
    6094 ^ 
    6095 program_source:345:1: error: implicit declaration of function 'simd_sum' is
    6096 invalid in OpenCL 
    6097 program_source:339:54: note: expanded from macro
    6098 'DECLARE_REDUCTION_NON_UNIFORM' 
    6099 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6100 ^ 
    6101 program_source:336:26: note: expanded from macro '\ 
    6102 DECLARE_F_REDUCTION_NON_UNIFORM' 
    6103 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6104 ^ 
    6105 :55:1: note: expanded from here 
    6106 simd_sum 
    6107 ^ 
    6108 program_source:346:1: error: implicit declaration of function
    6109 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    6110 DECLARE_REDUCTION_NON_UNIFORM(prefix_inclusive_sum, scan_inclusive_add) 
    6111 ^ 
    6112 program_source:338:60: note: expanded from macro
    6113 'DECLARE_REDUCTION_NON_UNIFORM' 
    6114 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6115 ^ 
    6116 program_source:333:26: note: expanded from macro '\ 
    6117 DECLARE_I_REDUCTION_NON_UNIFORM' 
    6118 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6119 ^ 
    6120 :57:1: note: expanded from here 
    6121 simd_prefix_inclusive_sum 
    6122 ^ 
    6123 program_source:346:1: error: implicit declaration of function
    6124 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    6125 program_source:338:60: note: expanded from macro
    6126 'DECLARE_REDUCTION_NON_UNIFORM' 
    6127 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6128 ^ 
    6129 program_source:333:26: note: expanded from macro '\ 
    6130 DECLARE_I_REDUCTION_NON_UNIFORM' 
    6131 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6132 ^ 
    6133 :57:1: note: expanded from here 
    6134 simd_prefix_inclusive_sum 
    6135 ^ 
    6136 program_source:346:1: error: implicit declaration of function
    6137 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    6138 program_source:339:54: note: expanded from macro
    6139 'DECLARE_REDUCTION_NON_UNIFORM' 
    6140 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6141 ^ 
    6142 program_source:336:26: note: expanded from macro '\ 
    6143 DECLARE_F_REDUCTION_NON_UNIFORM' 
    6144 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6145 ^ 
    6146 :59:1: note: expanded from here 
    6147 simd_prefix_inclusive_sum 
    6148 ^ 
    6149 program_source:347:1: error: implicit declaration of function
    6150 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    6151 DECLARE_REDUCTION_NON_UNIFORM(prefix_exclusive_sum, scan_exclusive_add) 
    6152 ^ 
    6153 program_source:338:60: note: expanded from macro
    6154 'DECLARE_REDUCTION_NON_UNIFORM' 
    6155 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6156 ^ 
    6157 program_source:333:26: note: expanded from macro '\ 
    6158 DECLARE_I_REDUCTION_NON_UNIFORM' 
    6159 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6160 ^ 
    6161 :61:1: note: expanded from here 
    6162 simd_prefix_exclusive_sum 
    6163 ^ 
    6164 program_source:347:1: error: implicit declaration of function
    6165 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    6166 program_source:338:60: note: expanded from macro
    6167 'DECLARE_REDUCTION_NON_UNIFORM' 
    6168 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6169 ^ 
    6170 program_source:333:26: note: expanded from macro '\ 
    6171 DECLARE_I_REDUCTION_NON_UNIFORM' 
    6172 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6173 ^ 
    6174 :61:1: note: expanded from here 
    6175 simd_prefix_exclusive_sum 
    6176 ^ 
    6177 program_source:347:1: error: implicit declaration of function
    6178 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    6179 program_source:339:54: note: expanded from macro
    6180 'DECLARE_REDUCTION_NON_UNIFORM' 
    6181 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6182 ^ 
    6183 program_source:336:26: note: expanded from macro '\ 
    6184 DECLARE_F_REDUCTION_NON_UNIFORM' 
    6185 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6186 ^ 
    6187 :63:1: note: expanded from here 
    6188 simd_prefix_exclusive_sum 
    6189 ^ 
    6190 program_source:348:1: error: implicit declaration of function 'simd_min' is
    6191 invalid in OpenCL 
    6192 DECLARE_REDUCTION_NON_UNIFORM(min, reduce_min) 
    6193 ^ 
    6194 program_source:338:60: note: expanded from macro
    6195 'DECLARE_REDUCTION_NON_UNIFORM' 
    6196 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6197 ^ 
    6198 program_source:333:26: note: expanded from macro '\ 
    6199 DECLARE_I_REDUCTION_NON_UNIFORM' 
    6200 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6201 ^ 
    6202 :65:1: note: expanded from here 
    6203 simd_min 
    6204 ^ 
    6205 program_source:348:1: error: implicit declaration of function 'simd_min' is
    6206 invalid in OpenCL 
    6207 program_source:338:60: note: expanded from macro
    6208 'DECLARE_REDUCTION_NON_UNIFORM' 
    6209 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6210 ^ 
    6211 program_source:333:26: note: expanded from macro '\ 
    6212 DECLARE_I_REDUCTION_NON_UNIFORM' 
    6213 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6214 ^ 
    6215 :65:1: note: expanded from here 
    6216 simd_min 
    6217 ^ 
    6218 program_source:348:1: error: implicit declaration of function 'simd_min' is
    6219 invalid in OpenCL 
    6220 program_source:339:54: note: expanded from macro
    6221 'DECLARE_REDUCTION_NON_UNIFORM' 
    6222 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6223 ^ 
    6224 program_source:336:26: note: expanded from macro '\ 
    6225 DECLARE_F_REDUCTION_NON_UNIFORM' 
    6226 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6227 ^ 
    6228 :67:1: note: expanded from here 
    6229 simd_min 
    6230 ^ 
    6231 program_source:349:1: error: implicit declaration of function 'simd_max' is
    6232 invalid in OpenCL 
    6233 DECLARE_REDUCTION_NON_UNIFORM(max, reduce_max) 
    6234 ^ 
    6235 program_source:338:60: note: expanded from macro
    6236 'DECLARE_REDUCTION_NON_UNIFORM' 
    6237 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6238 ^ 
    6239 program_source:333:26: note: expanded from macro '\ 
    6240 DECLARE_I_REDUCTION_NON_UNIFORM' 
    6241 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6242 ^ 
    6243 :69:1: note: expanded from here 
    6244 simd_max 
    6245 ^ 
     1664[deleted to fit within ticket limits]
     1665
    62461666program_source:349:1: error: implicit declaration of function 'simd_max' is
    62471667invalid in OpenCL