1// REQUIRES: amdgpu-registered-target 2// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -S -emit-llvm -o - %s | FileCheck %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 i32 27// CHECK: store i32 [[FLAGEXT]] 28void test_div_scale_f32(global float* out, global int* 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_fmas_f32 36// CHECK: call float @llvm.amdgcn.div.fmas.f32 37void test_div_fmas_f32(global float* out, float a, float b, float c, int d) 38{ 39 *out = __builtin_amdgcn_div_fmasf(a, b, c, d); 40} 41 42// CHECK-LABEL: @test_div_fmas_f64 43// CHECK: call double @llvm.amdgcn.div.fmas.f64 44void test_div_fmas_f64(global double* out, double a, double b, double c, int d) 45{ 46 *out = __builtin_amdgcn_div_fmas(a, b, c, d); 47} 48 49// CHECK-LABEL: @test_div_fixup_f32 50// CHECK: call float @llvm.amdgcn.div.fixup.f32 51void test_div_fixup_f32(global float* out, float a, float b, float c) 52{ 53 *out = __builtin_amdgcn_div_fixupf(a, b, c); 54} 55 56// CHECK-LABEL: @test_div_fixup_f64 57// CHECK: call double @llvm.amdgcn.div.fixup.f64 58void test_div_fixup_f64(global double* out, double a, double b, double c) 59{ 60 *out = __builtin_amdgcn_div_fixup(a, b, c); 61} 62 63// CHECK-LABEL: @test_trig_preop_f32 64// CHECK: call float @llvm.amdgcn.trig.preop.f32 65void test_trig_preop_f32(global float* out, float a, int b) 66{ 67 *out = __builtin_amdgcn_trig_preopf(a, b); 68} 69 70// CHECK-LABEL: @test_trig_preop_f64 71// CHECK: call double @llvm.amdgcn.trig.preop.f64 72void test_trig_preop_f64(global double* out, double a, int b) 73{ 74 *out = __builtin_amdgcn_trig_preop(a, b); 75} 76 77// CHECK-LABEL: @test_rcp_f32 78// CHECK: call float @llvm.amdgcn.rcp.f32 79void test_rcp_f32(global float* out, float a) 80{ 81 *out = __builtin_amdgcn_rcpf(a); 82} 83 84// CHECK-LABEL: @test_rcp_f64 85// CHECK: call double @llvm.amdgcn.rcp.f64 86void test_rcp_f64(global double* out, double a) 87{ 88 *out = __builtin_amdgcn_rcp(a); 89} 90 91// CHECK-LABEL: @test_rsq_f32 92// CHECK: call float @llvm.amdgcn.rsq.f32 93void test_rsq_f32(global float* out, float a) 94{ 95 *out = __builtin_amdgcn_rsqf(a); 96} 97 98// CHECK-LABEL: @test_rsq_f64 99// CHECK: call double @llvm.amdgcn.rsq.f64 100void test_rsq_f64(global double* out, double a) 101{ 102 *out = __builtin_amdgcn_rsq(a); 103} 104 105// CHECK-LABEL: @test_rsq_clamp_f32 106// CHECK: call float @llvm.amdgcn.rsq.clamp.f32 107void test_rsq_clamp_f32(global float* out, float a) 108{ 109 *out = __builtin_amdgcn_rsq_clampf(a); 110} 111 112// CHECK-LABEL: @test_rsq_clamp_f64 113// CHECK: call double @llvm.amdgcn.rsq.clamp.f64 114void test_rsq_clamp_f64(global double* out, double a) 115{ 116 *out = __builtin_amdgcn_rsq_clamp(a); 117} 118 119// CHECK-LABEL: @test_sin_f32 120// CHECK: call float @llvm.amdgcn.sin.f32 121void test_sin_f32(global float* out, float a) 122{ 123 *out = __builtin_amdgcn_sinf(a); 124} 125 126// CHECK-LABEL: @test_cos_f32 127// CHECK: call float @llvm.amdgcn.cos.f32 128void test_cos_f32(global float* out, float a) 129{ 130 *out = __builtin_amdgcn_cosf(a); 131} 132 133// CHECK-LABEL: @test_log_clamp_f32 134// CHECK: call float @llvm.amdgcn.log.clamp.f32 135void test_log_clamp_f32(global float* out, float a) 136{ 137 *out = __builtin_amdgcn_log_clampf(a); 138} 139 140// CHECK-LABEL: @test_ldexp_f32 141// CHECK: call float @llvm.amdgcn.ldexp.f32 142void test_ldexp_f32(global float* out, float a, int b) 143{ 144 *out = __builtin_amdgcn_ldexpf(a, b); 145} 146 147// CHECK-LABEL: @test_ldexp_f64 148// CHECK: call double @llvm.amdgcn.ldexp.f64 149void test_ldexp_f64(global double* out, double a, int b) 150{ 151 *out = __builtin_amdgcn_ldexp(a, b); 152} 153 154// CHECK-LABEL: @test_frexp_mant_f32 155// CHECK: call float @llvm.amdgcn.frexp.mant.f32 156void test_frexp_mant_f32(global float* out, float a) 157{ 158 *out = __builtin_amdgcn_frexp_mantf(a); 159} 160 161// CHECK-LABEL: @test_frexp_mant_f64 162// CHECK: call double @llvm.amdgcn.frexp.mant.f64 163void test_frexp_mant_f64(global double* out, double a) 164{ 165 *out = __builtin_amdgcn_frexp_mant(a); 166} 167 168// CHECK-LABEL: @test_frexp_exp_f32 169// CHECK: call i32 @llvm.amdgcn.frexp.exp.i32.f32 170void test_frexp_exp_f32(global int* out, float a) 171{ 172 *out = __builtin_amdgcn_frexp_expf(a); 173} 174 175// CHECK-LABEL: @test_frexp_exp_f64 176// CHECK: call i32 @llvm.amdgcn.frexp.exp.i32.f64 177void test_frexp_exp_f64(global int* out, double a) 178{ 179 *out = __builtin_amdgcn_frexp_exp(a); 180} 181 182// CHECK-LABEL: @test_fract_f32 183// CHECK: call float @llvm.amdgcn.fract.f32 184void test_fract_f32(global int* out, float a) 185{ 186 *out = __builtin_amdgcn_fractf(a); 187} 188 189// CHECK-LABEL: @test_fract_f64 190// CHECK: call double @llvm.amdgcn.fract.f64 191void test_fract_f64(global int* out, double a) 192{ 193 *out = __builtin_amdgcn_fract(a); 194} 195 196// CHECK-LABEL: @test_lerp 197// CHECK: call i32 @llvm.amdgcn.lerp 198void test_lerp(global int* out, int a, int b, int c) 199{ 200 *out = __builtin_amdgcn_lerp(a, b, c); 201} 202 203// CHECK-LABEL: @test_sicmp_i32 204// CHECK: call i64 @llvm.amdgcn.icmp.i32(i32 %a, i32 %b, i32 32) 205void test_sicmp_i32(global ulong* out, int a, int b) 206{ 207 *out = __builtin_amdgcn_sicmp(a, b, 32); 208} 209 210// CHECK-LABEL: @test_uicmp_i32 211// CHECK: call i64 @llvm.amdgcn.icmp.i32(i32 %a, i32 %b, i32 32) 212void test_uicmp_i32(global ulong* out, uint a, uint b) 213{ 214 *out = __builtin_amdgcn_uicmp(a, b, 32); 215} 216 217// CHECK-LABEL: @test_sicmp_i64 218// CHECK: call i64 @llvm.amdgcn.icmp.i64(i64 %a, i64 %b, i32 38) 219void test_sicmp_i64(global ulong* out, long a, long b) 220{ 221 *out = __builtin_amdgcn_sicmpl(a, b, 39-1); 222} 223 224// CHECK-LABEL: @test_uicmp_i64 225// CHECK: call i64 @llvm.amdgcn.icmp.i64(i64 %a, i64 %b, i32 35) 226void test_uicmp_i64(global ulong* out, ulong a, ulong b) 227{ 228 *out = __builtin_amdgcn_uicmpl(a, b, 30+5); 229} 230 231// CHECK-LABEL: @test_ds_swizzle 232// CHECK: call i32 @llvm.amdgcn.ds.swizzle(i32 %a, i32 32) 233void test_ds_swizzle(global int* out, int a) 234{ 235 *out = __builtin_amdgcn_ds_swizzle(a, 32); 236} 237 238// CHECK-LABEL: @test_fcmp_f32 239// CHECK: call i64 @llvm.amdgcn.fcmp.f32(float %a, float %b, i32 5) 240void test_fcmp_f32(global ulong* out, float a, float b) 241{ 242 *out = __builtin_amdgcn_fcmpf(a, b, 5); 243} 244 245// CHECK-LABEL: @test_fcmp_f64 246// CHECK: call i64 @llvm.amdgcn.fcmp.f64(double %a, double %b, i32 6) 247void test_fcmp_f64(global ulong* out, double a, double b) 248{ 249 *out = __builtin_amdgcn_fcmp(a, b, 3+3); 250} 251 252// CHECK-LABEL: @test_class_f32 253// CHECK: call i1 @llvm.amdgcn.class.f32 254void test_class_f32(global float* out, float a, int b) 255{ 256 *out = __builtin_amdgcn_classf(a, b); 257} 258 259// CHECK-LABEL: @test_class_f64 260// CHECK: call i1 @llvm.amdgcn.class.f64 261void test_class_f64(global double* out, double a, int b) 262{ 263 *out = __builtin_amdgcn_class(a, b); 264} 265 266// CHECK-LABEL: @test_s_barrier 267// CHECK: call void @llvm.amdgcn.s.barrier( 268void test_s_barrier() 269{ 270 __builtin_amdgcn_s_barrier(); 271} 272 273// CHECK-LABEL: @test_wave_barrier 274// CHECK: call void @llvm.amdgcn.wave.barrier( 275void test_wave_barrier() 276{ 277 __builtin_amdgcn_wave_barrier(); 278} 279 280// CHECK-LABEL: @test_s_memtime 281// CHECK: call i64 @llvm.amdgcn.s.memtime() 282void test_s_memtime(global ulong* out) 283{ 284 *out = __builtin_amdgcn_s_memtime(); 285} 286 287// CHECK-LABEL: @test_s_sleep 288// CHECK: call void @llvm.amdgcn.s.sleep(i32 1) 289// CHECK: call void @llvm.amdgcn.s.sleep(i32 15) 290void test_s_sleep() 291{ 292 __builtin_amdgcn_s_sleep(1); 293 __builtin_amdgcn_s_sleep(15); 294} 295 296// CHECK-LABEL: @test_s_incperflevel 297// CHECK: call void @llvm.amdgcn.s.incperflevel(i32 1) 298// CHECK: call void @llvm.amdgcn.s.incperflevel(i32 15) 299void test_s_incperflevel() 300{ 301 __builtin_amdgcn_s_incperflevel(1); 302 __builtin_amdgcn_s_incperflevel(15); 303} 304 305// CHECK-LABEL: @test_s_decperflevel 306// CHECK: call void @llvm.amdgcn.s.decperflevel(i32 1) 307// CHECK: call void @llvm.amdgcn.s.decperflevel(i32 15) 308void test_s_decperflevel() 309{ 310 __builtin_amdgcn_s_decperflevel(1); 311 __builtin_amdgcn_s_decperflevel(15); 312} 313 314// CHECK-LABEL: @test_cubeid( 315// CHECK: call float @llvm.amdgcn.cubeid(float %a, float %b, float %c) 316void test_cubeid(global float* out, float a, float b, float c) { 317 *out = __builtin_amdgcn_cubeid(a, b, c); 318} 319 320// CHECK-LABEL: @test_cubesc( 321// CHECK: call float @llvm.amdgcn.cubesc(float %a, float %b, float %c) 322void test_cubesc(global float* out, float a, float b, float c) { 323 *out = __builtin_amdgcn_cubesc(a, b, c); 324} 325 326// CHECK-LABEL: @test_cubetc( 327// CHECK: call float @llvm.amdgcn.cubetc(float %a, float %b, float %c) 328void test_cubetc(global float* out, float a, float b, float c) { 329 *out = __builtin_amdgcn_cubetc(a, b, c); 330} 331 332// CHECK-LABEL: @test_cubema( 333// CHECK: call float @llvm.amdgcn.cubema(float %a, float %b, float %c) 334void test_cubema(global float* out, float a, float b, float c) { 335 *out = __builtin_amdgcn_cubema(a, b, c); 336} 337 338// CHECK-LABEL: @test_read_exec( 339// CHECK: call i64 @llvm.read_register.i64(metadata ![[EXEC:[0-9]+]]) #[[READ_EXEC_ATTRS:[0-9]+]] 340void test_read_exec(global ulong* out) { 341 *out = __builtin_amdgcn_read_exec(); 342} 343 344// CHECK: declare i64 @llvm.read_register.i64(metadata) #[[NOUNWIND_READONLY:[0-9]+]] 345 346// CHECK-LABEL: @test_kernarg_segment_ptr 347// CHECK: call i8 addrspace(2)* @llvm.amdgcn.kernarg.segment.ptr() 348void test_kernarg_segment_ptr(__attribute__((address_space(2))) unsigned char ** out) 349{ 350 *out = __builtin_amdgcn_kernarg_segment_ptr(); 351} 352 353// CHECK-LABEL: @test_implicitarg_ptr 354// CHECK: call i8 addrspace(2)* @llvm.amdgcn.implicitarg.ptr() 355void test_implicitarg_ptr(__attribute__((address_space(2))) unsigned char ** out) 356{ 357 *out = __builtin_amdgcn_implicitarg_ptr(); 358} 359 360// CHECK-LABEL: @test_get_group_id( 361// CHECK: tail call i32 @llvm.amdgcn.workgroup.id.x() 362// CHECK: tail call i32 @llvm.amdgcn.workgroup.id.y() 363// CHECK: tail call i32 @llvm.amdgcn.workgroup.id.z() 364void test_get_group_id(int d, global int *out) 365{ 366 switch (d) { 367 case 0: *out = __builtin_amdgcn_workgroup_id_x(); break; 368 case 1: *out = __builtin_amdgcn_workgroup_id_y(); break; 369 case 2: *out = __builtin_amdgcn_workgroup_id_z(); break; 370 default: *out = 0; 371 } 372} 373 374// CHECK-LABEL: @test_get_local_id( 375// CHECK: tail call i32 @llvm.amdgcn.workitem.id.x(), !range [[WI_RANGE:![0-9]*]] 376// CHECK: tail call i32 @llvm.amdgcn.workitem.id.y(), !range [[WI_RANGE]] 377// CHECK: tail call i32 @llvm.amdgcn.workitem.id.z(), !range [[WI_RANGE]] 378void test_get_local_id(int d, global int *out) 379{ 380 switch (d) { 381 case 0: *out = __builtin_amdgcn_workitem_id_x(); break; 382 case 1: *out = __builtin_amdgcn_workitem_id_y(); break; 383 case 2: *out = __builtin_amdgcn_workitem_id_z(); break; 384 default: *out = 0; 385 } 386} 387 388// CHECK-DAG: [[WI_RANGE]] = !{i32 0, i32 1024} 389// CHECK-DAG: attributes #[[NOUNWIND_READONLY:[0-9]+]] = { nounwind readonly } 390// CHECK-DAG: attributes #[[READ_EXEC_ATTRS]] = { convergent } 391// CHECK-DAG: ![[EXEC]] = !{!"exec"} 392