| /external/mesa3d/src/freedreno/ir3/ |
| D | ir3_parser.y | 79 static struct ir3_instruction * new_instr(opc_t opc) in new_instr() function
592 cat0_instr: T_OP_NOP { new_instr(OPC_NOP); }
593 | T_OP_BR { new_instr(OPC_B); } cat0_src ',' cat0_immed
594 | T_OP_JUMP { new_instr(OPC_JUMP); } cat0_immed
595 | T_OP_CALL { new_instr(OPC_CALL); } cat0_immed
596 | T_OP_RET { new_instr(OPC_RET); }
597 | T_OP_KILL { new_instr(OPC_KILL); } cat0_src
598 | T_OP_END { new_instr(OPC_END); }
599 | T_OP_EMIT { new_instr(OPC_EMIT); }
600 | T_OP_CUT { new_instr(OPC_CUT); }
[all …]
|
| D | ir3_sched.c | 728 struct ir3_instruction *new_instr = ir3_instr_clone(orig_instr); in split_instr() local
729 di(new_instr, "split instruction"); in split_instr()
730 sched_node_init(ctx, new_instr); in split_instr()
731 return new_instr; in split_instr()
1040 struct ir3_instruction *new_instr = NULL; in sched_block() local
1049 new_instr = split_addr(ctx, &ctx->addr0, in sched_block()
1052 new_instr = split_addr(ctx, &ctx->addr1, in sched_block()
1055 new_instr = split_pred(ctx); in sched_block()
1065 if (new_instr) { in sched_block()
1066 list_delinit(&new_instr->node); in sched_block()
[all …]
|
| D | ir3.c | 1078 struct ir3_instruction *new_instr = instr_create(instr->block, in ir3_instr_clone() local
1083 regs = new_instr->regs; in ir3_instr_clone()
1084 *new_instr = *instr; in ir3_instr_clone()
1085 new_instr->regs = regs; in ir3_instr_clone()
1087 insert_instr(instr->block, new_instr); in ir3_instr_clone()
1090 new_instr->regs_count = 0; in ir3_instr_clone()
1094 ir3_reg_create(new_instr, reg->num, reg->flags); in ir3_instr_clone()
1098 return new_instr; in ir3_instr_clone()
|
| /external/mesa3d/src/freedreno/afuc/ |
| D | parser.y | 56 new_instr(int tok) in new_instr() function
196 alu_msb_instr: T_OP_MSB reg ',' reg { new_instr($1); dst($2); src2($4); }
198 alu_not_instr: T_OP_NOT reg ',' reg { new_instr($1); dst($2); src2($4); }
199 | T_OP_NOT reg ',' immediate { new_instr($1); dst($2); immed($4); }
201 alu_mov_instr: T_OP_MOV reg ',' reg { new_instr($1); dst($2); src1($4); }
203 new_instr($1); dst($2); immed($4); shift($6);
205 | T_OP_MOV reg ',' immediate { new_instr($1); dst($2); immed($4); }
207 new_instr($1); dst($2); label($4); shift($6);
209 | T_OP_MOV reg ',' T_LABEL_REF { new_instr($1); dst($2); label($4); }
211 alu_2src_op: T_OP_ADD { new_instr($1); }
[all …]
|
| /external/mesa3d/src/compiler/nir/ |
| D | nir_lower_atomics_to_ssbo.c | 91 nir_intrinsic_instr *new_instr = in lower_instr() local
101 new_instr->src[0] = nir_src_for_ssa(buffer); in lower_instr()
102 nir_src_copy(&new_instr->src[1], &instr->src[0], new_instr); in lower_instr()
103 new_instr->src[2] = nir_src_for_ssa(temp); in lower_instr()
110 new_instr->src[0] = nir_src_for_ssa(buffer); in lower_instr()
111 nir_src_copy(&new_instr->src[1], &instr->src[0], new_instr); in lower_instr()
112 new_instr->src[2] = nir_src_for_ssa(temp); in lower_instr()
116 new_instr->src[0] = nir_src_for_ssa(buffer); in lower_instr()
117 nir_src_copy(&new_instr->src[1], &instr->src[0], new_instr); in lower_instr()
121 new_instr->src[0] = nir_src_for_ssa(buffer); in lower_instr()
[all …]
|
| D | nir_opt_vectorize.c | 347 nir_instr *new_instr = instr_try_combine(nir, *stack_instr, instr, in vec_instr_stack_push() local
349 if (new_instr) { in vec_instr_stack_push()
350 *stack_instr = new_instr; in vec_instr_stack_push()
|
| D | nir_opt_constant_folding.c | 207 nir_intrinsic_instr *new_instr = in try_fold_intrinsic() local
209 nir_builder_instr_insert(b, &new_instr->instr); in try_fold_intrinsic()
|
| D | nir_search.c | 661 nir_algebraic_update_automaton(nir_instr *new_instr, in nir_algebraic_update_automaton() argument
672 add_uses_to_worklist(new_instr, automaton_worklist); in nir_algebraic_update_automaton()
|
| /external/mesa3d/src/gallium/drivers/zink/ |
| D | nir_lower_dynamic_bo_access.c | 50 nir_intrinsic_instr *new_instr = nir_intrinsic_instr_create(b->shader, instr->intrinsic); in recursive_generate_bo_ssa_def() local
51 new_instr->src[block_idx] = nir_src_for_ssa(nir_imm_int(b, start)); in recursive_generate_bo_ssa_def()
54 nir_src_copy(&new_instr->src[i], &instr->src[i], &new_instr->instr); in recursive_generate_bo_ssa_def()
57 …nir_intrinsic_set_align(new_instr, nir_intrinsic_align_mul(instr), nir_intrinsic_align_offset(inst… in recursive_generate_bo_ssa_def()
59 nir_intrinsic_set_range(new_instr, nir_intrinsic_range(instr)); in recursive_generate_bo_ssa_def()
61 new_instr->num_components = instr->num_components; in recursive_generate_bo_ssa_def()
63 nir_ssa_dest_init(&new_instr->instr, &new_instr->dest, in recursive_generate_bo_ssa_def()
66 nir_builder_instr_insert(b, &new_instr->instr); in recursive_generate_bo_ssa_def()
67 return &new_instr->dest.ssa; in recursive_generate_bo_ssa_def()
|
| /external/mesa3d/src/amd/compiler/ |
| D | aco_optimizer.cpp | 1060 …SMEM_instruction *new_instr = create_instruction<SMEM_instruction>(smem->opcode, Format::SMEM, sme… in label_instruction() local
1061 new_instr->operands[0] = smem->operands[0]; in label_instruction()
1062 new_instr->operands[1] = Operand(offset); in label_instruction()
1064 new_instr->operands[2] = smem->operands[2]; in label_instruction()
1065 new_instr->operands.back() = Operand(base); in label_instruction()
1067 new_instr->definitions[0] = smem->definitions[0]; in label_instruction()
1068 new_instr->sync = smem->sync; in label_instruction()
1069 new_instr->glc = smem->glc; in label_instruction()
1070 new_instr->dlc = smem->dlc; in label_instruction()
1071 new_instr->nv = smem->nv; in label_instruction()
[all …]
|
| /external/tensorflow/tensorflow/compiler/xla/service/ |
| D | change_op_data_type.cc | 46 HloInstruction* new_instr = comp->AddInstruction( in Run() local
49 instr, MakeConvertToHlo(new_instr, from_ty_))); in Run()
|
| D | hlo_computation.cc | 213 HloInstruction* new_instr = in RemoveParameter() local
216 TF_RETURN_IF_ERROR(param_instruction->ReplaceAllUsesWith(new_instr)); in RemoveParameter()
217 param_instructions_[param_no] = new_instr; in RemoveParameter()
266 HloInstruction* new_instr = AddInstructionInternal( in RemoveUnusedParametersImpl() local
269 TF_RETURN_IF_ERROR(param_instruction->ReplaceAllUsesWith(new_instr)); in RemoveUnusedParametersImpl()
270 param_instructions_[param_no] = new_instr; in RemoveUnusedParametersImpl()
1233 const HloInstruction* new_instr = replace(instr.get()); in CloneWithReplacements() local
1234 if (!new_instr) { in CloneWithReplacements()
1237 dfs_stack.push_back(new_instr); in CloneWithReplacements()
1279 std::unique_ptr<HloInstruction> new_instr = in CloneWithReplacements() local
[all …]
|
| D | space_to_batch_converter.cc | 192 HloInstruction* new_instr, HloInstruction* old_instr,
1003 auto new_instr = old_to_new_instrs_[old_producer]; in CanPropagate() local
1004 auto permute_dims = instr_to_dim_permute_map_[new_instr]; in CanPropagate()
1016 new_instr->shape().dimensions(j)) { in CanPropagate()
1044 auto new_instr = old_to_new_instrs_[consumer->mutable_operand(i)]; in CanPropagate() local
1045 auto permute_dims = instr_to_dim_permute_map_[new_instr]; in CanPropagate()
1057 new_instr->shape().dimensions(j)) { in CanPropagate()
1798 HloInstruction* new_instr = in Propagate() local
1802 auto permute_dims = instr_to_dim_permute_map_[new_instr]; in Propagate()
1805 const int64_t batch_size = new_instr->shape().dimensions(batch_dim); in Propagate()
[all …]
|
| D | while_loop_constant_sinking.cc | 31 HloInstruction* new_instr, in ReplaceUsesWhileKeepingLoopInvariance() argument
44 TF_RETURN_IF_ERROR(user->ReplaceOperandWith(i, new_instr)); in ReplaceUsesWhileKeepingLoopInvariance()
|
| D | convert_mover.cc | 140 HloInstruction* new_instr = comp->AddInstruction( in MoveConvertPrecisionOps() local
143 instr, HloInstruction::CreateConvert(instr->shape(), new_instr))); in MoveConvertPrecisionOps()
|
| D | dynamic_dimension_inference.cc | 692 HloInstruction* new_instr = computation->AddInstruction( in HandleGetDimensionSize() local
694 TF_RETURN_IF_ERROR(gds->ReplaceAllUsesWith(new_instr)); in HandleGetDimensionSize()
695 parent_->ReplaceAllDynamicDimensionUsesWith(gds, new_instr); in HandleGetDimensionSize()
|
| D | dynamic_padder.cc | 193 HloInstruction* new_instr = computation->AddInstruction( in ReplaceGetSize() local
195 TF_RETURN_IF_ERROR(instr->ReplaceAllUsesWith(new_instr)); in ReplaceGetSize()
197 new_instr); in ReplaceGetSize()
|
| D | hlo_instructions.cc | 1746 HloInstruction* new_instr = parent()->AddInstruction( in CloneAndAppendInstructionIntoCalledComputation() local
1748 TF_CHECK_OK(ReplaceAllUsesWithDifferentShape(new_instr)); in CloneAndAppendInstructionIntoCalledComputation()
|
| /external/tensorflow/tensorflow/compiler/xla/service/gpu/ |
| D | cudnn_fused_conv_rewriter.cc | 325 TF_ASSIGN_OR_RETURN(HloInstruction * new_instr, in FuseBiasOrSideInput()
327 TF_RETURN_IF_ERROR(comp->ReplaceInstruction(instr, new_instr)); in FuseBiasOrSideInput()
535 TF_ASSIGN_OR_RETURN(HloInstruction * new_instr, in FuseConvertToF16()
537 TF_RETURN_IF_ERROR(comp->ReplaceInstruction(instr, new_instr)); in FuseConvertToF16()
616 TF_ASSIGN_OR_RETURN(HloInstruction * new_instr, in FuseConvertToS8()
618 TF_RETURN_IF_ERROR(comp->ReplaceInstruction(instr, new_instr)); in FuseConvertToS8()
|
| D | cudnn_pad_for_convolutions.cc | 91 auto add = [&](std::unique_ptr<HloInstruction> new_instr) { in PadConv() argument
92 return conv->parent()->AddInstruction(std::move(new_instr)); in PadConv()
|
| D | horizontal_loop_fusion.cc | 384 HloInstruction* new_instr = comp->AddInstruction( in CreateFusedComputation() local
386 clone_map.insert({old_instr, new_instr}); in CreateFusedComputation()
|
| /external/vixl/src/aarch64/ |
| D | instructions-aarch64.h | 232 void SetInstructionBits(Instr new_instr) { in SetInstructionBits() argument
233 *(reinterpret_cast<Instr*>(this)) = new_instr; in SetInstructionBits()
|