| /external/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()
118 (n <= dispatch_width() && i < dispatch_width() / n)); in group()
164 dispatch_width() const in dispatch_width() function
187 assert(dispatch_width() <= 32);
191 DIV_ROUND_UP(n * type_sz(type) * dispatch_width(),
239 group() + dispatch_width() <= 16); in sample_mask_reg()
264 return emit(instruction(opcode, dispatch_width())); in emit()
273 return emit(instruction(opcode, dispatch_width(), dst)); in emit()
290 return emit(instruction(opcode, dispatch_width(), dst, in emit()
[all …]
|
| D | brw_nir_lower_cs_intrinsics.c | 29 unsigned dispatch_width; member
61 if (state->local_workgroup_size <= state->dispatch_width) in lower_cs_intrinsics_convert_block()
67 nir_imul(b, subgroup_id, nir_imm_int(b, state->dispatch_width)); in lower_cs_intrinsics_convert_block()
134 unsigned dispatch_width) in brw_nir_lower_cs_intrinsics() argument
142 state.dispatch_width = dispatch_width; in brw_nir_lower_cs_intrinsics()
|
| 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_reg_allocate.cpp | 49 int reg_width = dispatch_width / 8; in assign_regs_trivial()
76 brw_alloc_reg_set(struct brw_compiler *compiler, int dispatch_width) in brw_alloc_reg_set() argument
80 const int index = _mesa_logbase2(dispatch_width / 8); in brw_alloc_reg_set()
82 if (dispatch_width > 8 && devinfo->gen >= 7) { in brw_alloc_reg_set()
118 if (devinfo->gen <= 5 && dispatch_width >= 16) { in brw_alloc_reg_set()
165 if (devinfo->gen <= 5 && dispatch_width >= 16) { in brw_alloc_reg_set()
211 if (devinfo->gen <= 5 && dispatch_width >= 16) { in brw_alloc_reg_set()
252 if (devinfo->has_pln && dispatch_width == 8 && devinfo->gen <= 6) { in brw_alloc_reg_set()
447 if (devinfo->gen <= 5 && dispatch_width >= 16) { in setup_payload_interference()
471 int reg_width = v->dispatch_width / 8; in get_used_mrfs()
[all …]
|
| D | brw_fs_visitor.cpp | 94 int reg_width = dispatch_width / 8; in emit_dummy_fs()
177 if (devinfo->has_pln && dispatch_width == 16) { in emit_interpolation_setup_gen4()
208 if (devinfo->gen >= 8 || dispatch_width == 8) { in emit_interpolation_setup_gen6()
219 fs_reg int_pixel_xy(VGRF, alloc.allocate(dispatch_width / 8), in emit_interpolation_setup_gen6()
222 const fs_builder dbld = abld.exec_all().group(dispatch_width * 2, 0); in emit_interpolation_setup_gen6()
291 if (dispatch_width == 16) { in emit_interpolation_setup_gen6()
299 assert(dispatch_width != 32); /* not implemented yet */ in emit_interpolation_setup_gen6()
818 unsigned dispatch_width, in fs_visitor() argument
824 dispatch_width(dispatch_width), in fs_visitor()
826 bld(fs_builder(this, dispatch_width).at_end()) in fs_visitor()
[all …]
|
| D | brw_fs.h | 50 return offset(reg, bld.dispatch_width(), delta); in offset()
69 unsigned dispatch_width,
370 const unsigned dispatch_width; /**< 8, 16 or 32 */ variable
397 int generate_code(const cfg_t *cfg, int dispatch_width);
483 unsigned dispatch_width; /**< 8, 16 or 32 */ variable
|
| D | brw_fs.cpp | 598 if (dispatch_width == 8) in SHADER_TIME_ADD()
650 if (dispatch_width > n) { in limit_dispatch_width()
966 return 1 * dispatch_width / 8; in implied_mrf_writes()
970 return 2 * dispatch_width / 8; in implied_mrf_writes()
1001 int reg_width = dispatch_width / 8; in vgrf()
1193 if (dispatch_width == 8) { in emit_samplepos_setup()
1203 if (dispatch_width == 8) { in emit_samplepos_setup()
3119 if (dispatch_width >= 16) in remove_duplicate_mrf_writes()
3571 fs_reg imm(VGRF, alloc.allocate(dispatch_width / 8), in lower_integer_multiplication()
3902 assert(bld.dispatch_width() != 16); in lower_fb_write_logical_send()
[all …]
|
| D | brw_nir.h | 103 unsigned dispatch_width);
|
| D | brw_fs_nir.cpp | 244 if (dispatch_width > 8) in nir_emit_system_values()
246 if (dispatch_width > 16) { in nir_emit_system_values()
1393 shift_count = fs_reg(VGRF, alloc.allocate(dispatch_width / 4), in nir_emit_alu()
1743 mlen = 2 * bld.dispatch_width() / 8; in emit_pixel_interpolater_send()
4323 if (dispatch_width == 32) { in nir_emit_intrinsic()
4340 set_predicate(dispatch_width == 8 ? BRW_PREDICATE_ALIGN1_ANY8H : in nir_emit_intrinsic()
4341 dispatch_width == 16 ? BRW_PREDICATE_ALIGN1_ANY16H : in nir_emit_intrinsic()
4355 if (dispatch_width == 32) { in nir_emit_intrinsic()
4372 set_predicate(dispatch_width == 8 ? BRW_PREDICATE_ALIGN1_ALL8H : in nir_emit_intrinsic()
4373 dispatch_width == 16 ? BRW_PREDICATE_ALIGN1_ALL16H : in nir_emit_intrinsic()
[all …]
|
| D | brw_fs_generator.cpp | 293 (prog_data->dual_src_blend && dispatch_width == 16); in fire_fb_write()
1625 fs_generator::generate_code(const cfg_t *cfg, int dispatch_width) in generate_code() argument
1631 this->dispatch_width = dispatch_width; in generate_code()
2226 shader_name, dispatch_width, before_size / 16, loop_count, cfg->cycle_count, in generate_code()
2240 dispatch_width, before_size / 16, in generate_code()
|