Home
last modified time | relevance | path

Searched refs:row_mask (Results 1 – 24 of 24) sorted by relevance

/third_party/mindspore/mindspore/ccsrc/backend/kernel_compiler/gpu/cuda_impl/
Dnms_with_mask_impl.cu40 __global__ void MaskInit(int numSq, bool *row_mask) { in MaskInit() argument
42 row_mask[mat_pos] = true; in MaskInit()
111 bool *row_mask) { in NmsPass() argument
120 …row_mask[mask_index] = IouDecision(output, box_i, box_j, box_i_start_index, box_j_start_index, IOU… in NmsPass()
126 __global__ void ReducePass(const int num, bool *sel_boxes, bool *row_mask) { in ReducePass() argument
134 sel_boxes[j] = sel_boxes[j] && row_mask[i * num + j]; in ReducePass()
175 bool *row_mask, cudaStream_t cuda_stream) { in CalPreprocess() argument
177 MaskInit<<<GET_BLOCKS(total_val), GET_THREADS, 0, cuda_stream>>>(total_val, row_mask); in CalPreprocess()
192 …Nms(const int num, const float IOU_value, T *output, bool *sel_boxes, int box_size, bool *row_mask, in CalNms() argument
197 row_mask); in CalNms()
[all …]
Dnms_with_mask_impl.cuh28 bool *row_mask, cudaStream_t cuda_stream);
31 …ms(const int num, const float IOU_value, T *output, bool *sel_boxes, int box_size_, bool *row_mask,
/third_party/mindspore/mindspore/ccsrc/backend/kernel_compiler/cpu/
Dnms_with_mask_cpu_kernel.cc76 void NMSWithMaskCPUKernel<T>::MaskInit(size_t numSq, bool *row_mask) { in MaskInit() argument
77 auto task = [this, &row_mask](int start, int end) { in MaskInit()
79 row_mask[mat_pos] = true; in MaskInit()
161 bool *row_mask) { in NmsPass() argument
162 auto task = [this, &row_mask, &output, num, box_size, IOU_value](int start, int end) { in NmsPass()
169 row_mask[mask_index] = IouDecision(output, box_i_start_index, box_j_start_index, IOU_value); in NmsPass()
178 void NMSWithMaskCPUKernel<T>::ReducePass(const int num, bool *sel_boxes, const bool *row_mask) { in ReducePass() argument
185 auto task = [this, &sel_boxes, &row_mask, i, num](int start, int end) { in ReducePass()
187 sel_boxes[j] = sel_boxes[j] && row_mask[i * num + j]; in ReducePass()
228 auto row_mask = reinterpret_cast<bool *>(workspace[ROW_MASK]->addr); in Launch() local
[all …]
Dnms_with_mask_cpu_kernel.h50 void MaskInit(size_t numSq, bool *row_mask);
59 void NmsPass(const int num, const float IOU_value, const T *output, int box_size, bool *row_mask);
61 void ReducePass(const int num, bool *sel_boxes, const bool *row_mask);
/third_party/mindspore/mindspore/ccsrc/backend/kernel_compiler/gpu/math/
Dnms_with_mask_gpu_kernel.h55 bool *row_mask = GetDeviceAddress<bool>(workspace, 2); in Launch() local
61 CalPreprocess(num_input_, sel_idx, sel_boxes, input, output, index_buff, box_size_, row_mask, in Launch()
63 …CalNms(num_input_, iou_value_, output, sel_boxes, box_size_, row_mask, reinterpret_cast<cudaStream… in Launch()
/third_party/mesa3d/src/amd/compiler/
Daco_lower_to_hw_instr.cpp200 PhysReg vtmp_reg, ReduceOp op, unsigned dpp_ctrl, unsigned row_mask, in emit_int64_dpp_op() argument
215 bld.vop1_dpp(aco_opcode::v_mov_b32, vtmp_def[0], src0[0], dpp_ctrl, row_mask, bank_mask, in emit_int64_dpp_op()
220 dpp_ctrl, row_mask, bank_mask, bound_ctrl); in emit_int64_dpp_op()
223 Operand(vcc, bld.lm), dpp_ctrl, row_mask, bank_mask, bound_ctrl); in emit_int64_dpp_op()
225 bld.vop2_dpp(aco_opcode::v_and_b32, dst[0], src0[0], src1[0], dpp_ctrl, row_mask, bank_mask, in emit_int64_dpp_op()
227 bld.vop2_dpp(aco_opcode::v_and_b32, dst[1], src0[1], src1[1], dpp_ctrl, row_mask, bank_mask, in emit_int64_dpp_op()
230 bld.vop2_dpp(aco_opcode::v_or_b32, dst[0], src0[0], src1[0], dpp_ctrl, row_mask, bank_mask, in emit_int64_dpp_op()
232 bld.vop2_dpp(aco_opcode::v_or_b32, dst[1], src0[1], src1[1], dpp_ctrl, row_mask, bank_mask, in emit_int64_dpp_op()
235 bld.vop2_dpp(aco_opcode::v_xor_b32, dst[0], src0[0], src1[0], dpp_ctrl, row_mask, bank_mask, in emit_int64_dpp_op()
237 bld.vop2_dpp(aco_opcode::v_xor_b32, dst[1], src0[1], src1[1], dpp_ctrl, row_mask, bank_mask, in emit_int64_dpp_op()
[all …]
Daco_opt_value_numbering.cpp179 aDPP.bank_mask == bDPP.bank_mask && aDPP.row_mask == bDPP.row_mask && in operator ()()
Daco_print_ir.cpp599 if (dpp.row_mask != 0xf) in print_instr_format_specific()
600 fprintf(output, " row_mask:0x%.1x", dpp.row_mask); in print_instr_format_specific()
Daco_optimizer_postRA.cpp416 assert(mov->dpp().row_mask == 0xf && mov->dpp().bank_mask == 0xf); in try_combine_dpp()
Daco_ir.cpp350 dpp->row_mask = 0xf; in convert_to_DPP()
Daco_assembler.cpp674 uint32_t encoding = (0xF & dpp.row_mask) << 28; in emit_instruction()
Daco_optimizer.cpp1430 assert(instr->dpp().row_mask == 0xf && instr->dpp().bank_mask == 0xf); in label_instruction()
2106 new_dpp->row_mask = cmp_dpp.row_mask; in combine_inverse_comparison()
Daco_ir.h1412 uint8_t row_mask : 4; member
/third_party/skia/third_party/externals/swiftshader/third_party/llvm-10.0/llvm/lib/Target/AMDGPU/
DVOP1Instructions.td290 dpp_ctrl:$dpp_ctrl, row_mask:$row_mask,
838 (i32 (int_amdgcn_mov_dpp i32:$src, timm:$dpp_ctrl, timm:$row_mask, timm:$bank_mask,
841 (as_i32imm $row_mask), (as_i32imm $bank_mask),
846 (i32 (int_amdgcn_update_dpp i32:$old, i32:$src, timm:$dpp_ctrl, timm:$row_mask,
849 (as_i32imm $row_mask), (as_i32imm $bank_mask),
DGCNDPPCombine.cpp251 DPPInst.add(*TII->getNamedOperand(MovMI, AMDGPU::OpName::row_mask)); in createDPPInst()
370 auto *RowMaskOpnd = TII->getNamedOperand(MovMI, AMDGPU::OpName::row_mask); in combineDPPMov()
DVOP2Instructions.td293 dpp_ctrl:$dpp_ctrl, row_mask:$row_mask,
348 let AsmDPP = "$vdst, vcc, $src0, $src1 $dpp_ctrl$row_mask$bank_mask$bound_ctrl";
362 let AsmDPP = "$vdst, vcc, $src0, $src1, vcc $dpp_ctrl$row_mask$bank_mask$bound_ctrl";
381 dpp_ctrl:$dpp_ctrl, row_mask:$row_mask,
397 let AsmDPP = "$vdst, $src0, $src1, vcc $dpp_ctrl$row_mask$bank_mask$bound_ctrl";
417 dpp_ctrl:$dpp_ctrl, row_mask:$row_mask,
DSIInstrInfo.td1083 def row_mask : NamedOperandU32<"RowMask", NamedMatchClass<"RowMask">>;
1748 (ins dpp_ctrl:$dpp_ctrl, row_mask:$row_mask,
1754 Src0RC:$src0, dpp_ctrl:$dpp_ctrl, row_mask:$row_mask,
1759 dpp_ctrl:$dpp_ctrl, row_mask:$row_mask,
1768 dpp_ctrl:$dpp_ctrl, row_mask:$row_mask,
1774 row_mask:$row_mask, bank_mask:$bank_mask,
1970 string ret = dst#args#" $dpp_ctrl$row_mask$bank_mask$bound_ctrl";
DVOPInstructions.td579 bits<4> row_mask;
591 let Inst{63-60} = row_mask;
DSIInstructions.td1876 (i64 (int_amdgcn_mov_dpp i64:$src, timm:$dpp_ctrl, timm:$row_mask, timm:$bank_mask,
1879 (as_i32imm $row_mask), (as_i32imm $bank_mask),
1884 (i64 (int_amdgcn_update_dpp i64:$old, i64:$src, timm:$dpp_ctrl, timm:$row_mask,
1887 (as_i32imm $row_mask), (as_i32imm $bank_mask),
/third_party/skia/third_party/externals/swiftshader/third_party/llvm-10.0/llvm/include/llvm/IR/
DIntrinsicsAMDGPU.td1525 // llvm.amdgcn.mov.dpp.i32 <src> <dpp_ctrl> <row_mask> <bank_mask> <bound_ctrl>
1532 // llvm.amdgcn.update.dpp.i32 <old> <src> <dpp_ctrl> <row_mask> <bank_mask> <bound_ctrl>
1535 // v_mov_b32 <dest> <src> <dpp_ctrl> <row_mask> <bank_mask> <bound_ctrl>
/third_party/mesa3d/src/amd/llvm/
Dac_llvm_build.c3519 enum dpp_ctrl dpp_ctrl, unsigned row_mask, unsigned bank_mask, in _ac_build_dpp() argument
3531 LLVMConstInt(ctx->i32, row_mask, 0), LLVMConstInt(ctx->i32, bank_mask, 0), in _ac_build_dpp()
3539 enum dpp_ctrl dpp_ctrl, unsigned row_mask, unsigned bank_mask, in ac_build_dpp() argument
3557 _ac_build_dpp(ctx, old, src, dpp_ctrl, row_mask, bank_mask, bound_ctrl); in ac_build_dpp()
3562 ret = _ac_build_dpp(ctx, old, src, dpp_ctrl, row_mask, bank_mask, bound_ctrl); in ac_build_dpp()
/third_party/flutter/skia/third_party/externals/libpng/
Dpngrutil.c3365 static const png_uint_32 row_mask[2/*PACKSWAP*/][3/*depth*/][6] = in png_combine_row() local
3386 row_mask[png][DEPTH_INDEX(depth)][pass]) in png_combine_row()
/third_party/skia/third_party/externals/libpng/
Dpngrutil.c3365 static const png_uint_32 row_mask[2/*PACKSWAP*/][3/*depth*/][6] = in png_combine_row() local
3386 row_mask[png][DEPTH_INDEX(depth)][pass]) in png_combine_row()
/third_party/libpng/
Dpngrutil.c3365 static const png_uint_32 row_mask[2/*PACKSWAP*/][3/*depth*/][6] = in png_combine_row() local
3386 row_mask[png][DEPTH_INDEX(depth)][pass]) in png_combine_row()