1// REQUIRES: amdgpu-registered-target 2// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu tahiti -emit-llvm -o - %s | FileCheck -enable-var-scope --check-prefixes=CHECK-AMDGCN,CHECK %s 3// RUN: %clang_cc1 -cl-std=CL2.0 -triple spirv64-amd-amdhsa -emit-llvm -o - %s | FileCheck -enable-var-scope --check-prefixes=CHECK,CHECK-SPIRV %s 4 5 6#pragma OPENCL EXTENSION cl_khr_fp64 : enable 7 8typedef unsigned long ulong; 9typedef unsigned int uint; 10typedef unsigned short ushort; 11typedef half __attribute__((ext_vector_type(2))) half2; 12typedef short __attribute__((ext_vector_type(2))) short2; 13typedef ushort __attribute__((ext_vector_type(2))) ushort2; 14typedef uint __attribute__((ext_vector_type(4))) uint4; 15 16// CHECK-LABEL: @test_div_scale_f64 17// CHECK: {{.*}}call{{.*}} { double, i1 } @llvm.amdgcn.div.scale.f64(double %a, double %b, i1 true) 18// CHECK-DAG: [[FLAG:%.+]] = extractvalue { double, i1 } %{{.+}}, 1 19// CHECK-DAG: [[VAL:%.+]] = extractvalue { double, i1 } %{{.+}}, 0 20// CHECK: [[FLAGEXT:%.+]] = zext i1 [[FLAG]] to i32 21// CHECK: store i32 [[FLAGEXT]] 22void test_div_scale_f64(global double* out, global int* flagout, double a, double b) 23{ 24 bool flag; 25 *out = __builtin_amdgcn_div_scale(a, b, true, &flag); 26 *flagout = flag; 27} 28 29// CHECK-LABEL: @test_div_scale_f32( 30// CHECK: {{.*}}call{{.*}} { float, i1 } @llvm.amdgcn.div.scale.f32(float %a, float %b, i1 true) 31// CHECK-DAG: [[FLAG:%.+]] = extractvalue { float, i1 } %{{.+}}, 1 32// CHECK-DAG: [[VAL:%.+]] = extractvalue { float, i1 } %{{.+}}, 0 33// CHECK: [[FLAGEXT:%.+]] = zext i1 [[FLAG]] to i8 34// CHECK: store i8 [[FLAGEXT]] 35void test_div_scale_f32(global float* out, global bool* flagout, float a, float b) 36{ 37 bool flag; 38 *out = __builtin_amdgcn_div_scalef(a, b, true, &flag); 39 *flagout = flag; 40} 41 42// CHECK-LABEL: @test_div_scale_f32_global_ptr( 43// CHECK: {{.*}}call{{.*}} { float, i1 } @llvm.amdgcn.div.scale.f32(float %a, float %b, i1 true) 44// CHECK-DAG: [[FLAG:%.+]] = extractvalue { float, i1 } %{{.+}}, 1 45// CHECK-DAG: [[VAL:%.+]] = extractvalue { float, i1 } %{{.+}}, 0 46// CHECK: [[FLAGEXT:%.+]] = zext i1 [[FLAG]] to i8 47// CHECK: store i8 [[FLAGEXT]] 48void test_div_scale_f32_global_ptr(global float* out, global int* flagout, float a, float b, global bool* flag) 49{ 50 *out = __builtin_amdgcn_div_scalef(a, b, true, flag); 51} 52 53// CHECK-LABEL: @test_div_scale_f32_generic_ptr( 54// CHECK: {{.*}}call{{.*}} { float, i1 } @llvm.amdgcn.div.scale.f32(float %a, float %b, i1 true) 55// CHECK-DAG: [[FLAG:%.+]] = extractvalue { float, i1 } %{{.+}}, 1 56// CHECK-DAG: [[VAL:%.+]] = extractvalue { float, i1 } %{{.+}}, 0 57// CHECK: [[FLAGEXT:%.+]] = zext i1 [[FLAG]] to i8 58// CHECK: store i8 [[FLAGEXT]] 59void test_div_scale_f32_generic_ptr(global float* out, global int* flagout, float a, float b, global bool* flag_arg) 60{ 61 generic bool* flag = flag_arg; 62 *out = __builtin_amdgcn_div_scalef(a, b, true, flag); 63} 64 65// CHECK-LABEL: @test_div_fmas_f32 66// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.div.fmas.f32 67void test_div_fmas_f32(global float* out, float a, float b, float c, int d) 68{ 69 *out = __builtin_amdgcn_div_fmasf(a, b, c, d); 70} 71 72// CHECK-LABEL: @test_div_fmas_f64 73// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.div.fmas.f64 74void test_div_fmas_f64(global double* out, double a, double b, double c, int d) 75{ 76 *out = __builtin_amdgcn_div_fmas(a, b, c, d); 77} 78 79// CHECK-LABEL: @test_div_fixup_f32 80// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.div.fixup.f32 81void test_div_fixup_f32(global float* out, float a, float b, float c) 82{ 83 *out = __builtin_amdgcn_div_fixupf(a, b, c); 84} 85 86// CHECK-LABEL: @test_div_fixup_f64 87// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.div.fixup.f64 88void test_div_fixup_f64(global double* out, double a, double b, double c) 89{ 90 *out = __builtin_amdgcn_div_fixup(a, b, c); 91} 92 93// CHECK-LABEL: @test_trig_preop_f32 94// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.trig.preop.f32 95void test_trig_preop_f32(global float* out, float a, int b) 96{ 97 *out = __builtin_amdgcn_trig_preopf(a, b); 98} 99 100// CHECK-LABEL: @test_trig_preop_f64 101// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.trig.preop.f64 102void test_trig_preop_f64(global double* out, double a, int b) 103{ 104 *out = __builtin_amdgcn_trig_preop(a, b); 105} 106 107// CHECK-LABEL: @test_rcp_f32 108// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.rcp.f32 109void test_rcp_f32(global float* out, float a) 110{ 111 *out = __builtin_amdgcn_rcpf(a); 112} 113 114// CHECK-LABEL: @test_rcp_f64 115// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.rcp.f64 116void test_rcp_f64(global double* out, double a) 117{ 118 *out = __builtin_amdgcn_rcp(a); 119} 120 121// CHECK-LABEL: @test_sqrt_f32 122// CHECK: {{.*}}call{{.*}} float @llvm.{{((amdgcn.){0,1})}}sqrt.f32 123void test_sqrt_f32(global float* out, float a) 124{ 125 *out = __builtin_amdgcn_sqrtf(a); 126} 127 128// CHECK-LABEL: @test_sqrt_f64 129// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.sqrt.f64 130void test_sqrt_f64(global double* out, double a) 131{ 132 *out = __builtin_amdgcn_sqrt(a); 133} 134 135// CHECK-LABEL: @test_rsq_f32 136// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.rsq.f32 137void test_rsq_f32(global float* out, float a) 138{ 139 *out = __builtin_amdgcn_rsqf(a); 140} 141 142// CHECK-LABEL: @test_rsq_f64 143// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.rsq.f64 144void test_rsq_f64(global double* out, double a) 145{ 146 *out = __builtin_amdgcn_rsq(a); 147} 148 149// CHECK-LABEL: @test_rsq_clamp_f32 150// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.rsq.clamp.f32 151void test_rsq_clamp_f32(global float* out, float a) 152{ 153 *out = __builtin_amdgcn_rsq_clampf(a); 154} 155 156// CHECK-LABEL: @test_rsq_clamp_f64 157// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.rsq.clamp.f64 158void test_rsq_clamp_f64(global double* out, double a) 159{ 160 *out = __builtin_amdgcn_rsq_clamp(a); 161} 162 163// CHECK-LABEL: @test_sin_f32 164// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.sin.f32 165void test_sin_f32(global float* out, float a) 166{ 167 *out = __builtin_amdgcn_sinf(a); 168} 169 170// CHECK-LABEL: @test_cos_f32 171// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.cos.f32 172void test_cos_f32(global float* out, float a) 173{ 174 *out = __builtin_amdgcn_cosf(a); 175} 176 177// CHECK-LABEL: @test_log_f32 178// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.log.f32 179void test_log_f32(global float* out, float a) 180{ 181 *out = __builtin_amdgcn_logf(a); 182} 183 184// CHECK-LABEL: @test_exp2_f32 185// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.exp2.f32 186void test_exp2_f32(global float* out, float a) 187{ 188 *out = __builtin_amdgcn_exp2f(a); 189} 190 191// CHECK-LABEL: @test_log_clamp_f32 192// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.log.clamp.f32 193void test_log_clamp_f32(global float* out, float a) 194{ 195 *out = __builtin_amdgcn_log_clampf(a); 196} 197 198// CHECK-LABEL: @test_ldexp_f32 199// CHECK: {{.*}}call{{.*}} float @llvm.ldexp.f32.i32 200void test_ldexp_f32(global float* out, float a, int b) 201{ 202 *out = __builtin_amdgcn_ldexpf(a, b); 203} 204 205// CHECK-LABEL: @test_ldexp_f64 206// CHECK: {{.*}}call{{.*}} double @llvm.ldexp.f64.i32 207void test_ldexp_f64(global double* out, double a, int b) 208{ 209 *out = __builtin_amdgcn_ldexp(a, b); 210} 211 212// CHECK-LABEL: @test_frexp_mant_f32 213// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.frexp.mant.f32 214void test_frexp_mant_f32(global float* out, float a) 215{ 216 *out = __builtin_amdgcn_frexp_mantf(a); 217} 218 219// CHECK-LABEL: @test_frexp_mant_f64 220// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.frexp.mant.f64 221void test_frexp_mant_f64(global double* out, double a) 222{ 223 *out = __builtin_amdgcn_frexp_mant(a); 224} 225 226// CHECK-LABEL: @test_frexp_exp_f32 227// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.frexp.exp.i32.f32 228void test_frexp_exp_f32(global int* out, float a) 229{ 230 *out = __builtin_amdgcn_frexp_expf(a); 231} 232 233// CHECK-LABEL: @test_frexp_exp_f64 234// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.frexp.exp.i32.f64 235void test_frexp_exp_f64(global int* out, double a) 236{ 237 *out = __builtin_amdgcn_frexp_exp(a); 238} 239 240// CHECK-LABEL: @test_fract_f32 241// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.fract.f32 242void test_fract_f32(global int* out, float a) 243{ 244 *out = __builtin_amdgcn_fractf(a); 245} 246 247// CHECK-LABEL: @test_fract_f64 248// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.fract.f64 249void test_fract_f64(global int* out, double a) 250{ 251 *out = __builtin_amdgcn_fract(a); 252} 253 254// CHECK-LABEL: @test_lerp 255// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.lerp 256void test_lerp(global int* out, int a, int b, int c) 257{ 258 *out = __builtin_amdgcn_lerp(a, b, c); 259} 260 261// CHECK-LABEL: @test_sicmp_i32 262// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.icmp.i64.i32(i32 %a, i32 %b, i32 32) 263void test_sicmp_i32(global ulong* out, int a, int b) 264{ 265 *out = __builtin_amdgcn_sicmp(a, b, 32); 266} 267 268// CHECK-LABEL: @test_uicmp_i32 269// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.icmp.i64.i32(i32 %a, i32 %b, i32 32) 270void test_uicmp_i32(global ulong* out, uint a, uint b) 271{ 272 *out = __builtin_amdgcn_uicmp(a, b, 32); 273} 274 275// CHECK-LABEL: @test_sicmp_i64 276// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.icmp.i64.i64(i64 %a, i64 %b, i32 38) 277void test_sicmp_i64(global ulong* out, long a, long b) 278{ 279 *out = __builtin_amdgcn_sicmpl(a, b, 39-1); 280} 281 282// CHECK-LABEL: @test_uicmp_i64 283// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.icmp.i64.i64(i64 %a, i64 %b, i32 35) 284void test_uicmp_i64(global ulong* out, ulong a, ulong b) 285{ 286 *out = __builtin_amdgcn_uicmpl(a, b, 30+5); 287} 288 289// CHECK-LABEL: @test_ds_swizzle 290// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.ds.swizzle(i32 %a, i32 32) 291void test_ds_swizzle(global int* out, int a) 292{ 293 *out = __builtin_amdgcn_ds_swizzle(a, 32); 294} 295 296// CHECK-LABEL: @test_ds_permute 297// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.ds.permute(i32 %a, i32 %b) 298void test_ds_permute(global int* out, int a, int b) 299{ 300 out[0] = __builtin_amdgcn_ds_permute(a, b); 301} 302 303// CHECK-LABEL: @test_ds_bpermute 304// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.ds.bpermute(i32 %a, i32 %b) 305void test_ds_bpermute(global int* out, int a, int b) 306{ 307 *out = __builtin_amdgcn_ds_bpermute(a, b); 308} 309 310// CHECK-LABEL: @test_readfirstlane 311// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.readfirstlane.i32(i32 %a) 312void test_readfirstlane(global int* out, int a) 313{ 314 *out = __builtin_amdgcn_readfirstlane(a); 315} 316 317// CHECK-LABEL: @test_readlane 318// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.readlane.i32(i32 %a, i32 %b) 319void test_readlane(global int* out, int a, int b) 320{ 321 *out = __builtin_amdgcn_readlane(a, b); 322} 323 324// CHECK-LABEL: @test_fcmp_f32 325// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.fcmp.i64.f32(float %a, float %b, i32 5) 326void test_fcmp_f32(global ulong* out, float a, float b) 327{ 328 *out = __builtin_amdgcn_fcmpf(a, b, 5); 329} 330 331// CHECK-LABEL: @test_fcmp_f64 332// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.fcmp.i64.f64(double %a, double %b, i32 6) 333void test_fcmp_f64(global ulong* out, double a, double b) 334{ 335 *out = __builtin_amdgcn_fcmp(a, b, 3+3); 336} 337 338// CHECK-LABEL: @test_class_f32 339// CHECK: {{.*}}call{{.*}} i1 @llvm.amdgcn.class.f32 340void test_class_f32(global float* out, float a, int b) 341{ 342 *out = __builtin_amdgcn_classf(a, b); 343} 344 345// CHECK-LABEL: @test_class_f64 346// CHECK: {{.*}}call{{.*}} i1 @llvm.amdgcn.class.f64 347void test_class_f64(global double* out, double a, int b) 348{ 349 *out = __builtin_amdgcn_class(a, b); 350} 351 352// CHECK-LABEL: @test_buffer_wbinvl1 353// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.buffer.wbinvl1( 354void test_buffer_wbinvl1() 355{ 356 __builtin_amdgcn_buffer_wbinvl1(); 357} 358 359// CHECK-LABEL: @test_s_dcache_inv 360// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.dcache.inv( 361void test_s_dcache_inv() 362{ 363 __builtin_amdgcn_s_dcache_inv(); 364} 365 366// CHECK-LABEL: @test_s_waitcnt 367// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.waitcnt( 368void test_s_waitcnt() 369{ 370 __builtin_amdgcn_s_waitcnt(0); 371} 372 373// CHECK-LABEL: @test_s_sendmsg 374// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.sendmsg( 375void test_s_sendmsg() 376{ 377 __builtin_amdgcn_s_sendmsg(1, 0); 378} 379 380// CHECK-LABEL: @test_s_sendmsg_var 381// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.sendmsg( 382void test_s_sendmsg_var(int in) 383{ 384 __builtin_amdgcn_s_sendmsg(1, in); 385} 386 387// CHECK-LABEL: @test_s_sendmsghalt 388// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.sendmsghalt( 389void test_s_sendmsghalt() 390{ 391 __builtin_amdgcn_s_sendmsghalt(1, 0); 392} 393 394// CHECK-LABEL: @test_s_sendmsghalt 395// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.sendmsghalt( 396void test_s_sendmsghalt_var(int in) 397{ 398 __builtin_amdgcn_s_sendmsghalt(1, in); 399} 400 401// CHECK-LABEL: @test_s_barrier 402// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.barrier( 403void test_s_barrier() 404{ 405 __builtin_amdgcn_s_barrier(); 406} 407 408// CHECK-LABEL: @test_wave_barrier 409// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.wave.barrier( 410void test_wave_barrier() 411{ 412 __builtin_amdgcn_wave_barrier(); 413} 414 415// CHECK-LABEL: @test_sched_barrier 416// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.sched.barrier(i32 0) 417// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.sched.barrier(i32 1) 418// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.sched.barrier(i32 4) 419// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.sched.barrier(i32 15) 420void test_sched_barrier() 421{ 422 __builtin_amdgcn_sched_barrier(0); 423 __builtin_amdgcn_sched_barrier(1); 424 __builtin_amdgcn_sched_barrier(4); 425 __builtin_amdgcn_sched_barrier(15); 426} 427 428// CHECK-LABEL: @test_sched_group_barrier 429// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.sched.group.barrier(i32 0, i32 1, i32 2) 430// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.sched.group.barrier(i32 1, i32 2, i32 4) 431// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.sched.group.barrier(i32 4, i32 8, i32 16) 432// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.sched.group.barrier(i32 15, i32 10000, i32 -1) 433void test_sched_group_barrier() 434{ 435 __builtin_amdgcn_sched_group_barrier(0, 1, 2); 436 __builtin_amdgcn_sched_group_barrier(1, 2, 4); 437 __builtin_amdgcn_sched_group_barrier(4, 8, 16); 438 __builtin_amdgcn_sched_group_barrier(15, 10000, -1); 439} 440 441// CHECK-LABEL: @test_iglp_opt 442// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.iglp.opt(i32 0) 443// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.iglp.opt(i32 1) 444// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.iglp.opt(i32 4) 445// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.iglp.opt(i32 15) 446void test_iglp_opt() 447{ 448 __builtin_amdgcn_iglp_opt(0); 449 __builtin_amdgcn_iglp_opt(1); 450 __builtin_amdgcn_iglp_opt(4); 451 __builtin_amdgcn_iglp_opt(15); 452} 453 454// CHECK-LABEL: @test_s_sleep 455// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.sleep(i32 1) 456// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.sleep(i32 15) 457void test_s_sleep() 458{ 459 __builtin_amdgcn_s_sleep(1); 460 __builtin_amdgcn_s_sleep(15); 461} 462 463// CHECK-LABEL: @test_s_incperflevel 464// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.incperflevel(i32 1) 465// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.incperflevel(i32 15) 466void test_s_incperflevel() 467{ 468 __builtin_amdgcn_s_incperflevel(1); 469 __builtin_amdgcn_s_incperflevel(15); 470} 471 472// CHECK-LABEL: @test_s_decperflevel 473// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.decperflevel(i32 1) 474// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.decperflevel(i32 15) 475void test_s_decperflevel() 476{ 477 __builtin_amdgcn_s_decperflevel(1); 478 __builtin_amdgcn_s_decperflevel(15); 479} 480 481// CHECK-LABEL: @test_s_setprio 482// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.setprio(i16 0) 483// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.setprio(i16 3) 484void test_s_setprio() 485{ 486 __builtin_amdgcn_s_setprio(0); 487 __builtin_amdgcn_s_setprio(3); 488} 489 490// CHECK-LABEL: @test_cubeid( 491// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.cubeid(float %a, float %b, float %c) 492void test_cubeid(global float* out, float a, float b, float c) { 493 *out = __builtin_amdgcn_cubeid(a, b, c); 494} 495 496// CHECK-LABEL: @test_cubesc( 497// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.cubesc(float %a, float %b, float %c) 498void test_cubesc(global float* out, float a, float b, float c) { 499 *out = __builtin_amdgcn_cubesc(a, b, c); 500} 501 502// CHECK-LABEL: @test_cubetc( 503// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.cubetc(float %a, float %b, float %c) 504void test_cubetc(global float* out, float a, float b, float c) { 505 *out = __builtin_amdgcn_cubetc(a, b, c); 506} 507 508// CHECK-LABEL: @test_cubema( 509// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.cubema(float %a, float %b, float %c) 510void test_cubema(global float* out, float a, float b, float c) { 511 *out = __builtin_amdgcn_cubema(a, b, c); 512} 513 514// CHECK-LABEL: @test_read_exec( 515// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.ballot.i64(i1 true) 516void test_read_exec(global ulong* out) { 517 *out = __builtin_amdgcn_read_exec(); 518} 519 520// CHECK: declare i64 @llvm.amdgcn.ballot.i64(i1){{.*}} #[[$NOUNWIND_READONLY:[0-9]+]] 521 522// CHECK-LABEL: @test_read_exec_lo( 523// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.ballot.i32(i1 true) 524void test_read_exec_lo(global uint* out) { 525 *out = __builtin_amdgcn_read_exec_lo(); 526} 527 528// CHECK: declare i32 @llvm.amdgcn.ballot.i32(i1){{.*}} #[[$NOUNWIND_READONLY:[0-9]+]] 529 530// CHECK-LABEL: @test_read_exec_hi( 531// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.ballot.i64(i1 true) 532// CHECK: lshr i64 [[A:%.*]], 32 533// CHECK: trunc nuw i64 [[B:%.*]] to i32 534void test_read_exec_hi(global uint* out) { 535 *out = __builtin_amdgcn_read_exec_hi(); 536} 537 538// CHECK-LABEL: @test_dispatch_ptr 539// CHECK: {{.*}}call align 4 dereferenceable(64){{.*}} ptr addrspace(4) @llvm.amdgcn.dispatch.ptr() 540#if !defined(__SPIRV__) 541void test_dispatch_ptr(__constant unsigned char ** out) 542#else 543void test_dispatch_ptr(__attribute__((address_space(4))) unsigned char ** out) 544#endif 545{ 546 *out = __builtin_amdgcn_dispatch_ptr(); 547} 548 549// CHECK-LABEL: @test_queue_ptr 550// CHECK: {{.*}}call{{.*}} ptr addrspace(4) @llvm.amdgcn.queue.ptr() 551#if !defined(__SPIRV__) 552void test_queue_ptr(__constant unsigned char ** out) 553#else 554void test_queue_ptr(__attribute__((address_space(4))) unsigned char ** out) 555#endif 556{ 557 *out = __builtin_amdgcn_queue_ptr(); 558} 559 560// CHECK-LABEL: @test_kernarg_segment_ptr 561// CHECK: {{.*}}call{{.*}} ptr addrspace(4) @llvm.amdgcn.kernarg.segment.ptr() 562#if !defined(__SPIRV__) 563void test_kernarg_segment_ptr(__constant unsigned char ** out) 564#else 565void test_kernarg_segment_ptr(__attribute__((address_space(4))) unsigned char ** out) 566#endif 567{ 568 *out = __builtin_amdgcn_kernarg_segment_ptr(); 569} 570 571// CHECK-LABEL: @test_implicitarg_ptr 572// CHECK: {{.*}}call{{.*}} ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() 573#if !defined(__SPIRV__) 574void test_implicitarg_ptr(__constant unsigned char ** out) 575#else 576void test_implicitarg_ptr(__attribute__((address_space(4))) unsigned char ** out) 577#endif 578{ 579 *out = __builtin_amdgcn_implicitarg_ptr(); 580} 581 582// CHECK-LABEL: @test_get_group_id( 583// CHECK: tail call{{.*}} i32 @llvm.amdgcn.workgroup.id.x() 584// CHECK: tail call{{.*}} i32 @llvm.amdgcn.workgroup.id.y() 585// CHECK: tail call{{.*}} i32 @llvm.amdgcn.workgroup.id.z() 586void test_get_group_id(int d, global int *out) 587{ 588 switch (d) { 589 case 0: *out = __builtin_amdgcn_workgroup_id_x(); break; 590 case 1: *out = __builtin_amdgcn_workgroup_id_y(); break; 591 case 2: *out = __builtin_amdgcn_workgroup_id_z(); break; 592 default: *out = 0; 593 } 594} 595 596// CHECK-LABEL: @test_s_getreg( 597// CHECK: tail call{{.*}} i32 @llvm.amdgcn.s.getreg(i32 0) 598// CHECK: tail call{{.*}} i32 @llvm.amdgcn.s.getreg(i32 1) 599// CHECK: tail call{{.*}} i32 @llvm.amdgcn.s.getreg(i32 65535) 600void test_s_getreg(volatile global uint *out) 601{ 602 *out = __builtin_amdgcn_s_getreg(0); 603 *out = __builtin_amdgcn_s_getreg(1); 604 *out = __builtin_amdgcn_s_getreg(65535); 605} 606 607// CHECK-LABEL: @test_get_local_id( 608// CHECK: tail call noundef range(i32 0, 1024){{.*}} i32 @llvm.amdgcn.workitem.id.x() 609// CHECK: tail call noundef range(i32 0, 1024){{.*}} i32 @llvm.amdgcn.workitem.id.y() 610// CHECK: tail call noundef range(i32 0, 1024){{.*}} i32 @llvm.amdgcn.workitem.id.z() 611void test_get_local_id(int d, global int *out) 612{ 613 switch (d) { 614 case 0: *out = __builtin_amdgcn_workitem_id_x(); break; 615 case 1: *out = __builtin_amdgcn_workitem_id_y(); break; 616 case 2: *out = __builtin_amdgcn_workitem_id_z(); break; 617 default: *out = 0; 618 } 619} 620 621// CHECK-LABEL: @test_get_workgroup_size( 622// CHECK: {{.*}}call align 8 dereferenceable(256){{.*}} ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() 623// CHECK: getelementptr inbounds nuw i8, ptr addrspace(4) %{{.*}}, i64 12 624// CHECK: load i16, ptr addrspace(4) %{{.*}}, align 4, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef 625// CHECK: getelementptr inbounds nuw i8, ptr addrspace(4) %{{.*}}, i64 14 626// CHECK: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef 627// CHECK: getelementptr inbounds nuw i8, ptr addrspace(4) %{{.*}}, i64 16 628// CHECK: load i16, ptr addrspace(4) %{{.*}}, align 8, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef 629void test_get_workgroup_size(int d, global int *out) 630{ 631 switch (d) { 632 case 0: *out = __builtin_amdgcn_workgroup_size_x() + 1; break; 633 case 1: *out = __builtin_amdgcn_workgroup_size_y(); break; 634 case 2: *out = __builtin_amdgcn_workgroup_size_z(); break; 635 default: *out = 0; 636 } 637} 638 639// CHECK-LABEL: @test_get_grid_size( 640// CHECK: {{.*}}call align 4 dereferenceable(64){{.*}} ptr addrspace(4) @llvm.amdgcn.dispatch.ptr() 641// CHECK: getelementptr inbounds nuw i8, ptr addrspace(4) %{{.*}}, i64 %{{.+}} 642// CHECK: load i32, ptr addrspace(4) %{{.*}}, align 4, !range [[$GRID_RANGE:![0-9]+]], !invariant.load 643void test_get_grid_size(int d, global int *out) 644{ 645 switch (d) { 646 case 0: *out = __builtin_amdgcn_grid_size_x(); break; 647 case 1: *out = __builtin_amdgcn_grid_size_y(); break; 648 case 2: *out = __builtin_amdgcn_grid_size_z(); break; 649 default: *out = 0; 650 } 651} 652 653// CHECK-LABEL: @test_fmed3_f32 654// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.fmed3.f32( 655void test_fmed3_f32(global float* out, float a, float b, float c) 656{ 657 *out = __builtin_amdgcn_fmed3f(a, b, c); 658} 659 660// CHECK-LABEL: @test_s_getpc 661// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.s.getpc() 662void test_s_getpc(global ulong* out) 663{ 664 *out = __builtin_amdgcn_s_getpc(); 665} 666 667// CHECK-LABEL: @test_ds_append_lds( 668// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.ds.append.p3(ptr addrspace(3) %{{.+}}, i1 false) 669#if !defined(__SPIRV__) 670kernel void test_ds_append_lds(global int* out, local int* ptr) { 671#else 672kernel void test_ds_append_lds(__attribute__((address_space(1))) int* out, __attribute__((address_space(3))) int* ptr) { 673#endif 674 *out = __builtin_amdgcn_ds_append(ptr); 675} 676 677// CHECK-LABEL: @test_ds_consume_lds( 678// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.ds.consume.p3(ptr addrspace(3) %{{.+}}, i1 false) 679 680#if !defined(__SPIRV__) 681kernel void test_ds_consume_lds(global int* out, local int* ptr) { 682#else 683kernel void test_ds_consume_lds(__attribute__((address_space(1))) int* out, __attribute__((address_space(3))) int* ptr) { 684#endif 685 *out = __builtin_amdgcn_ds_consume(ptr); 686} 687 688// CHECK-LABEL: @test_gws_init( 689// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.ds.gws.init(i32 %value, i32 %id) 690kernel void test_gws_init(uint value, uint id) { 691 __builtin_amdgcn_ds_gws_init(value, id); 692} 693 694// CHECK-LABEL: @test_gws_barrier( 695// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.ds.gws.barrier(i32 %value, i32 %id) 696kernel void test_gws_barrier(uint value, uint id) { 697 __builtin_amdgcn_ds_gws_barrier(value, id); 698} 699 700// CHECK-LABEL: @test_gws_sema_v( 701// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.ds.gws.sema.v(i32 %id) 702kernel void test_gws_sema_v(uint id) { 703 __builtin_amdgcn_ds_gws_sema_v(id); 704} 705 706// CHECK-LABEL: @test_gws_sema_br( 707// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.ds.gws.sema.br(i32 %value, i32 %id) 708kernel void test_gws_sema_br(uint value, uint id) { 709 __builtin_amdgcn_ds_gws_sema_br(value, id); 710} 711 712// CHECK-LABEL: @test_gws_sema_p( 713// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.ds.gws.sema.p(i32 %id) 714kernel void test_gws_sema_p(uint id) { 715 __builtin_amdgcn_ds_gws_sema_p(id); 716} 717 718// CHECK-LABEL: @test_mbcnt_lo( 719// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.mbcnt.lo(i32 %src0, i32 %src1) 720kernel void test_mbcnt_lo(global uint* out, uint src0, uint src1) { 721 *out = __builtin_amdgcn_mbcnt_lo(src0, src1); 722} 723 724// CHECK-LABEL: @test_mbcnt_hi( 725// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.mbcnt.hi(i32 %src0, i32 %src1) 726kernel void test_mbcnt_hi(global uint* out, uint src0, uint src1) { 727 *out = __builtin_amdgcn_mbcnt_hi(src0, src1); 728} 729 730// CHECK-LABEL: @test_alignbit( 731// CHECK: tail call{{.*}} i32 @llvm.fshr.i32(i32 %src0, i32 %src1, i32 %src2) 732kernel void test_alignbit(global uint* out, uint src0, uint src1, uint src2) { 733 *out = __builtin_amdgcn_alignbit(src0, src1, src2); 734} 735 736// CHECK-LABEL: @test_alignbyte( 737// CHECK: tail call{{.*}} i32 @llvm.amdgcn.alignbyte(i32 %src0, i32 %src1, i32 %src2) 738kernel void test_alignbyte(global uint* out, uint src0, uint src1, uint src2) { 739 *out = __builtin_amdgcn_alignbyte(src0, src1, src2); 740} 741 742// CHECK-LABEL: @test_ubfe( 743// CHECK: tail call{{.*}} i32 @llvm.amdgcn.ubfe.i32(i32 %src0, i32 %src1, i32 %src2) 744kernel void test_ubfe(global uint* out, uint src0, uint src1, uint src2) { 745 *out = __builtin_amdgcn_ubfe(src0, src1, src2); 746} 747 748// CHECK-LABEL: @test_sbfe( 749// CHECK: tail call{{.*}} i32 @llvm.amdgcn.sbfe.i32(i32 %src0, i32 %src1, i32 %src2) 750kernel void test_sbfe(global uint* out, uint src0, uint src1, uint src2) { 751 *out = __builtin_amdgcn_sbfe(src0, src1, src2); 752} 753 754// CHECK-LABEL: @test_cvt_pkrtz( 755// CHECK: tail call{{.*}} <2 x half> @llvm.amdgcn.cvt.pkrtz(float %src0, float %src1) 756kernel void test_cvt_pkrtz(global half2* out, float src0, float src1) { 757 *out = __builtin_amdgcn_cvt_pkrtz(src0, src1); 758} 759 760// CHECK-LABEL: @test_cvt_pknorm_i16( 761// CHECK: tail call{{.*}} <2 x i16> @llvm.amdgcn.cvt.pknorm.i16(float %src0, float %src1) 762kernel void test_cvt_pknorm_i16(global short2* out, float src0, float src1) { 763 *out = __builtin_amdgcn_cvt_pknorm_i16(src0, src1); 764} 765 766// CHECK-LABEL: @test_cvt_pknorm_u16( 767// CHECK: tail call{{.*}} <2 x i16> @llvm.amdgcn.cvt.pknorm.u16(float %src0, float %src1) 768kernel void test_cvt_pknorm_u16(global ushort2* out, float src0, float src1) { 769 *out = __builtin_amdgcn_cvt_pknorm_u16(src0, src1); 770} 771 772// CHECK-LABEL: @test_cvt_pk_i16( 773// CHECK: tail call{{.*}} <2 x i16> @llvm.amdgcn.cvt.pk.i16(i32 %src0, i32 %src1) 774kernel void test_cvt_pk_i16(global short2* out, int src0, int src1) { 775 *out = __builtin_amdgcn_cvt_pk_i16(src0, src1); 776} 777 778// CHECK-LABEL: @test_cvt_pk_u16( 779// CHECK: tail call{{.*}} <2 x i16> @llvm.amdgcn.cvt.pk.u16(i32 %src0, i32 %src1) 780kernel void test_cvt_pk_u16(global ushort2* out, uint src0, uint src1) { 781 *out = __builtin_amdgcn_cvt_pk_u16(src0, src1); 782} 783 784// CHECK-LABEL: @test_cvt_pk_u8_f32 785// CHECK: tail call{{.*}} i32 @llvm.amdgcn.cvt.pk.u8.f32(float %src0, i32 %src1, i32 %src2) 786kernel void test_cvt_pk_u8_f32(global uint* out, float src0, uint src1, uint src2) { 787 *out = __builtin_amdgcn_cvt_pk_u8_f32(src0, src1, src2); 788} 789 790// CHECK-LABEL: @test_sad_u8( 791// CHECK: tail call{{.*}} i32 @llvm.amdgcn.sad.u8(i32 %src0, i32 %src1, i32 %src2) 792kernel void test_sad_u8(global uint* out, uint src0, uint src1, uint src2) { 793 *out = __builtin_amdgcn_sad_u8(src0, src1, src2); 794} 795 796// CHECK-LABEL: test_msad_u8( 797// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.msad.u8(i32 %src0, i32 %src1, i32 %src2) 798kernel void test_msad_u8(global uint* out, uint src0, uint src1, uint src2) { 799 *out = __builtin_amdgcn_msad_u8(src0, src1, src2); 800} 801 802// CHECK-LABEL: test_sad_hi_u8( 803// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.sad.hi.u8(i32 %src0, i32 %src1, i32 %src2) 804kernel void test_sad_hi_u8(global uint* out, uint src0, uint src1, uint src2) { 805 *out = __builtin_amdgcn_sad_hi_u8(src0, src1, src2); 806} 807 808// CHECK-LABEL: @test_sad_u16( 809// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.sad.u16(i32 %src0, i32 %src1, i32 %src2) 810kernel void test_sad_u16(global uint* out, uint src0, uint src1, uint src2) { 811 *out = __builtin_amdgcn_sad_u16(src0, src1, src2); 812} 813 814// CHECK-LABEL: @test_qsad_pk_u16_u8( 815// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.qsad.pk.u16.u8(i64 %src0, i32 %src1, i64 %src2) 816kernel void test_qsad_pk_u16_u8(global ulong* out, ulong src0, uint src1, ulong src2) { 817 *out = __builtin_amdgcn_qsad_pk_u16_u8(src0, src1, src2); 818} 819 820// CHECK-LABEL: @test_mqsad_pk_u16_u8( 821// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.mqsad.pk.u16.u8(i64 %src0, i32 %src1, i64 %src2) 822kernel void test_mqsad_pk_u16_u8(global ulong* out, ulong src0, uint src1, ulong src2) { 823 *out = __builtin_amdgcn_mqsad_pk_u16_u8(src0, src1, src2); 824} 825 826// CHECK-LABEL: test_mqsad_u32_u8( 827// CHECK: {{.*}}call{{.*}} <4 x i32> @llvm.amdgcn.mqsad.u32.u8(i64 %src0, i32 %src1, <4 x i32> %src2) 828kernel void test_mqsad_u32_u8(global uint4* out, ulong src0, uint src1, uint4 src2) { 829 *out = __builtin_amdgcn_mqsad_u32_u8(src0, src1, src2); 830} 831 832// CHECK-LABEL: test_s_setreg( 833// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.setreg(i32 8193, i32 %val) 834kernel void test_s_setreg(uint val) { 835 __builtin_amdgcn_s_setreg(8193, val); 836} 837 838// CHECK-LABEL test_atomic_inc_dec( 839#if !defined(__SPIRV__) 840void test_atomic_inc_dec(local uint *lptr, global uint *gptr, uint val) { 841#else 842void test_atomic_inc_dec(__attribute__((address_space(3))) uint *lptr, __attribute__((address_space(1))) uint *gptr, uint val) { 843#endif 844 uint res; 845 846 // CHECK: atomicrmw uinc_wrap ptr addrspace(3) %lptr, i32 %val syncscope("workgroup") seq_cst, align 4 847 res = __builtin_amdgcn_atomic_inc32(lptr, val, __ATOMIC_SEQ_CST, "workgroup"); 848 849 // CHECK: atomicrmw udec_wrap ptr addrspace(3) %lptr, i32 %val syncscope("workgroup") seq_cst, align 4 850 res = __builtin_amdgcn_atomic_dec32(lptr, val, __ATOMIC_SEQ_CST, "workgroup"); 851 852 // CHECK: atomicrmw uinc_wrap ptr addrspace(1) %gptr, i32 %val syncscope("agent") seq_cst, align 4 853 res = __builtin_amdgcn_atomic_inc32(gptr, val, __ATOMIC_SEQ_CST, "agent"); 854 855 // CHECK: atomicrmw udec_wrap ptr addrspace(1) %gptr, i32 %val seq_cst, align 4 856 res = __builtin_amdgcn_atomic_dec32(gptr, val, __ATOMIC_SEQ_CST, ""); 857 858 // CHECK: atomicrmw volatile udec_wrap ptr addrspace(1) %gptr, i32 %val seq_cst, align 4 859 #if !defined(__SPIRV__) 860 res = __builtin_amdgcn_atomic_dec32((volatile global uint*)gptr, val, __ATOMIC_SEQ_CST, ""); 861 #else 862 res = __builtin_amdgcn_atomic_dec32((volatile __attribute__((address_space(1))) uint*)gptr, val, __ATOMIC_SEQ_CST, ""); 863 #endif 864} 865 866// CHECK-LABEL test_wavefrontsize( 867unsigned test_wavefrontsize() { 868 869 // CHECK-AMDGCN: ret i32 {{[0-9]+}} 870 // CHECK-SPIRV: {{.*}}call{{.*}} i32 @llvm.amdgcn.wavefrontsize() 871 return __builtin_amdgcn_wavefrontsize(); 872} 873 874// CHECK-LABEL test_flt_rounds( 875unsigned test_flt_rounds() { 876 877 // CHECK: {{.*}}call{{.*}} i32 @llvm.get.rounding() 878 unsigned mode = __builtin_flt_rounds(); 879 880#if !defined(__SPIRV__) 881 // CHECK-AMDGCN: call void @llvm.set.rounding(i32 %0) 882 __builtin_set_flt_rounds(mode); 883#endif 884 885 return mode; 886} 887 888// CHECK-LABEL test_get_fpenv( 889unsigned long test_get_fpenv() { 890 // CHECK: {{.*}}call{{.*}} i64 @llvm.get.fpenv.i64() 891 return __builtin_amdgcn_get_fpenv(); 892} 893 894// CHECK-LABEL test_set_fpenv( 895void test_set_fpenv(unsigned long env) { 896 // CHECK: {{.*}}call{{.*}} void @llvm.set.fpenv.i64(i64 %[[ENV:.+]]) 897 __builtin_amdgcn_set_fpenv(env); 898} 899 900// CHECK-DAG: [[$GRID_RANGE]] = !{i32 1, i32 0} 901// CHECK-DAG: [[$WS_RANGE]] = !{i16 1, i16 1025} 902// CHECK-DAG: attributes #[[$NOUNWIND_READONLY]] = { convergent mustprogress nocallback nofree nounwind willreturn memory(none) } 903