Changes between Initial Version and Version 1 of Ticket #20472


Ignore:
Timestamp:
Jun 8, 2026, 2:22:18 PM (4 hours ago)
Author:
Eric Pettersen
Comment:

Reported by Orhi Esarte

Legend:

Unmodified
Added
Removed
Modified

  • Ticket #20472

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

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