/* * Copyright © 2016-2017 Broadcom * * Permission is hereby granted, free of charge, to any person obtaining a * copy of this software and associated documentation files (the "Software"), * to deal in the Software without restriction, including without limitation * the rights to use, copy, modify, merge, publish, distribute, sublicense, * and/or sell copies of the Software, and to permit persons to whom the * Software is furnished to do so, subject to the following conditions: * * The above copyright notice and this permission notice (including the next * paragraph) shall be included in all copies or substantial portions of the * Software. * * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS * IN THE SOFTWARE. */ #include "broadcom/common/v3d_device_info.h" #include "v3d_compiler.h" #include "compiler/nir/nir_schedule.h" #include "compiler/nir/nir_builder.h" int vir_get_nsrc(struct qinst *inst) { switch (inst->qpu.type) { case V3D_QPU_INSTR_TYPE_BRANCH: return 0; case V3D_QPU_INSTR_TYPE_ALU: if (inst->qpu.alu.add.op != V3D_QPU_A_NOP) return v3d_qpu_add_op_num_src(inst->qpu.alu.add.op); else return v3d_qpu_mul_op_num_src(inst->qpu.alu.mul.op); } return 0; } /** * Returns whether the instruction has any side effects that must be * preserved. */ bool vir_has_side_effects(struct v3d_compile *c, struct qinst *inst) { switch (inst->qpu.type) { case V3D_QPU_INSTR_TYPE_BRANCH: return true; case V3D_QPU_INSTR_TYPE_ALU: switch (inst->qpu.alu.add.op) { case V3D_QPU_A_SETREVF: case V3D_QPU_A_SETMSF: case V3D_QPU_A_VPMSETUP: case V3D_QPU_A_STVPMV: case V3D_QPU_A_STVPMD: case V3D_QPU_A_STVPMP: case V3D_QPU_A_VPMWT: case V3D_QPU_A_TMUWT: return true; default: break; } switch (inst->qpu.alu.mul.op) { case V3D_QPU_M_MULTOP: return true; default: break; } } if (inst->qpu.sig.ldtmu || inst->qpu.sig.ldvary || inst->qpu.sig.ldtlbu || inst->qpu.sig.ldtlb || inst->qpu.sig.wrtmuc || inst->qpu.sig.thrsw) { return true; } /* ldunifa works like ldunif: it reads an element and advances the * pointer, so each read has a side effect (we don't care for ldunif * because we reconstruct the uniform stream buffer after compiling * with the surviving uniforms), so allowing DCE to remove * one would break follow-up loads. We could fix this by emitting a * unifa for each ldunifa, but each unifa requires 3 delay slots * before a ldunifa, so that would be quite expensive. */ if (inst->qpu.sig.ldunifa || inst->qpu.sig.ldunifarf) return true; return false; } bool vir_is_raw_mov(struct qinst *inst) { if (inst->qpu.type != V3D_QPU_INSTR_TYPE_ALU || (inst->qpu.alu.mul.op != V3D_QPU_M_FMOV && inst->qpu.alu.mul.op != V3D_QPU_M_MOV)) { return false; } if (inst->qpu.alu.add.output_pack != V3D_QPU_PACK_NONE || inst->qpu.alu.mul.output_pack != V3D_QPU_PACK_NONE) { return false; } if (inst->qpu.alu.add.a.unpack != V3D_QPU_UNPACK_NONE || inst->qpu.alu.add.b.unpack != V3D_QPU_UNPACK_NONE || inst->qpu.alu.mul.a.unpack != V3D_QPU_UNPACK_NONE || inst->qpu.alu.mul.b.unpack != V3D_QPU_UNPACK_NONE) { return false; } if (inst->qpu.flags.ac != V3D_QPU_COND_NONE || inst->qpu.flags.mc != V3D_QPU_COND_NONE) return false; return true; } bool vir_is_add(struct qinst *inst) { return (inst->qpu.type == V3D_QPU_INSTR_TYPE_ALU && inst->qpu.alu.add.op != V3D_QPU_A_NOP); } bool vir_is_mul(struct qinst *inst) { return (inst->qpu.type == V3D_QPU_INSTR_TYPE_ALU && inst->qpu.alu.mul.op != V3D_QPU_M_NOP); } bool vir_is_tex(const struct v3d_device_info *devinfo, struct qinst *inst) { if (inst->dst.file == QFILE_MAGIC) return v3d_qpu_magic_waddr_is_tmu(devinfo, inst->dst.index); if (inst->qpu.type == V3D_QPU_INSTR_TYPE_ALU && inst->qpu.alu.add.op == V3D_QPU_A_TMUWT) { return true; } return false; } bool vir_writes_r4_implicitly(const struct v3d_device_info *devinfo, struct qinst *inst) { if (!devinfo->has_accumulators) return false; switch (inst->dst.file) { case QFILE_MAGIC: switch (inst->dst.index) { case V3D_QPU_WADDR_RECIP: case V3D_QPU_WADDR_RSQRT: case V3D_QPU_WADDR_EXP: case V3D_QPU_WADDR_LOG: case V3D_QPU_WADDR_SIN: return true; } break; default: break; } return false; } void vir_set_unpack(struct qinst *inst, int src, enum v3d_qpu_input_unpack unpack) { assert(src == 0 || src == 1); if (vir_is_add(inst)) { if (src == 0) inst->qpu.alu.add.a.unpack = unpack; else inst->qpu.alu.add.b.unpack = unpack; } else { assert(vir_is_mul(inst)); if (src == 0) inst->qpu.alu.mul.a.unpack = unpack; else inst->qpu.alu.mul.b.unpack = unpack; } } void vir_set_pack(struct qinst *inst, enum v3d_qpu_output_pack pack) { if (vir_is_add(inst)) { inst->qpu.alu.add.output_pack = pack; } else { assert(vir_is_mul(inst)); inst->qpu.alu.mul.output_pack = pack; } } void vir_set_cond(struct qinst *inst, enum v3d_qpu_cond cond) { if (vir_is_add(inst)) { inst->qpu.flags.ac = cond; } else { assert(vir_is_mul(inst)); inst->qpu.flags.mc = cond; } } enum v3d_qpu_cond vir_get_cond(struct qinst *inst) { assert(inst->qpu.type == V3D_QPU_INSTR_TYPE_ALU); if (vir_is_add(inst)) return inst->qpu.flags.ac; else if (vir_is_mul(inst)) return inst->qpu.flags.mc; else /* NOP */ return V3D_QPU_COND_NONE; } void vir_set_pf(struct v3d_compile *c, struct qinst *inst, enum v3d_qpu_pf pf) { c->flags_temp = -1; if (vir_is_add(inst)) { inst->qpu.flags.apf = pf; } else { assert(vir_is_mul(inst)); inst->qpu.flags.mpf = pf; } } void vir_set_uf(struct v3d_compile *c, struct qinst *inst, enum v3d_qpu_uf uf) { c->flags_temp = -1; if (vir_is_add(inst)) { inst->qpu.flags.auf = uf; } else { assert(vir_is_mul(inst)); inst->qpu.flags.muf = uf; } } #if 0 uint8_t vir_channels_written(struct qinst *inst) { if (vir_is_mul(inst)) { switch (inst->dst.pack) { case QPU_PACK_MUL_NOP: case QPU_PACK_MUL_8888: return 0xf; case QPU_PACK_MUL_8A: return 0x1; case QPU_PACK_MUL_8B: return 0x2; case QPU_PACK_MUL_8C: return 0x4; case QPU_PACK_MUL_8D: return 0x8; } } else { switch (inst->dst.pack) { case QPU_PACK_A_NOP: case QPU_PACK_A_8888: case QPU_PACK_A_8888_SAT: case QPU_PACK_A_32_SAT: return 0xf; case QPU_PACK_A_8A: case QPU_PACK_A_8A_SAT: return 0x1; case QPU_PACK_A_8B: case QPU_PACK_A_8B_SAT: return 0x2; case QPU_PACK_A_8C: case QPU_PACK_A_8C_SAT: return 0x4; case QPU_PACK_A_8D: case QPU_PACK_A_8D_SAT: return 0x8; case QPU_PACK_A_16A: case QPU_PACK_A_16A_SAT: return 0x3; case QPU_PACK_A_16B: case QPU_PACK_A_16B_SAT: return 0xc; } } unreachable("Bad pack field"); } #endif struct qreg vir_get_temp(struct v3d_compile *c) { struct qreg reg; reg.file = QFILE_TEMP; reg.index = c->num_temps++; if (c->num_temps > c->defs_array_size) { uint32_t old_size = c->defs_array_size; c->defs_array_size = MAX2(old_size * 2, 16); c->defs = reralloc(c, c->defs, struct qinst *, c->defs_array_size); memset(&c->defs[old_size], 0, sizeof(c->defs[0]) * (c->defs_array_size - old_size)); c->spillable = reralloc(c, c->spillable, BITSET_WORD, BITSET_WORDS(c->defs_array_size)); for (int i = old_size; i < c->defs_array_size; i++) BITSET_SET(c->spillable, i); } return reg; } struct qinst * vir_add_inst(enum v3d_qpu_add_op op, struct qreg dst, struct qreg src0, struct qreg src1) { struct qinst *inst = calloc(1, sizeof(*inst)); inst->qpu = v3d_qpu_nop(); inst->qpu.alu.add.op = op; inst->dst = dst; inst->src[0] = src0; inst->src[1] = src1; inst->uniform = ~0; inst->ip = -1; return inst; } struct qinst * vir_mul_inst(enum v3d_qpu_mul_op op, struct qreg dst, struct qreg src0, struct qreg src1) { struct qinst *inst = calloc(1, sizeof(*inst)); inst->qpu = v3d_qpu_nop(); inst->qpu.alu.mul.op = op; inst->dst = dst; inst->src[0] = src0; inst->src[1] = src1; inst->uniform = ~0; inst->ip = -1; return inst; } struct qinst * vir_branch_inst(struct v3d_compile *c, enum v3d_qpu_branch_cond cond) { struct qinst *inst = calloc(1, sizeof(*inst)); inst->qpu = v3d_qpu_nop(); inst->qpu.type = V3D_QPU_INSTR_TYPE_BRANCH; inst->qpu.branch.cond = cond; inst->qpu.branch.msfign = V3D_QPU_MSFIGN_NONE; inst->qpu.branch.bdi = V3D_QPU_BRANCH_DEST_REL; inst->qpu.branch.ub = true; inst->qpu.branch.bdu = V3D_QPU_BRANCH_DEST_REL; inst->dst = vir_nop_reg(); inst->uniform = vir_get_uniform_index(c, QUNIFORM_CONSTANT, 0); inst->ip = -1; return inst; } static void vir_emit(struct v3d_compile *c, struct qinst *inst) { inst->ip = -1; switch (c->cursor.mode) { case vir_cursor_add: list_add(&inst->link, c->cursor.link); break; case vir_cursor_addtail: list_addtail(&inst->link, c->cursor.link); break; } c->cursor = vir_after_inst(inst); c->live_intervals_valid = false; } /* Updates inst to write to a new temporary, emits it, and notes the def. */ struct qreg vir_emit_def(struct v3d_compile *c, struct qinst *inst) { assert(inst->dst.file == QFILE_NULL); /* If we're emitting an instruction that's a def, it had better be * writing a register. */ if (inst->qpu.type == V3D_QPU_INSTR_TYPE_ALU) { assert(inst->qpu.alu.add.op == V3D_QPU_A_NOP || v3d_qpu_add_op_has_dst(inst->qpu.alu.add.op)); assert(inst->qpu.alu.mul.op == V3D_QPU_M_NOP || v3d_qpu_mul_op_has_dst(inst->qpu.alu.mul.op)); } inst->dst = vir_get_temp(c); if (inst->dst.file == QFILE_TEMP) c->defs[inst->dst.index] = inst; vir_emit(c, inst); return inst->dst; } struct qinst * vir_emit_nondef(struct v3d_compile *c, struct qinst *inst) { if (inst->dst.file == QFILE_TEMP) c->defs[inst->dst.index] = NULL; vir_emit(c, inst); return inst; } struct qblock * vir_new_block(struct v3d_compile *c) { struct qblock *block = rzalloc(c, struct qblock); list_inithead(&block->instructions); block->predecessors = _mesa_set_create(block, _mesa_hash_pointer, _mesa_key_pointer_equal); block->index = c->next_block_index++; return block; } void vir_set_emit_block(struct v3d_compile *c, struct qblock *block) { c->cur_block = block; c->cursor = vir_after_block(block); list_addtail(&block->link, &c->blocks); } struct qblock * vir_entry_block(struct v3d_compile *c) { return list_first_entry(&c->blocks, struct qblock, link); } struct qblock * vir_exit_block(struct v3d_compile *c) { return list_last_entry(&c->blocks, struct qblock, link); } void vir_link_blocks(struct qblock *predecessor, struct qblock *successor) { _mesa_set_add(successor->predecessors, predecessor); if (predecessor->successors[0]) { assert(!predecessor->successors[1]); predecessor->successors[1] = successor; } else { predecessor->successors[0] = successor; } } const struct v3d_compiler * v3d_compiler_init(const struct v3d_device_info *devinfo, uint32_t max_inline_uniform_buffers) { struct v3d_compiler *compiler = rzalloc(NULL, struct v3d_compiler); if (!compiler) return NULL; compiler->devinfo = devinfo; compiler->max_inline_uniform_buffers = max_inline_uniform_buffers; if (!vir_init_reg_sets(compiler)) { ralloc_free(compiler); return NULL; } return compiler; } void v3d_compiler_free(const struct v3d_compiler *compiler) { ralloc_free((void *)compiler); } struct v3d_compiler_strategy { const char *name; uint32_t max_threads; uint32_t min_threads; bool disable_general_tmu_sched; bool disable_gcm; bool disable_loop_unrolling; bool disable_ubo_load_sorting; bool move_buffer_loads; bool disable_tmu_pipelining; uint32_t max_tmu_spills; }; static struct v3d_compile * vir_compile_init(const struct v3d_compiler *compiler, struct v3d_key *key, nir_shader *s, void (*debug_output)(const char *msg, void *debug_output_data), void *debug_output_data, int program_id, int variant_id, uint32_t compile_strategy_idx, const struct v3d_compiler_strategy *strategy, bool fallback_scheduler) { struct v3d_compile *c = rzalloc(NULL, struct v3d_compile); c->compiler = compiler; c->devinfo = compiler->devinfo; c->key = key; c->program_id = program_id; c->variant_id = variant_id; c->compile_strategy_idx = compile_strategy_idx; c->threads = strategy->max_threads; c->debug_output = debug_output; c->debug_output_data = debug_output_data; c->compilation_result = V3D_COMPILATION_SUCCEEDED; c->min_threads_for_reg_alloc = strategy->min_threads; c->max_tmu_spills = strategy->max_tmu_spills; c->fallback_scheduler = fallback_scheduler; c->disable_general_tmu_sched = strategy->disable_general_tmu_sched; c->disable_tmu_pipelining = strategy->disable_tmu_pipelining; c->disable_constant_ubo_load_sorting = strategy->disable_ubo_load_sorting; c->move_buffer_loads = strategy->move_buffer_loads; c->disable_gcm = strategy->disable_gcm; c->disable_loop_unrolling = V3D_DBG(NO_LOOP_UNROLL) ? true : strategy->disable_loop_unrolling; s = nir_shader_clone(c, s); c->s = s; list_inithead(&c->blocks); vir_set_emit_block(c, vir_new_block(c)); c->output_position_index = -1; c->output_sample_mask_index = -1; c->def_ht = _mesa_hash_table_create(c, _mesa_hash_pointer, _mesa_key_pointer_equal); c->tmu.outstanding_regs = _mesa_pointer_set_create(c); c->flags_temp = -1; return c; } static int type_size_vec4(const struct glsl_type *type, bool bindless) { return glsl_count_attribute_slots(type, false); } static enum nir_lower_tex_packing lower_tex_packing_cb(const nir_tex_instr *tex, const void *data) { struct v3d_compile *c = (struct v3d_compile *) data; int sampler_index = nir_tex_instr_need_sampler(tex) ? tex->sampler_index : tex->backend_flags; assert(sampler_index < c->key->num_samplers_used); return c->key->sampler[sampler_index].return_size == 16 ? nir_lower_tex_packing_16 : nir_lower_tex_packing_none; } static bool v3d_nir_lower_null_pointers_cb(nir_builder *b, nir_intrinsic_instr *intr, void *_state) { uint32_t buffer_src_idx; switch (intr->intrinsic) { case nir_intrinsic_load_ubo: case nir_intrinsic_load_ssbo: buffer_src_idx = 0; break; case nir_intrinsic_store_ssbo: buffer_src_idx = 1; break; default: return false; } /* If index if constant we are good */ nir_src *src = &intr->src[buffer_src_idx]; if (nir_src_is_const(*src)) return false; /* Otherwise, see if it comes from a bcsel including a null pointer */ if (src->ssa->parent_instr->type != nir_instr_type_alu) return false; nir_alu_instr *alu = nir_instr_as_alu(src->ssa->parent_instr); if (alu->op != nir_op_bcsel) return false; /* A null pointer is specified using block index 0xffffffff */ int32_t null_src_idx = -1; for (int i = 1; i < 3; i++) { /* FIXME: since we are running this before optimization maybe * we need to also handle the case where we may have bcsel * chain that we need to recurse? */ if (!nir_src_is_const(alu->src[i].src)) continue; if (nir_src_comp_as_uint(alu->src[i].src, 0) != 0xffffffff) continue; /* One of the bcsel srcs is a null pointer reference */ null_src_idx = i; break; } if (null_src_idx < 0) return false; assert(null_src_idx == 1 || null_src_idx == 2); int32_t copy_src_idx = null_src_idx == 1 ? 2 : 1; /* Rewrite the null pointer reference so we use the same buffer index * as the other bcsel branch. This will allow optimization to remove * the bcsel and we should then end up with a constant buffer index * like we need. */ b->cursor = nir_before_instr(&alu->instr); nir_def *copy = nir_mov(b, alu->src[copy_src_idx].src.ssa); nir_src_rewrite(&alu->src[null_src_idx].src, copy); return true; } static bool v3d_nir_lower_null_pointers(nir_shader *s) { return nir_shader_intrinsics_pass(s, v3d_nir_lower_null_pointers_cb, nir_metadata_control_flow, NULL); } static unsigned lower_bit_size_cb(const nir_instr *instr, void *_data) { if (instr->type != nir_instr_type_alu) return 0; nir_alu_instr *alu = nir_instr_as_alu(instr); switch (alu->op) { case nir_op_mov: case nir_op_vec2: case nir_op_vec3: case nir_op_vec4: case nir_op_vec5: case nir_op_vec8: case nir_op_vec16: case nir_op_b2i8: case nir_op_b2f16: case nir_op_b2i16: case nir_op_b2f32: case nir_op_b2i32: case nir_op_f2f16: case nir_op_f2f16_rtne: case nir_op_f2f16_rtz: case nir_op_f2f32: case nir_op_f2i32: case nir_op_f2u32: case nir_op_i2i8: case nir_op_i2i16: case nir_op_i2f16: case nir_op_i2f32: case nir_op_i2i32: case nir_op_u2u8: case nir_op_u2u16: case nir_op_u2f16: case nir_op_u2f32: case nir_op_u2u32: case nir_op_pack_32_2x16_split: case nir_op_pack_32_4x8_split: case nir_op_pack_half_2x16_split: return 0; /* we need to handle those here as they only work with 32 bits */ default: if (alu->src[0].src.ssa->bit_size != 1 && alu->src[0].src.ssa->bit_size < 32) return 32; return 0; } } static void v3d_lower_nir(struct v3d_compile *c) { struct nir_lower_tex_options tex_options = { .lower_txd = true, .lower_tg4_offsets = true, .lower_tg4_broadcom_swizzle = true, .lower_rect = false, /* XXX: Use this on V3D 3.x */ .lower_txp = ~0, /* Apply swizzles to all samplers. */ .swizzle_result = ~0, .lower_invalid_implicit_lod = true, }; /* Lower the format swizzle and (for 32-bit returns) * ARB_texture_swizzle-style swizzle. */ assert(c->key->num_tex_used <= ARRAY_SIZE(c->key->tex)); for (int i = 0; i < c->key->num_tex_used; i++) { for (int j = 0; j < 4; j++) tex_options.swizzles[i][j] = c->key->tex[i].swizzle[j]; } tex_options.lower_tex_packing_cb = lower_tex_packing_cb; tex_options.lower_tex_packing_data = c; NIR_PASS(_, c->s, nir_lower_tex, &tex_options); NIR_PASS(_, c->s, nir_lower_system_values); if (c->s->info.zero_initialize_shared_memory && c->s->info.shared_size > 0) { /* All our BOs allocate full pages, so the underlying allocation * for shared memory will always be a multiple of 4KB. This * ensures that we can do an exact number of full chunk_size * writes to initialize the memory independently of the actual * shared_size used by the shader, which is a requirement of * the initialization pass. */ const unsigned chunk_size = 16; /* max single store size */ NIR_PASS(_, c->s, nir_zero_initialize_shared_memory, align(c->s->info.shared_size, chunk_size), chunk_size); } NIR_PASS(_, c->s, nir_lower_compute_system_values, NULL); NIR_PASS(_, c->s, nir_lower_is_helper_invocation); NIR_PASS(_, c->s, v3d_nir_lower_null_pointers); NIR_PASS(_, c->s, nir_lower_bit_size, lower_bit_size_cb, NULL); } static void v3d_set_prog_data_uniforms(struct v3d_compile *c, struct v3d_prog_data *prog_data) { int count = c->num_uniforms; struct v3d_uniform_list *ulist = &prog_data->uniforms; ulist->count = count; ulist->data = ralloc_array(prog_data, uint32_t, count); memcpy(ulist->data, c->uniform_data, count * sizeof(*ulist->data)); ulist->contents = ralloc_array(prog_data, enum quniform_contents, count); memcpy(ulist->contents, c->uniform_contents, count * sizeof(*ulist->contents)); } static void v3d_vs_set_prog_data(struct v3d_compile *c, struct v3d_vs_prog_data *prog_data) { /* The vertex data gets format converted by the VPM so that * each attribute channel takes up a VPM column. Precompute * the sizes for the shader record. */ for (int i = 0; i < ARRAY_SIZE(prog_data->vattr_sizes); i++) { prog_data->vattr_sizes[i] = c->vattr_sizes[i]; prog_data->vpm_input_size += c->vattr_sizes[i]; } memset(prog_data->driver_location_map, -1, sizeof(prog_data->driver_location_map)); nir_foreach_shader_in_variable(var, c->s) { prog_data->driver_location_map[var->data.location] = var->data.driver_location; } prog_data->uses_vid = BITSET_TEST(c->s->info.system_values_read, SYSTEM_VALUE_VERTEX_ID) || BITSET_TEST(c->s->info.system_values_read, SYSTEM_VALUE_VERTEX_ID_ZERO_BASE); prog_data->uses_biid = BITSET_TEST(c->s->info.system_values_read, SYSTEM_VALUE_BASE_INSTANCE); prog_data->uses_iid = BITSET_TEST(c->s->info.system_values_read, SYSTEM_VALUE_INSTANCE_ID) || BITSET_TEST(c->s->info.system_values_read, SYSTEM_VALUE_INSTANCE_INDEX); if (prog_data->uses_vid) prog_data->vpm_input_size++; if (prog_data->uses_biid) prog_data->vpm_input_size++; if (prog_data->uses_iid) prog_data->vpm_input_size++; prog_data->writes_psiz = c->s->info.outputs_written & (1 << VARYING_SLOT_PSIZ); /* Input/output segment size are in sectors (8 rows of 32 bits per * channel). */ prog_data->vpm_input_size = align(prog_data->vpm_input_size, 8) / 8; prog_data->vpm_output_size = align(c->vpm_output_size, 8) / 8; /* Set us up for shared input/output segments. This is apparently * necessary for our VCM setup to avoid varying corruption. * * FIXME: initial testing on V3D 7.1 seems to work fine when using * separate segments. So we could try to reevaluate in the future, if * there is any advantage of using separate segments. */ prog_data->separate_segments = false; prog_data->vpm_output_size = MAX2(prog_data->vpm_output_size, prog_data->vpm_input_size); prog_data->vpm_input_size = 0; /* Compute VCM cache size. We set up our program to take up less than * half of the VPM, so that any set of bin and render programs won't * run out of space. We need space for at least one input segment, * and then allocate the rest to output segments (one for the current * program, the rest to VCM). The valid range of the VCM cache size * field is 1-4 16-vertex batches, but GFXH-1744 limits us to 2-4 * batches. */ assert(c->devinfo->vpm_size); int sector_size = V3D_CHANNELS * sizeof(uint32_t) * 8; int vpm_size_in_sectors = c->devinfo->vpm_size / sector_size; int half_vpm = vpm_size_in_sectors / 2; int vpm_output_sectors = half_vpm - prog_data->vpm_input_size; int vpm_output_batches = vpm_output_sectors / prog_data->vpm_output_size; assert(vpm_output_batches >= 2); prog_data->vcm_cache_size = CLAMP(vpm_output_batches - 1, 2, 4); } static void v3d_gs_set_prog_data(struct v3d_compile *c, struct v3d_gs_prog_data *prog_data) { prog_data->num_inputs = c->num_inputs; memcpy(prog_data->input_slots, c->input_slots, c->num_inputs * sizeof(*c->input_slots)); /* gl_PrimitiveIdIn is written by the GBG into the first word of the * VPM output header automatically and the shader will overwrite * it after reading it if necessary, so it doesn't add to the VPM * size requirements. */ prog_data->uses_pid = BITSET_TEST(c->s->info.system_values_read, SYSTEM_VALUE_PRIMITIVE_ID); /* Output segment size is in sectors (8 rows of 32 bits per channel) */ prog_data->vpm_output_size = align(c->vpm_output_size, 8) / 8; /* Compute SIMD dispatch width and update VPM output size accordingly * to ensure we can fit our program in memory. Available widths are * 16, 8, 4, 1. * * Notice that at draw time we will have to consider VPM memory * requirements from other stages and choose a smaller dispatch * width if needed to fit the program in VPM memory. */ prog_data->simd_width = 16; while ((prog_data->simd_width > 1 && prog_data->vpm_output_size > 16) || prog_data->simd_width == 2) { prog_data->simd_width >>= 1; prog_data->vpm_output_size = align(prog_data->vpm_output_size, 2) / 2; } assert(prog_data->vpm_output_size <= 16); assert(prog_data->simd_width != 2); prog_data->out_prim_type = c->s->info.gs.output_primitive; prog_data->num_invocations = c->s->info.gs.invocations; prog_data->writes_psiz = c->s->info.outputs_written & (1 << VARYING_SLOT_PSIZ); } static void v3d_set_fs_prog_data_inputs(struct v3d_compile *c, struct v3d_fs_prog_data *prog_data) { prog_data->num_inputs = c->num_inputs; memcpy(prog_data->input_slots, c->input_slots, c->num_inputs * sizeof(*c->input_slots)); STATIC_ASSERT(ARRAY_SIZE(prog_data->flat_shade_flags) > (V3D_MAX_FS_INPUTS - 1) / 24); for (int i = 0; i < V3D_MAX_FS_INPUTS; i++) { if (BITSET_TEST(c->flat_shade_flags, i)) prog_data->flat_shade_flags[i / 24] |= 1 << (i % 24); if (BITSET_TEST(c->noperspective_flags, i)) prog_data->noperspective_flags[i / 24] |= 1 << (i % 24); if (BITSET_TEST(c->centroid_flags, i)) prog_data->centroid_flags[i / 24] |= 1 << (i % 24); } } static void v3d_fs_set_prog_data(struct v3d_compile *c, struct v3d_fs_prog_data *prog_data) { v3d_set_fs_prog_data_inputs(c, prog_data); prog_data->writes_z = c->writes_z; prog_data->writes_z_from_fep = c->writes_z_from_fep; prog_data->disable_ez = !c->s->info.fs.early_fragment_tests; prog_data->uses_center_w = c->uses_center_w; prog_data->uses_implicit_point_line_varyings = c->uses_implicit_point_line_varyings; prog_data->lock_scoreboard_on_first_thrsw = c->lock_scoreboard_on_first_thrsw; prog_data->force_per_sample_msaa = c->s->info.fs.uses_sample_shading; prog_data->uses_pid = c->fs_uses_primitive_id; } static void v3d_cs_set_prog_data(struct v3d_compile *c, struct v3d_compute_prog_data *prog_data) { prog_data->shared_size = c->s->info.shared_size; prog_data->local_size[0] = c->s->info.workgroup_size[0]; prog_data->local_size[1] = c->s->info.workgroup_size[1]; prog_data->local_size[2] = c->s->info.workgroup_size[2]; prog_data->has_subgroups = c->has_subgroups; } static void v3d_set_prog_data(struct v3d_compile *c, struct v3d_prog_data *prog_data) { prog_data->threads = c->threads; prog_data->single_seg = !c->last_thrsw; prog_data->spill_size = c->spill_size; prog_data->tmu_spills = c->spills; prog_data->tmu_fills = c->fills; prog_data->tmu_count = c->tmu.total_count; prog_data->qpu_read_stalls = c->qpu_inst_stalled_count; prog_data->compile_strategy_idx = c->compile_strategy_idx; prog_data->tmu_dirty_rcl = c->tmu_dirty_rcl; prog_data->has_control_barrier = c->s->info.uses_control_barrier; prog_data->has_global_address = c->has_global_address; v3d_set_prog_data_uniforms(c, prog_data); switch (c->s->info.stage) { case MESA_SHADER_VERTEX: v3d_vs_set_prog_data(c, (struct v3d_vs_prog_data *)prog_data); break; case MESA_SHADER_GEOMETRY: v3d_gs_set_prog_data(c, (struct v3d_gs_prog_data *)prog_data); break; case MESA_SHADER_FRAGMENT: v3d_fs_set_prog_data(c, (struct v3d_fs_prog_data *)prog_data); break; case MESA_SHADER_COMPUTE: v3d_cs_set_prog_data(c, (struct v3d_compute_prog_data *)prog_data); break; default: unreachable("unsupported shader stage"); } } static uint64_t * v3d_return_qpu_insts(struct v3d_compile *c, uint32_t *final_assembly_size) { *final_assembly_size = c->qpu_inst_count * sizeof(uint64_t); uint64_t *qpu_insts = malloc(*final_assembly_size); if (!qpu_insts) return NULL; memcpy(qpu_insts, c->qpu_insts, *final_assembly_size); vir_compile_destroy(c); return qpu_insts; } static void v3d_nir_lower_vs_early(struct v3d_compile *c) { /* Split our I/O vars and dead code eliminate the unused * components. */ NIR_PASS(_, c->s, nir_lower_io_to_scalar_early, nir_var_shader_in | nir_var_shader_out); uint64_t used_outputs[4] = {0}; for (int i = 0; i < c->vs_key->num_used_outputs; i++) { int slot = v3d_slot_get_slot(c->vs_key->used_outputs[i]); int comp = v3d_slot_get_component(c->vs_key->used_outputs[i]); used_outputs[comp] |= 1ull << slot; } NIR_PASS(_, c->s, nir_remove_unused_io_vars, nir_var_shader_out, used_outputs, NULL); /* demotes to globals */ NIR_PASS(_, c->s, nir_lower_global_vars_to_local); v3d_optimize_nir(c, c->s); NIR_PASS(_, c->s, nir_remove_dead_variables, nir_var_shader_in, NULL); /* This must go before nir_lower_io */ if (c->vs_key->per_vertex_point_size) NIR_PASS(_, c->s, nir_lower_point_size, 1.0f, 0.0f); NIR_PASS(_, c->s, nir_lower_io, nir_var_shader_in | nir_var_shader_out, type_size_vec4, (nir_lower_io_options)0); /* clean up nir_lower_io's deref_var remains and do a constant folding pass * on the code it generated. */ NIR_PASS(_, c->s, nir_opt_dce); NIR_PASS(_, c->s, nir_opt_constant_folding); } static void v3d_nir_lower_gs_early(struct v3d_compile *c) { /* Split our I/O vars and dead code eliminate the unused * components. */ NIR_PASS(_, c->s, nir_lower_io_to_scalar_early, nir_var_shader_in | nir_var_shader_out); uint64_t used_outputs[4] = {0}; for (int i = 0; i < c->gs_key->num_used_outputs; i++) { int slot = v3d_slot_get_slot(c->gs_key->used_outputs[i]); int comp = v3d_slot_get_component(c->gs_key->used_outputs[i]); used_outputs[comp] |= 1ull << slot; } NIR_PASS(_, c->s, nir_remove_unused_io_vars, nir_var_shader_out, used_outputs, NULL); /* demotes to globals */ NIR_PASS(_, c->s, nir_lower_global_vars_to_local); v3d_optimize_nir(c, c->s); NIR_PASS(_, c->s, nir_remove_dead_variables, nir_var_shader_in, NULL); /* This must go before nir_lower_io */ if (c->gs_key->per_vertex_point_size) NIR_PASS(_, c->s, nir_lower_point_size, 1.0f, 0.0f); NIR_PASS(_, c->s, nir_lower_io, nir_var_shader_in | nir_var_shader_out, type_size_vec4, (nir_lower_io_options)0); /* clean up nir_lower_io's deref_var remains and do a constant folding pass * on the code it generated. */ NIR_PASS(_, c->s, nir_opt_dce); NIR_PASS(_, c->s, nir_opt_constant_folding); } static void v3d_fixup_fs_output_types(struct v3d_compile *c) { nir_foreach_shader_out_variable(var, c->s) { uint32_t mask = 0; switch (var->data.location) { case FRAG_RESULT_COLOR: mask = ~0; break; case FRAG_RESULT_DATA0: case FRAG_RESULT_DATA1: case FRAG_RESULT_DATA2: case FRAG_RESULT_DATA3: mask = 1 << (var->data.location - FRAG_RESULT_DATA0); break; } if (c->fs_key->int_color_rb & mask) { var->type = glsl_vector_type(GLSL_TYPE_INT, glsl_get_components(var->type)); } else if (c->fs_key->uint_color_rb & mask) { var->type = glsl_vector_type(GLSL_TYPE_UINT, glsl_get_components(var->type)); } } } static void v3d_nir_lower_fs_early(struct v3d_compile *c) { if (c->fs_key->int_color_rb || c->fs_key->uint_color_rb) v3d_fixup_fs_output_types(c); NIR_PASS(_, c->s, v3d_nir_lower_logic_ops, c); if (c->fs_key->line_smoothing) { NIR_PASS(_, c->s, v3d_nir_lower_line_smooth); NIR_PASS(_, c->s, nir_lower_global_vars_to_local); /* The lowering pass can introduce new sysval reads */ nir_shader_gather_info(c->s, nir_shader_get_entrypoint(c->s)); } } static void v3d_nir_lower_gs_late(struct v3d_compile *c) { if (c->key->ucp_enables) { NIR_PASS(_, c->s, nir_lower_clip_gs, c->key->ucp_enables, true, NULL); } /* Note: GS output scalarizing must happen after nir_lower_clip_gs. */ NIR_PASS_V(c->s, nir_lower_io_to_scalar, nir_var_shader_out, NULL, NULL); } static void v3d_nir_lower_vs_late(struct v3d_compile *c) { if (c->key->ucp_enables) { NIR_PASS(_, c->s, nir_lower_clip_vs, c->key->ucp_enables, false, true, NULL); NIR_PASS_V(c->s, nir_lower_io_to_scalar, nir_var_shader_out, NULL, NULL); } /* Note: VS output scalarizing must happen after nir_lower_clip_vs. */ NIR_PASS_V(c->s, nir_lower_io_to_scalar, nir_var_shader_out, NULL, NULL); } static void v3d_nir_lower_fs_late(struct v3d_compile *c) { /* In OpenGL the fragment shader can't read gl_ClipDistance[], but * Vulkan allows it, in which case the SPIR-V compiler will declare * VARING_SLOT_CLIP_DIST0 as compact array variable. Pass true as * the last parameter to always operate with a compact array in both * OpenGL and Vulkan so we do't have to care about the API we * are using. */ if (c->key->ucp_enables) NIR_PASS(_, c->s, nir_lower_clip_fs, c->key->ucp_enables, true); NIR_PASS_V(c->s, nir_lower_io_to_scalar, nir_var_shader_in, NULL, NULL); } static uint32_t vir_get_max_temps(struct v3d_compile *c) { int max_ip = 0; vir_for_each_inst_inorder(inst, c) max_ip++; uint32_t *pressure = rzalloc_array(NULL, uint32_t, max_ip); for (int t = 0; t < c->num_temps; t++) { for (int i = c->temp_start[t]; (i < c->temp_end[t] && i < max_ip); i++) { if (i > max_ip) break; pressure[i]++; } } uint32_t max_temps = 0; for (int i = 0; i < max_ip; i++) max_temps = MAX2(max_temps, pressure[i]); ralloc_free(pressure); return max_temps; } enum v3d_dependency_class { V3D_DEPENDENCY_CLASS_GS_VPM_OUTPUT_0 }; static bool v3d_intrinsic_dependency_cb(nir_intrinsic_instr *intr, nir_schedule_dependency *dep, void *user_data) { struct v3d_compile *c = user_data; switch (intr->intrinsic) { case nir_intrinsic_store_output: /* Writing to location 0 overwrites the value passed in for * gl_PrimitiveID on geometry shaders */ if (c->s->info.stage != MESA_SHADER_GEOMETRY || nir_intrinsic_base(intr) != 0) break; nir_const_value *const_value = nir_src_as_const_value(intr->src[1]); if (const_value == NULL) break; uint64_t offset = nir_const_value_as_uint(*const_value, nir_src_bit_size(intr->src[1])); if (offset != 0) break; dep->klass = V3D_DEPENDENCY_CLASS_GS_VPM_OUTPUT_0; dep->type = NIR_SCHEDULE_WRITE_DEPENDENCY; return true; case nir_intrinsic_load_primitive_id: if (c->s->info.stage != MESA_SHADER_GEOMETRY) break; dep->klass = V3D_DEPENDENCY_CLASS_GS_VPM_OUTPUT_0; dep->type = NIR_SCHEDULE_READ_DEPENDENCY; return true; default: break; } return false; } static unsigned v3d_instr_delay_cb(nir_instr *instr, void *data) { struct v3d_compile *c = (struct v3d_compile *) data; switch (instr->type) { case nir_instr_type_undef: case nir_instr_type_load_const: case nir_instr_type_alu: case nir_instr_type_deref: case nir_instr_type_jump: case nir_instr_type_parallel_copy: case nir_instr_type_call: case nir_instr_type_phi: return 1; /* We should not use very large delays for TMU instructions. Typically, * thread switches will be sufficient to hide all or most of the latency, * so we typically only need a little bit of extra room. If we over-estimate * the latency here we may end up unnecessarily delaying the critical path in * the shader, which would have a negative effect in performance, so here * we are trying to strike a balance based on empirical testing. */ case nir_instr_type_intrinsic: { nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr); if (!c->disable_general_tmu_sched) { switch (intr->intrinsic) { case nir_intrinsic_decl_reg: case nir_intrinsic_load_reg: case nir_intrinsic_store_reg: return 0; case nir_intrinsic_load_ssbo: case nir_intrinsic_load_scratch: case nir_intrinsic_load_shared: case nir_intrinsic_image_load: return 3; case nir_intrinsic_load_ubo: if (nir_src_is_divergent(intr->src[1])) return 3; FALLTHROUGH; default: return 1; } } else { switch (intr->intrinsic) { case nir_intrinsic_decl_reg: case nir_intrinsic_load_reg: case nir_intrinsic_store_reg: return 0; default: return 1; } } break; } case nir_instr_type_tex: return 5; case nir_instr_type_debug_info: return 0; } return 0; } static bool should_split_wrmask(const nir_instr *instr, const void *data) { nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr); switch (intr->intrinsic) { case nir_intrinsic_store_ssbo: case nir_intrinsic_store_shared: case nir_intrinsic_store_global: case nir_intrinsic_store_scratch: return true; default: return false; } } static nir_intrinsic_instr * nir_instr_as_constant_ubo_load(nir_instr *inst) { if (inst->type != nir_instr_type_intrinsic) return NULL; nir_intrinsic_instr *intr = nir_instr_as_intrinsic(inst); if (intr->intrinsic != nir_intrinsic_load_ubo) return NULL; assert(nir_src_is_const(intr->src[0])); if (!nir_src_is_const(intr->src[1])) return NULL; return intr; } static bool v3d_nir_sort_constant_ubo_load(nir_block *block, nir_intrinsic_instr *ref) { bool progress = false; nir_instr *ref_inst = &ref->instr; uint32_t ref_offset = nir_src_as_uint(ref->src[1]); uint32_t ref_index = nir_src_as_uint(ref->src[0]); /* Go through all instructions after ref searching for constant UBO * loads for the same UBO index. */ bool seq_break = false; nir_instr *inst = &ref->instr; nir_instr *next_inst = NULL; while (true) { inst = next_inst ? next_inst : nir_instr_next(inst); if (!inst) break; next_inst = NULL; if (inst->type != nir_instr_type_intrinsic) continue; nir_intrinsic_instr *intr = nir_instr_as_intrinsic(inst); if (intr->intrinsic != nir_intrinsic_load_ubo) continue; /* We only produce unifa sequences for non-divergent loads */ if (nir_src_is_divergent(intr->src[1])) continue; /* If there are any UBO loads that are not constant or that * use a different UBO index in between the reference load and * any other constant load for the same index, they would break * the unifa sequence. We will flag that so we can then move * all constant UBO loads for the reference index before these * and not just the ones that are not ordered to avoid breaking * the sequence and reduce unifa writes. */ if (!nir_src_is_const(intr->src[1])) { seq_break = true; continue; } uint32_t offset = nir_src_as_uint(intr->src[1]); assert(nir_src_is_const(intr->src[0])); uint32_t index = nir_src_as_uint(intr->src[0]); if (index != ref_index) { seq_break = true; continue; } /* Only move loads with an offset that is close enough to the * reference offset, since otherwise we would not be able to * skip the unifa write for them. See ntq_emit_load_ubo_unifa. */ if (abs((int)(ref_offset - offset)) > MAX_UNIFA_SKIP_DISTANCE) continue; /* We will move this load if its offset is smaller than ref's * (in which case we will move it before ref) or if the offset * is larger than ref's but there are sequence breakers in * in between (in which case we will move it after ref and * before the sequence breakers). */ if (!seq_break && offset >= ref_offset) continue; /* Find where exactly we want to move this load: * * If we are moving it before ref, we want to check any other * UBO loads we placed before ref and make sure we insert this * one properly ordered with them. Likewise, if we are moving * it after ref. */ nir_instr *pos = ref_inst; nir_instr *tmp = pos; do { if (offset < ref_offset) tmp = nir_instr_prev(tmp); else tmp = nir_instr_next(tmp); if (!tmp || tmp == inst) break; /* Ignore non-unifa UBO loads */ if (tmp->type != nir_instr_type_intrinsic) continue; nir_intrinsic_instr *tmp_intr = nir_instr_as_intrinsic(tmp); if (tmp_intr->intrinsic != nir_intrinsic_load_ubo) continue; if (nir_src_is_divergent(tmp_intr->src[1])) continue; /* Stop if we find a unifa UBO load that breaks the * sequence. */ if (!nir_src_is_const(tmp_intr->src[1])) break; if (nir_src_as_uint(tmp_intr->src[0]) != index) break; uint32_t tmp_offset = nir_src_as_uint(tmp_intr->src[1]); if (offset < ref_offset) { if (tmp_offset < offset || tmp_offset >= ref_offset) { break; } else { pos = tmp; } } else { if (tmp_offset > offset || tmp_offset <= ref_offset) { break; } else { pos = tmp; } } } while (true); /* We can't move the UBO load before the instruction that * defines its constant offset. If that instruction is placed * in between the new location (pos) and the current location * of this load, we will have to move that instruction too. * * We don't care about the UBO index definition because that * is optimized to be reused by all UBO loads for the same * index and therefore is certain to be defined before the * first UBO load that uses it. */ nir_instr *offset_inst = NULL; tmp = inst; while ((tmp = nir_instr_prev(tmp)) != NULL) { if (pos == tmp) { /* We reached the target location without * finding the instruction that defines the * offset, so that instruction must be before * the new position and we don't have to fix it. */ break; } if (intr->src[1].ssa->parent_instr == tmp) { offset_inst = tmp; break; } } if (offset_inst) { exec_node_remove(&offset_inst->node); exec_node_insert_node_before(&pos->node, &offset_inst->node); } /* Since we are moving the instruction before its current * location, grab its successor before the move so that * we can continue the next iteration of the main loop from * that instruction. */ next_inst = nir_instr_next(inst); /* Move this load to the selected location */ exec_node_remove(&inst->node); if (offset < ref_offset) exec_node_insert_node_before(&pos->node, &inst->node); else exec_node_insert_after(&pos->node, &inst->node); progress = true; } return progress; } static bool v3d_nir_sort_constant_ubo_loads_block(struct v3d_compile *c, nir_block *block) { bool progress = false; bool local_progress; do { local_progress = false; nir_foreach_instr_safe(inst, block) { nir_intrinsic_instr *intr = nir_instr_as_constant_ubo_load(inst); if (intr) { local_progress |= v3d_nir_sort_constant_ubo_load(block, intr); } } progress |= local_progress; } while (local_progress); return progress; } /** * Sorts constant UBO loads in each block by offset to maximize chances of * skipping unifa writes when converting to VIR. This can increase register * pressure. */ static bool v3d_nir_sort_constant_ubo_loads(nir_shader *s, struct v3d_compile *c) { nir_foreach_function_impl(impl, s) { nir_foreach_block(block, impl) { c->sorted_any_ubo_loads |= v3d_nir_sort_constant_ubo_loads_block(c, block); } nir_metadata_preserve(impl, nir_metadata_control_flow); } return c->sorted_any_ubo_loads; } static void lower_load_num_subgroups(struct v3d_compile *c, nir_builder *b, nir_intrinsic_instr *intr) { assert(c->s->info.stage == MESA_SHADER_COMPUTE); assert(intr->intrinsic == nir_intrinsic_load_num_subgroups); b->cursor = nir_after_instr(&intr->instr); uint32_t num_subgroups = DIV_ROUND_UP(c->s->info.workgroup_size[0] * c->s->info.workgroup_size[1] * c->s->info.workgroup_size[2], V3D_CHANNELS); nir_def *result = nir_imm_int(b, num_subgroups); nir_def_replace(&intr->def, result); } static bool lower_subgroup_intrinsics(struct v3d_compile *c, nir_block *block, nir_builder *b) { bool progress = false; nir_foreach_instr_safe(inst, block) { if (inst->type != nir_instr_type_intrinsic) continue;; nir_intrinsic_instr *intr = nir_instr_as_intrinsic(inst); if (!intr) continue; switch (intr->intrinsic) { case nir_intrinsic_load_num_subgroups: lower_load_num_subgroups(c, b, intr); progress = true; FALLTHROUGH; case nir_intrinsic_load_subgroup_id: case nir_intrinsic_load_subgroup_size: case nir_intrinsic_load_subgroup_invocation: case nir_intrinsic_elect: case nir_intrinsic_ballot: case nir_intrinsic_inverse_ballot: case nir_intrinsic_ballot_bitfield_extract: case nir_intrinsic_ballot_bit_count_reduce: case nir_intrinsic_ballot_find_lsb: case nir_intrinsic_ballot_find_msb: case nir_intrinsic_ballot_bit_count_exclusive: case nir_intrinsic_ballot_bit_count_inclusive: case nir_intrinsic_reduce: case nir_intrinsic_inclusive_scan: case nir_intrinsic_exclusive_scan: case nir_intrinsic_read_invocation: case nir_intrinsic_read_first_invocation: case nir_intrinsic_load_subgroup_eq_mask: case nir_intrinsic_load_subgroup_ge_mask: case nir_intrinsic_load_subgroup_gt_mask: case nir_intrinsic_load_subgroup_le_mask: case nir_intrinsic_load_subgroup_lt_mask: case nir_intrinsic_shuffle: case nir_intrinsic_shuffle_xor: case nir_intrinsic_shuffle_up: case nir_intrinsic_shuffle_down: case nir_intrinsic_vote_all: case nir_intrinsic_vote_any: case nir_intrinsic_vote_feq: case nir_intrinsic_vote_ieq: case nir_intrinsic_quad_broadcast: case nir_intrinsic_quad_swap_horizontal: case nir_intrinsic_quad_swap_vertical: case nir_intrinsic_quad_swap_diagonal: c->has_subgroups = true; break; default: break; } } return progress; } static bool v3d_nir_lower_subgroup_intrinsics(nir_shader *s, struct v3d_compile *c) { bool progress = false; nir_foreach_function_impl(impl, s) { nir_builder b = nir_builder_create(impl); nir_foreach_block(block, impl) progress |= lower_subgroup_intrinsics(c, block, &b); nir_metadata_preserve(impl, nir_metadata_control_flow); } return progress; } static void v3d_attempt_compile(struct v3d_compile *c) { switch (c->s->info.stage) { case MESA_SHADER_VERTEX: c->vs_key = (struct v3d_vs_key *) c->key; break; case MESA_SHADER_GEOMETRY: c->gs_key = (struct v3d_gs_key *) c->key; break; case MESA_SHADER_FRAGMENT: c->fs_key = (struct v3d_fs_key *) c->key; break; case MESA_SHADER_COMPUTE: break; default: unreachable("unsupported shader stage"); } switch (c->s->info.stage) { case MESA_SHADER_VERTEX: v3d_nir_lower_vs_early(c); break; case MESA_SHADER_GEOMETRY: v3d_nir_lower_gs_early(c); break; case MESA_SHADER_FRAGMENT: v3d_nir_lower_fs_early(c); break; default: break; } v3d_lower_nir(c); switch (c->s->info.stage) { case MESA_SHADER_VERTEX: v3d_nir_lower_vs_late(c); break; case MESA_SHADER_GEOMETRY: v3d_nir_lower_gs_late(c); break; case MESA_SHADER_FRAGMENT: v3d_nir_lower_fs_late(c); break; default: break; } NIR_PASS(_, c->s, v3d_nir_lower_io, c); NIR_PASS(_, c->s, v3d_nir_lower_txf_ms); NIR_PASS(_, c->s, v3d_nir_lower_image_load_store, c); NIR_PASS(_, c->s, nir_opt_idiv_const, 8); nir_lower_idiv_options idiv_options = { .allow_fp16 = true, }; NIR_PASS(_, c->s, nir_lower_idiv, &idiv_options); NIR_PASS(_, c->s, nir_lower_alu); if (c->key->robust_uniform_access || c->key->robust_storage_access || c->key->robust_image_access) { /* nir_lower_robust_access assumes constant buffer * indices on ubo/ssbo intrinsics so run copy propagation and * constant folding passes before we run the lowering to warrant * this. We also want to run the lowering before v3d_optimize to * clean-up redundant get_buffer_size calls produced in the pass. */ NIR_PASS(_, c->s, nir_copy_prop); NIR_PASS(_, c->s, nir_opt_constant_folding); nir_lower_robust_access_options opts = { .lower_image = c->key->robust_image_access, .lower_ssbo = c->key->robust_storage_access, .lower_ubo = c->key->robust_uniform_access, }; NIR_PASS(_, c->s, nir_lower_robust_access, &opts); } NIR_PASS(_, c->s, nir_lower_vars_to_scratch, nir_var_function_temp, 0, glsl_get_natural_size_align_bytes, glsl_get_natural_size_align_bytes); NIR_PASS(_, c->s, v3d_nir_lower_global_2x32); NIR_PASS(_, c->s, nir_lower_wrmasks, should_split_wrmask, c->s); NIR_PASS(_, c->s, v3d_nir_lower_load_store_bitsize); NIR_PASS(_, c->s, v3d_nir_lower_scratch); /* needs to run after load_store_bitsize */ NIR_PASS(_, c->s, nir_lower_pack); NIR_PASS(_, c->s, v3d_nir_lower_subgroup_intrinsics, c); const nir_lower_subgroups_options subgroup_opts = { .subgroup_size = V3D_CHANNELS, .ballot_components = 1, .ballot_bit_size = 32, .lower_to_scalar = true, .lower_inverse_ballot = true, .lower_subgroup_masks = true, .lower_relative_shuffle = true, .lower_quad = true, }; NIR_PASS(_, c->s, nir_lower_subgroups, &subgroup_opts); v3d_optimize_nir(c, c->s); /* Do late algebraic optimization to turn add(a, neg(b)) back into * subs, then the mandatory cleanup after algebraic. Note that it may * produce fnegs, and if so then we need to keep running to squash * fneg(fneg(a)). */ bool more_late_algebraic = true; while (more_late_algebraic) { more_late_algebraic = false; NIR_PASS(more_late_algebraic, c->s, nir_opt_algebraic_late); NIR_PASS(_, c->s, nir_opt_constant_folding); NIR_PASS(_, c->s, nir_copy_prop); NIR_PASS(_, c->s, nir_opt_dce); NIR_PASS(_, c->s, nir_opt_cse); } NIR_PASS(_, c->s, nir_lower_bool_to_int32); NIR_PASS(_, c->s, nir_convert_to_lcssa, true, true); NIR_PASS_V(c->s, nir_divergence_analysis); NIR_PASS(_, c->s, nir_convert_from_ssa, true); struct nir_schedule_options schedule_options = { /* Schedule for about half our register space, to enable more * shaders to hit 4 threads. */ .threshold = c->threads == 4 ? 24 : 48, /* Vertex shaders share the same memory for inputs and outputs, * fragment and geometry shaders do not. */ .stages_with_shared_io_memory = (((1 << MESA_ALL_SHADER_STAGES) - 1) & ~((1 << MESA_SHADER_FRAGMENT) | (1 << MESA_SHADER_GEOMETRY))), .fallback = c->fallback_scheduler, .intrinsic_cb = v3d_intrinsic_dependency_cb, .intrinsic_cb_data = c, .instr_delay_cb = v3d_instr_delay_cb, .instr_delay_cb_data = c, }; NIR_PASS_V(c->s, nir_schedule, &schedule_options); if (!c->disable_constant_ubo_load_sorting) NIR_PASS(_, c->s, v3d_nir_sort_constant_ubo_loads, c); const nir_move_options buffer_opts = c->move_buffer_loads ? (nir_move_load_ubo | nir_move_load_ssbo) : 0; NIR_PASS(_, c->s, nir_opt_move, nir_move_load_uniform | nir_move_const_undef | buffer_opts); NIR_PASS_V(c->s, nir_trivialize_registers); v3d_nir_to_vir(c); } uint32_t v3d_prog_data_size(gl_shader_stage stage) { static const int prog_data_size[] = { [MESA_SHADER_VERTEX] = sizeof(struct v3d_vs_prog_data), [MESA_SHADER_GEOMETRY] = sizeof(struct v3d_gs_prog_data), [MESA_SHADER_FRAGMENT] = sizeof(struct v3d_fs_prog_data), [MESA_SHADER_COMPUTE] = sizeof(struct v3d_compute_prog_data), }; assert(stage >= 0 && stage < ARRAY_SIZE(prog_data_size) && prog_data_size[stage]); return prog_data_size[stage]; } int v3d_shaderdb_dump(struct v3d_compile *c, char **shaderdb_str) { if (c == NULL || c->compilation_result != V3D_COMPILATION_SUCCEEDED) return -1; return asprintf(shaderdb_str, "%s shader: %d inst, %d threads, %d loops, " "%d uniforms, %d max-temps, %d:%d spills:fills, " "%d sfu-stalls, %d inst-and-stalls, %d nops", vir_get_stage_name(c), c->qpu_inst_count, c->threads, c->loops, c->num_uniforms, vir_get_max_temps(c), c->spills, c->fills, c->qpu_inst_stalled_count, c->qpu_inst_count + c->qpu_inst_stalled_count, c->nop_count); } /* This is a list of incremental changes to the compilation strategy * that will be used to try to compile the shader successfully. The * default strategy is to enable all optimizations which will have * the highest register pressure but is expected to produce most * optimal code. Following strategies incrementally disable specific * optimizations that are known to contribute to register pressure * in order to be able to compile the shader successfully while meeting * thread count requirements. * * V3D 4.1+ has a min thread count of 2, but we can use 1 here to also * cover previous hardware as well (meaning that we are not limiting * register allocation to any particular thread count). This is fine * because v3d_nir_to_vir will cap this to the actual minimum. */ static const struct v3d_compiler_strategy strategies[] = { /*0*/ { "default", 4, 4, false, false, false, false, false, false, 0 }, /*1*/ { "disable general TMU sched", 4, 4, true, false, false, false, false, false, 0 }, /*2*/ { "disable gcm", 4, 4, true, true, false, false, false, false, 0 }, /*3*/ { "disable loop unrolling", 4, 4, true, true, true, false, false, false, 0 }, /*4*/ { "disable UBO load sorting", 4, 4, true, true, true, true, false, false, 0 }, /*5*/ { "disable TMU pipelining", 4, 4, true, true, true, true, false, true, 0 }, /*6*/ { "lower thread count", 2, 1, false, false, false, false, false, false, -1 }, /*7*/ { "disable general TMU sched (2t)", 2, 1, true, false, false, false, false, false, -1 }, /*8*/ { "disable gcm (2t)", 2, 1, true, true, false, false, false, false, -1 }, /*9*/ { "disable loop unrolling (2t)", 2, 1, true, true, true, false, false, false, -1 }, /*10*/ { "Move buffer loads (2t)", 2, 1, true, true, true, true, true, false, -1 }, /*11*/ { "disable TMU pipelining (2t)", 2, 1, true, true, true, true, true, true, -1 }, /*12*/ { "fallback scheduler", 2, 1, true, true, true, true, true, true, -1 } }; /** * If a particular optimization didn't make any progress during a compile * attempt disabling it alone won't allow us to compile the shader successfully, * since we'll end up with the same code. Detect these scenarios so we can * avoid wasting time with useless compiles. We should also consider if the * gy changes other aspects of the compilation process though, like * spilling, and not skip it in that case. */ static bool skip_compile_strategy(struct v3d_compile *c, uint32_t idx) { /* We decide if we can skip a strategy based on the optimizations that * were active in the previous strategy, so we should only be calling this * for strategies after the first. */ assert(idx > 0); /* Don't skip a strategy that changes spilling behavior */ if (strategies[idx].max_tmu_spills != strategies[idx - 1].max_tmu_spills) { return false; } switch (idx) { /* General TMU sched.: skip if we didn't emit any TMU loads */ case 1: case 7: return !c->has_general_tmu_load; /* Global code motion: skip if nir_opt_gcm didn't make any progress */ case 2: case 8: return !c->gcm_progress; /* Loop unrolling: skip if we didn't unroll any loops */ case 3: case 9: return !c->unrolled_any_loops; /* UBO load sorting: skip if we didn't sort any loads */ case 4: return !c->sorted_any_ubo_loads; /* Move buffer loads: we assume any shader with difficult RA * most likely has UBO / SSBO loads so we never try to skip. * For now, we only try this for 2-thread compiles since it * is expected to impact instruction counts and latency. */ case 10: assert(c->threads < 4); return false; /* TMU pipelining: skip if we didn't pipeline any TMU ops */ case 5: case 11: return !c->pipelined_any_tmu; /* Lower thread count: skip if we already tried less that 4 threads */ case 6: return c->threads < 4; default: return false; }; } static inline void set_best_compile(struct v3d_compile **best, struct v3d_compile *c) { if (*best) vir_compile_destroy(*best); *best = c; } uint64_t *v3d_compile(const struct v3d_compiler *compiler, struct v3d_key *key, struct v3d_prog_data **out_prog_data, nir_shader *s, void (*debug_output)(const char *msg, void *debug_output_data), void *debug_output_data, int program_id, int variant_id, uint32_t *final_assembly_size) { struct v3d_compile *c = NULL; uint32_t best_spill_fill_count = UINT32_MAX; struct v3d_compile *best_c = NULL; for (int32_t strat = 0; strat < ARRAY_SIZE(strategies); strat++) { /* Fallback strategy */ if (strat > 0) { assert(c); if (skip_compile_strategy(c, strat)) continue; char *debug_msg; int ret = asprintf(&debug_msg, "Falling back to strategy '%s' " "for %s prog %d/%d", strategies[strat].name, vir_get_stage_name(c), c->program_id, c->variant_id); if (ret >= 0) { if (V3D_DBG(PERF)) fprintf(stderr, "%s\n", debug_msg); c->debug_output(debug_msg, c->debug_output_data); free(debug_msg); } if (c != best_c) vir_compile_destroy(c); } c = vir_compile_init(compiler, key, s, debug_output, debug_output_data, program_id, variant_id, strat, &strategies[strat], strat == ARRAY_SIZE(strategies) - 1); v3d_attempt_compile(c); /* Broken shader or driver bug */ if (c->compilation_result == V3D_COMPILATION_FAILED) break; /* If we compiled without spills, choose this. * Otherwise if this is a 4-thread compile, choose this (these * have a very low cap on the allowed TMU spills so we assume * it will be better than a 2-thread compile without spills). * Otherwise, keep going while tracking the strategy with the * lowest spill count. */ if (c->compilation_result == V3D_COMPILATION_SUCCEEDED) { if (c->spills == 0 || strategies[strat].min_threads == 4 || V3D_DBG(OPT_COMPILE_TIME)) { set_best_compile(&best_c, c); break; } else if (c->spills + c->fills < best_spill_fill_count) { set_best_compile(&best_c, c); best_spill_fill_count = c->spills + c->fills; } if (V3D_DBG(PERF)) { char *debug_msg; int ret = asprintf(&debug_msg, "Compiled %s prog %d/%d with %d " "spills and %d fills. Will try " "more strategies.", vir_get_stage_name(c), c->program_id, c->variant_id, c->spills, c->fills); if (ret >= 0) { fprintf(stderr, "%s\n", debug_msg); c->debug_output(debug_msg, c->debug_output_data); free(debug_msg); } } } /* Only try next streategy if we failed to register allocate * or we had to spill. */ assert(c->compilation_result == V3D_COMPILATION_FAILED_REGISTER_ALLOCATION || c->spills > 0); } /* If the best strategy was not the last, choose that */ if (best_c && c != best_c) set_best_compile(&c, best_c); if (V3D_DBG(PERF) && c->compilation_result != V3D_COMPILATION_FAILED_REGISTER_ALLOCATION && c->spills > 0) { char *debug_msg; int ret = asprintf(&debug_msg, "Compiled %s prog %d/%d with %d " "spills and %d fills", vir_get_stage_name(c), c->program_id, c->variant_id, c->spills, c->fills); fprintf(stderr, "%s\n", debug_msg); if (ret >= 0) { c->debug_output(debug_msg, c->debug_output_data); free(debug_msg); } } if (c->compilation_result != V3D_COMPILATION_SUCCEEDED) { fprintf(stderr, "Failed to compile %s prog %d/%d " "with any strategy.\n", vir_get_stage_name(c), c->program_id, c->variant_id); vir_compile_destroy(c); return NULL; } struct v3d_prog_data *prog_data; prog_data = rzalloc_size(NULL, v3d_prog_data_size(c->s->info.stage)); v3d_set_prog_data(c, prog_data); *out_prog_data = prog_data; char *shaderdb; int ret = v3d_shaderdb_dump(c, &shaderdb); if (ret >= 0) { if (V3D_DBG(SHADERDB)) fprintf(stderr, "SHADER-DB-%s - %s\n", s->info.name, shaderdb); c->debug_output(shaderdb, c->debug_output_data); free(shaderdb); } return v3d_return_qpu_insts(c, final_assembly_size); } void vir_remove_instruction(struct v3d_compile *c, struct qinst *qinst) { if (qinst->dst.file == QFILE_TEMP) c->defs[qinst->dst.index] = NULL; assert(&qinst->link != c->cursor.link); list_del(&qinst->link); free(qinst); c->live_intervals_valid = false; } struct qreg vir_follow_movs(struct v3d_compile *c, struct qreg reg) { /* XXX int pack = reg.pack; while (reg.file == QFILE_TEMP && c->defs[reg.index] && (c->defs[reg.index]->op == QOP_MOV || c->defs[reg.index]->op == QOP_FMOV) && !c->defs[reg.index]->dst.pack && !c->defs[reg.index]->src[0].pack) { reg = c->defs[reg.index]->src[0]; } reg.pack = pack; */ return reg; } void vir_compile_destroy(struct v3d_compile *c) { /* Defuse the assert that we aren't removing the cursor's instruction. */ c->cursor.link = NULL; vir_for_each_block(block, c) { while (!list_is_empty(&block->instructions)) { struct qinst *qinst = list_first_entry(&block->instructions, struct qinst, link); vir_remove_instruction(c, qinst); } } ralloc_free(c); } uint32_t vir_get_uniform_index(struct v3d_compile *c, enum quniform_contents contents, uint32_t data) { for (int i = 0; i < c->num_uniforms; i++) { if (c->uniform_contents[i] == contents && c->uniform_data[i] == data) { return i; } } uint32_t uniform = c->num_uniforms++; if (uniform >= c->uniform_array_size) { c->uniform_array_size = MAX2(MAX2(16, uniform + 1), c->uniform_array_size * 2); c->uniform_data = reralloc(c, c->uniform_data, uint32_t, c->uniform_array_size); c->uniform_contents = reralloc(c, c->uniform_contents, enum quniform_contents, c->uniform_array_size); } c->uniform_contents[uniform] = contents; c->uniform_data[uniform] = data; return uniform; } /* Looks back into the current block to find the ldunif that wrote the uniform * at the requested index. If it finds it, it returns true and writes the * destination register of the ldunif instruction to 'unif'. * * This can impact register pressure and end up leading to worse code, so we * limit the number of instructions we are willing to look back through to * strike a good balance. */ static bool try_opt_ldunif(struct v3d_compile *c, uint32_t index, struct qreg *unif) { uint32_t count = 20; struct qinst *prev_inst = NULL; assert(c->cur_block); #if MESA_DEBUG /* We can only reuse a uniform if it was emitted in the same block, * so callers must make sure the current instruction is being emitted * in the current block. */ bool found = false; vir_for_each_inst(inst, c->cur_block) { if (&inst->link == c->cursor.link) { found = true; break; } } assert(found || &c->cur_block->instructions == c->cursor.link); #endif list_for_each_entry_from_rev(struct qinst, inst, c->cursor.link->prev, &c->cur_block->instructions, link) { if ((inst->qpu.sig.ldunif || inst->qpu.sig.ldunifrf) && inst->uniform == index) { prev_inst = inst; break; } if (--count == 0) break; } if (!prev_inst) return false; /* Only reuse the ldunif result if it was written to a temp register, * otherwise there may be special restrictions (for example, ldunif * may write directly to unifa, which is a write-only register). */ if (prev_inst->dst.file != QFILE_TEMP) return false; list_for_each_entry_from(struct qinst, inst, prev_inst->link.next, &c->cur_block->instructions, link) { if (inst->dst.file == prev_inst->dst.file && inst->dst.index == prev_inst->dst.index) { return false; } } *unif = prev_inst->dst; return true; } struct qreg vir_uniform(struct v3d_compile *c, enum quniform_contents contents, uint32_t data) { const int num_uniforms = c->num_uniforms; const int index = vir_get_uniform_index(c, contents, data); /* If this is not the first time we see this uniform try to reuse the * result of the last ldunif that loaded it. */ const bool is_new_uniform = num_uniforms != c->num_uniforms; if (!is_new_uniform && !c->disable_ldunif_opt) { struct qreg ldunif_dst; if (try_opt_ldunif(c, index, &ldunif_dst)) return ldunif_dst; } struct qinst *inst = vir_NOP(c); inst->qpu.sig.ldunif = true; inst->uniform = index; inst->dst = vir_get_temp(c); c->defs[inst->dst.index] = inst; return inst->dst; } #define OPTPASS(func) \ do { \ bool stage_progress = func(c); \ if (stage_progress) { \ progress = true; \ if (print_opt_debug) { \ fprintf(stderr, \ "VIR opt pass %2d: %s progress\n", \ pass, #func); \ } \ /*XXX vir_validate(c);*/ \ } \ } while (0) void vir_optimize(struct v3d_compile *c) { bool print_opt_debug = false; int pass = 1; while (true) { bool progress = false; OPTPASS(vir_opt_copy_propagate); OPTPASS(vir_opt_redundant_flags); OPTPASS(vir_opt_dead_code); OPTPASS(vir_opt_small_immediates); OPTPASS(vir_opt_constant_alu); if (!progress) break; pass++; } } const char * vir_get_stage_name(struct v3d_compile *c) { if (c->vs_key && c->vs_key->is_coord) return "MESA_SHADER_VERTEX_BIN"; else if (c->gs_key && c->gs_key->is_coord) return "MESA_SHADER_GEOMETRY_BIN"; else return gl_shader_stage_name(c->s->info.stage); } static inline uint32_t compute_vpm_size_in_sectors(const struct v3d_device_info *devinfo) { assert(devinfo->vpm_size > 0); const uint32_t sector_size = V3D_CHANNELS * sizeof(uint32_t) * 8; return devinfo->vpm_size / sector_size; } /* Computes various parameters affecting VPM memory configuration for programs * involving geometry shaders to ensure the program fits in memory and honors * requirements described in section "VPM usage" of the programming manual. */ static bool compute_vpm_config_gs(struct v3d_device_info *devinfo, struct v3d_vs_prog_data *vs, struct v3d_gs_prog_data *gs, struct vpm_config *vpm_cfg_out) { const uint32_t A = vs->separate_segments ? 1 : 0; const uint32_t Ad = vs->vpm_input_size; const uint32_t Vd = vs->vpm_output_size; const uint32_t vpm_size = compute_vpm_size_in_sectors(devinfo); /* Try to fit program into our VPM memory budget by adjusting * configurable parameters iteratively. We do this in two phases: * the first phase tries to fit the program into the total available * VPM memory. If we succeed at that, then the second phase attempts * to fit the program into half of that budget so we can run bin and * render programs in parallel. */ struct vpm_config vpm_cfg[2]; struct vpm_config *final_vpm_cfg = NULL; uint32_t phase = 0; vpm_cfg[phase].As = 1; vpm_cfg[phase].Gs = 1; vpm_cfg[phase].Gd = gs->vpm_output_size; vpm_cfg[phase].gs_width = gs->simd_width; /* While there is a requirement that Vc >= [Vn / 16], this is * always the case when tessellation is not present because in that * case Vn can only be 6 at most (when input primitive is triangles * with adjacency). * * We always choose Vc=2. We can't go lower than this due to GFXH-1744, * and Broadcom has not found it worth it to increase it beyond this * in general. Increasing Vc also increases VPM memory pressure which * can turn up being detrimental for performance in some scenarios. */ vpm_cfg[phase].Vc = 2; /* Gv is a constraint on the hardware to not exceed the * specified number of vertex segments per GS batch. If adding a * new primitive to a GS batch would result in a range of more * than Gv vertex segments being referenced by the batch, then * the hardware will flush the batch and start a new one. This * means that we can choose any value we want, we just need to * be aware that larger values improve GS batch utilization * at the expense of more VPM memory pressure (which can affect * other performance aspects, such as GS dispatch width). * We start with the largest value, and will reduce it if we * find that total memory pressure is too high. */ vpm_cfg[phase].Gv = 3; do { /* When GS is present in absence of TES, then we need to satisfy * that Ve >= Gv. We go with the smallest value of Ve to avoid * increasing memory pressure. */ vpm_cfg[phase].Ve = vpm_cfg[phase].Gv; uint32_t vpm_sectors = A * vpm_cfg[phase].As * Ad + (vpm_cfg[phase].Vc + vpm_cfg[phase].Ve) * Vd + vpm_cfg[phase].Gs * vpm_cfg[phase].Gd; /* Ideally we want to use no more than half of the available * memory so we can execute a bin and render program in parallel * without stalls. If we achieved that then we are done. */ if (vpm_sectors <= vpm_size / 2) { final_vpm_cfg = &vpm_cfg[phase]; break; } /* At the very least, we should not allocate more than the * total available VPM memory. If we have a configuration that * succeeds at this we save it and continue to see if we can * meet the half-memory-use criteria too. */ if (phase == 0 && vpm_sectors <= vpm_size) { vpm_cfg[1] = vpm_cfg[0]; phase = 1; } /* Try lowering Gv */ if (vpm_cfg[phase].Gv > 0) { vpm_cfg[phase].Gv--; continue; } /* Try lowering GS dispatch width */ if (vpm_cfg[phase].gs_width > 1) { do { vpm_cfg[phase].gs_width >>= 1; vpm_cfg[phase].Gd = align(vpm_cfg[phase].Gd, 2) / 2; } while (vpm_cfg[phase].gs_width == 2); /* Reset Gv to max after dropping dispatch width */ vpm_cfg[phase].Gv = 3; continue; } /* We ran out of options to reduce memory pressure. If we * are at phase 1 we have at least a valid configuration, so we * we use that. */ if (phase == 1) final_vpm_cfg = &vpm_cfg[0]; break; } while (true); if (!final_vpm_cfg) return false; assert(final_vpm_cfg); assert(final_vpm_cfg->Gd <= 16); assert(final_vpm_cfg->Gv < 4); assert(final_vpm_cfg->Ve < 4); assert(final_vpm_cfg->Vc >= 2 && final_vpm_cfg->Vc <= 4); assert(final_vpm_cfg->gs_width == 1 || final_vpm_cfg->gs_width == 4 || final_vpm_cfg->gs_width == 8 || final_vpm_cfg->gs_width == 16); *vpm_cfg_out = *final_vpm_cfg; return true; } bool v3d_compute_vpm_config(struct v3d_device_info *devinfo, struct v3d_vs_prog_data *vs_bin, struct v3d_vs_prog_data *vs, struct v3d_gs_prog_data *gs_bin, struct v3d_gs_prog_data *gs, struct vpm_config *vpm_cfg_bin, struct vpm_config *vpm_cfg) { assert(vs && vs_bin); assert((gs != NULL) == (gs_bin != NULL)); if (!gs) { vpm_cfg_bin->As = 1; vpm_cfg_bin->Ve = 0; vpm_cfg_bin->Vc = vs_bin->vcm_cache_size; vpm_cfg->As = 1; vpm_cfg->Ve = 0; vpm_cfg->Vc = vs->vcm_cache_size; } else { if (!compute_vpm_config_gs(devinfo, vs_bin, gs_bin, vpm_cfg_bin)) return false; if (!compute_vpm_config_gs(devinfo, vs, gs, vpm_cfg)) return false; } return true; }