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