KonstantinSeurer
/
mesa
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

panvk: Add support for storage image 

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>
This commit is contained in:
Boris Brezillon 2021-09-23 16:27:06 +02:00 committed by Marge Bot
parent eca0a0e29e
commit 1056b3e39e
9 changed files with 284 additions and 49 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

10
src/panfrost/vulkan/panvk_cmd_buffer.c
View File

@ -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);

28
src/panfrost/vulkan/panvk_descriptor_set.c
View File

@ -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);
}

19
src/panfrost/vulkan/panvk_private.h
View File

@ -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;
};

134
src/panfrost/vulkan/panvk_vX_cmd_buffer.c
View File

@ -661,51 +661,143 @@ panvk_draw_prepare_varyings(struct panvk_cmd_buffer *cmdbuf,
}
static void
panvk_draw_prepare_attributes(struct panvk_cmd_buffer *cmdbuf,
struct panvk_draw_info *draw)
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)
{
const struct panvk_pipeline *pipeline = panvk_cmd_get_pipeline(cmdbuf, GRAPHICS);
struct panvk_descriptor_state *desc_state = &bind_point_state->desc_state;
const struct panvk_pipeline *pipeline = bind_point_state->pipeline;
/* TODO: images */
if (!pipeline->attribs.buf_count)
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;
if (cmdbuf->state.vb.attribs) {
draw->stages[MESA_SHADER_VERTEX].attributes = cmdbuf->state.vb.attribs;
draw->attribute_bufs = cmdbuf->state.vb.attrib_bufs;
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)
{
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;
if (desc_state->vs_attribs || !attrib_count)
return;
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

4
src/panfrost/vulkan/panvk_vX_cs.c
View File

@ -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;

51
src/panfrost/vulkan/panvk_vX_descriptor_set.c
View File

@ -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:

35
src/panfrost/vulkan/panvk_vX_image.c
View File

@ -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;
}

3
src/panfrost/vulkan/panvk_vX_pipeline.c
View File

@ -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;

49
src/panfrost/vulkan/panvk_vX_shader.c
View File

@ -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];