301 lines · cpp
1// RUN: %check_clang_tidy -check-suffix=OLDCLOLDAOC %s altera-single-work-item-barrier %t -- -header-filter=.* "--" -cl-std=CL1.2 -c -DOLDCLOLDAOC2// RUN: %check_clang_tidy -check-suffix=NEWCLOLDAOC %s altera-single-work-item-barrier %t -- -header-filter=.* "--" -cl-std=CL2.0 -c -DNEWCLOLDAOC3// 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 -DOLDCLNEWAOC4// 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 -DNEWCLNEWAOC5 6#ifdef OLDCLOLDAOC // OpenCL 1.2 Altera Offline Compiler < 17.17void __kernel error_barrier_no_id(__global int * foo, int size) {8 // 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]9 for (int j = 0; j < 256; j++) {10 for (int i = 256; i < size; i+= 256) {11 foo[j] += foo[j+i];12 }13 }14 barrier(CLK_GLOBAL_MEM_FENCE);15 // CHECK-MESSAGES-OLDCLOLDAOC: :[[@LINE-1]]:3: note: barrier call is in a single work-item and may error out16 for (int i = 1; i < 256; i++) {17 foo[0] += foo[i];18 }19}20 21void __kernel success_barrier_global_id(__global int * foo, int size) {22 barrier(CLK_GLOBAL_MEM_FENCE);23 int tid = get_global_id(0);24}25 26void __kernel success_barrier_local_id(__global int * foo, int size) {27 barrier(CLK_GLOBAL_MEM_FENCE);28 int tid = get_local_id(0);29}30 31void __kernel success_barrier_both_ids(__global int * foo, int size) {32 barrier(CLK_GLOBAL_MEM_FENCE);33 int gid = get_global_id(0);34 int lid = get_local_id(0);35}36 37void success_nokernel_barrier_no_id(__global int * foo, int size) {38 for (int j = 0; j < 256; j++) {39 for (int i = 256; i < size; i+= 256) {40 foo[j] += foo[j+i];41 }42 }43 barrier(CLK_GLOBAL_MEM_FENCE);44 for (int i = 1; i < 256; i++) {45 foo[0] += foo[i];46 }47}48 49void success_nokernel_barrier_global_id(__global int * foo, int size) {50 barrier(CLK_GLOBAL_MEM_FENCE);51 int tid = get_global_id(0);52}53 54void success_nokernel_barrier_local_id(__global int * foo, int size) {55 barrier(CLK_GLOBAL_MEM_FENCE);56 int tid = get_local_id(0);57}58 59void success_nokernel_barrier_both_ids(__global int * foo, int size) {60 barrier(CLK_GLOBAL_MEM_FENCE);61 int gid = get_global_id(0);62 int lid = get_local_id(0);63}64#endif65 66#ifdef NEWCLOLDAOC // OpenCL 2.0 Altera Offline Compiler < 17.167void __kernel error_barrier_no_id(__global int * foo, int size) {68 // 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]69 for (int j = 0; j < 256; j++) {70 for (int i = 256; i < size; i+= 256) {71 foo[j] += foo[j+i];72 }73 }74 work_group_barrier(CLK_GLOBAL_MEM_FENCE);75 // CHECK-MESSAGES-NEWCLOLDAOC: :[[@LINE-1]]:3: note: barrier call is in a single work-item and may error out76 for (int i = 1; i < 256; i++) {77 foo[0] += foo[i];78 }79}80 81void __kernel success_barrier_global_id(__global int * foo, int size) {82 work_group_barrier(CLK_GLOBAL_MEM_FENCE);83 int tid = get_global_id(0);84}85 86void __kernel success_barrier_local_id(__global int * foo, int size) {87 work_group_barrier(CLK_GLOBAL_MEM_FENCE);88 int tid = get_local_id(0);89}90 91void __kernel success_barrier_both_ids(__global int * foo, int size) {92 work_group_barrier(CLK_GLOBAL_MEM_FENCE);93 int gid = get_global_id(0);94 int lid = get_local_id(0);95}96 97void success_nokernel_barrier_no_id(__global int * foo, int size) {98 for (int j = 0; j < 256; j++) {99 for (int i = 256; i < size; i+= 256) {100 foo[j] += foo[j+i];101 }102 }103 work_group_barrier(CLK_GLOBAL_MEM_FENCE);104 for (int i = 1; i < 256; i++) {105 foo[0] += foo[i];106 }107}108 109void success_nokernel_barrier_global_id(__global int * foo, int size) {110 work_group_barrier(CLK_GLOBAL_MEM_FENCE);111 int tid = get_global_id(0);112}113 114void success_nokernel_barrier_local_id(__global int * foo, int size) {115 work_group_barrier(CLK_GLOBAL_MEM_FENCE);116 int tid = get_local_id(0);117}118 119void success_nokernel_barrier_both_ids(__global int * foo, int size) {120 work_group_barrier(CLK_GLOBAL_MEM_FENCE);121 int gid = get_global_id(0);122 int lid = get_local_id(0);123}124#endif125 126#ifdef OLDCLNEWAOC // OpenCL 1.2 Altera Offline Compiler >= 17.1127void __kernel error_barrier_no_id(__global int * foo, int size) {128 // 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]129 for (int j = 0; j < 256; j++) {130 for (int i = 256; i < size; i+= 256) {131 foo[j] += foo[j+i];132 }133 }134 barrier(CLK_GLOBAL_MEM_FENCE);135 // 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 efficient136 for (int i = 1; i < 256; i++) {137 foo[0] += foo[i];138 }139}140 141__attribute__ ((reqd_work_group_size(1,1,1)))142void __kernel error_barrier_no_id_work_group_size(__global int * foo, int size) {143 // 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]144 for (int j = 0; j < 256; j++) {145 for (int i = 256; i < size; i+= 256) {146 foo[j] += foo[j+i];147 }148 }149 barrier(CLK_GLOBAL_MEM_FENCE);150 // 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 efficient151 for (int i = 1; i < 256; i++) {152 foo[0] += foo[i];153 }154}155 156__attribute__ ((reqd_work_group_size(2,1,1)))157void __kernel success_barrier_no_id_work_group_size(__global int * foo, int size) {158 for (int j = 0; j < 256; j++) {159 for (int i = 256; i < size; i+= 256) {160 foo[j] += foo[j+i];161 }162 }163 barrier(CLK_GLOBAL_MEM_FENCE);164 for (int i = 1; i < 256; i++) {165 foo[0] += foo[i];166 }167}168 169void __kernel success_barrier_global_id(__global int * foo, int size) {170 barrier(CLK_GLOBAL_MEM_FENCE);171 int tid = get_global_id(0);172}173 174void __kernel success_barrier_local_id(__global int * foo, int size) {175 barrier(CLK_GLOBAL_MEM_FENCE);176 int tid = get_local_id(0);177}178 179void __kernel success_barrier_both_ids(__global int * foo, int size) {180 barrier(CLK_GLOBAL_MEM_FENCE);181 int gid = get_global_id(0);182 int lid = get_local_id(0);183}184 185void success_nokernel_barrier_no_id(__global int * foo, int size) {186 for (int j = 0; j < 256; j++) {187 for (int i = 256; i < size; i+= 256) {188 foo[j] += foo[j+i];189 }190 }191 barrier(CLK_GLOBAL_MEM_FENCE);192 for (int i = 1; i < 256; i++) {193 foo[0] += foo[i];194 }195}196 197void success_nokernel_barrier_global_id(__global int * foo, int size) {198 barrier(CLK_GLOBAL_MEM_FENCE);199 int tid = get_global_id(0);200}201 202void success_nokernel_barrier_local_id(__global int * foo, int size) {203 barrier(CLK_GLOBAL_MEM_FENCE);204 int tid = get_local_id(0);205}206 207void success_nokernel_barrier_both_ids(__global int * foo, int size) {208 barrier(CLK_GLOBAL_MEM_FENCE);209 int gid = get_global_id(0);210 int lid = get_local_id(0);211}212#endif213 214#ifdef NEWCLNEWAOC // OpenCL 2.0 Altera Offline Compiler >= 17.1215void __kernel error_barrier_no_id(__global int * foo, int size) {216 // 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]217 for (int j = 0; j < 256; j++) {218 for (int i = 256; i < size; i+= 256) {219 foo[j] += foo[j+i];220 }221 }222 work_group_barrier(CLK_GLOBAL_MEM_FENCE);223 // 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 efficient224 for (int i = 1; i < 256; i++) {225 foo[0] += foo[i];226 }227}228 229__attribute__ ((reqd_work_group_size(1,1,1)))230void __kernel error_barrier_no_id_work_group_size(__global int * foo, int size) {231 // 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]232 for (int j = 0; j < 256; j++) {233 for (int i = 256; i < size; i+= 256) {234 foo[j] += foo[j+i];235 }236 }237 work_group_barrier(CLK_GLOBAL_MEM_FENCE);238 // 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 efficient239 for (int i = 1; i < 256; i++) {240 foo[0] += foo[i];241 }242}243 244__attribute__ ((reqd_work_group_size(2,1,1)))245void __kernel success_barrier_no_id_work_group_size(__global int * foo, int size) {246 for (int j = 0; j < 256; j++) {247 for (int i = 256; i < size; i+= 256) {248 foo[j] += foo[j+i];249 }250 }251 work_group_barrier(CLK_GLOBAL_MEM_FENCE);252 for (int i = 1; i < 256; i++) {253 foo[0] += foo[i];254 }255}256 257void __kernel success_barrier_global_id(__global int * foo, int size) {258 work_group_barrier(CLK_GLOBAL_MEM_FENCE);259 int tid = get_global_id(0);260}261 262void __kernel success_barrier_local_id(__global int * foo, int size) {263 work_group_barrier(CLK_GLOBAL_MEM_FENCE);264 int tid = get_local_id(0);265}266 267void __kernel success_barrier_both_ids(__global int * foo, int size) {268 work_group_barrier(CLK_GLOBAL_MEM_FENCE);269 int gid = get_global_id(0);270 int lid = get_local_id(0);271}272 273void success_nokernel_barrier_no_id(__global int * foo, int size) {274 for (int j = 0; j < 256; j++) {275 for (int i = 256; i < size; i+= 256) {276 foo[j] += foo[j+i];277 }278 }279 work_group_barrier(CLK_GLOBAL_MEM_FENCE);280 for (int i = 1; i < 256; i++) {281 foo[0] += foo[i];282 }283}284 285void success_nokernel_barrier_global_id(__global int * foo, int size) {286 work_group_barrier(CLK_GLOBAL_MEM_FENCE);287 int tid = get_global_id(0);288}289 290void success_nokernel_barrier_local_id(__global int * foo, int size) {291 work_group_barrier(CLK_GLOBAL_MEM_FENCE);292 int tid = get_local_id(0);293}294 295void success_nokernel_barrier_both_ids(__global int * foo, int size) {296 work_group_barrier(CLK_GLOBAL_MEM_FENCE);297 int gid = get_global_id(0);298 int lid = get_local_id(0);299}300#endif301