Changes between Initial Version and Version 1 of Ticket #19942


Ignore:
Timestamp:
Mar 6, 2026, 9:03:03 AM (22 hours ago)
Author:
Eric Pettersen
Comment:

Legend:

Unmodified
Added
Removed
Modified

  • Ticket #19942

    • Property Component Unassigned → Third Party
    • Property Owner set to Tristan Croll
    • Property Platform → all
    • Property Project → ChimeraX
    • Property Status new → assigned
    • Property Summary ChimeraX bug report submission → isolde sim start: various OpenMM/OpenCL errors on Tahoe

  • Ticket #19942 – Description

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