Home
last modified time | relevance | path

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

12

/llvm-project-15.0.7/llvm/test/CodeGen/NVPTX/
H A Dshfl.ll4 declare i32 @llvm.nvvm.shfl.down.i32(i32, i32, i32)
5 declare float @llvm.nvvm.shfl.down.f32(float, i32, i32)
6 declare i32 @llvm.nvvm.shfl.up.i32(i32, i32, i32)
7 declare float @llvm.nvvm.shfl.up.f32(float, i32, i32)
8 declare i32 @llvm.nvvm.shfl.bfly.i32(i32, i32, i32)
9 declare float @llvm.nvvm.shfl.bfly.f32(float, i32, i32)
10 declare i32 @llvm.nvvm.shfl.idx.i32(i32, i32, i32)
11 declare float @llvm.nvvm.shfl.idx.f32(float, i32, i32)
14 ; shfl.down.
19 ; CHECK: shfl.down.b32 [[OUT:%r[0-9]+]], [[IN]], 1, 2;
[all …]
H A Dshfl-p.ll4 declare {i32, i1} @llvm.nvvm.shfl.down.i32p(i32, i32, i32)
5 declare {float, i1} @llvm.nvvm.shfl.down.f32p(float, i32, i32)
6 declare {i32, i1} @llvm.nvvm.shfl.up.i32p(i32, i32, i32)
7 declare {float, i1} @llvm.nvvm.shfl.up.f32p(float, i32, i32)
8 declare {i32, i1} @llvm.nvvm.shfl.bfly.i32p(i32, i32, i32)
9 declare {float, i1} @llvm.nvvm.shfl.bfly.f32p(float, i32, i32)
10 declare {i32, i1} @llvm.nvvm.shfl.idx.i32p(i32, i32, i32)
11 declare {float, i1} @llvm.nvvm.shfl.idx.f32p(float, i32, i32)
20 %val = call {i32, i1} @llvm.nvvm.shfl.down.i32p(i32 %a, i32 %b, i32 %c)
31 %val = call {i32, i1} @llvm.nvvm.shfl.down.i32p(i32 %a, i32 %b, i32 %c)
[all …]
H A Dshfl-sync-p.ll4 declare {i32, i1} @llvm.nvvm.shfl.sync.down.i32p(i32, i32, i32, i32)
5 declare {float, i1} @llvm.nvvm.shfl.sync.down.f32p(i32, float, i32, i32)
6 declare {i32, i1} @llvm.nvvm.shfl.sync.up.i32p(i32, i32, i32, i32)
7 declare {float, i1} @llvm.nvvm.shfl.sync.up.f32p(i32, float, i32, i32)
8 declare {i32, i1} @llvm.nvvm.shfl.sync.bfly.i32p(i32, i32, i32, i32)
9 declare {float, i1} @llvm.nvvm.shfl.sync.bfly.f32p(i32, float, i32, i32)
10 declare {i32, i1} @llvm.nvvm.shfl.sync.idx.i32p(i32, i32, i32, i32)
11 declare {float, i1} @llvm.nvvm.shfl.sync.idx.f32p(i32, float, i32, i32)
32 %val = call {i32, i1} @llvm.nvvm.shfl.sync.down.i32p(i32 1, i32 %a, i32 %b, i32 %c)
91 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 2, 3, 1;
[all …]
H A Dshfl-sync.ll4 declare i32 @llvm.nvvm.shfl.sync.down.i32(i32, i32, i32, i32)
5 declare float @llvm.nvvm.shfl.sync.down.f32(float, i32, i32, i32)
6 declare i32 @llvm.nvvm.shfl.sync.up.i32(i32, i32, i32, i32)
7 declare float @llvm.nvvm.shfl.sync.up.f32(float, i32, i32, i32)
8 declare i32 @llvm.nvvm.shfl.sync.bfly.i32(i32, i32, i32, i32)
9 declare float @llvm.nvvm.shfl.sync.bfly.f32(float, i32, i32, i32)
10 declare i32 @llvm.nvvm.shfl.sync.idx.i32(i32, i32, i32, i32)
11 declare float @llvm.nvvm.shfl.sync.idx.f32(float, i32, i32, i32)
51 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], [[B]], 2, 1;
72 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], 2, [[C]], 1;
[all …]
/llvm-project-15.0.7/mlir/test/Target/LLVMIR/
H A Dnvvmir.mlir54 …// CHECK: call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}…
55 %6 = nvvm.shfl.sync bfly %0, %3, %1, %2 : i32 -> i32
57 %7 = nvvm.shfl.sync bfly %0, %4, %1, %2 : f32 -> f32
58 // CHECK: call i32 @llvm.nvvm.shfl.sync.up.i32(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}})
59 %8 = nvvm.shfl.sync up %0, %3, %1, %2 : i32 -> i32
61 %9 = nvvm.shfl.sync up %0, %4, %1, %2 : f32 -> f32
63 %10 = nvvm.shfl.sync down %0, %3, %1, %2 : i32 -> i32
65 %11 = nvvm.shfl.sync down %0, %4, %1, %2 : f32 -> f32
67 %12 = nvvm.shfl.sync idx %0, %3, %1, %2 : i32 -> i32
69 %13 = nvvm.shfl.sync idx %0, %4, %1, %2 : f32 -> f32
[all …]
/llvm-project-15.0.7/llvm/test/CodeGen/RISCV/
H A Drv32zbp-intrinsic.ll164 declare i32 @llvm.riscv.shfl.i32(i32 %a, i32 %b)
169 ; RV32ZBP-NEXT: shfl a0, a0, a1
171 %tmp = call i32 @llvm.riscv.shfl.i32(i32 %a, i32 %b)
178 ; RV32ZBP-NEXT: shfl a0, a0, a1
190 %tmp = call i32 @llvm.riscv.shfl.i32(i32 %a, i32 1)
199 %tmp = call i32 @llvm.riscv.shfl.i32(i32 %a, i32 2)
208 %tmp = call i32 @llvm.riscv.shfl.i32(i32 %a, i32 3)
217 %tmp = call i32 @llvm.riscv.shfl.i32(i32 %a, i32 4)
226 %tmp = call i32 @llvm.riscv.shfl.i32(i32 %a, i32 6)
235 %tmp = call i32 @llvm.riscv.shfl.i32(i32 %a, i32 7)
[all …]
H A Drv64zbp-intrinsic.ll87 declare i32 @llvm.riscv.shfl.i32(i32 %a, i32 %b)
94 %tmp = call i32 @llvm.riscv.shfl.i32(i32 %a, i32 %b)
115 %tmp = call i32 @llvm.riscv.shfl.i32(i32 %a, i32 1)
124 %tmp = call i32 @llvm.riscv.shfl.i32(i32 %a, i32 2)
133 %tmp = call i32 @llvm.riscv.shfl.i32(i32 %a, i32 3)
142 %tmp = call i32 @llvm.riscv.shfl.i32(i32 %a, i32 4)
151 %tmp = call i32 @llvm.riscv.shfl.i32(i32 %a, i32 6)
160 %tmp = call i32 @llvm.riscv.shfl.i32(i32 %a, i32 7)
650 declare i64 @llvm.riscv.shfl.i64(i64 %a, i64 %b)
655 ; RV64ZBP-NEXT: shfl a0, a0, a1
[all …]
/llvm-project-15.0.7/mlir/test/Dialect/LLVMIR/
H A Dnvvm.mlir50 // CHECK: nvvm.shfl.sync bfly %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} : i32 -> i32
51 %0 = nvvm.shfl.sync bfly %arg0, %arg3, %arg1, %arg2 : i32 -> i32
52 // CHECK: nvvm.shfl.sync bfly %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} : f32 -> f32
53 %1 = nvvm.shfl.sync bfly %arg0, %arg4, %arg1, %arg2 : f32 -> f32
54 // CHECK: nvvm.shfl.sync up %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} : f32 -> f32
55 %2 = nvvm.shfl.sync up %arg0, %arg4, %arg1, %arg2 : f32 -> f32
56 // CHECK: nvvm.shfl.sync down %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} : f32 -> f32
57 %3 = nvvm.shfl.sync down %arg0, %arg4, %arg1, %arg2 : f32 -> f32
58 // CHECK: nvvm.shfl.sync idx %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} : f32 -> f32
59 %4 = nvvm.shfl.sync idx %arg0, %arg4, %arg1, %arg2 : f32 -> f32
[all …]
H A Dinvalid.mlir519 %0 = nvvm.shfl.sync bfly %arg0, %arg3, %arg1, %arg2 {return_value_and_is_valid} : i32 -> i32
526 …%0 = nvvm.shfl.sync bfly %arg0, %arg3, %arg1, %arg2 {return_value_and_is_valid} : i32 -> !llvm.str…
533 …%0 = nvvm.shfl.sync bfly %arg0, %arg3, %arg1, %arg2 {return_value_and_is_valid} : i32 -> !llvm.str…
/llvm-project-15.0.7/mlir/test/Integration/GPU/CUDA/
H A Dshuffle.mlir26 %shfl, %valid = gpu.shuffle xor %val, %offset, %width : f32
27 cf.cond_br %valid, ^bb1(%shfl : f32), ^bb0
/llvm-project-15.0.7/llvm/test/MC/RISCV/
H A Drv32zbp-valid.s20 # CHECK-ASM-AND-OBJ: shfl t0, t1, t2
22 shfl t0, t1, t2 label
H A Drv64zbp-valid.s20 # CHECK-ASM-AND-OBJ: shfl t0, t1, t2
22 shfl t0, t1, t2 label
H A Drv32zbp-invalid.s18 shfl t0, t1 # CHECK: :[[@LINE]]:1: error: too few operands for instruction label
H A Drv32zbp-aliases-valid.s228 shfl x5, x6, 13 label
H A Drv64zbp-aliases-valid.s328 shfl x5, x6, 13 label
/llvm-project-15.0.7/mlir/test/Conversion/GPUToNVVM/
H A Dgpu-to-nvvm.mlir89 // CHECK: nvvm.shfl.sync bfly {{.*}}
105 // CHECK: nvvm.shfl.sync bfly {{.*}}
133 …// CHECK: %[[#SHFL:]] = nvvm.shfl.sync bfly %[[#MASK]], %[[#VALUE]], %[[#OFFSET]], %[[#CLAMP]] {re…
136 %shfl, %pred = gpu.shuffle xor %arg0, %arg1, %arg2 : f32
142 …// CHECK: %[[#SHFL:]] = nvvm.shfl.sync up %[[#MASK]], %[[#VALUE]], %[[#OFFSET]], %[[#NUM_LANES]] {…
146 … // CHECK: nvvm.shfl.sync down {{.*}} {return_value_and_is_valid} : f32 -> !llvm.struct<(f32, i1)>
148 // CHECK: nvvm.shfl.sync idx {{.*}} {return_value_and_is_valid} : f32 -> !llvm.struct<(f32, i1)>
151 func.return %shfl, %shflu, %shfld, %shfli : f32, f32,f32, f32
/llvm-project-15.0.7/mlir/lib/Conversion/GPUToNVVM/
H A DLowerGpuOpsToNVVMOps.cpp114 Value shfl = rewriter.create<NVVM::ShflOp>( in matchAndRewrite() local
118 loc, valueTy, shfl, rewriter.getIndexArrayAttr(0)); in matchAndRewrite()
120 loc, predTy, shfl, rewriter.getIndexArrayAttr(1)); in matchAndRewrite()
/llvm-project-15.0.7/clang/test/CodeGen/RISCV/rvb-intrinsics/
H A Driscv32-zbp.c80 long shfl(long rs1, long rs2) in shfl() function
H A Driscv64-zbp.c144 long shfl(long rs1, long rs2) in shfl() function
/llvm-project-15.0.7/llvm/lib/Target/RISCV/
H A DRISCVScheduleB.td58 def WriteSHFL : SchedWrite; // shfl
138 def ReadSHFL : SchedRead; // shfl
H A DRISCVInstrInfoZb.td415 def SHFL : ALU_rr<0b0000100, 0b001, "shfl">,
781 def : InstAlias<"shfl $rd, $rs1, $shamt",
/llvm-project-15.0.7/mlir/test/Dialect/GPU/
H A Dinvalid.mlir306 …%shfl, %pred = "gpu.shuffle"(%arg0, %arg1, %arg2) { mode = #gpu<shuffle_mode xor> } : (f32, i32, i…
314 %shfl, %pred = gpu.shuffle xor %arg0, %arg1, %arg2 : index
H A Dops.mlir91 %shfl, %pred = gpu.shuffle xor %arg0, %offset, %width : f32
/llvm-project-15.0.7/mlir/docs/Dialects/
H A DVector.md87 dialect adds operations such as `mma.sync`, `shfl` and tests. For `CPU` things
/llvm-project-15.0.7/mlir/include/mlir/Dialect/LLVMIR/
H A DNVVMOps.td141 NVVM_Op<"shfl.sync">,

12