Lines Matching defs:program
38 collect_presched_stats(Program* program)
41 for (Block& block : program->blocks)
43 program->statistics[statistic_sgpr_presched] = presched_demand.sgpr;
44 program->statistics[statistic_vgpr_presched] = presched_demand.vgpr;
61 BlockCycleEstimator(Program* program_) : program(program_) {}
63 Program* program;
108 get_perf_info(Program* program, aco_ptr<Instruction>& instr)
115 if (program->gfx_level >= GFX10) {
154 return program->dev.has_fast_fma32 ? perf_info{4, WAIT_USE(valu, 4)}
185 perf_info perf = get_perf_info(program, instr);
201 perf_info perf = get_perf_info(program, instr);
259 get_wait_imm(Program* program, aco_ptr<Instruction>& instr)
268 unsigned max_lgkm_cnt = program->gfx_level >= GFX10 ? 62 : 14;
270 unsigned max_vm_cnt = program->gfx_level >= GFX9 ? 62 : 14;
288 wait_imm imm = get_wait_imm(program, instr);
309 } else if (program->gfx_level >= GFX10) {
318 if (program->gfx_level < GFX10)
355 perf_info perf = get_perf_info(program, instr);
360 bool dual_issue = program->gfx_level >= GFX10 && program->wave_size == 64 &&
361 is_vector(instr->opcode) && program->workgroup_size > 32;
369 cur_cycle += program->gfx_level >= GFX10 ? 1 : perf.latency;
372 wait_imm imm = get_wait_imm(program, instr);
435 collect_preasm_stats(Program* program)
437 for (Block& block : program->blocks) {
441 program->statistics[statistic_instructions] += block.instructions.size();
445 program->statistics[statistic_branches]++;
448 program->statistics[statistic_instructions] += 2;
455 program->statistics[statistic_vmem_clauses]++;
465 program->statistics[statistic_smem_clauses]++;
475 std::vector<BlockCycleEstimator> blocks(program->blocks.size(), program);
477 if (program->stage.has(SWStage::VS) && program->info.vs.has_prolog) {
479 for (Definition def : program->vs_inputs) {
486 for (Block& block : program->blocks) {
519 program->blocks[block.linear_preds[0]].kind & (block_kind_branch | block_kind_invert);
534 double parallelism = program->num_waves;
540 double wave64_per_cycle = waves_per_cycle * (program->wave_size / 64.0);
543 if (program->workgroup_size != UINT_MAX)
545 program->workgroup_size / (double)align(program->workgroup_size, program->wave_size);
548 program->statistics[statistic_latency] = round(latency);
549 program->statistics[statistic_inv_throughput] = round(1.0 / wave64_per_cycle);
552 aco_print_program(program, stderr, print_no_ssa | print_perf_info);
554 fprintf(stderr, "num_waves: %u\n", program->num_waves);
572 collect_postasm_stats(Program* program, const std::vector<uint32_t>& code)
574 program->statistics[aco::statistic_hash] = util_hash_crc32(code.data(), code.size() * 4);