| alan-baker | 227e978 | 2020-06-02 15:35:37 -0400 | [diff] [blame] | 1 | // RUN: clspv %s -o %t.spv -enable-pre=1 -enable-load-pre=1 |
| Kévin Petit | 23d5f18 | 2019-08-13 16:21:29 +0100 | [diff] [blame] | 2 | // RUN: spirv-dis -o %t2.spvasm %t.spv |
| 3 | // RUN: FileCheck %s < %t2.spvasm |
| 4 | // RUN: spirv-val --target-env vulkan1.0 %t.spv |
| 5 | |
| 6 | // CHECK: OpUndef |
| 7 | // CHECK-NOT: OpVariable {{.*}} Workgroup % |
| 8 | // CHECK: OpVariable {{.*}} Workgroup |
| 9 | // CHECK-NOT: OpVariable {{.*}} Workgroup % |
| 10 | // CHECK: OpVariable {{.*}} Workgroup |
| 11 | |
| 12 | __kernel void |
| 13 | top_scan(__global uint * isums, |
| 14 | const int n, |
| 15 | __local uint * lmem) |
| 16 | { |
| 17 | __local int s_seed; |
| 18 | s_seed = 0; barrier(CLK_LOCAL_MEM_FENCE); |
| 19 | |
| 20 | int last_thread = (get_local_id(0) < n && |
| 21 | (get_local_id(0)+1) == n) ? 1 : 0; |
| 22 | |
| 23 | for (int d = 0; d < 16; d++) |
| 24 | { |
| 25 | int idx = get_local_id(0); |
| 26 | lmem[idx] = 0; |
| 27 | if (last_thread) |
| 28 | { |
| 29 | s_seed += 42; |
| 30 | } |
| 31 | barrier(CLK_LOCAL_MEM_FENCE); |
| 32 | } |
| 33 | } |
| 34 | |