| // RUN: %check_clang_tidy -check-suffix=OLDCLOLDAOC %s altera-single-work-item-barrier %t -- -header-filter=.* "--" -cl-std=CL1.2 -c -DOLDCLOLDAOC |
| // RUN: %check_clang_tidy -check-suffix=NEWCLOLDAOC %s altera-single-work-item-barrier %t -- -header-filter=.* "--" -cl-std=CL2.0 -c -DNEWCLOLDAOC |
| // RUN: %check_clang_tidy -check-suffix=OLDCLNEWAOC %s altera-single-work-item-barrier %t -- -config='{CheckOptions: { altera-single-work-item-barrier.AOCVersion: 1701}}' -header-filter=.* "--" -cl-std=CL1.2 -c -DOLDCLNEWAOC |
| // RUN: %check_clang_tidy -check-suffix=NEWCLNEWAOC %s altera-single-work-item-barrier %t -- -config='{CheckOptions: { altera-single-work-item-barrier.AOCVersion: 1701}}' -header-filter=.* "--" -cl-std=CL2.0 -c -DNEWCLNEWAOC |
| |
| #ifdef OLDCLOLDAOC // OpenCL 1.2 Altera Offline Compiler < 17.1 |
| void __kernel error_barrier_no_id(__global int * foo, int size) { |
| // CHECK-MESSAGES-OLDCLOLDAOC: :[[@LINE-1]]:15: warning: kernel function 'error_barrier_no_id' does not call 'get_global_id' or 'get_local_id' and will be treated as a single work-item [altera-single-work-item-barrier] |
| for (int j = 0; j < 256; j++) { |
| for (int i = 256; i < size; i+= 256) { |
| foo[j] += foo[j+i]; |
| } |
| } |
| barrier(CLK_GLOBAL_MEM_FENCE); |
| // CHECK-MESSAGES-OLDCLOLDAOC: :[[@LINE-1]]:3: note: barrier call is in a single work-item and may error out |
| for (int i = 1; i < 256; i++) { |
| foo[0] += foo[i]; |
| } |
| } |
| |
| void __kernel success_barrier_global_id(__global int * foo, int size) { |
| barrier(CLK_GLOBAL_MEM_FENCE); |
| int tid = get_global_id(0); |
| } |
| |
| void __kernel success_barrier_local_id(__global int * foo, int size) { |
| barrier(CLK_GLOBAL_MEM_FENCE); |
| int tid = get_local_id(0); |
| } |
| |
| void __kernel success_barrier_both_ids(__global int * foo, int size) { |
| barrier(CLK_GLOBAL_MEM_FENCE); |
| int gid = get_global_id(0); |
| int lid = get_local_id(0); |
| } |
| |
| void success_nokernel_barrier_no_id(__global int * foo, int size) { |
| for (int j = 0; j < 256; j++) { |
| for (int i = 256; i < size; i+= 256) { |
| foo[j] += foo[j+i]; |
| } |
| } |
| barrier(CLK_GLOBAL_MEM_FENCE); |
| for (int i = 1; i < 256; i++) { |
| foo[0] += foo[i]; |
| } |
| } |
| |
| void success_nokernel_barrier_global_id(__global int * foo, int size) { |
| barrier(CLK_GLOBAL_MEM_FENCE); |
| int tid = get_global_id(0); |
| } |
| |
| void success_nokernel_barrier_local_id(__global int * foo, int size) { |
| barrier(CLK_GLOBAL_MEM_FENCE); |
| int tid = get_local_id(0); |
| } |
| |
| void success_nokernel_barrier_both_ids(__global int * foo, int size) { |
| barrier(CLK_GLOBAL_MEM_FENCE); |
| int gid = get_global_id(0); |
| int lid = get_local_id(0); |
| } |
| #endif |
| |
| #ifdef NEWCLOLDAOC // OpenCL 2.0 Altera Offline Compiler < 17.1 |
| void __kernel error_barrier_no_id(__global int * foo, int size) { |
| // CHECK-MESSAGES-NEWCLOLDAOC: :[[@LINE-1]]:15: warning: kernel function 'error_barrier_no_id' does not call 'get_global_id' or 'get_local_id' and will be treated as a single work-item [altera-single-work-item-barrier] |
| for (int j = 0; j < 256; j++) { |
| for (int i = 256; i < size; i+= 256) { |
| foo[j] += foo[j+i]; |
| } |
| } |
| work_group_barrier(CLK_GLOBAL_MEM_FENCE); |
| // CHECK-MESSAGES-NEWCLOLDAOC: :[[@LINE-1]]:3: note: barrier call is in a single work-item and may error out |
| for (int i = 1; i < 256; i++) { |
| foo[0] += foo[i]; |
| } |
| } |
| |
| void __kernel success_barrier_global_id(__global int * foo, int size) { |
| work_group_barrier(CLK_GLOBAL_MEM_FENCE); |
| int tid = get_global_id(0); |
| } |
| |
| void __kernel success_barrier_local_id(__global int * foo, int size) { |
| work_group_barrier(CLK_GLOBAL_MEM_FENCE); |
| int tid = get_local_id(0); |
| } |
| |
| void __kernel success_barrier_both_ids(__global int * foo, int size) { |
| work_group_barrier(CLK_GLOBAL_MEM_FENCE); |
| int gid = get_global_id(0); |
| int lid = get_local_id(0); |
| } |
| |
| void success_nokernel_barrier_no_id(__global int * foo, int size) { |
| for (int j = 0; j < 256; j++) { |
| for (int i = 256; i < size; i+= 256) { |
| foo[j] += foo[j+i]; |
| } |
| } |
| work_group_barrier(CLK_GLOBAL_MEM_FENCE); |
| for (int i = 1; i < 256; i++) { |
| foo[0] += foo[i]; |
| } |
| } |
| |
| void success_nokernel_barrier_global_id(__global int * foo, int size) { |
| work_group_barrier(CLK_GLOBAL_MEM_FENCE); |
| int tid = get_global_id(0); |
| } |
| |
| void success_nokernel_barrier_local_id(__global int * foo, int size) { |
| work_group_barrier(CLK_GLOBAL_MEM_FENCE); |
| int tid = get_local_id(0); |
| } |
| |
| void success_nokernel_barrier_both_ids(__global int * foo, int size) { |
| work_group_barrier(CLK_GLOBAL_MEM_FENCE); |
| int gid = get_global_id(0); |
| int lid = get_local_id(0); |
| } |
| #endif |
| |
| #ifdef OLDCLNEWAOC // OpenCL 1.2 Altera Offline Compiler >= 17.1 |
| void __kernel error_barrier_no_id(__global int * foo, int size) { |
| // CHECK-MESSAGES-OLDCLNEWAOC: :[[@LINE-1]]:15: warning: kernel function 'error_barrier_no_id' does not call an ID function and may be a viable single work-item, but will be forced to execute as an NDRange [altera-single-work-item-barrier] |
| for (int j = 0; j < 256; j++) { |
| for (int i = 256; i < size; i+= 256) { |
| foo[j] += foo[j+i]; |
| } |
| } |
| barrier(CLK_GLOBAL_MEM_FENCE); |
| // CHECK-MESSAGES-OLDCLNEWAOC: :[[@LINE-1]]:3: note: barrier call will force NDRange execution; if single work-item semantics are desired a mem_fence may be more efficient |
| for (int i = 1; i < 256; i++) { |
| foo[0] += foo[i]; |
| } |
| } |
| |
| __attribute__ ((reqd_work_group_size(1,1,1))) |
| void __kernel error_barrier_no_id_work_group_size(__global int * foo, int size) { |
| // CHECK-MESSAGES-OLDCLNEWAOC: :[[@LINE-1]]:15: warning: kernel function 'error_barrier_no_id_work_group_size' does not call an ID function and may be a viable single work-item, but will be forced to execute as an NDRange [altera-single-work-item-barrier] |
| for (int j = 0; j < 256; j++) { |
| for (int i = 256; i < size; i+= 256) { |
| foo[j] += foo[j+i]; |
| } |
| } |
| barrier(CLK_GLOBAL_MEM_FENCE); |
| // CHECK-MESSAGES-OLDCLNEWAOC: :[[@LINE-1]]:3: note: barrier call will force NDRange execution; if single work-item semantics are desired a mem_fence may be more efficient |
| for (int i = 1; i < 256; i++) { |
| foo[0] += foo[i]; |
| } |
| } |
| |
| __attribute__ ((reqd_work_group_size(2,1,1))) |
| void __kernel success_barrier_no_id_work_group_size(__global int * foo, int size) { |
| for (int j = 0; j < 256; j++) { |
| for (int i = 256; i < size; i+= 256) { |
| foo[j] += foo[j+i]; |
| } |
| } |
| barrier(CLK_GLOBAL_MEM_FENCE); |
| for (int i = 1; i < 256; i++) { |
| foo[0] += foo[i]; |
| } |
| } |
| |
| void __kernel success_barrier_global_id(__global int * foo, int size) { |
| barrier(CLK_GLOBAL_MEM_FENCE); |
| int tid = get_global_id(0); |
| } |
| |
| void __kernel success_barrier_local_id(__global int * foo, int size) { |
| barrier(CLK_GLOBAL_MEM_FENCE); |
| int tid = get_local_id(0); |
| } |
| |
| void __kernel success_barrier_both_ids(__global int * foo, int size) { |
| barrier(CLK_GLOBAL_MEM_FENCE); |
| int gid = get_global_id(0); |
| int lid = get_local_id(0); |
| } |
| |
| void success_nokernel_barrier_no_id(__global int * foo, int size) { |
| for (int j = 0; j < 256; j++) { |
| for (int i = 256; i < size; i+= 256) { |
| foo[j] += foo[j+i]; |
| } |
| } |
| barrier(CLK_GLOBAL_MEM_FENCE); |
| for (int i = 1; i < 256; i++) { |
| foo[0] += foo[i]; |
| } |
| } |
| |
| void success_nokernel_barrier_global_id(__global int * foo, int size) { |
| barrier(CLK_GLOBAL_MEM_FENCE); |
| int tid = get_global_id(0); |
| } |
| |
| void success_nokernel_barrier_local_id(__global int * foo, int size) { |
| barrier(CLK_GLOBAL_MEM_FENCE); |
| int tid = get_local_id(0); |
| } |
| |
| void success_nokernel_barrier_both_ids(__global int * foo, int size) { |
| barrier(CLK_GLOBAL_MEM_FENCE); |
| int gid = get_global_id(0); |
| int lid = get_local_id(0); |
| } |
| #endif |
| |
| #ifdef NEWCLNEWAOC // OpenCL 2.0 Altera Offline Compiler >= 17.1 |
| void __kernel error_barrier_no_id(__global int * foo, int size) { |
| // CHECK-MESSAGES-NEWCLNEWAOC: :[[@LINE-1]]:15: warning: kernel function 'error_barrier_no_id' does not call an ID function and may be a viable single work-item, but will be forced to execute as an NDRange [altera-single-work-item-barrier] |
| for (int j = 0; j < 256; j++) { |
| for (int i = 256; i < size; i+= 256) { |
| foo[j] += foo[j+i]; |
| } |
| } |
| work_group_barrier(CLK_GLOBAL_MEM_FENCE); |
| // CHECK-MESSAGES-NEWCLNEWAOC: :[[@LINE-1]]:3: note: barrier call will force NDRange execution; if single work-item semantics are desired a mem_fence may be more efficient |
| for (int i = 1; i < 256; i++) { |
| foo[0] += foo[i]; |
| } |
| } |
| |
| __attribute__ ((reqd_work_group_size(1,1,1))) |
| void __kernel error_barrier_no_id_work_group_size(__global int * foo, int size) { |
| // CHECK-MESSAGES-NEWCLNEWAOC: :[[@LINE-1]]:15: warning: kernel function 'error_barrier_no_id_work_group_size' does not call an ID function and may be a viable single work-item, but will be forced to execute as an NDRange [altera-single-work-item-barrier] |
| for (int j = 0; j < 256; j++) { |
| for (int i = 256; i < size; i+= 256) { |
| foo[j] += foo[j+i]; |
| } |
| } |
| work_group_barrier(CLK_GLOBAL_MEM_FENCE); |
| // CHECK-MESSAGES-NEWCLNEWAOC: :[[@LINE-1]]:3: note: barrier call will force NDRange execution; if single work-item semantics are desired a mem_fence may be more efficient |
| for (int i = 1; i < 256; i++) { |
| foo[0] += foo[i]; |
| } |
| } |
| |
| __attribute__ ((reqd_work_group_size(2,1,1))) |
| void __kernel success_barrier_no_id_work_group_size(__global int * foo, int size) { |
| for (int j = 0; j < 256; j++) { |
| for (int i = 256; i < size; i+= 256) { |
| foo[j] += foo[j+i]; |
| } |
| } |
| work_group_barrier(CLK_GLOBAL_MEM_FENCE); |
| for (int i = 1; i < 256; i++) { |
| foo[0] += foo[i]; |
| } |
| } |
| |
| void __kernel success_barrier_global_id(__global int * foo, int size) { |
| work_group_barrier(CLK_GLOBAL_MEM_FENCE); |
| int tid = get_global_id(0); |
| } |
| |
| void __kernel success_barrier_local_id(__global int * foo, int size) { |
| work_group_barrier(CLK_GLOBAL_MEM_FENCE); |
| int tid = get_local_id(0); |
| } |
| |
| void __kernel success_barrier_both_ids(__global int * foo, int size) { |
| work_group_barrier(CLK_GLOBAL_MEM_FENCE); |
| int gid = get_global_id(0); |
| int lid = get_local_id(0); |
| } |
| |
| void success_nokernel_barrier_no_id(__global int * foo, int size) { |
| for (int j = 0; j < 256; j++) { |
| for (int i = 256; i < size; i+= 256) { |
| foo[j] += foo[j+i]; |
| } |
| } |
| work_group_barrier(CLK_GLOBAL_MEM_FENCE); |
| for (int i = 1; i < 256; i++) { |
| foo[0] += foo[i]; |
| } |
| } |
| |
| void success_nokernel_barrier_global_id(__global int * foo, int size) { |
| work_group_barrier(CLK_GLOBAL_MEM_FENCE); |
| int tid = get_global_id(0); |
| } |
| |
| void success_nokernel_barrier_local_id(__global int * foo, int size) { |
| work_group_barrier(CLK_GLOBAL_MEM_FENCE); |
| int tid = get_local_id(0); |
| } |
| |
| void success_nokernel_barrier_both_ids(__global int * foo, int size) { |
| work_group_barrier(CLK_GLOBAL_MEM_FENCE); |
| int gid = get_global_id(0); |
| int lid = get_local_id(0); |
| } |
| #endif |