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