| /third_party/mesa3d/src/intel/compiler/ |
| D | brw_fs_builder.h | 56 unsigned dispatch_width) : in fs_builder() argument
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
201 assert(dispatch_width() <= 32);
205 DIV_ROUND_UP(n * type_sz(type) * dispatch_width(),
260 return emit(instruction(opcode, dispatch_width())); in emit()
269 return emit(instruction(opcode, dispatch_width(), dst)); in emit()
286 return emit(instruction(opcode, dispatch_width(), dst, in emit()
290 return emit(instruction(opcode, dispatch_width(), dst, src0)); in emit()
[all …]
|
| D | brw_vec4_builder.h | 54 vec4_builder(backend_shader *shader, unsigned dispatch_width = 8) :
56 _dispatch_width(dispatch_width), _group(0), in shader()
114 (n <= dispatch_width() && i < dispatch_width() / n)); in group()
151 dispatch_width() const in dispatch_width() function
174 assert(dispatch_width() <= 32);
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()
313 inst->exec_size = dispatch_width(); in emit()
|
| D | brw_fs.h | 78 return offset(reg, bld.dispatch_width(), delta); in offset()
101 unsigned dispatch_width,
427 const unsigned dispatch_width; /**< 8, 16 or 32 */ variable
478 int generate_code(const cfg_t *cfg, int dispatch_width,
574 unsigned dispatch_width; /**< 8, 16 or 32 */ variable
592 if (bld.dispatch_width() > 16) {
595 const unsigned m = bld.dispatch_width() / hbld.dispatch_width();
619 const unsigned m = bld.dispatch_width() / hbld.dispatch_width(); in fetch_barycentric_reg()
660 bool brw_nir_lower_simd(nir_shader *nir, unsigned dispatch_width);
|
| D | brw_mesh.cpp | 230 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() argument
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, true /* is_scalar */); in brw_compile_mesh()
[all …]
|
| D | brw_fs_visitor.cpp | 67 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 / 8; q++) { in emit_interpolation_setup_gfx6()
[all …]
|
| D | brw_fs_reg_allocate.cpp | 50 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()
315 int reg_width = fs->dispatch_width / 8; in fs_reg_alloc()
[all …]
|
| D | brw_fs.cpp | 591 dispatch_width, stage_abbrev, msg); in vfail()
624 if (dispatch_width > n) { in limit_dispatch_width()
1118 int reg_width = dispatch_width / 8; in vgrf()
1363 for (unsigned i = 0; i < DIV_ROUND_UP(dispatch_width, 16); i++) { in emit_sampleid_setup()
1364 const fs_builder hbld = abld.group(MIN2(16, dispatch_width), i); in emit_sampleid_setup()
3421 if (dispatch_width >= 16) in remove_duplicate_mrf_writes()
3938 fs_reg imm(VGRF, alloc.allocate(dispatch_width / 8), inst->dst.type); in lower_mul_dword_inst()
4383 assert(bld.dispatch_width() <= 16); in brw_sample_mask_reg()
4386 assert(v->devinfo->ver >= 6 && bld.dispatch_width() <= 16); in brw_sample_mask_reg()
4432 bld.dispatch_width() == inst->exec_size); in brw_emit_predicate_on_sample_mask()
[all …]
|
| D | brw_fs_nir.cpp | 46 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) { in nir_emit_cs_intrinsic()
3878 inst->size_written = 3 * dispatch_width * 4; in nir_emit_cs_intrinsic()
3924 inst->size_written = instr->num_components * dispatch_width * 4; in nir_emit_cs_intrinsic()
4182 const unsigned chan_index_bits = ffs(dispatch_width) - 1; in swizzle_nir_scratch_addr()
4362 inst->size_written = instr->num_components * dispatch_width * 4; in nir_emit_intrinsic()
[all …]
|
| D | brw_compiler.h | 1085 unsigned dispatch_width) in brw_cs_prog_data_prog_offset() argument
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
|
| D | brw_lower_logical_sends.cpp | 281 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 …]
|
| D | brw_fs_generator.cpp | 195 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 = dispatch_width; in generate_code()
|
| D | brw_wm_iz.cpp | 125 assert(dispatch_width <= 16); in setup_fs_payload_gfx4()
|
| D | brw_ir_performance.cpp | 1576 unsigned dispatch_width) in calculate_performance() argument
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()
|
| D | brw_vec4_generator.cpp | 2259 stats->dispatch_width = 0; in generate_code()
|
| /third_party/mesa3d/src/intel/vulkan/ |
| D | anv_pipeline.c | 3007 unsigned simd_width = exe->stats.dispatch_width; in anv_GetPipelineExecutablePropertiesKHR()
|