maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity
llvm-project/clang/test/CodeGenOpenCL/builtins-amdgcn.cl

431 lines
12 KiB
Common Lisp
Raw Normal View History

AMDGPU: Rename builtins to use amdgcn prefix Keep the ones still used by libclc around for now. Emit the new amdgcn intrinsic name if not targeting r600, in which case the old AMDGPU name is still used. llvm-svn: 258560
2016-01-23 05:30:53 +08:00
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -S -emit-llvm -o - %s | FileCheck %s
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
AMDGPU: Add builtins for recently added intrinsics llvm-svn: 262126
2016-02-27 17:54:43 +08:00
typedef unsigned long ulong;
AMDGPU : Add Clang builtin intrinsics for compare with the full wavefront result. Differential Revision: http://reviews.llvm.org/D22934 llvm-svn: 277824
2016-08-05 23:38:46 +08:00
typedef unsigned int uint;
AMDGPU: Add builtins for recently added intrinsics llvm-svn: 262126
2016-02-27 17:54:43 +08:00
AMDGPU: Rename builtins to use amdgcn prefix Keep the ones still used by libclc around for now. Emit the new amdgcn intrinsic name if not targeting r600, in which case the old AMDGPU name is still used. llvm-svn: 258560
2016-01-23 05:30:53 +08:00
// CHECK-LABEL: @test_div_scale_f64
// CHECK: call { double, i1 } @llvm.amdgcn.div.scale.f64(double %a, double %b, i1 true)
// CHECK-DAG: [[FLAG:%.+]] = extractvalue { double, i1 } %{{.+}}, 1
// CHECK-DAG: [[VAL:%.+]] = extractvalue { double, i1 } %{{.+}}, 0
// CHECK: [[FLAGEXT:%.+]] = zext i1 [[FLAG]] to i32
// CHECK: store i32 [[FLAGEXT]]
void test_div_scale_f64(global double* out, global int* flagout, double a, double b)
{
bool flag;
*out = __builtin_amdgcn_div_scale(a, b, true, &flag);
*flagout = flag;
}
// CHECK-LABEL: @test_div_scale_f32
// CHECK: call { float, i1 } @llvm.amdgcn.div.scale.f32(float %a, float %b, i1 true)
// CHECK-DAG: [[FLAG:%.+]] = extractvalue { float, i1 } %{{.+}}, 1
// CHECK-DAG: [[VAL:%.+]] = extractvalue { float, i1 } %{{.+}}, 0
// CHECK: [[FLAGEXT:%.+]] = zext i1 [[FLAG]] to i32
// CHECK: store i32 [[FLAGEXT]]
void test_div_scale_f32(global float* out, global int* flagout, float a, float b)
{
bool flag;
*out = __builtin_amdgcn_div_scalef(a, b, true, &flag);
*flagout = flag;
}
// CHECK-LABEL: @test_div_fmas_f32
// CHECK: call float @llvm.amdgcn.div.fmas.f32
void test_div_fmas_f32(global float* out, float a, float b, float c, int d)
{
*out = __builtin_amdgcn_div_fmasf(a, b, c, d);
}
// CHECK-LABEL: @test_div_fmas_f64
// CHECK: call double @llvm.amdgcn.div.fmas.f64
void test_div_fmas_f64(global double* out, double a, double b, double c, int d)
{
*out = __builtin_amdgcn_div_fmas(a, b, c, d);
}
// CHECK-LABEL: @test_div_fixup_f32
// CHECK: call float @llvm.amdgcn.div.fixup.f32
void test_div_fixup_f32(global float* out, float a, float b, float c)
{
*out = __builtin_amdgcn_div_fixupf(a, b, c);
}
// CHECK-LABEL: @test_div_fixup_f64
// CHECK: call double @llvm.amdgcn.div.fixup.f64
void test_div_fixup_f64(global double* out, double a, double b, double c)
{
*out = __builtin_amdgcn_div_fixup(a, b, c);
}
// CHECK-LABEL: @test_trig_preop_f32
// CHECK: call float @llvm.amdgcn.trig.preop.f32
void test_trig_preop_f32(global float* out, float a, int b)
{
*out = __builtin_amdgcn_trig_preopf(a, b);
}
// CHECK-LABEL: @test_trig_preop_f64
// CHECK: call double @llvm.amdgcn.trig.preop.f64
void test_trig_preop_f64(global double* out, double a, int b)
{
*out = __builtin_amdgcn_trig_preop(a, b);
}
// CHECK-LABEL: @test_rcp_f32
// CHECK: call float @llvm.amdgcn.rcp.f32
void test_rcp_f32(global float* out, float a)
{
*out = __builtin_amdgcn_rcpf(a);
}
// CHECK-LABEL: @test_rcp_f64
// CHECK: call double @llvm.amdgcn.rcp.f64
void test_rcp_f64(global double* out, double a)
{
*out = __builtin_amdgcn_rcp(a);
}
// CHECK-LABEL: @test_rsq_f32
// CHECK: call float @llvm.amdgcn.rsq.f32
void test_rsq_f32(global float* out, float a)
{
*out = __builtin_amdgcn_rsqf(a);
}
// CHECK-LABEL: @test_rsq_f64
// CHECK: call double @llvm.amdgcn.rsq.f64
void test_rsq_f64(global double* out, double a)
{
*out = __builtin_amdgcn_rsq(a);
}
AMDGPU: Update builtin for intrinsic change llvm-svn: 260781
2016-02-13 09:03:09 +08:00
// CHECK-LABEL: @test_rsq_clamp_f32
// CHECK: call float @llvm.amdgcn.rsq.clamp.f32
void test_rsq_clamp_f32(global float* out, float a)
AMDGPU: Rename builtins to use amdgcn prefix Keep the ones still used by libclc around for now. Emit the new amdgcn intrinsic name if not targeting r600, in which case the old AMDGPU name is still used. llvm-svn: 258560
2016-01-23 05:30:53 +08:00
{
AMDGPU: Update builtin for intrinsic change llvm-svn: 260781
2016-02-13 09:03:09 +08:00
*out = __builtin_amdgcn_rsq_clampf(a);
AMDGPU: Rename builtins to use amdgcn prefix Keep the ones still used by libclc around for now. Emit the new amdgcn intrinsic name if not targeting r600, in which case the old AMDGPU name is still used. llvm-svn: 258560
2016-01-23 05:30:53 +08:00
}
AMDGPU: Update builtin for intrinsic change llvm-svn: 260781
2016-02-13 09:03:09 +08:00
// CHECK-LABEL: @test_rsq_clamp_f64
// CHECK: call double @llvm.amdgcn.rsq.clamp.f64
void test_rsq_clamp_f64(global double* out, double a)
AMDGPU: Rename builtins to use amdgcn prefix Keep the ones still used by libclc around for now. Emit the new amdgcn intrinsic name if not targeting r600, in which case the old AMDGPU name is still used. llvm-svn: 258560
2016-01-23 05:30:53 +08:00
{
AMDGPU: Update builtin for intrinsic change llvm-svn: 260781
2016-02-13 09:03:09 +08:00
*out = __builtin_amdgcn_rsq_clamp(a);
AMDGPU: Rename builtins to use amdgcn prefix Keep the ones still used by libclc around for now. Emit the new amdgcn intrinsic name if not targeting r600, in which case the old AMDGPU name is still used. llvm-svn: 258560
2016-01-23 05:30:53 +08:00
}
AMDGPU: Add sin/cos builtins llvm-svn: 260783
2016-02-13 09:21:09 +08:00
// CHECK-LABEL: @test_sin_f32
// CHECK: call float @llvm.amdgcn.sin.f32
void test_sin_f32(global float* out, float a)
{
*out = __builtin_amdgcn_sinf(a);
}
// CHECK-LABEL: @test_cos_f32
// CHECK: call float @llvm.amdgcn.cos.f32
void test_cos_f32(global float* out, float a)
{
*out = __builtin_amdgcn_cosf(a);
}
// CHECK-LABEL: @test_log_clamp_f32
// CHECK: call float @llvm.amdgcn.log.clamp.f32
void test_log_clamp_f32(global float* out, float a)
{
*out = __builtin_amdgcn_log_clampf(a);
}
AMDGPU: Rename builtins to use amdgcn prefix Keep the ones still used by libclc around for now. Emit the new amdgcn intrinsic name if not targeting r600, in which case the old AMDGPU name is still used. llvm-svn: 258560
2016-01-23 05:30:53 +08:00
// CHECK-LABEL: @test_ldexp_f32
// CHECK: call float @llvm.amdgcn.ldexp.f32
void test_ldexp_f32(global float* out, float a, int b)
{
*out = __builtin_amdgcn_ldexpf(a, b);
}
// CHECK-LABEL: @test_ldexp_f64
// CHECK: call double @llvm.amdgcn.ldexp.f64
void test_ldexp_f64(global double* out, double a, int b)
{
*out = __builtin_amdgcn_ldexp(a, b);
}
AMDGPU: Add frexp_mant + frexp_exp builtins llvm-svn: 264960
2016-03-31 06:57:40 +08:00
// CHECK-LABEL: @test_frexp_mant_f32
// CHECK: call float @llvm.amdgcn.frexp.mant.f32
void test_frexp_mant_f32(global float* out, float a)
{
*out = __builtin_amdgcn_frexp_mantf(a);
}
// CHECK-LABEL: @test_frexp_mant_f64
// CHECK: call double @llvm.amdgcn.frexp.mant.f64
void test_frexp_mant_f64(global double* out, double a)
{
*out = __builtin_amdgcn_frexp_mant(a);
}
// CHECK-LABEL: @test_frexp_exp_f32
[AMDGPU] Change frexp.exp builtin to return i16 for f16 input Differential Revision: https://reviews.llvm.org/D26863 llvm-svn: 287390
2016-11-19 06:31:51 +08:00
// CHECK: call i32 @llvm.amdgcn.frexp.exp.i32.f32
AMDGPU: Add frexp_mant + frexp_exp builtins llvm-svn: 264960
2016-03-31 06:57:40 +08:00
void test_frexp_exp_f32(global int* out, float a)
{
*out = __builtin_amdgcn_frexp_expf(a);
}
// CHECK-LABEL: @test_frexp_exp_f64
[AMDGPU] Change frexp.exp builtin to return i16 for f16 input Differential Revision: https://reviews.llvm.org/D26863 llvm-svn: 287390
2016-11-19 06:31:51 +08:00
// CHECK: call i32 @llvm.amdgcn.frexp.exp.i32.f64
AMDGPU: Add frexp_mant + frexp_exp builtins llvm-svn: 264960
2016-03-31 06:57:40 +08:00
void test_frexp_exp_f64(global int* out, double a)
{
*out = __builtin_amdgcn_frexp_exp(a);
}
AMDGPU: Add fract builtin llvm-svn: 271080
2016-05-28 08:43:27 +08:00
// CHECK-LABEL: @test_fract_f32
// CHECK: call float @llvm.amdgcn.fract.f32
void test_fract_f32(global int* out, float a)
{
*out = __builtin_amdgcn_fractf(a);
}
// CHECK-LABEL: @test_fract_f64
// CHECK: call double @llvm.amdgcn.fract.f64
void test_fract_f64(global int* out, double a)
{
*out = __builtin_amdgcn_fract(a);
}
AMDGPU: Add Clang Builtin for v_lerp_u8 Differential Revision: http://reviews.llvm.org/D22380 llvm-svn: 275577
2016-07-16 00:43:03 +08:00
// CHECK-LABEL: @test_lerp
// CHECK: call i32 @llvm.amdgcn.lerp
void test_lerp(global int* out, int a, int b, int c)
{
*out = __builtin_amdgcn_lerp(a, b, c);
}
AMDGPU : Add Clang builtin intrinsics for compare with the full wavefront result. Differential Revision: http://reviews.llvm.org/D22934 llvm-svn: 277824
2016-08-05 23:38:46 +08:00
// CHECK-LABEL: @test_sicmp_i32
// CHECK: call i64 @llvm.amdgcn.icmp.i32(i32 %a, i32 %b, i32 32)
void test_sicmp_i32(global ulong* out, int a, int b)
{
*out = __builtin_amdgcn_sicmp(a, b, 32);
}
// CHECK-LABEL: @test_uicmp_i32
// CHECK: call i64 @llvm.amdgcn.icmp.i32(i32 %a, i32 %b, i32 32)
void test_uicmp_i32(global ulong* out, uint a, uint b)
{
*out = __builtin_amdgcn_uicmp(a, b, 32);
}
// CHECK-LABEL: @test_sicmp_i64
// CHECK: call i64 @llvm.amdgcn.icmp.i64(i64 %a, i64 %b, i32 38)
void test_sicmp_i64(global ulong* out, long a, long b)
{
*out = __builtin_amdgcn_sicmpl(a, b, 39-1);
}
// CHECK-LABEL: @test_uicmp_i64
// CHECK: call i64 @llvm.amdgcn.icmp.i64(i64 %a, i64 %b, i32 35)
void test_uicmp_i64(global ulong* out, ulong a, ulong b)
{
*out = __builtin_amdgcn_uicmpl(a, b, 30+5);
}
AMDGPU: Add clang builtin for ds_swizzle. Summary: int __builtin_amdgcn_ds_swizzle (int a, int imm); while imm is a constant. Differential Revision: http://reviews.llvm.org/D23682 llvm-svn: 279165
2016-08-19 06:04:54 +08:00
// CHECK-LABEL: @test_ds_swizzle
// CHECK: call i32 @llvm.amdgcn.ds.swizzle(i32 %a, i32 32)
void test_ds_swizzle(global int* out, int a)
{
*out = __builtin_amdgcn_ds_swizzle(a, 32);
}
AMDGPU : Add Clang builtin intrinsics for compare with the full wavefront result. Differential Revision: http://reviews.llvm.org/D22934 llvm-svn: 277824
2016-08-05 23:38:46 +08:00
// CHECK-LABEL: @test_fcmp_f32
// CHECK: call i64 @llvm.amdgcn.fcmp.f32(float %a, float %b, i32 5)
void test_fcmp_f32(global ulong* out, float a, float b)
{
*out = __builtin_amdgcn_fcmpf(a, b, 5);
}
// CHECK-LABEL: @test_fcmp_f64
// CHECK: call i64 @llvm.amdgcn.fcmp.f64(double %a, double %b, i32 6)
void test_fcmp_f64(global ulong* out, double a, double b)
{
*out = __builtin_amdgcn_fcmp(a, b, 3+3);
}
AMDGPU: Rename builtins to use amdgcn prefix Keep the ones still used by libclc around for now. Emit the new amdgcn intrinsic name if not targeting r600, in which case the old AMDGPU name is still used. llvm-svn: 258560
2016-01-23 05:30:53 +08:00
// CHECK-LABEL: @test_class_f32
// CHECK: call i1 @llvm.amdgcn.class.f32
void test_class_f32(global float* out, float a, int b)
{
*out = __builtin_amdgcn_classf(a, b);
}
// CHECK-LABEL: @test_class_f64
// CHECK: call i1 @llvm.amdgcn.class.f64
void test_class_f64(global double* out, double a, int b)
{
*out = __builtin_amdgcn_class(a, b);
}
AMDGPU: export l1 cache invalidation intrinsics Differential Revision: https://reviews.llvm.org/D30360 llvm-svn: 296240
2017-02-25 12:20:22 +08:00
// CHECK-LABEL: @test_buffer_wbinvl1
// CHECK: call void @llvm.amdgcn.buffer.wbinvl1(
void test_buffer_wbinvl1()
{
__builtin_amdgcn_buffer_wbinvl1();
}
// CHECK-LABEL: @test_s_dcache_inv
// CHECK: call void @llvm.amdgcn.s.dcache.inv(
void test_s_dcache_inv()
{
__builtin_amdgcn_s_dcache_inv();
}
AMDGPU: export s_waitcnt builtin Differential Revision: https://reviews.llvm.org/D30359 llvm-svn: 296239
2017-02-25 12:20:20 +08:00
// CHECK-LABEL: @test_s_waitcnt
// CHECK: call void @llvm.amdgcn.s.waitcnt(
void test_s_waitcnt()
{
__builtin_amdgcn_s_waitcnt(0);
}
AMDGPU: Add barrier builtin llvm-svn: 258564
2016-01-23 05:56:30 +08:00
// CHECK-LABEL: @test_s_barrier
// CHECK: call void @llvm.amdgcn.s.barrier(
void test_s_barrier()
{
__builtin_amdgcn_s_barrier();
}
AMDGPU: Rename builtins to use amdgcn prefix Keep the ones still used by libclc around for now. Emit the new amdgcn intrinsic name if not targeting r600, in which case the old AMDGPU name is still used. llvm-svn: 258560
2016-01-23 05:30:53 +08:00
[AMDGPU] Add wave barrier builtin The wave barrier represents the discardable barrier. Its main purpose is to carry convergent attribute, thus preventing illegal CFG optimizations. All lanes in a wave come to convergence point simultaneously with SIMT, thus no special instruction is needed in the ISA. The barrier is discarded during code generation. Differential Revision: https://reviews.llvm.org/D26584 llvm-svn: 287006
2016-11-16 02:58:03 +08:00
// CHECK-LABEL: @test_wave_barrier
// CHECK: call void @llvm.amdgcn.wave.barrier(
void test_wave_barrier()
{
__builtin_amdgcn_wave_barrier();
}
AMDGPU: Add builtins for recently added intrinsics llvm-svn: 262126
2016-02-27 17:54:43 +08:00
// CHECK-LABEL: @test_s_memtime
// CHECK: call i64 @llvm.amdgcn.s.memtime()
void test_s_memtime(global ulong* out)
{
*out = __builtin_amdgcn_s_memtime();
}
// CHECK-LABEL: @test_s_sleep
// CHECK: call void @llvm.amdgcn.s.sleep(i32 1)
// CHECK: call void @llvm.amdgcn.s.sleep(i32 15)
void test_s_sleep()
{
__builtin_amdgcn_s_sleep(1);
__builtin_amdgcn_s_sleep(15);
}
[AMDGPU] add s_incperflevel/s_decperflevel builtins Differential revision: https://reviews.llvm.org/D23668 llvm-svn: 279235
2016-08-19 20:54:31 +08:00
// CHECK-LABEL: @test_s_incperflevel
// CHECK: call void @llvm.amdgcn.s.incperflevel(i32 1)
// CHECK: call void @llvm.amdgcn.s.incperflevel(i32 15)
void test_s_incperflevel()
{
__builtin_amdgcn_s_incperflevel(1);
__builtin_amdgcn_s_incperflevel(15);
}
// CHECK-LABEL: @test_s_decperflevel
// CHECK: call void @llvm.amdgcn.s.decperflevel(i32 1)
// CHECK: call void @llvm.amdgcn.s.decperflevel(i32 15)
void test_s_decperflevel()
{
__builtin_amdgcn_s_decperflevel(1);
__builtin_amdgcn_s_decperflevel(15);
}
AMDGPU: Add amdgcn cube builtins llvm-svn: 258794
2016-01-26 14:37:54 +08:00
// CHECK-LABEL: @test_cubeid(
// CHECK: call float @llvm.amdgcn.cubeid(float %a, float %b, float %c)
void test_cubeid(global float* out, float a, float b, float c) {
*out = __builtin_amdgcn_cubeid(a, b, c);
}
// CHECK-LABEL: @test_cubesc(
// CHECK: call float @llvm.amdgcn.cubesc(float %a, float %b, float %c)
void test_cubesc(global float* out, float a, float b, float c) {
*out = __builtin_amdgcn_cubesc(a, b, c);
}
// CHECK-LABEL: @test_cubetc(
// CHECK: call float @llvm.amdgcn.cubetc(float %a, float %b, float %c)
void test_cubetc(global float* out, float a, float b, float c) {
*out = __builtin_amdgcn_cubetc(a, b, c);
}
// CHECK-LABEL: @test_cubema(
// CHECK: call float @llvm.amdgcn.cubema(float %a, float %b, float %c)
void test_cubema(global float* out, float a, float b, float c) {
*out = __builtin_amdgcn_cubema(a, b, c);
}
AMDGPU: Add builtin to read exec mask llvm-svn: 273965
2016-06-28 08:13:17 +08:00
// CHECK-LABEL: @test_read_exec(
// CHECK: call i64 @llvm.read_register.i64(metadata ![[EXEC:[0-9]+]]) #[[READ_EXEC_ATTRS:[0-9]+]]
void test_read_exec(global ulong* out) {
*out = __builtin_amdgcn_read_exec();
}
// CHECK: declare i64 @llvm.read_register.i64(metadata) #[[NOUNWIND_READONLY:[0-9]+]]
AMDGPU: Export workitem builtins Reviewers: tstellardAMD Differential Revision: http://reviews.llvm.org/D20299 llvm-svn: 275030
2016-07-11 06:38:04 +08:00
// CHECK-LABEL: @test_kernarg_segment_ptr
// CHECK: call i8 addrspace(2)* @llvm.amdgcn.kernarg.segment.ptr()
void test_kernarg_segment_ptr(__attribute__((address_space(2))) unsigned char ** out)
{
*out = __builtin_amdgcn_kernarg_segment_ptr();
}
// CHECK-LABEL: @test_implicitarg_ptr
// CHECK: call i8 addrspace(2)* @llvm.amdgcn.implicitarg.ptr()
void test_implicitarg_ptr(__attribute__((address_space(2))) unsigned char ** out)
{
*out = __builtin_amdgcn_implicitarg_ptr();
}
// CHECK-LABEL: @test_get_group_id(
// CHECK: tail call i32 @llvm.amdgcn.workgroup.id.x()
// CHECK: tail call i32 @llvm.amdgcn.workgroup.id.y()
// CHECK: tail call i32 @llvm.amdgcn.workgroup.id.z()
void test_get_group_id(int d, global int *out)
{
switch (d) {
case 0: *out = __builtin_amdgcn_workgroup_id_x(); break;
case 1: *out = __builtin_amdgcn_workgroup_id_y(); break;
case 2: *out = __builtin_amdgcn_workgroup_id_z(); break;
default: *out = 0;
}
}
AMDGPU: Add builtin for getreg intrinsic llvm-svn: 292636
2017-01-21 03:24:22 +08:00
// CHECK-LABEL: @test_s_getreg(
// CHECK: tail call i32 @llvm.amdgcn.s.getreg(i32 0)
// CHECK: tail call i32 @llvm.amdgcn.s.getreg(i32 1)
// CHECK: tail call i32 @llvm.amdgcn.s.getreg(i32 65535)
void test_s_getreg(volatile global uint *out)
{
*out = __builtin_amdgcn_s_getreg(0);
*out = __builtin_amdgcn_s_getreg(1);
*out = __builtin_amdgcn_s_getreg(65535);
}
AMDGPU: Export workitem builtins Reviewers: tstellardAMD Differential Revision: http://reviews.llvm.org/D20299 llvm-svn: 275030
2016-07-11 06:38:04 +08:00
// CHECK-LABEL: @test_get_local_id(
// CHECK: tail call i32 @llvm.amdgcn.workitem.id.x(), !range [[WI_RANGE:![0-9]*]]
// CHECK: tail call i32 @llvm.amdgcn.workitem.id.y(), !range [[WI_RANGE]]
// CHECK: tail call i32 @llvm.amdgcn.workitem.id.z(), !range [[WI_RANGE]]
void test_get_local_id(int d, global int *out)
{
switch (d) {
case 0: *out = __builtin_amdgcn_workitem_id_x(); break;
case 1: *out = __builtin_amdgcn_workitem_id_y(); break;
case 2: *out = __builtin_amdgcn_workitem_id_z(); break;
default: *out = 0;
}
}
AMDGPU: Add builtin for fmed3 intrinsic llvm-svn: 293600
2017-01-31 11:42:07 +08:00
// CHECK-LABEL: @test_fmed3_f32
// CHECK: call float @llvm.amdgcn.fmed3.f32(
void test_fmed3_f32(global float* out, float a, float b, float c)
{
*out = __builtin_amdgcn_fmed3f(a, b, c);
}
AMDGPU: Export workitem builtins Reviewers: tstellardAMD Differential Revision: http://reviews.llvm.org/D20299 llvm-svn: 275030
2016-07-11 06:38:04 +08:00
// CHECK-DAG: [[WI_RANGE]] = !{i32 0, i32 1024}
AMDGPU: Add builtin to read exec mask llvm-svn: 273965
2016-06-28 08:13:17 +08:00
// CHECK-DAG: attributes #[[NOUNWIND_READONLY:[0-9]+]] = { nounwind readonly }
// CHECK-DAG: attributes #[[READ_EXEC_ATTRS]] = { convergent }
AMDGPU: Export workitem builtins Reviewers: tstellardAMD Differential Revision: http://reviews.llvm.org/D20299 llvm-svn: 275030
2016-07-11 06:38:04 +08:00
// CHECK-DAG: ![[EXEC]] = !{!"exec"}