1// REQUIRES: amdgpu-registered-target 2// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -S -emit-llvm -o - %s | FileCheck -enable-var-scope %s 3 4#pragma OPENCL EXTENSION cl_khr_fp64 : enable 5 6typedef unsigned long ulong; 7typedef unsigned int uint; 8 9// CHECK-LABEL: @test_div_scale_f64 10// CHECK: call { double, i1 } @llvm.amdgcn.div.scale.f64(double %a, double %b, i1 true) 11// CHECK-DAG: [[FLAG:%.+]] = extractvalue { double, i1 } %{{.+}}, 1 12// CHECK-DAG: [[VAL:%.+]] = extractvalue { double, i1 } %{{.+}}, 0 13// CHECK: [[FLAGEXT:%.+]] = zext i1 [[FLAG]] to i32 14// CHECK: store i32 [[FLAGEXT]] 15void test_div_scale_f64(global double* out, global int* flagout, double a, double b) 16{ 17 bool flag; 18 *out = __builtin_amdgcn_div_scale(a, b, true, &flag); 19 *flagout = flag; 20} 21 22// CHECK-LABEL: @test_div_scale_f32( 23// CHECK: call { float, i1 } @llvm.amdgcn.div.scale.f32(float %a, float %b, i1 true) 24// CHECK-DAG: [[FLAG:%.+]] = extractvalue { float, i1 } %{{.+}}, 1 25// CHECK-DAG: [[VAL:%.+]] = extractvalue { float, i1 } %{{.+}}, 0 26// CHECK: [[FLAGEXT:%.+]] = zext i1 [[FLAG]] to i8 27// CHECK: store i8 [[FLAGEXT]] 28void test_div_scale_f32(global float* out, global bool* flagout, float a, float b) 29{ 30 bool flag; 31 *out = __builtin_amdgcn_div_scalef(a, b, true, &flag); 32 *flagout = flag; 33} 34 35// CHECK-LABEL: @test_div_scale_f32_global_ptr( 36// CHECK: call { float, i1 } @llvm.amdgcn.div.scale.f32(float %a, float %b, i1 true) 37// CHECK-DAG: [[FLAG:%.+]] = extractvalue { float, i1 } %{{.+}}, 1 38// CHECK-DAG: [[VAL:%.+]] = extractvalue { float, i1 } %{{.+}}, 0 39// CHECK: [[FLAGEXT:%.+]] = zext i1 [[FLAG]] to i8 40// CHECK: store i8 [[FLAGEXT]] 41void test_div_scale_f32_global_ptr(global float* out, global int* flagout, float a, float b, global bool* flag) 42{ 43 *out = __builtin_amdgcn_div_scalef(a, b, true, flag); 44} 45 46// CHECK-LABEL: @test_div_scale_f32_generic_ptr( 47// CHECK: call { float, i1 } @llvm.amdgcn.div.scale.f32(float %a, float %b, i1 true) 48// CHECK-DAG: [[FLAG:%.+]] = extractvalue { float, i1 } %{{.+}}, 1 49// CHECK-DAG: [[VAL:%.+]] = extractvalue { float, i1 } %{{.+}}, 0 50// CHECK: [[FLAGEXT:%.+]] = zext i1 [[FLAG]] to i8 51// CHECK: store i8 [[FLAGEXT]] 52void test_div_scale_f32_generic_ptr(global float* out, global int* flagout, float a, float b, global bool* flag_arg) 53{ 54 generic bool* flag = flag_arg; 55 *out = __builtin_amdgcn_div_scalef(a, b, true, flag); 56} 57 58// CHECK-LABEL: @test_div_fmas_f32 59// CHECK: call float @llvm.amdgcn.div.fmas.f32 60void test_div_fmas_f32(global float* out, float a, float b, float c, int d) 61{ 62 *out = __builtin_amdgcn_div_fmasf(a, b, c, d); 63} 64 65// CHECK-LABEL: @test_div_fmas_f64 66// CHECK: call double @llvm.amdgcn.div.fmas.f64 67void test_div_fmas_f64(global double* out, double a, double b, double c, int d) 68{ 69 *out = __builtin_amdgcn_div_fmas(a, b, c, d); 70} 71 72// CHECK-LABEL: @test_div_fixup_f32 73// CHECK: call float @llvm.amdgcn.div.fixup.f32 74void test_div_fixup_f32(global float* out, float a, float b, float c) 75{ 76 *out = __builtin_amdgcn_div_fixupf(a, b, c); 77} 78 79// CHECK-LABEL: @test_div_fixup_f64 80// CHECK: call double @llvm.amdgcn.div.fixup.f64 81void test_div_fixup_f64(global double* out, double a, double b, double c) 82{ 83 *out = __builtin_amdgcn_div_fixup(a, b, c); 84} 85 86// CHECK-LABEL: @test_trig_preop_f32 87// CHECK: call float @llvm.amdgcn.trig.preop.f32 88void test_trig_preop_f32(global float* out, float a, int b) 89{ 90 *out = __builtin_amdgcn_trig_preopf(a, b); 91} 92 93// CHECK-LABEL: @test_trig_preop_f64 94// CHECK: call double @llvm.amdgcn.trig.preop.f64 95void test_trig_preop_f64(global double* out, double a, int b) 96{ 97 *out = __builtin_amdgcn_trig_preop(a, b); 98} 99 100// CHECK-LABEL: @test_rcp_f32 101// CHECK: call float @llvm.amdgcn.rcp.f32 102void test_rcp_f32(global float* out, float a) 103{ 104 *out = __builtin_amdgcn_rcpf(a); 105} 106 107// CHECK-LABEL: @test_rcp_f64 108// CHECK: call double @llvm.amdgcn.rcp.f64 109void test_rcp_f64(global double* out, double a) 110{ 111 *out = __builtin_amdgcn_rcp(a); 112} 113 114// CHECK-LABEL: @test_rsq_f32 115// CHECK: call float @llvm.amdgcn.rsq.f32 116void test_rsq_f32(global float* out, float a) 117{ 118 *out = __builtin_amdgcn_rsqf(a); 119} 120 121// CHECK-LABEL: @test_rsq_f64 122// CHECK: call double @llvm.amdgcn.rsq.f64 123void test_rsq_f64(global double* out, double a) 124{ 125 *out = __builtin_amdgcn_rsq(a); 126} 127 128// CHECK-LABEL: @test_rsq_clamp_f32 129// CHECK: call float @llvm.amdgcn.rsq.clamp.f32 130void test_rsq_clamp_f32(global float* out, float a) 131{ 132 *out = __builtin_amdgcn_rsq_clampf(a); 133} 134 135// CHECK-LABEL: @test_rsq_clamp_f64 136// CHECK: call double @llvm.amdgcn.rsq.clamp.f64 137void test_rsq_clamp_f64(global double* out, double a) 138{ 139 *out = __builtin_amdgcn_rsq_clamp(a); 140} 141 142// CHECK-LABEL: @test_sin_f32 143// CHECK: call float @llvm.amdgcn.sin.f32 144void test_sin_f32(global float* out, float a) 145{ 146 *out = __builtin_amdgcn_sinf(a); 147} 148 149// CHECK-LABEL: @test_cos_f32 150// CHECK: call float @llvm.amdgcn.cos.f32 151void test_cos_f32(global float* out, float a) 152{ 153 *out = __builtin_amdgcn_cosf(a); 154} 155 156// CHECK-LABEL: @test_log_clamp_f32 157// CHECK: call float @llvm.amdgcn.log.clamp.f32 158void test_log_clamp_f32(global float* out, float a) 159{ 160 *out = __builtin_amdgcn_log_clampf(a); 161} 162 163// CHECK-LABEL: @test_ldexp_f32 164// CHECK: call float @llvm.amdgcn.ldexp.f32 165void test_ldexp_f32(global float* out, float a, int b) 166{ 167 *out = __builtin_amdgcn_ldexpf(a, b); 168} 169 170// CHECK-LABEL: @test_ldexp_f64 171// CHECK: call double @llvm.amdgcn.ldexp.f64 172void test_ldexp_f64(global double* out, double a, int b) 173{ 174 *out = __builtin_amdgcn_ldexp(a, b); 175} 176 177// CHECK-LABEL: @test_frexp_mant_f32 178// CHECK: call float @llvm.amdgcn.frexp.mant.f32 179void test_frexp_mant_f32(global float* out, float a) 180{ 181 *out = __builtin_amdgcn_frexp_mantf(a); 182} 183 184// CHECK-LABEL: @test_frexp_mant_f64 185// CHECK: call double @llvm.amdgcn.frexp.mant.f64 186void test_frexp_mant_f64(global double* out, double a) 187{ 188 *out = __builtin_amdgcn_frexp_mant(a); 189} 190 191// CHECK-LABEL: @test_frexp_exp_f32 192// CHECK: call i32 @llvm.amdgcn.frexp.exp.i32.f32 193void test_frexp_exp_f32(global int* out, float a) 194{ 195 *out = __builtin_amdgcn_frexp_expf(a); 196} 197 198// CHECK-LABEL: @test_frexp_exp_f64 199// CHECK: call i32 @llvm.amdgcn.frexp.exp.i32.f64 200void test_frexp_exp_f64(global int* out, double a) 201{ 202 *out = __builtin_amdgcn_frexp_exp(a); 203} 204 205// CHECK-LABEL: @test_fract_f32 206// CHECK: call float @llvm.amdgcn.fract.f32 207void test_fract_f32(global int* out, float a) 208{ 209 *out = __builtin_amdgcn_fractf(a); 210} 211 212// CHECK-LABEL: @test_fract_f64 213// CHECK: call double @llvm.amdgcn.fract.f64 214void test_fract_f64(global int* out, double a) 215{ 216 *out = __builtin_amdgcn_fract(a); 217} 218 219// CHECK-LABEL: @test_lerp 220// CHECK: call i32 @llvm.amdgcn.lerp 221void test_lerp(global int* out, int a, int b, int c) 222{ 223 *out = __builtin_amdgcn_lerp(a, b, c); 224} 225 226// CHECK-LABEL: @test_sicmp_i32 227// CHECK: call i64 @llvm.amdgcn.icmp.i32(i32 %a, i32 %b, i32 32) 228void test_sicmp_i32(global ulong* out, int a, int b) 229{ 230 *out = __builtin_amdgcn_sicmp(a, b, 32); 231} 232 233// CHECK-LABEL: @test_uicmp_i32 234// CHECK: call i64 @llvm.amdgcn.icmp.i32(i32 %a, i32 %b, i32 32) 235void test_uicmp_i32(global ulong* out, uint a, uint b) 236{ 237 *out = __builtin_amdgcn_uicmp(a, b, 32); 238} 239 240// CHECK-LABEL: @test_sicmp_i64 241// CHECK: call i64 @llvm.amdgcn.icmp.i64(i64 %a, i64 %b, i32 38) 242void test_sicmp_i64(global ulong* out, long a, long b) 243{ 244 *out = __builtin_amdgcn_sicmpl(a, b, 39-1); 245} 246 247// CHECK-LABEL: @test_uicmp_i64 248// CHECK: call i64 @llvm.amdgcn.icmp.i64(i64 %a, i64 %b, i32 35) 249void test_uicmp_i64(global ulong* out, ulong a, ulong b) 250{ 251 *out = __builtin_amdgcn_uicmpl(a, b, 30+5); 252} 253 254// CHECK-LABEL: @test_ds_swizzle 255// CHECK: call i32 @llvm.amdgcn.ds.swizzle(i32 %a, i32 32) 256void test_ds_swizzle(global int* out, int a) 257{ 258 *out = __builtin_amdgcn_ds_swizzle(a, 32); 259} 260 261// CHECK-LABEL: @test_ds_permute 262// CHECK: call i32 @llvm.amdgcn.ds.permute(i32 %a, i32 %b) 263void test_ds_permute(global int* out, int a, int b) 264{ 265 out[0] = __builtin_amdgcn_ds_permute(a, b); 266} 267 268// CHECK-LABEL: @test_ds_bpermute 269// CHECK: call i32 @llvm.amdgcn.ds.bpermute(i32 %a, i32 %b) 270void test_ds_bpermute(global int* out, int a, int b) 271{ 272 *out = __builtin_amdgcn_ds_bpermute(a, b); 273} 274 275// CHECK-LABEL: @test_readfirstlane 276// CHECK: call i32 @llvm.amdgcn.readfirstlane(i32 %a) 277void test_readfirstlane(global int* out, int a) 278{ 279 *out = __builtin_amdgcn_readfirstlane(a); 280} 281 282// CHECK-LABEL: @test_readlane 283// CHECK: call i32 @llvm.amdgcn.readlane(i32 %a, i32 %b) 284void test_readlane(global int* out, int a, int b) 285{ 286 *out = __builtin_amdgcn_readlane(a, b); 287} 288 289// CHECK-LABEL: @test_fcmp_f32 290// CHECK: call i64 @llvm.amdgcn.fcmp.f32(float %a, float %b, i32 5) 291void test_fcmp_f32(global ulong* out, float a, float b) 292{ 293 *out = __builtin_amdgcn_fcmpf(a, b, 5); 294} 295 296// CHECK-LABEL: @test_fcmp_f64 297// CHECK: call i64 @llvm.amdgcn.fcmp.f64(double %a, double %b, i32 6) 298void test_fcmp_f64(global ulong* out, double a, double b) 299{ 300 *out = __builtin_amdgcn_fcmp(a, b, 3+3); 301} 302 303// CHECK-LABEL: @test_class_f32 304// CHECK: call i1 @llvm.amdgcn.class.f32 305void test_class_f32(global float* out, float a, int b) 306{ 307 *out = __builtin_amdgcn_classf(a, b); 308} 309 310// CHECK-LABEL: @test_class_f64 311// CHECK: call i1 @llvm.amdgcn.class.f64 312void test_class_f64(global double* out, double a, int b) 313{ 314 *out = __builtin_amdgcn_class(a, b); 315} 316 317// CHECK-LABEL: @test_buffer_wbinvl1 318// CHECK: call void @llvm.amdgcn.buffer.wbinvl1( 319void test_buffer_wbinvl1() 320{ 321 __builtin_amdgcn_buffer_wbinvl1(); 322} 323 324// CHECK-LABEL: @test_s_dcache_inv 325// CHECK: call void @llvm.amdgcn.s.dcache.inv( 326void test_s_dcache_inv() 327{ 328 __builtin_amdgcn_s_dcache_inv(); 329} 330 331// CHECK-LABEL: @test_s_waitcnt 332// CHECK: call void @llvm.amdgcn.s.waitcnt( 333void test_s_waitcnt() 334{ 335 __builtin_amdgcn_s_waitcnt(0); 336} 337 338// CHECK-LABEL: @test_s_sendmsg 339// CHECK: call void @llvm.amdgcn.s.sendmsg( 340void test_s_sendmsg() 341{ 342 __builtin_amdgcn_s_sendmsg(1, 0); 343} 344 345// CHECK-LABEL: @test_s_sendmsg_var 346// CHECK: call void @llvm.amdgcn.s.sendmsg( 347void test_s_sendmsg_var(int in) 348{ 349 __builtin_amdgcn_s_sendmsg(1, in); 350} 351 352// CHECK-LABEL: @test_s_sendmsghalt 353// CHECK: call void @llvm.amdgcn.s.sendmsghalt( 354void test_s_sendmsghalt() 355{ 356 __builtin_amdgcn_s_sendmsghalt(1, 0); 357} 358 359// CHECK-LABEL: @test_s_sendmsghalt 360// CHECK: call void @llvm.amdgcn.s.sendmsghalt( 361void test_s_sendmsghalt_var(int in) 362{ 363 __builtin_amdgcn_s_sendmsghalt(1, in); 364} 365 366// CHECK-LABEL: @test_s_barrier 367// CHECK: call void @llvm.amdgcn.s.barrier( 368void test_s_barrier() 369{ 370 __builtin_amdgcn_s_barrier(); 371} 372 373// CHECK-LABEL: @test_wave_barrier 374// CHECK: call void @llvm.amdgcn.wave.barrier( 375void test_wave_barrier() 376{ 377 __builtin_amdgcn_wave_barrier(); 378} 379 380// CHECK-LABEL: @test_s_memtime 381// CHECK: call i64 @llvm.amdgcn.s.memtime() 382void test_s_memtime(global ulong* out) 383{ 384 *out = __builtin_amdgcn_s_memtime(); 385} 386 387// CHECK-LABEL: @test_s_sleep 388// CHECK: call void @llvm.amdgcn.s.sleep(i32 1) 389// CHECK: call void @llvm.amdgcn.s.sleep(i32 15) 390void test_s_sleep() 391{ 392 __builtin_amdgcn_s_sleep(1); 393 __builtin_amdgcn_s_sleep(15); 394} 395 396// CHECK-LABEL: @test_s_incperflevel 397// CHECK: call void @llvm.amdgcn.s.incperflevel(i32 1) 398// CHECK: call void @llvm.amdgcn.s.incperflevel(i32 15) 399void test_s_incperflevel() 400{ 401 __builtin_amdgcn_s_incperflevel(1); 402 __builtin_amdgcn_s_incperflevel(15); 403} 404 405// CHECK-LABEL: @test_s_decperflevel 406// CHECK: call void @llvm.amdgcn.s.decperflevel(i32 1) 407// CHECK: call void @llvm.amdgcn.s.decperflevel(i32 15) 408void test_s_decperflevel() 409{ 410 __builtin_amdgcn_s_decperflevel(1); 411 __builtin_amdgcn_s_decperflevel(15); 412} 413 414// CHECK-LABEL: @test_cubeid( 415// CHECK: call float @llvm.amdgcn.cubeid(float %a, float %b, float %c) 416void test_cubeid(global float* out, float a, float b, float c) { 417 *out = __builtin_amdgcn_cubeid(a, b, c); 418} 419 420// CHECK-LABEL: @test_cubesc( 421// CHECK: call float @llvm.amdgcn.cubesc(float %a, float %b, float %c) 422void test_cubesc(global float* out, float a, float b, float c) { 423 *out = __builtin_amdgcn_cubesc(a, b, c); 424} 425 426// CHECK-LABEL: @test_cubetc( 427// CHECK: call float @llvm.amdgcn.cubetc(float %a, float %b, float %c) 428void test_cubetc(global float* out, float a, float b, float c) { 429 *out = __builtin_amdgcn_cubetc(a, b, c); 430} 431 432// CHECK-LABEL: @test_cubema( 433// CHECK: call float @llvm.amdgcn.cubema(float %a, float %b, float %c) 434void test_cubema(global float* out, float a, float b, float c) { 435 *out = __builtin_amdgcn_cubema(a, b, c); 436} 437 438// CHECK-LABEL: @test_read_exec( 439// CHECK: call i64 @llvm.read_register.i64(metadata ![[$EXEC:[0-9]+]]) #[[$READ_EXEC_ATTRS:[0-9]+]] 440void test_read_exec(global ulong* out) { 441 *out = __builtin_amdgcn_read_exec(); 442} 443 444// CHECK: declare i64 @llvm.read_register.i64(metadata) #[[$NOUNWIND_READONLY:[0-9]+]] 445 446// CHECK-LABEL: @test_read_exec_lo( 447// CHECK: call i32 @llvm.read_register.i32(metadata ![[$EXEC_LO:[0-9]+]]) #[[$READ_EXEC_ATTRS]] 448void test_read_exec_lo(global uint* out) { 449 *out = __builtin_amdgcn_read_exec_lo(); 450} 451 452// CHECK-LABEL: @test_read_exec_hi( 453// CHECK: call i32 @llvm.read_register.i32(metadata ![[$EXEC_HI:[0-9]+]]) #[[$READ_EXEC_ATTRS]] 454void test_read_exec_hi(global uint* out) { 455 *out = __builtin_amdgcn_read_exec_hi(); 456} 457 458// CHECK-LABEL: @test_dispatch_ptr 459// CHECK: call i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr() 460void test_dispatch_ptr(__constant unsigned char ** out) 461{ 462 *out = __builtin_amdgcn_dispatch_ptr(); 463} 464 465// CHECK-LABEL: @test_queue_ptr 466// CHECK: call i8 addrspace(4)* @llvm.amdgcn.queue.ptr() 467void test_queue_ptr(__constant unsigned char ** out) 468{ 469 *out = __builtin_amdgcn_queue_ptr(); 470} 471 472// CHECK-LABEL: @test_kernarg_segment_ptr 473// CHECK: call i8 addrspace(4)* @llvm.amdgcn.kernarg.segment.ptr() 474void test_kernarg_segment_ptr(__constant unsigned char ** out) 475{ 476 *out = __builtin_amdgcn_kernarg_segment_ptr(); 477} 478 479// CHECK-LABEL: @test_implicitarg_ptr 480// CHECK: call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr() 481void test_implicitarg_ptr(__constant unsigned char ** out) 482{ 483 *out = __builtin_amdgcn_implicitarg_ptr(); 484} 485 486// CHECK-LABEL: @test_get_group_id( 487// CHECK: tail call i32 @llvm.amdgcn.workgroup.id.x() 488// CHECK: tail call i32 @llvm.amdgcn.workgroup.id.y() 489// CHECK: tail call i32 @llvm.amdgcn.workgroup.id.z() 490void test_get_group_id(int d, global int *out) 491{ 492 switch (d) { 493 case 0: *out = __builtin_amdgcn_workgroup_id_x(); break; 494 case 1: *out = __builtin_amdgcn_workgroup_id_y(); break; 495 case 2: *out = __builtin_amdgcn_workgroup_id_z(); break; 496 default: *out = 0; 497 } 498} 499 500// CHECK-LABEL: @test_s_getreg( 501// CHECK: tail call i32 @llvm.amdgcn.s.getreg(i32 0) 502// CHECK: tail call i32 @llvm.amdgcn.s.getreg(i32 1) 503// CHECK: tail call i32 @llvm.amdgcn.s.getreg(i32 65535) 504void test_s_getreg(volatile global uint *out) 505{ 506 *out = __builtin_amdgcn_s_getreg(0); 507 *out = __builtin_amdgcn_s_getreg(1); 508 *out = __builtin_amdgcn_s_getreg(65535); 509} 510 511// CHECK-LABEL: @test_get_local_id( 512// CHECK: tail call i32 @llvm.amdgcn.workitem.id.x(), !range [[$WI_RANGE:![0-9]*]] 513// CHECK: tail call i32 @llvm.amdgcn.workitem.id.y(), !range [[$WI_RANGE]] 514// CHECK: tail call i32 @llvm.amdgcn.workitem.id.z(), !range [[$WI_RANGE]] 515void test_get_local_id(int d, global int *out) 516{ 517 switch (d) { 518 case 0: *out = __builtin_amdgcn_workitem_id_x(); break; 519 case 1: *out = __builtin_amdgcn_workitem_id_y(); break; 520 case 2: *out = __builtin_amdgcn_workitem_id_z(); break; 521 default: *out = 0; 522 } 523} 524 525// CHECK-LABEL: @test_fmed3_f32 526// CHECK: call float @llvm.amdgcn.fmed3.f32( 527void test_fmed3_f32(global float* out, float a, float b, float c) 528{ 529 *out = __builtin_amdgcn_fmed3f(a, b, c); 530} 531 532// CHECK-LABEL: @test_s_getpc 533// CHECK: call i64 @llvm.amdgcn.s.getpc() 534void test_s_getpc(global ulong* out) 535{ 536 *out = __builtin_amdgcn_s_getpc(); 537} 538 539// CHECK-LABEL: @test_ds_append_lds( 540// CHECK: call i32 @llvm.amdgcn.ds.append.p3i32(i32 addrspace(3)* %ptr, i1 false) 541kernel void test_ds_append_lds(global int* out, local int* ptr) { 542 *out = __builtin_amdgcn_ds_append(ptr); 543} 544 545// CHECK-LABEL: @test_ds_consume_lds( 546// CHECK: call i32 @llvm.amdgcn.ds.consume.p3i32(i32 addrspace(3)* %ptr, i1 false) 547kernel void test_ds_consume_lds(global int* out, local int* ptr) { 548 *out = __builtin_amdgcn_ds_consume(ptr); 549} 550 551// CHECK-DAG: [[$WI_RANGE]] = !{i32 0, i32 1024} 552// CHECK-DAG: attributes #[[$NOUNWIND_READONLY:[0-9]+]] = { nounwind readonly } 553// CHECK-DAG: attributes #[[$READ_EXEC_ATTRS]] = { convergent } 554// CHECK-DAG: ![[$EXEC]] = !{!"exec"} 555// CHECK-DAG: ![[$EXEC_LO]] = !{!"exec_lo"} 556// CHECK-DAG: ![[$EXEC_HI]] = !{!"exec_hi"} 557