Module: Mesa Branch: main Commit: 1056b3e39ecfc8d755ef81b60b106ed85d3b28ab URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=1056b3e39ecfc8d755ef81b60b106ed85d3b28ab
Author: Boris Brezillon <[email protected]> Date: Thu Sep 23 16:27:06 2021 +0200 panvk: Add support for storage image Signed-off-by: Boris Brezillon <[email protected]> Reviewed-by: Jason Ekstrand <[email protected]> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/15248> --- src/panfrost/vulkan/panvk_cmd_buffer.c | 10 +- src/panfrost/vulkan/panvk_descriptor_set.c | 28 ++++-- src/panfrost/vulkan/panvk_private.h | 19 +++- src/panfrost/vulkan/panvk_vX_cmd_buffer.c | 132 ++++++++++++++++++++++---- src/panfrost/vulkan/panvk_vX_cs.c | 4 +- src/panfrost/vulkan/panvk_vX_descriptor_set.c | 51 +++++++--- src/panfrost/vulkan/panvk_vX_image.c | 35 +++++++ src/panfrost/vulkan/panvk_vX_pipeline.c | 3 + src/panfrost/vulkan/panvk_vX_shader.c | 49 +++++++++- 9 files changed, 283 insertions(+), 48 deletions(-) diff --git a/src/panfrost/vulkan/panvk_cmd_buffer.c b/src/panfrost/vulkan/panvk_cmd_buffer.c index 0a4b3871a4d..ba20cb94cd3 100644 --- a/src/panfrost/vulkan/panvk_cmd_buffer.c +++ b/src/panfrost/vulkan/panvk_cmd_buffer.c @@ -41,6 +41,8 @@ panvk_CmdBindVertexBuffers(VkCommandBuffer commandBuffer, const VkDeviceSize *pOffsets) { VK_FROM_HANDLE(panvk_cmd_buffer, cmdbuf, commandBuffer); + struct panvk_descriptor_state *desc_state = + panvk_cmd_get_desc_state(cmdbuf, GRAPHICS); assert(firstBinding + bindingCount <= MAX_VBS); @@ -50,8 +52,9 @@ panvk_CmdBindVertexBuffers(VkCommandBuffer commandBuffer, cmdbuf->state.vb.bufs[firstBinding + i].address = buf->bo->ptr.gpu + pOffsets[i]; cmdbuf->state.vb.bufs[firstBinding + i].size = buf->size - pOffsets[i]; } + cmdbuf->state.vb.count = MAX2(cmdbuf->state.vb.count, firstBinding + bindingCount); - cmdbuf->state.vb.attrib_bufs = cmdbuf->state.vb.attribs = 0; + desc_state->vs_attrib_bufs = desc_state->vs_attribs = 0; } void @@ -121,6 +124,11 @@ panvk_CmdBindDescriptorSets(VkCommandBuffer commandBuffer, if (set->layout->num_samplers) descriptors_state->samplers = 0; + + if (set->layout->num_imgs) { + descriptors_state->vs_attrib_bufs = descriptors_state->non_vs_attrib_bufs = 0; + descriptors_state->vs_attribs = descriptors_state->non_vs_attribs = 0; + } } assert(dynoffset_idx == dynamicOffsetCount); diff --git a/src/panfrost/vulkan/panvk_descriptor_set.c b/src/panfrost/vulkan/panvk_descriptor_set.c index 7c2957fe687..7240dd16cc5 100644 --- a/src/panfrost/vulkan/panvk_descriptor_set.c +++ b/src/panfrost/vulkan/panvk_descriptor_set.c @@ -87,7 +87,7 @@ panvk_CreateDescriptorSetLayout(VkDevice _device, set_layout->binding_count = num_bindings; unsigned sampler_idx = 0, tex_idx = 0, ubo_idx = 0, ssbo_idx = 0; - unsigned dyn_ubo_idx = 0, dyn_ssbo_idx = 0, desc_idx = 0; + unsigned dyn_ubo_idx = 0, dyn_ssbo_idx = 0, desc_idx = 0, img_idx = 0; for (unsigned i = 0; i < pCreateInfo->bindingCount; i++) { const VkDescriptorSetLayoutBinding *binding = &bindings[i]; @@ -120,7 +120,6 @@ panvk_CreateDescriptorSetLayout(VkDevice _device, tex_idx += binding_layout->array_size; break; case VK_DESCRIPTOR_TYPE_SAMPLED_IMAGE: - case VK_DESCRIPTOR_TYPE_STORAGE_IMAGE: case VK_DESCRIPTOR_TYPE_UNIFORM_TEXEL_BUFFER: case VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER: case VK_DESCRIPTOR_TYPE_INPUT_ATTACHMENT: @@ -143,6 +142,10 @@ panvk_CreateDescriptorSetLayout(VkDevice _device, binding_layout->ssbo_idx = ssbo_idx; ssbo_idx += binding_layout->array_size; break; + case VK_DESCRIPTOR_TYPE_STORAGE_IMAGE: + binding_layout->img_idx = img_idx; + img_idx += binding_layout->array_size; + break; default: unreachable("Invalid descriptor type"); } @@ -155,6 +158,7 @@ panvk_CreateDescriptorSetLayout(VkDevice _device, set_layout->num_dyn_ubos = dyn_ubo_idx; set_layout->num_ssbos = ssbo_idx; set_layout->num_dyn_ssbos = dyn_ssbo_idx; + set_layout->num_imgs = img_idx; free(bindings); *pSetLayout = panvk_descriptor_set_layout_to_handle(set_layout); @@ -181,6 +185,7 @@ panvk_DestroyDescriptorSetLayout(VkDevice _device, /* FIXME: make sure those values are correct */ #define PANVK_MAX_TEXTURES (1 << 16) +#define PANVK_MAX_IMAGES (1 << 8) #define PANVK_MAX_SAMPLERS (1 << 16) #define PANVK_MAX_UBOS 255 @@ -203,8 +208,9 @@ panvk_GetDescriptorSetLayoutSupport(VkDevice _device, return; } - unsigned sampler_idx = 0, tex_idx = 0, ubo_idx = 0, UNUSED ssbo_idx = 0, - UNUSED dynoffset_idx = 0; + unsigned sampler_idx = 0, tex_idx = 0, ubo_idx = 0; + unsigned ssbo_idx = 0, dynoffset_idx = 0, img_idx = 0; + for (unsigned i = 0; i < pCreateInfo->bindingCount; i++) { const VkDescriptorSetLayoutBinding *binding = &bindings[i]; @@ -217,7 +223,6 @@ panvk_GetDescriptorSetLayoutSupport(VkDevice _device, tex_idx += binding->descriptorCount; break; case VK_DESCRIPTOR_TYPE_SAMPLED_IMAGE: - case VK_DESCRIPTOR_TYPE_STORAGE_IMAGE: case VK_DESCRIPTOR_TYPE_UNIFORM_TEXEL_BUFFER: case VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER: case VK_DESCRIPTOR_TYPE_INPUT_ATTACHMENT: @@ -235,6 +240,9 @@ panvk_GetDescriptorSetLayoutSupport(VkDevice _device, case VK_DESCRIPTOR_TYPE_STORAGE_BUFFER: ssbo_idx += binding->descriptorCount; break; + case VK_DESCRIPTOR_TYPE_STORAGE_IMAGE: + img_idx += binding->descriptorCount; + break; default: unreachable("Invalid descriptor type"); } @@ -245,7 +253,8 @@ panvk_GetDescriptorSetLayoutSupport(VkDevice _device, */ if (tex_idx > PANVK_MAX_TEXTURES / MAX_SETS || sampler_idx > PANVK_MAX_SAMPLERS / MAX_SETS || - ubo_idx > PANVK_MAX_UBOS / MAX_SETS) + ubo_idx > PANVK_MAX_UBOS / MAX_SETS || + img_idx > PANVK_MAX_IMAGES / MAX_SETS) return; pSupport->supported = true; @@ -275,7 +284,7 @@ panvk_CreatePipelineLayout(VkDevice _device, _mesa_sha1_init(&ctx); unsigned sampler_idx = 0, tex_idx = 0, ssbo_idx = 0, ubo_idx = 0; - unsigned dyn_ubo_idx = 0, dyn_ssbo_idx = 0; + unsigned dyn_ubo_idx = 0, dyn_ssbo_idx = 0, img_idx = 0; for (unsigned set = 0; set < pCreateInfo->setLayoutCount; set++) { VK_FROM_HANDLE(panvk_descriptor_set_layout, set_layout, pCreateInfo->pSetLayouts[set]); @@ -286,12 +295,14 @@ panvk_CreatePipelineLayout(VkDevice _device, layout->sets[set].dyn_ubo_offset = dyn_ubo_idx; layout->sets[set].ssbo_offset = ssbo_idx; layout->sets[set].dyn_ssbo_offset = dyn_ssbo_idx; + layout->sets[set].img_offset = img_idx; sampler_idx += set_layout->num_samplers; tex_idx += set_layout->num_textures; ubo_idx += set_layout->num_ubos; dyn_ubo_idx += set_layout->num_dyn_ubos; ssbo_idx += set_layout->num_ssbos; dyn_ssbo_idx += set_layout->num_dyn_ssbos; + img_idx += set_layout->num_imgs; for (unsigned b = 0; b < set_layout->binding_count; b++) { struct panvk_descriptor_set_binding_layout *binding_layout = @@ -327,6 +338,7 @@ panvk_CreatePipelineLayout(VkDevice _device, layout->num_dyn_ubos = dyn_ubo_idx; layout->num_ssbos = ssbo_idx; layout->num_dyn_ssbos = dyn_ssbo_idx; + layout->num_imgs = img_idx; _mesa_sha1_final(&ctx, layout->sha1); @@ -444,6 +456,8 @@ panvk_descriptor_set_destroy(struct panvk_device *device, vk_free(&device->vk.alloc, set->dyn_ubos); vk_free(&device->vk.alloc, set->ssbos); vk_free(&device->vk.alloc, set->dyn_ssbos); + vk_free(&device->vk.alloc, set->img_fmts); + vk_free(&device->vk.alloc, set->img_attrib_bufs); vk_free(&device->vk.alloc, set->descs); vk_object_free(&device->vk, NULL, set); } diff --git a/src/panfrost/vulkan/panvk_private.h b/src/panfrost/vulkan/panvk_private.h index 4648d4f9e60..66566142b27 100644 --- a/src/panfrost/vulkan/panvk_private.h +++ b/src/panfrost/vulkan/panvk_private.h @@ -362,6 +362,8 @@ struct panvk_descriptor_set { struct panvk_buffer_desc *dyn_ubos; void *samplers; void *textures; + void *img_attrib_bufs; + uint32_t *img_fmts; }; #define MAX_SETS 4 @@ -408,6 +410,7 @@ struct panvk_descriptor_set_layout { unsigned num_dyn_ubos; unsigned num_ssbos; unsigned num_dyn_ssbos; + unsigned num_imgs; /* Number of bindings in this descriptor set */ uint32_t binding_count; @@ -426,6 +429,7 @@ struct panvk_pipeline_layout { unsigned num_dyn_ubos; unsigned num_ssbos; unsigned num_dyn_ssbos; + uint32_t num_imgs; uint32_t num_sets; struct { @@ -441,6 +445,7 @@ struct panvk_pipeline_layout { unsigned dyn_ubo_offset; unsigned ssbo_offset; unsigned dyn_ssbo_offset; + unsigned img_offset; } sets[MAX_SETS]; }; @@ -504,6 +509,10 @@ struct panvk_descriptor_state { mali_ptr textures; mali_ptr samplers; mali_ptr push_constants; + mali_ptr vs_attribs; + mali_ptr vs_attrib_bufs; + mali_ptr non_vs_attribs; + mali_ptr non_vs_attrib_bufs; }; #define INVOCATION_DESC_WORDS 2 @@ -522,10 +531,10 @@ struct panvk_draw_info { struct { mali_ptr varyings; mali_ptr attributes; + mali_ptr attribute_bufs; mali_ptr push_constants; } stages[MESA_SHADER_STAGES]; mali_ptr varying_bufs; - mali_ptr attribute_bufs; mali_ptr textures; mali_ptr samplers; mali_ptr ubos; @@ -596,8 +605,6 @@ struct panvk_cmd_state { struct { struct panvk_attrib_buf bufs[MAX_VBS]; unsigned count; - mali_ptr attribs; - mali_ptr attrib_bufs; } vb; /* Index buffer */ @@ -730,6 +737,7 @@ struct panvk_shader { struct util_dynarray binary; unsigned sysval_ubo; struct pan_compute_dim local_size; + bool has_img_access; }; struct panvk_shader * @@ -775,6 +783,9 @@ struct panvk_pipeline { mali_ptr vpd; mali_ptr rsds[MESA_SHADER_STAGES]; + /* shader stage bit is set of the stage accesses storage images */ + uint32_t img_access_mask; + unsigned num_ubos; unsigned num_sysvals; @@ -931,6 +942,7 @@ unsigned panvk_image_get_total_size(const struct panvk_image *image); #define TEXTURE_DESC_WORDS 8 +#define ATTRIB_BUF_DESC_WORDS 4 struct panvk_image_view { struct vk_object_base base; @@ -940,6 +952,7 @@ struct panvk_image_view { struct panfrost_bo *bo; struct { uint32_t tex[TEXTURE_DESC_WORDS]; + uint32_t img_attrib_buf[ATTRIB_BUF_DESC_WORDS * 2]; } descs; }; diff --git a/src/panfrost/vulkan/panvk_vX_cmd_buffer.c b/src/panfrost/vulkan/panvk_vX_cmd_buffer.c index 3bc413f99bd..0ec3f407361 100644 --- a/src/panfrost/vulkan/panvk_vX_cmd_buffer.c +++ b/src/panfrost/vulkan/panvk_vX_cmd_buffer.c @@ -661,51 +661,143 @@ panvk_draw_prepare_varyings(struct panvk_cmd_buffer *cmdbuf, } static void -panvk_draw_prepare_attributes(struct panvk_cmd_buffer *cmdbuf, +panvk_fill_non_vs_attribs(struct panvk_cmd_buffer *cmdbuf, + struct panvk_cmd_bind_point_state *bind_point_state, + void *attrib_bufs, void *attribs, + unsigned first_buf) +{ + struct panvk_descriptor_state *desc_state = &bind_point_state->desc_state; + const struct panvk_pipeline *pipeline = bind_point_state->pipeline; + + for (unsigned s = 0; s < pipeline->layout->num_sets; s++) { + const struct panvk_descriptor_set *set = desc_state->sets[s]; + + if (!set) continue; + + const struct panvk_descriptor_set_layout *layout = set->layout; + unsigned img_idx = pipeline->layout->sets[s].img_offset; + unsigned offset = img_idx * pan_size(ATTRIBUTE_BUFFER) * 2; + unsigned size = layout->num_imgs * pan_size(ATTRIBUTE_BUFFER) * 2; + + memcpy(attrib_bufs + offset, desc_state->sets[s]->img_attrib_bufs, size); + + offset = img_idx * pan_size(ATTRIBUTE); + for (unsigned i = 0; i < layout->num_imgs; i++) { + pan_pack(attribs + offset, ATTRIBUTE, cfg) { + cfg.buffer_index = first_buf + (img_idx + i) * 2; + cfg.format = desc_state->sets[s]->img_fmts[i]; + cfg.offset_enable = PAN_ARCH <= 5; + } + offset += pan_size(ATTRIBUTE); + } + } +} + +static void +panvk_prepare_non_vs_attribs(struct panvk_cmd_buffer *cmdbuf, + struct panvk_cmd_bind_point_state *bind_point_state) +{ + struct panvk_descriptor_state *desc_state = &bind_point_state->desc_state; + const struct panvk_pipeline *pipeline = bind_point_state->pipeline; + + if (desc_state->non_vs_attribs || !pipeline->img_access_mask) + return; + + unsigned attrib_count = pipeline->layout->num_imgs; + unsigned attrib_buf_count = (pipeline->layout->num_imgs * 2); + struct panfrost_ptr bufs = + pan_pool_alloc_desc_array(&cmdbuf->desc_pool.base, + attrib_buf_count + (PAN_ARCH >= 6 ? 1 : 0), + ATTRIBUTE_BUFFER); + struct panfrost_ptr attribs = + pan_pool_alloc_desc_array(&cmdbuf->desc_pool.base, attrib_count, + ATTRIBUTE); + + panvk_fill_non_vs_attribs(cmdbuf, bind_point_state, bufs.cpu, attribs.cpu, 0); + + desc_state->non_vs_attrib_bufs = bufs.gpu; + desc_state->non_vs_attribs = attribs.gpu; +} + +static void +panvk_draw_prepare_vs_attribs(struct panvk_cmd_buffer *cmdbuf, struct panvk_draw_info *draw) { - const struct panvk_pipeline *pipeline = panvk_cmd_get_pipeline(cmdbuf, GRAPHICS); + struct panvk_cmd_bind_point_state *bind_point_state = + panvk_cmd_get_bind_point_state(cmdbuf, GRAPHICS); + struct panvk_descriptor_state *desc_state = &bind_point_state->desc_state; + const struct panvk_pipeline *pipeline = bind_point_state->pipeline; + unsigned num_imgs = + pipeline->img_access_mask & BITFIELD_BIT(MESA_SHADER_VERTEX) ? + pipeline->layout->num_imgs : 0; + unsigned attrib_count = pipeline->attribs.buf_count + num_imgs; - /* TODO: images */ - if (!pipeline->attribs.buf_count) + if (desc_state->vs_attribs || !attrib_count) return; - if (cmdbuf->state.vb.attribs) { - draw->stages[MESA_SHADER_VERTEX].attributes = cmdbuf->state.vb.attribs; - draw->attribute_bufs = cmdbuf->state.vb.attrib_bufs; + if (!pipeline->attribs.buf_count) { + panvk_prepare_non_vs_attribs(cmdbuf, bind_point_state); + desc_state->vs_attrib_bufs = desc_state->non_vs_attrib_bufs; + desc_state->vs_attribs = desc_state->non_vs_attribs; return; } - unsigned buf_count = pipeline->attribs.buf_count * 2; + unsigned attrib_buf_count = attrib_count * 2; struct panfrost_ptr bufs = pan_pool_alloc_desc_array(&cmdbuf->desc_pool.base, - buf_count + (PAN_ARCH >= 6 ? 1 : 0), + attrib_buf_count + (PAN_ARCH >= 6 ? 1 : 0), ATTRIBUTE_BUFFER); + struct panfrost_ptr attribs = + pan_pool_alloc_desc_array(&cmdbuf->desc_pool.base, attrib_count, + ATTRIBUTE); panvk_per_arch(emit_attrib_bufs)(&pipeline->attribs, cmdbuf->state.vb.bufs, cmdbuf->state.vb.count, draw, bufs.cpu); - cmdbuf->state.vb.attrib_bufs = bufs.gpu; - - struct panfrost_ptr attribs = - pan_pool_alloc_desc_array(&cmdbuf->desc_pool.base, - pipeline->attribs.attrib_count, - ATTRIBUTE); - panvk_per_arch(emit_attribs)(cmdbuf->device, &pipeline->attribs, cmdbuf->state.vb.bufs, cmdbuf->state.vb.count, attribs.cpu); + if (attrib_count > pipeline->attribs.buf_count) { + unsigned bufs_offset = pipeline->attribs.buf_count * pan_size(ATTRIBUTE_BUFFER) * 2; + unsigned attribs_offset = pipeline->attribs.buf_count * pan_size(ATTRIBUTE); + + panvk_fill_non_vs_attribs(cmdbuf, bind_point_state, + bufs.cpu + bufs_offset, attribs.cpu + attribs_offset, + pipeline->attribs.buf_count * 2); + } + /* A NULL entry is needed to stop prefecting on Bifrost */ #if PAN_ARCH >= 6 - memset(bufs.cpu + (pan_size(ATTRIBUTE_BUFFER) * buf_count), 0, + memset(bufs.cpu + (pan_size(ATTRIBUTE_BUFFER) * attrib_buf_count), 0, pan_size(ATTRIBUTE_BUFFER)); #endif - cmdbuf->state.vb.attribs = attribs.gpu; - draw->stages[MESA_SHADER_VERTEX].attributes = cmdbuf->state.vb.attribs; - draw->attribute_bufs = cmdbuf->state.vb.attrib_bufs; + desc_state->vs_attrib_bufs = bufs.gpu; + desc_state->vs_attribs = attribs.gpu; +} + +static void +panvk_draw_prepare_attributes(struct panvk_cmd_buffer *cmdbuf, + struct panvk_draw_info *draw) +{ + struct panvk_cmd_bind_point_state *bind_point_state = + panvk_cmd_get_bind_point_state(cmdbuf, GRAPHICS); + struct panvk_descriptor_state *desc_state = &bind_point_state->desc_state; + const struct panvk_pipeline *pipeline = bind_point_state->pipeline; + + for (unsigned i = 0; i < ARRAY_SIZE(draw->stages); i++) { + if (i == MESA_SHADER_VERTEX) { + panvk_draw_prepare_vs_attribs(cmdbuf, draw); + draw->stages[i].attributes = desc_state->vs_attribs; + draw->stages[i].attribute_bufs = desc_state->vs_attrib_bufs; + } else if (pipeline->img_access_mask & BITFIELD_BIT(i)) { + panvk_prepare_non_vs_attribs(cmdbuf, bind_point_state); + draw->stages[i].attributes = desc_state->non_vs_attribs; + draw->stages[i].attribute_bufs = desc_state->non_vs_attrib_bufs; + } + } } static void diff --git a/src/panfrost/vulkan/panvk_vX_cs.c b/src/panfrost/vulkan/panvk_vX_cs.c index ad5ceeff359..36c0bdbce79 100644 --- a/src/panfrost/vulkan/panvk_vX_cs.c +++ b/src/panfrost/vulkan/panvk_vX_cs.c @@ -402,7 +402,7 @@ panvk_per_arch(emit_vertex_job)(const struct panvk_pipeline *pipeline, cfg.draw_descriptor_is_64b = true; cfg.state = pipeline->rsds[MESA_SHADER_VERTEX]; cfg.attributes = draw->stages[MESA_SHADER_VERTEX].attributes; - cfg.attribute_buffers = draw->attribute_bufs; + cfg.attribute_buffers = draw->stages[MESA_SHADER_VERTEX].attribute_bufs; cfg.varyings = draw->stages[MESA_SHADER_VERTEX].varyings; cfg.varying_buffers = draw->varying_bufs; cfg.thread_storage = draw->tls; @@ -461,7 +461,7 @@ panvk_emit_tiler_dcd(const struct panvk_pipeline *pipeline, cfg.position = draw->position; cfg.state = draw->fs_rsd; cfg.attributes = draw->stages[MESA_SHADER_FRAGMENT].attributes; - cfg.attribute_buffers = draw->attribute_bufs; + cfg.attribute_buffers = draw->stages[MESA_SHADER_FRAGMENT].attribute_bufs; cfg.viewport = draw->viewport; cfg.varyings = draw->stages[MESA_SHADER_FRAGMENT].varyings; cfg.varying_buffers = cfg.varyings ? draw->varying_bufs : 0; diff --git a/src/panfrost/vulkan/panvk_vX_descriptor_set.c b/src/panfrost/vulkan/panvk_vX_descriptor_set.c index f1ac690ace2..37ef98eae46 100644 --- a/src/panfrost/vulkan/panvk_vX_descriptor_set.c +++ b/src/panfrost/vulkan/panvk_vX_descriptor_set.c @@ -114,6 +114,22 @@ panvk_per_arch(descriptor_set_create)(struct panvk_device *device, goto err_free_set; } + if (layout->num_imgs) { + set->img_fmts = + vk_zalloc(&device->vk.alloc, + sizeof(*set->img_fmts) * layout->num_imgs, + 8, VK_OBJECT_TYPE_DESCRIPTOR_SET); + if (!set->img_fmts) + goto err_free_set; + + set->img_attrib_bufs = + vk_zalloc(&device->vk.alloc, + pan_size(ATTRIBUTE_BUFFER) * 2 * layout->num_imgs, + 8, VK_OBJECT_TYPE_DESCRIPTOR_SET); + if (!set->img_attrib_bufs) + goto err_free_set; + } + for (unsigned i = 0; i < layout->binding_count; i++) { if (!layout->bindings[i].immutable_samplers) continue; @@ -134,6 +150,8 @@ err_free_set: vk_free(&device->vk.alloc, set->dyn_ssbos); vk_free(&device->vk.alloc, set->ubos); vk_free(&device->vk.alloc, set->dyn_ubos); + vk_free(&device->vk.alloc, set->img_fmts); + vk_free(&device->vk.alloc, set->img_attrib_bufs); vk_free(&device->vk.alloc, set->descs); vk_object_free(&device->vk, NULL, set); return vk_error(device, VK_ERROR_OUT_OF_HOST_MEMORY); @@ -171,17 +189,6 @@ err_free_sets: return result; } -static void -panvk_set_image_desc(struct panvk_descriptor *desc, - const VkDescriptorImageInfo *pImageInfo) -{ - VK_FROM_HANDLE(panvk_sampler, sampler, pImageInfo->sampler); - VK_FROM_HANDLE(panvk_image_view, image_view, pImageInfo->imageView); - desc->image.sampler = sampler; - desc->image.view = image_view; - desc->image.layout = pImageInfo->imageLayout; -} - static void panvk_set_texel_buffer_view_desc(struct panvk_descriptor *desc, const VkBufferView *pTexelBufferView) @@ -237,6 +244,20 @@ panvk_per_arch(set_texture_desc)(struct panvk_descriptor_set *set, #endif } +static void +panvk_set_img_desc(struct panvk_device *dev, + struct panvk_descriptor_set *set, + unsigned idx, + const VkDescriptorImageInfo *pImageInfo) +{ + const struct panfrost_device *pdev = &dev->physical_device->pdev; + VK_FROM_HANDLE(panvk_image_view, view, pImageInfo->imageView); + void *attrib_buf = (uint8_t *)set->img_attrib_bufs + (pan_size(ATTRIBUTE_BUFFER) * 2 * idx); + + set->img_fmts[idx] = pdev->formats[view->pview.format].hw; + memcpy(attrib_buf, view->descs.img_attrib_buf, pan_size(ATTRIBUTE_BUFFER) * 2); +} + static void panvk_per_arch(write_descriptor_set)(struct panvk_device *dev, const VkWriteDescriptorSet *pDescriptorWrite) @@ -292,8 +313,12 @@ panvk_per_arch(write_descriptor_set)(struct panvk_device *dev, case VK_DESCRIPTOR_TYPE_STORAGE_IMAGE: case VK_DESCRIPTOR_TYPE_INPUT_ATTACHMENT: - for (unsigned i = 0; i < ndescs; i++) - panvk_set_image_desc(&descs[i], &pDescriptorWrite->pImageInfo[src_offset + i]); + for (unsigned i = 0; i < ndescs; i++) { + const VkDescriptorImageInfo *info = &pDescriptorWrite->pImageInfo[src_offset + i]; + unsigned img = binding_layout->img_idx + dest_offset + i; + + panvk_set_img_desc(dev, set, img, info); + } break; case VK_DESCRIPTOR_TYPE_UNIFORM_TEXEL_BUFFER: diff --git a/src/panfrost/vulkan/panvk_vX_image.c b/src/panfrost/vulkan/panvk_vX_image.c index f11c2f4b5fe..dc2cb31849f 100644 --- a/src/panfrost/vulkan/panvk_vX_image.c +++ b/src/panfrost/vulkan/panvk_vX_image.c @@ -154,6 +154,41 @@ panvk_per_arch(CreateImageView)(VkDevice _device, GENX(panfrost_new_texture)(pdev, &view->pview, tex_desc, &surf_descs); } + if (image->usage & VK_IMAGE_USAGE_STORAGE_BIT) { + uint8_t *attrib_buf = (uint8_t *)view->descs.img_attrib_buf; + bool is_3d = image->pimage.layout.dim == MALI_TEXTURE_DIMENSION_3D; + unsigned offset = image->pimage.data.offset; + offset += panfrost_texture_offset(&image->pimage.layout, + view->pview.first_level, + is_3d ? 0 : view->pview.first_layer, + is_3d ? view->pview.first_layer : 0); + + pan_pack(attrib_buf, ATTRIBUTE_BUFFER, cfg) { + cfg.type = image->pimage.layout.modifier == DRM_FORMAT_MOD_LINEAR ? + MALI_ATTRIBUTE_TYPE_3D_LINEAR : MALI_ATTRIBUTE_TYPE_3D_INTERLEAVED; + cfg.pointer = image->pimage.data.bo->ptr.gpu + offset; + cfg.stride = util_format_get_blocksize(view->pview.format); + cfg.size = image->pimage.data.bo->size - offset; + } + + attrib_buf += pan_size(ATTRIBUTE_BUFFER); + pan_pack(attrib_buf, ATTRIBUTE_BUFFER_CONTINUATION_3D, cfg) { + unsigned level = view->pview.first_level; + + cfg.s_dimension = u_minify(image->pimage.layout.width, level); + cfg.t_dimension = u_minify(image->pimage.layout.height, level); + cfg.r_dimension = + view->pview.dim == MALI_TEXTURE_DIMENSION_3D ? + u_minify(image->pimage.layout.depth, level) : + (view->pview.last_layer - view->pview.first_layer + 1); + cfg.row_stride = image->pimage.layout.slices[level].row_stride; + if (cfg.r_dimension > 1) { + cfg.slice_stride = + panfrost_get_layer_stride(&image->pimage.layout, level); + } + } + } + *pView = panvk_image_view_to_handle(view); return VK_SUCCESS; } diff --git a/src/panfrost/vulkan/panvk_vX_pipeline.c b/src/panfrost/vulkan/panvk_vX_pipeline.c index 971f00e7c9a..de8e7787be3 100644 --- a/src/panfrost/vulkan/panvk_vX_pipeline.c +++ b/src/panfrost/vulkan/panvk_vX_pipeline.c @@ -326,6 +326,9 @@ panvk_pipeline_builder_init_shaders(struct panvk_pipeline_builder *builder, pipeline->tls_size = MAX2(pipeline->tls_size, shader->info.tls_size); pipeline->wls_size = MAX2(pipeline->wls_size, shader->info.wls_size); + if (shader->has_img_access) + pipeline->img_access_mask |= BITFIELD_BIT(i); + if (i == MESA_SHADER_VERTEX && shader->info.vs.writes_point_size) pipeline->ia.writes_point_size = true; diff --git a/src/panfrost/vulkan/panvk_vX_shader.c b/src/panfrost/vulkan/panvk_vX_shader.c index 8ebfc93c3b0..82b8cc6223e 100644 --- a/src/panfrost/vulkan/panvk_vX_shader.c +++ b/src/panfrost/vulkan/panvk_vX_shader.c @@ -29,6 +29,7 @@ #include "panvk_private.h" #include "nir_builder.h" +#include "nir_deref.h" #include "nir_lower_blend.h" #include "nir_conversion_builder.h" #include "spirv/nir_spirv.h" @@ -81,6 +82,7 @@ panvk_spirv_to_nir(const void *code, struct panvk_lower_misc_ctx { struct panvk_shader *shader; const struct panvk_pipeline_layout *layout; + bool has_img_access; }; static unsigned @@ -187,9 +189,40 @@ lower_load_vulkan_descriptor(nir_builder *b, nir_intrinsic_instr *intrin) nir_instr_remove(&intrin->instr); } +static void +type_size_align_1(const struct glsl_type *type, unsigned *size, unsigned *align) +{ + unsigned s; + + if (glsl_type_is_array(type)) + s = glsl_get_aoa_size(type); + else + s = 1; + + *size = s; + *align = s; +} + +static nir_ssa_def * +get_img_index(nir_builder *b, nir_deref_instr *deref, + const struct panvk_lower_misc_ctx *ctx) +{ + nir_variable *var = nir_deref_instr_get_variable(deref); + unsigned set = var->data.descriptor_set; + unsigned binding = var->data.binding; + const struct panvk_descriptor_set_binding_layout *bind_layout = + &ctx->layout->sets[set].layout->bindings[binding]; + assert(bind_layout->type == VK_DESCRIPTOR_TYPE_STORAGE_IMAGE || + bind_layout->type == VK_DESCRIPTOR_TYPE_UNIFORM_TEXEL_BUFFER || + bind_layout->type == VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER); + + return nir_iadd_imm(b, nir_build_deref_offset(b, deref, type_size_align_1), + bind_layout->img_idx + ctx->layout->sets[set].img_offset); +} + static bool lower_intrinsic(nir_builder *b, nir_intrinsic_instr *intr, - const struct panvk_lower_misc_ctx *ctx) + struct panvk_lower_misc_ctx *ctx) { switch (intr->intrinsic) { case nir_intrinsic_vulkan_resource_index: @@ -198,6 +231,15 @@ lower_intrinsic(nir_builder *b, nir_intrinsic_instr *intr, case nir_intrinsic_load_vulkan_descriptor: lower_load_vulkan_descriptor(b, intr); return true; + case nir_intrinsic_image_deref_store: + case nir_intrinsic_image_deref_load: { + nir_deref_instr *deref = nir_src_as_deref(intr->src[0]); + + b->cursor = nir_before_instr(&intr->instr); + nir_rewrite_image_intrinsic(intr, get_img_index(b, deref, ctx), false); + ctx->has_img_access = true; + return true; + } default: return false; } @@ -209,7 +251,7 @@ panvk_lower_misc_instr(nir_builder *b, nir_instr *instr, void *data) { - const struct panvk_lower_misc_ctx *ctx = data; + struct panvk_lower_misc_ctx *ctx = data; switch (instr->type) { case nir_instr_type_tex: @@ -569,6 +611,7 @@ panvk_per_arch(shader_create)(struct panvk_device *dev, .layout = layout, }; NIR_PASS_V(nir, panvk_lower_misc, &ctx); + shader->has_img_access = ctx.has_img_access; nir_shader_gather_info(nir, nir_shader_get_entrypoint(nir)); if (unlikely(dev->physical_device->instance->debug_flags & PANVK_DEBUG_NIR)) { @@ -583,6 +626,8 @@ panvk_per_arch(shader_create)(struct panvk_device *dev, shader->info.sysvals.sysval_count ? sysval_ubo + 1 : layout->num_ubos; shader->info.sampler_count = layout->num_samplers; shader->info.texture_count = layout->num_textures; + if (ctx.has_img_access) + shader->info.attribute_count += layout->num_imgs; shader->sysval_ubo = sysval_ubo; shader->local_size.x = nir->info.workgroup_size[0];
