diff options
| author | Matt Arsenault <Matthew.Arsenault@amd.com> | 2019-01-28 23:59:18 +0000 |
|---|---|---|
| committer | Matt Arsenault <Matthew.Arsenault@amd.com> | 2019-01-28 23:59:18 +0000 |
| commit | ac340cfc24f78b2804b02bd9edd49d2d172c25c4 (patch) | |
| tree | e739b743037a8c749e604962fea4fe14076939ea /test/CodeGenOpenCL | |
| parent | 0fcbf6da336a994cad2e6c0962e779d47a3ed79c (diff) | |
AMDGPU: Add ds append/consume builtins
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@352443 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'test/CodeGenOpenCL')
| -rw-r--r-- | test/CodeGenOpenCL/builtins-amdgcn.cl | 12 |
1 files changed, 12 insertions, 0 deletions
diff --git a/test/CodeGenOpenCL/builtins-amdgcn.cl b/test/CodeGenOpenCL/builtins-amdgcn.cl index dc7f480209..6b7ea52dab 100644 --- a/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -536,6 +536,18 @@ void test_s_getpc(global ulong* out) *out = __builtin_amdgcn_s_getpc(); } +// CHECK-LABEL: @test_ds_append_lds( +// CHECK: call i32 @llvm.amdgcn.ds.append.p3i32(i32 addrspace(3)* %ptr, i1 false) +kernel void test_ds_append_lds(global int* out, local int* ptr) { + *out = __builtin_amdgcn_ds_append(ptr); +} + +// CHECK-LABEL: @test_ds_consume_lds( +// CHECK: call i32 @llvm.amdgcn.ds.consume.p3i32(i32 addrspace(3)* %ptr, i1 false) +kernel void test_ds_consume_lds(global int* out, local int* ptr) { + *out = __builtin_amdgcn_ds_consume(ptr); +} + // CHECK-DAG: [[$WI_RANGE]] = !{i32 0, i32 1024} // CHECK-DAG: attributes #[[$NOUNWIND_READONLY:[0-9]+]] = { nounwind readonly } // CHECK-DAG: attributes #[[$READ_EXEC_ATTRS]] = { convergent } |