diff options
| author | Jan Vesely <jan.vesely@rutgers.edu> | 2016-07-10 22:38:04 +0000 |
|---|---|---|
| committer | Jan Vesely <jan.vesely@rutgers.edu> | 2016-07-10 22:38:04 +0000 |
| commit | 72c8091697858422eefc75c89842fb246da2cb59 (patch) | |
| tree | 27d990719aea0dc3cb948376fb8e403b9e512bc7 /test/CodeGenOpenCL/builtins-r600.cl | |
| parent | b09c538d507889ab9c2bc4634124d21d04a4a94d (diff) | |
AMDGPU: Export workitem builtins
Reviewers: tstellardAMD
Differential Revision: http://reviews.llvm.org/D20299
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@275030 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'test/CodeGenOpenCL/builtins-r600.cl')
| -rw-r--r-- | test/CodeGenOpenCL/builtins-r600.cl | 37 |
1 files changed, 37 insertions, 0 deletions
diff --git a/test/CodeGenOpenCL/builtins-r600.cl b/test/CodeGenOpenCL/builtins-r600.cl index 9ebcb1fe9d..0af663e6db 100644 --- a/test/CodeGenOpenCL/builtins-r600.cl +++ b/test/CodeGenOpenCL/builtins-r600.cl @@ -32,3 +32,40 @@ void test_legacy_ldexp_f64(global double* out, double a, int b) *out = __builtin_amdgpu_ldexp(a, b); } #endif + +// CHECK-LABEL: @test_implicitarg_ptr +// CHECK: call i8 addrspace(7)* @llvm.r600.implicitarg.ptr() +void test_implicitarg_ptr(__attribute__((address_space(7))) unsigned char ** out) +{ + *out = __builtin_r600_implicitarg_ptr(); +} + +// CHECK-LABEL: @test_get_group_id( +// CHECK: tail call i32 @llvm.r600.read.tgid.x() +// CHECK: tail call i32 @llvm.r600.read.tgid.y() +// CHECK: tail call i32 @llvm.r600.read.tgid.z() +void test_get_group_id(int d, global int *out) +{ + switch (d) { + case 0: *out = __builtin_r600_read_tgid_x(); break; + case 1: *out = __builtin_r600_read_tgid_y(); break; + case 2: *out = __builtin_r600_read_tgid_z(); break; + default: *out = 0; + } +} + +// CHECK-LABEL: @test_get_local_id( +// CHECK: tail call i32 @llvm.r600.read.tidig.x(), !range [[WI_RANGE:![0-9]*]] +// CHECK: tail call i32 @llvm.r600.read.tidig.y(), !range [[WI_RANGE]] +// CHECK: tail call i32 @llvm.r600.read.tidig.z(), !range [[WI_RANGE]] +void test_get_local_id(int d, global int *out) +{ + switch (d) { + case 0: *out = __builtin_r600_read_tidig_x(); break; + case 1: *out = __builtin_r600_read_tidig_y(); break; + case 2: *out = __builtin_r600_read_tidig_z(); break; + default: *out = 0; + } +} + +// CHECK-DAG: [[WI_RANGE]] = !{i32 0, i32 1024} |