From 928f5f54349902c497e9293adeae2580123afbd9 Mon Sep 17 00:00:00 2001 From: Caio Marcelo de Oliveira Filho Date: Fri, 27 Mar 2020 08:18:00 -0700 Subject: [PATCH] anv: Stop using cs_prog_data->threads Move the calculation to helper functions -- similar to what GL already needs to do. This is a preparation for dropping this field since this value is expected to be calculated by the drivers now for variable group size case. And also the field would get in the way of brw_compile_cs producing multiple SIMD variants (like FS). Reviewed-by: Jordan Justen Reviewed-by: Paulo Zanoni Part-of: --- src/intel/vulkan/anv_cmd_buffer.c | 5 +++-- src/intel/vulkan/anv_pipeline.c | 17 +++++++++++++++++ src/intel/vulkan/anv_private.h | 6 ++++++ src/intel/vulkan/genX_cmd_buffer.c | 4 ++-- src/intel/vulkan/genX_pipeline.c | 6 ++++-- 5 files changed, 32 insertions(+), 6 deletions(-) diff --git a/src/intel/vulkan/anv_cmd_buffer.c b/src/intel/vulkan/anv_cmd_buffer.c index 188aff6be74..8f94715c0d0 100644 --- a/src/intel/vulkan/anv_cmd_buffer.c +++ b/src/intel/vulkan/anv_cmd_buffer.c @@ -834,8 +834,9 @@ anv_cmd_buffer_cs_push_constants(struct anv_cmd_buffer *cmd_buffer) const struct brw_cs_prog_data *cs_prog_data = get_cs_prog_data(pipeline); const struct anv_push_range *range = &pipeline->cs->bind_map.push_ranges[0]; + const uint32_t threads = anv_cs_threads(pipeline); const unsigned total_push_constants_size = - brw_cs_push_const_total_size(cs_prog_data, cs_prog_data->threads); + brw_cs_push_const_total_size(cs_prog_data, threads); if (total_push_constants_size == 0) return (struct anv_state) { .offset = 0 }; @@ -858,7 +859,7 @@ anv_cmd_buffer_cs_push_constants(struct anv_cmd_buffer *cmd_buffer) } if (cs_prog_data->push.per_thread.size > 0) { - for (unsigned t = 0; t < cs_prog_data->threads; t++) { + for (unsigned t = 0; t < threads; t++) { memcpy(dst, src, cs_prog_data->push.per_thread.size); uint32_t *subgroup_id = dst + diff --git a/src/intel/vulkan/anv_pipeline.c b/src/intel/vulkan/anv_pipeline.c index 9ccf638ed40..b4f6077f0b9 100644 --- a/src/intel/vulkan/anv_pipeline.c +++ b/src/intel/vulkan/anv_pipeline.c @@ -1710,6 +1710,23 @@ anv_pipeline_compile_cs(struct anv_compute_pipeline *pipeline, return VK_SUCCESS; } +uint32_t +anv_cs_workgroup_size(const struct anv_compute_pipeline *pipeline) +{ + const struct brw_cs_prog_data *cs_prog_data = get_cs_prog_data(pipeline); + return cs_prog_data->local_size[0] * + cs_prog_data->local_size[1] * + cs_prog_data->local_size[2]; +} + +uint32_t +anv_cs_threads(const struct anv_compute_pipeline *pipeline) +{ + const struct brw_cs_prog_data *cs_prog_data = get_cs_prog_data(pipeline); + return DIV_ROUND_UP(anv_cs_workgroup_size(pipeline), + cs_prog_data->simd_size); +} + /** * Copy pipeline state not marked as dynamic. * Dynamic state is pipeline state which hasn't been provided at pipeline diff --git a/src/intel/vulkan/anv_private.h b/src/intel/vulkan/anv_private.h index 51f1ae823fc..af07a1d203a 100644 --- a/src/intel/vulkan/anv_private.h +++ b/src/intel/vulkan/anv_private.h @@ -3315,6 +3315,12 @@ anv_pipeline_compile_cs(struct anv_compute_pipeline *pipeline, const char *entrypoint, const VkSpecializationInfo *spec_info); +uint32_t +anv_cs_workgroup_size(const struct anv_compute_pipeline *pipeline); + +uint32_t +anv_cs_threads(const struct anv_compute_pipeline *pipeline); + struct anv_format_plane { enum isl_format isl_format:16; struct isl_swizzle swizzle; diff --git a/src/intel/vulkan/genX_cmd_buffer.c b/src/intel/vulkan/genX_cmd_buffer.c index 7af1da0f5e4..13ad1ced3bc 100644 --- a/src/intel/vulkan/genX_cmd_buffer.c +++ b/src/intel/vulkan/genX_cmd_buffer.c @@ -4295,7 +4295,7 @@ void genX(CmdDispatchBase)( ggw.SIMDSize = prog_data->simd_size / 16; ggw.ThreadDepthCounterMaximum = 0; ggw.ThreadHeightCounterMaximum = 0; - ggw.ThreadWidthCounterMaximum = prog_data->threads - 1; + ggw.ThreadWidthCounterMaximum = anv_cs_threads(pipeline) - 1; ggw.ThreadGroupIDXDimension = groupCountX; ggw.ThreadGroupIDYDimension = groupCountY; ggw.ThreadGroupIDZDimension = groupCountZ; @@ -4411,7 +4411,7 @@ void genX(CmdDispatchIndirect)( ggw.SIMDSize = prog_data->simd_size / 16; ggw.ThreadDepthCounterMaximum = 0; ggw.ThreadHeightCounterMaximum = 0; - ggw.ThreadWidthCounterMaximum = prog_data->threads - 1; + ggw.ThreadWidthCounterMaximum = anv_cs_threads(pipeline) - 1; ggw.RightExecutionMask = pipeline->cs_right_mask; ggw.BottomExecutionMask = 0xffffffff; } diff --git a/src/intel/vulkan/genX_pipeline.c b/src/intel/vulkan/genX_pipeline.c index 2c1d7545b72..c6f479168b6 100644 --- a/src/intel/vulkan/genX_pipeline.c +++ b/src/intel/vulkan/genX_pipeline.c @@ -2321,8 +2321,10 @@ compute_pipeline_create( else pipeline->cs_right_mask = ~0u >> (32 - cs_prog_data->simd_size); + const uint32_t threads = anv_cs_threads(pipeline); + const uint32_t vfe_curbe_allocation = - ALIGN(cs_prog_data->push.per_thread.regs * cs_prog_data->threads + + ALIGN(cs_prog_data->push.per_thread.regs * threads + cs_prog_data->push.cross_thread.regs, 2); const uint32_t subslices = MAX2(device->physical->subslice_total, 1); @@ -2405,7 +2407,7 @@ compute_pipeline_create( .ThreadPreemptionDisable = true, #endif - .NumberofThreadsinGPGPUThreadGroup = cs_prog_data->threads, + .NumberofThreadsinGPGPUThreadGroup = threads, }; GENX(INTERFACE_DESCRIPTOR_DATA_pack)(NULL, pipeline->interface_descriptor_data, -- 2.30.2