1303 lines · plain
1// REQUIRES: amdgpu-registered-target2// 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 %s3// RUN: %clang_cc1 -cl-std=CL2.0 -triple spirv64-amd-amdhsa -emit-llvm -o - %s | FileCheck -enable-var-scope --check-prefixes=CHECK,CHECK-SPIRV %s4 5 6#pragma OPENCL EXTENSION cl_khr_fp64 : enable7 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_f6417// CHECK: {{.*}}call{{.*}} { double, i1 } @llvm.amdgcn.div.scale.f64(double %a, double %b, i1 true)18// CHECK-DAG: [[FLAG:%.+]] = extractvalue { double, i1 } %{{.+}}, 119// CHECK-DAG: [[VAL:%.+]] = extractvalue { double, i1 } %{{.+}}, 020// CHECK: [[FLAGEXT:%.+]] = zext i1 [[FLAG]] to i3221// 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 } %{{.+}}, 132// CHECK-DAG: [[VAL:%.+]] = extractvalue { float, i1 } %{{.+}}, 033// CHECK: [[FLAGEXT:%.+]] = zext i1 [[FLAG]] to i834// 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 } %{{.+}}, 145// CHECK-DAG: [[VAL:%.+]] = extractvalue { float, i1 } %{{.+}}, 046// CHECK: [[FLAGEXT:%.+]] = zext i1 [[FLAG]] to i847// 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 } %{{.+}}, 156// CHECK-DAG: [[VAL:%.+]] = extractvalue { float, i1 } %{{.+}}, 057// CHECK: [[FLAGEXT:%.+]] = zext i1 [[FLAG]] to i858// 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_f3266// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.div.fmas.f3267void 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_f6473// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.div.fmas.f6474void 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_f3280// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.div.fixup.f3281void 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_f6487// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.div.fixup.f6488void 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_f3294// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.trig.preop.f3295void 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_f64101// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.trig.preop.f64102void 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_f32108// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.rcp.f32109void test_rcp_f32(global float* out, float a)110{111 *out = __builtin_amdgcn_rcpf(a);112}113 114// CHECK-LABEL: @test_rcp_f64115// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.rcp.f64116void test_rcp_f64(global double* out, double a)117{118 *out = __builtin_amdgcn_rcp(a);119}120 121// CHECK-LABEL: @test_sqrt_f32122// CHECK: {{.*}}call{{.*}} float @llvm.{{((amdgcn.){0,1})}}sqrt.f32123void test_sqrt_f32(global float* out, float a)124{125 *out = __builtin_amdgcn_sqrtf(a);126}127 128// CHECK-LABEL: @test_sqrt_f64129// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.sqrt.f64130void test_sqrt_f64(global double* out, double a)131{132 *out = __builtin_amdgcn_sqrt(a);133}134 135// CHECK-LABEL: @test_rsq_f32136// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.rsq.f32137void test_rsq_f32(global float* out, float a)138{139 *out = __builtin_amdgcn_rsqf(a);140}141 142// CHECK-LABEL: @test_rsq_f64143// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.rsq.f64144void test_rsq_f64(global double* out, double a)145{146 *out = __builtin_amdgcn_rsq(a);147}148 149// CHECK-LABEL: @test_rsq_clamp_f32150// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.rsq.clamp.f32151void 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_f64157// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.rsq.clamp.f64158void test_rsq_clamp_f64(global double* out, double a)159{160 *out = __builtin_amdgcn_rsq_clamp(a);161}162 163// CHECK-LABEL: @test_sin_f32164// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.sin.f32165void test_sin_f32(global float* out, float a)166{167 *out = __builtin_amdgcn_sinf(a);168}169 170// CHECK-LABEL: @test_cos_f32171// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.cos.f32172void test_cos_f32(global float* out, float a)173{174 *out = __builtin_amdgcn_cosf(a);175}176 177// CHECK-LABEL: @test_log_f32178// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.log.f32179void test_log_f32(global float* out, float a)180{181 *out = __builtin_amdgcn_logf(a);182}183 184// CHECK-LABEL: @test_exp2_f32185// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.exp2.f32186void test_exp2_f32(global float* out, float a)187{188 *out = __builtin_amdgcn_exp2f(a);189}190 191// CHECK-LABEL: @test_log_clamp_f32192// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.log.clamp.f32193void test_log_clamp_f32(global float* out, float a)194{195 *out = __builtin_amdgcn_log_clampf(a);196}197 198// CHECK-LABEL: @test_ldexp_f32199// CHECK: {{.*}}call{{.*}} float @llvm.ldexp.f32.i32200void 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_f64206// CHECK: {{.*}}call{{.*}} double @llvm.ldexp.f64.i32207void 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_f32213// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.frexp.mant.f32214void 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_f64220// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.frexp.mant.f64221void 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_f32227// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.frexp.exp.i32.f32228void 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_f64234// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.frexp.exp.i32.f64235void test_frexp_exp_f64(global int* out, double a)236{237 *out = __builtin_amdgcn_frexp_exp(a);238}239 240// CHECK-LABEL: @test_fract_f32241// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.fract.f32242void test_fract_f32(global int* out, float a)243{244 *out = __builtin_amdgcn_fractf(a);245}246 247// CHECK-LABEL: @test_fract_f64248// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.fract.f64249void test_fract_f64(global int* out, double a)250{251 *out = __builtin_amdgcn_fract(a);252}253 254// CHECK-LABEL: @test_sicmp_i32255// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.icmp.i64.i32(i32 %a, i32 %b, i32 32)256void test_sicmp_i32(global ulong* out, int a, int b)257{258 *out = __builtin_amdgcn_sicmp(a, b, 32);259}260 261// CHECK-LABEL: @test_uicmp_i32262// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.icmp.i64.i32(i32 %a, i32 %b, i32 32)263void test_uicmp_i32(global ulong* out, uint a, uint b)264{265 *out = __builtin_amdgcn_uicmp(a, b, 32);266}267 268// CHECK-LABEL: @test_sicmp_i64269// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.icmp.i64.i64(i64 %a, i64 %b, i32 38)270void test_sicmp_i64(global ulong* out, long a, long b)271{272 *out = __builtin_amdgcn_sicmpl(a, b, 39-1);273}274 275// CHECK-LABEL: @test_uicmp_i64276// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.icmp.i64.i64(i64 %a, i64 %b, i32 35)277void test_uicmp_i64(global ulong* out, ulong a, ulong b)278{279 *out = __builtin_amdgcn_uicmpl(a, b, 30+5);280}281 282// CHECK-LABEL: @test_ds_swizzle283// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.ds.swizzle(i32 %a, i32 32)284void test_ds_swizzle(global int* out, int a)285{286 *out = __builtin_amdgcn_ds_swizzle(a, 32);287}288 289// CHECK-LABEL: @test_ds_permute290// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.ds.permute(i32 %a, i32 %b)291void test_ds_permute(global int* out, int a, int b)292{293 out[0] = __builtin_amdgcn_ds_permute(a, b);294}295 296// CHECK-LABEL: @test_ds_bpermute297// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.ds.bpermute(i32 %a, i32 %b)298void test_ds_bpermute(global int* out, int a, int b)299{300 *out = __builtin_amdgcn_ds_bpermute(a, b);301}302 303// CHECK-LABEL: @test_readfirstlane304// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.readfirstlane.i32(i32 %a)305void test_readfirstlane(global int* out, int a)306{307 *out = __builtin_amdgcn_readfirstlane(a);308}309 310// CHECK-LABEL: @test_readlane311// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.readlane.i32(i32 %a, i32 %b)312void test_readlane(global int* out, int a, int b)313{314 *out = __builtin_amdgcn_readlane(a, b);315}316 317// CHECK-LABEL: @test_fcmp_f32318// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.fcmp.i64.f32(float %a, float %b, i32 5)319void test_fcmp_f32(global ulong* out, float a, float b)320{321 *out = __builtin_amdgcn_fcmpf(a, b, 5);322}323 324// CHECK-LABEL: @test_fcmp_f64325// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.fcmp.i64.f64(double %a, double %b, i32 6)326void test_fcmp_f64(global ulong* out, double a, double b)327{328 *out = __builtin_amdgcn_fcmp(a, b, 3+3);329}330 331// CHECK-LABEL: @test_class_f32332// CHECK: {{.*}}call{{.*}} i1 @llvm.amdgcn.class.f32333void test_class_f32(global float* out, float a, int b)334{335 *out = __builtin_amdgcn_classf(a, b);336}337 338// CHECK-LABEL: @test_class_f64339// CHECK: {{.*}}call{{.*}} i1 @llvm.amdgcn.class.f64340void test_class_f64(global double* out, double a, int b)341{342 *out = __builtin_amdgcn_class(a, b);343}344 345// CHECK-LABEL: @test_buffer_wbinvl1346// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.buffer.wbinvl1(347void test_buffer_wbinvl1()348{349 __builtin_amdgcn_buffer_wbinvl1();350}351 352// CHECK-LABEL: @test_s_dcache_inv353// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.dcache.inv(354void test_s_dcache_inv()355{356 __builtin_amdgcn_s_dcache_inv();357}358 359// CHECK-LABEL: @test_s_waitcnt360// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.waitcnt(361void test_s_waitcnt()362{363 __builtin_amdgcn_s_waitcnt(0);364}365 366// CHECK-LABEL: @test_s_sendmsg367// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.sendmsg(368void test_s_sendmsg()369{370 __builtin_amdgcn_s_sendmsg(1, 0);371}372 373// CHECK-LABEL: @test_s_sendmsg_var374// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.sendmsg(375void test_s_sendmsg_var(int in)376{377 __builtin_amdgcn_s_sendmsg(1, in);378}379 380// CHECK-LABEL: @test_s_sendmsghalt381// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.sendmsghalt(382void test_s_sendmsghalt()383{384 __builtin_amdgcn_s_sendmsghalt(1, 0);385}386 387// CHECK-LABEL: @test_s_sendmsghalt388// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.sendmsghalt(389void test_s_sendmsghalt_var(int in)390{391 __builtin_amdgcn_s_sendmsghalt(1, in);392}393 394// CHECK-LABEL: @test_wave_reduce_add_u32_default395// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.add.i32(396void test_wave_reduce_add_u32_default(global int* out, int in)397{398 *out = __builtin_amdgcn_wave_reduce_add_u32(in, 0);399}400 401// CHECK-LABEL: @test_wave_reduce_add_u64_default402// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.add.i64(403void test_wave_reduce_add_u64_default(global int* out, long in)404{405 *out = __builtin_amdgcn_wave_reduce_add_u64(in, 0);406}407 408// CHECK-LABEL: @test_wave_reduce_fadd_f32_default409// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.wave.reduce.fadd.f32(410void test_wave_reduce_fadd_f32_default(global float* out, float in)411{412 *out = __builtin_amdgcn_wave_reduce_fadd_f32(in, 0);413}414 415// CHECK-LABEL: @test_wave_reduce_add_u32_iterative416// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.add.i32(417void test_wave_reduce_add_u32_iterative(global int* out, int in)418{419 *out = __builtin_amdgcn_wave_reduce_add_u32(in, 1);420}421 422// CHECK-LABEL: @test_wave_reduce_add_u64_iterative423// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.add.i64(424void test_wave_reduce_add_u64_iterative(global int* out, long in)425{426 *out = __builtin_amdgcn_wave_reduce_add_u64(in, 1);427}428 429// CHECK-LABEL: @test_wave_reduce_fadd_f32_iterative430// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.wave.reduce.fadd.f32(431void test_wave_reduce_fadd_f32_iterative(global float* out, float in)432{433 *out = __builtin_amdgcn_wave_reduce_fadd_f32(in, 0);434}435 436// CHECK-LABEL: @test_wave_reduce_add_u32_dpp437// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.add.i32(438void test_wave_reduce_add_u32_dpp(global int* out, int in)439{440 *out = __builtin_amdgcn_wave_reduce_add_u32(in, 2);441}442 443// CHECK-LABEL: @test_wave_reduce_add_u64_dpp444// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.add.i64(445void test_wave_reduce_add_u64_dpp(global int* out, long in)446{447 *out = __builtin_amdgcn_wave_reduce_add_u64(in, 2);448}449 450// CHECK-LABEL: @test_wave_reduce_fadd_f32_dpp451// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.wave.reduce.fadd.f32(452void test_wave_reduce_fadd_f32_dpp(global float* out, float in)453{454 *out = __builtin_amdgcn_wave_reduce_fadd_f32(in, 0);455}456 457// CHECK-LABEL: @test_wave_reduce_sub_u32_default458// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.sub.i32(459void test_wave_reduce_sub_u32_default(global int* out, int in)460{461 *out = __builtin_amdgcn_wave_reduce_sub_u32(in, 0);462}463 464// CHECK-LABEL: @test_wave_reduce_sub_u64_default465// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.sub.i64(466void test_wave_reduce_sub_u64_default(global int* out, long in)467{468 *out = __builtin_amdgcn_wave_reduce_sub_u64(in, 0);469}470 471// CHECK-LABEL: @test_wave_reduce_fsub_f32_default472// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.wave.reduce.fsub.f32(473void test_wave_reduce_fsub_f32_default(global float* out, float in)474{475 *out = __builtin_amdgcn_wave_reduce_fsub_f32(in, 0);476}477 478// CHECK-LABEL: @test_wave_reduce_sub_u32_iterative479// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.sub.i32(480void test_wave_reduce_sub_u32_iterative(global int* out, int in)481{482 *out = __builtin_amdgcn_wave_reduce_sub_u32(in, 1);483}484 485// CHECK-LABEL: @test_wave_reduce_sub_u64_iterative486// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.sub.i64(487void test_wave_reduce_sub_u64_iterative(global int* out, long in)488{489 *out = __builtin_amdgcn_wave_reduce_sub_u64(in, 1);490}491 492// CHECK-LABEL: @test_wave_reduce_fsub_f32_iterative493// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.wave.reduce.fsub.f32(494void test_wave_reduce_fsub_f32_iterative(global float* out, float in)495{496 *out = __builtin_amdgcn_wave_reduce_fsub_f32(in, 0);497}498 499// CHECK-LABEL: @test_wave_reduce_sub_u32_dpp500// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.sub.i32(501void test_wave_reduce_sub_u32_dpp(global int* out, int in)502{503 *out = __builtin_amdgcn_wave_reduce_sub_u32(in, 2);504}505 506// CHECK-LABEL: @test_wave_reduce_sub_u64_dpp507// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.sub.i64(508void test_wave_reduce_sub_u64_dpp(global int* out, long in)509{510 *out = __builtin_amdgcn_wave_reduce_sub_u64(in, 2);511}512 513// CHECK-LABEL: @test_wave_reduce_fsub_f32_dpp514// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.wave.reduce.fsub.f32(515void test_wave_reduce_fsub_f32_dpp(global float* out, float in)516{517 *out = __builtin_amdgcn_wave_reduce_fsub_f32(in, 0);518}519 520// CHECK-LABEL: @test_wave_reduce_and_b32_default521// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.and.i32(522void test_wave_reduce_and_b32_default(global int* out, int in)523{524 *out = __builtin_amdgcn_wave_reduce_and_b32(in, 0);525}526 527// CHECK-LABEL: @test_wave_reduce_and_b64_default528// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.and.i64(529void test_wave_reduce_and_b64_default(global int* out, long in)530{531 *out = __builtin_amdgcn_wave_reduce_and_b64(in, 0);532}533 534// CHECK-LABEL: @test_wave_reduce_and_b32_iterative535// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.and.i32(536void test_wave_reduce_and_b32_iterative(global int* out, int in)537{538 *out = __builtin_amdgcn_wave_reduce_and_b32(in, 1);539}540 541// CHECK-LABEL: @test_wave_reduce_and_b64_iterative542// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.and.i64(543void test_wave_reduce_and_b64_iterative(global int* out, long in)544{545 *out = __builtin_amdgcn_wave_reduce_and_b64(in, 1);546}547 548// CHECK-LABEL: @test_wave_reduce_and_b32_dpp549// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.and.i32(550void test_wave_reduce_and_b32_dpp(global int* out, int in)551{552 *out = __builtin_amdgcn_wave_reduce_and_b32(in, 2);553}554 555// CHECK-LABEL: @test_wave_reduce_and_b64_dpp556// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.and.i64(557void test_wave_reduce_and_b64_dpp(global int* out, long in)558{559 *out = __builtin_amdgcn_wave_reduce_and_b64(in, 2);560}561 562// CHECK-LABEL: @test_wave_reduce_or_b32_default563// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.or.i32(564void test_wave_reduce_or_b32_default(global int* out, int in)565{566 *out = __builtin_amdgcn_wave_reduce_or_b32(in, 0);567}568 569// CHECK-LABEL: @test_wave_reduce_or_b64_default570// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.or.i64(571void test_wave_reduce_or_b64_default(global int* out, long in)572{573 *out = __builtin_amdgcn_wave_reduce_or_b64(in, 0);574}575 576// CHECK-LABEL: @test_wave_reduce_or_b32_iterative577// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.or.i32(578void test_wave_reduce_or_b32_iterative(global int* out, int in)579{580 *out = __builtin_amdgcn_wave_reduce_or_b32(in, 1);581}582 583// CHECK-LABEL: @test_wave_reduce_or_b64_iterative584// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.or.i64(585void test_wave_reduce_or_b64_iterative(global int* out, long in)586{587 *out = __builtin_amdgcn_wave_reduce_or_b64(in, 1);588}589 590// CHECK-LABEL: @test_wave_reduce_or_b32_dpp591// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.or.i32(592void test_wave_reduce_or_b32_dpp(global int* out, int in)593{594 *out = __builtin_amdgcn_wave_reduce_or_b32(in, 2);595}596 597// CHECK-LABEL: @test_wave_reduce_or_b64_dpp598// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.or.i64(599void test_wave_reduce_or_b64_dpp(global int* out, long in)600{601 *out = __builtin_amdgcn_wave_reduce_or_b64(in, 2);602}603 604// CHECK-LABEL: @test_wave_reduce_xor_b32_default605// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.xor.i32(606void test_wave_reduce_xor_b32_default(global int* out, int in)607{608 *out = __builtin_amdgcn_wave_reduce_xor_b32(in, 0);609}610 611// CHECK-LABEL: @test_wave_reduce_xor_b64_default612// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.xor.i64(613void test_wave_reduce_xor_b64_default(global int* out, long in)614{615 *out = __builtin_amdgcn_wave_reduce_xor_b64(in, 0);616}617 618// CHECK-LABEL: @test_wave_reduce_xor_b32_iterative619// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.xor.i32(620void test_wave_reduce_xor_b32_iterative(global int* out, int in)621{622 *out = __builtin_amdgcn_wave_reduce_xor_b32(in, 1);623}624 625// CHECK-LABEL: @test_wave_reduce_xor_b64_iterative626// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.xor.i64(627void test_wave_reduce_xor_b64_iterative(global int* out, long in)628{629 *out = __builtin_amdgcn_wave_reduce_xor_b64(in, 1);630}631 632// CHECK-LABEL: @test_wave_reduce_xor_b32_dpp633// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.xor.i32(634void test_wave_reduce_xor_b32_dpp(global int* out, int in)635{636 *out = __builtin_amdgcn_wave_reduce_xor_b32(in, 2);637}638 639// CHECK-LABEL: @test_wave_reduce_xor_b64_dpp640// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.xor.i64(641void test_wave_reduce_xor_b64_dpp(global int* out, long in)642{643 *out = __builtin_amdgcn_wave_reduce_xor_b64(in, 2);644}645 646// CHECK-LABEL: @test_wave_reduce_min_i32_default647// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.min.i32(648void test_wave_reduce_min_i32_default(global int* out, int in)649{650 *out = __builtin_amdgcn_wave_reduce_min_i32(in, 0);651}652 653// CHECK-LABEL: @test_wave_reduce_min_i64_default654// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.min.i64(655void test_wave_reduce_min_i64_default(global int* out, long in)656{657 *out = __builtin_amdgcn_wave_reduce_min_i64(in, 0);658}659 660// CHECK-LABEL: @test_wave_reduce_fmin_f32_default661// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.wave.reduce.fmin.f32(662void test_wave_reduce_fmin_f32_default(global float* out, float in)663{664 *out = __builtin_amdgcn_wave_reduce_fmin_f32(in, 0);665}666 667// CHECK-LABEL: @test_wave_reduce_min_i32_iterative668// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.min.i32(669void test_wave_reduce_min_i32_iterative(global int* out, int in)670{671 *out = __builtin_amdgcn_wave_reduce_min_i32(in, 1);672}673 674// CHECK-LABEL: @test_wave_reduce_min_i64_iterative675// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.min.i64(676void test_wave_reduce_min_i64_iterative(global int* out, long in)677{678 *out = __builtin_amdgcn_wave_reduce_min_i64(in, 1);679}680 681// CHECK-LABEL: @test_wave_reduce_fmin_f32_iterative682// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.wave.reduce.fmin.f32(683void test_wave_reduce_fmin_f32_iterative(global float* out, float in)684{685 *out = __builtin_amdgcn_wave_reduce_fmin_f32(in, 0);686}687 688// CHECK-LABEL: @test_wave_reduce_min_i32_dpp689// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.min.i32(690void test_wave_reduce_min_i32_dpp(global int* out, int in)691{692 *out = __builtin_amdgcn_wave_reduce_min_i32(in, 2);693}694 695// CHECK-LABEL: @test_wave_reduce_min_i64_dpp696// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.min.i64(697void test_wave_reduce_min_i64_dpp(global int* out, long in)698{699 *out = __builtin_amdgcn_wave_reduce_min_i64(in, 2);700}701 702// CHECK-LABEL: @test_wave_reduce_fmin_f32_dpp703// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.wave.reduce.fmin.f32(704void test_wave_reduce_fmin_f32_dpp(global float* out, float in)705{706 *out = __builtin_amdgcn_wave_reduce_fmin_f32(in, 0);707}708 709// CHECK-LABEL: @test_wave_reduce_min_u32_default710// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.umin.i32(711void test_wave_reduce_min_u32_default(global int* out, int in)712{713 *out = __builtin_amdgcn_wave_reduce_min_u32(in, 0);714}715 716// CHECK-LABEL: @test_wave_reduce_min_u64_default717// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.umin.i64(718void test_wave_reduce_min_u64_default(global int* out, long in)719{720 *out = __builtin_amdgcn_wave_reduce_min_u64(in, 0);721}722 723// CHECK-LABEL: @test_wave_reduce_min_u32_iterative724// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.umin.i32(725void test_wave_reduce_min_u32_iterative(global int* out, int in)726{727 *out = __builtin_amdgcn_wave_reduce_min_u32(in, 1);728}729 730// CHECK-LABEL: @test_wave_reduce_min_u64_iterative731// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.umin.i64(732void test_wave_reduce_min_u64_iterative(global int* out, long in)733{734 *out = __builtin_amdgcn_wave_reduce_min_u64(in, 1);735}736 737// CHECK-LABEL: @test_wave_reduce_min_u32_dpp738// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.umin.i32(739void test_wave_reduce_min_u32_dpp(global int* out, int in)740{741 *out = __builtin_amdgcn_wave_reduce_min_u32(in, 2);742}743 744// CHECK-LABEL: @test_wave_reduce_min_u64_dpp745// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.umin.i64(746void test_wave_reduce_min_u64_dpp(global int* out, long in)747{748 *out = __builtin_amdgcn_wave_reduce_min_u64(in, 2);749}750 751// CHECK-LABEL: @test_wave_reduce_max_i32_default752// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.max.i32(753void test_wave_reduce_max_i32_default(global int* out, int in)754{755 *out = __builtin_amdgcn_wave_reduce_max_i32(in, 0);756}757 758// CHECK-LABEL: @test_wave_reduce_max_i64_default759// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.max.i64(760void test_wave_reduce_max_i64_default(global int* out, long in)761{762 *out = __builtin_amdgcn_wave_reduce_max_i64(in, 0);763}764 765// CHECK-LABEL: @test_wave_reduce_fmax_f32_default766// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.wave.reduce.fmax.f32(767void test_wave_reduce_fmax_f32_default(global float* out, float in)768{769 *out = __builtin_amdgcn_wave_reduce_fmax_f32(in, 0);770}771 772// CHECK-LABEL: @test_wave_reduce_max_i32_iterative773// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.max.i32(774void test_wave_reduce_max_i32_iterative(global int* out, int in)775{776 *out = __builtin_amdgcn_wave_reduce_max_i32(in, 1);777}778 779// CHECK-LABEL: @test_wave_reduce_max_i64_iterative780// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.max.i64(781void test_wave_reduce_max_i64_iterative(global int* out, long in)782{783 *out = __builtin_amdgcn_wave_reduce_max_i64(in, 1);784}785 786// CHECK-LABEL: @test_wave_reduce_fmax_f32_iterative787// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.wave.reduce.fmax.f32(788void test_wave_reduce_fmax_f32_iterative(global float* out, float in)789{790 *out = __builtin_amdgcn_wave_reduce_fmax_f32(in, 0);791}792 793// CHECK-LABEL: @test_wave_reduce_max_i32_dpp794// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.max.i32(795void test_wave_reduce_max_i32_dpp(global int* out, int in)796{797 *out = __builtin_amdgcn_wave_reduce_max_i32(in, 2);798}799 800// CHECK-LABEL: @test_wave_reduce_max_i64_dpp801// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.max.i64(802void test_wave_reduce_max_i64_dpp(global int* out, long in)803{804 *out = __builtin_amdgcn_wave_reduce_max_i64(in, 2);805}806 807// CHECK-LABEL: @test_wave_reduce_fmax_f32_dpp808// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.wave.reduce.fmax.f32(809void test_wave_reduce_fmax_f32_dpp(global float* out, float in)810{811 *out = __builtin_amdgcn_wave_reduce_fmax_f32(in, 0);812}813 814// CHECK-LABEL: @test_wave_reduce_max_u32_default815// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.umax.i32(816void test_wave_reduce_max_u32_default(global int* out, int in)817{818 *out = __builtin_amdgcn_wave_reduce_max_u32(in, 0);819}820 821// CHECK-LABEL: @test_wave_reduce_max_u64_default822// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.umax.i64(823void test_wave_reduce_max_u64_default(global int* out, long in)824{825 *out = __builtin_amdgcn_wave_reduce_max_u64(in, 0);826}827 828// CHECK-LABEL: @test_wave_reduce_max_u32_iterative829// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.umax.i32(830void test_wave_reduce_max_u32_iterative(global int* out, int in)831{832 *out = __builtin_amdgcn_wave_reduce_max_u32(in, 1);833}834 835// CHECK-LABEL: @test_wave_reduce_max_u64_iterative836// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.umax.i64(837void test_wave_reduce_max_u64_iterative(global int* out, long in)838{839 *out = __builtin_amdgcn_wave_reduce_max_u64(in, 1);840}841 842// CHECK-LABEL: @test_wave_reduce_max_u32_dpp843// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.umax.i32(844void test_wave_reduce_max_u32_dpp(global int* out, int in)845{846 *out = __builtin_amdgcn_wave_reduce_max_u32(in, 2);847}848 849// CHECK-LABEL: @test_wave_reduce_max_u64_dpp850// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.umax.i64(851void test_wave_reduce_max_u64_dpp(global int* out, long in)852{853 *out = __builtin_amdgcn_wave_reduce_max_u64(in, 2);854}855 856// CHECK-LABEL: @test_s_barrier857// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.barrier(858void test_s_barrier()859{860 __builtin_amdgcn_s_barrier();861}862 863// CHECK-LABEL: @test_wave_barrier864// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.wave.barrier(865void test_wave_barrier()866{867 __builtin_amdgcn_wave_barrier();868}869 870// CHECK-LABEL: @test_sched_barrier871// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.sched.barrier(i32 0)872// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.sched.barrier(i32 1)873// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.sched.barrier(i32 4)874// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.sched.barrier(i32 15)875void test_sched_barrier()876{877 __builtin_amdgcn_sched_barrier(0);878 __builtin_amdgcn_sched_barrier(1);879 __builtin_amdgcn_sched_barrier(4);880 __builtin_amdgcn_sched_barrier(15);881}882 883// CHECK-LABEL: @test_sched_group_barrier884// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.sched.group.barrier(i32 0, i32 1, i32 2)885// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.sched.group.barrier(i32 1, i32 2, i32 4)886// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.sched.group.barrier(i32 4, i32 8, i32 16)887// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.sched.group.barrier(i32 15, i32 10000, i32 -1)888void test_sched_group_barrier()889{890 __builtin_amdgcn_sched_group_barrier(0, 1, 2);891 __builtin_amdgcn_sched_group_barrier(1, 2, 4);892 __builtin_amdgcn_sched_group_barrier(4, 8, 16);893 __builtin_amdgcn_sched_group_barrier(15, 10000, -1);894}895 896// CHECK-LABEL: @test_iglp_opt897// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.iglp.opt(i32 0)898// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.iglp.opt(i32 1)899// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.iglp.opt(i32 4)900// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.iglp.opt(i32 15)901void test_iglp_opt()902{903 __builtin_amdgcn_iglp_opt(0);904 __builtin_amdgcn_iglp_opt(1);905 __builtin_amdgcn_iglp_opt(4);906 __builtin_amdgcn_iglp_opt(15);907}908 909// CHECK-LABEL: @test_s_sleep910// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.sleep(i32 1)911// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.sleep(i32 15)912void test_s_sleep()913{914 __builtin_amdgcn_s_sleep(1);915 __builtin_amdgcn_s_sleep(15);916}917 918// CHECK-LABEL: @test_s_incperflevel919// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.incperflevel(i32 1)920// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.incperflevel(i32 15)921void test_s_incperflevel()922{923 __builtin_amdgcn_s_incperflevel(1);924 __builtin_amdgcn_s_incperflevel(15);925}926 927// CHECK-LABEL: @test_s_decperflevel928// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.decperflevel(i32 1)929// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.decperflevel(i32 15)930void test_s_decperflevel()931{932 __builtin_amdgcn_s_decperflevel(1);933 __builtin_amdgcn_s_decperflevel(15);934}935 936// CHECK-LABEL: @test_s_setprio937// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.setprio(i16 0)938// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.setprio(i16 3)939void test_s_setprio()940{941 __builtin_amdgcn_s_setprio(0);942 __builtin_amdgcn_s_setprio(3);943}944 945// CHECK-LABEL: @test_read_exec(946// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.ballot.i64(i1 true)947void test_read_exec(global ulong* out) {948 *out = __builtin_amdgcn_read_exec();949}950 951// CHECK: declare i64 @llvm.amdgcn.ballot.i64(i1){{.*}} #[[$NOUNWIND_READONLY:[0-9]+]]952 953// CHECK-LABEL: @test_read_exec_lo(954// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.ballot.i32(i1 true)955void test_read_exec_lo(global uint* out) {956 *out = __builtin_amdgcn_read_exec_lo();957}958 959// CHECK: declare i32 @llvm.amdgcn.ballot.i32(i1){{.*}} #[[$NOUNWIND_READONLY:[0-9]+]]960 961// CHECK-LABEL: @test_read_exec_hi(962// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.ballot.i64(i1 true)963// CHECK: lshr i64 [[A:%.*]], 32964// CHECK: trunc nuw i64 [[B:%.*]] to i32965void test_read_exec_hi(global uint* out) {966 *out = __builtin_amdgcn_read_exec_hi();967}968 969// CHECK-LABEL: @test_dispatch_ptr970// CHECK: {{.*}}call align 4 dereferenceable(64){{.*}} ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()971#if !defined(__SPIRV__)972void test_dispatch_ptr(__constant unsigned char ** out)973#else974void test_dispatch_ptr(__attribute__((address_space(4))) unsigned char ** out)975#endif976{977 *out = __builtin_amdgcn_dispatch_ptr();978}979 980// CHECK-LABEL: @test_queue_ptr981// CHECK: {{.*}}call{{.*}} ptr addrspace(4) @llvm.amdgcn.queue.ptr()982#if !defined(__SPIRV__)983void test_queue_ptr(__constant unsigned char ** out)984#else985void test_queue_ptr(__attribute__((address_space(4))) unsigned char ** out)986#endif987{988 *out = __builtin_amdgcn_queue_ptr();989}990 991// CHECK-LABEL: @test_kernarg_segment_ptr992// CHECK: {{.*}}call{{.*}} ptr addrspace(4) @llvm.amdgcn.kernarg.segment.ptr()993#if !defined(__SPIRV__)994void test_kernarg_segment_ptr(__constant unsigned char ** out)995#else996void test_kernarg_segment_ptr(__attribute__((address_space(4))) unsigned char ** out)997#endif998{999 *out = __builtin_amdgcn_kernarg_segment_ptr();1000}1001 1002// CHECK-LABEL: @test_implicitarg_ptr1003// CHECK: {{.*}}call{{.*}} ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()1004#if !defined(__SPIRV__)1005void test_implicitarg_ptr(__constant unsigned char ** out)1006#else1007void test_implicitarg_ptr(__attribute__((address_space(4))) unsigned char ** out)1008#endif1009{1010 *out = __builtin_amdgcn_implicitarg_ptr();1011}1012 1013// CHECK-LABEL: @test_get_group_id(1014// CHECK: tail call{{.*}} i32 @llvm.amdgcn.workgroup.id.x()1015// CHECK: tail call{{.*}} i32 @llvm.amdgcn.workgroup.id.y()1016// CHECK: tail call{{.*}} i32 @llvm.amdgcn.workgroup.id.z()1017void test_get_group_id(int d, global int *out)1018{1019 switch (d) {1020 case 0: *out = __builtin_amdgcn_workgroup_id_x(); break;1021 case 1: *out = __builtin_amdgcn_workgroup_id_y(); break;1022 case 2: *out = __builtin_amdgcn_workgroup_id_z(); break;1023 default: *out = 0;1024 }1025}1026 1027// CHECK-LABEL: @test_s_getreg(1028// CHECK: tail call{{.*}} i32 @llvm.amdgcn.s.getreg(i32 0)1029// CHECK: tail call{{.*}} i32 @llvm.amdgcn.s.getreg(i32 1)1030// CHECK: tail call{{.*}} i32 @llvm.amdgcn.s.getreg(i32 65535)1031void test_s_getreg(volatile global uint *out)1032{1033 *out = __builtin_amdgcn_s_getreg(0);1034 *out = __builtin_amdgcn_s_getreg(1);1035 *out = __builtin_amdgcn_s_getreg(65535);1036}1037 1038// CHECK-LABEL: @test_get_local_id(1039// CHECK: tail call{{.*}} i32 @llvm.amdgcn.workitem.id.x()1040// CHECK: tail call{{.*}} i32 @llvm.amdgcn.workitem.id.y()1041// CHECK: tail call{{.*}} i32 @llvm.amdgcn.workitem.id.z()1042void test_get_local_id(int d, global int *out)1043{1044 switch (d) {1045 case 0: *out = __builtin_amdgcn_workitem_id_x(); break;1046 case 1: *out = __builtin_amdgcn_workitem_id_y(); break;1047 case 2: *out = __builtin_amdgcn_workitem_id_z(); break;1048 default: *out = 0;1049 }1050}1051 1052// CHECK: declare noundef range(i32 0, 1024) i32 @llvm.amdgcn.workitem.id.x()1053// CHECK: declare noundef range(i32 0, 1024) i32 @llvm.amdgcn.workitem.id.y()1054// CHECK: declare noundef range(i32 0, 1024) i32 @llvm.amdgcn.workitem.id.z()1055 1056// CHECK-LABEL: @test_get_workgroup_size(1057// CHECK: {{.*}}call align 8 dereferenceable(256){{.*}} ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()1058// CHECK: getelementptr inbounds nuw i8, ptr addrspace(4) %{{.*}}, i64 121059// CHECK: load i16, ptr addrspace(4) %{{.*}}, align 4, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef1060// CHECK: getelementptr inbounds nuw i8, ptr addrspace(4) %{{.*}}, i64 141061// CHECK: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef1062// CHECK: getelementptr inbounds nuw i8, ptr addrspace(4) %{{.*}}, i64 161063// CHECK: load i16, ptr addrspace(4) %{{.*}}, align 8, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef1064void test_get_workgroup_size(int d, global int *out)1065{1066 switch (d) {1067 case 0: *out = __builtin_amdgcn_workgroup_size_x() + 1; break;1068 case 1: *out = __builtin_amdgcn_workgroup_size_y(); break;1069 case 2: *out = __builtin_amdgcn_workgroup_size_z(); break;1070 default: *out = 0;1071 }1072}1073 1074// CHECK-LABEL: @test_get_grid_size(1075// CHECK: {{.*}}call align 4 dereferenceable(64){{.*}} ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()1076// CHECK: getelementptr inbounds nuw i8, ptr addrspace(4) %{{.*}}, i64 %{{.+}}1077// CHECK: load i32, ptr addrspace(4) %{{.*}}, align 4, !range [[$GRID_RANGE:![0-9]+]], !invariant.load1078void test_get_grid_size(int d, global int *out)1079{1080 switch (d) {1081 case 0: *out = __builtin_amdgcn_grid_size_x(); break;1082 case 1: *out = __builtin_amdgcn_grid_size_y(); break;1083 case 2: *out = __builtin_amdgcn_grid_size_z(); break;1084 default: *out = 0;1085 }1086}1087 1088// CHECK-LABEL: @test_fmed3_f321089// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.fmed3.f32(1090void test_fmed3_f32(global float* out, float a, float b, float c)1091{1092 *out = __builtin_amdgcn_fmed3f(a, b, c);1093}1094 1095// CHECK-LABEL: @test_s_getpc1096// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.s.getpc()1097void test_s_getpc(global ulong* out)1098{1099 *out = __builtin_amdgcn_s_getpc();1100}1101 1102// CHECK-LABEL: @test_ds_append_lds(1103// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.ds.append.p3(ptr addrspace(3) %{{.+}}, i1 false)1104#if !defined(__SPIRV__)1105kernel void test_ds_append_lds(global int* out, local int* ptr) {1106#else1107kernel void test_ds_append_lds(__attribute__((address_space(1))) int* out, __attribute__((address_space(3))) int* ptr) {1108#endif1109 *out = __builtin_amdgcn_ds_append(ptr);1110}1111 1112// CHECK-LABEL: @test_ds_consume_lds(1113// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.ds.consume.p3(ptr addrspace(3) %{{.+}}, i1 false)1114 1115#if !defined(__SPIRV__)1116kernel void test_ds_consume_lds(global int* out, local int* ptr) {1117#else1118kernel void test_ds_consume_lds(__attribute__((address_space(1))) int* out, __attribute__((address_space(3))) int* ptr) {1119#endif1120 *out = __builtin_amdgcn_ds_consume(ptr);1121}1122 1123// CHECK-LABEL: @test_gws_init(1124// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.ds.gws.init(i32 %value, i32 %id)1125kernel void test_gws_init(uint value, uint id) {1126 __builtin_amdgcn_ds_gws_init(value, id);1127}1128 1129// CHECK-LABEL: @test_gws_barrier(1130// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.ds.gws.barrier(i32 %value, i32 %id)1131kernel void test_gws_barrier(uint value, uint id) {1132 __builtin_amdgcn_ds_gws_barrier(value, id);1133}1134 1135// CHECK-LABEL: @test_gws_sema_v(1136// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.ds.gws.sema.v(i32 %id)1137kernel void test_gws_sema_v(uint id) {1138 __builtin_amdgcn_ds_gws_sema_v(id);1139}1140 1141// CHECK-LABEL: @test_gws_sema_br(1142// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.ds.gws.sema.br(i32 %value, i32 %id)1143kernel void test_gws_sema_br(uint value, uint id) {1144 __builtin_amdgcn_ds_gws_sema_br(value, id);1145}1146 1147// CHECK-LABEL: @test_gws_sema_p(1148// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.ds.gws.sema.p(i32 %id)1149kernel void test_gws_sema_p(uint id) {1150 __builtin_amdgcn_ds_gws_sema_p(id);1151}1152 1153// CHECK-LABEL: @test_mbcnt_lo(1154// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.mbcnt.lo(i32 %src0, i32 %src1)1155kernel void test_mbcnt_lo(global uint* out, uint src0, uint src1) {1156 *out = __builtin_amdgcn_mbcnt_lo(src0, src1);1157}1158 1159// CHECK-LABEL: @test_mbcnt_hi(1160// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.mbcnt.hi(i32 %src0, i32 %src1)1161kernel void test_mbcnt_hi(global uint* out, uint src0, uint src1) {1162 *out = __builtin_amdgcn_mbcnt_hi(src0, src1);1163}1164 1165// CHECK-LABEL: @test_alignbit(1166// CHECK: tail call{{.*}} i32 @llvm.fshr.i32(i32 %src0, i32 %src1, i32 %src2)1167kernel void test_alignbit(global uint* out, uint src0, uint src1, uint src2) {1168 *out = __builtin_amdgcn_alignbit(src0, src1, src2);1169}1170 1171// CHECK-LABEL: @test_alignbyte(1172// CHECK: tail call{{.*}} i32 @llvm.amdgcn.alignbyte(i32 %src0, i32 %src1, i32 %src2)1173kernel void test_alignbyte(global uint* out, uint src0, uint src1, uint src2) {1174 *out = __builtin_amdgcn_alignbyte(src0, src1, src2);1175}1176 1177// CHECK-LABEL: @test_ubfe(1178// CHECK: tail call{{.*}} i32 @llvm.amdgcn.ubfe.i32(i32 %src0, i32 %src1, i32 %src2)1179kernel void test_ubfe(global uint* out, uint src0, uint src1, uint src2) {1180 *out = __builtin_amdgcn_ubfe(src0, src1, src2);1181}1182 1183// CHECK-LABEL: @test_sbfe(1184// CHECK: tail call{{.*}} i32 @llvm.amdgcn.sbfe.i32(i32 %src0, i32 %src1, i32 %src2)1185kernel void test_sbfe(global uint* out, uint src0, uint src1, uint src2) {1186 *out = __builtin_amdgcn_sbfe(src0, src1, src2);1187}1188 1189// CHECK-LABEL: @test_cvt_pkrtz(1190// CHECK: tail call{{.*}} <2 x half> @llvm.amdgcn.cvt.pkrtz(float %src0, float %src1)1191kernel void test_cvt_pkrtz(global half2* out, float src0, float src1) {1192 *out = __builtin_amdgcn_cvt_pkrtz(src0, src1);1193}1194 1195// CHECK-LABEL: @test_cvt_pk_i16(1196// CHECK: tail call{{.*}} <2 x i16> @llvm.amdgcn.cvt.pk.i16(i32 %src0, i32 %src1)1197kernel void test_cvt_pk_i16(global short2* out, int src0, int src1) {1198 *out = __builtin_amdgcn_cvt_pk_i16(src0, src1);1199}1200 1201// CHECK-LABEL: @test_cvt_pk_u16(1202// CHECK: tail call{{.*}} <2 x i16> @llvm.amdgcn.cvt.pk.u16(i32 %src0, i32 %src1)1203kernel void test_cvt_pk_u16(global ushort2* out, uint src0, uint src1) {1204 *out = __builtin_amdgcn_cvt_pk_u16(src0, src1);1205}1206 1207// CHECK-LABEL: @test_cvt_pk_u8_f321208// CHECK: tail call{{.*}} i32 @llvm.amdgcn.cvt.pk.u8.f32(float %src0, i32 %src1, i32 %src2)1209kernel void test_cvt_pk_u8_f32(global uint* out, float src0, uint src1, uint src2) {1210 *out = __builtin_amdgcn_cvt_pk_u8_f32(src0, src1, src2);1211}1212 1213// CHECK-LABEL: test_msad_u8(1214// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.msad.u8(i32 %src0, i32 %src1, i32 %src2)1215kernel void test_msad_u8(global uint* out, uint src0, uint src1, uint src2) {1216 *out = __builtin_amdgcn_msad_u8(src0, src1, src2);1217}1218 1219// CHECK-LABEL: @test_mqsad_pk_u16_u8(1220// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.mqsad.pk.u16.u8(i64 %src0, i32 %src1, i64 %src2)1221kernel void test_mqsad_pk_u16_u8(global ulong* out, ulong src0, uint src1, ulong src2) {1222 *out = __builtin_amdgcn_mqsad_pk_u16_u8(src0, src1, src2);1223}1224 1225// CHECK-LABEL: test_mqsad_u32_u8(1226// CHECK: {{.*}}call{{.*}} <4 x i32> @llvm.amdgcn.mqsad.u32.u8(i64 %src0, i32 %src1, <4 x i32> %src2)1227kernel void test_mqsad_u32_u8(global uint4* out, ulong src0, uint src1, uint4 src2) {1228 *out = __builtin_amdgcn_mqsad_u32_u8(src0, src1, src2);1229}1230 1231// CHECK-LABEL: test_s_setreg(1232// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.setreg(i32 8193, i32 %val)1233kernel void test_s_setreg(uint val) {1234 __builtin_amdgcn_s_setreg(8193, val);1235}1236 1237// CHECK-LABEL test_atomic_inc_dec(1238#if !defined(__SPIRV__)1239void test_atomic_inc_dec(local uint *lptr, global uint *gptr, uint val) {1240#else1241void test_atomic_inc_dec(__attribute__((address_space(3))) uint *lptr, __attribute__((address_space(1))) uint *gptr, uint val) {1242#endif1243 uint res;1244 1245 // CHECK: atomicrmw uinc_wrap ptr addrspace(3) %lptr, i32 %val syncscope("workgroup") seq_cst, align 41246 res = __builtin_amdgcn_atomic_inc32(lptr, val, __ATOMIC_SEQ_CST, "workgroup");1247 1248 // CHECK: atomicrmw udec_wrap ptr addrspace(3) %lptr, i32 %val syncscope("workgroup") seq_cst, align 41249 res = __builtin_amdgcn_atomic_dec32(lptr, val, __ATOMIC_SEQ_CST, "workgroup");1250 1251 // CHECK-AMDGCN: atomicrmw uinc_wrap ptr addrspace(1) %gptr, i32 %val syncscope("agent") seq_cst, align 41252 // CHECK-SPIRV: atomicrmw uinc_wrap ptr addrspace(1) %gptr, i32 %val syncscope("device") seq_cst, align 41253 res = __builtin_amdgcn_atomic_inc32(gptr, val, __ATOMIC_SEQ_CST, "agent");1254 1255 // CHECK: atomicrmw udec_wrap ptr addrspace(1) %gptr, i32 %val seq_cst, align 41256 res = __builtin_amdgcn_atomic_dec32(gptr, val, __ATOMIC_SEQ_CST, "");1257 1258 // CHECK: atomicrmw volatile udec_wrap ptr addrspace(1) %gptr, i32 %val seq_cst, align 41259 #if !defined(__SPIRV__)1260 res = __builtin_amdgcn_atomic_dec32((volatile global uint*)gptr, val, __ATOMIC_SEQ_CST, "");1261 #else1262 res = __builtin_amdgcn_atomic_dec32((volatile __attribute__((address_space(1))) uint*)gptr, val, __ATOMIC_SEQ_CST, "");1263 #endif1264}1265 1266// CHECK-LABEL test_wavefrontsize(1267unsigned test_wavefrontsize() {1268 1269 // CHECK-AMDGCN: ret i32 {{[0-9]+}}1270 // CHECK-SPIRV: {{.*}}call{{.*}} i32 @llvm.amdgcn.wavefrontsize()1271 return __builtin_amdgcn_wavefrontsize();1272}1273 1274// CHECK-LABEL test_flt_rounds(1275unsigned test_flt_rounds() {1276 1277 // CHECK: {{.*}}call{{.*}} i32 @llvm.get.rounding()1278 unsigned mode = __builtin_flt_rounds();1279 1280#if !defined(__SPIRV__)1281 // CHECK-AMDGCN: call void @llvm.set.rounding(i32 %0)1282 __builtin_set_flt_rounds(mode);1283#endif1284 1285 return mode;1286}1287 1288// CHECK-LABEL test_get_fpenv(1289unsigned long test_get_fpenv() {1290 // CHECK: {{.*}}call{{.*}} i64 @llvm.get.fpenv.i64()1291 return __builtin_amdgcn_get_fpenv();1292}1293 1294// CHECK-LABEL test_set_fpenv(1295void test_set_fpenv(unsigned long env) {1296 // CHECK: {{.*}}call{{.*}} void @llvm.set.fpenv.i64(i64 %[[ENV:.+]])1297 __builtin_amdgcn_set_fpenv(env);1298}1299 1300// CHECK-DAG: [[$GRID_RANGE]] = !{i32 1, i32 0}1301// CHECK-DAG: [[$WS_RANGE]] = !{i16 1, i16 1025}1302// CHECK-DAG: attributes #[[$NOUNWIND_READONLY]] = { convergent mustprogress nocallback nofree nounwind willreturn memory(none) }1303