Lines Matching +full:a630 +full:- +full:test
34 /* clang-format off */
47 {"spillall", IR3_DBG_SPILLALL, "Spill as much as possible to test the spiller"},
50 /* DEBUG-only options: */
52 {"ramsgs", IR3_DBG_RAMSGS, "Enable register-allocation debug messages"},
55 /* clang-format on */
69 disk_cache_destroy(compiler->disk_cache); in ir3_compiler_destroy()
165 compiler->dev = dev; in ir3_compiler_create()
166 compiler->dev_id = dev_id; in ir3_compiler_create()
167 compiler->gen = fd_dev_gen(dev_id); in ir3_compiler_create()
168 compiler->robust_buffer_access2 = options->robust_buffer_access2; in ir3_compiler_create()
171 compiler->local_mem_size = 32 * 1024; in ir3_compiler_create()
173 compiler->branchstack_size = 64; in ir3_compiler_create()
174 compiler->wave_granularity = 2; in ir3_compiler_create()
175 compiler->max_waves = 16; in ir3_compiler_create()
177 compiler->max_variable_workgroup_size = 1024; in ir3_compiler_create()
179 const struct fd_dev_info *dev_info = fd_dev_info(compiler->dev_id); in ir3_compiler_create()
181 if (compiler->gen >= 6) { in ir3_compiler_create()
182 compiler->samgq_workaround = true; in ir3_compiler_create()
189 * Also, according to the observation on a630/a650/a660, max_const_pipeline in ir3_compiler_create()
199 compiler->max_const_pipeline = 512; in ir3_compiler_create()
200 compiler->max_const_frag = 512; in ir3_compiler_create()
201 compiler->max_const_geom = 512; in ir3_compiler_create()
202 compiler->max_const_safe = 100; in ir3_compiler_create()
209 compiler->max_const_compute = 256; in ir3_compiler_create()
212 compiler->has_clip_cull = true; in ir3_compiler_create()
215 compiler->has_pvtmem = true; in ir3_compiler_create()
217 compiler->has_preamble = true; in ir3_compiler_create()
219 compiler->tess_use_shared = dev_info->a6xx.tess_use_shared; in ir3_compiler_create()
221 compiler->storage_16bit = dev_info->a6xx.storage_16bit; in ir3_compiler_create()
223 compiler->has_getfiberid = dev_info->a6xx.has_getfiberid; in ir3_compiler_create()
225 compiler->has_dp2acc = dev_info->a6xx.has_dp2acc; in ir3_compiler_create()
226 compiler->has_dp4acc = dev_info->a6xx.has_dp4acc; in ir3_compiler_create()
228 compiler->shared_consts_base_offset = 504; in ir3_compiler_create()
229 compiler->shared_consts_size = 8; in ir3_compiler_create()
230 compiler->geom_shared_consts_size_quirk = 16; in ir3_compiler_create()
232 compiler->max_const_pipeline = 512; in ir3_compiler_create()
233 compiler->max_const_geom = 512; in ir3_compiler_create()
234 compiler->max_const_frag = 512; in ir3_compiler_create()
235 compiler->max_const_compute = 512; in ir3_compiler_create()
240 compiler->max_const_safe = 256; in ir3_compiler_create()
243 if (compiler->gen >= 6) { in ir3_compiler_create()
244 compiler->reg_size_vec4 = dev_info->a6xx.reg_size_vec4; in ir3_compiler_create()
245 } else if (compiler->gen >= 4) { in ir3_compiler_create()
246 /* On a4xx-a5xx, using r24.x and above requires using the smallest in ir3_compiler_create()
249 compiler->reg_size_vec4 = 48; in ir3_compiler_create()
252 compiler->reg_size_vec4 = 96; in ir3_compiler_create()
255 if (compiler->gen >= 6) { in ir3_compiler_create()
256 compiler->threadsize_base = 64; in ir3_compiler_create()
257 } else if (compiler->gen >= 4) { in ir3_compiler_create()
261 compiler->threadsize_base = 32; in ir3_compiler_create()
263 compiler->threadsize_base = 8; in ir3_compiler_create()
266 if (compiler->gen >= 4) { in ir3_compiler_create()
268 compiler->flat_bypass = true; in ir3_compiler_create()
269 compiler->levels_add_one = false; in ir3_compiler_create()
270 compiler->unminify_coords = false; in ir3_compiler_create()
271 compiler->txf_ms_with_isaml = false; in ir3_compiler_create()
272 compiler->array_index_add_half = true; in ir3_compiler_create()
273 compiler->instr_align = 16; in ir3_compiler_create()
274 compiler->const_upload_unit = 4; in ir3_compiler_create()
277 compiler->flat_bypass = false; in ir3_compiler_create()
278 compiler->levels_add_one = true; in ir3_compiler_create()
279 compiler->unminify_coords = true; in ir3_compiler_create()
280 compiler->txf_ms_with_isaml = true; in ir3_compiler_create()
281 compiler->array_index_add_half = false; in ir3_compiler_create()
282 compiler->instr_align = 4; in ir3_compiler_create()
283 compiler->const_upload_unit = 8; in ir3_compiler_create()
286 compiler->bool_type = (compiler->gen >= 5) ? TYPE_U16 : TYPE_U32; in ir3_compiler_create()
287 compiler->has_shared_regfile = compiler->gen >= 5; in ir3_compiler_create()
289 compiler->push_ubo_with_preamble = options->push_ubo_with_preamble; in ir3_compiler_create()
292 if (options->push_ubo_with_preamble) in ir3_compiler_create()
293 assert(compiler->has_preamble); in ir3_compiler_create()
295 if (compiler->gen >= 6) { in ir3_compiler_create()
296 compiler->nir_options = nir_options_a6xx; in ir3_compiler_create()
297 compiler->nir_options.has_udot_4x8 = dev_info->a6xx.has_dp2acc; in ir3_compiler_create()
298 compiler->nir_options.has_sudot_4x8 = dev_info->a6xx.has_dp2acc; in ir3_compiler_create()
300 compiler->nir_options = nir_options; in ir3_compiler_create()
302 if (compiler->gen <= 2) in ir3_compiler_create()
303 compiler->nir_options.force_indirect_unrolling = nir_var_all; in ir3_compiler_create()
306 /* 16-bit ALU op generation is mostly controlled by frontend compiler options, but in ir3_compiler_create()
307 * this core NIR option enables some optimizations of 16-bit operations. in ir3_compiler_create()
309 if (compiler->gen >= 5 && !(ir3_shader_debug & IR3_DBG_NOFP16)) in ir3_compiler_create()
310 compiler->nir_options.support_16bit_alu = true; in ir3_compiler_create()
312 if (!options->disable_cache) in ir3_compiler_create()
321 return &compiler->nir_options; in ir3_get_compiler_options()