maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity
llvm-project/clang-tools-extra/test/clang-tidy/checkers/altera-single-work-item-bar...

301 lines
10 KiB
C++
Raw Blame History

// RUN: %check_clang_tidy -check-suffix=OLDCLOLDAOC %s altera-single-work-item-barrier %t -- -header-filter=.* "--" -cl-std=CL1.2 -c --include opencl-c.h -DOLDCLOLDAOC
// RUN: %check_clang_tidy -check-suffix=NEWCLOLDAOC %s altera-single-work-item-barrier %t -- -header-filter=.* "--" -cl-std=CL2.0 -c --include opencl-c.h -DNEWCLOLDAOC
// RUN: %check_clang_tidy -check-suffix=OLDCLNEWAOC %s altera-single-work-item-barrier %t -- -config='{CheckOptions: [{key: altera-single-work-item-barrier.AOCVersion, value: 1701}]}' -header-filter=.* "--" -cl-std=CL1.2 -c --include opencl-c.h -DOLDCLNEWAOC
// RUN: %check_clang_tidy -check-suffix=NEWCLNEWAOC %s altera-single-work-item-barrier %t -- -config='{CheckOptions: [{key: altera-single-work-item-barrier.AOCVersion, value: 1701}]}' -header-filter=.* "--" -cl-std=CL2.0 -c --include opencl-c.h -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
Reference in New Issue View Git Blame Copy Permalink