Home
last modified time | relevance | path

Searched refs:dispatch_width (Results 1 - 15 of 15) sorted by relevance

/third_party/mesa3d/src/intel/compiler/
H A Dbrw_fs_builder.h53 * \p dispatch_width gives the native execution width of the program.
56 unsigned dispatch_width) : in fs_builder()
58 _dispatch_width(dispatch_width), in fs_builder()
119 if (n <= dispatch_width() && i < dispatch_width() / n) { in group()
178 dispatch_width() const in dispatch_width() function in brw::fs_builder
195 * dispatch_width units (which is just enough space for one logical
201 assert(dispatch_width() <= 32); in vgrf()
205 DIV_ROUND_UP(n * type_sz(type) * dispatch_width(), in vgrf()
260 return emit(instruction(opcode, dispatch_width())); in emit()
55 fs_builder(backend_shader *shader, unsigned dispatch_width) fs_builder() argument
[all...]
H A Dbrw_vec4_builder.h54 vec4_builder(backend_shader *shader, unsigned dispatch_width = 8) : in vec4_builder()
56 _dispatch_width(dispatch_width), _group(0), in vec4_builder()
114 (n <= dispatch_width() && i < dispatch_width() / n)); in group()
151 dispatch_width() const in dispatch_width() function in brw::vec4_builder
168 * dispatch_width units (which is just enough space for four logical
174 assert(dispatch_width() <= 32); in vgrf()
190 return dst_reg(retype(brw_null_vec(dispatch_width()), in null_reg_f()
200 return dst_reg(retype(brw_null_vec(dispatch_width()), in null_reg_d()
210 return dst_reg(retype(brw_null_vec(dispatch_width()), in null_reg_ud()
[all...]
H A Dbrw_fs.h78 return offset(reg, bld.dispatch_width(), delta); in offset()
101 unsigned dispatch_width,
427 const unsigned dispatch_width; /**< 8, 16 or 32 */ member in fs_visitor
478 int generate_code(const cfg_t *cfg, int dispatch_width,
574 unsigned dispatch_width; /**< 8, 16 or 32 */ member in fs_generator
592 if (bld.dispatch_width() > 16) { in fetch_payload_reg()
595 const unsigned m = bld.dispatch_width() / hbld.dispatch_width(); in fetch_payload_reg()
619 const unsigned m = bld.dispatch_width() / hbld.dispatch_width(); in fetch_barycentric_reg()
[all...]
H A Dbrw_mesh.cpp230 const unsigned dispatch_width = 8 << simd; in brw_compile_task() local
233 brw_nir_apply_key(shader, compiler, &key->base, dispatch_width, true /* is_scalar */); in brw_compile_task()
236 NIR_PASS(_, shader, brw_nir_lower_simd, dispatch_width); in brw_compile_task()
244 &prog_data->base.base, shader, dispatch_width, in brw_compile_task()
285 g.generate_code(selected->cfg, selected->dispatch_width, selected->shader_stats, in brw_compile_task()
538 unsigned dispatch_width) in brw_nir_initialize_mue()
609 if (workgroup_size > dispatch_width) { in brw_nir_initialize_mue()
731 const unsigned dispatch_width = 8 << simd; in brw_compile_mesh() local
740 NIR_PASS_V(shader, brw_nir_initialize_mue, &prog_data->map, dispatch_width); in brw_compile_mesh()
742 brw_nir_apply_key(shader, compiler, &key->base, dispatch_width, tru in brw_compile_mesh()
536 brw_nir_initialize_mue(nir_shader *nir, const struct brw_mue_map *map, unsigned dispatch_width) brw_nir_initialize_mue() argument
[all...]
H A Dbrw_fs_visitor.cpp67 int reg_width = dispatch_width / 8; in emit_dummy_fs()
186 for (unsigned i = 0; i < dispatch_width / 8; i++) { in emit_interpolation_setup_gfx4()
350 abld.exec_all().group(MIN2(16, dispatch_width) * 2, 0); in emit_interpolation_setup_gfx6()
386 for (unsigned i = 0; i < DIV_ROUND_UP(dispatch_width, 16); i++) { in emit_interpolation_setup_gfx6()
387 const fs_builder hbld = abld.group(MIN2(16, dispatch_width), i); in emit_interpolation_setup_gfx6()
392 abld.exec_all().group(hbld.dispatch_width() * 2, 0); in emit_interpolation_setup_gfx6()
413 } else if (devinfo->ver >= 8 || dispatch_width == 8) { in emit_interpolation_setup_gfx6()
425 abld.exec_all().group(hbld.dispatch_width() * 2, 0); in emit_interpolation_setup_gfx6()
534 for (unsigned i = 0; i < DIV_ROUND_UP(dispatch_width, 16); i++) { in emit_interpolation_setup_gfx6()
550 for (unsigned q = 0; q < dispatch_width / in emit_interpolation_setup_gfx6()
1133 fs_visitor(const struct brw_compiler *compiler, void *log_data, void *mem_ctx, const brw_base_prog_key *key, struct brw_stage_prog_data *prog_data, const nir_shader *shader, unsigned dispatch_width, bool debug_enabled) fs_visitor() argument
[all...]
H A Dbrw_fs_reg_allocate.cpp50 int reg_width = dispatch_width / 8; in assign_regs_trivial()
80 aligned_bary_size(unsigned dispatch_width) in aligned_bary_size() argument
82 return (dispatch_width == 8 ? 2 : 4); in aligned_bary_size()
86 brw_alloc_reg_set(struct brw_compiler *compiler, int dispatch_width) in brw_alloc_reg_set() argument
90 const int index = util_logbase2(dispatch_width / 8); in brw_alloc_reg_set()
92 if (dispatch_width > 8 && devinfo->ver >= 7) { in brw_alloc_reg_set()
133 if (devinfo->ver <= 5 && dispatch_width >= 16) { in brw_alloc_reg_set()
156 (dispatch_width == 8 && devinfo->ver <= 5))) { in brw_alloc_reg_set()
157 int contig_len = aligned_bary_size(dispatch_width); in brw_alloc_reg_set()
310 * (dispatch_width in fs_reg_alloc()
[all...]
H A Dbrw_compiler.h1085 unsigned dispatch_width) in brw_cs_prog_data_prog_offset()
1087 assert(dispatch_width == 8 || in brw_cs_prog_data_prog_offset()
1088 dispatch_width == 16 || in brw_cs_prog_data_prog_offset()
1089 dispatch_width == 32); in brw_cs_prog_data_prog_offset()
1090 const unsigned index = dispatch_width / 16; in brw_cs_prog_data_prog_offset()
1539 uint32_t dispatch_width; /**< 0 for vec4 */ member
1084 brw_cs_prog_data_prog_offset(const struct brw_cs_prog_data *prog_data, unsigned dispatch_width) brw_cs_prog_data_prog_offset() argument
H A Dbrw_fs.cpp591 dispatch_width, stage_abbrev, msg); in vfail()
624 if (dispatch_width > n) { in limit_dispatch_width()
1118 int reg_width = dispatch_width / 8;
1363 for (unsigned i = 0; i < DIV_ROUND_UP(dispatch_width, 16); i++) {
1364 const fs_builder hbld = abld.group(MIN2(16, dispatch_width), i);
3421 if (dispatch_width >= 16)
3938 fs_reg imm(VGRF, alloc.allocate(dispatch_width / 8), inst->dst.type);
4372 * thread payload, \p bld is required to have a dispatch_width() not greater
4383 assert(bld.dispatch_width() <= 16);
4386 assert(v->devinfo->ver >= 6 && bld.dispatch_width() <
[all...]
H A Dbrw_fs_nir.cpp46 last_scratch = ALIGN(nir->scratch_size, 4) * dispatch_width; in emit_nir_code()
230 for (unsigned i = 0; i < DIV_ROUND_UP(v->dispatch_width, 16); i++) { in emit_system_values_block()
231 const fs_builder hbld = abld.group(MIN2(16, v->dispatch_width), i); in emit_system_values_block()
297 if (dispatch_width > 8) in nir_emit_system_values()
299 if (dispatch_width > 16) { in nir_emit_system_values()
3835 workgroup_size() <= dispatch_width) {
3878 inst->size_written = 3 * dispatch_width * 4;
3924 inst->size_written = instr->num_components * dispatch_width * 4;
4182 const unsigned chan_index_bits = ffs(dispatch_width) - 1;
4362 inst->size_written = instr->num_components * dispatch_width *
[all...]
H A Dbrw_lower_logical_sends.cpp281 for (unsigned i = 0; i < bld.dispatch_width() / 8; i++) { in lower_fb_write_logical_send()
334 assert(bld.dispatch_width() == 8); in lower_fb_write_logical_send()
395 if (devinfo->ver < 6 && bld.dispatch_width() == 16) in lower_fb_write_logical_send()
495 (op == SHADER_OPCODE_TEX && bld.dispatch_width() == 8))) { in lower_sampler_logical_send_gfx4()
505 assert(bld.dispatch_width() == 8); in lower_sampler_logical_send_gfx4()
540 assert(shadow_c.file != BAD_FILE ? bld.dispatch_width() == 8 : in lower_sampler_logical_send_gfx4()
541 bld.dispatch_width() == 16); in lower_sampler_logical_send_gfx4()
551 if (op == SHADER_OPCODE_TEX && bld.dispatch_width() == 8) { in lower_sampler_logical_send_gfx4()
751 sources * DIV_ROUND_UP(requested_alignment_sz, bld.dispatch_width()); in emit_load_payload_with_padding()
759 retype(dst, src[i].type).component_size(bld.dispatch_width()); in emit_load_payload_with_padding()
[all...]
H A Dbrw_wm_iz.cpp125 assert(dispatch_width <= 16); in setup_fs_payload_gfx4()
H A Dbrw_fs_generator.cpp195 prog_data(prog_data), dispatch_width(0), in fs_generator()
493 inst->exec_size == dispatch_width; in generate_mov_indirect()
679 lower_width == dispatch_width; in generate_shuffle()
1004 assert(!inst->eot || inst->exec_size == dispatch_width); in generate_tex()
1764 fs_generator::generate_code(const cfg_t *cfg, int dispatch_width, in generate_code() argument
1772 this->dispatch_width = dispatch_width; in generate_code()
2575 dispatch_width, before_size / 16, in generate_code()
2609 dispatch_width, before_size / 16 - nop_count, in generate_code()
2618 stats->dispatch_width in generate_code()
[all...]
H A Dbrw_ir_performance.cpp1576 unsigned dispatch_width) in calculate_performance()
1603 const float discard_weight = (dispatch_width > 16 || s->devinfo->ver < 12 ? in calculate_performance()
1635 p.throughput = dispatch_width * calculate_thread_throughput(st, elapsed); in calculate_performance()
1642 calculate_performance(*this, v, issue_fs_inst, v->dispatch_width); in performance()
1572 calculate_performance(performance &p, const backend_shader *s, void (*issue_instruction)( state &, const struct brw_isa_info *, const backend_instruction *), unsigned dispatch_width) calculate_performance() argument
H A Dbrw_vec4_generator.cpp2259 stats->dispatch_width = 0; in generate_code()
/third_party/mesa3d/src/intel/vulkan/
H A Danv_pipeline.c3007 unsigned simd_width = exe->stats.dispatch_width; in anv_GetPipelineExecutablePropertiesKHR()

Completed in 32 milliseconds