Home
last modified time | relevance | path

Searched refs:shfl (Results 1 – 15 of 15) sorted by relevance

/external/swiftshader/third_party/llvm-7.0/llvm/test/CodeGen/NVPTX/
Dshfl.ll3 declare i32 @llvm.nvvm.shfl.down.i32(i32, i32, i32)
4 declare float @llvm.nvvm.shfl.down.f32(float, i32, i32)
5 declare i32 @llvm.nvvm.shfl.up.i32(i32, i32, i32)
6 declare float @llvm.nvvm.shfl.up.f32(float, i32, i32)
7 declare i32 @llvm.nvvm.shfl.bfly.i32(i32, i32, i32)
8 declare float @llvm.nvvm.shfl.bfly.f32(float, i32, i32)
9 declare i32 @llvm.nvvm.shfl.idx.i32(i32, i32, i32)
10 declare float @llvm.nvvm.shfl.idx.f32(float, i32, i32)
13 ; shfl.down.
15 ; CHECK-LABEL: .func{{.*}}shfl.down1
[all …]
Dshfl-sync.ll3 declare i32 @llvm.nvvm.shfl.sync.down.i32(i32, i32, i32, i32)
4 declare float @llvm.nvvm.shfl.sync.down.f32(float, i32, i32, i32)
5 declare i32 @llvm.nvvm.shfl.sync.up.i32(i32, i32, i32, i32)
6 declare float @llvm.nvvm.shfl.sync.up.f32(float, i32, i32, i32)
7 declare i32 @llvm.nvvm.shfl.sync.bfly.i32(i32, i32, i32, i32)
8 declare float @llvm.nvvm.shfl.sync.bfly.f32(float, i32, i32, i32)
9 declare i32 @llvm.nvvm.shfl.sync.idx.i32(i32, i32, i32, i32)
10 declare float @llvm.nvvm.shfl.sync.idx.f32(float, i32, i32, i32)
12 ; CHECK-LABEL: .func{{.*}}shfl.sync.rrr
13 define i32 @shfl.sync.rrr(i32 %mask, i32 %a, i32 %b, i32 %c) {
[all …]
/external/llvm/test/CodeGen/NVPTX/
Dshfl.ll3 declare i32 @llvm.nvvm.shfl.down.i32(i32, i32, i32)
4 declare float @llvm.nvvm.shfl.down.f32(float, i32, i32)
5 declare i32 @llvm.nvvm.shfl.up.i32(i32, i32, i32)
6 declare float @llvm.nvvm.shfl.up.f32(float, i32, i32)
7 declare i32 @llvm.nvvm.shfl.bfly.i32(i32, i32, i32)
8 declare float @llvm.nvvm.shfl.bfly.f32(float, i32, i32)
9 declare i32 @llvm.nvvm.shfl.idx.i32(i32, i32, i32)
10 declare float @llvm.nvvm.shfl.idx.f32(float, i32, i32)
13 ; shfl.down.
15 ; CHECK-LABEL: .func{{.*}}shfl.down1
[all …]
/external/mesa3d/src/gallium/drivers/nouveau/codegen/
Dnv50_ir_lowering_gm107.cpp207 Instruction *shfl; in handleDFDX() local
224 shfl = bld.mkOp3(OP_SHFL, TYPE_F32, bld.getScratch(), insn->getSrc(0), in handleDFDX()
226 shfl->subOp = NV50_IR_SUBOP_SHFL_BFLY; in handleDFDX()
231 insn->setSrc(0, shfl->getDef(0)); in handleDFDX()
/external/swiftshader/third_party/llvm-7.0/llvm/include/llvm/IR/
DIntrinsicsNVVM.td3721 // shfl.down.b32 dest, val, offset, mask_and_clamp
3724 [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.shfl.down.i32">,
3728 [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.shfl.down.f32">,
3731 // shfl.up.b32 dest, val, offset, mask_and_clamp
3734 [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.shfl.up.i32">,
3738 [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.shfl.up.f32">,
3741 // shfl.bfly.b32 dest, val, offset, mask_and_clamp
3744 [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.shfl.bfly.i32">,
3748 [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.shfl.bfly.f32">,
3751 // shfl.idx.b32 dest, val, lane, mask_and_clamp
[all …]
/external/llvm/include/llvm/IR/
DIntrinsicsNVVM.td3695 // shfl.down.b32 dest, val, offset, mask_and_clamp
3698 [IntrNoMem, IntrConvergent], "llvm.nvvm.shfl.down.i32">,
3702 [IntrNoMem, IntrConvergent], "llvm.nvvm.shfl.down.f32">,
3705 // shfl.up.b32 dest, val, offset, mask_and_clamp
3708 [IntrNoMem, IntrConvergent], "llvm.nvvm.shfl.up.i32">,
3712 [IntrNoMem, IntrConvergent], "llvm.nvvm.shfl.up.f32">,
3715 // shfl.bfly.b32 dest, val, offset, mask_and_clamp
3718 [IntrNoMem, IntrConvergent], "llvm.nvvm.shfl.bfly.i32">,
3722 [IntrNoMem, IntrConvergent], "llvm.nvvm.shfl.bfly.f32">,
3725 // shfl.idx.b32 dest, val, lane, mask_and_clamp
[all …]
/external/swiftshader/third_party/llvm-7.0/llvm/lib/Target/NVPTX/
DNVPTXIntrinsics.td106 // shfl.{up,down,bfly,idx}.b32
108 // The last two parameters to shfl can be regs or imms. ptxas is smart
115 !strconcat("shfl.", mode, ".b32 $dst, $src, $offset, $mask;"),
121 !strconcat("shfl.", mode, ".b32 $dst, $src, $offset, $mask;"),
127 !strconcat("shfl.", mode, ".b32 $dst, $src, $offset, $mask;"),
133 !strconcat("shfl.", mode, ".b32 $dst, $src, $offset, $mask;"),
147 // Threadmask and the last two parameters to shfl.sync can be regs or imms.
154 !strconcat("shfl.sync.", mode, ".b32 $dst, $src, $offset, $mask, $threadmask;"),
161 !strconcat("shfl.sync.", mode, ".b32 $dst, $src, $offset, $mask, $threadmask;"),
168 !strconcat("shfl.sync.", mode, ".b32 $dst, $src, $offset, $mask, $threadmask;"),
[all …]
/external/swiftshader/third_party/llvm-7.0/configs/common/include/llvm/IR/
DIntrinsicEnums.inc3946 nvvm_shfl_bfly_f32, // llvm.nvvm.shfl.bfly.f32
3947 nvvm_shfl_bfly_i32, // llvm.nvvm.shfl.bfly.i32
3948 nvvm_shfl_down_f32, // llvm.nvvm.shfl.down.f32
3949 nvvm_shfl_down_i32, // llvm.nvvm.shfl.down.i32
3950 nvvm_shfl_idx_f32, // llvm.nvvm.shfl.idx.f32
3951 nvvm_shfl_idx_i32, // llvm.nvvm.shfl.idx.i32
3952 nvvm_shfl_sync_bfly_f32, // llvm.nvvm.shfl.sync.bfly.f32
3953 nvvm_shfl_sync_bfly_i32, // llvm.nvvm.shfl.sync.bfly.i32
3954 nvvm_shfl_sync_down_f32, // llvm.nvvm.shfl.sync.down.f32
3955 nvvm_shfl_sync_down_i32, // llvm.nvvm.shfl.sync.down.i32
[all …]
DIntrinsicImpl.inc3972 "llvm.nvvm.shfl.bfly.f32",
3973 "llvm.nvvm.shfl.bfly.i32",
3974 "llvm.nvvm.shfl.down.f32",
3975 "llvm.nvvm.shfl.down.i32",
3976 "llvm.nvvm.shfl.idx.f32",
3977 "llvm.nvvm.shfl.idx.i32",
3978 "llvm.nvvm.shfl.sync.bfly.f32",
3979 "llvm.nvvm.shfl.sync.bfly.i32",
3980 "llvm.nvvm.shfl.sync.down.f32",
3981 "llvm.nvvm.shfl.sync.down.i32",
[all …]
/external/llvm/lib/Target/NVPTX/
DNVPTXIntrinsics.td68 // shfl.{up,down,bfly,idx}.b32
70 // The last two parameters to shfl can be regs or imms. ptxas is smart
77 !strconcat("shfl.", mode, ".b32 $dst, $src, $offset, $mask;"),
83 !strconcat("shfl.", mode, ".b32 $dst, $src, $offset, $mask;"),
89 !strconcat("shfl.", mode, ".b32 $dst, $src, $offset, $mask;"),
95 !strconcat("shfl.", mode, ".b32 $dst, $src, $offset, $mask;"),
/external/swiftshader/third_party/llvm-subzero/build/Fuchsia/include/llvm/IR/
DIntrinsics.gen3193 nvvm_shfl_bfly_f32, // llvm.nvvm.shfl.bfly.f32
3194 nvvm_shfl_bfly_i32, // llvm.nvvm.shfl.bfly.i32
3195 nvvm_shfl_down_f32, // llvm.nvvm.shfl.down.f32
3196 nvvm_shfl_down_i32, // llvm.nvvm.shfl.down.i32
3197 nvvm_shfl_idx_f32, // llvm.nvvm.shfl.idx.f32
3198 nvvm_shfl_idx_i32, // llvm.nvvm.shfl.idx.i32
3199 nvvm_shfl_up_f32, // llvm.nvvm.shfl.up.f32
3200 nvvm_shfl_up_i32, // llvm.nvvm.shfl.up.i32
9251 "llvm.nvvm.shfl.bfly.f32",
9252 "llvm.nvvm.shfl.bfly.i32",
[all …]
/external/swiftshader/third_party/llvm-subzero/build/Windows/include/llvm/IR/
DIntrinsics.gen3193 nvvm_shfl_bfly_f32, // llvm.nvvm.shfl.bfly.f32
3194 nvvm_shfl_bfly_i32, // llvm.nvvm.shfl.bfly.i32
3195 nvvm_shfl_down_f32, // llvm.nvvm.shfl.down.f32
3196 nvvm_shfl_down_i32, // llvm.nvvm.shfl.down.i32
3197 nvvm_shfl_idx_f32, // llvm.nvvm.shfl.idx.f32
3198 nvvm_shfl_idx_i32, // llvm.nvvm.shfl.idx.i32
3199 nvvm_shfl_up_f32, // llvm.nvvm.shfl.up.f32
3200 nvvm_shfl_up_i32, // llvm.nvvm.shfl.up.i32
9251 "llvm.nvvm.shfl.bfly.f32",
9252 "llvm.nvvm.shfl.bfly.i32",
[all …]
/external/swiftshader/third_party/llvm-subzero/build/Linux/include/llvm/IR/
DIntrinsics.gen3193 nvvm_shfl_bfly_f32, // llvm.nvvm.shfl.bfly.f32
3194 nvvm_shfl_bfly_i32, // llvm.nvvm.shfl.bfly.i32
3195 nvvm_shfl_down_f32, // llvm.nvvm.shfl.down.f32
3196 nvvm_shfl_down_i32, // llvm.nvvm.shfl.down.i32
3197 nvvm_shfl_idx_f32, // llvm.nvvm.shfl.idx.f32
3198 nvvm_shfl_idx_i32, // llvm.nvvm.shfl.idx.i32
3199 nvvm_shfl_up_f32, // llvm.nvvm.shfl.up.f32
3200 nvvm_shfl_up_i32, // llvm.nvvm.shfl.up.i32
9251 "llvm.nvvm.shfl.bfly.f32",
9252 "llvm.nvvm.shfl.bfly.i32",
[all …]
/external/swiftshader/third_party/llvm-subzero/build/MacOS/include/llvm/IR/
DIntrinsics.gen3187 nvvm_shfl_bfly_f32, // llvm.nvvm.shfl.bfly.f32
3188 nvvm_shfl_bfly_i32, // llvm.nvvm.shfl.bfly.i32
3189 nvvm_shfl_down_f32, // llvm.nvvm.shfl.down.f32
3190 nvvm_shfl_down_i32, // llvm.nvvm.shfl.down.i32
3191 nvvm_shfl_idx_f32, // llvm.nvvm.shfl.idx.f32
3192 nvvm_shfl_idx_i32, // llvm.nvvm.shfl.idx.i32
3193 nvvm_shfl_up_f32, // llvm.nvvm.shfl.up.f32
3194 nvvm_shfl_up_i32, // llvm.nvvm.shfl.up.i32
9211 "llvm.nvvm.shfl.bfly.f32",
9212 "llvm.nvvm.shfl.bfly.i32",
[all …]
/external/swiftshader/third_party/llvm-subzero/build/Android/include/llvm/IR/
DIntrinsics.gen3193 nvvm_shfl_bfly_f32, // llvm.nvvm.shfl.bfly.f32
3194 nvvm_shfl_bfly_i32, // llvm.nvvm.shfl.bfly.i32
3195 nvvm_shfl_down_f32, // llvm.nvvm.shfl.down.f32
3196 nvvm_shfl_down_i32, // llvm.nvvm.shfl.down.i32
3197 nvvm_shfl_idx_f32, // llvm.nvvm.shfl.idx.f32
3198 nvvm_shfl_idx_i32, // llvm.nvvm.shfl.idx.i32
3199 nvvm_shfl_up_f32, // llvm.nvvm.shfl.up.f32
3200 nvvm_shfl_up_i32, // llvm.nvvm.shfl.up.i32
9251 "llvm.nvvm.shfl.bfly.f32",
9252 "llvm.nvvm.shfl.bfly.i32",
[all …]