fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2025-12-24 21:50:12 +01:00
Code Issues Projects Releases Packages Wiki Activity Actions 1

radv: Move buffer related NIR meta shaders to radv_meta_nir.c

Browse source 
Reviewed-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/33494>
This commit is contained in:
Timur Kristóf 2025-02-11 13:10:30 +01:00 • committed by Marge Bot
parent 85b5eec159
commit c8842d19ed
4 changed files with 62 additions and 49 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

1
src/amd/vulkan/meson.build
View file

@ -62,6 +62,7 @@ libradv_files = files(
'meta/radv_meta_resolve.c',
'meta/radv_meta_resolve_cs.c',
'meta/radv_meta_resolve_fs.c',
'nir/radv_meta_nir.c',
'nir/radv_nir.h',
'nir/radv_nir_apply_pipeline_layout.c',
'nir/radv_nir_export_multiview.c',

3
src/amd/vulkan/meta/radv_meta.h
View file

@ -275,6 +275,9 @@ nir_def *get_global_ids(nir_builder *b, unsigned num_components);
void radv_break_on_count(nir_builder *b, nir_variable *var, nir_def *count);
nir_shader *radv_meta_nir_build_buffer_fill_shader(struct radv_device *dev);
nir_shader *radv_meta_nir_build_buffer_copy_shader(struct radv_device *dev);
uint32_t radv_fill_buffer(struct radv_cmd_buffer *cmd_buffer, const struct radv_image *image,
struct radeon_winsys_bo *bo, uint64_t va, uint64_t size, uint32_t value);

51
src/amd/vulkan/meta/radv_meta_buffer.c
View file

@ -1,35 +1,11 @@
#include "nir/nir_builder.h"
#include "radv_cp_dma.h"
#include "radv_debug.h"
#include "radv_meta.h"
#include "radv_sdma.h"
#include "radv_cs.h"
#include "sid.h"
#include "vk_common_entrypoints.h"
static nir_shader *
build_buffer_fill_shader(struct radv_device *dev)
{
nir_builder b = radv_meta_init_shader(dev, MESA_SHADER_COMPUTE, "meta_buffer_fill");
b.shader->info.workgroup_size[0] = 64;
nir_def *pconst = nir_load_push_constant(&b, 4, 32, nir_imm_int(&b, 0), .range = 16);
nir_def *buffer_addr = nir_pack_64_2x32(&b, nir_channels(&b, pconst, 0b0011));
nir_def *max_offset = nir_channel(&b, pconst, 2);
nir_def *data = nir_swizzle(&b, nir_channel(&b, pconst, 3), (unsigned[]){0, 0, 0, 0}, 4);
nir_def *global_id =
nir_iadd(&b, nir_imul_imm(&b, nir_channel(&b, nir_load_workgroup_id(&b), 0), b.shader->info.workgroup_size[0]),
nir_load_local_invocation_index(&b));
nir_def *offset = nir_imin(&b, nir_imul_imm(&b, global_id, 16), max_offset);
nir_def *dst_addr = nir_iadd(&b, buffer_addr, nir_u2u64(&b, offset));
nir_build_store_global(&b, data, dst_addr, .align_mul = 4);
return b.shader;
}
struct fill_constants {
uint64_t addr;
uint32_t max_offset;
@ -58,7 +34,7 @@ get_fill_pipeline(struct radv_device *device, VkPipeline *pipeline_out, VkPipeli
return VK_SUCCESS;
}
nir_shader *cs = build_buffer_fill_shader(device);
nir_shader *cs = radv_meta_nir_build_buffer_fill_shader(device);
const VkPipelineShaderStageCreateInfo stage_info = {
.sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
@ -82,29 +58,6 @@ get_fill_pipeline(struct radv_device *device, VkPipeline *pipeline_out, VkPipeli
return result;
}
static nir_shader *
build_buffer_copy_shader(struct radv_device *dev)
{
nir_builder b = radv_meta_init_shader(dev, MESA_SHADER_COMPUTE, "meta_buffer_copy");
b.shader->info.workgroup_size[0] = 64;
nir_def *pconst = nir_load_push_constant(&b, 4, 32, nir_imm_int(&b, 0), .range = 16);
nir_def *max_offset = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 0), .base = 16, .range = 4);
nir_def *src_addr = nir_pack_64_2x32(&b, nir_channels(&b, pconst, 0b0011));
nir_def *dst_addr = nir_pack_64_2x32(&b, nir_channels(&b, pconst, 0b1100));
nir_def *global_id =
nir_iadd(&b, nir_imul_imm(&b, nir_channel(&b, nir_load_workgroup_id(&b), 0), b.shader->info.workgroup_size[0]),
nir_load_local_invocation_index(&b));
nir_def *offset = nir_u2u64(&b, nir_imin(&b, nir_imul_imm(&b, global_id, 16), max_offset));
nir_def *data = nir_build_load_global(&b, 4, 32, nir_iadd(&b, src_addr, offset), .align_mul = 4);
nir_build_store_global(&b, data, nir_iadd(&b, dst_addr, offset), .align_mul = 4);
return b.shader;
}
struct copy_constants {
uint64_t src_addr;
uint64_t dst_addr;
@ -133,7 +86,7 @@ get_copy_pipeline(struct radv_device *device, VkPipeline *pipeline_out, VkPipeli
return VK_SUCCESS;
}
nir_shader *cs = build_buffer_copy_shader(device);
nir_shader *cs = radv_meta_nir_build_buffer_copy_shader(device);
const VkPipelineShaderStageCreateInfo stage_info = {
.sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,

56
src/amd/vulkan/nir/radv_meta_nir.c Normal file
View file

@ -0,0 +1,56 @@
/* Based on anv:
* Copyright © 2015 Intel Corporation
*
* Copyright © 2016 Red Hat Inc.
* Copyright © 2018 Valve Corporation
*
* SPDX-License-Identifier: MIT
*/
#include "../meta/radv_meta.h"
#include "nir_builder.h"
nir_shader *
radv_meta_nir_build_buffer_fill_shader(struct radv_device *dev)
{
nir_builder b = radv_meta_init_shader(dev, MESA_SHADER_COMPUTE, "meta_buffer_fill");
b.shader->info.workgroup_size[0] = 64;
nir_def *pconst = nir_load_push_constant(&b, 4, 32, nir_imm_int(&b, 0), .range = 16);
nir_def *buffer_addr = nir_pack_64_2x32(&b, nir_channels(&b, pconst, 0b0011));
nir_def *max_offset = nir_channel(&b, pconst, 2);
nir_def *data = nir_swizzle(&b, nir_channel(&b, pconst, 3), (unsigned[]){0, 0, 0, 0}, 4);
nir_def *global_id =
nir_iadd(&b, nir_imul_imm(&b, nir_channel(&b, nir_load_workgroup_id(&b), 0), b.shader->info.workgroup_size[0]),
nir_load_local_invocation_index(&b));
nir_def *offset = nir_imin(&b, nir_imul_imm(&b, global_id, 16), max_offset);
nir_def *dst_addr = nir_iadd(&b, buffer_addr, nir_u2u64(&b, offset));
nir_build_store_global(&b, data, dst_addr, .align_mul = 4);
return b.shader;
}
nir_shader *
radv_meta_nir_build_buffer_copy_shader(struct radv_device *dev)
{
nir_builder b = radv_meta_init_shader(dev, MESA_SHADER_COMPUTE, "meta_buffer_copy");
b.shader->info.workgroup_size[0] = 64;
nir_def *pconst = nir_load_push_constant(&b, 4, 32, nir_imm_int(&b, 0), .range = 16);
nir_def *max_offset = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 0), .base = 16, .range = 4);
nir_def *src_addr = nir_pack_64_2x32(&b, nir_channels(&b, pconst, 0b0011));
nir_def *dst_addr = nir_pack_64_2x32(&b, nir_channels(&b, pconst, 0b1100));
nir_def *global_id =
nir_iadd(&b, nir_imul_imm(&b, nir_channel(&b, nir_load_workgroup_id(&b), 0), b.shader->info.workgroup_size[0]),
nir_load_local_invocation_index(&b));
nir_def *offset = nir_u2u64(&b, nir_imin(&b, nir_imul_imm(&b, global_id, 16), max_offset));
nir_def *data = nir_build_load_global(&b, 4, 32, nir_iadd(&b, src_addr, offset), .align_mul = 4);
nir_build_store_global(&b, data, nir_iadd(&b, dst_addr, offset), .align_mul = 4);
return b.shader;
}