diff options
| author | Matt Arsenault <Matthew.Arsenault@amd.com> | 2019-07-17 00:01:03 +0000 |
|---|---|---|
| committer | Matt Arsenault <Matthew.Arsenault@amd.com> | 2019-07-17 00:01:03 +0000 |
| commit | e56865d40c1691fdd215179c827be66f0cd5bae0 (patch) | |
| tree | 634662df9c300bbc390bd0ff21713dbef8897ae6 /clang/test/CodeGenOpenCL/builtins-amdgcn-gfx10.cl | |
| parent | 98a48794ecdfb131c73abc7ef0dd45435863ff4b (diff) | |
| download | bcm5719-llvm-e56865d40c1691fdd215179c827be66f0cd5bae0.tar.gz bcm5719-llvm-e56865d40c1691fdd215179c827be66f0cd5bae0.zip | |
AMDGPU: Add some missing builtins
llvm-svn: 366286
Diffstat (limited to 'clang/test/CodeGenOpenCL/builtins-amdgcn-gfx10.cl')
| -rw-r--r-- | clang/test/CodeGenOpenCL/builtins-amdgcn-gfx10.cl | 24 |
1 files changed, 24 insertions, 0 deletions
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx10.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx10.cl new file mode 100644 index 00000000000..3921cb90c3a --- /dev/null +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx10.cl @@ -0,0 +1,24 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1011 -S -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1012 -S -emit-llvm -o - %s | FileCheck %s + +typedef unsigned int uint; + +// CHECK-LABEL: @test_permlane16( +// CHECK: call i32 @llvm.amdgcn.permlane16(i32 %a, i32 %b, i32 %c, i32 %d, i1 true, i1 true) +void test_permlane16(global uint* out, uint a, uint b, uint c, uint d) { + *out = __builtin_amdgcn_permlane16(a, b, c, d, 1, 1); +} + +// CHECK-LABEL: @test_permlanex16( +// CHECK: call i32 @llvm.amdgcn.permlanex16(i32 %a, i32 %b, i32 %c, i32 %d, i1 true, i1 true) +void test_permlanex16(global uint* out, uint a, uint b, uint c, uint d) { + *out = __builtin_amdgcn_permlanex16(a, b, c, d, 1, 1); +} + +// CHECK-LABEL: @test_mov_dpp8( +// CHECK: call i32 @llvm.amdgcn.mov.dpp8.i32(i32 %a, i32 1) +void test_mov_dpp8(global uint* out, uint a) { + *out = __builtin_amdgcn_mov_dpp8(a, 1); +} |
