Lines Matching defs:instr
74 unsigned predict_cost(aco_ptr<Instruction>& instr);
75 void add(aco_ptr<Instruction>& instr);
80 unsigned get_dependency_cost(aco_ptr<Instruction>& instr);
82 void use_resources(aco_ptr<Instruction>& instr);
83 int32_t cycles_until_res_available(aco_ptr<Instruction>& instr);
108 get_perf_info(Program* program, aco_ptr<Instruction>& instr)
110 instr_class cls = instr_info.classes[(int)instr->opcode];
138 return instr->ds().gds ? perf_info{0, WAIT_USE(export_gds, 1)}
167 return instr->ds().gds ? perf_info{4, WAIT_USE(export_gds, 4)}
183 BlockCycleEstimator::use_resources(aco_ptr<Instruction>& instr)
185 perf_info perf = get_perf_info(program, instr);
199 BlockCycleEstimator::cycles_until_res_available(aco_ptr<Instruction>& instr)
201 perf_info perf = get_perf_info(program, instr);
213 get_wait_counter_info(aco_ptr<Instruction>& instr)
218 if (instr->isEXP())
221 if (instr->isFlatLike()) {
222 unsigned lgkm = instr->isFlat() ? 20 : 0;
223 if (!instr->definitions.empty())
229 if (instr->isSMEM()) {
230 if (instr->definitions.empty())
232 if (instr->operands.empty()) /* s_memtime and s_memrealtime */
235 bool likely_desc_load = instr->operands[0].size() == 2;
236 bool soe = instr->operands.size() >= (!instr->definitions.empty() ? 3 : 4);
238 instr->operands[1].isConstant() && (!soe || instr->operands.back().isConstant());
246 if (instr->format == Format::DS)
249 if (instr->isVMEM() && !instr->definitions.empty())
252 if (instr->isVMEM() && instr->definitions.empty())
259 get_wait_imm(Program* program, aco_ptr<Instruction>& instr)
261 if (instr->opcode == aco_opcode::s_endpgm) {
263 } else if (instr->opcode == aco_opcode::s_waitcnt) {
264 return wait_imm(GFX10_3, instr->sopp().imm);
265 } else if (instr->opcode == aco_opcode::s_waitcnt_vscnt) {
266 return wait_imm(0, 0, 0, instr->sopk().imm);
273 wait_counter_info wait_info = get_wait_counter_info(instr);
284 BlockCycleEstimator::get_dependency_cost(aco_ptr<Instruction>& instr)
288 wait_imm imm = get_wait_imm(program, instr);
306 if (instr->opcode == aco_opcode::s_endpgm) {
310 for (Operand& op : instr->operands) {
325 BlockCycleEstimator::predict_cost(aco_ptr<Instruction>& instr)
327 int32_t dep = get_dependency_cost(instr);
328 return dep + std::max(cycles_until_res_available(instr) - dep, 0);
353 BlockCycleEstimator::add(aco_ptr<Instruction>& instr)
355 perf_info perf = get_perf_info(program, instr);
357 cur_cycle += get_dependency_cost(instr);
361 is_vector(instr->opcode) && program->workgroup_size > 32;
363 cur_cycle += cycles_until_res_available(instr);
366 use_resources(instr);
372 wait_imm imm = get_wait_imm(program, instr);
382 wait_counter_info wait_info = get_wait_counter_info(instr);
398 for (Definition& def : instr->definitions) {
443 for (aco_ptr<Instruction>& instr : block.instructions) {
444 if (instr->isSOPP() && instr->sopp().block != -1)
447 if (instr->opcode == aco_opcode::p_constaddr)
450 if ((instr->isVMEM() || instr->isScratch() || instr->isGlobal()) &&
451 !instr->operands.empty()) {
454 { return should_form_clause(instr.get(), other); }))
456 vmem_clause.insert(instr.get());
461 if (instr->isSMEM() && !instr->operands.empty()) {
464 { return should_form_clause(instr.get(), other); }))
466 smem_clause.insert(instr.get());
491 for (aco_ptr<Instruction>& instr : block.instructions) {
493 block_est.add(instr);
494 instr->pass_flags = block_est.cur_cycle - before;