panvk: Support creation of compute pipelines

Signed-off-by: Boris Brezillon <boris.brezillon@collabora.com>
Reviewed-by: Jason Ekstrand <jason.ekstrand@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/15248>

diff --git a/src/panfrost/vulkan/panvk_pipeline.c b/src/panfrost/vulkan/panvk_pipeline.c
index 4dd1676..369dd6b 100644 (file)
--- a/src/panfrost/vulkan/panvk_pipeline.c
+++ b/src/panfrost/vulkan/panvk_pipeline.c
@@ -41,18 +41,6 @@
 #include "vk_util.h"
 
 
-VkResult
-panvk_CreateComputePipelines(VkDevice _device,
-                             VkPipelineCache pipelineCache,
-                             uint32_t count,
-                             const VkComputePipelineCreateInfo *pCreateInfos,
-                             const VkAllocationCallbacks *pAllocator,
-                             VkPipeline *pPipelines)
-{
-   panvk_stub();
-   return VK_SUCCESS;
-}
-
 void
 panvk_DestroyPipeline(VkDevice _device,
                       VkPipeline _pipeline,

diff --git a/src/panfrost/vulkan/panvk_private.h b/src/panfrost/vulkan/panvk_private.h
index 38510f1..e67b354 100644 (file)
--- a/src/panfrost/vulkan/panvk_private.h
+++ b/src/panfrost/vulkan/panvk_private.h
@@ -725,6 +725,7 @@ struct panvk_shader {
    struct pan_shader_info info;
    struct util_dynarray binary;
    unsigned sysval_ubo;
+   struct pan_compute_dim local_size;
 };
 
 struct panvk_shader *
@@ -793,6 +794,10 @@ struct panvk_pipeline {
    } fs;
 
    struct {
+      struct pan_compute_dim local_size;
+   } cs;
+
+   struct {
       unsigned topology;
       bool writes_point_size;
       bool primitive_restart;

diff --git a/src/panfrost/vulkan/panvk_vX_pipeline.c b/src/panfrost/vulkan/panvk_vX_pipeline.c
index 089de93..79ed2ce 100644 (file)
--- a/src/panfrost/vulkan/panvk_vX_pipeline.c
+++ b/src/panfrost/vulkan/panvk_vX_pipeline.c
@@ -48,7 +48,10 @@ struct panvk_pipeline_builder
    struct panvk_device *device;
    struct panvk_pipeline_cache *cache;
    const VkAllocationCallbacks *alloc;
-   const VkGraphicsPipelineCreateInfo *create_info;
+   struct {
+      const VkGraphicsPipelineCreateInfo *gfx;
+      const VkComputePipelineCreateInfo *compute;
+   } create_info;
    const struct panvk_pipeline_layout *layout;
 
    struct panvk_shader *shaders[MESA_SHADER_STAGES];
@@ -110,9 +113,16 @@ panvk_pipeline_builder_compile_shaders(struct panvk_pipeline_builder *builder,
    const VkPipelineShaderStageCreateInfo *stage_infos[MESA_SHADER_STAGES] = {
       NULL
    };
-   for (uint32_t i = 0; i < builder->create_info->stageCount; i++) {
-      gl_shader_stage stage = vk_to_mesa_shader_stage(builder->create_info->pStages[i].stage);
-      stage_infos[stage] = &builder->create_info->pStages[i];
+   const VkPipelineShaderStageCreateInfo *stages =
+      builder->create_info.gfx ?
+      builder->create_info.gfx->pStages :
+      &builder->create_info.compute->stage;
+   unsigned stage_count =
+      builder->create_info.gfx ? builder->create_info.gfx->stageCount : 1;
+
+   for (uint32_t i = 0; i < stage_count; i++) {
+      gl_shader_stage stage = vk_to_mesa_shader_stage(stages[i].stage);
+      stage_infos[stage] = &stages[i];
    }
 
    /* compile shaders in reverse order */
@@ -209,7 +219,8 @@ panvk_pipeline_builder_alloc_static_state_bo(struct panvk_pipeline_builder *buil
          bo_size += pan_size(BLEND) * MAX2(pipeline->blend.state.rt_count, 1);
    }
 
-   if (panvk_pipeline_static_state(pipeline, VK_DYNAMIC_STATE_VIEWPORT) &&
+   if (builder->create_info.gfx &&
+       panvk_pipeline_static_state(pipeline, VK_DYNAMIC_STATE_VIEWPORT) &&
        panvk_pipeline_static_state(pipeline, VK_DYNAMIC_STATE_SCISSOR)) {
       bo_size = ALIGN_POT(bo_size, pan_alignment(VIEWPORT));
       builder->vpd_offset = bo_size;
@@ -262,11 +273,11 @@ panvk_pipeline_builder_upload_sysval(struct panvk_pipeline_builder *builder,
 {
    switch (PAN_SYSVAL_TYPE(id)) {
    case PAN_SYSVAL_VIEWPORT_SCALE:
-      panvk_sysval_upload_viewport_scale(builder->create_info->pViewportState->pViewports,
+      panvk_sysval_upload_viewport_scale(builder->create_info.gfx->pViewportState->pViewports,
                                          data);
       break;
    case PAN_SYSVAL_VIEWPORT_OFFSET:
-      panvk_sysval_upload_viewport_offset(builder->create_info->pViewportState->pViewports,
+      panvk_sysval_upload_viewport_offset(builder->create_info.gfx->pViewportState->pViewports,
                                           data);
       break;
    default:
@@ -345,6 +356,9 @@ panvk_pipeline_builder_init_shaders(struct panvk_pipeline_builder *builder,
 
       pipeline->rsds[i] = gpu_rsd;
       panvk_pipeline_builder_init_sysvals(builder, pipeline, i);
+
+      if (i == MESA_SHADER_COMPUTE)
+         pipeline->cs.local_size = shader->local_size;
    }
 
    pipeline->num_ubos = builder->layout->num_ubos;
@@ -373,17 +387,17 @@ panvk_pipeline_builder_parse_viewport(struct panvk_pipeline_builder *builder,
        panvk_pipeline_static_state(pipeline, VK_DYNAMIC_STATE_VIEWPORT) &&
        panvk_pipeline_static_state(pipeline, VK_DYNAMIC_STATE_SCISSOR)) {
       void *vpd = pipeline->state_bo->ptr.cpu + builder->vpd_offset;
-      panvk_per_arch(emit_viewport)(builder->create_info->pViewportState->pViewports,
-                                    builder->create_info->pViewportState->pScissors,
+      panvk_per_arch(emit_viewport)(builder->create_info.gfx->pViewportState->pViewports,
+                                    builder->create_info.gfx->pViewportState->pScissors,
                                     vpd);
       pipeline->vpd = pipeline->state_bo->ptr.gpu +
                       builder->vpd_offset;
    }
    if (panvk_pipeline_static_state(pipeline, VK_DYNAMIC_STATE_VIEWPORT))
-      pipeline->viewport = builder->create_info->pViewportState->pViewports[0];
+      pipeline->viewport = builder->create_info.gfx->pViewportState->pViewports[0];
 
    if (panvk_pipeline_static_state(pipeline, VK_DYNAMIC_STATE_SCISSOR))
-      pipeline->scissor = builder->create_info->pViewportState->pScissors[0];
+      pipeline->scissor = builder->create_info.gfx->pViewportState->pScissors[0];
 }
 
 static void
@@ -391,7 +405,7 @@ panvk_pipeline_builder_parse_dynamic(struct panvk_pipeline_builder *builder,
                                      struct panvk_pipeline *pipeline)
 {
    const VkPipelineDynamicStateCreateInfo *dynamic_info =
-      builder->create_info->pDynamicState;
+      builder->create_info.gfx->pDynamicState;
 
    if (!dynamic_info)
       return;
@@ -440,9 +454,9 @@ panvk_pipeline_builder_parse_input_assembly(struct panvk_pipeline_builder *build
                                             struct panvk_pipeline *pipeline)
 {
    pipeline->ia.primitive_restart =
-      builder->create_info->pInputAssemblyState->primitiveRestartEnable;
+      builder->create_info.gfx->pInputAssemblyState->primitiveRestartEnable;
    pipeline->ia.topology =
-      translate_prim_topology(builder->create_info->pInputAssemblyState->topology);
+      translate_prim_topology(builder->create_info.gfx->pInputAssemblyState->topology);
 }
 
 static enum pipe_logicop
@@ -576,24 +590,24 @@ panvk_pipeline_builder_parse_color_blend(struct panvk_pipeline_builder *builder,
 {
    struct panfrost_device *pdev = &builder->device->physical_device->pdev;
    pipeline->blend.state.logicop_enable =
-      builder->create_info->pColorBlendState->logicOpEnable;
+      builder->create_info.gfx->pColorBlendState->logicOpEnable;
    pipeline->blend.state.logicop_func =
-      translate_logicop(builder->create_info->pColorBlendState->logicOp);
+      translate_logicop(builder->create_info.gfx->pColorBlendState->logicOp);
    pipeline->blend.state.rt_count = util_last_bit(builder->active_color_attachments);
    memcpy(pipeline->blend.state.constants,
-          builder->create_info->pColorBlendState->blendConstants,
+          builder->create_info.gfx->pColorBlendState->blendConstants,
           sizeof(pipeline->blend.state.constants));
 
    for (unsigned i = 0; i < pipeline->blend.state.rt_count; i++) {
       const VkPipelineColorBlendAttachmentState *in =
-         &builder->create_info->pColorBlendState->pAttachments[i];
+         &builder->create_info.gfx->pColorBlendState->pAttachments[i];
       struct pan_blend_rt_state *out = &pipeline->blend.state.rts[i];
 
       out->format = builder->color_attachment_formats[i];
 
       bool dest_has_alpha = util_format_has_alpha(out->format);
 
-      out->nr_samples = builder->create_info->pMultisampleState->rasterizationSamples;
+      out->nr_samples = builder->create_info.gfx->pMultisampleState->rasterizationSamples;
       out->equation.blend_enable = in->blendEnable;
       out->equation.color_mask = in->colorWriteMask;
       out->equation.rgb_func = translate_blend_op(in->colorBlendOp);
@@ -637,15 +651,15 @@ panvk_pipeline_builder_parse_multisample(struct panvk_pipeline_builder *builder,
                                          struct panvk_pipeline *pipeline)
 {
    unsigned nr_samples =
-      MAX2(builder->create_info->pMultisampleState->rasterizationSamples, 1);
+      MAX2(builder->create_info.gfx->pMultisampleState->rasterizationSamples, 1);
 
    pipeline->ms.rast_samples =
-      builder->create_info->pMultisampleState->rasterizationSamples;
+      builder->create_info.gfx->pMultisampleState->rasterizationSamples;
    pipeline->ms.sample_mask =
-      builder->create_info->pMultisampleState->pSampleMask ?
-      builder->create_info->pMultisampleState->pSampleMask[0] : UINT16_MAX;
+      builder->create_info.gfx->pMultisampleState->pSampleMask ?
+      builder->create_info.gfx->pMultisampleState->pSampleMask[0] : UINT16_MAX;
    pipeline->ms.min_samples =
-      MAX2(builder->create_info->pMultisampleState->minSampleShading * nr_samples, 1);
+      MAX2(builder->create_info.gfx->pMultisampleState->minSampleShading * nr_samples, 1);
 }
 
 static enum mali_stencil_op
@@ -668,54 +682,54 @@ static void
 panvk_pipeline_builder_parse_zs(struct panvk_pipeline_builder *builder,
                                 struct panvk_pipeline *pipeline)
 {
-   pipeline->zs.z_test = builder->create_info->pDepthStencilState->depthTestEnable;
-   pipeline->zs.z_write = builder->create_info->pDepthStencilState->depthWriteEnable;
+   pipeline->zs.z_test = builder->create_info.gfx->pDepthStencilState->depthTestEnable;
+   pipeline->zs.z_write = builder->create_info.gfx->pDepthStencilState->depthWriteEnable;
    pipeline->zs.z_compare_func =
-      panvk_per_arch(translate_compare_func)(builder->create_info->pDepthStencilState->depthCompareOp);
-   pipeline->zs.s_test = builder->create_info->pDepthStencilState->stencilTestEnable;
+      panvk_per_arch(translate_compare_func)(builder->create_info.gfx->pDepthStencilState->depthCompareOp);
+   pipeline->zs.s_test = builder->create_info.gfx->pDepthStencilState->stencilTestEnable;
    pipeline->zs.s_front.fail_op =
-      translate_stencil_op(builder->create_info->pDepthStencilState->front.failOp);
+      translate_stencil_op(builder->create_info.gfx->pDepthStencilState->front.failOp);
    pipeline->zs.s_front.pass_op =
-      translate_stencil_op(builder->create_info->pDepthStencilState->front.passOp);
+      translate_stencil_op(builder->create_info.gfx->pDepthStencilState->front.passOp);
    pipeline->zs.s_front.z_fail_op =
-      translate_stencil_op(builder->create_info->pDepthStencilState->front.depthFailOp);
+      translate_stencil_op(builder->create_info.gfx->pDepthStencilState->front.depthFailOp);
    pipeline->zs.s_front.compare_func =
-      panvk_per_arch(translate_compare_func)(builder->create_info->pDepthStencilState->front.compareOp);
+      panvk_per_arch(translate_compare_func)(builder->create_info.gfx->pDepthStencilState->front.compareOp);
    pipeline->zs.s_front.compare_mask =
-      builder->create_info->pDepthStencilState->front.compareMask;
+      builder->create_info.gfx->pDepthStencilState->front.compareMask;
    pipeline->zs.s_front.write_mask =
-      builder->create_info->pDepthStencilState->front.writeMask;
+      builder->create_info.gfx->pDepthStencilState->front.writeMask;
    pipeline->zs.s_front.ref =
-      builder->create_info->pDepthStencilState->front.reference;
+      builder->create_info.gfx->pDepthStencilState->front.reference;
    pipeline->zs.s_back.fail_op =
-      translate_stencil_op(builder->create_info->pDepthStencilState->back.failOp);
+      translate_stencil_op(builder->create_info.gfx->pDepthStencilState->back.failOp);
    pipeline->zs.s_back.pass_op =
-      translate_stencil_op(builder->create_info->pDepthStencilState->back.passOp);
+      translate_stencil_op(builder->create_info.gfx->pDepthStencilState->back.passOp);
    pipeline->zs.s_back.z_fail_op =
-      translate_stencil_op(builder->create_info->pDepthStencilState->back.depthFailOp);
+      translate_stencil_op(builder->create_info.gfx->pDepthStencilState->back.depthFailOp);
    pipeline->zs.s_back.compare_func =
-      panvk_per_arch(translate_compare_func)(builder->create_info->pDepthStencilState->back.compareOp);
+      panvk_per_arch(translate_compare_func)(builder->create_info.gfx->pDepthStencilState->back.compareOp);
    pipeline->zs.s_back.compare_mask =
-      builder->create_info->pDepthStencilState->back.compareMask;
+      builder->create_info.gfx->pDepthStencilState->back.compareMask;
    pipeline->zs.s_back.write_mask =
-      builder->create_info->pDepthStencilState->back.writeMask;
+      builder->create_info.gfx->pDepthStencilState->back.writeMask;
    pipeline->zs.s_back.ref =
-      builder->create_info->pDepthStencilState->back.reference;
+      builder->create_info.gfx->pDepthStencilState->back.reference;
 }
 
 static void
 panvk_pipeline_builder_parse_rast(struct panvk_pipeline_builder *builder,
                                   struct panvk_pipeline *pipeline)
 {
-   pipeline->rast.clamp_depth = builder->create_info->pRasterizationState->depthClampEnable;
-   pipeline->rast.depth_bias.enable = builder->create_info->pRasterizationState->depthBiasEnable;
+   pipeline->rast.clamp_depth = builder->create_info.gfx->pRasterizationState->depthClampEnable;
+   pipeline->rast.depth_bias.enable = builder->create_info.gfx->pRasterizationState->depthBiasEnable;
    pipeline->rast.depth_bias.constant_factor =
-      builder->create_info->pRasterizationState->depthBiasConstantFactor;
-   pipeline->rast.depth_bias.clamp = builder->create_info->pRasterizationState->depthBiasClamp;
-   pipeline->rast.depth_bias.slope_factor = builder->create_info->pRasterizationState->depthBiasSlopeFactor;
-   pipeline->rast.front_ccw = builder->create_info->pRasterizationState->frontFace == VK_FRONT_FACE_COUNTER_CLOCKWISE;
-   pipeline->rast.cull_front_face = builder->create_info->pRasterizationState->cullMode & VK_CULL_MODE_FRONT_BIT;
-   pipeline->rast.cull_back_face = builder->create_info->pRasterizationState->cullMode & VK_CULL_MODE_BACK_BIT;
+      builder->create_info.gfx->pRasterizationState->depthBiasConstantFactor;
+   pipeline->rast.depth_bias.clamp = builder->create_info.gfx->pRasterizationState->depthBiasClamp;
+   pipeline->rast.depth_bias.slope_factor = builder->create_info.gfx->pRasterizationState->depthBiasSlopeFactor;
+   pipeline->rast.front_ccw = builder->create_info.gfx->pRasterizationState->frontFace == VK_FRONT_FACE_COUNTER_CLOCKWISE;
+   pipeline->rast.cull_front_face = builder->create_info.gfx->pRasterizationState->cullMode & VK_CULL_MODE_FRONT_BIT;
+   pipeline->rast.cull_back_face = builder->create_info.gfx->pRasterizationState->cullMode & VK_CULL_MODE_BACK_BIT;
 }
 
 static bool
@@ -850,7 +864,7 @@ panvk_pipeline_builder_parse_vertex_input(struct panvk_pipeline_builder *builder
 {
    struct panvk_attribs_info *attribs = &pipeline->attribs;
    const VkPipelineVertexInputStateCreateInfo *info =
-      builder->create_info->pVertexInputState;
+      builder->create_info.gfx->pVertexInputState;
 
    for (unsigned i = 0; i < info->vertexBindingDescriptionCount; i++) {
       const VkVertexInputBindingDescription *desc =
@@ -898,22 +912,27 @@ panvk_pipeline_builder_build(struct panvk_pipeline_builder *builder,
       return result;
 
    /* TODO: make those functions return a result and handle errors */
-   panvk_pipeline_builder_parse_dynamic(builder, *pipeline);
-   panvk_pipeline_builder_parse_color_blend(builder, *pipeline);
-   panvk_pipeline_builder_compile_shaders(builder, *pipeline);
-   panvk_pipeline_builder_collect_varyings(builder, *pipeline);
-   panvk_pipeline_builder_parse_input_assembly(builder, *pipeline);
-   panvk_pipeline_builder_parse_multisample(builder, *pipeline);
-   panvk_pipeline_builder_parse_zs(builder, *pipeline);
-   panvk_pipeline_builder_parse_rast(builder, *pipeline);
-   panvk_pipeline_builder_parse_vertex_input(builder, *pipeline);
-
-
-   panvk_pipeline_builder_upload_shaders(builder, *pipeline);
-   panvk_pipeline_builder_init_fs_state(builder, *pipeline);
-   panvk_pipeline_builder_alloc_static_state_bo(builder, *pipeline);
-   panvk_pipeline_builder_init_shaders(builder, *pipeline);
-   panvk_pipeline_builder_parse_viewport(builder, *pipeline);
+   if (builder->create_info.gfx) {
+      panvk_pipeline_builder_parse_dynamic(builder, *pipeline);
+      panvk_pipeline_builder_parse_color_blend(builder, *pipeline);
+      panvk_pipeline_builder_compile_shaders(builder, *pipeline);
+      panvk_pipeline_builder_collect_varyings(builder, *pipeline);
+      panvk_pipeline_builder_parse_input_assembly(builder, *pipeline);
+      panvk_pipeline_builder_parse_multisample(builder, *pipeline);
+      panvk_pipeline_builder_parse_zs(builder, *pipeline);
+      panvk_pipeline_builder_parse_rast(builder, *pipeline);
+      panvk_pipeline_builder_parse_vertex_input(builder, *pipeline);
+      panvk_pipeline_builder_upload_shaders(builder, *pipeline);
+      panvk_pipeline_builder_init_fs_state(builder, *pipeline);
+      panvk_pipeline_builder_alloc_static_state_bo(builder, *pipeline);
+      panvk_pipeline_builder_init_shaders(builder, *pipeline);
+      panvk_pipeline_builder_parse_viewport(builder, *pipeline);
+   } else {
+      panvk_pipeline_builder_compile_shaders(builder, *pipeline);
+      panvk_pipeline_builder_upload_shaders(builder, *pipeline);
+      panvk_pipeline_builder_alloc_static_state_bo(builder, *pipeline);
+      panvk_pipeline_builder_init_shaders(builder, *pipeline);
+   }
 
    return VK_SUCCESS;
 }
@@ -931,7 +950,7 @@ panvk_pipeline_builder_init_graphics(struct panvk_pipeline_builder *builder,
       .device = dev,
       .cache = cache,
       .layout = layout,
-      .create_info = create_info,
+      .create_info.gfx = create_info,
       .alloc = alloc,
    };
 
@@ -996,3 +1015,56 @@ panvk_per_arch(CreateGraphicsPipelines)(VkDevice device,
 
    return VK_SUCCESS;
 }
+
+static void
+panvk_pipeline_builder_init_compute(struct panvk_pipeline_builder *builder,
+                                    struct panvk_device *dev,
+                                    struct panvk_pipeline_cache *cache,
+                                    const VkComputePipelineCreateInfo *create_info,
+                                    const VkAllocationCallbacks *alloc)
+{
+   VK_FROM_HANDLE(panvk_pipeline_layout, layout, create_info->layout);
+   assert(layout);
+   *builder = (struct panvk_pipeline_builder) {
+      .device = dev,
+      .cache = cache,
+      .layout = layout,
+      .create_info.compute = create_info,
+      .alloc = alloc,
+   };
+}
+
+VkResult
+panvk_per_arch(CreateComputePipelines)(VkDevice device,
+                                       VkPipelineCache pipelineCache,
+                                       uint32_t count,
+                                       const VkComputePipelineCreateInfo *pCreateInfos,
+                                       const VkAllocationCallbacks *pAllocator,
+                                       VkPipeline *pPipelines)
+{
+   VK_FROM_HANDLE(panvk_device, dev, device);
+   VK_FROM_HANDLE(panvk_pipeline_cache, cache, pipelineCache);
+
+   for (uint32_t i = 0; i < count; i++) {
+      struct panvk_pipeline_builder builder;
+      panvk_pipeline_builder_init_compute(&builder, dev, cache,
+                                          &pCreateInfos[i], pAllocator);
+
+      struct panvk_pipeline *pipeline;
+      VkResult result = panvk_pipeline_builder_build(&builder, &pipeline);
+      panvk_pipeline_builder_finish(&builder);
+
+      if (result != VK_SUCCESS) {
+         for (uint32_t j = 0; j < i; j++) {
+            panvk_DestroyPipeline(device, pPipelines[j], pAllocator);
+            pPipelines[j] = VK_NULL_HANDLE;
+         }
+
+         return result;
+      }
+
+      pPipelines[i] = panvk_pipeline_to_handle(pipeline);
+   }
+
+   return VK_SUCCESS;
+}

diff --git a/src/panfrost/vulkan/panvk_vX_shader.c b/src/panfrost/vulkan/panvk_vX_shader.c
index a3f3c58..44cda54 100644 (file)
--- a/src/panfrost/vulkan/panvk_vX_shader.c
+++ b/src/panfrost/vulkan/panvk_vX_shader.c
@@ -579,6 +579,9 @@ panvk_per_arch(shader_create)(struct panvk_device *dev,
    shader->info.texture_count = layout->num_textures;
 
    shader->sysval_ubo = sysval_ubo;
+   shader->local_size.x = nir->info.workgroup_size[0];
+   shader->local_size.y = nir->info.workgroup_size[1];
+   shader->local_size.z = nir->info.workgroup_size[2];
 
    ralloc_free(nir);