Lines Matching full:compiler
57 ir3_compiler_destroy(struct ir3_compiler *compiler) in ir3_compiler_destroy() argument
59 disk_cache_destroy(compiler->disk_cache); in ir3_compiler_destroy()
60 ralloc_free(compiler); in ir3_compiler_destroy()
133 struct ir3_compiler *compiler = rzalloc(NULL, struct ir3_compiler); in ir3_compiler_create() local
143 compiler->dev = dev; in ir3_compiler_create()
144 compiler->dev_id = dev_id; in ir3_compiler_create()
145 compiler->gen = fd_dev_gen(dev_id); in ir3_compiler_create()
146 compiler->is_64bit = fd_dev_64b(dev_id); in ir3_compiler_create()
147 compiler->options = *options; in ir3_compiler_create()
150 compiler->branchstack_size = 64; in ir3_compiler_create()
151 compiler->wave_granularity = dev_info->wave_granularity; in ir3_compiler_create()
152 compiler->max_waves = dev_info->max_waves; in ir3_compiler_create()
154 compiler->max_variable_workgroup_size = 1024; in ir3_compiler_create()
156 compiler->local_mem_size = dev_info->cs_shared_mem_size; in ir3_compiler_create()
158 compiler->num_predicates = 1; in ir3_compiler_create()
159 compiler->bitops_can_write_predicates = false; in ir3_compiler_create()
160 compiler->has_branch_and_or = false; in ir3_compiler_create()
161 compiler->has_rpt_bary_f = false; in ir3_compiler_create()
162 compiler->has_alias_tex = false; in ir3_compiler_create()
163 compiler->delay_slots.alu_to_alu = 3; in ir3_compiler_create()
164 compiler->delay_slots.non_alu = 6; in ir3_compiler_create()
165 compiler->delay_slots.cat3_src2_read = 2; in ir3_compiler_create()
167 if (compiler->gen >= 6) { in ir3_compiler_create()
168 compiler->samgq_workaround = true; in ir3_compiler_create()
185 compiler->max_const_pipeline = 512; in ir3_compiler_create()
186 compiler->max_const_frag = 512; in ir3_compiler_create()
187 compiler->max_const_geom = 512; in ir3_compiler_create()
188 compiler->max_const_safe = 100; in ir3_compiler_create()
196 compiler->max_const_compute = in ir3_compiler_create()
197 (compiler->gen >= 7 && !dev_info->a7xx.compute_constlen_quirk) ? 512 : 256; in ir3_compiler_create()
200 compiler->has_clip_cull = true; in ir3_compiler_create()
202 compiler->has_preamble = true; in ir3_compiler_create()
204 compiler->tess_use_shared = dev_info->a6xx.tess_use_shared; in ir3_compiler_create()
206 compiler->has_getfiberid = dev_info->a6xx.has_getfiberid; in ir3_compiler_create()
208 compiler->has_dp2acc = dev_info->a6xx.has_dp2acc; in ir3_compiler_create()
209 compiler->has_dp4acc = dev_info->a6xx.has_dp4acc; in ir3_compiler_create()
210 compiler->has_compliant_dp4acc = dev_info->a7xx.has_compliant_dp4acc; in ir3_compiler_create()
212 if (compiler->gen == 6 && options->shared_push_consts) { in ir3_compiler_create()
213 compiler->shared_consts_base_offset = 504; in ir3_compiler_create()
214 compiler->shared_consts_size = 8; in ir3_compiler_create()
215 compiler->geom_shared_consts_size_quirk = 16; in ir3_compiler_create()
217 compiler->shared_consts_base_offset = -1; in ir3_compiler_create()
218 compiler->shared_consts_size = 0; in ir3_compiler_create()
219 compiler->geom_shared_consts_size_quirk = 0; in ir3_compiler_create()
222 compiler->has_fs_tex_prefetch = dev_info->a6xx.has_fs_tex_prefetch; in ir3_compiler_create()
223 compiler->stsc_duplication_quirk = dev_info->a7xx.stsc_duplication_quirk; in ir3_compiler_create()
224 compiler->load_shader_consts_via_preamble = dev_info->a7xx.load_shader_consts_via_preamble; in ir3_compiler_create()
225 …compiler->load_inline_uniforms_via_preamble_ldgk = dev_info->a7xx.load_inline_uniforms_via_preambl… in ir3_compiler_create()
226 compiler->num_predicates = 4; in ir3_compiler_create()
227 compiler->bitops_can_write_predicates = true; in ir3_compiler_create()
228 compiler->has_branch_and_or = true; in ir3_compiler_create()
229 compiler->has_predication = true; in ir3_compiler_create()
230 compiler->predtf_nop_quirk = dev_info->a6xx.predtf_nop_quirk; in ir3_compiler_create()
231 compiler->prede_nop_quirk = dev_info->a6xx.prede_nop_quirk; in ir3_compiler_create()
232 compiler->has_scalar_alu = dev_info->a6xx.has_scalar_alu; in ir3_compiler_create()
233 compiler->has_isam_v = dev_info->a6xx.has_isam_v; in ir3_compiler_create()
234 compiler->has_ssbo_imm_offsets = dev_info->a6xx.has_ssbo_imm_offsets; in ir3_compiler_create()
235 …compiler->fs_must_have_non_zero_constlen_quirk = dev_info->a7xx.fs_must_have_non_zero_constlen_qui… in ir3_compiler_create()
236 compiler->has_early_preamble = dev_info->a6xx.has_early_preamble; in ir3_compiler_create()
237 compiler->has_rpt_bary_f = true; in ir3_compiler_create()
238 compiler->has_shfl = true; in ir3_compiler_create()
239 compiler->reading_shading_rate_requires_smask_quirk = in ir3_compiler_create()
241 compiler->has_alias_rt = dev_info->a7xx.has_alias_rt; in ir3_compiler_create()
243 if (compiler->gen >= 7) { in ir3_compiler_create()
244 compiler->has_alias_tex = true; in ir3_compiler_create()
245 compiler->delay_slots.alu_to_alu = 2; in ir3_compiler_create()
246 compiler->delay_slots.non_alu = 5; in ir3_compiler_create()
247 compiler->delay_slots.cat3_src2_read = 1; in ir3_compiler_create()
250 compiler->max_const_pipeline = 512; in ir3_compiler_create()
251 compiler->max_const_geom = 512; in ir3_compiler_create()
252 compiler->max_const_frag = 512; in ir3_compiler_create()
253 compiler->max_const_compute = 512; in ir3_compiler_create()
258 compiler->max_const_safe = 256; in ir3_compiler_create()
260 compiler->has_scalar_alu = false; in ir3_compiler_create()
261 compiler->has_isam_v = false; in ir3_compiler_create()
262 compiler->has_ssbo_imm_offsets = false; in ir3_compiler_create()
263 compiler->has_early_preamble = false; in ir3_compiler_create()
267 compiler->pvtmem_per_fiber_align = compiler->gen >= 4 ? 512 : 128; in ir3_compiler_create()
269 compiler->has_pvtmem = compiler->gen >= 5; in ir3_compiler_create()
271 compiler->has_isam_ssbo = compiler->gen >= 6; in ir3_compiler_create()
273 if (compiler->gen >= 6) { in ir3_compiler_create()
274 compiler->reg_size_vec4 = dev_info->a6xx.reg_size_vec4; in ir3_compiler_create()
275 } else if (compiler->gen >= 4) { in ir3_compiler_create()
279 compiler->reg_size_vec4 = 48; in ir3_compiler_create()
282 compiler->reg_size_vec4 = 96; in ir3_compiler_create()
285 compiler->threadsize_base = dev_info->threadsize_base; in ir3_compiler_create()
287 if (compiler->gen >= 4) { in ir3_compiler_create()
289 compiler->flat_bypass = true; in ir3_compiler_create()
290 compiler->levels_add_one = false; in ir3_compiler_create()
291 compiler->unminify_coords = false; in ir3_compiler_create()
292 compiler->txf_ms_with_isaml = false; in ir3_compiler_create()
293 compiler->array_index_add_half = true; in ir3_compiler_create()
294 compiler->instr_align = 16; in ir3_compiler_create()
295 compiler->const_upload_unit = 4; in ir3_compiler_create()
298 compiler->flat_bypass = false; in ir3_compiler_create()
299 compiler->levels_add_one = true; in ir3_compiler_create()
300 compiler->unminify_coords = true; in ir3_compiler_create()
301 compiler->txf_ms_with_isaml = true; in ir3_compiler_create()
302 compiler->array_index_add_half = false; in ir3_compiler_create()
303 compiler->instr_align = 4; in ir3_compiler_create()
304 compiler->const_upload_unit = 8; in ir3_compiler_create()
307 compiler->bool_type = (compiler->gen >= 5) ? TYPE_U16 : TYPE_U32; in ir3_compiler_create()
308 compiler->has_shared_regfile = compiler->gen >= 5; in ir3_compiler_create()
309 compiler->has_bitwise_triops = compiler->gen >= 5; in ir3_compiler_create()
313 assert(compiler->has_preamble); in ir3_compiler_create()
315 /* Set up nir shader compiler options, using device-specific overrides of our base settings. */ in ir3_compiler_create()
316 compiler->nir_options = ir3_base_options; in ir3_compiler_create()
317 compiler->nir_options.has_iadd3 = dev_info->a6xx.has_sad; in ir3_compiler_create()
319 if (compiler->gen >= 6) { in ir3_compiler_create()
320 compiler->nir_options.force_indirect_unrolling = nir_var_all, in ir3_compiler_create()
321 compiler->nir_options.lower_device_index_to_zero = true; in ir3_compiler_create()
324 compiler->nir_options.has_udot_4x8 = in ir3_compiler_create()
325 compiler->nir_options.has_udot_4x8_sat = true; in ir3_compiler_create()
326 compiler->nir_options.has_sudot_4x8 = in ir3_compiler_create()
327 compiler->nir_options.has_sudot_4x8_sat = true; in ir3_compiler_create()
331 compiler->nir_options.has_sdot_4x8 = in ir3_compiler_create()
332 compiler->nir_options.has_sdot_4x8_sat = true; in ir3_compiler_create()
334 } else if (compiler->gen >= 3 && compiler->gen <= 5) { in ir3_compiler_create()
335 compiler->nir_options.vertex_id_zero_based = true; in ir3_compiler_create()
336 } else if (compiler->gen <= 2) { in ir3_compiler_create()
337 /* a2xx compiler doesn't handle indirect: */ in ir3_compiler_create()
338 compiler->nir_options.force_indirect_unrolling = nir_var_all; in ir3_compiler_create()
342 compiler->nir_options.lower_base_vertex = true; in ir3_compiler_create()
345 /* 16-bit ALU op generation is mostly controlled by frontend compiler options, but in ir3_compiler_create()
348 if (compiler->gen >= 5 && !(ir3_shader_debug & IR3_DBG_NOFP16)) in ir3_compiler_create()
349 compiler->nir_options.support_16bit_alu = true; in ir3_compiler_create()
351 compiler->nir_options.support_indirect_inputs = (uint8_t)BITFIELD_MASK(PIPE_SHADER_TYPES); in ir3_compiler_create()
352 compiler->nir_options.support_indirect_outputs = (uint8_t)BITFIELD_MASK(PIPE_SHADER_TYPES); in ir3_compiler_create()
355 ir3_disk_cache_init(compiler); in ir3_compiler_create()
357 return compiler; in ir3_compiler_create()
361 ir3_get_compiler_options(struct ir3_compiler *compiler) in ir3_get_compiler_options() argument
363 return &compiler->nir_options; in ir3_get_compiler_options()