diff options
| author | Matt Arsenault <Matthew.Arsenault@amd.com> | 2017-10-09 20:06:37 +0000 |
|---|---|---|
| committer | Matt Arsenault <Matthew.Arsenault@amd.com> | 2017-10-09 20:06:37 +0000 |
| commit | f12e3b848afe8aec6b7bbabb5e26875e8b5636fc (patch) | |
| tree | 491b9e77ff73ee4000dc605b9d2e677f0486f9ab /clang/test/CodeGenOpenCL/builtins-amdgcn.cl | |
| parent | c1d5955684dba9a31cb0ff3b4b61fe2a84e392e8 (diff) | |
| download | bcm5719-llvm-f12e3b848afe8aec6b7bbabb5e26875e8b5636fc.tar.gz bcm5719-llvm-f12e3b848afe8aec6b7bbabb5e26875e8b5636fc.zip | |
AMDGPU: Add read_exec_lo/hi builtins
llvm-svn: 315238
Diffstat (limited to 'clang/test/CodeGenOpenCL/builtins-amdgcn.cl')
| -rw-r--r-- | clang/test/CodeGenOpenCL/builtins-amdgcn.cl | 14 |
1 files changed, 14 insertions, 0 deletions
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl index 65b0666ad1f..9f036547bf4 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -421,6 +421,18 @@ void test_read_exec(global ulong* out) { // CHECK: declare i64 @llvm.read_register.i64(metadata) #[[NOUNWIND_READONLY:[0-9]+]] +// CHECK-LABEL: @test_read_exec_lo( +// CHECK: call i32 @llvm.read_register.i32(metadata ![[EXEC_LO:[0-9]+]]) #[[READ_EXEC_ATTRS]] +void test_read_exec_lo(global uint* out) { + *out = __builtin_amdgcn_read_exec_lo(); +} + +// CHECK-LABEL: @test_read_exec_hi( +// CHECK: call i32 @llvm.read_register.i32(metadata ![[EXEC_HI:[0-9]+]]) #[[READ_EXEC_ATTRS]] +void test_read_exec_hi(global uint* out) { + *out = __builtin_amdgcn_read_exec_hi(); +} + // CHECK-LABEL: @test_dispatch_ptr // CHECK: call i8 addrspace(2)* @llvm.amdgcn.dispatch.ptr() void test_dispatch_ptr(__attribute__((address_space(2))) unsigned char ** out) @@ -499,3 +511,5 @@ void test_s_getpc(global ulong* out) // CHECK-DAG: attributes #[[NOUNWIND_READONLY:[0-9]+]] = { nounwind readonly } // CHECK-DAG: attributes #[[READ_EXEC_ATTRS]] = { convergent } // CHECK-DAG: ![[EXEC]] = !{!"exec"} +// CHECK-DAG: ![[EXEC_LO]] = !{!"exec_lo"} +// CHECK-DAG: ![[EXEC_HI]] = !{!"exec_hi"} |
