Changes between Initial Version and Version 1 of Ticket #19017


Ignore:
Timestamp:
Oct 6, 2025, 11:15:36 AM (2 months ago)
Author:
Eric Pettersen
Comment:

Legend:

Unmodified
Added
Removed
Modified

  • Ticket #19017

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

    initial v1  
    15091509DECLARE_I_REDUCTION_BASE(xor) 
    15101510^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1511 program_source:263:44: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    1512 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    1513 ^ 
    1514 program_source:218:51: note: expanded from macro '\ 
    1515 EXPOSE_FUNCTION_I_ARGS_1' 
    1516 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    1517 ^ 
    1518 program_source:164:9: note: expanded from macro '\ 
    1519 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    1520 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1521 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1522 program_source:290:1: error: illegal string literal in 'asm' 
    1523 DECLARE_I_REDUCTION_BASE(xor) 
    1524 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1525 program_source:264:43: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    1526 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    1527 ^ 
    1528 program_source:217:40: note: expanded from macro '\ 
    1529 EXPOSE_FUNCTION_I_ARGS_1' 
    1530 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    1531 ^ 
    1532 program_source:164:9: note: expanded from macro '\ 
    1533 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    1534 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1535 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1536 program_source:290:1: error: illegal string literal in 'asm' 
    1537 DECLARE_I_REDUCTION_BASE(xor) 
    1538 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1539 program_source:264:43: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    1540 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    1541 ^ 
    1542 program_source:218:51: note: expanded from macro '\ 
    1543 EXPOSE_FUNCTION_I_ARGS_1' 
    1544 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    1545 ^ 
    1546 program_source:164:9: note: expanded from macro '\ 
    1547 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    1548 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1549 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1550 program_source:292:1: error: illegal string literal in 'asm' 
    1551 DECLARE_SHUFFLE_BASE(broadcast) 
    1552 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1553 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1554 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    1555 ^ 
    1556 program_source:224:38: note: expanded from macro '\ 
    1557 EXPOSE_FUNCTION_ARGS_2' 
    1558 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    1559 ^ 
    1560 program_source:168:9: note: expanded from macro '\ 
    1561 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1562 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1563 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1564 program_source:292:1: error: illegal string literal in 'asm' 
    1565 DECLARE_SHUFFLE_BASE(broadcast) 
    1566 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1567 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1568 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    1569 ^ 
    1570 program_source:225:51: note: expanded from macro '\ 
    1571 EXPOSE_FUNCTION_ARGS_2' 
    1572 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    1573 ^ 
    1574 program_source:168:9: note: expanded from macro '\ 
    1575 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1576 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1577 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1578 program_source:292:1: error: illegal string literal in 'asm' 
    1579 DECLARE_SHUFFLE_BASE(broadcast) 
    1580 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1581 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1582 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    1583 ^ 
    1584 program_source:226:52: note: expanded from macro '\ 
    1585 EXPOSE_FUNCTION_ARGS_2' 
    1586 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    1587 ^ 
    1588 program_source:168:9: note: expanded from macro '\ 
    1589 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1590 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1591 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1592 program_source:292:1: error: illegal string literal in 'asm' 
    1593 DECLARE_SHUFFLE_BASE(broadcast) 
    1594 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1595 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1596 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    1597 ^ 
    1598 program_source:224:38: note: expanded from macro '\ 
    1599 EXPOSE_FUNCTION_ARGS_2' 
    1600 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    1601 ^ 
    1602 program_source:168:9: note: expanded from macro '\ 
    1603 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1604 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1605 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1606 program_source:292:1: error: illegal string literal in 'asm' 
    1607 DECLARE_SHUFFLE_BASE(broadcast) 
    1608 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1609 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1610 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    1611 ^ 
    1612 program_source:225:51: note: expanded from macro '\ 
    1613 EXPOSE_FUNCTION_ARGS_2' 
    1614 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    1615 ^ 
    1616 program_source:168:9: note: expanded from macro '\ 
    1617 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1618 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1619 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1620 program_source:292:1: error: illegal string literal in 'asm' 
    1621 DECLARE_SHUFFLE_BASE(broadcast) 
    1622 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1623 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1624 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    1625 ^ 
    1626 program_source:226:52: note: expanded from macro '\ 
    1627 EXPOSE_FUNCTION_ARGS_2' 
    1628 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    1629 ^ 
    1630 program_source:168:9: note: expanded from macro '\ 
    1631 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1632 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1633 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1634 program_source:293:1: error: illegal string literal in 'asm' 
    1635 DECLARE_REDUCTION_BASE(broadcast_first) 
    1636 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1637 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    1638 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    1639 ^ 
    1640 program_source:263:44: note: expanded from macro '\ 
    1641 DECLARE_I_REDUCTION_BASE' 
    1642 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    1643 ^ 
    1644 program_source:217:40: note: expanded from macro '\ 
    1645 EXPOSE_FUNCTION_I_ARGS_1' 
    1646 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    1647 ^ 
    1648 program_source:164:9: note: expanded from macro '\ 
    1649 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    1650 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1651 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1652 program_source:293:1: error: illegal string literal in 'asm' 
    1653 DECLARE_REDUCTION_BASE(broadcast_first) 
    1654 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1655 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    1656 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    1657 ^ 
    1658 program_source:263:44: note: expanded from macro '\ 
    1659 DECLARE_I_REDUCTION_BASE' 
    1660 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    1661 ^ 
    1662 program_source:218:51: note: expanded from macro '\ 
    1663 EXPOSE_FUNCTION_I_ARGS_1' 
    1664 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    1665 ^ 
    1666 program_source:164:9: note: expanded from macro '\ 
    1667 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    1668 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1669 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1670 program_source:293:1: error: illegal string literal in 'asm' 
    1671 DECLARE_REDUCTION_BASE(broadcast_first) 
    1672 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1673 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    1674 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    1675 ^ 
    1676 program_source:264:43: note: expanded from macro '\ 
    1677 DECLARE_I_REDUCTION_BASE' 
    1678 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    1679 ^ 
    1680 program_source:217:40: note: expanded from macro '\ 
    1681 EXPOSE_FUNCTION_I_ARGS_1' 
    1682 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    1683 ^ 
    1684 program_source:164:9: note: expanded from macro '\ 
    1685 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    1686 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1687 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1688 program_source:293:1: error: illegal string literal in 'asm' 
    1689 DECLARE_REDUCTION_BASE(broadcast_first) 
    1690 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1691 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    1692 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    1693 ^ 
    1694 program_source:264:43: note: expanded from macro '\ 
    1695 DECLARE_I_REDUCTION_BASE' 
    1696 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    1697 ^ 
    1698 program_source:218:51: note: expanded from macro '\ 
    1699 EXPOSE_FUNCTION_I_ARGS_1' 
    1700 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    1701 ^ 
    1702 program_source:164:9: note: expanded from macro '\ 
    1703 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    1704 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1705 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1706 program_source:293:1: error: illegal string literal in 'asm' 
    1707 DECLARE_REDUCTION_BASE(broadcast_first) 
    1708 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1709 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    1710 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    1711 ^ 
    1712 program_source:267:44: note: expanded from macro '\ 
    1713 DECLARE_F_REDUCTION_BASE' 
    1714 #define DECLARE_F_REDUCTION_BASE(METAL_OP) \ 
    1715 ^ 
    1716 program_source:221:40: note: expanded from macro '\ 
    1717 EXPOSE_FUNCTION_F_ARGS_1' 
    1718 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    1719 ^ 
    1720 program_source:164:9: note: expanded from macro '\ 
    1721 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    1722 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1723 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1724 program_source:293:1: error: illegal string literal in 'asm' 
    1725 DECLARE_REDUCTION_BASE(broadcast_first) 
    1726 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1727 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    1728 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    1729 ^ 
    1730 program_source:268:43: note: expanded from macro '\ 
    1731 DECLARE_F_REDUCTION_BASE' 
    1732 EXPOSE_FUNCTION_F_ARGS_1(quad_##METAL_OP) \ 
    1733 ^ 
    1734 program_source:221:40: note: expanded from macro '\ 
    1735 EXPOSE_FUNCTION_F_ARGS_1' 
    1736 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    1737 ^ 
    1738 program_source:164:9: note: expanded from macro '\ 
    1739 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    1740 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1741 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1742 program_source:295:1: error: illegal string literal in 'asm' 
    1743 DECLARE_SHUFFLE_BASE(shuffle) 
    1744 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1745 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1746 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    1747 ^ 
    1748 program_source:224:38: note: expanded from macro '\ 
    1749 EXPOSE_FUNCTION_ARGS_2' 
    1750 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    1751 ^ 
    1752 program_source:168:9: note: expanded from macro '\ 
    1753 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1754 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1755 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1756 program_source:295:1: error: illegal string literal in 'asm' 
    1757 DECLARE_SHUFFLE_BASE(shuffle) 
    1758 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1759 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1760 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    1761 ^ 
    1762 program_source:225:51: note: expanded from macro '\ 
    1763 EXPOSE_FUNCTION_ARGS_2' 
    1764 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    1765 ^ 
    1766 program_source:168:9: note: expanded from macro '\ 
    1767 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1768 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1769 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1770 program_source:295:1: error: illegal string literal in 'asm' 
    1771 DECLARE_SHUFFLE_BASE(shuffle) 
    1772 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1773 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1774 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    1775 ^ 
    1776 program_source:226:52: note: expanded from macro '\ 
    1777 EXPOSE_FUNCTION_ARGS_2' 
    1778 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    1779 ^ 
    1780 program_source:168:9: note: expanded from macro '\ 
    1781 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1782 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1783 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1784 program_source:295:1: error: illegal string literal in 'asm' 
    1785 DECLARE_SHUFFLE_BASE(shuffle) 
    1786 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1787 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1788 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    1789 ^ 
    1790 program_source:224:38: note: expanded from macro '\ 
    1791 EXPOSE_FUNCTION_ARGS_2' 
    1792 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    1793 ^ 
    1794 program_source:168:9: note: expanded from macro '\ 
    1795 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1796 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1797 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1798 program_source:295:1: error: illegal string literal in 'asm' 
    1799 DECLARE_SHUFFLE_BASE(shuffle) 
    1800 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1801 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1802 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    1803 ^ 
    1804 program_source:225:51: note: expanded from macro '\ 
    1805 EXPOSE_FUNCTION_ARGS_2' 
    1806 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    1807 ^ 
    1808 program_source:168:9: note: expanded from macro '\ 
    1809 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1810 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1811 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1812 program_source:295:1: error: illegal string literal in 'asm' 
    1813 DECLARE_SHUFFLE_BASE(shuffle) 
    1814 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1815 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1816 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    1817 ^ 
    1818 program_source:226:52: note: expanded from macro '\ 
    1819 EXPOSE_FUNCTION_ARGS_2' 
    1820 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    1821 ^ 
    1822 program_source:168:9: note: expanded from macro '\ 
    1823 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1824 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1825 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1826 program_source:296:1: error: illegal string literal in 'asm' 
    1827 DECLARE_SHUFFLE_BASE(shuffle_xor) 
    1828 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1829 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1830 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    1831 ^ 
    1832 program_source:224:38: note: expanded from macro '\ 
    1833 EXPOSE_FUNCTION_ARGS_2' 
    1834 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    1835 ^ 
    1836 program_source:168:9: note: expanded from macro '\ 
    1837 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1838 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1839 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1840 program_source:296:1: error: illegal string literal in 'asm' 
    1841 DECLARE_SHUFFLE_BASE(shuffle_xor) 
    1842 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1843 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1844 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    1845 ^ 
    1846 program_source:225:51: note: expanded from macro '\ 
    1847 EXPOSE_FUNCTION_ARGS_2' 
    1848 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    1849 ^ 
    1850 program_source:168:9: note: expanded from macro '\ 
    1851 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1852 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1853 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1854 program_source:296:1: error: illegal string literal in 'asm' 
    1855 DECLARE_SHUFFLE_BASE(shuffle_xor) 
    1856 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1857 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1858 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    1859 ^ 
    1860 program_source:226:52: note: expanded from macro '\ 
    1861 EXPOSE_FUNCTION_ARGS_2' 
    1862 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    1863 ^ 
    1864 program_source:168:9: note: expanded from macro '\ 
    1865 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1866 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1867 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1868 program_source:296:1: error: illegal string literal in 'asm' 
    1869 DECLARE_SHUFFLE_BASE(shuffle_xor) 
    1870 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1871 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1872 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    1873 ^ 
    1874 program_source:224:38: note: expanded from macro '\ 
    1875 EXPOSE_FUNCTION_ARGS_2' 
    1876 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    1877 ^ 
    1878 program_source:168:9: note: expanded from macro '\ 
    1879 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1880 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1881 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1882 program_source:296:1: error: illegal string literal in 'asm' 
    1883 DECLARE_SHUFFLE_BASE(shuffle_xor) 
    1884 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1885 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1886 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    1887 ^ 
    1888 program_source:225:51: note: expanded from macro '\ 
    1889 EXPOSE_FUNCTION_ARGS_2' 
    1890 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    1891 ^ 
    1892 program_source:168:9: note: expanded from macro '\ 
    1893 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1894 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1895 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1896 program_source:296:1: error: illegal string literal in 'asm' 
    1897 DECLARE_SHUFFLE_BASE(shuffle_xor) 
    1898 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1899 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1900 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    1901 ^ 
    1902 program_source:226:52: note: expanded from macro '\ 
    1903 EXPOSE_FUNCTION_ARGS_2' 
    1904 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    1905 ^ 
    1906 program_source:168:9: note: expanded from macro '\ 
    1907 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1908 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1909 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1910 program_source:297:1: error: illegal string literal in 'asm' 
    1911 DECLARE_SHUFFLE_BASE(shuffle_up) 
    1912 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1913 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1914 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    1915 ^ 
    1916 program_source:224:38: note: expanded from macro '\ 
    1917 EXPOSE_FUNCTION_ARGS_2' 
    1918 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    1919 ^ 
    1920 program_source:168:9: note: expanded from macro '\ 
    1921 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1922 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1923 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1924 program_source:297:1: error: illegal string literal in 'asm' 
    1925 DECLARE_SHUFFLE_BASE(shuffle_up) 
    1926 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1927 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1928 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    1929 ^ 
    1930 program_source:225:51: note: expanded from macro '\ 
    1931 EXPOSE_FUNCTION_ARGS_2' 
    1932 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    1933 ^ 
    1934 program_source:168:9: note: expanded from macro '\ 
    1935 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1936 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1937 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1938 program_source:297:1: error: illegal string literal in 'asm' 
    1939 DECLARE_SHUFFLE_BASE(shuffle_up) 
    1940 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1941 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1942 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    1943 ^ 
    1944 program_source:226:52: note: expanded from macro '\ 
    1945 EXPOSE_FUNCTION_ARGS_2' 
    1946 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    1947 ^ 
    1948 program_source:168:9: note: expanded from macro '\ 
    1949 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1950 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1951 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1952 program_source:297:1: error: illegal string literal in 'asm' 
    1953 DECLARE_SHUFFLE_BASE(shuffle_up) 
    1954 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1955 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1956 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    1957 ^ 
    1958 program_source:224:38: note: expanded from macro '\ 
    1959 EXPOSE_FUNCTION_ARGS_2' 
    1960 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    1961 ^ 
    1962 program_source:168:9: note: expanded from macro '\ 
    1963 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1964 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1965 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1966 program_source:297:1: error: illegal string literal in 'asm' 
    1967 DECLARE_SHUFFLE_BASE(shuffle_up) 
    1968 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1969 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1970 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    1971 ^ 
    1972 program_source:225:51: note: expanded from macro '\ 
    1973 EXPOSE_FUNCTION_ARGS_2' 
    1974 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    1975 ^ 
    1976 program_source:168:9: note: expanded from macro '\ 
    1977 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1978 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1979 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1980 program_source:297:1: error: illegal string literal in 'asm' 
    1981 DECLARE_SHUFFLE_BASE(shuffle_up) 
    1982 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1983 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1984 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    1985 ^ 
    1986 program_source:226:52: note: expanded from macro '\ 
    1987 EXPOSE_FUNCTION_ARGS_2' 
    1988 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    1989 ^ 
    1990 program_source:168:9: note: expanded from macro '\ 
    1991 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    1992 __asm("air." #EXPR "." #AIR_TYPE); \ 
    1993 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1994 program_source:298:1: error: illegal string literal in 'asm' 
    1995 DECLARE_SHUFFLE_BASE(shuffle_down) 
    1996 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    1997 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    1998 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    1999 ^ 
    2000 program_source:224:38: note: expanded from macro '\ 
    2001 EXPOSE_FUNCTION_ARGS_2' 
    2002 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    2003 ^ 
    2004 program_source:168:9: note: expanded from macro '\ 
    2005 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    2006 __asm("air." #EXPR "." #AIR_TYPE); \ 
    2007 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2008 program_source:298:1: error: illegal string literal in 'asm' 
    2009 DECLARE_SHUFFLE_BASE(shuffle_down) 
    2010 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2011 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    2012 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    2013 ^ 
    2014 program_source:225:51: note: expanded from macro '\ 
    2015 EXPOSE_FUNCTION_ARGS_2' 
    2016 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    2017 ^ 
    2018 program_source:168:9: note: expanded from macro '\ 
    2019 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    2020 __asm("air." #EXPR "." #AIR_TYPE); \ 
    2021 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2022 program_source:298:1: error: illegal string literal in 'asm' 
    2023 DECLARE_SHUFFLE_BASE(shuffle_down) 
    2024 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2025 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    2026 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    2027 ^ 
    2028 program_source:226:52: note: expanded from macro '\ 
    2029 EXPOSE_FUNCTION_ARGS_2' 
    2030 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    2031 ^ 
    2032 program_source:168:9: note: expanded from macro '\ 
    2033 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    2034 __asm("air." #EXPR "." #AIR_TYPE); \ 
    2035 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2036 program_source:298:1: error: illegal string literal in 'asm' 
    2037 DECLARE_SHUFFLE_BASE(shuffle_down) 
    2038 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2039 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    2040 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    2041 ^ 
    2042 program_source:224:38: note: expanded from macro '\ 
    2043 EXPOSE_FUNCTION_ARGS_2' 
    2044 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    2045 ^ 
    2046 program_source:168:9: note: expanded from macro '\ 
    2047 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    2048 __asm("air." #EXPR "." #AIR_TYPE); \ 
    2049 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2050 program_source:298:1: error: illegal string literal in 'asm' 
    2051 DECLARE_SHUFFLE_BASE(shuffle_down) 
    2052 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2053 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    2054 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    2055 ^ 
    2056 program_source:225:51: note: expanded from macro '\ 
    2057 EXPOSE_FUNCTION_ARGS_2' 
    2058 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    2059 ^ 
    2060 program_source:168:9: note: expanded from macro '\ 
    2061 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    2062 __asm("air." #EXPR "." #AIR_TYPE); \ 
    2063 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2064 program_source:298:1: error: illegal string literal in 'asm' 
    2065 DECLARE_SHUFFLE_BASE(shuffle_down) 
    2066 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2067 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    2068 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    2069 ^ 
    2070 program_source:226:52: note: expanded from macro '\ 
    2071 EXPOSE_FUNCTION_ARGS_2' 
    2072 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    2073 ^ 
    2074 program_source:168:9: note: expanded from macro '\ 
    2075 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    2076 __asm("air." #EXPR "." #AIR_TYPE); \ 
    2077 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2078 program_source:299:1: error: illegal string literal in 'asm' 
    2079 DECLARE_SHUFFLE_BASE(shuffle_rotate_up) 
    2080 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2081 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    2082 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    2083 ^ 
    2084 program_source:224:38: note: expanded from macro '\ 
    2085 EXPOSE_FUNCTION_ARGS_2' 
    2086 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    2087 ^ 
    2088 program_source:168:9: note: expanded from macro '\ 
    2089 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    2090 __asm("air." #EXPR "." #AIR_TYPE); \ 
    2091 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2092 program_source:299:1: error: illegal string literal in 'asm' 
    2093 DECLARE_SHUFFLE_BASE(shuffle_rotate_up) 
    2094 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2095 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    2096 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    2097 ^ 
    2098 program_source:225:51: note: expanded from macro '\ 
    2099 EXPOSE_FUNCTION_ARGS_2' 
    2100 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    2101 ^ 
    2102 program_source:168:9: note: expanded from macro '\ 
    2103 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    2104 __asm("air." #EXPR "." #AIR_TYPE); \ 
    2105 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2106 program_source:299:1: error: illegal string literal in 'asm' 
    2107 DECLARE_SHUFFLE_BASE(shuffle_rotate_up) 
    2108 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2109 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    2110 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    2111 ^ 
    2112 program_source:226:52: note: expanded from macro '\ 
    2113 EXPOSE_FUNCTION_ARGS_2' 
    2114 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    2115 ^ 
    2116 program_source:168:9: note: expanded from macro '\ 
    2117 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    2118 __asm("air." #EXPR "." #AIR_TYPE); \ 
    2119 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2120 program_source:299:1: error: illegal string literal in 'asm' 
    2121 DECLARE_SHUFFLE_BASE(shuffle_rotate_up) 
    2122 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2123 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    2124 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    2125 ^ 
    2126 program_source:224:38: note: expanded from macro '\ 
    2127 EXPOSE_FUNCTION_ARGS_2' 
    2128 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    2129 ^ 
    2130 program_source:168:9: note: expanded from macro '\ 
    2131 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    2132 __asm("air." #EXPR "." #AIR_TYPE); \ 
    2133 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2134 program_source:299:1: error: illegal string literal in 'asm' 
    2135 DECLARE_SHUFFLE_BASE(shuffle_rotate_up) 
    2136 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2137 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    2138 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    2139 ^ 
    2140 program_source:225:51: note: expanded from macro '\ 
    2141 EXPOSE_FUNCTION_ARGS_2' 
    2142 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    2143 ^ 
    2144 program_source:168:9: note: expanded from macro '\ 
    2145 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    2146 __asm("air." #EXPR "." #AIR_TYPE); \ 
    2147 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2148 program_source:299:1: error: illegal string literal in 'asm' 
    2149 DECLARE_SHUFFLE_BASE(shuffle_rotate_up) 
    2150 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2151 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    2152 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    2153 ^ 
    2154 program_source:226:52: note: expanded from macro '\ 
    2155 EXPOSE_FUNCTION_ARGS_2' 
    2156 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    2157 ^ 
    2158 program_source:168:9: note: expanded from macro '\ 
    2159 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    2160 __asm("air." #EXPR "." #AIR_TYPE); \ 
    2161 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2162 program_source:300:1: error: illegal string literal in 'asm' 
    2163 DECLARE_SHUFFLE_BASE(shuffle_rotate_down) 
    2164 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2165 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    2166 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    2167 ^ 
    2168 program_source:224:38: note: expanded from macro '\ 
    2169 EXPOSE_FUNCTION_ARGS_2' 
    2170 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    2171 ^ 
    2172 program_source:168:9: note: expanded from macro '\ 
    2173 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    2174 __asm("air." #EXPR "." #AIR_TYPE); \ 
    2175 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2176 program_source:300:1: error: illegal string literal in 'asm' 
    2177 DECLARE_SHUFFLE_BASE(shuffle_rotate_down) 
    2178 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2179 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    2180 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    2181 ^ 
    2182 program_source:225:51: note: expanded from macro '\ 
    2183 EXPOSE_FUNCTION_ARGS_2' 
    2184 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    2185 ^ 
    2186 program_source:168:9: note: expanded from macro '\ 
    2187 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    2188 __asm("air." #EXPR "." #AIR_TYPE); \ 
    2189 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2190 program_source:300:1: error: illegal string literal in 'asm' 
    2191 DECLARE_SHUFFLE_BASE(shuffle_rotate_down) 
    2192 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2193 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    2194 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    2195 ^ 
    2196 program_source:226:52: note: expanded from macro '\ 
    2197 EXPOSE_FUNCTION_ARGS_2' 
    2198 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    2199 ^ 
    2200 program_source:168:9: note: expanded from macro '\ 
    2201 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    2202 __asm("air." #EXPR "." #AIR_TYPE); \ 
    2203 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2204 program_source:300:1: error: illegal string literal in 'asm' 
    2205 DECLARE_SHUFFLE_BASE(shuffle_rotate_down) 
    2206 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2207 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    2208 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    2209 ^ 
    2210 program_source:224:38: note: expanded from macro '\ 
    2211 EXPOSE_FUNCTION_ARGS_2' 
    2212 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    2213 ^ 
    2214 program_source:168:9: note: expanded from macro '\ 
    2215 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    2216 __asm("air." #EXPR "." #AIR_TYPE); \ 
    2217 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2218 program_source:300:1: error: illegal string literal in 'asm' 
    2219 DECLARE_SHUFFLE_BASE(shuffle_rotate_down) 
    2220 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2221 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    2222 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    2223 ^ 
    2224 program_source:225:51: note: expanded from macro '\ 
    2225 EXPOSE_FUNCTION_ARGS_2' 
    2226 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    2227 ^ 
    2228 program_source:168:9: note: expanded from macro '\ 
    2229 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    2230 __asm("air." #EXPR "." #AIR_TYPE); \ 
    2231 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2232 program_source:300:1: error: illegal string literal in 'asm' 
    2233 DECLARE_SHUFFLE_BASE(shuffle_rotate_down) 
    2234 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2235 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    2236 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    2237 ^ 
    2238 program_source:226:52: note: expanded from macro '\ 
    2239 EXPOSE_FUNCTION_ARGS_2' 
    2240 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    2241 ^ 
    2242 program_source:168:9: note: expanded from macro '\ 
    2243 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    2244 __asm("air." #EXPR "." #AIR_TYPE); \ 
    2245 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2246 program_source:302:1: error: illegal string literal in 'asm' 
    2247 EXPOSE_FUNCTION_I_ARGS_3(ctz) 
    2248 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2249 program_source:229:40: note: expanded from macro 'EXPOSE_FUNCTION_I_ARGS_3' 
    2250 #define EXPOSE_FUNCTION_I_ARGS_3(EXPR) \ 
    2251 ^ 
    2252 program_source:172:9: note: expanded from macro '\ 
    2253 EXPOSE_FUNCTION_OVERLOAD_ARGS_3' 
    2254 __asm("air." #EXPR "." #AIR_TYPE); \ 
    2255 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2256 program_source:302:1: error: implicit declaration of function '___metal_ctz'
    2257 is invalid in OpenCL 
    2258 program_source:229:40: note: expanded from macro 'EXPOSE_FUNCTION_I_ARGS_3' 
    2259 #define EXPOSE_FUNCTION_I_ARGS_3(EXPR) \ 
    2260 ^ 
    2261 program_source:175:10: note: expanded from macro '\ 
    2262 EXPOSE_FUNCTION_OVERLOAD_ARGS_3' 
    2263 return ___metal_##EXPR(x, false); \ 
    2264 ^ 
    2265 <scratch space>:12:1: note: expanded from here 
    2266 ___metal_ctz 
    2267 ^ 
    2268 program_source:302:1: error: illegal string literal in 'asm' 
    2269 EXPOSE_FUNCTION_I_ARGS_3(ctz) 
    2270 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2271 program_source:230:49: note: expanded from macro 'EXPOSE_FUNCTION_I_ARGS_3' 
    2272 EXPOSE_FUNCTION_OVERLOAD_ARGS_3(EXPR, int, i32) \ 
    2273 ^ 
    2274 program_source:172:9: note: expanded from macro '\ 
    2275 EXPOSE_FUNCTION_OVERLOAD_ARGS_3' 
    2276 __asm("air." #EXPR "." #AIR_TYPE); \ 
    2277 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    2278 program_source:302:1: error: implicit declaration of function '___metal_ctz'
    2279 is invalid in OpenCL 
    2280 program_source:230:49: note: expanded from macro 'EXPOSE_FUNCTION_I_ARGS_3' 
    2281 EXPOSE_FUNCTION_OVERLOAD_ARGS_3(EXPR, int, i32) \ 
    2282 ^ 
    2283 program_source:175:10: note: expanded from macro '\ 
    2284 EXPOSE_FUNCTION_OVERLOAD_ARGS_3' 
    2285 return ___metal_##EXPR(x, false); \ 
    2286 ^ 
    2287 <scratch space>:16:1: note: expanded from here 
    2288 ___metal_ctz 
    2289 ^ 
    2290 program_source:317:1: error: implicit declaration of function 'simd_sum' is
    2291 invalid in OpenCL 
    2292 DECLARE_REDUCTION_UNIFORM(sum, reduce_add) 
    2293 ^ 
    2294 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2295 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2296 ^ 
    2297 program_source:305:26: note: expanded from macro '\ 
    2298 DECLARE_I_REDUCTION_UNIFORM' 
    2299 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2300 ^ 
    2301 <scratch space>:17:1: note: expanded from here 
    2302 simd_sum 
    2303 ^ 
    2304 program_source:317:1: error: implicit declaration of function 'simd_sum' is
    2305 invalid in OpenCL 
    2306 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2307 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2308 ^ 
    2309 program_source:305:26: note: expanded from macro '\ 
    2310 DECLARE_I_REDUCTION_UNIFORM' 
    2311 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2312 ^ 
    2313 <scratch space>:17:1: note: expanded from here 
    2314 simd_sum 
    2315 ^ 
    2316 program_source:317:1: error: implicit declaration of function 'simd_sum' is
    2317 invalid in OpenCL 
    2318 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2319 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2320 ^ 
    2321 program_source:308:26: note: expanded from macro '\ 
    2322 DECLARE_F_REDUCTION_UNIFORM' 
    2323 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2324 ^ 
    2325 <scratch space>:19:1: note: expanded from here 
    2326 simd_sum 
    2327 ^ 
    2328 program_source:318:1: error: implicit declaration of function
    2329 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    2330 DECLARE_REDUCTION_UNIFORM(prefix_inclusive_sum, scan_inclusive_add) 
    2331 ^ 
    2332 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2333 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2334 ^ 
    2335 program_source:305:26: note: expanded from macro '\ 
    2336 DECLARE_I_REDUCTION_UNIFORM' 
    2337 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2338 ^ 
    2339 <scratch space>:21:1: note: expanded from here 
    2340 simd_prefix_inclusive_sum 
    2341 ^ 
    2342 program_source:318:1: error: implicit declaration of function
    2343 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    2344 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2345 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2346 ^ 
    2347 program_source:305:26: note: expanded from macro '\ 
    2348 DECLARE_I_REDUCTION_UNIFORM' 
    2349 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2350 ^ 
    2351 <scratch space>:21:1: note: expanded from here 
    2352 simd_prefix_inclusive_sum 
    2353 ^ 
    2354 program_source:318:1: error: implicit declaration of function
    2355 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    2356 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2357 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2358 ^ 
    2359 program_source:308:26: note: expanded from macro '\ 
    2360 DECLARE_F_REDUCTION_UNIFORM' 
    2361 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2362 ^ 
    2363 <scratch space>:23:1: note: expanded from here 
    2364 simd_prefix_inclusive_sum 
    2365 ^ 
    2366 program_source:319:1: error: implicit declaration of function
    2367 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    2368 DECLARE_REDUCTION_UNIFORM(prefix_exclusive_sum, scan_exclusive_add) 
    2369 ^ 
    2370 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2371 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2372 ^ 
    2373 program_source:305:26: note: expanded from macro '\ 
    2374 DECLARE_I_REDUCTION_UNIFORM' 
    2375 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2376 ^ 
    2377 <scratch space>:25:1: note: expanded from here 
    2378 simd_prefix_exclusive_sum 
    2379 ^ 
    2380 program_source:319:1: error: implicit declaration of function
    2381 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    2382 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2383 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2384 ^ 
    2385 program_source:305:26: note: expanded from macro '\ 
    2386 DECLARE_I_REDUCTION_UNIFORM' 
    2387 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2388 ^ 
    2389 <scratch space>:25:1: note: expanded from here 
    2390 simd_prefix_exclusive_sum 
    2391 ^ 
    2392 program_source:319:1: error: implicit declaration of function
    2393 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    2394 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2395 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2396 ^ 
    2397 program_source:308:26: note: expanded from macro '\ 
    2398 DECLARE_F_REDUCTION_UNIFORM' 
    2399 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2400 ^ 
    2401 <scratch space>:27:1: note: expanded from here 
    2402 simd_prefix_exclusive_sum 
    2403 ^ 
    2404 program_source:320:1: error: implicit declaration of function 'simd_min' is
    2405 invalid in OpenCL 
    2406 DECLARE_REDUCTION_UNIFORM(min, reduce_min) 
    2407 ^ 
    2408 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2409 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2410 ^ 
    2411 program_source:305:26: note: expanded from macro '\ 
    2412 DECLARE_I_REDUCTION_UNIFORM' 
    2413 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2414 ^ 
    2415 <scratch space>:29:1: note: expanded from here 
    2416 simd_min 
    2417 ^ 
    2418 program_source:320:1: error: implicit declaration of function 'simd_min' is
    2419 invalid in OpenCL 
    2420 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2421 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2422 ^ 
    2423 program_source:305:26: note: expanded from macro '\ 
    2424 DECLARE_I_REDUCTION_UNIFORM' 
    2425 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2426 ^ 
    2427 <scratch space>:29:1: note: expanded from here 
    2428 simd_min 
    2429 ^ 
    2430 program_source:320:1: error: implicit declaration of function 'simd_min' is
    2431 invalid in OpenCL 
    2432 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2433 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2434 ^ 
    2435 program_source:308:26: note: expanded from macro '\ 
    2436 DECLARE_F_REDUCTION_UNIFORM' 
    2437 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2438 ^ 
    2439 <scratch space>:31:1: note: expanded from here 
    2440 simd_min 
    2441 ^ 
    2442 program_source:321:1: error: implicit declaration of function 'simd_max' is
    2443 invalid in OpenCL 
    2444 DECLARE_REDUCTION_UNIFORM(max, reduce_max) 
    2445 ^ 
    2446 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2447 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2448 ^ 
    2449 program_source:305:26: note: expanded from macro '\ 
    2450 DECLARE_I_REDUCTION_UNIFORM' 
    2451 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2452 ^ 
    2453 <scratch space>:33:1: note: expanded from here 
    2454 simd_max 
    2455 ^ 
    2456 program_source:321:1: error: implicit declaration of function 'simd_max' is
    2457 invalid in OpenCL 
    2458 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2459 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2460 ^ 
    2461 program_source:305:26: note: expanded from macro '\ 
    2462 DECLARE_I_REDUCTION_UNIFORM' 
    2463 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2464 ^ 
    2465 <scratch space>:33:1: note: expanded from here 
    2466 simd_max 
    2467 ^ 
    2468 program_source:321:1: error: implicit declaration of function 'simd_max' is
    2469 invalid in OpenCL 
    2470 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2471 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2472 ^ 
    2473 program_source:308:26: note: expanded from macro '\ 
    2474 DECLARE_F_REDUCTION_UNIFORM' 
    2475 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2476 ^ 
    2477 <scratch space>:35:1: note: expanded from here 
    2478 simd_max 
    2479 ^ 
    2480 program_source:323:1: error: implicit declaration of function 'simd_shuffle'
    2481 is invalid in OpenCL 
    2482 DECLARE_SHUFFLE_UNIFORM(shuffle, shuffle) 
    2483 ^ 
    2484 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2485 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2486 ^ 
    2487 <scratch space>:37:1: note: expanded from here 
    2488 simd_shuffle 
    2489 ^ 
    2490 program_source:323:1: error: implicit declaration of function 'simd_shuffle'
    2491 is invalid in OpenCL 
    2492 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2493 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2494 ^ 
    2495 <scratch space>:37:1: note: expanded from here 
    2496 simd_shuffle 
    2497 ^ 
    2498 program_source:323:1: error: implicit declaration of function 'simd_shuffle'
    2499 is invalid in OpenCL 
    2500 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2501 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2502 ^ 
    2503 <scratch space>:37:1: note: expanded from here 
    2504 simd_shuffle 
    2505 ^ 
    2506 program_source:324:1: error: implicit declaration of function
    2507 'simd_shuffle_xor' is invalid in OpenCL 
    2508 DECLARE_SHUFFLE_UNIFORM(shuffle_xor, shuffle_xor) 
    2509 ^ 
    2510 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2511 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2512 ^ 
    2513 <scratch space>:39:1: note: expanded from here 
    2514 simd_shuffle_xor 
    2515 ^ 
    2516 program_source:324:1: error: implicit declaration of function
    2517 'simd_shuffle_xor' is invalid in OpenCL 
    2518 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2519 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2520 ^ 
    2521 <scratch space>:39:1: note: expanded from here 
    2522 simd_shuffle_xor 
    2523 ^ 
    2524 program_source:324:1: error: implicit declaration of function
    2525 'simd_shuffle_xor' is invalid in OpenCL 
    2526 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2527 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2528 ^ 
    2529 <scratch space>:39:1: note: expanded from here 
    2530 simd_shuffle_xor 
    2531 ^ 
    2532 program_source:325:1: error: implicit declaration of function
    2533 'simd_shuffle_up' is invalid in OpenCL 
    2534 DECLARE_SHUFFLE_UNIFORM(shuffle_up, shuffle_up) 
    2535 ^ 
    2536 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2537 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2538 ^ 
    2539 <scratch space>:41:1: note: expanded from here 
    2540 simd_shuffle_up 
    2541 ^ 
    2542 program_source:325:1: error: implicit declaration of function
    2543 'simd_shuffle_up' is invalid in OpenCL 
    2544 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2545 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2546 ^ 
    2547 <scratch space>:41:1: note: expanded from here 
    2548 simd_shuffle_up 
    2549 ^ 
    2550 program_source:325:1: error: implicit declaration of function
    2551 'simd_shuffle_up' is invalid in OpenCL 
    2552 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2553 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2554 ^ 
    2555 <scratch space>:41:1: note: expanded from here 
    2556 simd_shuffle_up 
    2557 ^ 
    2558 program_source:326:1: error: implicit declaration of function
    2559 'simd_shuffle_down' is invalid in OpenCL 
    2560 DECLARE_SHUFFLE_UNIFORM(shuffle_down, shuffle_down) 
    2561 ^ 
    2562 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2563 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2564 ^ 
    2565 <scratch space>:43:1: note: expanded from here 
    2566 simd_shuffle_down 
    2567 ^ 
    2568 program_source:326:1: error: implicit declaration of function
    2569 'simd_shuffle_down' is invalid in OpenCL 
    2570 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2571 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2572 ^ 
    2573 <scratch space>:43:1: note: expanded from here 
    2574 simd_shuffle_down 
    2575 ^ 
    2576 program_source:326:1: error: implicit declaration of function
    2577 'simd_shuffle_down' is invalid in OpenCL 
    2578 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2579 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2580 ^ 
    2581 <scratch space>:43:1: note: expanded from here 
    2582 simd_shuffle_down 
    2583 ^ 
    2584 program_source:327:1: error: implicit declaration of function
    2585 'simd_shuffle_rotate_down' is invalid in OpenCL 
    2586 DECLARE_SHUFFLE_UNIFORM(shuffle_rotate_down, rotate) 
    2587 ^ 
    2588 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2589 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2590 ^ 
    2591 <scratch space>:45:1: note: expanded from here 
    2592 simd_shuffle_rotate_down 
    2593 ^ 
    2594 program_source:327:1: error: implicit declaration of function
    2595 'simd_shuffle_rotate_down' is invalid in OpenCL 
    2596 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2597 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2598 ^ 
    2599 <scratch space>:45:1: note: expanded from here 
    2600 simd_shuffle_rotate_down 
    2601 ^ 
    2602 program_source:327:1: error: implicit declaration of function
    2603 'simd_shuffle_rotate_down' is invalid in OpenCL 
    2604 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2605 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2606 ^ 
    2607 <scratch space>:45:1: note: expanded from here 
    2608 simd_shuffle_rotate_down 
    2609 ^ 
    2610 program_source:329:1: error: implicit declaration of function 'simd_broadcast'
    2611 is invalid in OpenCL 
    2612 DECLARE_SHUFFLE_UNIFORM(broadcast, broadcast) 
    2613 ^ 
    2614 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2615 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2616 ^ 
    2617 <scratch space>:47:1: note: expanded from here 
    2618 simd_broadcast 
    2619 ^ 
    2620 program_source:329:1: error: implicit declaration of function 'simd_broadcast'
    2621 is invalid in OpenCL 
    2622 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2623 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2624 ^ 
    2625 <scratch space>:47:1: note: expanded from here 
    2626 simd_broadcast 
    2627 ^ 
    2628 program_source:329:1: error: implicit declaration of function 'simd_broadcast'
    2629 is invalid in OpenCL 
    2630 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    2631 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2632 ^ 
    2633 <scratch space>:47:1: note: expanded from here 
    2634 simd_broadcast 
    2635 ^ 
    2636 program_source:330:1: error: implicit declaration of function
    2637 'simd_broadcast_first' is invalid in OpenCL 
    2638 DECLARE_REDUCTION_UNIFORM(broadcast_first, broadcast_first) 
    2639 ^ 
    2640 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2641 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2642 ^ 
    2643 program_source:305:26: note: expanded from macro '\ 
    2644 DECLARE_I_REDUCTION_UNIFORM' 
    2645 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2646 ^ 
    2647 <scratch space>:49:1: note: expanded from here 
    2648 simd_broadcast_first 
    2649 ^ 
    2650 program_source:330:1: error: implicit declaration of function
    2651 'simd_broadcast_first' is invalid in OpenCL 
    2652 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2653 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2654 ^ 
    2655 program_source:305:26: note: expanded from macro '\ 
    2656 DECLARE_I_REDUCTION_UNIFORM' 
    2657 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2658 ^ 
    2659 <scratch space>:49:1: note: expanded from here 
    2660 simd_broadcast_first 
    2661 ^ 
    2662 program_source:330:1: error: implicit declaration of function
    2663 'simd_broadcast_first' is invalid in OpenCL 
    2664 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    2665 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2666 ^ 
    2667 program_source:308:26: note: expanded from macro '\ 
    2668 DECLARE_F_REDUCTION_UNIFORM' 
    2669 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    2670 ^ 
    2671 <scratch space>:51:1: note: expanded from here 
    2672 simd_broadcast_first 
    2673 ^ 
    2674 program_source:345:1: error: implicit declaration of function 'simd_sum' is
    2675 invalid in OpenCL 
    2676 DECLARE_REDUCTION_NON_UNIFORM(sum, reduce_add) 
    2677 ^ 
    2678 program_source:338:60: note: expanded from macro
    2679 'DECLARE_REDUCTION_NON_UNIFORM' 
    2680 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2681 ^ 
    2682 program_source:333:26: note: expanded from macro '\ 
    2683 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2684 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2685 ^ 
    2686 <scratch space>:53:1: note: expanded from here 
    2687 simd_sum 
    2688 ^ 
    2689 program_source:345:1: error: implicit declaration of function 'simd_sum' is
    2690 invalid in OpenCL 
    2691 program_source:338:60: note: expanded from macro
    2692 'DECLARE_REDUCTION_NON_UNIFORM' 
    2693 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2694 ^ 
    2695 program_source:333:26: note: expanded from macro '\ 
    2696 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2697 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2698 ^ 
    2699 <scratch space>:53:1: note: expanded from here 
    2700 simd_sum 
    2701 ^ 
    2702 program_source:345:1: error: implicit declaration of function 'simd_sum' is
    2703 invalid in OpenCL 
    2704 program_source:339:54: note: expanded from macro
    2705 'DECLARE_REDUCTION_NON_UNIFORM' 
    2706 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2707 ^ 
    2708 program_source:336:26: note: expanded from macro '\ 
    2709 DECLARE_F_REDUCTION_NON_UNIFORM' 
    2710 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2711 ^ 
    2712 <scratch space>:55:1: note: expanded from here 
    2713 simd_sum 
    2714 ^ 
    2715 program_source:346:1: error: implicit declaration of function
    2716 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    2717 DECLARE_REDUCTION_NON_UNIFORM(prefix_inclusive_sum, scan_inclusive_add) 
    2718 ^ 
    2719 program_source:338:60: note: expanded from macro
    2720 'DECLARE_REDUCTION_NON_UNIFORM' 
    2721 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2722 ^ 
    2723 program_source:333:26: note: expanded from macro '\ 
    2724 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2725 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2726 ^ 
    2727 <scratch space>:57:1: note: expanded from here 
    2728 simd_prefix_inclusive_sum 
    2729 ^ 
    2730 program_source:346:1: error: implicit declaration of function
    2731 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    2732 program_source:338:60: note: expanded from macro
    2733 'DECLARE_REDUCTION_NON_UNIFORM' 
    2734 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2735 ^ 
    2736 program_source:333:26: note: expanded from macro '\ 
    2737 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2738 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2739 ^ 
    2740 <scratch space>:57:1: note: expanded from here 
    2741 simd_prefix_inclusive_sum 
    2742 ^ 
    2743 program_source:346:1: error: implicit declaration of function
    2744 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    2745 program_source:339:54: note: expanded from macro
    2746 'DECLARE_REDUCTION_NON_UNIFORM' 
    2747 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2748 ^ 
    2749 program_source:336:26: note: expanded from macro '\ 
    2750 DECLARE_F_REDUCTION_NON_UNIFORM' 
    2751 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2752 ^ 
    2753 <scratch space>:59:1: note: expanded from here 
    2754 simd_prefix_inclusive_sum 
    2755 ^ 
    2756 program_source:347:1: error: implicit declaration of function
    2757 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    2758 DECLARE_REDUCTION_NON_UNIFORM(prefix_exclusive_sum, scan_exclusive_add) 
    2759 ^ 
    2760 program_source:338:60: note: expanded from macro
    2761 'DECLARE_REDUCTION_NON_UNIFORM' 
    2762 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2763 ^ 
    2764 program_source:333:26: note: expanded from macro '\ 
    2765 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2766 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2767 ^ 
    2768 <scratch space>:61:1: note: expanded from here 
    2769 simd_prefix_exclusive_sum 
    2770 ^ 
    2771 program_source:347:1: error: implicit declaration of function
    2772 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    2773 program_source:338:60: note: expanded from macro
    2774 'DECLARE_REDUCTION_NON_UNIFORM' 
    2775 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2776 ^ 
    2777 program_source:333:26: note: expanded from macro '\ 
    2778 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2779 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2780 ^ 
    2781 <scratch space>:61:1: note: expanded from here 
    2782 simd_prefix_exclusive_sum 
    2783 ^ 
    2784 program_source:347:1: error: implicit declaration of function
    2785 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    2786 program_source:339:54: note: expanded from macro
    2787 'DECLARE_REDUCTION_NON_UNIFORM' 
    2788 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2789 ^ 
    2790 program_source:336:26: note: expanded from macro '\ 
    2791 DECLARE_F_REDUCTION_NON_UNIFORM' 
    2792 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2793 ^ 
    2794 <scratch space>:63:1: note: expanded from here 
    2795 simd_prefix_exclusive_sum 
    2796 ^ 
    2797 program_source:348:1: error: implicit declaration of function 'simd_min' is
    2798 invalid in OpenCL 
    2799 DECLARE_REDUCTION_NON_UNIFORM(min, reduce_min) 
    2800 ^ 
    2801 program_source:338:60: note: expanded from macro
    2802 'DECLARE_REDUCTION_NON_UNIFORM' 
    2803 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2804 ^ 
    2805 program_source:333:26: note: expanded from macro '\ 
    2806 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2807 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2808 ^ 
    2809 <scratch space>:65:1: note: expanded from here 
    2810 simd_min 
    2811 ^ 
    2812 program_source:348:1: error: implicit declaration of function 'simd_min' is
    2813 invalid in OpenCL 
    2814 program_source:338:60: note: expanded from macro
    2815 'DECLARE_REDUCTION_NON_UNIFORM' 
    2816 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2817 ^ 
    2818 program_source:333:26: note: expanded from macro '\ 
    2819 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2820 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2821 ^ 
    2822 <scratch space>:65:1: note: expanded from here 
    2823 simd_min 
    2824 ^ 
    2825 program_source:348:1: error: implicit declaration of function 'simd_min' is
    2826 invalid in OpenCL 
    2827 program_source:339:54: note: expanded from macro
    2828 'DECLARE_REDUCTION_NON_UNIFORM' 
    2829 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2830 ^ 
    2831 program_source:336:26: note: expanded from macro '\ 
    2832 DECLARE_F_REDUCTION_NON_UNIFORM' 
    2833 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2834 ^ 
    2835 <scratch space>:67:1: note: expanded from here 
    2836 simd_min 
    2837 ^ 
    2838 program_source:349:1: error: implicit declaration of function 'simd_max' is
    2839 invalid in OpenCL 
    2840 DECLARE_REDUCTION_NON_UNIFORM(max, reduce_max) 
    2841 ^ 
    2842 program_source:338:60: note: expanded from macro
    2843 'DECLARE_REDUCTION_NON_UNIFORM' 
    2844 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2845 ^ 
    2846 program_source:333:26: note: expanded from macro '\ 
    2847 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2848 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2849 ^ 
    2850 <scratch space>:69:1: note: expanded from here 
    2851 simd_max 
    2852 ^ 
    2853 program_source:349:1: error: implicit declaration of function 'simd_max' is
    2854 invalid in OpenCL 
    2855 program_source:338:60: note: expanded from macro
    2856 'DECLARE_REDUCTION_NON_UNIFORM' 
    2857 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2858 ^ 
    2859 program_source:333:26: note: expanded from macro '\ 
    2860 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2861 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2862 ^ 
    2863 <scratch space>:69:1: note: expanded from here 
    2864 simd_max 
    2865 ^ 
    2866 program_source:349:1: error: implicit declaration of function 'simd_max' is
    2867 invalid in OpenCL 
    2868 program_source:339:54: note: expanded from macro
    2869 'DECLARE_REDUCTION_NON_UNIFORM' 
    2870 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2871 ^ 
    2872 program_source:336:26: note: expanded from macro '\ 
    2873 DECLARE_F_REDUCTION_NON_UNIFORM' 
    2874 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2875 ^ 
    2876 <scratch space>:71:1: note: expanded from here 
    2877 simd_max 
    2878 ^ 
    2879 program_source:351:1: error: implicit declaration of function 'simd_product'
    2880 is invalid in OpenCL 
    2881 DECLARE_REDUCTION_NON_UNIFORM(product, reduce_mul) 
    2882 ^ 
    2883 program_source:338:60: note: expanded from macro
    2884 'DECLARE_REDUCTION_NON_UNIFORM' 
    2885 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2886 ^ 
    2887 program_source:333:26: note: expanded from macro '\ 
    2888 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2889 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2890 ^ 
    2891 <scratch space>:73:1: note: expanded from here 
    2892 simd_product 
    2893 ^ 
    2894 program_source:351:1: error: implicit declaration of function 'simd_product'
    2895 is invalid in OpenCL 
    2896 program_source:338:60: note: expanded from macro
    2897 'DECLARE_REDUCTION_NON_UNIFORM' 
    2898 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2899 ^ 
    2900 program_source:333:26: note: expanded from macro '\ 
    2901 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2902 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2903 ^ 
    2904 <scratch space>:73:1: note: expanded from here 
    2905 simd_product 
    2906 ^ 
    2907 program_source:351:1: error: implicit declaration of function 'simd_product'
    2908 is invalid in OpenCL 
    2909 program_source:339:54: note: expanded from macro
    2910 'DECLARE_REDUCTION_NON_UNIFORM' 
    2911 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2912 ^ 
    2913 program_source:336:26: note: expanded from macro '\ 
    2914 DECLARE_F_REDUCTION_NON_UNIFORM' 
    2915 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2916 ^ 
    2917 <scratch space>:75:1: note: expanded from here 
    2918 simd_product 
    2919 ^ 
    2920 program_source:352:1: error: implicit declaration of function
    2921 'simd_prefix_inclusive_product' is invalid in OpenCL 
    2922 DECLARE_REDUCTION_NON_UNIFORM(prefix_inclusive_product, scan_inclusive_mul) 
    2923 ^ 
    2924 program_source:338:60: note: expanded from macro
    2925 'DECLARE_REDUCTION_NON_UNIFORM' 
    2926 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2927 ^ 
    2928 program_source:333:26: note: expanded from macro '\ 
    2929 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2930 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2931 ^ 
    2932 <scratch space>:77:1: note: expanded from here 
    2933 simd_prefix_inclusive_product 
    2934 ^ 
    2935 program_source:352:1: error: implicit declaration of function
    2936 'simd_prefix_inclusive_product' is invalid in OpenCL 
    2937 program_source:338:60: note: expanded from macro
    2938 'DECLARE_REDUCTION_NON_UNIFORM' 
    2939 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2940 ^ 
    2941 program_source:333:26: note: expanded from macro '\ 
    2942 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2943 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2944 ^ 
    2945 <scratch space>:77:1: note: expanded from here 
    2946 simd_prefix_inclusive_product 
    2947 ^ 
    2948 program_source:352:1: error: implicit declaration of function
    2949 'simd_prefix_inclusive_product' is invalid in OpenCL 
    2950 program_source:339:54: note: expanded from macro
    2951 'DECLARE_REDUCTION_NON_UNIFORM' 
    2952 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2953 ^ 
    2954 program_source:336:26: note: expanded from macro '\ 
    2955 DECLARE_F_REDUCTION_NON_UNIFORM' 
    2956 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2957 ^ 
    2958 <scratch space>:79:1: note: expanded from here 
    2959 simd_prefix_inclusive_product 
    2960 ^ 
    2961 program_source:353:1: error: implicit declaration of function
    2962 'simd_prefix_exclusive_product' is invalid in OpenCL 
    2963 DECLARE_REDUCTION_NON_UNIFORM(prefix_exclusive_product, scan_exclusive_mul) 
    2964 ^ 
    2965 program_source:338:60: note: expanded from macro
    2966 'DECLARE_REDUCTION_NON_UNIFORM' 
    2967 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2968 ^ 
    2969 program_source:333:26: note: expanded from macro '\ 
    2970 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2971 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2972 ^ 
    2973 <scratch space>:81:1: note: expanded from here 
    2974 simd_prefix_exclusive_product 
    2975 ^ 
    2976 program_source:353:1: error: implicit declaration of function
    2977 'simd_prefix_exclusive_product' is invalid in OpenCL 
    2978 program_source:338:60: note: expanded from macro
    2979 'DECLARE_REDUCTION_NON_UNIFORM' 
    2980 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2981 ^ 
    2982 program_source:333:26: note: expanded from macro '\ 
    2983 DECLARE_I_REDUCTION_NON_UNIFORM' 
    2984 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2985 ^ 
    2986 <scratch space>:81:1: note: expanded from here 
    2987 simd_prefix_exclusive_product 
    2988 ^ 
    2989 program_source:353:1: error: implicit declaration of function
    2990 'simd_prefix_exclusive_product' is invalid in OpenCL 
    2991 program_source:339:54: note: expanded from macro
    2992 'DECLARE_REDUCTION_NON_UNIFORM' 
    2993 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    2994 ^ 
    2995 program_source:336:26: note: expanded from macro '\ 
    2996 DECLARE_F_REDUCTION_NON_UNIFORM' 
    2997 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    2998 ^ 
    2999 <scratch space>:83:1: note: expanded from here 
    3000 simd_prefix_exclusive_product 
    3001 ^ 
    3002 program_source:354:1: error: implicit declaration of function 'simd_and' is
    3003 invalid in OpenCL 
    3004 DECLARE_I_REDUCTION_NON_UNIFORM(and, reduce_and) 
    3005 ^ 
    3006 program_source:333:26: note: expanded from macro
    3007 'DECLARE_I_REDUCTION_NON_UNIFORM' 
    3008 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    3009 ^ 
    3010 <scratch space>:85:1: note: expanded from here 
    3011 simd_and 
    3012 ^ 
    3013 program_source:354:1: error: implicit declaration of function 'simd_and' is
    3014 invalid in OpenCL 
    3015 program_source:333:26: note: expanded from macro
    3016 'DECLARE_I_REDUCTION_NON_UNIFORM' 
    3017 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    3018 ^ 
    3019 <scratch space>:85:1: note: expanded from here 
    3020 simd_and 
    3021 ^ 
    3022 program_source:355:1: error: implicit declaration of function 'simd_or' is
    3023 invalid in OpenCL 
    3024 DECLARE_I_REDUCTION_NON_UNIFORM(or, reduce_or) 
    3025 ^ 
    3026 program_source:333:26: note: expanded from macro
    3027 'DECLARE_I_REDUCTION_NON_UNIFORM' 
    3028 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    3029 ^ 
    3030 <scratch space>:87:1: note: expanded from here 
    3031 simd_or 
    3032 ^ 
    3033 program_source:355:1: error: implicit declaration of function 'simd_or' is
    3034 invalid in OpenCL 
    3035 program_source:333:26: note: expanded from macro
    3036 'DECLARE_I_REDUCTION_NON_UNIFORM' 
    3037 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    3038 ^ 
    3039 <scratch space>:87:1: note: expanded from here 
    3040 simd_or 
    3041 ^ 
    3042 program_source:356:1: error: implicit declaration of function 'simd_xor' is
    3043 invalid in OpenCL 
    3044 DECLARE_I_REDUCTION_NON_UNIFORM(xor, reduce_xor) 
    3045 ^ 
    3046 program_source:333:26: note: expanded from macro
    3047 'DECLARE_I_REDUCTION_NON_UNIFORM' 
    3048 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    3049 ^ 
    3050 <scratch space>:89:1: note: expanded from here 
    3051 simd_xor 
    3052 ^ 
    3053 program_source:356:1: error: implicit declaration of function 'simd_xor' is
    3054 invalid in OpenCL 
    3055 program_source:333:26: note: expanded from macro
    3056 'DECLARE_I_REDUCTION_NON_UNIFORM' 
    3057 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    3058 ^ 
    3059 <scratch space>:89:1: note: expanded from here 
    3060 simd_xor 
    3061 ^ 
    3062 program_source:358:1: error: implicit declaration of function 'simd_broadcast'
    3063 is invalid in OpenCL 
    3064 DECLARE_SHUFFLE_NON_UNIFORM(broadcast, broadcast) 
    3065 ^ 
    3066 program_source:343:24: note: expanded from macro 'DECLARE_SHUFFLE_NON_UNIFORM' 
    3067 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    3068 ^ 
    3069 <scratch space>:91:1: note: expanded from here 
    3070 simd_broadcast 
    3071 ^ 
    3072 program_source:358:1: error: implicit declaration of function 'simd_broadcast'
    3073 is invalid in OpenCL 
    3074 program_source:343:24: note: expanded from macro 'DECLARE_SHUFFLE_NON_UNIFORM' 
    3075 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    3076 ^ 
    3077 <scratch space>:91:1: note: expanded from here 
    3078 simd_broadcast 
    3079 ^ 
    3080 program_source:358:1: error: implicit declaration of function 'simd_broadcast'
    3081 is invalid in OpenCL 
    3082 program_source:343:24: note: expanded from macro 'DECLARE_SHUFFLE_NON_UNIFORM' 
    3083 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    3084 ^ 
    3085 <scratch space>:91:1: note: expanded from here 
    3086 simd_broadcast 
    3087 ^ 
    3088 program_source:391:1: error: implicit declaration of function 'quad_sum' is
    3089 invalid in OpenCL 
    3090 DECLARE_REDUCTION_CLUSTERED(sum, reduce_add) 
    3091 ^ 
    3092 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3093 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3094 ^ 
    3095 program_source:360:60: note: expanded from macro '\ 
    3096 DECLARE_I_REDUCTION_CLUSTERED' 
    3097 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3098 ^ 
    3099 program_source:245:71: note: expanded from macro '\ 
    3100 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3101 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3102 ^ 
    3103 program_source:196:12: note: expanded from macro '\ 
    3104 OVERLOAD_CLUSTERED_ARGS_1' 
    3105 return quad_##METAL_SUFFIX(x); \ 
    3106 ^ 
    3107 <scratch space>:94:1: note: expanded from here 
    3108 quad_sum 
    3109 ^ 
    3110 program_source:391:1: error: implicit declaration of function 'simd_sum' is
    3111 invalid in OpenCL 
    3112 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3113 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3114 ^ 
    3115 program_source:360:60: note: expanded from macro '\ 
    3116 DECLARE_I_REDUCTION_CLUSTERED' 
    3117 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3118 ^ 
    3119 program_source:245:71: note: expanded from macro '\ 
    3120 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3121 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3122 ^ 
    3123 program_source:198:12: note: expanded from macro '\ 
    3124 OVERLOAD_CLUSTERED_ARGS_1' 
    3125 return simd_##METAL_SUFFIX(x); \ 
    3126 ^ 
    3127 <scratch space>:95:1: note: expanded from here 
    3128 simd_sum 
    3129 ^ 
    3130 program_source:391:1: error: implicit declaration of function 'quad_sum' is
    3131 invalid in OpenCL 
    3132 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3133 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3134 ^ 
    3135 program_source:360:60: note: expanded from macro '\ 
    3136 DECLARE_I_REDUCTION_CLUSTERED' 
    3137 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3138 ^ 
    3139 program_source:246:59: note: expanded from macro '\ 
    3140 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3141 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3142 ^ 
    3143 program_source:196:12: note: expanded from macro '\ 
    3144 OVERLOAD_CLUSTERED_ARGS_1' 
    3145 return quad_##METAL_SUFFIX(x); \ 
    3146 ^ 
    3147 <scratch space>:96:1: note: expanded from here 
    3148 quad_sum 
    3149 ^ 
    3150 program_source:391:1: error: implicit declaration of function 'simd_sum' is
    3151 invalid in OpenCL 
    3152 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3153 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3154 ^ 
    3155 program_source:360:60: note: expanded from macro '\ 
    3156 DECLARE_I_REDUCTION_CLUSTERED' 
    3157 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3158 ^ 
    3159 program_source:246:59: note: expanded from macro '\ 
    3160 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3161 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3162 ^ 
    3163 program_source:198:12: note: expanded from macro '\ 
    3164 OVERLOAD_CLUSTERED_ARGS_1' 
    3165 return simd_##METAL_SUFFIX(x); \ 
    3166 ^ 
    3167 <scratch space>:97:1: note: expanded from here 
    3168 simd_sum 
    3169 ^ 
    3170 program_source:391:1: error: implicit declaration of function 'quad_sum' is
    3171 invalid in OpenCL 
    3172 program_source:385:52: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3173 DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3174 ^ 
    3175 program_source:363:60: note: expanded from macro '\ 
    3176 DECLARE_F_REDUCTION_CLUSTERED' 
    3177 #define DECLARE_F_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3178 ^ 
    3179 program_source:249:71: note: expanded from macro '\ 
    3180 BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1' 
    3181 #define BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3182 ^ 
    3183 program_source:196:12: note: expanded from macro '\ 
    3184 OVERLOAD_CLUSTERED_ARGS_1' 
    3185 return quad_##METAL_SUFFIX(x); \ 
    3186 ^ 
    3187 <scratch space>:99:1: note: expanded from here 
    3188 quad_sum 
    3189 ^ 
    3190 program_source:391:1: error: implicit declaration of function 'simd_sum' is
    3191 invalid in OpenCL 
    3192 program_source:385:52: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3193 DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3194 ^ 
    3195 program_source:363:60: note: expanded from macro '\ 
    3196 DECLARE_F_REDUCTION_CLUSTERED' 
    3197 #define DECLARE_F_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3198 ^ 
    3199 program_source:249:71: note: expanded from macro '\ 
    3200 BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1' 
    3201 #define BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3202 ^ 
    3203 program_source:198:12: note: expanded from macro '\ 
    3204 OVERLOAD_CLUSTERED_ARGS_1' 
    3205 return simd_##METAL_SUFFIX(x); \ 
    3206 ^ 
    3207 <scratch space>:100:1: note: expanded from here 
    3208 simd_sum 
    3209 ^ 
    3210 program_source:392:1: error: implicit declaration of function 'quad_min' is
    3211 invalid in OpenCL 
    3212 DECLARE_REDUCTION_CLUSTERED(min, reduce_min) 
    3213 ^ 
    3214 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3215 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3216 ^ 
    3217 program_source:360:60: note: expanded from macro '\ 
    3218 DECLARE_I_REDUCTION_CLUSTERED' 
    3219 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3220 ^ 
    3221 program_source:245:71: note: expanded from macro '\ 
    3222 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3223 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3224 ^ 
    3225 program_source:196:12: note: expanded from macro '\ 
    3226 OVERLOAD_CLUSTERED_ARGS_1' 
    3227 return quad_##METAL_SUFFIX(x); \ 
    3228 ^ 
    3229 <scratch space>:102:1: note: expanded from here 
    3230 quad_min 
    3231 ^ 
    3232 program_source:392:1: error: implicit declaration of function 'simd_min' is
    3233 invalid in OpenCL 
    3234 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3235 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3236 ^ 
    3237 program_source:360:60: note: expanded from macro '\ 
    3238 DECLARE_I_REDUCTION_CLUSTERED' 
    3239 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3240 ^ 
    3241 program_source:245:71: note: expanded from macro '\ 
    3242 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3243 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3244 ^ 
    3245 program_source:198:12: note: expanded from macro '\ 
    3246 OVERLOAD_CLUSTERED_ARGS_1' 
    3247 return simd_##METAL_SUFFIX(x); \ 
    3248 ^ 
    3249 <scratch space>:103:1: note: expanded from here 
    3250 simd_min 
    3251 ^ 
    3252 program_source:392:1: error: implicit declaration of function 'quad_min' is
    3253 invalid in OpenCL 
    3254 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3255 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3256 ^ 
    3257 program_source:360:60: note: expanded from macro '\ 
    3258 DECLARE_I_REDUCTION_CLUSTERED' 
    3259 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3260 ^ 
    3261 program_source:246:59: note: expanded from macro '\ 
    3262 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3263 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3264 ^ 
    3265 program_source:196:12: note: expanded from macro '\ 
    3266 OVERLOAD_CLUSTERED_ARGS_1' 
    3267 return quad_##METAL_SUFFIX(x); \ 
    3268 ^ 
    3269 <scratch space>:104:1: note: expanded from here 
    3270 quad_min 
    3271 ^ 
    3272 program_source:392:1: error: implicit declaration of function 'simd_min' is
    3273 invalid in OpenCL 
    3274 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3275 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3276 ^ 
    3277 program_source:360:60: note: expanded from macro '\ 
    3278 DECLARE_I_REDUCTION_CLUSTERED' 
    3279 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3280 ^ 
    3281 program_source:246:59: note: expanded from macro '\ 
    3282 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3283 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3284 ^ 
    3285 program_source:198:12: note: expanded from macro '\ 
    3286 OVERLOAD_CLUSTERED_ARGS_1' 
    3287 return simd_##METAL_SUFFIX(x); \ 
    3288 ^ 
    3289 <scratch space>:105:1: note: expanded from here 
    3290 simd_min 
    3291 ^ 
    3292 program_source:392:1: error: implicit declaration of function 'quad_min' is
    3293 invalid in OpenCL 
    3294 program_source:385:52: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3295 DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3296 ^ 
    3297 program_source:363:60: note: expanded from macro '\ 
    3298 DECLARE_F_REDUCTION_CLUSTERED' 
    3299 #define DECLARE_F_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3300 ^ 
    3301 program_source:249:71: note: expanded from macro '\ 
    3302 BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1' 
    3303 #define BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3304 ^ 
    3305 program_source:196:12: note: expanded from macro '\ 
    3306 OVERLOAD_CLUSTERED_ARGS_1' 
    3307 return quad_##METAL_SUFFIX(x); \ 
    3308 ^ 
    3309 <scratch space>:107:1: note: expanded from here 
    3310 quad_min 
    3311 ^ 
    3312 program_source:392:1: error: implicit declaration of function 'simd_min' is
    3313 invalid in OpenCL 
    3314 program_source:385:52: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3315 DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3316 ^ 
    3317 program_source:363:60: note: expanded from macro '\ 
    3318 DECLARE_F_REDUCTION_CLUSTERED' 
    3319 #define DECLARE_F_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3320 ^ 
    3321 program_source:249:71: note: expanded from macro '\ 
    3322 BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1' 
    3323 #define BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3324 ^ 
    3325 program_source:198:12: note: expanded from macro '\ 
    3326 OVERLOAD_CLUSTERED_ARGS_1' 
    3327 return simd_##METAL_SUFFIX(x); \ 
    3328 ^ 
    3329 <scratch space>:108:1: note: expanded from here 
    3330 simd_min 
    3331 ^ 
    3332 program_source:393:1: error: implicit declaration of function 'quad_max' is
    3333 invalid in OpenCL 
    3334 DECLARE_REDUCTION_CLUSTERED(max, reduce_max) 
    3335 ^ 
    3336 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3337 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3338 ^ 
    3339 program_source:360:60: note: expanded from macro '\ 
    3340 DECLARE_I_REDUCTION_CLUSTERED' 
    3341 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3342 ^ 
    3343 program_source:245:71: note: expanded from macro '\ 
    3344 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3345 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3346 ^ 
    3347 program_source:196:12: note: expanded from macro '\ 
    3348 OVERLOAD_CLUSTERED_ARGS_1' 
    3349 return quad_##METAL_SUFFIX(x); \ 
    3350 ^ 
    3351 <scratch space>:110:1: note: expanded from here 
    3352 quad_max 
    3353 ^ 
    3354 program_source:393:1: error: implicit declaration of function 'simd_max' is
    3355 invalid in OpenCL 
    3356 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3357 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3358 ^ 
    3359 program_source:360:60: note: expanded from macro '\ 
    3360 DECLARE_I_REDUCTION_CLUSTERED' 
    3361 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3362 ^ 
    3363 program_source:245:71: note: expanded from macro '\ 
    3364 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3365 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3366 ^ 
    3367 program_source:198:12: note: expanded from macro '\ 
    3368 OVERLOAD_CLUSTERED_ARGS_1' 
    3369 return simd_##METAL_SUFFIX(x); \ 
    3370 ^ 
    3371 <scratch space>:111:1: note: expanded from here 
    3372 simd_max 
    3373 ^ 
    3374 program_source:393:1: error: implicit declaration of function 'quad_max' is
    3375 invalid in OpenCL 
    3376 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3377 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3378 ^ 
    3379 program_source:360:60: note: expanded from macro '\ 
    3380 DECLARE_I_REDUCTION_CLUSTERED' 
    3381 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3382 ^ 
    3383 program_source:246:59: note: expanded from macro '\ 
    3384 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3385 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3386 ^ 
    3387 program_source:196:12: note: expanded from macro '\ 
    3388 OVERLOAD_CLUSTERED_ARGS_1' 
    3389 return quad_##METAL_SUFFIX(x); \ 
    3390 ^ 
    3391 <scratch space>:112:1: note: expanded from here 
    3392 quad_max 
    3393 ^ 
    3394 program_source:393:1: error: implicit declaration of function 'simd_max' is
    3395 invalid in OpenCL 
    3396 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3397 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3398 ^ 
    3399 program_source:360:60: note: expanded from macro '\ 
    3400 DECLARE_I_REDUCTION_CLUSTERED' 
    3401 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3402 ^ 
    3403 program_source:246:59: note: expanded from macro '\ 
    3404 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3405 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3406 ^ 
    3407 program_source:198:12: note: expanded from macro '\ 
    3408 OVERLOAD_CLUSTERED_ARGS_1' 
    3409 return simd_##METAL_SUFFIX(x); \ 
    3410 ^ 
    3411 <scratch space>:113:1: note: expanded from here 
    3412 simd_max 
    3413 ^ 
    3414 program_source:393:1: error: implicit declaration of function 'quad_max' is
    3415 invalid in OpenCL 
    3416 program_source:385:52: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3417 DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3418 ^ 
    3419 program_source:363:60: note: expanded from macro '\ 
    3420 DECLARE_F_REDUCTION_CLUSTERED' 
    3421 #define DECLARE_F_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3422 ^ 
    3423 program_source:249:71: note: expanded from macro '\ 
    3424 BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1' 
    3425 #define BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3426 ^ 
    3427 program_source:196:12: note: expanded from macro '\ 
    3428 OVERLOAD_CLUSTERED_ARGS_1' 
    3429 return quad_##METAL_SUFFIX(x); \ 
    3430 ^ 
    3431 <scratch space>:115:1: note: expanded from here 
    3432 quad_max 
    3433 ^ 
    3434 program_source:393:1: error: implicit declaration of function 'simd_max' is
    3435 invalid in OpenCL 
    3436 program_source:385:52: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3437 DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3438 ^ 
    3439 program_source:363:60: note: expanded from macro '\ 
    3440 DECLARE_F_REDUCTION_CLUSTERED' 
    3441 #define DECLARE_F_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3442 ^ 
    3443 program_source:249:71: note: expanded from macro '\ 
    3444 BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1' 
    3445 #define BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3446 ^ 
    3447 program_source:198:12: note: expanded from macro '\ 
    3448 OVERLOAD_CLUSTERED_ARGS_1' 
    3449 return simd_##METAL_SUFFIX(x); \ 
    3450 ^ 
    3451 <scratch space>:116:1: note: expanded from here 
    3452 simd_max 
    3453 ^ 
    3454 program_source:395:1: error: implicit declaration of function 'quad_product'
    3455 is invalid in OpenCL 
    3456 DECLARE_REDUCTION_CLUSTERED(product, reduce_mul) 
    3457 ^ 
    3458 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3459 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3460 ^ 
    3461 program_source:360:60: note: expanded from macro '\ 
    3462 DECLARE_I_REDUCTION_CLUSTERED' 
    3463 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3464 ^ 
    3465 program_source:245:71: note: expanded from macro '\ 
    3466 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3467 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3468 ^ 
    3469 program_source:196:12: note: expanded from macro '\ 
    3470 OVERLOAD_CLUSTERED_ARGS_1' 
    3471 return quad_##METAL_SUFFIX(x); \ 
    3472 ^ 
    3473 <scratch space>:118:1: note: expanded from here 
    3474 quad_product 
    3475 ^ 
    3476 program_source:395:1: error: implicit declaration of function 'simd_product'
    3477 is invalid in OpenCL 
    3478 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3479 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3480 ^ 
    3481 program_source:360:60: note: expanded from macro '\ 
    3482 DECLARE_I_REDUCTION_CLUSTERED' 
    3483 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3484 ^ 
    3485 program_source:245:71: note: expanded from macro '\ 
    3486 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3487 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3488 ^ 
    3489 program_source:198:12: note: expanded from macro '\ 
    3490 OVERLOAD_CLUSTERED_ARGS_1' 
    3491 return simd_##METAL_SUFFIX(x); \ 
    3492 ^ 
    3493 <scratch space>:119:1: note: expanded from here 
    3494 simd_product 
    3495 ^ 
    3496 program_source:395:1: error: implicit declaration of function 'quad_product'
    3497 is invalid in OpenCL 
    3498 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3499 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3500 ^ 
    3501 program_source:360:60: note: expanded from macro '\ 
    3502 DECLARE_I_REDUCTION_CLUSTERED' 
    3503 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3504 ^ 
    3505 program_source:246:59: note: expanded from macro '\ 
    3506 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3507 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3508 ^ 
    3509 program_source:196:12: note: expanded from macro '\ 
    3510 OVERLOAD_CLUSTERED_ARGS_1' 
    3511 return quad_##METAL_SUFFIX(x); \ 
    3512 ^ 
    3513 <scratch space>:120:1: note: expanded from here 
    3514 quad_product 
    3515 ^ 
    3516 program_source:395:1: error: implicit declaration of function 'simd_product'
    3517 is invalid in OpenCL 
    3518 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3519 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3520 ^ 
    3521 program_source:360:60: note: expanded from macro '\ 
    3522 DECLARE_I_REDUCTION_CLUSTERED' 
    3523 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3524 ^ 
    3525 program_source:246:59: note: expanded from macro '\ 
    3526 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3527 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3528 ^ 
    3529 program_source:198:12: note: expanded from macro '\ 
    3530 OVERLOAD_CLUSTERED_ARGS_1' 
    3531 return simd_##METAL_SUFFIX(x); \ 
    3532 ^ 
    3533 <scratch space>:121:1: note: expanded from here 
    3534 simd_product 
    3535 ^ 
    3536 program_source:395:1: error: implicit declaration of function 'quad_product'
    3537 is invalid in OpenCL 
    3538 program_source:385:52: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3539 DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3540 ^ 
    3541 program_source:363:60: note: expanded from macro '\ 
    3542 DECLARE_F_REDUCTION_CLUSTERED' 
    3543 #define DECLARE_F_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3544 ^ 
    3545 program_source:249:71: note: expanded from macro '\ 
    3546 BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1' 
    3547 #define BRIDGE_FUNCTION_CLUSTERED_F_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>:123:1: note: expanded from here 
    3554 quad_product 
    3555 ^ 
    3556 program_source:395:1: error: implicit declaration of function 'simd_product'
    3557 is invalid in OpenCL 
    3558 program_source:385:52: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    3559 DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3560 ^ 
    3561 program_source:363:60: note: expanded from macro '\ 
    3562 DECLARE_F_REDUCTION_CLUSTERED' 
    3563 #define DECLARE_F_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3564 ^ 
    3565 program_source:249:71: note: expanded from macro '\ 
    3566 BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1' 
    3567 #define BRIDGE_FUNCTION_CLUSTERED_F_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3568 ^ 
    3569 program_source:198:12: note: expanded from macro '\ 
    3570 OVERLOAD_CLUSTERED_ARGS_1' 
    3571 return simd_##METAL_SUFFIX(x); \ 
    3572 ^ 
    3573 <scratch space>:124:1: note: expanded from here 
    3574 simd_product 
    3575 ^ 
    3576 program_source:396:1: error: implicit declaration of function 'quad_and' is
    3577 invalid in OpenCL 
    3578 DECLARE_I_REDUCTION_CLUSTERED(and, reduce_and) 
    3579 ^ 
    3580 program_source:360:60: note: expanded from macro
    3581 'DECLARE_I_REDUCTION_CLUSTERED' 
    3582 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3583 ^ 
    3584 program_source:245:71: note: expanded from macro '\ 
    3585 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3586 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3587 ^ 
    3588 program_source:196:12: note: expanded from macro '\ 
    3589 OVERLOAD_CLUSTERED_ARGS_1' 
    3590 return quad_##METAL_SUFFIX(x); \ 
    3591 ^ 
    3592 <scratch space>:126:1: note: expanded from here 
    3593 quad_and 
    3594 ^ 
    3595 program_source:396:1: error: implicit declaration of function 'simd_and' is
    3596 invalid in OpenCL 
    3597 program_source:360:60: note: expanded from macro
    3598 'DECLARE_I_REDUCTION_CLUSTERED' 
    3599 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3600 ^ 
    3601 program_source:245:71: note: expanded from macro '\ 
    3602 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3603 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3604 ^ 
    3605 program_source:198:12: note: expanded from macro '\ 
    3606 OVERLOAD_CLUSTERED_ARGS_1' 
    3607 return simd_##METAL_SUFFIX(x); \ 
    3608 ^ 
    3609 <scratch space>:127:1: note: expanded from here 
    3610 simd_and 
    3611 ^ 
    3612 program_source:396:1: error: implicit declaration of function 'quad_and' is
    3613 invalid in OpenCL 
    3614 program_source:360:60: note: expanded from macro
    3615 'DECLARE_I_REDUCTION_CLUSTERED' 
    3616 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3617 ^ 
    3618 program_source:246:59: note: expanded from macro '\ 
    3619 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3620 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3621 ^ 
    3622 program_source:196:12: note: expanded from macro '\ 
    3623 OVERLOAD_CLUSTERED_ARGS_1' 
    3624 return quad_##METAL_SUFFIX(x); \ 
    3625 ^ 
    3626 <scratch space>:128:1: note: expanded from here 
    3627 quad_and 
    3628 ^ 
    3629 program_source:396:1: error: implicit declaration of function 'simd_and' is
    3630 invalid in OpenCL 
    3631 program_source:360:60: note: expanded from macro
    3632 'DECLARE_I_REDUCTION_CLUSTERED' 
    3633 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3634 ^ 
    3635 program_source:246:59: note: expanded from macro '\ 
    3636 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3637 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3638 ^ 
    3639 program_source:198:12: note: expanded from macro '\ 
    3640 OVERLOAD_CLUSTERED_ARGS_1' 
    3641 return simd_##METAL_SUFFIX(x); \ 
    3642 ^ 
    3643 <scratch space>:129:1: note: expanded from here 
    3644 simd_and 
    3645 ^ 
    3646 program_source:397:1: error: implicit declaration of function 'quad_or' is
    3647 invalid in OpenCL 
    3648 DECLARE_I_REDUCTION_CLUSTERED(or, reduce_or) 
    3649 ^ 
    3650 program_source:360:60: note: expanded from macro
    3651 'DECLARE_I_REDUCTION_CLUSTERED' 
    3652 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3653 ^ 
    3654 program_source:245:71: note: expanded from macro '\ 
    3655 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3656 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3657 ^ 
    3658 program_source:196:12: note: expanded from macro '\ 
    3659 OVERLOAD_CLUSTERED_ARGS_1' 
    3660 return quad_##METAL_SUFFIX(x); \ 
    3661 ^ 
    3662 <scratch space>:131:1: note: expanded from here 
    3663 quad_or 
    3664 ^ 
    3665 program_source:397:1: error: implicit declaration of function 'simd_or' is
    3666 invalid in OpenCL 
    3667 program_source:360:60: note: expanded from macro
    3668 'DECLARE_I_REDUCTION_CLUSTERED' 
    3669 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3670 ^ 
    3671 program_source:245:71: note: expanded from macro '\ 
    3672 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3673 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3674 ^ 
    3675 program_source:198:12: note: expanded from macro '\ 
    3676 OVERLOAD_CLUSTERED_ARGS_1' 
    3677 return simd_##METAL_SUFFIX(x); \ 
    3678 ^ 
    3679 <scratch space>:132:1: note: expanded from here 
    3680 simd_or 
    3681 ^ 
    3682 program_source:397:1: error: implicit declaration of function 'quad_or' is
    3683 invalid in OpenCL 
    3684 program_source:360:60: note: expanded from macro
    3685 'DECLARE_I_REDUCTION_CLUSTERED' 
    3686 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3687 ^ 
    3688 program_source:246:59: note: expanded from macro '\ 
    3689 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3690 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3691 ^ 
    3692 program_source:196:12: note: expanded from macro '\ 
    3693 OVERLOAD_CLUSTERED_ARGS_1' 
    3694 return quad_##METAL_SUFFIX(x); \ 
    3695 ^ 
    3696 <scratch space>:133:1: note: expanded from here 
    3697 quad_or 
    3698 ^ 
    3699 program_source:397:1: error: implicit declaration of function 'simd_or' is
    3700 invalid in OpenCL 
    3701 program_source:360:60: note: expanded from macro
    3702 'DECLARE_I_REDUCTION_CLUSTERED' 
    3703 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3704 ^ 
    3705 program_source:246:59: note: expanded from macro '\ 
    3706 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3707 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3708 ^ 
    3709 program_source:198:12: note: expanded from macro '\ 
    3710 OVERLOAD_CLUSTERED_ARGS_1' 
    3711 return simd_##METAL_SUFFIX(x); \ 
    3712 ^ 
    3713 <scratch space>:134:1: note: expanded from here 
    3714 simd_or 
    3715 ^ 
    3716 program_source:398:1: error: implicit declaration of function 'quad_xor' is
    3717 invalid in OpenCL 
    3718 DECLARE_I_REDUCTION_CLUSTERED(xor, reduce_xor) 
    3719 ^ 
    3720 program_source:360:60: note: expanded from macro
    3721 'DECLARE_I_REDUCTION_CLUSTERED' 
    3722 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3723 ^ 
    3724 program_source:245:71: note: expanded from macro '\ 
    3725 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3726 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3727 ^ 
    3728 program_source:196:12: note: expanded from macro '\ 
    3729 OVERLOAD_CLUSTERED_ARGS_1' 
    3730 return quad_##METAL_SUFFIX(x); \ 
    3731 ^ 
    3732 <scratch space>:136:1: note: expanded from here 
    3733 quad_xor 
    3734 ^ 
    3735 program_source:398:1: error: implicit declaration of function 'simd_xor' is
    3736 invalid in OpenCL 
    3737 program_source:360:60: note: expanded from macro
    3738 'DECLARE_I_REDUCTION_CLUSTERED' 
    3739 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3740 ^ 
    3741 program_source:245:71: note: expanded from macro '\ 
    3742 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3743 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    3744 ^ 
    3745 program_source:198:12: note: expanded from macro '\ 
    3746 OVERLOAD_CLUSTERED_ARGS_1' 
    3747 return simd_##METAL_SUFFIX(x); \ 
    3748 ^ 
    3749 <scratch space>:137:1: note: expanded from here 
    3750 simd_xor 
    3751 ^ 
    3752 program_source:398:1: error: implicit declaration of function 'quad_xor' is
    3753 invalid in OpenCL 
    3754 program_source:360:60: note: expanded from macro
    3755 'DECLARE_I_REDUCTION_CLUSTERED' 
    3756 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3757 ^ 
    3758 program_source:246:59: note: expanded from macro '\ 
    3759 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3760 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3761 ^ 
    3762 program_source:196:12: note: expanded from macro '\ 
    3763 OVERLOAD_CLUSTERED_ARGS_1' 
    3764 return quad_##METAL_SUFFIX(x); \ 
    3765 ^ 
    3766 <scratch space>:138:1: note: expanded from here 
    3767 quad_xor 
    3768 ^ 
    3769 program_source:398:1: error: implicit declaration of function 'simd_xor' is
    3770 invalid in OpenCL 
    3771 program_source:360:60: note: expanded from macro
    3772 'DECLARE_I_REDUCTION_CLUSTERED' 
    3773 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3774 ^ 
    3775 program_source:246:59: note: expanded from macro '\ 
    3776 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    3777 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3778 ^ 
    3779 program_source:198:12: note: expanded from macro '\ 
    3780 OVERLOAD_CLUSTERED_ARGS_1' 
    3781 return simd_##METAL_SUFFIX(x); \ 
    3782 ^ 
    3783 <scratch space>:139:1: note: expanded from here 
    3784 simd_xor 
    3785 ^ 
    3786 program_source:399:1: error: implicit declaration of function 'simd_and' is
    3787 invalid in OpenCL 
    3788 DECLARE_B_REDUCTION_CLUSTERED(and) 
    3789 ^ 
    3790 program_source:369:8: note: expanded from macro
    3791 'DECLARE_B_REDUCTION_CLUSTERED' 
    3792 return simd_##OP(select(0, 1, predicate != 0)); \ 
    3793 ^ 
    3794 <scratch space>:141:1: note: expanded from here 
    3795 simd_and 
    3796 ^ 
    3797 program_source:399:1: error: implicit declaration of function 'quad_and' is
    3798 invalid in OpenCL 
    3799 program_source:378:12: note: expanded from macro
    3800 'DECLARE_B_REDUCTION_CLUSTERED' 
    3801 return quad_##OP(x); \ 
    3802 ^ 
    3803 <scratch space>:143:1: note: expanded from here 
    3804 quad_and 
    3805 ^ 
    3806 program_source:399:1: error: implicit declaration of function 'simd_and' is
    3807 invalid in OpenCL 
    3808 program_source:380:12: note: expanded from macro
    3809 'DECLARE_B_REDUCTION_CLUSTERED' 
    3810 return simd_##OP(x); \ 
    3811 ^ 
    3812 <scratch space>:144:1: note: expanded from here 
    3813 simd_and 
    3814 ^ 
    3815 program_source:400:1: error: implicit declaration of function 'simd_or' is
    3816 invalid in OpenCL 
    3817 DECLARE_B_REDUCTION_CLUSTERED(or) 
    3818 ^ 
    3819 program_source:369:8: note: expanded from macro
    3820 'DECLARE_B_REDUCTION_CLUSTERED' 
    3821 return simd_##OP(select(0, 1, predicate != 0)); \ 
    3822 ^ 
    3823 <scratch space>:146:1: note: expanded from here 
    3824 simd_or 
    3825 ^ 
    3826 program_source:400:1: error: implicit declaration of function 'quad_or' is
    3827 invalid in OpenCL 
    3828 program_source:378:12: note: expanded from macro
    3829 'DECLARE_B_REDUCTION_CLUSTERED' 
    3830 return quad_##OP(x); \ 
    3831 ^ 
    3832 <scratch space>:148:1: note: expanded from here 
    3833 quad_or 
    3834 ^ 
    3835 program_source:400:1: error: implicit declaration of function 'simd_or' is
    3836 invalid in OpenCL 
    3837 program_source:380:12: note: expanded from macro
    3838 'DECLARE_B_REDUCTION_CLUSTERED' 
    3839 return simd_##OP(x); \ 
    3840 ^ 
    3841 <scratch space>:149:1: note: expanded from here 
    3842 simd_or 
    3843 ^ 
    3844 program_source:401:1: error: implicit declaration of function 'simd_xor' is
    3845 invalid in OpenCL 
    3846 DECLARE_B_REDUCTION_CLUSTERED(xor) 
    3847 ^ 
    3848 program_source:369:8: note: expanded from macro
    3849 'DECLARE_B_REDUCTION_CLUSTERED' 
    3850 return simd_##OP(select(0, 1, predicate != 0)); \ 
    3851 ^ 
    3852 <scratch space>:151:1: note: expanded from here 
    3853 simd_xor 
    3854 ^ 
    3855 program_source:401:1: error: implicit declaration of function 'quad_xor' is
    3856 invalid in OpenCL 
    3857 program_source:378:12: note: expanded from macro
    3858 'DECLARE_B_REDUCTION_CLUSTERED' 
    3859 return quad_##OP(x); \ 
    3860 ^ 
    3861 <scratch space>:153:1: note: expanded from here 
    3862 quad_xor 
    3863 ^ 
    3864 program_source:401:1: error: implicit declaration of function 'simd_xor' is
    3865 invalid in OpenCL 
    3866 program_source:380:12: note: expanded from macro
    3867 'DECLARE_B_REDUCTION_CLUSTERED' 
    3868 return simd_##OP(x); \ 
    3869 ^ 
    3870 <scratch space>:154:1: note: expanded from here 
    3871 simd_xor 
    3872 ^ 
    3873 program_source:403:1: error: implicit declaration of function
    3874 'quad_shuffle_rotate_down' is invalid in OpenCL 
    3875 DECLARE_SHUFFLE_CLUSTERED(shuffle_rotate_down, rotate) 
    3876 ^ 
    3877 program_source:388:56: note: expanded from macro 'DECLARE_SHUFFLE_CLUSTERED' 
    3878 #define DECLARE_SHUFFLE_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3879 ^ 
    3880 program_source:252:69: note: expanded from macro '\ 
    3881 BRIDGE_FUNCTION_CLUSTERED_ARGS_2' 
    3882 #define BRIDGE_FUNCTION_CLUSTERED_ARGS_2(METAL_SUFFIX, OPENCL_EXPR) \ 
    3883 ^ 
    3884 program_source:208:12: note: expanded from macro '\ 
    3885 OVERLOAD_CLUSTERED_ARGS_2' 
    3886 return quad_##METAL_SUFFIX(x, ushort(delta)); \ 
    3887 ^ 
    3888 <scratch space>:156:1: note: expanded from here 
    3889 quad_shuffle_rotate_down 
    3890 ^ 
    3891 program_source:403:1: error: implicit declaration of function
    3892 'simd_shuffle_rotate_down' is invalid in OpenCL 
    3893 program_source:388:56: note: expanded from macro 'DECLARE_SHUFFLE_CLUSTERED' 
    3894 #define DECLARE_SHUFFLE_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3895 ^ 
    3896 program_source:252:69: note: expanded from macro '\ 
    3897 BRIDGE_FUNCTION_CLUSTERED_ARGS_2' 
    3898 #define BRIDGE_FUNCTION_CLUSTERED_ARGS_2(METAL_SUFFIX, OPENCL_EXPR) \ 
    3899 ^ 
    3900 program_source:210:12: note: expanded from macro '\ 
    3901 OVERLOAD_CLUSTERED_ARGS_2' 
    3902 return simd_##METAL_SUFFIX(x, ushort(delta)); \ 
    3903 ^ 
    3904 <scratch space>:157:1: note: expanded from here 
    3905 simd_shuffle_rotate_down 
    3906 ^ 
    3907 program_source:403:1: error: implicit declaration of function
    3908 'quad_shuffle_rotate_down' is invalid in OpenCL 
    3909 program_source:388:56: note: expanded from macro 'DECLARE_SHUFFLE_CLUSTERED' 
    3910 #define DECLARE_SHUFFLE_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3911 ^ 
    3912 program_source:253:59: note: expanded from macro '\ 
    3913 BRIDGE_FUNCTION_CLUSTERED_ARGS_2' 
    3914 OVERLOAD_CLUSTERED_ARGS_2(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3915 ^ 
    3916 program_source:208:12: note: expanded from macro '\ 
    3917 OVERLOAD_CLUSTERED_ARGS_2' 
    3918 return quad_##METAL_SUFFIX(x, ushort(delta)); \ 
    3919 ^ 
    3920 <scratch space>:158:1: note: expanded from here 
    3921 quad_shuffle_rotate_down 
    3922 ^ 
    3923 program_source:403:1: error: implicit declaration of function
    3924 'simd_shuffle_rotate_down' is invalid in OpenCL 
    3925 program_source:388:56: note: expanded from macro 'DECLARE_SHUFFLE_CLUSTERED' 
    3926 #define DECLARE_SHUFFLE_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3927 ^ 
    3928 program_source:253:59: note: expanded from macro '\ 
    3929 BRIDGE_FUNCTION_CLUSTERED_ARGS_2' 
    3930 OVERLOAD_CLUSTERED_ARGS_2(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    3931 ^ 
    3932 program_source:210:12: note: expanded from macro '\ 
    3933 OVERLOAD_CLUSTERED_ARGS_2' 
    3934 return simd_##METAL_SUFFIX(x, ushort(delta)); \ 
    3935 ^ 
    3936 <scratch space>:159:1: note: expanded from here 
    3937 simd_shuffle_rotate_down 
    3938 ^ 
    3939 program_source:403:1: error: implicit declaration of function
    3940 'quad_shuffle_rotate_down' is invalid in OpenCL 
    3941 program_source:388:56: note: expanded from macro 'DECLARE_SHUFFLE_CLUSTERED' 
    3942 #define DECLARE_SHUFFLE_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3943 ^ 
    3944 program_source:254:60: note: expanded from macro '\ 
    3945 BRIDGE_FUNCTION_CLUSTERED_ARGS_2' 
    3946 OVERLOAD_CLUSTERED_ARGS_2(METAL_SUFFIX, OPENCL_EXPR, uint) \ 
    3947 ^ 
    3948 program_source:208:12: note: expanded from macro '\ 
    3949 OVERLOAD_CLUSTERED_ARGS_2' 
    3950 return quad_##METAL_SUFFIX(x, ushort(delta)); \ 
    3951 ^ 
    3952 <scratch space>:160:1: note: expanded from here 
    3953 quad_shuffle_rotate_down 
    3954 ^ 
    3955 program_source:403:1: error: implicit declaration of function
    3956 'simd_shuffle_rotate_down' is invalid in OpenCL 
    3957 program_source:388:56: note: expanded from macro 'DECLARE_SHUFFLE_CLUSTERED' 
    3958 #define DECLARE_SHUFFLE_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    3959 ^ 
    3960 program_source:254:60: note: expanded from macro '\ 
    3961 BRIDGE_FUNCTION_CLUSTERED_ARGS_2' 
    3962 OVERLOAD_CLUSTERED_ARGS_2(METAL_SUFFIX, OPENCL_EXPR, uint) \ 
    3963 ^ 
    3964 program_source:210:12: note: expanded from macro '\ 
    3965 OVERLOAD_CLUSTERED_ARGS_2' 
    3966 return simd_##METAL_SUFFIX(x, ushort(delta)); \ 
    3967 ^ 
    3968 <scratch space>:161:1: note: expanded from here 
    3969 simd_shuffle_rotate_down 
    3970 ^ 
    3971 program_source:409:1: error: illegal string literal in 'asm' 
    3972 EXPOSE_BALLOT(simd_is_first, , bool, ) 
    3973 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3974 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    3975 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    3976 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3977 program_source:410:1: error: illegal string literal in 'asm' 
    3978 EXPOSE_BALLOT(simd_all, bool expr, bool, ) 
    3979 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3980 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    3981 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    3982 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3983 program_source:411:1: error: illegal string literal in 'asm' 
    3984 EXPOSE_BALLOT(simd_any, bool expr, bool, ) 
    3985 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3986 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    3987 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    3988 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3989 program_source:412:1: error: illegal string literal in 'asm' 
    3990 EXPOSE_BALLOT(simd_ballot, bool expr, ulong, .i64) 
    3991 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3992 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    3993 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    3994 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3995 program_source:413:1: error: illegal string literal in 'asm' 
    3996 EXPOSE_BALLOT(simd_active_threads_mask, , ulong, .i64) 
    3997 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    3998 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    3999 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    4000 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4001 program_source:415:1: error: illegal string literal in 'asm' 
    4002 EXPOSE_BALLOT(quad_is_first, , bool, ) 
    4003 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4004 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    4005 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    4006 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4007 program_source:416:1: error: illegal string literal in 'asm' 
    4008 EXPOSE_BALLOT(quad_all, bool expr, bool, ) 
    4009 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4010 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    4011 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    4012 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4013 program_source:417:1: error: illegal string literal in 'asm' 
    4014 EXPOSE_BALLOT(quad_any, bool expr, bool, ) 
    4015 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4016 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    4017 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    4018 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4019 program_source:418:1: error: illegal string literal in 'asm' 
    4020 EXPOSE_BALLOT(quad_ballot, bool expr, ushort, ) 
    4021 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4022 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    4023 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    4024 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4025 program_source:419:1: error: illegal string literal in 'asm' 
    4026 EXPOSE_BALLOT(quad_active_threads_mask, , ushort, ) 
    4027 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4028 program_source:407:9: note: expanded from macro 'EXPOSE_BALLOT' 
    4029 __asm("air." #FUNC_EXPR #AIR_EXPR); \ 
    4030 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4031 program_source:422:23: error: implicit declaration of function 'simd_is_first'
    4032 is invalid in OpenCL 
    4033 return select(0, 1, simd_is_first()); 
    4034 ^ 
    4035 program_source:426:23: error: implicit declaration of function 'simd_all' is
    4036 invalid in OpenCL 
    4037 return select(0, 1, simd_all(predicate != 0)); 
    4038 ^ 
    4039 program_source:430:23: error: implicit declaration of function 'simd_any' is
    4040 invalid in OpenCL 
    4041 return select(0, 1, simd_any(predicate != 0)); 
    4042 ^ 
    4043 program_source:434:23: error: implicit declaration of function 'simd_all' is
    4044 invalid in OpenCL 
    4045 return select(0, 1, simd_all(predicate != 0)); 
    4046 ^ 
    4047 program_source:438:23: error: implicit declaration of function 'simd_any' is
    4048 invalid in OpenCL 
    4049 return select(0, 1, simd_any(predicate != 0)); 
    4050 ^ 
    4051 program_source:443:14: error: implicit declaration of function 'simd_ballot'
    4052 is invalid in OpenCL 
    4053 output.x = simd_ballot(predicate != 0); 
    4054 ^ 
    4055 program_source:480:9: error: illegal string literal in 'asm' 
    4056 __asm("air.simdgroup.barrier"); 
    4057 ^~~~~~~~~~~~~~~~~~~~~~~ 
    4058 program_source:487:3: error: implicit declaration of function
    4059 '__metal_simdgroup_barrier' is invalid in OpenCL 
    4060 __metal_simdgroup_barrier(flags, 1); 
    4061 ^ 
    4062 program_source:487:3: note: did you mean 'simdgroup_barrier'? 
    4063 program_source:483:6: note: 'simdgroup_barrier' declared here 
    4064 void simdgroup_barrier(int flags) { 
    4065 ^ 
    4066 program_source:681:9: error: implicit declaration of function 'simd_is_first'
    4067 is invalid in OpenCL 
    4068 if (simd_is_first()) { 
    4069 ^ 
    4070 program_source:673:45: warning: comparison of integers of different signs:
    4071 '__private unsigned int' and '__private int' [-Wsign-compare] 
    4072 for (unsigned int index = thread; index < bufferSize; index +=
    4073 get_local_size(0)) 
    4074 ~~~~~ ^ ~~~~~~~~~~ 
    4075 program_source:685:45: warning: comparison of integers of different signs:
    4076 'unsigned int' and '__private int' [-Wsign-compare] 
    4077 if (thread%(32*2) == 0 && thread+32 < workGroupSize) 
    4078 ~~~~~~~~~ ^ ~~~~~~~~~~~~~ 
    4079 program_source:689:45: warning: comparison of integers of different signs:
    4080 'unsigned int' and '__private int' [-Wsign-compare] 
    4081 if (thread%(64*2) == 0 && thread+64 < workGroupSize) 
    4082 ~~~~~~~~~ ^ ~~~~~~~~~~~~~ 
    4083 program_source:693:47: warning: comparison of integers of different signs:
    4084 'unsigned int' and '__private int' [-Wsign-compare] 
    4085 if (thread%(128*2) == 0 && thread+128 < workGroupSize) 
    4086 ~~~~~~~~~~ ^ ~~~~~~~~~~~~~ 
    4087  
    4088  
    4089 openmm.OpenMMException: Error compiling kernel: program_source:279:1: error:
    4090 illegal string literal in 'asm' 
    4091 DECLARE_REDUCTION_BASE(sum) 
    4092 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4093 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4094 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4095 ^ 
    4096 program_source:263:44: note: expanded from macro '\ 
    4097 DECLARE_I_REDUCTION_BASE' 
    4098 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4099 ^ 
    4100 program_source:217:40: note: expanded from macro '\ 
    4101 EXPOSE_FUNCTION_I_ARGS_1' 
    4102 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4103 ^ 
    4104 program_source:164:9: note: expanded from macro '\ 
    4105 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4106 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4107 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4108 program_source:279:1: error: illegal string literal in 'asm' 
    4109 DECLARE_REDUCTION_BASE(sum) 
    4110 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4111 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4112 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4113 ^ 
    4114 program_source:263:44: note: expanded from macro '\ 
    4115 DECLARE_I_REDUCTION_BASE' 
    4116 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4117 ^ 
    4118 program_source:218:51: note: expanded from macro '\ 
    4119 EXPOSE_FUNCTION_I_ARGS_1' 
    4120 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4121 ^ 
    4122 program_source:164:9: note: expanded from macro '\ 
    4123 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4124 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4125 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4126 program_source:279:1: error: illegal string literal in 'asm' 
    4127 DECLARE_REDUCTION_BASE(sum) 
    4128 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4129 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4130 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4131 ^ 
    4132 program_source:264:43: note: expanded from macro '\ 
    4133 DECLARE_I_REDUCTION_BASE' 
    4134 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4135 ^ 
    4136 program_source:217:40: note: expanded from macro '\ 
    4137 EXPOSE_FUNCTION_I_ARGS_1' 
    4138 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4139 ^ 
    4140 program_source:164:9: note: expanded from macro '\ 
    4141 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4142 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4143 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4144 program_source:279:1: error: illegal string literal in 'asm' 
    4145 DECLARE_REDUCTION_BASE(sum) 
    4146 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4147 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4148 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4149 ^ 
    4150 program_source:264:43: note: expanded from macro '\ 
    4151 DECLARE_I_REDUCTION_BASE' 
    4152 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4153 ^ 
    4154 program_source:218:51: note: expanded from macro '\ 
    4155 EXPOSE_FUNCTION_I_ARGS_1' 
    4156 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4157 ^ 
    4158 program_source:164:9: note: expanded from macro '\ 
    4159 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4160 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4161 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4162 program_source:279:1: error: illegal string literal in 'asm' 
    4163 DECLARE_REDUCTION_BASE(sum) 
    4164 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4165 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4166 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4167 ^ 
    4168 program_source:267:44: note: expanded from macro '\ 
    4169 DECLARE_F_REDUCTION_BASE' 
    4170 #define DECLARE_F_REDUCTION_BASE(METAL_OP) \ 
    4171 ^ 
    4172 program_source:221:40: note: expanded from macro '\ 
    4173 EXPOSE_FUNCTION_F_ARGS_1' 
    4174 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4175 ^ 
    4176 program_source:164:9: note: expanded from macro '\ 
    4177 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4178 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4179 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4180 program_source:279:1: error: illegal string literal in 'asm' 
    4181 DECLARE_REDUCTION_BASE(sum) 
    4182 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4183 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4184 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4185 ^ 
    4186 program_source:268:43: note: expanded from macro '\ 
    4187 DECLARE_F_REDUCTION_BASE' 
    4188 EXPOSE_FUNCTION_F_ARGS_1(quad_##METAL_OP) \ 
    4189 ^ 
    4190 program_source:221:40: note: expanded from macro '\ 
    4191 EXPOSE_FUNCTION_F_ARGS_1' 
    4192 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4193 ^ 
    4194 program_source:164:9: note: expanded from macro '\ 
    4195 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4196 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4197 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4198 program_source:280:1: error: illegal string literal in 'asm' 
    4199 DECLARE_REDUCTION_BASE(prefix_inclusive_sum) 
    4200 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4201 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4202 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4203 ^ 
    4204 program_source:263:44: note: expanded from macro '\ 
    4205 DECLARE_I_REDUCTION_BASE' 
    4206 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4207 ^ 
    4208 program_source:217:40: note: expanded from macro '\ 
    4209 EXPOSE_FUNCTION_I_ARGS_1' 
    4210 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4211 ^ 
    4212 program_source:164:9: note: expanded from macro '\ 
    4213 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4214 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4215 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4216 program_source:280:1: error: illegal string literal in 'asm' 
    4217 DECLARE_REDUCTION_BASE(prefix_inclusive_sum) 
    4218 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4219 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4220 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4221 ^ 
    4222 program_source:263:44: note: expanded from macro '\ 
    4223 DECLARE_I_REDUCTION_BASE' 
    4224 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4225 ^ 
    4226 program_source:218:51: note: expanded from macro '\ 
    4227 EXPOSE_FUNCTION_I_ARGS_1' 
    4228 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4229 ^ 
    4230 program_source:164:9: note: expanded from macro '\ 
    4231 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4232 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4233 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4234 program_source:280:1: error: illegal string literal in 'asm' 
    4235 DECLARE_REDUCTION_BASE(prefix_inclusive_sum) 
    4236 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4237 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4238 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4239 ^ 
    4240 program_source:264:43: note: expanded from macro '\ 
    4241 DECLARE_I_REDUCTION_BASE' 
    4242 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4243 ^ 
    4244 program_source:217:40: note: expanded from macro '\ 
    4245 EXPOSE_FUNCTION_I_ARGS_1' 
    4246 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4247 ^ 
    4248 program_source:164:9: note: expanded from macro '\ 
    4249 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4250 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4251 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4252 program_source:280:1: error: illegal string literal in 'asm' 
    4253 DECLARE_REDUCTION_BASE(prefix_inclusive_sum) 
    4254 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4255 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4256 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4257 ^ 
    4258 program_source:264:43: note: expanded from macro '\ 
    4259 DECLARE_I_REDUCTION_BASE' 
    4260 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4261 ^ 
    4262 program_source:218:51: note: expanded from macro '\ 
    4263 EXPOSE_FUNCTION_I_ARGS_1' 
    4264 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4265 ^ 
    4266 program_source:164:9: note: expanded from macro '\ 
    4267 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4268 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4269 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4270 program_source:280:1: error: illegal string literal in 'asm' 
    4271 DECLARE_REDUCTION_BASE(prefix_inclusive_sum) 
    4272 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4273 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4274 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4275 ^ 
    4276 program_source:267:44: note: expanded from macro '\ 
    4277 DECLARE_F_REDUCTION_BASE' 
    4278 #define DECLARE_F_REDUCTION_BASE(METAL_OP) \ 
    4279 ^ 
    4280 program_source:221:40: note: expanded from macro '\ 
    4281 EXPOSE_FUNCTION_F_ARGS_1' 
    4282 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4283 ^ 
    4284 program_source:164:9: note: expanded from macro '\ 
    4285 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4286 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4287 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4288 program_source:280:1: error: illegal string literal in 'asm' 
    4289 DECLARE_REDUCTION_BASE(prefix_inclusive_sum) 
    4290 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4291 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4292 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4293 ^ 
    4294 program_source:268:43: note: expanded from macro '\ 
    4295 DECLARE_F_REDUCTION_BASE' 
    4296 EXPOSE_FUNCTION_F_ARGS_1(quad_##METAL_OP) \ 
    4297 ^ 
    4298 program_source:221:40: note: expanded from macro '\ 
    4299 EXPOSE_FUNCTION_F_ARGS_1' 
    4300 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4301 ^ 
    4302 program_source:164:9: note: expanded from macro '\ 
    4303 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4304 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4305 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4306 program_source:281:1: error: illegal string literal in 'asm' 
    4307 DECLARE_REDUCTION_BASE(prefix_exclusive_sum) 
    4308 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4309 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4310 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4311 ^ 
    4312 program_source:263:44: note: expanded from macro '\ 
    4313 DECLARE_I_REDUCTION_BASE' 
    4314 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4315 ^ 
    4316 program_source:217:40: note: expanded from macro '\ 
    4317 EXPOSE_FUNCTION_I_ARGS_1' 
    4318 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4319 ^ 
    4320 program_source:164:9: note: expanded from macro '\ 
    4321 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4322 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4323 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4324 program_source:281:1: error: illegal string literal in 'asm' 
    4325 DECLARE_REDUCTION_BASE(prefix_exclusive_sum) 
    4326 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4327 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4328 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4329 ^ 
    4330 program_source:263:44: note: expanded from macro '\ 
    4331 DECLARE_I_REDUCTION_BASE' 
    4332 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4333 ^ 
    4334 program_source:218:51: note: expanded from macro '\ 
    4335 EXPOSE_FUNCTION_I_ARGS_1' 
    4336 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4337 ^ 
    4338 program_source:164:9: note: expanded from macro '\ 
    4339 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4340 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4341 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4342 program_source:281:1: error: illegal string literal in 'asm' 
    4343 DECLARE_REDUCTION_BASE(prefix_exclusive_sum) 
    4344 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4345 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4346 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4347 ^ 
    4348 program_source:264:43: note: expanded from macro '\ 
    4349 DECLARE_I_REDUCTION_BASE' 
    4350 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4351 ^ 
    4352 program_source:217:40: note: expanded from macro '\ 
    4353 EXPOSE_FUNCTION_I_ARGS_1' 
    4354 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4355 ^ 
    4356 program_source:164:9: note: expanded from macro '\ 
    4357 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4358 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4359 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4360 program_source:281:1: error: illegal string literal in 'asm' 
    4361 DECLARE_REDUCTION_BASE(prefix_exclusive_sum) 
    4362 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4363 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4364 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4365 ^ 
    4366 program_source:264:43: note: expanded from macro '\ 
    4367 DECLARE_I_REDUCTION_BASE' 
    4368 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4369 ^ 
    4370 program_source:218:51: note: expanded from macro '\ 
    4371 EXPOSE_FUNCTION_I_ARGS_1' 
    4372 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4373 ^ 
    4374 program_source:164:9: note: expanded from macro '\ 
    4375 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4376 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4377 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4378 program_source:281:1: error: illegal string literal in 'asm' 
    4379 DECLARE_REDUCTION_BASE(prefix_exclusive_sum) 
    4380 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4381 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4382 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4383 ^ 
    4384 program_source:267:44: note: expanded from macro '\ 
    4385 DECLARE_F_REDUCTION_BASE' 
    4386 #define DECLARE_F_REDUCTION_BASE(METAL_OP) \ 
    4387 ^ 
    4388 program_source:221:40: note: expanded from macro '\ 
    4389 EXPOSE_FUNCTION_F_ARGS_1' 
    4390 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4391 ^ 
    4392 program_source:164:9: note: expanded from macro '\ 
    4393 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4394 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4395 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4396 program_source:281:1: error: illegal string literal in 'asm' 
    4397 DECLARE_REDUCTION_BASE(prefix_exclusive_sum) 
    4398 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4399 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4400 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4401 ^ 
    4402 program_source:268:43: note: expanded from macro '\ 
    4403 DECLARE_F_REDUCTION_BASE' 
    4404 EXPOSE_FUNCTION_F_ARGS_1(quad_##METAL_OP) \ 
    4405 ^ 
    4406 program_source:221:40: note: expanded from macro '\ 
    4407 EXPOSE_FUNCTION_F_ARGS_1' 
    4408 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4409 ^ 
    4410 program_source:164:9: note: expanded from macro '\ 
    4411 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4412 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4413 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4414 program_source:282:1: error: illegal string literal in 'asm' 
    4415 DECLARE_REDUCTION_BASE(min) 
    4416 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4417 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4418 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4419 ^ 
    4420 program_source:263:44: note: expanded from macro '\ 
    4421 DECLARE_I_REDUCTION_BASE' 
    4422 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4423 ^ 
    4424 program_source:217:40: note: expanded from macro '\ 
    4425 EXPOSE_FUNCTION_I_ARGS_1' 
    4426 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4427 ^ 
    4428 program_source:164:9: note: expanded from macro '\ 
    4429 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4430 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4431 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4432 program_source:282:1: error: illegal string literal in 'asm' 
    4433 DECLARE_REDUCTION_BASE(min) 
    4434 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4435 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4436 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4437 ^ 
    4438 program_source:263:44: note: expanded from macro '\ 
    4439 DECLARE_I_REDUCTION_BASE' 
    4440 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4441 ^ 
    4442 program_source:218:51: note: expanded from macro '\ 
    4443 EXPOSE_FUNCTION_I_ARGS_1' 
    4444 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4445 ^ 
    4446 program_source:164:9: note: expanded from macro '\ 
    4447 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4448 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4449 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4450 program_source:282:1: error: illegal string literal in 'asm' 
    4451 DECLARE_REDUCTION_BASE(min) 
    4452 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4453 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4454 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4455 ^ 
    4456 program_source:264:43: note: expanded from macro '\ 
    4457 DECLARE_I_REDUCTION_BASE' 
    4458 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4459 ^ 
    4460 program_source:217:40: note: expanded from macro '\ 
    4461 EXPOSE_FUNCTION_I_ARGS_1' 
    4462 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4463 ^ 
    4464 program_source:164:9: note: expanded from macro '\ 
    4465 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4466 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4467 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4468 program_source:282:1: error: illegal string literal in 'asm' 
    4469 DECLARE_REDUCTION_BASE(min) 
    4470 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4471 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4472 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4473 ^ 
    4474 program_source:264:43: note: expanded from macro '\ 
    4475 DECLARE_I_REDUCTION_BASE' 
    4476 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4477 ^ 
    4478 program_source:218:51: note: expanded from macro '\ 
    4479 EXPOSE_FUNCTION_I_ARGS_1' 
    4480 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4481 ^ 
    4482 program_source:164:9: note: expanded from macro '\ 
    4483 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4484 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4485 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4486 program_source:282:1: error: illegal string literal in 'asm' 
    4487 DECLARE_REDUCTION_BASE(min) 
    4488 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4489 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4490 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4491 ^ 
    4492 program_source:267:44: note: expanded from macro '\ 
    4493 DECLARE_F_REDUCTION_BASE' 
    4494 #define DECLARE_F_REDUCTION_BASE(METAL_OP) \ 
    4495 ^ 
    4496 program_source:221:40: note: expanded from macro '\ 
    4497 EXPOSE_FUNCTION_F_ARGS_1' 
    4498 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4499 ^ 
    4500 program_source:164:9: note: expanded from macro '\ 
    4501 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4502 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4503 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4504 program_source:282:1: error: illegal string literal in 'asm' 
    4505 DECLARE_REDUCTION_BASE(min) 
    4506 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4507 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4508 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4509 ^ 
    4510 program_source:268:43: note: expanded from macro '\ 
    4511 DECLARE_F_REDUCTION_BASE' 
    4512 EXPOSE_FUNCTION_F_ARGS_1(quad_##METAL_OP) \ 
    4513 ^ 
    4514 program_source:221:40: note: expanded from macro '\ 
    4515 EXPOSE_FUNCTION_F_ARGS_1' 
    4516 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4517 ^ 
    4518 program_source:164:9: note: expanded from macro '\ 
    4519 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4520 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4521 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4522 program_source:283:1: error: illegal string literal in 'asm' 
    4523 DECLARE_REDUCTION_BASE(max) 
    4524 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4525 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4526 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4527 ^ 
    4528 program_source:263:44: note: expanded from macro '\ 
    4529 DECLARE_I_REDUCTION_BASE' 
    4530 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4531 ^ 
    4532 program_source:217:40: note: expanded from macro '\ 
    4533 EXPOSE_FUNCTION_I_ARGS_1' 
    4534 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4535 ^ 
    4536 program_source:164:9: note: expanded from macro '\ 
    4537 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4538 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4539 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4540 program_source:283:1: error: illegal string literal in 'asm' 
    4541 DECLARE_REDUCTION_BASE(max) 
    4542 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4543 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4544 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4545 ^ 
    4546 program_source:263:44: note: expanded from macro '\ 
    4547 DECLARE_I_REDUCTION_BASE' 
    4548 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4549 ^ 
    4550 program_source:218:51: note: expanded from macro '\ 
    4551 EXPOSE_FUNCTION_I_ARGS_1' 
    4552 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4553 ^ 
    4554 program_source:164:9: note: expanded from macro '\ 
    4555 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4556 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4557 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4558 program_source:283:1: error: illegal string literal in 'asm' 
    4559 DECLARE_REDUCTION_BASE(max) 
    4560 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4561 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4562 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4563 ^ 
    4564 program_source:264:43: note: expanded from macro '\ 
    4565 DECLARE_I_REDUCTION_BASE' 
    4566 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4567 ^ 
    4568 program_source:217:40: note: expanded from macro '\ 
    4569 EXPOSE_FUNCTION_I_ARGS_1' 
    4570 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4571 ^ 
    4572 program_source:164:9: note: expanded from macro '\ 
    4573 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4574 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4575 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4576 program_source:283:1: error: illegal string literal in 'asm' 
    4577 DECLARE_REDUCTION_BASE(max) 
    4578 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4579 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4580 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4581 ^ 
    4582 program_source:264:43: note: expanded from macro '\ 
    4583 DECLARE_I_REDUCTION_BASE' 
    4584 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4585 ^ 
    4586 program_source:218:51: note: expanded from macro '\ 
    4587 EXPOSE_FUNCTION_I_ARGS_1' 
    4588 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4589 ^ 
    4590 program_source:164:9: note: expanded from macro '\ 
    4591 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4592 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4593 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4594 program_source:283:1: error: illegal string literal in 'asm' 
    4595 DECLARE_REDUCTION_BASE(max) 
    4596 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4597 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4598 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4599 ^ 
    4600 program_source:267:44: note: expanded from macro '\ 
    4601 DECLARE_F_REDUCTION_BASE' 
    4602 #define DECLARE_F_REDUCTION_BASE(METAL_OP) \ 
    4603 ^ 
    4604 program_source:221:40: note: expanded from macro '\ 
    4605 EXPOSE_FUNCTION_F_ARGS_1' 
    4606 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4607 ^ 
    4608 program_source:164:9: note: expanded from macro '\ 
    4609 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4610 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4611 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4612 program_source:283:1: error: illegal string literal in 'asm' 
    4613 DECLARE_REDUCTION_BASE(max) 
    4614 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4615 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4616 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4617 ^ 
    4618 program_source:268:43: note: expanded from macro '\ 
    4619 DECLARE_F_REDUCTION_BASE' 
    4620 EXPOSE_FUNCTION_F_ARGS_1(quad_##METAL_OP) \ 
    4621 ^ 
    4622 program_source:221:40: note: expanded from macro '\ 
    4623 EXPOSE_FUNCTION_F_ARGS_1' 
    4624 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4625 ^ 
    4626 program_source:164:9: note: expanded from macro '\ 
    4627 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4628 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4629 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4630 program_source:285:1: error: illegal string literal in 'asm' 
    4631 DECLARE_REDUCTION_BASE(product) 
    4632 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4633 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4634 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4635 ^ 
    4636 program_source:263:44: note: expanded from macro '\ 
    4637 DECLARE_I_REDUCTION_BASE' 
    4638 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4639 ^ 
    4640 program_source:217:40: note: expanded from macro '\ 
    4641 EXPOSE_FUNCTION_I_ARGS_1' 
    4642 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4643 ^ 
    4644 program_source:164:9: note: expanded from macro '\ 
    4645 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4646 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4647 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4648 program_source:285:1: error: illegal string literal in 'asm' 
    4649 DECLARE_REDUCTION_BASE(product) 
    4650 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4651 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4652 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4653 ^ 
    4654 program_source:263:44: note: expanded from macro '\ 
    4655 DECLARE_I_REDUCTION_BASE' 
    4656 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4657 ^ 
    4658 program_source:218:51: note: expanded from macro '\ 
    4659 EXPOSE_FUNCTION_I_ARGS_1' 
    4660 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4661 ^ 
    4662 program_source:164:9: note: expanded from macro '\ 
    4663 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4664 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4665 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4666 program_source:285:1: error: illegal string literal in 'asm' 
    4667 DECLARE_REDUCTION_BASE(product) 
    4668 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4669 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4670 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4671 ^ 
    4672 program_source:264:43: note: expanded from macro '\ 
    4673 DECLARE_I_REDUCTION_BASE' 
    4674 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4675 ^ 
    4676 program_source:217:40: note: expanded from macro '\ 
    4677 EXPOSE_FUNCTION_I_ARGS_1' 
    4678 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4679 ^ 
    4680 program_source:164:9: note: expanded from macro '\ 
    4681 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4682 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4683 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4684 program_source:285:1: error: illegal string literal in 'asm' 
    4685 DECLARE_REDUCTION_BASE(product) 
    4686 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4687 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4688 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4689 ^ 
    4690 program_source:264:43: note: expanded from macro '\ 
    4691 DECLARE_I_REDUCTION_BASE' 
    4692 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4693 ^ 
    4694 program_source:218:51: note: expanded from macro '\ 
    4695 EXPOSE_FUNCTION_I_ARGS_1' 
    4696 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4697 ^ 
    4698 program_source:164:9: note: expanded from macro '\ 
    4699 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4700 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4701 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4702 program_source:285:1: error: illegal string literal in 'asm' 
    4703 DECLARE_REDUCTION_BASE(product) 
    4704 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4705 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4706 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4707 ^ 
    4708 program_source:267:44: note: expanded from macro '\ 
    4709 DECLARE_F_REDUCTION_BASE' 
    4710 #define DECLARE_F_REDUCTION_BASE(METAL_OP) \ 
    4711 ^ 
    4712 program_source:221:40: note: expanded from macro '\ 
    4713 EXPOSE_FUNCTION_F_ARGS_1' 
    4714 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4715 ^ 
    4716 program_source:164:9: note: expanded from macro '\ 
    4717 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4718 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4719 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4720 program_source:285:1: error: illegal string literal in 'asm' 
    4721 DECLARE_REDUCTION_BASE(product) 
    4722 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4723 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4724 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4725 ^ 
    4726 program_source:268:43: note: expanded from macro '\ 
    4727 DECLARE_F_REDUCTION_BASE' 
    4728 EXPOSE_FUNCTION_F_ARGS_1(quad_##METAL_OP) \ 
    4729 ^ 
    4730 program_source:221:40: note: expanded from macro '\ 
    4731 EXPOSE_FUNCTION_F_ARGS_1' 
    4732 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4733 ^ 
    4734 program_source:164:9: note: expanded from macro '\ 
    4735 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4736 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4737 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4738 program_source:286:1: error: illegal string literal in 'asm' 
    4739 DECLARE_REDUCTION_BASE(prefix_inclusive_product) 
    4740 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4741 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4742 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4743 ^ 
    4744 program_source:263:44: note: expanded from macro '\ 
    4745 DECLARE_I_REDUCTION_BASE' 
    4746 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4747 ^ 
    4748 program_source:217:40: note: expanded from macro '\ 
    4749 EXPOSE_FUNCTION_I_ARGS_1' 
    4750 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4751 ^ 
    4752 program_source:164:9: note: expanded from macro '\ 
    4753 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4754 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4755 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4756 program_source:286:1: error: illegal string literal in 'asm' 
    4757 DECLARE_REDUCTION_BASE(prefix_inclusive_product) 
    4758 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4759 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4760 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4761 ^ 
    4762 program_source:263:44: note: expanded from macro '\ 
    4763 DECLARE_I_REDUCTION_BASE' 
    4764 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4765 ^ 
    4766 program_source:218:51: note: expanded from macro '\ 
    4767 EXPOSE_FUNCTION_I_ARGS_1' 
    4768 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4769 ^ 
    4770 program_source:164:9: note: expanded from macro '\ 
    4771 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4772 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4773 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4774 program_source:286:1: error: illegal string literal in 'asm' 
    4775 DECLARE_REDUCTION_BASE(prefix_inclusive_product) 
    4776 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4777 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4778 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4779 ^ 
    4780 program_source:264:43: note: expanded from macro '\ 
    4781 DECLARE_I_REDUCTION_BASE' 
    4782 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4783 ^ 
    4784 program_source:217:40: note: expanded from macro '\ 
    4785 EXPOSE_FUNCTION_I_ARGS_1' 
    4786 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4787 ^ 
    4788 program_source:164:9: note: expanded from macro '\ 
    4789 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4790 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4791 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4792 program_source:286:1: error: illegal string literal in 'asm' 
    4793 DECLARE_REDUCTION_BASE(prefix_inclusive_product) 
    4794 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4795 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4796 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4797 ^ 
    4798 program_source:264:43: note: expanded from macro '\ 
    4799 DECLARE_I_REDUCTION_BASE' 
    4800 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4801 ^ 
    4802 program_source:218:51: note: expanded from macro '\ 
    4803 EXPOSE_FUNCTION_I_ARGS_1' 
    4804 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4805 ^ 
    4806 program_source:164:9: note: expanded from macro '\ 
    4807 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4808 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4809 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4810 program_source:286:1: error: illegal string literal in 'asm' 
    4811 DECLARE_REDUCTION_BASE(prefix_inclusive_product) 
    4812 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4813 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4814 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4815 ^ 
    4816 program_source:267:44: note: expanded from macro '\ 
    4817 DECLARE_F_REDUCTION_BASE' 
    4818 #define DECLARE_F_REDUCTION_BASE(METAL_OP) \ 
    4819 ^ 
    4820 program_source:221:40: note: expanded from macro '\ 
    4821 EXPOSE_FUNCTION_F_ARGS_1' 
    4822 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4823 ^ 
    4824 program_source:164:9: note: expanded from macro '\ 
    4825 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4826 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4827 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4828 program_source:286:1: error: illegal string literal in 'asm' 
    4829 DECLARE_REDUCTION_BASE(prefix_inclusive_product) 
    4830 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4831 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4832 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4833 ^ 
    4834 program_source:268:43: note: expanded from macro '\ 
    4835 DECLARE_F_REDUCTION_BASE' 
    4836 EXPOSE_FUNCTION_F_ARGS_1(quad_##METAL_OP) \ 
    4837 ^ 
    4838 program_source:221:40: note: expanded from macro '\ 
    4839 EXPOSE_FUNCTION_F_ARGS_1' 
    4840 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4841 ^ 
    4842 program_source:164:9: note: expanded from macro '\ 
    4843 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4844 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4845 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4846 program_source:287:1: error: illegal string literal in 'asm' 
    4847 DECLARE_REDUCTION_BASE(prefix_exclusive_product) 
    4848 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4849 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4850 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4851 ^ 
    4852 program_source:263:44: note: expanded from macro '\ 
    4853 DECLARE_I_REDUCTION_BASE' 
    4854 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4855 ^ 
    4856 program_source:217:40: note: expanded from macro '\ 
    4857 EXPOSE_FUNCTION_I_ARGS_1' 
    4858 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4859 ^ 
    4860 program_source:164:9: note: expanded from macro '\ 
    4861 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4862 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4863 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4864 program_source:287:1: error: illegal string literal in 'asm' 
    4865 DECLARE_REDUCTION_BASE(prefix_exclusive_product) 
    4866 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4867 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4868 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4869 ^ 
    4870 program_source:263:44: note: expanded from macro '\ 
    4871 DECLARE_I_REDUCTION_BASE' 
    4872 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4873 ^ 
    4874 program_source:218:51: note: expanded from macro '\ 
    4875 EXPOSE_FUNCTION_I_ARGS_1' 
    4876 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4877 ^ 
    4878 program_source:164:9: note: expanded from macro '\ 
    4879 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4880 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4881 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4882 program_source:287:1: error: illegal string literal in 'asm' 
    4883 DECLARE_REDUCTION_BASE(prefix_exclusive_product) 
    4884 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4885 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4886 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4887 ^ 
    4888 program_source:264:43: note: expanded from macro '\ 
    4889 DECLARE_I_REDUCTION_BASE' 
    4890 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4891 ^ 
    4892 program_source:217:40: note: expanded from macro '\ 
    4893 EXPOSE_FUNCTION_I_ARGS_1' 
    4894 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4895 ^ 
    4896 program_source:164:9: note: expanded from macro '\ 
    4897 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4898 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4899 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4900 program_source:287:1: error: illegal string literal in 'asm' 
    4901 DECLARE_REDUCTION_BASE(prefix_exclusive_product) 
    4902 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4903 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4904 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    4905 ^ 
    4906 program_source:264:43: note: expanded from macro '\ 
    4907 DECLARE_I_REDUCTION_BASE' 
    4908 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4909 ^ 
    4910 program_source:218:51: note: expanded from macro '\ 
    4911 EXPOSE_FUNCTION_I_ARGS_1' 
    4912 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4913 ^ 
    4914 program_source:164:9: note: expanded from macro '\ 
    4915 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4916 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4917 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4918 program_source:287:1: error: illegal string literal in 'asm' 
    4919 DECLARE_REDUCTION_BASE(prefix_exclusive_product) 
    4920 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4921 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4922 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4923 ^ 
    4924 program_source:267:44: note: expanded from macro '\ 
    4925 DECLARE_F_REDUCTION_BASE' 
    4926 #define DECLARE_F_REDUCTION_BASE(METAL_OP) \ 
    4927 ^ 
    4928 program_source:221:40: note: expanded from macro '\ 
    4929 EXPOSE_FUNCTION_F_ARGS_1' 
    4930 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4931 ^ 
    4932 program_source:164:9: note: expanded from macro '\ 
    4933 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4934 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4935 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4936 program_source:287:1: error: illegal string literal in 'asm' 
    4937 DECLARE_REDUCTION_BASE(prefix_exclusive_product) 
    4938 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4939 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    4940 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4941 ^ 
    4942 program_source:268:43: note: expanded from macro '\ 
    4943 DECLARE_F_REDUCTION_BASE' 
    4944 EXPOSE_FUNCTION_F_ARGS_1(quad_##METAL_OP) \ 
    4945 ^ 
    4946 program_source:221:40: note: expanded from macro '\ 
    4947 EXPOSE_FUNCTION_F_ARGS_1' 
    4948 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    4949 ^ 
    4950 program_source:164:9: note: expanded from macro '\ 
    4951 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4952 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4953 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4954 program_source:288:1: error: illegal string literal in 'asm' 
    4955 DECLARE_I_REDUCTION_BASE(and) 
    4956 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4957 program_source:263:44: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    4958 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4959 ^ 
    4960 program_source:217:40: note: expanded from macro '\ 
    4961 EXPOSE_FUNCTION_I_ARGS_1' 
    4962 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4963 ^ 
    4964 program_source:164:9: note: expanded from macro '\ 
    4965 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4966 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4967 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4968 program_source:288:1: error: illegal string literal in 'asm' 
    4969 DECLARE_I_REDUCTION_BASE(and) 
    4970 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4971 program_source:263:44: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    4972 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    4973 ^ 
    4974 program_source:218:51: note: expanded from macro '\ 
    4975 EXPOSE_FUNCTION_I_ARGS_1' 
    4976 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    4977 ^ 
    4978 program_source:164:9: note: expanded from macro '\ 
    4979 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4980 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4981 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4982 program_source:288:1: error: illegal string literal in 'asm' 
    4983 DECLARE_I_REDUCTION_BASE(and) 
    4984 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4985 program_source:264:43: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    4986 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    4987 ^ 
    4988 program_source:217:40: note: expanded from macro '\ 
    4989 EXPOSE_FUNCTION_I_ARGS_1' 
    4990 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    4991 ^ 
    4992 program_source:164:9: note: expanded from macro '\ 
    4993 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    4994 __asm("air." #EXPR "." #AIR_TYPE); \ 
    4995 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4996 program_source:288:1: error: illegal string literal in 'asm' 
    4997 DECLARE_I_REDUCTION_BASE(and) 
    4998 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    4999 program_source:264:43: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    5000 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    5001 ^ 
    5002 program_source:218:51: note: expanded from macro '\ 
    5003 EXPOSE_FUNCTION_I_ARGS_1' 
    5004 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    5005 ^ 
    5006 program_source:164:9: note: expanded from macro '\ 
    5007 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    5008 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5009 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5010 program_source:289:1: error: illegal string literal in 'asm' 
    5011 DECLARE_I_REDUCTION_BASE(or) 
    5012 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5013 program_source:263:44: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    5014 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    5015 ^ 
    5016 program_source:217:40: note: expanded from macro '\ 
    5017 EXPOSE_FUNCTION_I_ARGS_1' 
    5018 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    5019 ^ 
    5020 program_source:164:9: note: expanded from macro '\ 
    5021 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    5022 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5023 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5024 program_source:289:1: error: illegal string literal in 'asm' 
    5025 DECLARE_I_REDUCTION_BASE(or) 
    5026 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5027 program_source:263:44: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    5028 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    5029 ^ 
    5030 program_source:218:51: note: expanded from macro '\ 
    5031 EXPOSE_FUNCTION_I_ARGS_1' 
    5032 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    5033 ^ 
    5034 program_source:164:9: note: expanded from macro '\ 
    5035 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    5036 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5037 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5038 program_source:289:1: error: illegal string literal in 'asm' 
    5039 DECLARE_I_REDUCTION_BASE(or) 
    5040 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5041 program_source:264:43: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    5042 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    5043 ^ 
    5044 program_source:217:40: note: expanded from macro '\ 
    5045 EXPOSE_FUNCTION_I_ARGS_1' 
    5046 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    5047 ^ 
    5048 program_source:164:9: note: expanded from macro '\ 
    5049 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    5050 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5051 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5052 program_source:289:1: error: illegal string literal in 'asm' 
    5053 DECLARE_I_REDUCTION_BASE(or) 
    5054 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5055 program_source:264:43: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    5056 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    5057 ^ 
    5058 program_source:218:51: note: expanded from macro '\ 
    5059 EXPOSE_FUNCTION_I_ARGS_1' 
    5060 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    5061 ^ 
    5062 program_source:164:9: note: expanded from macro '\ 
    5063 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    5064 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5065 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5066 program_source:290:1: error: illegal string literal in 'asm' 
    5067 DECLARE_I_REDUCTION_BASE(xor) 
    5068 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5069 program_source:263:44: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    5070 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    5071 ^ 
    5072 program_source:217:40: note: expanded from macro '\ 
    5073 EXPOSE_FUNCTION_I_ARGS_1' 
    5074 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    5075 ^ 
    5076 program_source:164:9: note: expanded from macro '\ 
    5077 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    5078 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5079 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5080 program_source:290:1: error: illegal string literal in 'asm' 
    5081 DECLARE_I_REDUCTION_BASE(xor) 
    5082 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5083 program_source:263:44: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    5084 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    5085 ^ 
    5086 program_source:218:51: note: expanded from macro '\ 
    5087 EXPOSE_FUNCTION_I_ARGS_1' 
    5088 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    5089 ^ 
    5090 program_source:164:9: note: expanded from macro '\ 
    5091 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    5092 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5093 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5094 program_source:290:1: error: illegal string literal in 'asm' 
    5095 DECLARE_I_REDUCTION_BASE(xor) 
    5096 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5097 program_source:264:43: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    5098 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    5099 ^ 
    5100 program_source:217:40: note: expanded from macro '\ 
    5101 EXPOSE_FUNCTION_I_ARGS_1' 
    5102 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    5103 ^ 
    5104 program_source:164:9: note: expanded from macro '\ 
    5105 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    5106 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5107 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5108 program_source:290:1: error: illegal string literal in 'asm' 
    5109 DECLARE_I_REDUCTION_BASE(xor) 
    5110 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5111 program_source:264:43: note: expanded from macro 'DECLARE_I_REDUCTION_BASE' 
    5112 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    5113 ^ 
    5114 program_source:218:51: note: expanded from macro '\ 
    5115 EXPOSE_FUNCTION_I_ARGS_1' 
    5116 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    5117 ^ 
    5118 program_source:164:9: note: expanded from macro '\ 
    5119 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    5120 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5121 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5122 program_source:292:1: error: illegal string literal in 'asm' 
    5123 DECLARE_SHUFFLE_BASE(broadcast) 
    5124 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5125 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5126 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5127 ^ 
    5128 program_source:224:38: note: expanded from macro '\ 
    5129 EXPOSE_FUNCTION_ARGS_2' 
    5130 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    5131 ^ 
    5132 program_source:168:9: note: expanded from macro '\ 
    5133 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5134 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5135 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5136 program_source:292:1: error: illegal string literal in 'asm' 
    5137 DECLARE_SHUFFLE_BASE(broadcast) 
    5138 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5139 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5140 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5141 ^ 
    5142 program_source:225:51: note: expanded from macro '\ 
    5143 EXPOSE_FUNCTION_ARGS_2' 
    5144 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    5145 ^ 
    5146 program_source:168:9: note: expanded from macro '\ 
    5147 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5148 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5149 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5150 program_source:292:1: error: illegal string literal in 'asm' 
    5151 DECLARE_SHUFFLE_BASE(broadcast) 
    5152 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5153 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5154 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5155 ^ 
    5156 program_source:226:52: note: expanded from macro '\ 
    5157 EXPOSE_FUNCTION_ARGS_2' 
    5158 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    5159 ^ 
    5160 program_source:168:9: note: expanded from macro '\ 
    5161 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5162 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5163 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5164 program_source:292:1: error: illegal string literal in 'asm' 
    5165 DECLARE_SHUFFLE_BASE(broadcast) 
    5166 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5167 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5168 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5169 ^ 
    5170 program_source:224:38: note: expanded from macro '\ 
    5171 EXPOSE_FUNCTION_ARGS_2' 
    5172 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    5173 ^ 
    5174 program_source:168:9: note: expanded from macro '\ 
    5175 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5176 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5177 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5178 program_source:292:1: error: illegal string literal in 'asm' 
    5179 DECLARE_SHUFFLE_BASE(broadcast) 
    5180 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5181 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5182 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5183 ^ 
    5184 program_source:225:51: note: expanded from macro '\ 
    5185 EXPOSE_FUNCTION_ARGS_2' 
    5186 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    5187 ^ 
    5188 program_source:168:9: note: expanded from macro '\ 
    5189 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5190 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5191 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5192 program_source:292:1: error: illegal string literal in 'asm' 
    5193 DECLARE_SHUFFLE_BASE(broadcast) 
    5194 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5195 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5196 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5197 ^ 
    5198 program_source:226:52: note: expanded from macro '\ 
    5199 EXPOSE_FUNCTION_ARGS_2' 
    5200 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    5201 ^ 
    5202 program_source:168:9: note: expanded from macro '\ 
    5203 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5204 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5205 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5206 program_source:293:1: error: illegal string literal in 'asm' 
    5207 DECLARE_REDUCTION_BASE(broadcast_first) 
    5208 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5209 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    5210 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    5211 ^ 
    5212 program_source:263:44: note: expanded from macro '\ 
    5213 DECLARE_I_REDUCTION_BASE' 
    5214 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    5215 ^ 
    5216 program_source:217:40: note: expanded from macro '\ 
    5217 EXPOSE_FUNCTION_I_ARGS_1' 
    5218 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    5219 ^ 
    5220 program_source:164:9: note: expanded from macro '\ 
    5221 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    5222 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5223 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5224 program_source:293:1: error: illegal string literal in 'asm' 
    5225 DECLARE_REDUCTION_BASE(broadcast_first) 
    5226 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5227 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    5228 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    5229 ^ 
    5230 program_source:263:44: note: expanded from macro '\ 
    5231 DECLARE_I_REDUCTION_BASE' 
    5232 #define DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    5233 ^ 
    5234 program_source:218:51: note: expanded from macro '\ 
    5235 EXPOSE_FUNCTION_I_ARGS_1' 
    5236 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    5237 ^ 
    5238 program_source:164:9: note: expanded from macro '\ 
    5239 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    5240 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5241 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5242 program_source:293:1: error: illegal string literal in 'asm' 
    5243 DECLARE_REDUCTION_BASE(broadcast_first) 
    5244 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5245 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    5246 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    5247 ^ 
    5248 program_source:264:43: note: expanded from macro '\ 
    5249 DECLARE_I_REDUCTION_BASE' 
    5250 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    5251 ^ 
    5252 program_source:217:40: note: expanded from macro '\ 
    5253 EXPOSE_FUNCTION_I_ARGS_1' 
    5254 #define EXPOSE_FUNCTION_I_ARGS_1(EXPR) \ 
    5255 ^ 
    5256 program_source:164:9: note: expanded from macro '\ 
    5257 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    5258 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5259 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5260 program_source:293:1: error: illegal string literal in 'asm' 
    5261 DECLARE_REDUCTION_BASE(broadcast_first) 
    5262 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5263 program_source:271:42: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    5264 #define DECLARE_REDUCTION_BASE(METAL_OP) \ 
    5265 ^ 
    5266 program_source:264:43: note: expanded from macro '\ 
    5267 DECLARE_I_REDUCTION_BASE' 
    5268 EXPOSE_FUNCTION_I_ARGS_1(quad_##METAL_OP) \ 
    5269 ^ 
    5270 program_source:218:51: note: expanded from macro '\ 
    5271 EXPOSE_FUNCTION_I_ARGS_1' 
    5272 EXPOSE_FUNCTION_OVERLOAD_ARGS_1(EXPR, int, s.i32) \ 
    5273 ^ 
    5274 program_source:164:9: note: expanded from macro '\ 
    5275 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    5276 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5277 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5278 program_source:293:1: error: illegal string literal in 'asm' 
    5279 DECLARE_REDUCTION_BASE(broadcast_first) 
    5280 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5281 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    5282 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    5283 ^ 
    5284 program_source:267:44: note: expanded from macro '\ 
    5285 DECLARE_F_REDUCTION_BASE' 
    5286 #define DECLARE_F_REDUCTION_BASE(METAL_OP) \ 
    5287 ^ 
    5288 program_source:221:40: note: expanded from macro '\ 
    5289 EXPOSE_FUNCTION_F_ARGS_1' 
    5290 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    5291 ^ 
    5292 program_source:164:9: note: expanded from macro '\ 
    5293 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    5294 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5295 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5296 program_source:293:1: error: illegal string literal in 'asm' 
    5297 DECLARE_REDUCTION_BASE(broadcast_first) 
    5298 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5299 program_source:272:36: note: expanded from macro 'DECLARE_REDUCTION_BASE' 
    5300 DECLARE_I_REDUCTION_BASE(METAL_OP) \ 
    5301 ^ 
    5302 program_source:268:43: note: expanded from macro '\ 
    5303 DECLARE_F_REDUCTION_BASE' 
    5304 EXPOSE_FUNCTION_F_ARGS_1(quad_##METAL_OP) \ 
    5305 ^ 
    5306 program_source:221:40: note: expanded from macro '\ 
    5307 EXPOSE_FUNCTION_F_ARGS_1' 
    5308 #define EXPOSE_FUNCTION_F_ARGS_1(EXPR) \ 
    5309 ^ 
    5310 program_source:164:9: note: expanded from macro '\ 
    5311 EXPOSE_FUNCTION_OVERLOAD_ARGS_1' 
    5312 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5313 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5314 program_source:295:1: error: illegal string literal in 'asm' 
    5315 DECLARE_SHUFFLE_BASE(shuffle) 
    5316 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5317 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5318 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5319 ^ 
    5320 program_source:224:38: note: expanded from macro '\ 
    5321 EXPOSE_FUNCTION_ARGS_2' 
    5322 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    5323 ^ 
    5324 program_source:168:9: note: expanded from macro '\ 
    5325 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5326 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5327 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5328 program_source:295:1: error: illegal string literal in 'asm' 
    5329 DECLARE_SHUFFLE_BASE(shuffle) 
    5330 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5331 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5332 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5333 ^ 
    5334 program_source:225:51: note: expanded from macro '\ 
    5335 EXPOSE_FUNCTION_ARGS_2' 
    5336 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    5337 ^ 
    5338 program_source:168:9: note: expanded from macro '\ 
    5339 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5340 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5341 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5342 program_source:295:1: error: illegal string literal in 'asm' 
    5343 DECLARE_SHUFFLE_BASE(shuffle) 
    5344 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5345 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5346 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5347 ^ 
    5348 program_source:226:52: note: expanded from macro '\ 
    5349 EXPOSE_FUNCTION_ARGS_2' 
    5350 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    5351 ^ 
    5352 program_source:168:9: note: expanded from macro '\ 
    5353 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5354 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5355 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5356 program_source:295:1: error: illegal string literal in 'asm' 
    5357 DECLARE_SHUFFLE_BASE(shuffle) 
    5358 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5359 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5360 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5361 ^ 
    5362 program_source:224:38: note: expanded from macro '\ 
    5363 EXPOSE_FUNCTION_ARGS_2' 
    5364 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    5365 ^ 
    5366 program_source:168:9: note: expanded from macro '\ 
    5367 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5368 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5369 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5370 program_source:295:1: error: illegal string literal in 'asm' 
    5371 DECLARE_SHUFFLE_BASE(shuffle) 
    5372 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5373 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5374 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5375 ^ 
    5376 program_source:225:51: note: expanded from macro '\ 
    5377 EXPOSE_FUNCTION_ARGS_2' 
    5378 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    5379 ^ 
    5380 program_source:168:9: note: expanded from macro '\ 
    5381 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5382 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5383 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5384 program_source:295:1: error: illegal string literal in 'asm' 
    5385 DECLARE_SHUFFLE_BASE(shuffle) 
    5386 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5387 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5388 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5389 ^ 
    5390 program_source:226:52: note: expanded from macro '\ 
    5391 EXPOSE_FUNCTION_ARGS_2' 
    5392 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    5393 ^ 
    5394 program_source:168:9: note: expanded from macro '\ 
    5395 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5396 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5397 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5398 program_source:296:1: error: illegal string literal in 'asm' 
    5399 DECLARE_SHUFFLE_BASE(shuffle_xor) 
    5400 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5401 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5402 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5403 ^ 
    5404 program_source:224:38: note: expanded from macro '\ 
    5405 EXPOSE_FUNCTION_ARGS_2' 
    5406 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    5407 ^ 
    5408 program_source:168:9: note: expanded from macro '\ 
    5409 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5410 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5411 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5412 program_source:296:1: error: illegal string literal in 'asm' 
    5413 DECLARE_SHUFFLE_BASE(shuffle_xor) 
    5414 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5415 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5416 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5417 ^ 
    5418 program_source:225:51: note: expanded from macro '\ 
    5419 EXPOSE_FUNCTION_ARGS_2' 
    5420 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    5421 ^ 
    5422 program_source:168:9: note: expanded from macro '\ 
    5423 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5424 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5425 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5426 program_source:296:1: error: illegal string literal in 'asm' 
    5427 DECLARE_SHUFFLE_BASE(shuffle_xor) 
    5428 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5429 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5430 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5431 ^ 
    5432 program_source:226:52: note: expanded from macro '\ 
    5433 EXPOSE_FUNCTION_ARGS_2' 
    5434 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    5435 ^ 
    5436 program_source:168:9: note: expanded from macro '\ 
    5437 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5438 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5439 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5440 program_source:296:1: error: illegal string literal in 'asm' 
    5441 DECLARE_SHUFFLE_BASE(shuffle_xor) 
    5442 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5443 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5444 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5445 ^ 
    5446 program_source:224:38: note: expanded from macro '\ 
    5447 EXPOSE_FUNCTION_ARGS_2' 
    5448 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    5449 ^ 
    5450 program_source:168:9: note: expanded from macro '\ 
    5451 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5452 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5453 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5454 program_source:296:1: error: illegal string literal in 'asm' 
    5455 DECLARE_SHUFFLE_BASE(shuffle_xor) 
    5456 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5457 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5458 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5459 ^ 
    5460 program_source:225:51: note: expanded from macro '\ 
    5461 EXPOSE_FUNCTION_ARGS_2' 
    5462 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    5463 ^ 
    5464 program_source:168:9: note: expanded from macro '\ 
    5465 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5466 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5467 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5468 program_source:296:1: error: illegal string literal in 'asm' 
    5469 DECLARE_SHUFFLE_BASE(shuffle_xor) 
    5470 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5471 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5472 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5473 ^ 
    5474 program_source:226:52: note: expanded from macro '\ 
    5475 EXPOSE_FUNCTION_ARGS_2' 
    5476 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    5477 ^ 
    5478 program_source:168:9: note: expanded from macro '\ 
    5479 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5480 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5481 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5482 program_source:297:1: error: illegal string literal in 'asm' 
    5483 DECLARE_SHUFFLE_BASE(shuffle_up) 
    5484 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5485 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5486 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5487 ^ 
    5488 program_source:224:38: note: expanded from macro '\ 
    5489 EXPOSE_FUNCTION_ARGS_2' 
    5490 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    5491 ^ 
    5492 program_source:168:9: note: expanded from macro '\ 
    5493 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5494 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5495 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5496 program_source:297:1: error: illegal string literal in 'asm' 
    5497 DECLARE_SHUFFLE_BASE(shuffle_up) 
    5498 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5499 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5500 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5501 ^ 
    5502 program_source:225:51: note: expanded from macro '\ 
    5503 EXPOSE_FUNCTION_ARGS_2' 
    5504 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    5505 ^ 
    5506 program_source:168:9: note: expanded from macro '\ 
    5507 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5508 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5509 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5510 program_source:297:1: error: illegal string literal in 'asm' 
    5511 DECLARE_SHUFFLE_BASE(shuffle_up) 
    5512 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5513 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5514 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5515 ^ 
    5516 program_source:226:52: note: expanded from macro '\ 
    5517 EXPOSE_FUNCTION_ARGS_2' 
    5518 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    5519 ^ 
    5520 program_source:168:9: note: expanded from macro '\ 
    5521 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5522 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5523 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5524 program_source:297:1: error: illegal string literal in 'asm' 
    5525 DECLARE_SHUFFLE_BASE(shuffle_up) 
    5526 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5527 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5528 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5529 ^ 
    5530 program_source:224:38: note: expanded from macro '\ 
    5531 EXPOSE_FUNCTION_ARGS_2' 
    5532 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    5533 ^ 
    5534 program_source:168:9: note: expanded from macro '\ 
    5535 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5536 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5537 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5538 program_source:297:1: error: illegal string literal in 'asm' 
    5539 DECLARE_SHUFFLE_BASE(shuffle_up) 
    5540 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5541 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5542 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5543 ^ 
    5544 program_source:225:51: note: expanded from macro '\ 
    5545 EXPOSE_FUNCTION_ARGS_2' 
    5546 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    5547 ^ 
    5548 program_source:168:9: note: expanded from macro '\ 
    5549 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5550 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5551 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5552 program_source:297:1: error: illegal string literal in 'asm' 
    5553 DECLARE_SHUFFLE_BASE(shuffle_up) 
    5554 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5555 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5556 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5557 ^ 
    5558 program_source:226:52: note: expanded from macro '\ 
    5559 EXPOSE_FUNCTION_ARGS_2' 
    5560 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    5561 ^ 
    5562 program_source:168:9: note: expanded from macro '\ 
    5563 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5564 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5565 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5566 program_source:298:1: error: illegal string literal in 'asm' 
    5567 DECLARE_SHUFFLE_BASE(shuffle_down) 
    5568 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5569 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5570 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5571 ^ 
    5572 program_source:224:38: note: expanded from macro '\ 
    5573 EXPOSE_FUNCTION_ARGS_2' 
    5574 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    5575 ^ 
    5576 program_source:168:9: note: expanded from macro '\ 
    5577 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5578 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5579 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5580 program_source:298:1: error: illegal string literal in 'asm' 
    5581 DECLARE_SHUFFLE_BASE(shuffle_down) 
    5582 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5583 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5584 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5585 ^ 
    5586 program_source:225:51: note: expanded from macro '\ 
    5587 EXPOSE_FUNCTION_ARGS_2' 
    5588 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    5589 ^ 
    5590 program_source:168:9: note: expanded from macro '\ 
    5591 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5592 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5593 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5594 program_source:298:1: error: illegal string literal in 'asm' 
    5595 DECLARE_SHUFFLE_BASE(shuffle_down) 
    5596 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5597 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5598 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5599 ^ 
    5600 program_source:226:52: note: expanded from macro '\ 
    5601 EXPOSE_FUNCTION_ARGS_2' 
    5602 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    5603 ^ 
    5604 program_source:168:9: note: expanded from macro '\ 
    5605 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5606 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5607 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5608 program_source:298:1: error: illegal string literal in 'asm' 
    5609 DECLARE_SHUFFLE_BASE(shuffle_down) 
    5610 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5611 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5612 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5613 ^ 
    5614 program_source:224:38: note: expanded from macro '\ 
    5615 EXPOSE_FUNCTION_ARGS_2' 
    5616 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    5617 ^ 
    5618 program_source:168:9: note: expanded from macro '\ 
    5619 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5620 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5621 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5622 program_source:298:1: error: illegal string literal in 'asm' 
    5623 DECLARE_SHUFFLE_BASE(shuffle_down) 
    5624 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5625 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5626 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5627 ^ 
    5628 program_source:225:51: note: expanded from macro '\ 
    5629 EXPOSE_FUNCTION_ARGS_2' 
    5630 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    5631 ^ 
    5632 program_source:168:9: note: expanded from macro '\ 
    5633 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5634 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5635 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5636 program_source:298:1: error: illegal string literal in 'asm' 
    5637 DECLARE_SHUFFLE_BASE(shuffle_down) 
    5638 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5639 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5640 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5641 ^ 
    5642 program_source:226:52: note: expanded from macro '\ 
    5643 EXPOSE_FUNCTION_ARGS_2' 
    5644 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    5645 ^ 
    5646 program_source:168:9: note: expanded from macro '\ 
    5647 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5648 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5649 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5650 program_source:299:1: error: illegal string literal in 'asm' 
    5651 DECLARE_SHUFFLE_BASE(shuffle_rotate_up) 
    5652 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5653 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5654 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5655 ^ 
    5656 program_source:224:38: note: expanded from macro '\ 
    5657 EXPOSE_FUNCTION_ARGS_2' 
    5658 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    5659 ^ 
    5660 program_source:168:9: note: expanded from macro '\ 
    5661 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5662 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5663 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5664 program_source:299:1: error: illegal string literal in 'asm' 
    5665 DECLARE_SHUFFLE_BASE(shuffle_rotate_up) 
    5666 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5667 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5668 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5669 ^ 
    5670 program_source:225:51: note: expanded from macro '\ 
    5671 EXPOSE_FUNCTION_ARGS_2' 
    5672 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    5673 ^ 
    5674 program_source:168:9: note: expanded from macro '\ 
    5675 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5676 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5677 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5678 program_source:299:1: error: illegal string literal in 'asm' 
    5679 DECLARE_SHUFFLE_BASE(shuffle_rotate_up) 
    5680 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5681 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5682 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5683 ^ 
    5684 program_source:226:52: note: expanded from macro '\ 
    5685 EXPOSE_FUNCTION_ARGS_2' 
    5686 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    5687 ^ 
    5688 program_source:168:9: note: expanded from macro '\ 
    5689 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5690 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5691 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5692 program_source:299:1: error: illegal string literal in 'asm' 
    5693 DECLARE_SHUFFLE_BASE(shuffle_rotate_up) 
    5694 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5695 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5696 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5697 ^ 
    5698 program_source:224:38: note: expanded from macro '\ 
    5699 EXPOSE_FUNCTION_ARGS_2' 
    5700 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    5701 ^ 
    5702 program_source:168:9: note: expanded from macro '\ 
    5703 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5704 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5705 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5706 program_source:299:1: error: illegal string literal in 'asm' 
    5707 DECLARE_SHUFFLE_BASE(shuffle_rotate_up) 
    5708 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5709 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5710 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5711 ^ 
    5712 program_source:225:51: note: expanded from macro '\ 
    5713 EXPOSE_FUNCTION_ARGS_2' 
    5714 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    5715 ^ 
    5716 program_source:168:9: note: expanded from macro '\ 
    5717 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5718 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5719 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5720 program_source:299:1: error: illegal string literal in 'asm' 
    5721 DECLARE_SHUFFLE_BASE(shuffle_rotate_up) 
    5722 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5723 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5724 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5725 ^ 
    5726 program_source:226:52: note: expanded from macro '\ 
    5727 EXPOSE_FUNCTION_ARGS_2' 
    5728 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    5729 ^ 
    5730 program_source:168:9: note: expanded from macro '\ 
    5731 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5732 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5733 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5734 program_source:300:1: error: illegal string literal in 'asm' 
    5735 DECLARE_SHUFFLE_BASE(shuffle_rotate_down) 
    5736 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5737 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5738 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5739 ^ 
    5740 program_source:224:38: note: expanded from macro '\ 
    5741 EXPOSE_FUNCTION_ARGS_2' 
    5742 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    5743 ^ 
    5744 program_source:168:9: note: expanded from macro '\ 
    5745 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5746 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5747 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5748 program_source:300:1: error: illegal string literal in 'asm' 
    5749 DECLARE_SHUFFLE_BASE(shuffle_rotate_down) 
    5750 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5751 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5752 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5753 ^ 
    5754 program_source:225:51: note: expanded from macro '\ 
    5755 EXPOSE_FUNCTION_ARGS_2' 
    5756 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    5757 ^ 
    5758 program_source:168:9: note: expanded from macro '\ 
    5759 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5760 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5761 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5762 program_source:300:1: error: illegal string literal in 'asm' 
    5763 DECLARE_SHUFFLE_BASE(shuffle_rotate_down) 
    5764 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5765 program_source:275:40: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5766 #define DECLARE_SHUFFLE_BASE(METAL_OP) \ 
    5767 ^ 
    5768 program_source:226:52: note: expanded from macro '\ 
    5769 EXPOSE_FUNCTION_ARGS_2' 
    5770 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    5771 ^ 
    5772 program_source:168:9: note: expanded from macro '\ 
    5773 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5774 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5775 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5776 program_source:300:1: error: illegal string literal in 'asm' 
    5777 DECLARE_SHUFFLE_BASE(shuffle_rotate_down) 
    5778 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5779 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5780 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5781 ^ 
    5782 program_source:224:38: note: expanded from macro '\ 
    5783 EXPOSE_FUNCTION_ARGS_2' 
    5784 #define EXPOSE_FUNCTION_ARGS_2(EXPR) \ 
    5785 ^ 
    5786 program_source:168:9: note: expanded from macro '\ 
    5787 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5788 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5789 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5790 program_source:300:1: error: illegal string literal in 'asm' 
    5791 DECLARE_SHUFFLE_BASE(shuffle_rotate_down) 
    5792 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5793 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5794 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5795 ^ 
    5796 program_source:225:51: note: expanded from macro '\ 
    5797 EXPOSE_FUNCTION_ARGS_2' 
    5798 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, int, s.i32) \ 
    5799 ^ 
    5800 program_source:168:9: note: expanded from macro '\ 
    5801 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5802 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5803 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5804 program_source:300:1: error: illegal string literal in 'asm' 
    5805 DECLARE_SHUFFLE_BASE(shuffle_rotate_down) 
    5806 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5807 program_source:276:41: note: expanded from macro 'DECLARE_SHUFFLE_BASE' 
    5808 EXPOSE_FUNCTION_ARGS_2(quad_##METAL_OP) \ 
    5809 ^ 
    5810 program_source:226:52: note: expanded from macro '\ 
    5811 EXPOSE_FUNCTION_ARGS_2' 
    5812 EXPOSE_FUNCTION_OVERLOAD_ARGS_2(EXPR, uint, u.i32) \ 
    5813 ^ 
    5814 program_source:168:9: note: expanded from macro '\ 
    5815 EXPOSE_FUNCTION_OVERLOAD_ARGS_2' 
    5816 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5817 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5818 program_source:302:1: error: illegal string literal in 'asm' 
    5819 EXPOSE_FUNCTION_I_ARGS_3(ctz) 
    5820 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5821 program_source:229:40: note: expanded from macro 'EXPOSE_FUNCTION_I_ARGS_3' 
    5822 #define EXPOSE_FUNCTION_I_ARGS_3(EXPR) \ 
    5823 ^ 
    5824 program_source:172:9: note: expanded from macro '\ 
    5825 EXPOSE_FUNCTION_OVERLOAD_ARGS_3' 
    5826 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5827 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5828 program_source:302:1: error: implicit declaration of function '___metal_ctz'
    5829 is invalid in OpenCL 
    5830 program_source:229:40: note: expanded from macro 'EXPOSE_FUNCTION_I_ARGS_3' 
    5831 #define EXPOSE_FUNCTION_I_ARGS_3(EXPR) \ 
    5832 ^ 
    5833 program_source:175:10: note: expanded from macro '\ 
    5834 EXPOSE_FUNCTION_OVERLOAD_ARGS_3' 
    5835 return ___metal_##EXPR(x, false); \ 
    5836 ^ 
    5837 :12:1: note: expanded from here 
    5838 ___metal_ctz 
    5839 ^ 
    5840 program_source:302:1: error: illegal string literal in 'asm' 
    5841 EXPOSE_FUNCTION_I_ARGS_3(ctz) 
    5842 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5843 program_source:230:49: note: expanded from macro 'EXPOSE_FUNCTION_I_ARGS_3' 
    5844 EXPOSE_FUNCTION_OVERLOAD_ARGS_3(EXPR, int, i32) \ 
    5845 ^ 
    5846 program_source:172:9: note: expanded from macro '\ 
    5847 EXPOSE_FUNCTION_OVERLOAD_ARGS_3' 
    5848 __asm("air." #EXPR "." #AIR_TYPE); \ 
    5849 ^~~~~~~~~~~~~~~~~~~~~~~~~~ 
    5850 program_source:302:1: error: implicit declaration of function '___metal_ctz'
    5851 is invalid in OpenCL 
    5852 program_source:230:49: note: expanded from macro 'EXPOSE_FUNCTION_I_ARGS_3' 
    5853 EXPOSE_FUNCTION_OVERLOAD_ARGS_3(EXPR, int, i32) \ 
    5854 ^ 
    5855 program_source:175:10: note: expanded from macro '\ 
    5856 EXPOSE_FUNCTION_OVERLOAD_ARGS_3' 
    5857 return ___metal_##EXPR(x, false); \ 
    5858 ^ 
    5859 :16:1: note: expanded from here 
    5860 ___metal_ctz 
    5861 ^ 
    5862 program_source:317:1: error: implicit declaration of function 'simd_sum' is
    5863 invalid in OpenCL 
    5864 DECLARE_REDUCTION_UNIFORM(sum, reduce_add) 
    5865 ^ 
    5866 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5867 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5868 ^ 
    5869 program_source:305:26: note: expanded from macro '\ 
    5870 DECLARE_I_REDUCTION_UNIFORM' 
    5871 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5872 ^ 
    5873 :17:1: note: expanded from here 
    5874 simd_sum 
    5875 ^ 
    5876 program_source:317:1: error: implicit declaration of function 'simd_sum' is
    5877 invalid in OpenCL 
    5878 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5879 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5880 ^ 
    5881 program_source:305:26: note: expanded from macro '\ 
    5882 DECLARE_I_REDUCTION_UNIFORM' 
    5883 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5884 ^ 
    5885 :17:1: note: expanded from here 
    5886 simd_sum 
    5887 ^ 
    5888 program_source:317:1: error: implicit declaration of function 'simd_sum' is
    5889 invalid in OpenCL 
    5890 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5891 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5892 ^ 
    5893 program_source:308:26: note: expanded from macro '\ 
    5894 DECLARE_F_REDUCTION_UNIFORM' 
    5895 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5896 ^ 
    5897 :19:1: note: expanded from here 
    5898 simd_sum 
    5899 ^ 
    5900 program_source:318:1: error: implicit declaration of function
    5901 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    5902 DECLARE_REDUCTION_UNIFORM(prefix_inclusive_sum, scan_inclusive_add) 
    5903 ^ 
    5904 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5905 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5906 ^ 
    5907 program_source:305:26: note: expanded from macro '\ 
    5908 DECLARE_I_REDUCTION_UNIFORM' 
    5909 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5910 ^ 
    5911 :21:1: note: expanded from here 
    5912 simd_prefix_inclusive_sum 
    5913 ^ 
    5914 program_source:318:1: error: implicit declaration of function
    5915 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    5916 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5917 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5918 ^ 
    5919 program_source:305:26: note: expanded from macro '\ 
    5920 DECLARE_I_REDUCTION_UNIFORM' 
    5921 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5922 ^ 
    5923 :21:1: note: expanded from here 
    5924 simd_prefix_inclusive_sum 
    5925 ^ 
    5926 program_source:318:1: error: implicit declaration of function
    5927 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    5928 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5929 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5930 ^ 
    5931 program_source:308:26: note: expanded from macro '\ 
    5932 DECLARE_F_REDUCTION_UNIFORM' 
    5933 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5934 ^ 
    5935 :23:1: note: expanded from here 
    5936 simd_prefix_inclusive_sum 
    5937 ^ 
    5938 program_source:319:1: error: implicit declaration of function
    5939 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    5940 DECLARE_REDUCTION_UNIFORM(prefix_exclusive_sum, scan_exclusive_add) 
    5941 ^ 
    5942 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5943 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5944 ^ 
    5945 program_source:305:26: note: expanded from macro '\ 
    5946 DECLARE_I_REDUCTION_UNIFORM' 
    5947 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5948 ^ 
    5949 :25:1: note: expanded from here 
    5950 simd_prefix_exclusive_sum 
    5951 ^ 
    5952 program_source:319:1: error: implicit declaration of function
    5953 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    5954 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5955 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5956 ^ 
    5957 program_source:305:26: note: expanded from macro '\ 
    5958 DECLARE_I_REDUCTION_UNIFORM' 
    5959 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5960 ^ 
    5961 :25:1: note: expanded from here 
    5962 simd_prefix_exclusive_sum 
    5963 ^ 
    5964 program_source:319:1: error: implicit declaration of function
    5965 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    5966 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5967 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5968 ^ 
    5969 program_source:308:26: note: expanded from macro '\ 
    5970 DECLARE_F_REDUCTION_UNIFORM' 
    5971 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5972 ^ 
    5973 :27:1: note: expanded from here 
    5974 simd_prefix_exclusive_sum 
    5975 ^ 
    5976 program_source:320:1: error: implicit declaration of function 'simd_min' is
    5977 invalid in OpenCL 
    5978 DECLARE_REDUCTION_UNIFORM(min, reduce_min) 
    5979 ^ 
    5980 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5981 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5982 ^ 
    5983 program_source:305:26: note: expanded from macro '\ 
    5984 DECLARE_I_REDUCTION_UNIFORM' 
    5985 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5986 ^ 
    5987 :29:1: note: expanded from here 
    5988 simd_min 
    5989 ^ 
    5990 program_source:320:1: error: implicit declaration of function 'simd_min' is
    5991 invalid in OpenCL 
    5992 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    5993 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    5994 ^ 
    5995 program_source:305:26: note: expanded from macro '\ 
    5996 DECLARE_I_REDUCTION_UNIFORM' 
    5997 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    5998 ^ 
    5999 :29:1: note: expanded from here 
    6000 simd_min 
    6001 ^ 
    6002 program_source:320:1: error: implicit declaration of function 'simd_min' is
    6003 invalid in OpenCL 
    6004 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    6005 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6006 ^ 
    6007 program_source:308:26: note: expanded from macro '\ 
    6008 DECLARE_F_REDUCTION_UNIFORM' 
    6009 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6010 ^ 
    6011 :31:1: note: expanded from here 
    6012 simd_min 
    6013 ^ 
    6014 program_source:321:1: error: implicit declaration of function 'simd_max' is
    6015 invalid in OpenCL 
    6016 DECLARE_REDUCTION_UNIFORM(max, reduce_max) 
    6017 ^ 
    6018 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    6019 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6020 ^ 
    6021 program_source:305:26: note: expanded from macro '\ 
    6022 DECLARE_I_REDUCTION_UNIFORM' 
    6023 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6024 ^ 
    6025 :33:1: note: expanded from here 
    6026 simd_max 
    6027 ^ 
    6028 program_source:321:1: error: implicit declaration of function 'simd_max' is
    6029 invalid in OpenCL 
    6030 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    6031 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6032 ^ 
    6033 program_source:305:26: note: expanded from macro '\ 
    6034 DECLARE_I_REDUCTION_UNIFORM' 
    6035 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6036 ^ 
    6037 :33:1: note: expanded from here 
    6038 simd_max 
    6039 ^ 
    6040 program_source:321:1: error: implicit declaration of function 'simd_max' is
    6041 invalid in OpenCL 
    6042 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    6043 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6044 ^ 
    6045 program_source:308:26: note: expanded from macro '\ 
    6046 DECLARE_F_REDUCTION_UNIFORM' 
    6047 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6048 ^ 
    6049 :35:1: note: expanded from here 
    6050 simd_max 
    6051 ^ 
    6052 program_source:323:1: error: implicit declaration of function 'simd_shuffle'
    6053 is invalid in OpenCL 
    6054 DECLARE_SHUFFLE_UNIFORM(shuffle, shuffle) 
    6055 ^ 
    6056 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    6057 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6058 ^ 
    6059 :37:1: note: expanded from here 
    6060 simd_shuffle 
    6061 ^ 
    6062 program_source:323:1: error: implicit declaration of function 'simd_shuffle'
    6063 is invalid in OpenCL 
    6064 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    6065 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6066 ^ 
    6067 :37:1: note: expanded from here 
    6068 simd_shuffle 
    6069 ^ 
    6070 program_source:323:1: error: implicit declaration of function 'simd_shuffle'
    6071 is invalid in OpenCL 
    6072 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    6073 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6074 ^ 
    6075 :37:1: note: expanded from here 
    6076 simd_shuffle 
    6077 ^ 
    6078 program_source:324:1: error: implicit declaration of function
    6079 'simd_shuffle_xor' is invalid in OpenCL 
    6080 DECLARE_SHUFFLE_UNIFORM(shuffle_xor, shuffle_xor) 
    6081 ^ 
    6082 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    6083 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6084 ^ 
    6085 :39:1: note: expanded from here 
    6086 simd_shuffle_xor 
    6087 ^ 
    6088 program_source:324:1: error: implicit declaration of function
    6089 'simd_shuffle_xor' is invalid in OpenCL 
    6090 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    6091 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6092 ^ 
    6093 :39:1: note: expanded from here 
    6094 simd_shuffle_xor 
    6095 ^ 
    6096 program_source:324:1: error: implicit declaration of function
    6097 'simd_shuffle_xor' is invalid in OpenCL 
    6098 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    6099 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6100 ^ 
    6101 :39:1: note: expanded from here 
    6102 simd_shuffle_xor 
    6103 ^ 
    6104 program_source:325:1: error: implicit declaration of function
    6105 'simd_shuffle_up' is invalid in OpenCL 
    6106 DECLARE_SHUFFLE_UNIFORM(shuffle_up, shuffle_up) 
    6107 ^ 
    6108 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    6109 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6110 ^ 
    6111 :41:1: note: expanded from here 
    6112 simd_shuffle_up 
    6113 ^ 
    6114 program_source:325:1: error: implicit declaration of function
    6115 'simd_shuffle_up' is invalid in OpenCL 
    6116 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    6117 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6118 ^ 
    6119 :41:1: note: expanded from here 
    6120 simd_shuffle_up 
    6121 ^ 
    6122 program_source:325:1: error: implicit declaration of function
    6123 'simd_shuffle_up' is invalid in OpenCL 
    6124 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    6125 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6126 ^ 
    6127 :41:1: note: expanded from here 
    6128 simd_shuffle_up 
    6129 ^ 
    6130 program_source:326:1: error: implicit declaration of function
    6131 'simd_shuffle_down' is invalid in OpenCL 
    6132 DECLARE_SHUFFLE_UNIFORM(shuffle_down, shuffle_down) 
    6133 ^ 
    6134 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    6135 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6136 ^ 
    6137 :43:1: note: expanded from here 
    6138 simd_shuffle_down 
    6139 ^ 
    6140 program_source:326:1: error: implicit declaration of function
    6141 'simd_shuffle_down' is invalid in OpenCL 
    6142 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    6143 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6144 ^ 
    6145 :43:1: note: expanded from here 
    6146 simd_shuffle_down 
    6147 ^ 
    6148 program_source:326:1: error: implicit declaration of function
    6149 'simd_shuffle_down' is invalid in OpenCL 
    6150 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    6151 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6152 ^ 
    6153 :43:1: note: expanded from here 
    6154 simd_shuffle_down 
    6155 ^ 
    6156 program_source:327:1: error: implicit declaration of function
    6157 'simd_shuffle_rotate_down' is invalid in OpenCL 
    6158 DECLARE_SHUFFLE_UNIFORM(shuffle_rotate_down, rotate) 
    6159 ^ 
    6160 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    6161 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6162 ^ 
    6163 :45:1: note: expanded from here 
    6164 simd_shuffle_rotate_down 
    6165 ^ 
    6166 program_source:327:1: error: implicit declaration of function
    6167 'simd_shuffle_rotate_down' is invalid in OpenCL 
    6168 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    6169 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6170 ^ 
    6171 :45:1: note: expanded from here 
    6172 simd_shuffle_rotate_down 
    6173 ^ 
    6174 program_source:327:1: error: implicit declaration of function
    6175 'simd_shuffle_rotate_down' is invalid in OpenCL 
    6176 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    6177 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6178 ^ 
    6179 :45:1: note: expanded from here 
    6180 simd_shuffle_rotate_down 
    6181 ^ 
    6182 program_source:329:1: error: implicit declaration of function 'simd_broadcast'
    6183 is invalid in OpenCL 
    6184 DECLARE_SHUFFLE_UNIFORM(broadcast, broadcast) 
    6185 ^ 
    6186 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    6187 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6188 ^ 
    6189 :47:1: note: expanded from here 
    6190 simd_broadcast 
    6191 ^ 
    6192 program_source:329:1: error: implicit declaration of function 'simd_broadcast'
    6193 is invalid in OpenCL 
    6194 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    6195 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6196 ^ 
    6197 :47:1: note: expanded from here 
    6198 simd_broadcast 
    6199 ^ 
    6200 program_source:329:1: error: implicit declaration of function 'simd_broadcast'
    6201 is invalid in OpenCL 
    6202 program_source:315:24: note: expanded from macro 'DECLARE_SHUFFLE_UNIFORM' 
    6203 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6204 ^ 
    6205 :47:1: note: expanded from here 
    6206 simd_broadcast 
    6207 ^ 
    6208 program_source:330:1: error: implicit declaration of function
    6209 'simd_broadcast_first' is invalid in OpenCL 
    6210 DECLARE_REDUCTION_UNIFORM(broadcast_first, broadcast_first) 
    6211 ^ 
    6212 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    6213 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6214 ^ 
    6215 program_source:305:26: note: expanded from macro '\ 
    6216 DECLARE_I_REDUCTION_UNIFORM' 
    6217 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6218 ^ 
    6219 :49:1: note: expanded from here 
    6220 simd_broadcast_first 
    6221 ^ 
    6222 program_source:330:1: error: implicit declaration of function
    6223 'simd_broadcast_first' is invalid in OpenCL 
    6224 program_source:310:56: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    6225 #define DECLARE_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6226 ^ 
    6227 program_source:305:26: note: expanded from macro '\ 
    6228 DECLARE_I_REDUCTION_UNIFORM' 
    6229 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6230 ^ 
    6231 :49:1: note: expanded from here 
    6232 simd_broadcast_first 
    6233 ^ 
    6234 program_source:330:1: error: implicit declaration of function
    6235 'simd_broadcast_first' is invalid in OpenCL 
    6236 program_source:311:50: note: expanded from macro 'DECLARE_REDUCTION_UNIFORM' 
    6237 DECLARE_I_REDUCTION_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6238 ^ 
    6239 program_source:308:26: note: expanded from macro '\ 
    6240 DECLARE_F_REDUCTION_UNIFORM' 
    6241 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_##OPENCL_OP) \ 
    6242 ^ 
    6243 :51:1: note: expanded from here 
    6244 simd_broadcast_first 
    6245 ^ 
    6246 program_source:345:1: error: implicit declaration of function 'simd_sum' is
    6247 invalid in OpenCL 
    6248 DECLARE_REDUCTION_NON_UNIFORM(sum, reduce_add) 
    6249 ^ 
    6250 program_source:338:60: note: expanded from macro
    6251 'DECLARE_REDUCTION_NON_UNIFORM' 
    6252 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6253 ^ 
    6254 program_source:333:26: note: expanded from macro '\ 
    6255 DECLARE_I_REDUCTION_NON_UNIFORM' 
    6256 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6257 ^ 
    6258 :53:1: note: expanded from here 
    6259 simd_sum 
    6260 ^ 
    6261 program_source:345:1: error: implicit declaration of function 'simd_sum' is
    6262 invalid in OpenCL 
    6263 program_source:338:60: note: expanded from macro
    6264 'DECLARE_REDUCTION_NON_UNIFORM' 
    6265 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6266 ^ 
    6267 program_source:333:26: note: expanded from macro '\ 
    6268 DECLARE_I_REDUCTION_NON_UNIFORM' 
    6269 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6270 ^ 
    6271 :53:1: note: expanded from here 
    6272 simd_sum 
    6273 ^ 
    6274 program_source:345:1: error: implicit declaration of function 'simd_sum' is
    6275 invalid in OpenCL 
    6276 program_source:339:54: note: expanded from macro
    6277 'DECLARE_REDUCTION_NON_UNIFORM' 
    6278 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6279 ^ 
    6280 program_source:336:26: note: expanded from macro '\ 
    6281 DECLARE_F_REDUCTION_NON_UNIFORM' 
    6282 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6283 ^ 
    6284 :55:1: note: expanded from here 
    6285 simd_sum 
    6286 ^ 
    6287 program_source:346:1: error: implicit declaration of function
    6288 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    6289 DECLARE_REDUCTION_NON_UNIFORM(prefix_inclusive_sum, scan_inclusive_add) 
    6290 ^ 
    6291 program_source:338:60: note: expanded from macro
    6292 'DECLARE_REDUCTION_NON_UNIFORM' 
    6293 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6294 ^ 
    6295 program_source:333:26: note: expanded from macro '\ 
    6296 DECLARE_I_REDUCTION_NON_UNIFORM' 
    6297 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6298 ^ 
    6299 :57:1: note: expanded from here 
    6300 simd_prefix_inclusive_sum 
    6301 ^ 
    6302 program_source:346:1: error: implicit declaration of function
    6303 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    6304 program_source:338:60: note: expanded from macro
    6305 'DECLARE_REDUCTION_NON_UNIFORM' 
    6306 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6307 ^ 
    6308 program_source:333:26: note: expanded from macro '\ 
    6309 DECLARE_I_REDUCTION_NON_UNIFORM' 
    6310 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6311 ^ 
    6312 :57:1: note: expanded from here 
    6313 simd_prefix_inclusive_sum 
    6314 ^ 
    6315 program_source:346:1: error: implicit declaration of function
    6316 'simd_prefix_inclusive_sum' is invalid in OpenCL 
    6317 program_source:339:54: note: expanded from macro
    6318 'DECLARE_REDUCTION_NON_UNIFORM' 
    6319 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6320 ^ 
    6321 program_source:336:26: note: expanded from macro '\ 
    6322 DECLARE_F_REDUCTION_NON_UNIFORM' 
    6323 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6324 ^ 
    6325 :59:1: note: expanded from here 
    6326 simd_prefix_inclusive_sum 
    6327 ^ 
    6328 program_source:347:1: error: implicit declaration of function
    6329 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    6330 DECLARE_REDUCTION_NON_UNIFORM(prefix_exclusive_sum, scan_exclusive_add) 
    6331 ^ 
    6332 program_source:338:60: note: expanded from macro
    6333 'DECLARE_REDUCTION_NON_UNIFORM' 
    6334 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6335 ^ 
    6336 program_source:333:26: note: expanded from macro '\ 
    6337 DECLARE_I_REDUCTION_NON_UNIFORM' 
    6338 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6339 ^ 
    6340 :61:1: note: expanded from here 
    6341 simd_prefix_exclusive_sum 
    6342 ^ 
    6343 program_source:347:1: error: implicit declaration of function
    6344 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    6345 program_source:338:60: note: expanded from macro
    6346 'DECLARE_REDUCTION_NON_UNIFORM' 
    6347 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6348 ^ 
    6349 program_source:333:26: note: expanded from macro '\ 
    6350 DECLARE_I_REDUCTION_NON_UNIFORM' 
    6351 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6352 ^ 
    6353 :61:1: note: expanded from here 
    6354 simd_prefix_exclusive_sum 
    6355 ^ 
    6356 program_source:347:1: error: implicit declaration of function
    6357 'simd_prefix_exclusive_sum' is invalid in OpenCL 
    6358 program_source:339:54: note: expanded from macro
    6359 'DECLARE_REDUCTION_NON_UNIFORM' 
    6360 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6361 ^ 
    6362 program_source:336:26: note: expanded from macro '\ 
    6363 DECLARE_F_REDUCTION_NON_UNIFORM' 
    6364 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6365 ^ 
    6366 :63:1: note: expanded from here 
    6367 simd_prefix_exclusive_sum 
    6368 ^ 
    6369 program_source:348:1: error: implicit declaration of function 'simd_min' is
    6370 invalid in OpenCL 
    6371 DECLARE_REDUCTION_NON_UNIFORM(min, reduce_min) 
    6372 ^ 
    6373 program_source:338:60: note: expanded from macro
    6374 'DECLARE_REDUCTION_NON_UNIFORM' 
    6375 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6376 ^ 
    6377 program_source:333:26: note: expanded from macro '\ 
    6378 DECLARE_I_REDUCTION_NON_UNIFORM' 
    6379 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6380 ^ 
    6381 :65:1: note: expanded from here 
    6382 simd_min 
    6383 ^ 
    6384 program_source:348:1: error: implicit declaration of function 'simd_min' is
    6385 invalid in OpenCL 
    6386 program_source:338:60: note: expanded from macro
    6387 'DECLARE_REDUCTION_NON_UNIFORM' 
    6388 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6389 ^ 
    6390 program_source:333:26: note: expanded from macro '\ 
    6391 DECLARE_I_REDUCTION_NON_UNIFORM' 
    6392 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6393 ^ 
    6394 :65:1: note: expanded from here 
    6395 simd_min 
    6396 ^ 
    6397 program_source:348:1: error: implicit declaration of function 'simd_min' is
    6398 invalid in OpenCL 
    6399 program_source:339:54: note: expanded from macro
    6400 'DECLARE_REDUCTION_NON_UNIFORM' 
    6401 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6402 ^ 
    6403 program_source:336:26: note: expanded from macro '\ 
    6404 DECLARE_F_REDUCTION_NON_UNIFORM' 
    6405 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6406 ^ 
    6407 :67:1: note: expanded from here 
    6408 simd_min 
    6409 ^ 
    6410 program_source:349:1: error: implicit declaration of function 'simd_max' is
    6411 invalid in OpenCL 
    6412 DECLARE_REDUCTION_NON_UNIFORM(max, reduce_max) 
    6413 ^ 
    6414 program_source:338:60: note: expanded from macro
    6415 'DECLARE_REDUCTION_NON_UNIFORM' 
    6416 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6417 ^ 
    6418 program_source:333:26: note: expanded from macro '\ 
    6419 DECLARE_I_REDUCTION_NON_UNIFORM' 
    6420 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6421 ^ 
    6422 :69:1: note: expanded from here 
    6423 simd_max 
    6424 ^ 
    6425 program_source:349:1: error: implicit declaration of function 'simd_max' is
    6426 invalid in OpenCL 
    6427 program_source:338:60: note: expanded from macro
    6428 'DECLARE_REDUCTION_NON_UNIFORM' 
    6429 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6430 ^ 
    6431 program_source:333:26: note: expanded from macro '\ 
    6432 DECLARE_I_REDUCTION_NON_UNIFORM' 
    6433 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6434 ^ 
    6435 :69:1: note: expanded from here 
    6436 simd_max 
    6437 ^ 
    6438 program_source:349:1: error: implicit declaration of function 'simd_max' is
    6439 invalid in OpenCL 
    6440 program_source:339:54: note: expanded from macro
    6441 'DECLARE_REDUCTION_NON_UNIFORM' 
    6442 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6443 ^ 
    6444 program_source:336:26: note: expanded from macro '\ 
    6445 DECLARE_F_REDUCTION_NON_UNIFORM' 
    6446 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6447 ^ 
    6448 :71:1: note: expanded from here 
    6449 simd_max 
    6450 ^ 
    6451 program_source:351:1: error: implicit declaration of function 'simd_product'
    6452 is invalid in OpenCL 
    6453 DECLARE_REDUCTION_NON_UNIFORM(product, reduce_mul) 
    6454 ^ 
    6455 program_source:338:60: note: expanded from macro
    6456 'DECLARE_REDUCTION_NON_UNIFORM' 
    6457 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6458 ^ 
    6459 program_source:333:26: note: expanded from macro '\ 
    6460 DECLARE_I_REDUCTION_NON_UNIFORM' 
    6461 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6462 ^ 
    6463 :73:1: note: expanded from here 
    6464 simd_product 
    6465 ^ 
    6466 program_source:351:1: error: implicit declaration of function 'simd_product'
    6467 is invalid in OpenCL 
    6468 program_source:338:60: note: expanded from macro
    6469 'DECLARE_REDUCTION_NON_UNIFORM' 
    6470 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6471 ^ 
    6472 program_source:333:26: note: expanded from macro '\ 
    6473 DECLARE_I_REDUCTION_NON_UNIFORM' 
    6474 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6475 ^ 
    6476 :73:1: note: expanded from here 
    6477 simd_product 
    6478 ^ 
    6479 program_source:351:1: error: implicit declaration of function 'simd_product'
    6480 is invalid in OpenCL 
    6481 program_source:339:54: note: expanded from macro
    6482 'DECLARE_REDUCTION_NON_UNIFORM' 
    6483 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6484 ^ 
    6485 program_source:336:26: note: expanded from macro '\ 
    6486 DECLARE_F_REDUCTION_NON_UNIFORM' 
    6487 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6488 ^ 
    6489 :75:1: note: expanded from here 
    6490 simd_product 
    6491 ^ 
    6492 program_source:352:1: error: implicit declaration of function
    6493 'simd_prefix_inclusive_product' is invalid in OpenCL 
    6494 DECLARE_REDUCTION_NON_UNIFORM(prefix_inclusive_product, scan_inclusive_mul) 
    6495 ^ 
    6496 program_source:338:60: note: expanded from macro
    6497 'DECLARE_REDUCTION_NON_UNIFORM' 
    6498 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6499 ^ 
    6500 program_source:333:26: note: expanded from macro '\ 
    6501 DECLARE_I_REDUCTION_NON_UNIFORM' 
    6502 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6503 ^ 
    6504 :77:1: note: expanded from here 
    6505 simd_prefix_inclusive_product 
    6506 ^ 
    6507 program_source:352:1: error: implicit declaration of function
    6508 'simd_prefix_inclusive_product' is invalid in OpenCL 
    6509 program_source:338:60: note: expanded from macro
    6510 'DECLARE_REDUCTION_NON_UNIFORM' 
    6511 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6512 ^ 
    6513 program_source:333:26: note: expanded from macro '\ 
    6514 DECLARE_I_REDUCTION_NON_UNIFORM' 
    6515 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6516 ^ 
    6517 :77:1: note: expanded from here 
    6518 simd_prefix_inclusive_product 
    6519 ^ 
    6520 program_source:352:1: error: implicit declaration of function
    6521 'simd_prefix_inclusive_product' is invalid in OpenCL 
    6522 program_source:339:54: note: expanded from macro
    6523 'DECLARE_REDUCTION_NON_UNIFORM' 
    6524 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6525 ^ 
    6526 program_source:336:26: note: expanded from macro '\ 
    6527 DECLARE_F_REDUCTION_NON_UNIFORM' 
    6528 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6529 ^ 
    6530 :79:1: note: expanded from here 
    6531 simd_prefix_inclusive_product 
    6532 ^ 
    6533 program_source:353:1: error: implicit declaration of function
    6534 'simd_prefix_exclusive_product' is invalid in OpenCL 
    6535 DECLARE_REDUCTION_NON_UNIFORM(prefix_exclusive_product, scan_exclusive_mul) 
    6536 ^ 
    6537 program_source:338:60: note: expanded from macro
    6538 'DECLARE_REDUCTION_NON_UNIFORM' 
    6539 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6540 ^ 
    6541 program_source:333:26: note: expanded from macro '\ 
    6542 DECLARE_I_REDUCTION_NON_UNIFORM' 
    6543 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6544 ^ 
    6545 :81:1: note: expanded from here 
    6546 simd_prefix_exclusive_product 
    6547 ^ 
    6548 program_source:353:1: error: implicit declaration of function
    6549 'simd_prefix_exclusive_product' is invalid in OpenCL 
    6550 program_source:338:60: note: expanded from macro
    6551 'DECLARE_REDUCTION_NON_UNIFORM' 
    6552 #define DECLARE_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6553 ^ 
    6554 program_source:333:26: note: expanded from macro '\ 
    6555 DECLARE_I_REDUCTION_NON_UNIFORM' 
    6556 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6557 ^ 
    6558 :81:1: note: expanded from here 
    6559 simd_prefix_exclusive_product 
    6560 ^ 
    6561 program_source:353:1: error: implicit declaration of function
    6562 'simd_prefix_exclusive_product' is invalid in OpenCL 
    6563 program_source:339:54: note: expanded from macro
    6564 'DECLARE_REDUCTION_NON_UNIFORM' 
    6565 DECLARE_I_REDUCTION_NON_UNIFORM(METAL_OP, OPENCL_OP) \ 
    6566 ^ 
    6567 program_source:336:26: note: expanded from macro '\ 
    6568 DECLARE_F_REDUCTION_NON_UNIFORM' 
    6569 BRIDGE_FUNCTION_F_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6570 ^ 
    6571 :83:1: note: expanded from here 
    6572 simd_prefix_exclusive_product 
    6573 ^ 
    6574 program_source:354:1: error: implicit declaration of function 'simd_and' is
    6575 invalid in OpenCL 
    6576 DECLARE_I_REDUCTION_NON_UNIFORM(and, reduce_and) 
    6577 ^ 
    6578 program_source:333:26: note: expanded from macro
    6579 'DECLARE_I_REDUCTION_NON_UNIFORM' 
    6580 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6581 ^ 
    6582 :85:1: note: expanded from here 
    6583 simd_and 
    6584 ^ 
    6585 program_source:354:1: error: implicit declaration of function 'simd_and' is
    6586 invalid in OpenCL 
    6587 program_source:333:26: note: expanded from macro
    6588 'DECLARE_I_REDUCTION_NON_UNIFORM' 
    6589 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6590 ^ 
    6591 :85:1: note: expanded from here 
    6592 simd_and 
    6593 ^ 
    6594 program_source:355:1: error: implicit declaration of function 'simd_or' is
    6595 invalid in OpenCL 
    6596 DECLARE_I_REDUCTION_NON_UNIFORM(or, reduce_or) 
    6597 ^ 
    6598 program_source:333:26: note: expanded from macro
    6599 'DECLARE_I_REDUCTION_NON_UNIFORM' 
    6600 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6601 ^ 
    6602 :87:1: note: expanded from here 
    6603 simd_or 
    6604 ^ 
    6605 program_source:355:1: error: implicit declaration of function 'simd_or' is
    6606 invalid in OpenCL 
    6607 program_source:333:26: note: expanded from macro
    6608 'DECLARE_I_REDUCTION_NON_UNIFORM' 
    6609 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6610 ^ 
    6611 :87:1: note: expanded from here 
    6612 simd_or 
    6613 ^ 
    6614 program_source:356:1: error: implicit declaration of function 'simd_xor' is
    6615 invalid in OpenCL 
    6616 DECLARE_I_REDUCTION_NON_UNIFORM(xor, reduce_xor) 
    6617 ^ 
    6618 program_source:333:26: note: expanded from macro
    6619 'DECLARE_I_REDUCTION_NON_UNIFORM' 
    6620 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6621 ^ 
    6622 :89:1: note: expanded from here 
    6623 simd_xor 
    6624 ^ 
    6625 program_source:356:1: error: implicit declaration of function 'simd_xor' is
    6626 invalid in OpenCL 
    6627 program_source:333:26: note: expanded from macro
    6628 'DECLARE_I_REDUCTION_NON_UNIFORM' 
    6629 BRIDGE_FUNCTION_I_ARGS_1(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6630 ^ 
    6631 :89:1: note: expanded from here 
    6632 simd_xor 
    6633 ^ 
    6634 program_source:358:1: error: implicit declaration of function 'simd_broadcast'
    6635 is invalid in OpenCL 
    6636 DECLARE_SHUFFLE_NON_UNIFORM(broadcast, broadcast) 
    6637 ^ 
    6638 program_source:343:24: note: expanded from macro 'DECLARE_SHUFFLE_NON_UNIFORM' 
    6639 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6640 ^ 
    6641 :91:1: note: expanded from here 
    6642 simd_broadcast 
    6643 ^ 
    6644 program_source:358:1: error: implicit declaration of function 'simd_broadcast'
    6645 is invalid in OpenCL 
    6646 program_source:343:24: note: expanded from macro 'DECLARE_SHUFFLE_NON_UNIFORM' 
    6647 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6648 ^ 
    6649 :91:1: note: expanded from here 
    6650 simd_broadcast 
    6651 ^ 
    6652 program_source:358:1: error: implicit declaration of function 'simd_broadcast'
    6653 is invalid in OpenCL 
    6654 program_source:343:24: note: expanded from macro 'DECLARE_SHUFFLE_NON_UNIFORM' 
    6655 BRIDGE_FUNCTION_ARGS_2(simd_##METAL_OP, sub_group_non_uniform_##OPENCL_OP) \ 
    6656 ^ 
    6657 :91:1: note: expanded from here 
    6658 simd_broadcast 
    6659 ^ 
    6660 program_source:391:1: error: implicit declaration of function 'quad_sum' is
    6661 invalid in OpenCL 
    6662 DECLARE_REDUCTION_CLUSTERED(sum, reduce_add) 
    6663 ^ 
    6664 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    6665 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    6666 ^ 
    6667 program_source:360:60: note: expanded from macro '\ 
    6668 DECLARE_I_REDUCTION_CLUSTERED' 
    6669 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    6670 ^ 
    6671 program_source:245:71: note: expanded from macro '\ 
    6672 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    6673 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    6674 ^ 
    6675 program_source:196:12: note: expanded from macro '\ 
    6676 OVERLOAD_CLUSTERED_ARGS_1' 
    6677 return quad_##METAL_SUFFIX(x); \ 
    6678 ^ 
    6679 :94:1: note: expanded from here 
    6680 quad_sum 
    6681 ^ 
    6682 program_source:391:1: error: implicit declaration of function 'simd_sum' is
    6683 invalid in OpenCL 
    6684 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    6685 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    6686 ^ 
    6687 program_source:360:60: note: expanded from macro '\ 
    6688 DECLARE_I_REDUCTION_CLUSTERED' 
    6689 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    6690 ^ 
    6691 program_source:245:71: note: expanded from macro '\ 
    6692 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    6693 #define BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1(METAL_SUFFIX, OPENCL_EXPR) \ 
    6694 ^ 
    6695 program_source:198:12: note: expanded from macro '\ 
    6696 OVERLOAD_CLUSTERED_ARGS_1' 
    6697 return simd_##METAL_SUFFIX(x); \ 
    6698 ^ 
    6699 :95:1: note: expanded from here 
    6700 simd_sum 
    6701 ^ 
    6702 program_source:391:1: error: implicit declaration of function 'quad_sum' is
    6703 invalid in OpenCL 
    6704 program_source:384:58: note: expanded from macro 'DECLARE_REDUCTION_CLUSTERED' 
    6705 #define DECLARE_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    6706 ^ 
    6707 program_source:360:60: note: expanded from macro '\ 
    6708 DECLARE_I_REDUCTION_CLUSTERED' 
    6709 #define DECLARE_I_REDUCTION_CLUSTERED(METAL_OP, OPENCL_OP) \ 
    6710 ^ 
    6711 program_source:246:59: note: expanded from macro '\ 
    6712 BRIDGE_FUNCTION_CLUSTERED_I_ARGS_1' 
    6713 OVERLOAD_CLUSTERED_ARGS_1(METAL_SUFFIX, OPENCL_EXPR, int) \ 
    6714 ^ 
     1511[deleted to fit within ticket limits]
     1512
    67151513program_source:196:12: note: expanded from macro '\ 
    67161514OVERLOAD_CLUSTERED_ARGS_1'