Home
last modified time | relevance | path

Searched refs:cs_prog_data (Results 1 – 25 of 27) sorted by relevance

12

/aosp_15_r20/external/mesa3d/src/intel/compiler/
H A Dbrw_simd_selection.cpp87 const auto cs_prog_data = get_cs_prog_data(state); in brw_simd_should_compile() local
95 const bool workgroup_size_variable = cs_prog_data && cs_prog_data->local_size[0] == 0; in brw_simd_should_compile()
108 if (cs_prog_data) { in brw_simd_should_compile()
109 const unsigned workgroup_size = cs_prog_data->local_size[0] * in brw_simd_should_compile()
110 cs_prog_data->local_size[1] * in brw_simd_should_compile()
111 cs_prog_data->local_size[2]; in brw_simd_should_compile()
145 if (width == 32 && cs_prog_data && cs_prog_data->base.ray_queries > 0) { in brw_simd_should_compile()
150 if (width == 32 && cs_prog_data && cs_prog_data->uses_btd_stack_ids) { in brw_simd_should_compile()
200 auto cs_prog_data = get_cs_prog_data(state); in brw_simd_mark_compiled() local
203 if (cs_prog_data) in brw_simd_mark_compiled()
[all …]
H A Dintel_clc.c148 const struct brw_cs_prog_data *cs_prog_data) in print_cs_prog_data_fields() argument
151 fprintf(fp, "%s." #field " = " fmt ",\n", pad, cs_prog_data->field) in print_cs_prog_data_fields()
155 cs_prog_data->field ? "true" : "false") in print_cs_prog_data_fields()
158 assert(cs_prog_data->base.stage == MESA_SHADER_COMPUTE); in print_cs_prog_data_fields()
160 assert(cs_prog_data->base.zero_push_reg == 0); in print_cs_prog_data_fields()
161 assert(cs_prog_data->base.push_reg_mask_param == 0); in print_cs_prog_data_fields()
172 assert(!cs_prog_data->base.has_ubo_pull); in print_cs_prog_data_fields()
173 assert(cs_prog_data->base.dispatch_grf_start_reg == 0); in print_cs_prog_data_fields()
174 assert(!cs_prog_data->base.use_alt_mode); in print_cs_prog_data_fields()
175 assert(cs_prog_data->base.param == 0); in print_cs_prog_data_fields()
[all …]
H A Dbrw_compile_cs.cpp31 struct brw_cs_prog_data *cs_prog_data) in cs_fill_push_const_info() argument
33 const struct brw_stage_prog_data *prog_data = &cs_prog_data->base; in cs_fill_push_const_info()
52 fill_push_const_block_info(&cs_prog_data->push.cross_thread, cross_thread_dwords); in cs_fill_push_const_info()
53 fill_push_const_block_info(&cs_prog_data->push.per_thread, per_thread_dwords); in cs_fill_push_const_info()
55 assert(cs_prog_data->push.cross_thread.dwords % 8 == 0 || in cs_fill_push_const_info()
56 cs_prog_data->push.per_thread.size == 0); in cs_fill_push_const_info()
57 assert(cs_prog_data->push.cross_thread.dwords + in cs_fill_push_const_info()
58 cs_prog_data->push.per_thread.dwords == in cs_fill_push_const_info()
H A Dbrw_fs.cpp1793 brw_cs_push_const_total_size(const struct brw_cs_prog_data *cs_prog_data, in brw_cs_push_const_total_size() argument
1796 assert(cs_prog_data->push.per_thread.size % REG_SIZE == 0); in brw_cs_push_const_total_size()
1797 assert(cs_prog_data->push.cross_thread.size % REG_SIZE == 0); in brw_cs_push_const_total_size()
1798 return cs_prog_data->push.per_thread.size * threads + in brw_cs_push_const_total_size()
1799 cs_prog_data->push.cross_thread.size; in brw_cs_push_const_total_size()
H A Dbrw_compiler.h1556 brw_cs_push_const_total_size(const struct brw_cs_prog_data *cs_prog_data,
H A Dbrw_fs_nir.cpp4491 struct brw_cs_prog_data *cs_prog_data = brw_cs_prog_data(s.prog_data); in fs_nir_emit_cs_intrinsic() local
4513 cs_prog_data->uses_barrier = true; in fs_nir_emit_cs_intrinsic()
4523 assert(cs_prog_data->generate_local_id); in fs_nir_emit_cs_intrinsic()
4543 cs_prog_data->uses_num_work_groups = true; in fs_nir_emit_cs_intrinsic()
4595 cs_prog_data->uses_systolic = true; in fs_nir_emit_cs_intrinsic()
/aosp_15_r20/external/mesa3d/src/intel/compiler/elk/
H A Delk_simd_selection.cpp77 const auto cs_prog_data = get_cs_prog_data(state); in elk_simd_should_compile() local
85 const bool workgroup_size_variable = cs_prog_data && cs_prog_data->local_size[0] == 0; in elk_simd_should_compile()
98 if (cs_prog_data) { in elk_simd_should_compile()
99 const unsigned workgroup_size = cs_prog_data->local_size[0] * in elk_simd_should_compile()
100 cs_prog_data->local_size[1] * in elk_simd_should_compile()
101 cs_prog_data->local_size[2]; in elk_simd_should_compile()
160 auto cs_prog_data = get_cs_prog_data(state); in elk_simd_mark_compiled() local
163 if (cs_prog_data) in elk_simd_mark_compiled()
164 cs_prog_data->prog_mask |= 1u << simd; in elk_simd_mark_compiled()
170 if (cs_prog_data) in elk_simd_mark_compiled()
[all …]
H A Delk_fs.cpp6811 elk_cs_push_const_total_size(const struct elk_cs_prog_data *cs_prog_data, in elk_cs_push_const_total_size() argument
6814 assert(cs_prog_data->push.per_thread.size % REG_SIZE == 0); in elk_cs_push_const_total_size()
6815 assert(cs_prog_data->push.cross_thread.size % REG_SIZE == 0); in elk_cs_push_const_total_size()
6816 return cs_prog_data->push.per_thread.size * threads + in elk_cs_push_const_total_size()
6817 cs_prog_data->push.cross_thread.size; in elk_cs_push_const_total_size()
6830 struct elk_cs_prog_data *cs_prog_data) in cs_fill_push_const_info() argument
6832 const struct elk_stage_prog_data *prog_data = &cs_prog_data->base; in cs_fill_push_const_info()
6855 fill_push_const_block_info(&cs_prog_data->push.cross_thread, cross_thread_dwords); in cs_fill_push_const_info()
6856 fill_push_const_block_info(&cs_prog_data->push.per_thread, per_thread_dwords); in cs_fill_push_const_info()
6858 assert(cs_prog_data->push.cross_thread.dwords % 8 == 0 || in cs_fill_push_const_info()
[all …]
H A Delk_compiler.h1658 elk_cs_push_const_total_size(const struct elk_cs_prog_data *cs_prog_data,
H A Delk_fs_nir.cpp4004 struct elk_cs_prog_data *cs_prog_data = elk_cs_prog_data(s.prog_data); in fs_nir_emit_cs_intrinsic() local
4026 cs_prog_data->uses_barrier = true; in fs_nir_emit_cs_intrinsic()
4046 cs_prog_data->uses_num_work_groups = true; in fs_nir_emit_cs_intrinsic()
/aosp_15_r20/external/mesa3d/src/intel/blorp/
H A Dblorp_genX_exec_brw.h1579 const struct brw_cs_prog_data *cs_prog_data = params->cs_prog_data; in blorp_get_compute_push_const() local
1581 ALIGN(brw_cs_push_const_total_size(cs_prog_data, threads), 64); in blorp_get_compute_push_const()
1582 assert(cs_prog_data->push.cross_thread.size + in blorp_get_compute_push_const()
1583 cs_prog_data->push.per_thread.size == sizeof(params->wm_inputs)); in blorp_get_compute_push_const()
1608 if (cs_prog_data->push.cross_thread.size > 0) { in blorp_get_compute_push_const()
1609 memcpy(dst, src, cs_prog_data->push.cross_thread.size); in blorp_get_compute_push_const()
1610 dst += cs_prog_data->push.cross_thread.size; in blorp_get_compute_push_const()
1611 src += cs_prog_data->push.cross_thread.size; in blorp_get_compute_push_const()
1614 assert(GFX_VERx10 < 125 || cs_prog_data->push.per_thread.size == 0); in blorp_get_compute_push_const()
1616 if (cs_prog_data->push.per_thread.size > 0) { in blorp_get_compute_push_const()
[all …]
H A Dblorp_brw.c136 struct brw_cs_prog_data *cs_prog_data = rzalloc(mem_ctx, struct brw_cs_prog_data); in blorp_compile_cs_brw() local
137 cs_prog_data->base.nr_params = nr_params; in blorp_compile_cs_brw()
138 cs_prog_data->base.param = rzalloc_array(NULL, uint32_t, nr_params); in blorp_compile_cs_brw()
141 cs_prog_data); in blorp_compile_cs_brw()
156 .prog_data = cs_prog_data, in blorp_compile_cs_brw()
161 ralloc_free(cs_prog_data->base.param); in blorp_compile_cs_brw()
162 cs_prog_data->base.param = NULL; in blorp_compile_cs_brw()
166 .kernel_size = cs_prog_data->base.program_size, in blorp_compile_cs_brw()
167 .prog_data = cs_prog_data, in blorp_compile_cs_brw()
168 .prog_data_size = sizeof(*cs_prog_data), in blorp_compile_cs_brw()
H A Dblorp_elk.c144 struct elk_cs_prog_data *cs_prog_data = rzalloc(mem_ctx, struct elk_cs_prog_data); in blorp_compile_cs_elk() local
145 cs_prog_data->base.nr_params = nr_params; in blorp_compile_cs_elk()
146 cs_prog_data->base.param = rzalloc_array(NULL, uint32_t, nr_params); in blorp_compile_cs_elk()
149 cs_prog_data); in blorp_compile_cs_elk()
164 .prog_data = cs_prog_data, in blorp_compile_cs_elk()
169 ralloc_free(cs_prog_data->base.param); in blorp_compile_cs_elk()
170 cs_prog_data->base.param = NULL; in blorp_compile_cs_elk()
174 .kernel_size = cs_prog_data->base.program_size, in blorp_compile_cs_elk()
175 .prog_data = cs_prog_data, in blorp_compile_cs_elk()
176 .prog_data_size = sizeof(*cs_prog_data), in blorp_compile_cs_elk()
H A Dblorp_genX_exec_elk.h1887 const struct elk_cs_prog_data *cs_prog_data = params->cs_prog_data; in blorp_get_compute_push_const() local
1889 ALIGN(elk_cs_push_const_total_size(cs_prog_data, threads), 64); in blorp_get_compute_push_const()
1890 assert(cs_prog_data->push.cross_thread.size + in blorp_get_compute_push_const()
1891 cs_prog_data->push.per_thread.size == sizeof(params->wm_inputs)); in blorp_get_compute_push_const()
1908 if (cs_prog_data->push.cross_thread.size > 0) { in blorp_get_compute_push_const()
1909 memcpy(dst, src, cs_prog_data->push.cross_thread.size); in blorp_get_compute_push_const()
1910 dst += cs_prog_data->push.cross_thread.size; in blorp_get_compute_push_const()
1911 src += cs_prog_data->push.cross_thread.size; in blorp_get_compute_push_const()
1914 if (cs_prog_data->push.per_thread.size > 0) { in blorp_get_compute_push_const()
1916 memcpy(dst, src, (cs_prog_data->push.per_thread.dwords - 1) * 4); in blorp_get_compute_push_const()
[all …]
H A Dblorp_priv.h266 void *cs_prog_data; member
H A Dblorp_clear.c130 &params->cs_prog_kernel, &params->cs_prog_data)) in blorp_params_get_clear_kernel_cs()
177 &params->cs_prog_kernel, &params->cs_prog_data); in blorp_params_get_clear_kernel_cs()
H A Dblorp_blit.c1538 &params->cs_prog_kernel, &params->cs_prog_data)) in blorp_get_blit_kernel_cs()
1558 &params->cs_prog_kernel, &params->cs_prog_data); in blorp_get_blit_kernel_cs()
/aosp_15_r20/external/mesa3d/src/intel/vulkan_hasvk/
H A Danv_cmd_buffer.c579 const struct elk_cs_prog_data *cs_prog_data = get_cs_prog_data(pipeline); in anv_cmd_buffer_cs_push_constants() local
583 elk_cs_get_dispatch_info(devinfo, cs_prog_data, NULL); in anv_cmd_buffer_cs_push_constants()
585 elk_cs_push_const_total_size(cs_prog_data, dispatch.threads); in anv_cmd_buffer_cs_push_constants()
601 if (cs_prog_data->push.cross_thread.size > 0) { in anv_cmd_buffer_cs_push_constants()
602 memcpy(dst, src, cs_prog_data->push.cross_thread.size); in anv_cmd_buffer_cs_push_constants()
603 dst += cs_prog_data->push.cross_thread.size; in anv_cmd_buffer_cs_push_constants()
604 src += cs_prog_data->push.cross_thread.size; in anv_cmd_buffer_cs_push_constants()
607 if (cs_prog_data->push.per_thread.size > 0) { in anv_cmd_buffer_cs_push_constants()
609 memcpy(dst, src, cs_prog_data->push.per_thread.size); in anv_cmd_buffer_cs_push_constants()
613 (range->start * 32 + cs_prog_data->push.cross_thread.size); in anv_cmd_buffer_cs_push_constants()
[all …]
H A DgenX_pipeline.c1882 const struct elk_cs_prog_data *cs_prog_data = get_cs_prog_data(pipeline); in genX() local
1884 anv_pipeline_setup_l3_config(&pipeline->base, cs_prog_data->base.total_shared > 0); in genX()
1887 elk_cs_get_dispatch_info(devinfo, cs_prog_data, NULL); in genX()
1889 ALIGN(cs_prog_data->push.per_thread.regs * dispatch.threads + in genX()
1890 cs_prog_data->push.cross_thread.regs, 2); in genX()
1936 elk_cs_prog_data_prog_offset(cs_prog_data, dispatch.simd_size), in genX()
1942 .BarrierEnable = cs_prog_data->uses_barrier, in genX()
1943 … .SharedLocalMemorySize = intel_compute_slm_encode_size(GFX_VER, cs_prog_data->base.total_shared), in genX()
1948 .ConstantURBEntryReadLength = cs_prog_data->push.per_thread.regs, in genX()
1951 cs_prog_data->push.cross_thread.regs, in genX()
/aosp_15_r20/external/mesa3d/src/intel/vulkan/
H A Danv_cmd_buffer.c1239 const struct brw_cs_prog_data *cs_prog_data = get_cs_prog_data(pipeline); in anv_cmd_buffer_cs_push_constants() local
1243 brw_cs_get_dispatch_info(devinfo, cs_prog_data, NULL); in anv_cmd_buffer_cs_push_constants()
1245 brw_cs_push_const_total_size(cs_prog_data, dispatch.threads); in anv_cmd_buffer_cs_push_constants()
1268 if (cs_prog_data->push.cross_thread.size > 0) { in anv_cmd_buffer_cs_push_constants()
1269 memcpy(dst, src, cs_prog_data->push.cross_thread.size); in anv_cmd_buffer_cs_push_constants()
1270 dst += cs_prog_data->push.cross_thread.size; in anv_cmd_buffer_cs_push_constants()
1271 src += cs_prog_data->push.cross_thread.size; in anv_cmd_buffer_cs_push_constants()
1274 if (cs_prog_data->push.per_thread.size > 0) { in anv_cmd_buffer_cs_push_constants()
1276 memcpy(dst, src, cs_prog_data->push.per_thread.size); in anv_cmd_buffer_cs_push_constants()
1280 (range->start * 32 + cs_prog_data->push.cross_thread.size); in anv_cmd_buffer_cs_push_constants()
[all …]
H A DgenX_pipeline.c2148 const struct brw_cs_prog_data *cs_prog_data = get_cs_prog_data(pipeline); in genX() local
2149 anv_pipeline_setup_l3_config(&pipeline->base, cs_prog_data->base.total_shared > 0); in genX()
2159 const struct brw_cs_prog_data *cs_prog_data = get_cs_prog_data(pipeline); in genX() local
2161 anv_pipeline_setup_l3_config(&pipeline->base, cs_prog_data->base.total_shared > 0); in genX()
2164 brw_cs_get_dispatch_info(devinfo, cs_prog_data, NULL); in genX()
2166 ALIGN(cs_prog_data->push.per_thread.regs * dispatch.threads + in genX()
2167 cs_prog_data->push.cross_thread.regs, 2); in genX()
2182 if (cs_prog_data->base.total_scratch) { in genX()
2186 vfe.PerThreadScratchSpace = ffs(cs_prog_data->base.total_scratch) - 11; in genX()
2195 brw_cs_prog_data_prog_offset(cs_prog_data, dispatch.simd_size), in genX()
[all …]
H A DgenX_cmd_compute.c601 const struct brw_cs_prog_data *cs_prog_data = in genX() local
653 brw_cs_get_dispatch_info(devinfo, cs_prog_data, NULL); in genX()
661 cw.LocalXMaximum = cs_prog_data->local_size[0] - 1; in genX()
662 cw.LocalYMaximum = cs_prog_data->local_size[1] - 1; in genX()
663 cw.LocalZMaximum = cs_prog_data->local_size[2] - 1; in genX()
678 cs_prog_data, in genX()
1058 const struct brw_cs_prog_data *cs_prog_data = in cmd_buffer_trace_rays() local
1061 brw_cs_get_dispatch_info(device->info, cs_prog_data, NULL); in cmd_buffer_trace_rays()
/aosp_15_r20/external/mesa3d/src/gallium/drivers/crocus/
H A Dcrocus_program.c2505 struct elk_cs_prog_data *cs_prog_data = in crocus_compile_cs() local
2507 struct elk_stage_prog_data *prog_data = &cs_prog_data->base; in crocus_compile_cs()
2515 NIR_PASS_V(nir, elk_nir_lower_cs_intrinsics, devinfo, cs_prog_data); in crocus_compile_cs()
2531 .prog_data = cs_prog_data, in crocus_compile_cs()
2551 prog_data, sizeof(*cs_prog_data), NULL, in crocus_compile_cs()
2608 crocus_fill_cs_push_const_buffer(struct elk_cs_prog_data *cs_prog_data, in crocus_fill_cs_push_const_buffer() argument
2612 assert(elk_cs_push_const_total_size(cs_prog_data, threads) > 0); in crocus_fill_cs_push_const_buffer()
2613 assert(cs_prog_data->push.cross_thread.size == 0); in crocus_fill_cs_push_const_buffer()
2614 assert(cs_prog_data->push.per_thread.dwords == 1); in crocus_fill_cs_push_const_buffer()
2615 assert(cs_prog_data->base.param[0] == ELK_PARAM_BUILTIN_SUBGROUP_ID); in crocus_fill_cs_push_const_buffer()
H A Dcrocus_state.c8054 struct elk_cs_prog_data *cs_prog_data = (void *) prog_data; local
8056 elk_cs_get_dispatch_info(devinfo, cs_prog_data, grid->block);
8074 cs_prog_data->local_size[0] == 0 /* Variable local group size */) {
8126 ALIGN(cs_prog_data->push.per_thread.regs * dispatch.threads +
8127 cs_prog_data->push.cross_thread.regs, 2);
8133 cs_prog_data->local_size[0] == 0 /* Variable local group size */) {
8135 assert(cs_prog_data->push.cross_thread.dwords == 0 &&
8136 cs_prog_data->push.per_thread.dwords == 1 &&
8137 cs_prog_data->base.param[0] == ELK_PARAM_BUILTIN_SUBGROUP_ID);
8139 elk_cs_push_const_total_size(cs_prog_data, dispatch.threads);
[all …]
H A Dcrocus_context.h776 void crocus_fill_cs_push_const_buffer(struct elk_cs_prog_data *cs_prog_data,

12