fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-05-07 04:58:05 +02:00
Code Issues Projects Releases Packages Wiki Activity Actions 4

radv/rt: declare shader arguments for resource/sampler heaps

Browse source 
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/39483>
This commit is contained in:
Samuel Pitoiset 2025-06-26 12:47:27 +02:00 • committed by Marge Bot
parent cdfb9a24ba
commit cebac5a427
14 changed files with 146 additions and 52 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

2
src/amd/compiler/aco_nir_call_attribs.h
View file

@ -43,6 +43,8 @@ enum aco_nir_rt_function_arg {
RT_ARG_LAUNCH_SIZE,
RT_ARG_DESCRIPTORS,
RT_ARG_DYNAMIC_DESCRIPTORS,
RT_ARG_HEAP_RESOURCE = RT_ARG_DESCRIPTORS,
RT_ARG_HEAP_SAMPLER = RT_ARG_DYNAMIC_DESCRIPTORS,
RT_ARG_PUSH_CONSTANTS,
RT_ARG_SBT_DESCRIPTORS,
RT_ARG_COUNT,

1
src/amd/compiler/aco_shader_info.h
View file

@ -114,6 +114,7 @@ struct aco_shader_info {
unsigned workgroup_size;
unsigned lds_size;
bool merged_shader_compiled_separately; /* GFX9+ */
bool descriptor_heap;
struct ac_arg next_stage_pc;
struct ac_arg epilog_pc; /* Vulkan only */
struct {

3
src/amd/compiler/instruction_selection/aco_instruction_selection.h
View file

@ -306,7 +306,8 @@ void finish_program(isel_context* ctx);
ABI nir_abi_to_aco(unsigned nir_abi_mask);
param_assignment_hints get_ahit_isec_param_hints(const struct callee_info& traversal_info);
param_assignment_hints get_ahit_isec_param_hints(const struct callee_info& traversal_info,
bool uses_descriptor_heap);
struct callee_info get_callee_info(amd_gfx_level gfx_level, unsigned wave_size, const ABI& abi,
unsigned param_count, const nir_parameter* parameters,

11
src/amd/compiler/instruction_selection/aco_isel_helpers.cpp
View file

@ -902,7 +902,7 @@ param_hint_map(param_assignment_hints& hints, const struct callee_info& traversa
}
param_assignment_hints
get_ahit_isec_param_hints(const struct callee_info& traversal_info)
get_ahit_isec_param_hints(const struct callee_info& traversal_info, bool uses_descriptor_heap)
{
param_assignment_hints hints;
hints.stack_pointer_affinity = traversal_info.stack_ptr;
@ -914,8 +914,13 @@ get_ahit_isec_param_hints(const struct callee_info& traversal_info)
param_hint_map(hints, traversal_info, RT_ARG_LAUNCH_ID, RT_ARG_LAUNCH_ID);
param_hint_map(hints, traversal_info, RT_ARG_LAUNCH_SIZE, RT_ARG_LAUNCH_SIZE);
param_hint_map(hints, traversal_info, RT_ARG_DESCRIPTORS, RT_ARG_DESCRIPTORS);
param_hint_map(hints, traversal_info, RT_ARG_DYNAMIC_DESCRIPTORS, RT_ARG_DYNAMIC_DESCRIPTORS);
if (uses_descriptor_heap) {
param_hint_map(hints, traversal_info, RT_ARG_HEAP_RESOURCE, RT_ARG_HEAP_RESOURCE);
param_hint_map(hints, traversal_info, RT_ARG_HEAP_SAMPLER, RT_ARG_HEAP_SAMPLER);
} else {
param_hint_map(hints, traversal_info, RT_ARG_DESCRIPTORS, RT_ARG_DESCRIPTORS);
param_hint_map(hints, traversal_info, RT_ARG_DYNAMIC_DESCRIPTORS, RT_ARG_DYNAMIC_DESCRIPTORS);
}
param_hint_map(hints, traversal_info, RT_ARG_PUSH_CONSTANTS, RT_ARG_PUSH_CONSTANTS);
param_hint_map(hints, traversal_info, RT_ARG_SBT_DESCRIPTORS, RT_ARG_SBT_DESCRIPTORS);
param_hint_map(hints, traversal_info, AHIT_ISEC_ARG_SHADER_RECORD_PTR,

5
src/amd/compiler/instruction_selection/aco_select_nir.cpp
View file

@ -754,7 +754,7 @@ visit_call(isel_context* ctx, nir_call_instr* instr)
param_assignment_hints hints;
if (nir_abi == ACO_NIR_CALL_ABI_AHIT_ISEC)
hints = get_ahit_isec_param_hints(ctx->callee_info);
hints = get_ahit_isec_param_hints(ctx->callee_info, ctx->program->info.descriptor_heap);
ABI abi = nir_abi_to_aco(instr->callee->driver_attributes);
@ -1390,7 +1390,8 @@ select_program_rt(isel_context& ctx, unsigned shader_count, struct nir_shader* c
callee_info traversal_info = get_callee_info(
ctx.program->gfx_level, ctx.program->wave_size, rtTraversalABI,
traversal_function->num_params, traversal_function->params, NULL, limit);
callee_hints = get_ahit_isec_param_hints(traversal_info);
callee_hints =
get_ahit_isec_param_hints(traversal_info, ctx.program->info.descriptor_heap);
}
/* TODO: callable abi? */

25
src/amd/vulkan/nir/radv_nir_rt_stage_common.c
View file

@ -257,7 +257,7 @@ radv_nir_return_param_from_type(nir_parameter *param, const glsl_type *type, boo
}
void
radv_build_rt_prolog(struct radv_device *device, struct radv_shader_stage *stage)
radv_build_rt_prolog(struct radv_device *device, struct radv_shader_stage *stage, bool uses_descriptor_heap)
{
const struct radv_physical_device *pdev = radv_device_physical(device);
@ -268,6 +268,7 @@ radv_build_rt_prolog(struct radv_device *device, struct radv_shader_stage *stage
stage->info.loads_push_constants = true;
stage->info.loads_dynamic_offsets = true;
stage->info.force_indirect_descriptors = true;
stage->info.descriptor_heap = uses_descriptor_heap;
stage->info.wave_size = pdev->rt_wave_size;
stage->info.workgroup_size = stage->info.wave_size;
stage->info.user_data_0 = R_00B900_COMPUTE_USER_DATA_0;
@ -288,11 +289,18 @@ radv_build_rt_prolog(struct radv_device *device, struct radv_shader_stage *stage
b.shader->info.min_subgroup_size = pdev->rt_wave_size;
nir_function *raygen_function = nir_function_create(b.shader, "raygen_func");
radv_nir_init_rt_function_params(raygen_function, MESA_SHADER_RAYGEN, 0, 0);
radv_nir_init_rt_function_params(raygen_function, MESA_SHADER_RAYGEN, 0, 0, uses_descriptor_heap);
nir_def *descriptors, *dynamic_descriptors, *heap_resource, *heap_sampler;
if (uses_descriptor_heap) {
heap_resource = ac_nir_load_arg(&b, &stage->args.ac, stage->args.descriptors[RADV_HEAP_RESOURCE]);
heap_sampler = ac_nir_load_arg(&b, &stage->args.ac, stage->args.descriptors[RADV_HEAP_SAMPLER]);
} else {
descriptors = ac_nir_load_arg(&b, &stage->args.ac, stage->args.descriptors[0]);
dynamic_descriptors = ac_nir_load_arg(&b, &stage->args.ac, stage->args.ac.dynamic_descriptors);
}
nir_def *descriptors = ac_nir_load_arg(&b, &stage->args.ac, stage->args.descriptors[0]);
nir_def *push_constants = ac_nir_load_arg(&b, &stage->args.ac, stage->args.ac.push_constants);
nir_def *dynamic_descriptors = ac_nir_load_arg(&b, &stage->args.ac, stage->args.ac.dynamic_descriptors);
nir_def *sbt_desc = nir_pack_64_2x32(&b, ac_nir_load_arg(&b, &stage->args.ac, stage->args.ac.rt.sbt_descriptors));
nir_def *launch_size_addr = nir_pack_64_2x32(&b, ac_nir_load_arg(&b, &stage->args.ac, stage->args.ac.rt.launch_size_addr));
nir_def *traversal_addr =
@ -399,8 +407,13 @@ radv_build_rt_prolog(struct radv_device *device, struct radv_shader_stage *stage
nir_def *params[RAYGEN_ARG_COUNT];
params[RT_ARG_LAUNCH_ID] = nir_vec3(&b, id_x, id_y, wg_ids[2]);
params[RT_ARG_LAUNCH_SIZE] = launch_sizes;
params[RT_ARG_DESCRIPTORS] = descriptors;
params[RT_ARG_DYNAMIC_DESCRIPTORS] = dynamic_descriptors;
if (uses_descriptor_heap) {
params[RT_ARG_HEAP_RESOURCE] = heap_resource;
params[RT_ARG_HEAP_SAMPLER] = heap_sampler;
} else {
params[RT_ARG_DESCRIPTORS] = descriptors;
params[RT_ARG_DYNAMIC_DESCRIPTORS] = dynamic_descriptors;
}
params[RT_ARG_PUSH_CONSTANTS] = push_constants;
params[RT_ARG_SBT_DESCRIPTORS] = sbt_desc;
params[RAYGEN_ARG_SHADER_RECORD_PTR] = shader_record_ptr;

3
src/amd/vulkan/nir/radv_nir_rt_stage_common.h
View file

@ -159,5 +159,6 @@ struct radv_nir_rt_traversal_result radv_build_traversal(struct radv_device *dev
struct radv_nir_rt_traversal_params *params,
struct radv_ray_tracing_stage_info *info);
void radv_build_rt_prolog(struct radv_device *device, struct radv_shader_stage *stage);
void radv_build_rt_prolog(struct radv_device *device, struct radv_shader_stage *stage, bool uses_descriptor_heap);
#endif // MESA_RADV_NIR_RT_STAGE_COMMON_H

27
src/amd/vulkan/nir/radv_nir_rt_stage_cps.c
View file

@ -507,15 +507,20 @@ radv_nir_lower_rt_io_cps(nir_shader *nir)
}
static void
init_cps_function(nir_function *function, bool has_position_fetch)
init_cps_function(nir_function *function, bool has_position_fetch, bool uses_descriptor_heap)
{
function->num_params = has_position_fetch ? CPS_ARG_COUNT : CPS_ARG_COUNT - 1;
function->params = rzalloc_array_size(function->shader, sizeof(nir_parameter), function->num_params);
radv_nir_param_from_type(function->params + RT_ARG_LAUNCH_ID, glsl_vector_type(GLSL_TYPE_UINT, 3), false, 0);
radv_nir_param_from_type(function->params + RT_ARG_LAUNCH_SIZE, glsl_vector_type(GLSL_TYPE_UINT, 3), true, 0);
radv_nir_param_from_type(function->params + RT_ARG_DESCRIPTORS, glsl_uint_type(), true, 0);
radv_nir_param_from_type(function->params + RT_ARG_DYNAMIC_DESCRIPTORS, glsl_uint_type(), true, 0);
if (uses_descriptor_heap) {
radv_nir_param_from_type(function->params + RT_ARG_HEAP_RESOURCE, glsl_uint_type(), true, 0);
radv_nir_param_from_type(function->params + RT_ARG_HEAP_SAMPLER, glsl_uint_type(), true, 0);
} else {
radv_nir_param_from_type(function->params + RT_ARG_DESCRIPTORS, glsl_uint_type(), true, 0);
radv_nir_param_from_type(function->params + RT_ARG_DYNAMIC_DESCRIPTORS, glsl_uint_type(), true, 0);
}
radv_nir_param_from_type(function->params + RT_ARG_PUSH_CONSTANTS, glsl_uint_type(), true, 0);
radv_nir_param_from_type(function->params + RT_ARG_SBT_DESCRIPTORS, glsl_uint64_t_type(), true, 0);
radv_nir_param_from_type(function->params + RAYGEN_ARG_TRAVERSAL_ADDR, glsl_uint64_t_type(), true, 0);
@ -552,15 +557,16 @@ radv_nir_lower_rt_abi_cps(nir_shader *shader, const struct radv_shader_info *inf
struct radv_device *device, struct radv_ray_tracing_pipeline *pipeline,
bool has_position_fetch, const struct radv_ray_tracing_stage_info *traversal_info)
{
const bool uses_descriptor_heap = pipeline->base.base.create_flags & VK_PIPELINE_CREATE_2_DESCRIPTOR_HEAP_BIT_EXT;
nir_function_impl *impl = nir_shader_get_entrypoint(shader);
/* The first raygen shader gets called by the prolog with the standard raygen signature. Only shaders called by the
* first shader can use the CPS function signature.
*/
if (shader->info.stage != MESA_SHADER_RAYGEN || resume_shader)
init_cps_function(impl->function, has_position_fetch);
init_cps_function(impl->function, has_position_fetch, uses_descriptor_heap);
else
radv_nir_init_rt_function_params(impl->function, MESA_SHADER_RAYGEN, 0, 0);
radv_nir_init_rt_function_params(impl->function, MESA_SHADER_RAYGEN, 0, 0, uses_descriptor_heap);
if (traversal_info) {
unsigned idx;
@ -624,14 +630,19 @@ radv_nir_lower_rt_abi_cps(nir_shader *shader, const struct radv_shader_info *inf
/* tail-call next shader */
nir_def *shader_addr = nir_load_var(&b, vars.shader_addr);
nir_function *continuation_func = nir_function_create(shader, "continuation_func");
init_cps_function(continuation_func, has_position_fetch);
init_cps_function(continuation_func, has_position_fetch, uses_descriptor_heap);
unsigned param_count = continuation_func->num_params;
nir_def **next_args = rzalloc_array_size(b.shader, sizeof(nir_def *), param_count);
next_args[RT_ARG_LAUNCH_ID] = nir_load_param(&b, RT_ARG_LAUNCH_ID);
next_args[RT_ARG_LAUNCH_SIZE] = nir_load_param(&b, RT_ARG_LAUNCH_SIZE);
next_args[RT_ARG_DESCRIPTORS] = nir_load_param(&b, RT_ARG_DESCRIPTORS);
next_args[RT_ARG_DYNAMIC_DESCRIPTORS] = nir_load_param(&b, RT_ARG_DYNAMIC_DESCRIPTORS);
if (uses_descriptor_heap) {
next_args[RT_ARG_HEAP_RESOURCE] = nir_load_param(&b, RT_ARG_HEAP_RESOURCE);
next_args[RT_ARG_HEAP_SAMPLER] = nir_load_param(&b, RT_ARG_HEAP_SAMPLER);
} else {
next_args[RT_ARG_DESCRIPTORS] = nir_load_param(&b, RT_ARG_DESCRIPTORS);
next_args[RT_ARG_DYNAMIC_DESCRIPTORS] = nir_load_param(&b, RT_ARG_DYNAMIC_DESCRIPTORS);
}
next_args[RT_ARG_PUSH_CONSTANTS] = nir_load_param(&b, RT_ARG_PUSH_CONSTANTS);
next_args[RT_ARG_SBT_DESCRIPTORS] = nir_load_param(&b, RT_ARG_SBT_DESCRIPTORS);
next_args[RAYGEN_ARG_TRAVERSAL_ADDR] = nir_load_var(&b, vars.traversal_addr);

87
src/amd/vulkan/nir/radv_nir_rt_stage_functions.c
View file

@ -22,22 +22,27 @@
#include "vk_pipeline.h"
static void
radv_nir_init_common_rt_params(nir_function *function)
radv_nir_init_common_rt_params(nir_function *function, bool uses_descriptor_heap)
{
radv_nir_param_from_type(function->params + RT_ARG_LAUNCH_ID, glsl_vector_type(GLSL_TYPE_UINT, 3), false, 0);
radv_nir_param_from_type(function->params + RT_ARG_LAUNCH_SIZE, glsl_vector_type(GLSL_TYPE_UINT, 3), true, 0);
radv_nir_param_from_type(function->params + RT_ARG_DESCRIPTORS, glsl_uint_type(), true, 0);
radv_nir_param_from_type(function->params + RT_ARG_DYNAMIC_DESCRIPTORS, glsl_uint_type(), true, 0);
if (uses_descriptor_heap) {
radv_nir_param_from_type(function->params + RT_ARG_HEAP_RESOURCE, glsl_uint_type(), true, 0);
radv_nir_param_from_type(function->params + RT_ARG_HEAP_SAMPLER, glsl_uint_type(), true, 0);
} else {
radv_nir_param_from_type(function->params + RT_ARG_DESCRIPTORS, glsl_uint_type(), true, 0);
radv_nir_param_from_type(function->params + RT_ARG_DYNAMIC_DESCRIPTORS, glsl_uint_type(), true, 0);
}
radv_nir_param_from_type(function->params + RT_ARG_PUSH_CONSTANTS, glsl_uint_type(), true, 0);
radv_nir_param_from_type(function->params + RT_ARG_SBT_DESCRIPTORS, glsl_uint64_t_type(), true, 0);
}
static void
radv_nir_init_traversal_params(nir_function *function, unsigned payload_size)
radv_nir_init_traversal_params(nir_function *function, unsigned payload_size, bool uses_descriptor_heap)
{
function->num_params = TRAVERSAL_ARG_PAYLOAD_BASE + DIV_ROUND_UP(payload_size, 4);
function->params = rzalloc_array_size(function->shader, sizeof(nir_parameter), function->num_params);
radv_nir_init_common_rt_params(function);
radv_nir_init_common_rt_params(function, uses_descriptor_heap);
radv_nir_param_from_type(function->params + TRAVERSAL_ARG_TRAVERSAL_ADDR, glsl_uint64_t_type(), true, 0);
radv_nir_param_from_type(function->params + TRAVERSAL_ARG_SHADER_RECORD_PTR, glsl_uint64_t_type(), false, ACO_NIR_PARAM_ATTRIB_DISCARDABLE);
radv_nir_param_from_type(function->params + TRAVERSAL_ARG_ACCEL_STRUCT, glsl_uint64_t_type(), false, 0);
@ -68,7 +73,7 @@ radv_nir_init_traversal_params(nir_function *function, unsigned payload_size)
void
radv_nir_init_rt_function_params(nir_function *function, mesa_shader_stage stage, unsigned payload_size,
unsigned hit_attrib_size)
unsigned hit_attrib_size, bool uses_descriptor_heap)
{
unsigned payload_base = -1u;
@ -76,7 +81,7 @@ radv_nir_init_rt_function_params(nir_function *function, mesa_shader_stage stage
case MESA_SHADER_RAYGEN:
function->num_params = RAYGEN_ARG_COUNT;
function->params = rzalloc_array_size(function->shader, sizeof(nir_parameter), function->num_params);
radv_nir_init_common_rt_params(function);
radv_nir_init_common_rt_params(function, uses_descriptor_heap);
radv_nir_param_from_type(function->params + RAYGEN_ARG_TRAVERSAL_ADDR, glsl_uint64_t_type(), true, 0);
radv_nir_param_from_type(function->params + RAYGEN_ARG_SHADER_RECORD_PTR, glsl_uint64_t_type(), false, 0);
function->driver_attributes = (uint32_t)ACO_NIR_CALL_ABI_RT_RECURSIVE | ACO_NIR_FUNCTION_ATTRIB_NORETURN;
@ -84,7 +89,7 @@ radv_nir_init_rt_function_params(nir_function *function, mesa_shader_stage stage
case MESA_SHADER_CALLABLE:
function->num_params = RAYGEN_ARG_COUNT + DIV_ROUND_UP(payload_size, 4);
function->params = rzalloc_array_size(function->shader, sizeof(nir_parameter), function->num_params);
radv_nir_init_common_rt_params(function);
radv_nir_init_common_rt_params(function, uses_descriptor_heap);
radv_nir_param_from_type(function->params + RAYGEN_ARG_TRAVERSAL_ADDR, glsl_uint64_t_type(), true, 0);
radv_nir_param_from_type(function->params + RAYGEN_ARG_SHADER_RECORD_PTR, glsl_uint64_t_type(), false, 0);
@ -96,7 +101,7 @@ radv_nir_init_rt_function_params(nir_function *function, mesa_shader_stage stage
function->num_params =
AHIT_ISEC_ARG_HIT_ATTRIB_PAYLOAD_BASE + DIV_ROUND_UP(hit_attrib_size, 4) + DIV_ROUND_UP(payload_size, 4);
function->params = rzalloc_array_size(function->shader, sizeof(nir_parameter), function->num_params);
radv_nir_init_common_rt_params(function);
radv_nir_init_common_rt_params(function, uses_descriptor_heap);
radv_nir_param_from_type(function->params + AHIT_ISEC_ARG_SHADER_RECORD_PTR, glsl_uint64_t_type(), false, 0);
radv_nir_param_from_type(function->params + AHIT_ISEC_ARG_CULL_MASK_AND_FLAGS, glsl_uint_type(), false, 0);
radv_nir_param_from_type(function->params + AHIT_ISEC_ARG_SBT_INDEX, glsl_uint_type(), false, 0);
@ -126,7 +131,7 @@ radv_nir_init_rt_function_params(nir_function *function, mesa_shader_stage stage
case MESA_SHADER_MISS:
function->num_params = CHIT_MISS_ARG_PAYLOAD_BASE + DIV_ROUND_UP(payload_size, 4);
function->params = rzalloc_array_size(function->shader, sizeof(nir_parameter), function->num_params);
radv_nir_init_common_rt_params(function);
radv_nir_init_common_rt_params(function, uses_descriptor_heap);
radv_nir_param_from_type(function->params + CHIT_MISS_ARG_TRAVERSAL_ADDR, glsl_uint64_t_type(), true, 0);
radv_nir_param_from_type(function->params + CHIT_MISS_ARG_SHADER_RECORD_PTR, glsl_uint64_t_type(), false, 0);
radv_nir_param_from_type(function->params + CHIT_MISS_ARG_ACCEL_STRUCT, glsl_uint64_t_type(), false, 0);
@ -212,6 +217,7 @@ static struct rt_variables
create_rt_variables(nir_shader *shader, struct radv_device *device, const VkPipelineCreateFlags2 flags,
unsigned max_payload_size, unsigned max_hit_attrib_size)
{
const bool uses_descriptor_heap = flags & VK_PIPELINE_CREATE_2_DESCRIPTOR_HEAP_BIT_EXT;
struct rt_variables vars = {
.device = device,
.flags = flags,
@ -227,16 +233,19 @@ create_rt_variables(nir_shader *shader, struct radv_device *device, const VkPipe
}
nir_function *trace_ray_func = nir_function_create(shader, "trace_ray_func");
radv_nir_init_traversal_params(trace_ray_func, max_payload_size);
radv_nir_init_traversal_params(trace_ray_func, max_payload_size, uses_descriptor_heap);
vars.trace_ray_func = trace_ray_func;
nir_function *ahit_isec_func = nir_function_create(shader, "ahit_isec_func");
radv_nir_init_rt_function_params(ahit_isec_func, MESA_SHADER_ANY_HIT, max_payload_size, max_hit_attrib_size);
radv_nir_init_rt_function_params(ahit_isec_func, MESA_SHADER_ANY_HIT, max_payload_size, max_hit_attrib_size,
uses_descriptor_heap);
vars.ahit_isec_func = ahit_isec_func;
nir_function *chit_miss_func = nir_function_create(shader, "chit_miss_func");
radv_nir_init_rt_function_params(chit_miss_func, MESA_SHADER_CLOSEST_HIT, max_payload_size, max_hit_attrib_size);
radv_nir_init_rt_function_params(chit_miss_func, MESA_SHADER_CLOSEST_HIT, max_payload_size, max_hit_attrib_size,
uses_descriptor_heap);
vars.chit_miss_func = chit_miss_func;
nir_function *callable_func = nir_function_create(shader, "callable_func");
radv_nir_init_rt_function_params(callable_func, MESA_SHADER_CALLABLE, max_payload_size, max_hit_attrib_size);
radv_nir_init_rt_function_params(callable_func, MESA_SHADER_CALLABLE, max_payload_size, max_hit_attrib_size,
uses_descriptor_heap);
vars.callable_func = callable_func;
vars.shader_record_ptr_param = -1u;
@ -339,6 +348,7 @@ lower_rt_instruction(nir_builder *b, nir_instr *instr, void *_vars)
nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
struct rt_variables *vars = _vars;
const bool uses_descriptor_heap = vars->flags & VK_PIPELINE_CREATE_2_DESCRIPTOR_HEAP_BIT_EXT;
b->cursor = nir_before_instr(&intr->instr);
@ -352,8 +362,13 @@ lower_rt_instruction(nir_builder *b, nir_instr *instr, void *_vars)
nir_def **args = rzalloc_array_size(b->shader, sizeof(nir_def *), param_count);
args[RT_ARG_LAUNCH_ID] = nir_load_param(b, RT_ARG_LAUNCH_ID);
args[RT_ARG_LAUNCH_SIZE] = nir_load_param(b, RT_ARG_LAUNCH_SIZE);
args[RT_ARG_DESCRIPTORS] = nir_load_param(b, RT_ARG_DESCRIPTORS);
args[RT_ARG_DYNAMIC_DESCRIPTORS] = nir_load_param(b, RT_ARG_DYNAMIC_DESCRIPTORS);
if (uses_descriptor_heap) {
args[RT_ARG_HEAP_RESOURCE] = nir_load_param(b, RT_ARG_HEAP_RESOURCE);
args[RT_ARG_HEAP_SAMPLER] = nir_load_param(b, RT_ARG_HEAP_SAMPLER);
} else {
args[RT_ARG_DESCRIPTORS] = nir_load_param(b, RT_ARG_DESCRIPTORS);
args[RT_ARG_DYNAMIC_DESCRIPTORS] = nir_load_param(b, RT_ARG_DYNAMIC_DESCRIPTORS);
}
args[RT_ARG_PUSH_CONSTANTS] = nir_load_param(b, RT_ARG_PUSH_CONSTANTS);
args[RT_ARG_SBT_DESCRIPTORS] = nir_load_param(b, RT_ARG_SBT_DESCRIPTORS);
args[RAYGEN_ARG_TRAVERSAL_ADDR] = nir_undef(b, 1, 64);
@ -375,7 +390,13 @@ lower_rt_instruction(nir_builder *b, nir_instr *instr, void *_vars)
args[RT_ARG_LAUNCH_ID] = nir_load_param(b, RT_ARG_LAUNCH_ID);
args[RT_ARG_LAUNCH_SIZE] = nir_load_param(b, RT_ARG_LAUNCH_SIZE);
args[RT_ARG_DESCRIPTORS] = nir_load_param(b, RT_ARG_DESCRIPTORS);
args[RT_ARG_DYNAMIC_DESCRIPTORS] = nir_load_param(b, RT_ARG_DYNAMIC_DESCRIPTORS);
if (uses_descriptor_heap) {
args[RT_ARG_HEAP_RESOURCE] = nir_load_param(b, RT_ARG_HEAP_RESOURCE);
args[RT_ARG_HEAP_SAMPLER] = nir_load_param(b, RT_ARG_HEAP_SAMPLER);
} else {
args[RT_ARG_DESCRIPTORS] = nir_load_param(b, RT_ARG_DESCRIPTORS);
args[RT_ARG_DYNAMIC_DESCRIPTORS] = nir_load_param(b, RT_ARG_DYNAMIC_DESCRIPTORS);
}
args[RT_ARG_PUSH_CONSTANTS] = nir_load_param(b, RT_ARG_PUSH_CONSTANTS);
args[RT_ARG_SBT_DESCRIPTORS] = nir_load_param(b, RT_ARG_SBT_DESCRIPTORS);
args[TRAVERSAL_ARG_TRAVERSAL_ADDR] = traversal_addr;
@ -503,6 +524,14 @@ lower_rt_instruction(nir_builder *b, nir_instr *instr, void *_vars)
ret = nir_load_param(b, RT_ARG_DESCRIPTORS);
break;
}
case nir_intrinsic_load_rt_heap_resource_amd: {
ret = nir_load_param(b, RT_ARG_HEAP_RESOURCE);
break;
}
case nir_intrinsic_load_rt_heap_sampler_amd: {
ret = nir_load_param(b, RT_ARG_HEAP_SAMPLER);
break;
}
case nir_intrinsic_load_rt_dynamic_descriptors_amd: {
ret = nir_load_param(b, RT_ARG_DYNAMIC_DESCRIPTORS);
break;
@ -553,8 +582,13 @@ lower_rt_instruction(nir_builder *b, nir_instr *instr, void *_vars)
nir_def **args = rzalloc_array_size(b->shader, sizeof(nir_def *), param_count);
args[RT_ARG_LAUNCH_ID] = nir_load_param(b, RT_ARG_LAUNCH_ID);
args[RT_ARG_LAUNCH_SIZE] = nir_load_param(b, RT_ARG_LAUNCH_SIZE);
args[RT_ARG_DESCRIPTORS] = nir_load_param(b, RT_ARG_DESCRIPTORS);
args[RT_ARG_DYNAMIC_DESCRIPTORS] = nir_load_param(b, RT_ARG_DYNAMIC_DESCRIPTORS);
if (uses_descriptor_heap) {
args[RT_ARG_HEAP_RESOURCE] = nir_load_param(b, RT_ARG_HEAP_RESOURCE);
args[RT_ARG_HEAP_SAMPLER] = nir_load_param(b, RT_ARG_HEAP_SAMPLER);
} else {
args[RT_ARG_DESCRIPTORS] = nir_load_param(b, RT_ARG_DESCRIPTORS);
args[RT_ARG_DYNAMIC_DESCRIPTORS] = nir_load_param(b, RT_ARG_DYNAMIC_DESCRIPTORS);
}
args[RT_ARG_PUSH_CONSTANTS] = nir_load_param(b, RT_ARG_PUSH_CONSTANTS);
args[RT_ARG_SBT_DESCRIPTORS] = nir_load_param(b, RT_ARG_SBT_DESCRIPTORS);
args[CHIT_MISS_ARG_TRAVERSAL_ADDR] = nir_load_param(b, vars->traversal_addr_param);
@ -598,8 +632,13 @@ lower_rt_instruction(nir_builder *b, nir_instr *instr, void *_vars)
nir_def **args = rzalloc_array_size(b->shader, sizeof(nir_def *), param_count);
args[RT_ARG_LAUNCH_ID] = nir_load_param(b, RT_ARG_LAUNCH_ID);
args[RT_ARG_LAUNCH_SIZE] = nir_load_param(b, RT_ARG_LAUNCH_SIZE);
args[RT_ARG_DESCRIPTORS] = nir_load_param(b, RT_ARG_DESCRIPTORS);
args[RT_ARG_DYNAMIC_DESCRIPTORS] = nir_load_param(b, RT_ARG_DYNAMIC_DESCRIPTORS);
if (uses_descriptor_heap) {
args[RT_ARG_HEAP_RESOURCE] = nir_load_param(b, RT_ARG_HEAP_RESOURCE);
args[RT_ARG_HEAP_SAMPLER] = nir_load_param(b, RT_ARG_HEAP_SAMPLER);
} else {
args[RT_ARG_DESCRIPTORS] = nir_load_param(b, RT_ARG_DESCRIPTORS);
args[RT_ARG_DYNAMIC_DESCRIPTORS] = nir_load_param(b, RT_ARG_DYNAMIC_DESCRIPTORS);
}
args[RT_ARG_PUSH_CONSTANTS] = nir_load_param(b, RT_ARG_PUSH_CONSTANTS);
args[RT_ARG_SBT_DESCRIPTORS] = nir_load_param(b, RT_ARG_SBT_DESCRIPTORS);
args[CHIT_MISS_ARG_TRAVERSAL_ADDR] = nir_load_param(b, vars->traversal_addr_param);
@ -816,13 +855,15 @@ radv_nir_lower_rt_abi_functions(nir_shader *shader, const struct radv_shader_inf
uint32_t hit_attrib_size, struct radv_device *device,
struct radv_ray_tracing_pipeline *pipeline)
{
const bool uses_descriptor_heap = pipeline->base.base.create_flags & VK_PIPELINE_CREATE_2_DESCRIPTOR_HEAP_BIT_EXT;
nir_function_impl *impl = nir_shader_get_entrypoint(shader);
nir_function *entrypoint_function = impl->function;
if (radv_is_traversal_shader(shader))
radv_nir_init_traversal_params(entrypoint_function, payload_size);
radv_nir_init_traversal_params(entrypoint_function, payload_size, uses_descriptor_heap);
else
radv_nir_init_rt_function_params(entrypoint_function, shader->info.stage, payload_size, hit_attrib_size);
radv_nir_init_rt_function_params(entrypoint_function, shader->info.stage, payload_size, hit_attrib_size,
uses_descriptor_heap);
struct rt_variables vars =
create_rt_variables(shader, device, pipeline->base.base.create_flags, payload_size, hit_attrib_size);

2
src/amd/vulkan/nir/radv_nir_rt_stage_functions.h
View file

@ -14,7 +14,7 @@
nir_function_impl *radv_get_rt_shader_entrypoint(nir_shader *shader);
void radv_nir_init_rt_function_params(nir_function *function, mesa_shader_stage stage, unsigned payload_size,
unsigned hit_attrib_size);
unsigned hit_attrib_size, bool uses_descriptor_heap);
void radv_nir_lower_rt_abi_functions(nir_shader *shader, const struct radv_shader_info *info, uint32_t payload_size,
uint32_t hit_attrib_size, struct radv_device *device,

3
src/amd/vulkan/nir/radv_nir_rt_stage_monolithic.c
View file

@ -458,8 +458,9 @@ void
radv_nir_lower_rt_abi_monolithic(nir_shader *shader, struct radv_device *device,
struct radv_ray_tracing_pipeline *pipeline)
{
const bool uses_descriptor_heap = pipeline->base.base.create_flags & VK_PIPELINE_CREATE_2_DESCRIPTOR_HEAP_BIT_EXT;
nir_function_impl *impl = nir_shader_get_entrypoint(shader);
radv_nir_init_rt_function_params(impl->function, MESA_SHADER_RAYGEN, 0, 0);
radv_nir_init_rt_function_params(impl->function, MESA_SHADER_RAYGEN, 0, 0, uses_descriptor_heap);
nir_builder b = nir_builder_at(nir_before_impl(impl));

25
src/amd/vulkan/nir/radv_nir_rt_traversal_shader.c
View file

@ -834,6 +834,8 @@ handle_candidate_triangle(nir_builder *b, struct radv_triangle_intersection *int
const struct radv_ray_traversal_args *args, const struct radv_ray_flags *ray_flags)
{
struct traversal_data *data = args->data;
const bool uses_descriptor_heap =
data->pipeline->base.base.create_flags & VK_PIPELINE_CREATE_2_DESCRIPTOR_HEAP_BIT_EXT;
nir_def *geometry_id = nir_iand_imm(b, intersection->base.geometry_id_and_flags, 0xfffffff);
nir_def *sbt_idx =
@ -925,8 +927,13 @@ handle_candidate_triangle(nir_builder *b, struct radv_triangle_intersection *int
nir_def **params = rzalloc_array_size(b->shader, sizeof(nir_def *), param_count);
params[RT_ARG_LAUNCH_ID] = nir_load_ray_launch_id(b);
params[RT_ARG_LAUNCH_SIZE] = nir_load_ray_launch_size(b);
params[RT_ARG_DESCRIPTORS] = nir_load_rt_descriptors_amd(b);
params[RT_ARG_DYNAMIC_DESCRIPTORS] = nir_load_rt_dynamic_descriptors_amd(b);
if (uses_descriptor_heap) {
params[RT_ARG_HEAP_RESOURCE] = nir_load_rt_heap_resource_amd(b);
params[RT_ARG_HEAP_SAMPLER] = nir_load_rt_heap_sampler_amd(b);
} else {
params[RT_ARG_DESCRIPTORS] = nir_load_rt_descriptors_amd(b);
params[RT_ARG_DYNAMIC_DESCRIPTORS] = nir_load_rt_dynamic_descriptors_amd(b);
}
params[RT_ARG_PUSH_CONSTANTS] = nir_load_rt_push_constants_amd(b);
params[RT_ARG_SBT_DESCRIPTORS] = nir_load_sbt_base_amd(b);
params[AHIT_ISEC_ARG_SHADER_RECORD_PTR] = sbt_data.shader_record_ptr;
@ -988,6 +995,8 @@ handle_candidate_aabb(nir_builder *b, struct radv_leaf_intersection *intersectio
const struct radv_ray_traversal_args *args)
{
struct traversal_data *data = args->data;
const bool uses_descriptor_heap =
data->pipeline->base.base.create_flags & VK_PIPELINE_CREATE_2_DESCRIPTOR_HEAP_BIT_EXT;
nir_def *geometry_id = nir_iand_imm(b, intersection->geometry_id_and_flags, 0xfffffff);
nir_def *sbt_idx =
@ -1075,8 +1084,13 @@ handle_candidate_aabb(nir_builder *b, struct radv_leaf_intersection *intersectio
nir_def **params = rzalloc_array_size(b->shader, sizeof(nir_def *), param_count);
params[RT_ARG_LAUNCH_ID] = nir_load_ray_launch_id(b);
params[RT_ARG_LAUNCH_SIZE] = nir_load_ray_launch_size(b);
params[RT_ARG_DESCRIPTORS] = nir_load_rt_descriptors_amd(b);
params[RT_ARG_DYNAMIC_DESCRIPTORS] = nir_load_rt_dynamic_descriptors_amd(b);
if (uses_descriptor_heap) {
params[RT_ARG_HEAP_RESOURCE] = nir_load_rt_heap_resource_amd(b);
params[RT_ARG_HEAP_SAMPLER] = nir_load_rt_heap_sampler_amd(b);
} else {
params[RT_ARG_DESCRIPTORS] = nir_load_rt_descriptors_amd(b);
params[RT_ARG_DYNAMIC_DESCRIPTORS] = nir_load_rt_dynamic_descriptors_amd(b);
}
params[RT_ARG_PUSH_CONSTANTS] = nir_load_rt_push_constants_amd(b);
params[RT_ARG_SBT_DESCRIPTORS] = nir_load_sbt_base_amd(b);
params[AHIT_ISEC_ARG_SHADER_RECORD_PTR] = sbt_data.shader_record_ptr;
@ -1137,6 +1151,7 @@ struct radv_nir_rt_traversal_result
radv_build_traversal(struct radv_device *device, struct radv_ray_tracing_pipeline *pipeline, nir_builder *b,
struct radv_nir_rt_traversal_params *params, struct radv_ray_tracing_stage_info *info)
{
const bool uses_descriptor_heap = pipeline->base.base.create_flags & VK_PIPELINE_CREATE_2_DESCRIPTOR_HEAP_BIT_EXT;
const struct radv_physical_device *pdev = radv_device_physical(device);
nir_variable *barycentrics =
nir_variable_create(b->shader, nir_var_ray_hit_attrib, glsl_vector_type(GLSL_TYPE_FLOAT, 2), "barycentrics");
@ -1153,7 +1168,7 @@ radv_build_traversal(struct radv_device *device, struct radv_ray_tracing_pipelin
if (!params->preprocess_ahit_isec) {
nir_function *ahit_isec_func = nir_function_create(b->shader, "ahit_isec_func");
radv_nir_init_rt_function_params(ahit_isec_func, MESA_SHADER_ANY_HIT, params->payload_size,
params->hit_attrib_size);
params->hit_attrib_size, uses_descriptor_heap);
data.ahit_isec_func = ahit_isec_func;
}

1
src/amd/vulkan/radv_aco_shader_info.h
View file

@ -57,6 +57,7 @@ radv_aco_convert_shader_info(struct aco_shader_info *aco_info, const struct radv
ASSIGN_FIELD(vs.has_prolog);
ASSIGN_FIELD(ps.num_inputs);
ASSIGN_FIELD(cs.uses_full_subgroups);
ASSIGN_FIELD(descriptor_heap);
aco_info->vs.any_tcs_inputs_via_lds = radv->vs.tcs_inputs_via_lds != 0;
/* S2 must not be modified for correct hang recovery when NGG_WAVE_ID_EN=1. */
aco_info->vs.preserve_s2 = ngg_wave_id_en && gfx_level < GFX12;

3
src/amd/vulkan/radv_pipeline_rt.c
View file

@ -1044,11 +1044,12 @@ postprocess_rt_config(struct ac_shader_config *config, const struct radeon_info
static void
compile_rt_prolog(struct radv_device *device, struct radv_ray_tracing_pipeline *pipeline)
{
const bool uses_descriptor_heap = pipeline->base.base.create_flags & VK_PIPELINE_CREATE_2_DESCRIPTOR_HEAP_BIT_EXT;
const struct radv_physical_device *pdev = radv_device_physical(device);
uint32_t push_constant_size = 0;
struct radv_shader_stage prolog_stage = {0};
radv_build_rt_prolog(device, &prolog_stage);
radv_build_rt_prolog(device, &prolog_stage, uses_descriptor_heap);
prolog_stage.nir->options = &pdev->nir_options[MESA_SHADER_COMPUTE];
radv_optimize_nir(prolog_stage.nir, false);
radv_postprocess_nir(device, NULL, &prolog_stage);