diff options
| author | Artem Belevich <tra@google.com> | 2017-09-20 21:23:07 +0000 |
|---|---|---|
| committer | Artem Belevich <tra@google.com> | 2017-09-20 21:23:07 +0000 |
| commit | 4654dc89beb0884dd707f402a06c1884f88196ad (patch) | |
| tree | 004a34e5afeaaae1565dccb57f5907aa777b889a /clang/test/CodeGen/builtins-nvptx.c | |
| parent | 562bf99ee61dbc28901253590780981c45e38e72 (diff) | |
| download | bcm5719-llvm-4654dc89beb0884dd707f402a06c1884f88196ad.tar.gz bcm5719-llvm-4654dc89beb0884dd707f402a06c1884f88196ad.zip | |
[NVPTX] Implemented shfl.sync instruction and supporting intrinsics/builtins.
Differential Revision: https://reviews.llvm.org/D38090
llvm-svn: 313820
Diffstat (limited to 'clang/test/CodeGen/builtins-nvptx.c')
| -rw-r--r-- | clang/test/CodeGen/builtins-nvptx.c | 21 |
1 files changed, 21 insertions, 0 deletions
diff --git a/clang/test/CodeGen/builtins-nvptx.c b/clang/test/CodeGen/builtins-nvptx.c index b0d646a51fe..c97b549cbe0 100644 --- a/clang/test/CodeGen/builtins-nvptx.c +++ b/clang/test/CodeGen/builtins-nvptx.c @@ -636,3 +636,24 @@ __device__ void nvvm_ldg(const void *p) { typedef double double2 __attribute__((ext_vector_type(2))); __nvvm_ldg_d2((const double2 *)p); } + +// CHECK-LABEL: nvvm_shfl +__device__ void nvvm_shfl(int i, float f, int a, int b) { + // CHECK: call i32 @llvm.nvvm.shfl.down.i32(i32 + __nvvm_shfl_down_i32(i, a, b); + // CHECK: call float @llvm.nvvm.shfl.down.f32(float + __nvvm_shfl_down_f32(f, a, b); + // CHECK: call i32 @llvm.nvvm.shfl.up.i32(i32 + __nvvm_shfl_up_i32(i, a, b); + // CHECK: call float @llvm.nvvm.shfl.up.f32(float + __nvvm_shfl_up_f32(f, a, b); + // CHECK: call i32 @llvm.nvvm.shfl.bfly.i32(i32 + __nvvm_shfl_bfly_i32(i, a, b); + // CHECK: call float @llvm.nvvm.shfl.bfly.f32(float + __nvvm_shfl_bfly_f32(f, a, b); + // CHECK: call i32 @llvm.nvvm.shfl.idx.i32(i32 + __nvvm_shfl_idx_i32(i, a, b); + // CHECK: call float @llvm.nvvm.shfl.idx.f32(float + __nvvm_shfl_idx_f32(f, a, b); + // CHECK: ret void +} |
