From 1342f43ec920c5e226affe57ebe23601a5a1a903 Mon Sep 17 00:00:00 2001 From: Tim Shen Date: Mon, 24 Jun 2019 23:29:20 +0000 Subject: [PATCH] Revert "[NVPTX][NFC] Fix documentation for shfl instructions." The original documentation is correct as it matches the C++ builtins. git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@364250 91177308-0d34-0410-b5e6-96231b3b80d8 --- include/llvm/IR/IntrinsicsNVVM.td | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/include/llvm/IR/IntrinsicsNVVM.td b/include/llvm/IR/IntrinsicsNVVM.td index d5e9dfb17c6..0301e374583 100644 --- a/include/llvm/IR/IntrinsicsNVVM.td +++ b/include/llvm/IR/IntrinsicsNVVM.td @@ -3964,7 +3964,7 @@ def int_nvvm_read_ptx_sreg_warpsize : PTXReadSRegIntrinsic_r32<"warpsize">; // SHUFFLE // -// shfl.down.b32 dest, val, lane_or_offset, mask_and_clamp +// shfl.down.b32 dest, val, offset, mask_and_clamp def int_nvvm_shfl_down_i32 : Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.shfl.down.i32">, @@ -3974,7 +3974,7 @@ def int_nvvm_shfl_down_f32 : [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.shfl.down.f32">, GCCBuiltin<"__nvvm_shfl_down_f32">; -// shfl.up.b32 dest, val, lane_or_offset, mask_and_clamp +// shfl.up.b32 dest, val, offset, mask_and_clamp def int_nvvm_shfl_up_i32 : Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.shfl.up.i32">, @@ -3984,7 +3984,7 @@ def int_nvvm_shfl_up_f32 : [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.shfl.up.f32">, GCCBuiltin<"__nvvm_shfl_up_f32">; -// shfl.bfly.b32 dest, val, lane_or_offset, mask_and_clamp +// shfl.bfly.b32 dest, val, offset, mask_and_clamp def int_nvvm_shfl_bfly_i32 : Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.shfl.bfly.i32">, @@ -3994,7 +3994,7 @@ def int_nvvm_shfl_bfly_f32 : [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.shfl.bfly.f32">, GCCBuiltin<"__nvvm_shfl_bfly_f32">; -// shfl.idx.b32 dest, val, lane_or_offset, mask_and_clamp +// shfl.idx.b32 dest, val, lane, mask_and_clamp def int_nvvm_shfl_idx_i32 : Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.shfl.idx.i32">, @@ -4008,7 +4008,7 @@ def int_nvvm_shfl_idx_f32 : // On sm_70 these don't have to be convergent, so we may eventually want to // implement non-convergent variant of this intrinsic. -// shfl.sync.down.b32 dest, val, lane_or_offset, mask_and_clamp, threadmask +// shfl.sync.down.b32 dest, threadmask, val, offset , mask_and_clamp def int_nvvm_shfl_sync_down_i32 : Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.shfl.sync.down.i32">, @@ -4018,7 +4018,7 @@ def int_nvvm_shfl_sync_down_f32 : [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.shfl.sync.down.f32">, GCCBuiltin<"__nvvm_shfl_sync_down_f32">; -// shfl.sync.up.b32 dest, val, lane_or_offset, mask_and_clamp, threadmask +// shfl.sync.up.b32 dest, threadmask, val, offset, mask_and_clamp def int_nvvm_shfl_sync_up_i32 : Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.shfl.sync.up.i32">, @@ -4028,7 +4028,7 @@ def int_nvvm_shfl_sync_up_f32 : [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.shfl.sync.up.f32">, GCCBuiltin<"__nvvm_shfl_sync_up_f32">; -// shfl.sync.bfly.b32 dest, val, lane_or_offset, mask_and_clamp, threadmask +// shfl.sync.bfly.b32 dest, threadmask, val, offset, mask_and_clamp def int_nvvm_shfl_sync_bfly_i32 : Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.shfl.sync.bfly.i32">, @@ -4038,7 +4038,7 @@ def int_nvvm_shfl_sync_bfly_f32 : [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.shfl.sync.bfly.f32">, GCCBuiltin<"__nvvm_shfl_sync_bfly_f32">; -// shfl.sync.idx.b32 dest, val, lane_or_offset, mask_and_clamp, threadmask +// shfl.sync.idx.b32 dest, threadmask, val, lane, mask_and_clamp def int_nvvm_shfl_sync_idx_i32 : Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.shfl.sync.idx.i32">, -- 2.40.0