| SKIP: FAILED |
| |
| #include <metal_stdlib> |
| using namespace metal; |
| |
| threadgroup int a; |
| threadgroup int b; |
| threadgroup int c; |
| void uses_a() { |
| a = (a + 1); |
| } |
| void uses_b() { |
| b = (b * 2); |
| } |
| void uses_a_and_b() { |
| b = a; |
| } |
| void no_uses() { |
| } |
| void outer() { |
| a = 0; |
| uses_a(); |
| uses_a_and_b(); |
| uses_b(); |
| no_uses(); |
| } |
| kernel void main1(uint tint_local_index [[thread_index_in_threadgroup]]) { |
| if ((tint_local_index == 0u)) { |
| a = 0; |
| } |
| threadgroup_barrier(mem_flags::mem_threadgroup); |
| a = 42; |
| uses_a(); |
| } |
| kernel void main2(uint tint_local_index [[thread_index_in_threadgroup]]) { |
| if ((tint_local_index == 0u)) { |
| b = 0; |
| } |
| threadgroup_barrier(mem_flags::mem_threadgroup); |
| b = 7; |
| uses_b(); |
| } |
| kernel void main3(uint tint_local_index [[thread_index_in_threadgroup]]) { |
| if ((tint_local_index == 0u)) { |
| a = 0; |
| b = 0; |
| } |
| threadgroup_barrier(mem_flags::mem_threadgroup); |
| outer(); |
| no_uses(); |
| } |
| kernel void main4() { |
| no_uses(); |
| } |
| program_source:4:17: error: program scope variable must reside in constant address space |
| threadgroup int a; |
| ^ |
| program_source:5:17: error: program scope variable must reside in constant address space |
| threadgroup int b; |
| ^ |
| program_source:6:17: error: program scope variable must reside in constant address space |
| threadgroup int c; |
| ^ |
| program_source:26:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality] |
| if ((tint_local_index == 0u)) { |
| ~~~~~~~~~~~~~~~~~^~~~~ |
| program_source:26:25: note: remove extraneous parentheses around the comparison to silence this warning |
| if ((tint_local_index == 0u)) { |
| ~ ^ ~ |
| program_source:26:25: note: use '=' to turn this equality comparison into an assignment |
| if ((tint_local_index == 0u)) { |
| ^~ |
| = |
| program_source:34:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality] |
| if ((tint_local_index == 0u)) { |
| ~~~~~~~~~~~~~~~~~^~~~~ |
| program_source:34:25: note: remove extraneous parentheses around the comparison to silence this warning |
| if ((tint_local_index == 0u)) { |
| ~ ^ ~ |
| program_source:34:25: note: use '=' to turn this equality comparison into an assignment |
| if ((tint_local_index == 0u)) { |
| ^~ |
| = |
| program_source:42:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality] |
| if ((tint_local_index == 0u)) { |
| ~~~~~~~~~~~~~~~~~^~~~~ |
| program_source:42:25: note: remove extraneous parentheses around the comparison to silence this warning |
| if ((tint_local_index == 0u)) { |
| ~ ^ ~ |
| program_source:42:25: note: use '=' to turn this equality comparison into an assignment |
| if ((tint_local_index == 0u)) { |
| ^~ |
| = |
| |
