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

radv/rt: Refactor and split radv_nir_rt_shader.c

Browse source 
This splits up radv_nir_rt_shader.c into several parts.
The first part is all ray traversal lowering for RT pipelines, located
at radv_nir_rt_traversal_shader.c. It implements building the traversal
loop, including inlined any-hit/intersection shaders (optionally as a
completely separate shader).

The second part is lowering for individual RT stages (right now,
monolithic vs. CPS-style separate compilation). Each lowering technique
lives in its own file (radv_nir_rt_stage_{monolithic,cps}.c).

Code shared between RT lowering techniques (shader inlining helpers and
storage lowering passes) gets moved into radv_nir_rt_stage_common.c.

One header, radv_nir_rt_stage.h, is the public interface for RT pipeline
stage lowering. Functions exposed to users (really just
radv_pipeline_rt.c) go there. The header for internal shared helpers is
radv_nir_rt_stage_common.c.

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38809>
This commit is contained in:
Natalie Vock 2025-12-04 00:08:22 +01:00 • committed by Marge Bot
parent 5681fabdc2
commit a488203e85
12 changed files with 2881 additions and 2222 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

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

@ -88,7 +88,10 @@ libradv_files = files(
'nir/radv_nir_opt_tid_function.c',
'nir/radv_nir_remap_color_attachment.c',
'nir/radv_nir_rt_common.c',
'nir/radv_nir_rt_shader.c',
'nir/radv_nir_rt_stage_common.c',
'nir/radv_nir_rt_stage_cps.c',
'nir/radv_nir_rt_stage_monolithic.c',
'nir/radv_nir_rt_traversal_shader.c',
'radv_acceleration_structure.c',
'radv_android.c',
'radv_android.h',

2194
src/amd/vulkan/nir/radv_nir_rt_shader.c
View file

File diff suppressed because it is too large Load diff

246
src/amd/vulkan/nir/radv_nir_rt_stage_common.c Normal file
View file

@ -0,0 +1,246 @@
/*
* Copyright © 2025 Valve Corporation
* Copyright © 2021 Google
*
* SPDX-License-Identifier: MIT
*/
#include "nir/radv_nir_rt_stage_common.h"
#include "nir_builder.h"
struct radv_nir_sbt_data
radv_nir_load_sbt_entry(nir_builder *b, nir_def *idx, enum radv_nir_sbt_type binding, enum radv_nir_sbt_entry offset)
{
struct radv_nir_sbt_data data;
nir_def *desc_base_addr = nir_load_sbt_base_amd(b);
nir_def *desc = nir_pack_64_2x32(b, ac_nir_load_smem(b, 2, desc_base_addr, nir_imm_int(b, binding), 4, 0));
nir_def *stride_offset = nir_imm_int(b, binding + (binding == SBT_RAYGEN ? 8 : 16));
nir_def *stride = ac_nir_load_smem(b, 1, desc_base_addr, stride_offset, 4, 0);
nir_def *addr = nir_iadd(b, desc, nir_u2u64(b, nir_iadd_imm(b, nir_imul(b, idx, stride), offset)));
unsigned load_size = offset == SBT_RECURSIVE_PTR ? 64 : 32;
data.shader_addr = nir_load_global(b, 1, load_size, addr, .access = ACCESS_CAN_REORDER | ACCESS_NON_WRITEABLE);
data.shader_record_ptr = nir_iadd_imm(b, addr, RADV_RT_HANDLE_SIZE - offset);
return data;
}
void
radv_nir_inline_constants(nir_shader *dst, nir_shader *src)
{
if (!src->constant_data_size)
return;
uint32_t old_constant_data_size = dst->constant_data_size;
uint32_t base_offset = align(dst->constant_data_size, 64);
dst->constant_data_size = base_offset + src->constant_data_size;
dst->constant_data = rerzalloc_size(dst, dst->constant_data, old_constant_data_size, dst->constant_data_size);
memcpy((char *)dst->constant_data + base_offset, src->constant_data, src->constant_data_size);
if (!base_offset)
return;
uint32_t base_align_mul = base_offset ? 1 << (ffs(base_offset) - 1) : NIR_ALIGN_MUL_MAX;
nir_foreach_block (block, nir_shader_get_entrypoint(src)) {
nir_foreach_instr (instr, block) {
if (instr->type != nir_instr_type_intrinsic)
continue;
nir_intrinsic_instr *intrinsic = nir_instr_as_intrinsic(instr);
if (intrinsic->intrinsic == nir_intrinsic_load_constant) {
nir_intrinsic_set_base(intrinsic, base_offset + nir_intrinsic_base(intrinsic));
uint32_t align_mul = nir_intrinsic_align_mul(intrinsic);
uint32_t align_offset = nir_intrinsic_align_offset(intrinsic);
align_mul = MIN2(align_mul, base_align_mul);
nir_intrinsic_set_align(intrinsic, align_mul, align_offset % align_mul);
}
}
}
}
struct inlined_shader_case {
struct radv_ray_tracing_group *group;
uint32_t call_idx;
};
static int
compare_inlined_shader_case(const void *a, const void *b)
{
const struct inlined_shader_case *visit_a = a;
const struct inlined_shader_case *visit_b = b;
return visit_a->call_idx > visit_b->call_idx ? 1 : visit_a->call_idx < visit_b->call_idx ? -1 : 0;
}
static void
insert_inlined_range(nir_builder *b, nir_def *sbt_idx, radv_insert_shader_case shader_case,
struct radv_rt_case_data *data, struct inlined_shader_case *cases, uint32_t length)
{
if (length >= INLINED_SHADER_BSEARCH_THRESHOLD) {
nir_push_if(b, nir_ige_imm(b, sbt_idx, cases[length / 2].call_idx));
{
insert_inlined_range(b, sbt_idx, shader_case, data, cases + (length / 2), length - (length / 2));
}
nir_push_else(b, NULL);
{
insert_inlined_range(b, sbt_idx, shader_case, data, cases, length / 2);
}
nir_pop_if(b, NULL);
} else {
for (uint32_t i = 0; i < length; ++i)
shader_case(b, sbt_idx, cases[i].group, data);
}
}
void
radv_visit_inlined_shaders(nir_builder *b, nir_def *sbt_idx, bool can_have_null_shaders, struct radv_rt_case_data *data,
radv_get_group_info group_info, radv_insert_shader_case shader_case)
{
struct inlined_shader_case *cases = calloc(data->pipeline->group_count, sizeof(struct inlined_shader_case));
uint32_t case_count = 0;
for (unsigned i = 0; i < data->pipeline->group_count; i++) {
struct radv_ray_tracing_group *group = &data->pipeline->groups[i];
uint32_t shader_index = VK_SHADER_UNUSED_KHR;
uint32_t handle_index = VK_SHADER_UNUSED_KHR;
group_info(group, &shader_index, &handle_index, data);
if (shader_index == VK_SHADER_UNUSED_KHR)
continue;
/* Avoid emitting stages with the same shaders/handles multiple times. */
bool duplicate = false;
for (unsigned j = 0; j < i; j++) {
uint32_t other_shader_index = VK_SHADER_UNUSED_KHR;
uint32_t other_handle_index = VK_SHADER_UNUSED_KHR;
group_info(&data->pipeline->groups[j], &other_shader_index, &other_handle_index, data);
if (handle_index == other_handle_index) {
duplicate = true;
break;
}
}
if (!duplicate) {
cases[case_count++] = (struct inlined_shader_case){
.group = group,
.call_idx = handle_index,
};
}
}
qsort(cases, case_count, sizeof(struct inlined_shader_case), compare_inlined_shader_case);
/* Do not emit 'if (sbt_idx != 0) { ... }' is there are only a few cases. */
can_have_null_shaders &= case_count >= RADV_RT_SWITCH_NULL_CHECK_THRESHOLD;
if (can_have_null_shaders)
nir_push_if(b, nir_ine_imm(b, sbt_idx, 0));
insert_inlined_range(b, sbt_idx, shader_case, data, cases, case_count);
if (can_have_null_shaders)
nir_pop_if(b, NULL);
free(cases);
}
bool
radv_nir_lower_rt_derefs(nir_shader *shader)
{
nir_function_impl *impl = nir_shader_get_entrypoint(shader);
bool progress = false;
nir_builder b;
nir_def *arg_offset = NULL;
nir_foreach_block (block, impl) {
nir_foreach_instr_safe (instr, block) {
if (instr->type != nir_instr_type_deref)
continue;
nir_deref_instr *deref = nir_instr_as_deref(instr);
if (!nir_deref_mode_is(deref, nir_var_shader_call_data))
continue;
deref->modes = nir_var_function_temp;
progress = true;
if (deref->deref_type == nir_deref_type_var) {
if (!arg_offset) {
b = nir_builder_at(nir_before_impl(impl));
arg_offset = nir_load_rt_arg_scratch_offset_amd(&b);
}
b.cursor = nir_before_instr(&deref->instr);
nir_deref_instr *replacement =
nir_build_deref_cast(&b, arg_offset, nir_var_function_temp, deref->var->type, 0);
nir_def_replace(&deref->def, &replacement->def);
}
}
}
return nir_progress(progress, impl, nir_metadata_control_flow);
}
/* Lowers hit attributes to registers or shared memory. If hit_attribs is NULL, attributes are
* lowered to shared memory. */
bool
radv_nir_lower_hit_attribs(nir_shader *shader, nir_variable **hit_attribs, uint32_t workgroup_size)
{
bool progress = false;
nir_function_impl *impl = nir_shader_get_entrypoint(shader);
nir_foreach_variable_with_modes (attrib, shader, nir_var_ray_hit_attrib) {
attrib->data.mode = nir_var_shader_temp;
progress = true;
}
nir_builder b = nir_builder_create(impl);
nir_foreach_block (block, impl) {
nir_foreach_instr_safe (instr, block) {
if (instr->type != nir_instr_type_intrinsic)
continue;
nir_intrinsic_instr *intrin = nir_instr_as_intrinsic(instr);
if (intrin->intrinsic != nir_intrinsic_load_hit_attrib_amd &&
intrin->intrinsic != nir_intrinsic_store_hit_attrib_amd)
continue;
progress = true;
b.cursor = nir_after_instr(instr);
nir_def *offset;
if (!hit_attribs)
offset = nir_imul_imm(
&b, nir_iadd_imm(&b, nir_load_local_invocation_index(&b), nir_intrinsic_base(intrin) * workgroup_size),
sizeof(uint32_t));
if (intrin->intrinsic == nir_intrinsic_load_hit_attrib_amd) {
nir_def *ret;
if (hit_attribs)
ret = nir_load_var(&b, hit_attribs[nir_intrinsic_base(intrin)]);
else
ret = nir_load_shared(&b, 1, 32, offset, .base = 0, .align_mul = 4);
nir_def_rewrite_uses(nir_instr_def(instr), ret);
} else {
if (hit_attribs)
nir_store_var(&b, hit_attribs[nir_intrinsic_base(intrin)], intrin->src->ssa, 0x1);
else
nir_store_shared(&b, intrin->src->ssa, offset, .base = 0, .align_mul = 4);
}
nir_instr_remove(instr);
}
}
if (!hit_attribs)
shader->info.shared_size = MAX2(shader->info.shared_size, workgroup_size * RADV_MAX_HIT_ATTRIB_SIZE);
return nir_progress(progress, impl, nir_metadata_control_flow);
}

153
src/amd/vulkan/nir/radv_nir_rt_stage_common.h Normal file
View file

@ -0,0 +1,153 @@
/*
* Copyright © 2025 Valve Corporation
* Copyright © 2021 Google
*
* SPDX-License-Identifier: MIT
*/
/* This file contains internal helpers for RT lowering shared between different lowering implementations. */
#ifndef MESA_RADV_NIR_RT_STAGE_COMMON_H
#define MESA_RADV_NIR_RT_STAGE_COMMON_H
#include "nir/radv_nir.h"
#include "ac_nir.h"
#include "radv_pipeline_cache.h"
#include "radv_pipeline_rt.h"
/*
*
* Common Constants
*
*/
/* Traversal stack size. This stack is put in LDS and experimentally 16 entries results in best
* performance. */
#define MAX_STACK_ENTRY_COUNT 16
#define RADV_RT_SWITCH_NULL_CHECK_THRESHOLD 3
/* Minimum number of inlined shaders to use binary search to select which shader to run. */
#define INLINED_SHADER_BSEARCH_THRESHOLD 16
/*
*
* Shader Inlining
*
*/
struct radv_rt_case_data {
struct radv_device *device;
struct radv_ray_tracing_pipeline *pipeline;
void *param_data;
};
typedef void (*radv_get_group_info)(struct radv_ray_tracing_group *, uint32_t *, uint32_t *,
struct radv_rt_case_data *);
typedef void (*radv_insert_shader_case)(nir_builder *, nir_def *, struct radv_ray_tracing_group *,
struct radv_rt_case_data *);
void radv_visit_inlined_shaders(nir_builder *b, nir_def *sbt_idx, bool can_have_null_shaders,
struct radv_rt_case_data *data, radv_get_group_info group_info,
radv_insert_shader_case shader_case);
/* Transfer inline constant data from src to dst, to prepare inlining src into dst */
void radv_nir_inline_constants(nir_shader *dst, nir_shader *src);
/*
*
* SBT Helpers
*
*/
struct radv_nir_sbt_data {
/* For inlined shaders, the index/ID of the shader to be executed.
* For separately-compiled shaders, an address to jump execution to.
*/
nir_def *shader_addr;
nir_def *shader_record_ptr;
};
enum radv_nir_sbt_type {
SBT_RAYGEN = offsetof(VkTraceRaysIndirectCommand2KHR, raygenShaderRecordAddress),
SBT_MISS = offsetof(VkTraceRaysIndirectCommand2KHR, missShaderBindingTableAddress),
SBT_HIT = offsetof(VkTraceRaysIndirectCommand2KHR, hitShaderBindingTableAddress),
SBT_CALLABLE = offsetof(VkTraceRaysIndirectCommand2KHR, callableShaderBindingTableAddress),
};
enum radv_nir_sbt_entry {
SBT_RECURSIVE_PTR = offsetof(struct radv_pipeline_group_handle, recursive_shader_ptr),
SBT_GENERAL_IDX = offsetof(struct radv_pipeline_group_handle, general_index),
SBT_CLOSEST_HIT_IDX = offsetof(struct radv_pipeline_group_handle, closest_hit_index),
SBT_INTERSECTION_IDX = offsetof(struct radv_pipeline_group_handle, intersection_index),
SBT_ANY_HIT_IDX = offsetof(struct radv_pipeline_group_handle, any_hit_index),
};
struct radv_nir_sbt_data radv_nir_load_sbt_entry(nir_builder *b, nir_def *idx, enum radv_nir_sbt_type binding,
enum radv_nir_sbt_entry offset);
/*
*
* Common lowering passes
*
*/
bool radv_nir_lower_rt_derefs(nir_shader *shader);
bool radv_nir_lower_hit_attribs(nir_shader *shader, nir_variable **hit_attribs, uint32_t workgroup_size);
/*
*
* Ray Traversal Helpers
*
*/
typedef void (*radv_nir_ahit_isec_preprocess_cb)(nir_shader *shader, void *data);
/* All parameters for performing ray traversal. */
struct radv_nir_rt_traversal_params {
nir_def *accel_struct;
nir_def *origin;
nir_def *direction;
nir_def *tmin;
nir_def *tmax;
nir_def *sbt_offset;
nir_def *sbt_stride;
nir_def *cull_mask_and_flags;
nir_def *miss_index;
bool ignore_cull_mask;
radv_nir_ahit_isec_preprocess_cb preprocess_ahit_isec;
/* User data passed to the inlining callback */
void *cb_data;
};
/* Variables describing the result of the traversal loop. */
struct radv_nir_rt_traversal_result {
nir_variable *sbt_index;
nir_variable *tmax;
nir_variable *hit;
nir_variable *primitive_addr;
nir_variable *primitive_id;
nir_variable *geometry_id_and_flags;
nir_variable *instance_addr;
nir_variable *hit_kind;
/* barycentrics are a bit special, because they're hit attributes (specifically, the first two hit attributes in
* attribute storage) under the hood.
* They're not considered in the init_traversal_result/copy_traversal_result helpers and need manual initialization
* wherever used.
*/
nir_variable *barycentrics;
};
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);
#endif // MESA_RADV_NIR_RT_STAGE_COMMON_H

687
src/amd/vulkan/nir/radv_nir_rt_stage_cps.c Normal file
View file

@ -0,0 +1,687 @@
/*
* Copyright © 2025 Valve Corporation
* Copyright © 2021 Google
*
* SPDX-License-Identifier: MIT
*/
#include "nir/nir.h"
#include "nir/nir_builder.h"
#include "nir/radv_nir.h"
#include "nir/radv_nir_rt_common.h"
#include "nir/radv_nir_rt_stage_common.h"
#include "nir/radv_nir_rt_stage_cps.h"
#include "ac_nir.h"
#include "radv_device.h"
#include "radv_physical_device.h"
#include "radv_pipeline_rt.h"
#include "radv_shader.h"
static bool
radv_arg_def_is_unused(nir_def *def)
{
nir_foreach_use (use, def) {
nir_instr *use_instr = nir_src_parent_instr(use);
if (use_instr->type == nir_instr_type_intrinsic) {
nir_intrinsic_instr *use_intr = nir_instr_as_intrinsic(use_instr);
if (use_intr->intrinsic == nir_intrinsic_store_scalar_arg_amd ||
use_intr->intrinsic == nir_intrinsic_store_vector_arg_amd)
continue;
} else if (use_instr->type == nir_instr_type_phi) {
nir_cf_node *prev_node = nir_cf_node_prev(&use_instr->block->cf_node);
if (!prev_node)
return false;
nir_phi_instr *phi = nir_instr_as_phi(use_instr);
if (radv_arg_def_is_unused(&phi->def))
continue;
}
return false;
}
return true;
}
static bool
radv_gather_unused_args_instr(nir_builder *b, nir_intrinsic_instr *instr, void *data)
{
if (instr->intrinsic != nir_intrinsic_load_scalar_arg_amd && instr->intrinsic != nir_intrinsic_load_vector_arg_amd)
return false;
if (!radv_arg_def_is_unused(&instr->def)) {
/* This arg is used for more than passing data to the next stage. */
struct radv_ray_tracing_stage_info *info = data;
BITSET_CLEAR(info->unused_args, nir_intrinsic_base(instr));
}
return false;
}
void
radv_gather_unused_args(struct radv_ray_tracing_stage_info *info, nir_shader *nir)
{
nir_shader_intrinsics_pass(nir, radv_gather_unused_args_instr, nir_metadata_all, info);
}
/*
* Global variables for an RT pipeline
*/
struct rt_variables {
struct radv_device *device;
const VkPipelineCreateFlags2 flags;
nir_variable *shader_addr;
nir_variable *traversal_addr;
/* scratch offset of the argument area relative to stack_ptr */
nir_variable *arg;
nir_variable *stack_ptr;
nir_variable *launch_sizes[3];
nir_variable *launch_ids[3];
/* global address of the SBT entry used for the shader */
nir_variable *shader_record_ptr;
/* trace_ray arguments */
nir_variable *accel_struct;
nir_variable *cull_mask_and_flags;
nir_variable *sbt_offset;
nir_variable *sbt_stride;
nir_variable *miss_index;
nir_variable *origin;
nir_variable *tmin;
nir_variable *direction;
nir_variable *tmax;
/* Properties of the primitive currently being visited. */
nir_variable *primitive_addr;
nir_variable *primitive_id;
nir_variable *geometry_id_and_flags;
nir_variable *instance_addr;
nir_variable *hit_kind;
unsigned stack_size;
};
static struct rt_variables
create_rt_variables(nir_shader *shader, struct radv_device *device, const VkPipelineCreateFlags2 flags)
{
struct rt_variables vars = {
.device = device,
.flags = flags,
};
vars.shader_addr = nir_variable_create(shader, nir_var_shader_temp, glsl_uint64_t_type(), "shader_addr");
vars.traversal_addr = nir_variable_create(shader, nir_var_shader_temp, glsl_uint64_t_type(), "traversal_addr");
vars.arg = nir_variable_create(shader, nir_var_shader_temp, glsl_uint_type(), "arg");
vars.stack_ptr = nir_variable_create(shader, nir_var_shader_temp, glsl_uint_type(), "stack_ptr");
vars.shader_record_ptr = nir_variable_create(shader, nir_var_shader_temp, glsl_uint64_t_type(), "shader_record_ptr");
vars.launch_sizes[0] = nir_variable_create(shader, nir_var_shader_temp, glsl_uint_type(), "launch_size_x");
vars.launch_sizes[1] = nir_variable_create(shader, nir_var_shader_temp, glsl_uint_type(), "launch_size_y");
vars.launch_sizes[2] = nir_variable_create(shader, nir_var_shader_temp, glsl_uint_type(), "launch_size_z");
vars.launch_ids[0] = nir_variable_create(shader, nir_var_shader_temp, glsl_uint_type(), "launch_id_x");
vars.launch_ids[1] = nir_variable_create(shader, nir_var_shader_temp, glsl_uint_type(), "launch_id_y");
vars.launch_ids[2] = nir_variable_create(shader, nir_var_shader_temp, glsl_uint_type(), "launch_id_z");
const struct glsl_type *vec3_type = glsl_vector_type(GLSL_TYPE_FLOAT, 3);
vars.accel_struct = nir_variable_create(shader, nir_var_shader_temp, glsl_uint64_t_type(), "accel_struct");
vars.cull_mask_and_flags = nir_variable_create(shader, nir_var_shader_temp, glsl_uint_type(), "cull_mask_and_flags");
vars.sbt_offset = nir_variable_create(shader, nir_var_shader_temp, glsl_uint_type(), "sbt_offset");
vars.sbt_stride = nir_variable_create(shader, nir_var_shader_temp, glsl_uint_type(), "sbt_stride");
vars.miss_index = nir_variable_create(shader, nir_var_shader_temp, glsl_uint_type(), "miss_index");
vars.origin = nir_variable_create(shader, nir_var_shader_temp, vec3_type, "ray_origin");
vars.tmin = nir_variable_create(shader, nir_var_shader_temp, glsl_float_type(), "ray_tmin");
vars.direction = nir_variable_create(shader, nir_var_shader_temp, vec3_type, "ray_direction");
vars.tmax = nir_variable_create(shader, nir_var_shader_temp, glsl_float_type(), "ray_tmax");
vars.primitive_addr = nir_variable_create(shader, nir_var_shader_temp, glsl_uint64_t_type(), "primitive_addr");
vars.primitive_id = nir_variable_create(shader, nir_var_shader_temp, glsl_uint_type(), "primitive_id");
vars.geometry_id_and_flags =
nir_variable_create(shader, nir_var_shader_temp, glsl_uint_type(), "geometry_id_and_flags");
vars.instance_addr = nir_variable_create(shader, nir_var_shader_temp, glsl_uint64_t_type(), "instance_addr");
vars.hit_kind = nir_variable_create(shader, nir_var_shader_temp, glsl_uint_type(), "hit_kind");
return vars;
}
static void
insert_rt_return(nir_builder *b, const struct rt_variables *vars)
{
nir_store_var(b, vars->stack_ptr, nir_iadd_imm(b, nir_load_var(b, vars->stack_ptr), -16), 1);
nir_store_var(b, vars->shader_addr, nir_load_scratch(b, 1, 64, nir_load_var(b, vars->stack_ptr), .align_mul = 16),
1);
}
struct radv_rt_shader_info {
bool uses_launch_id;
bool uses_launch_size;
};
struct radv_lower_rt_instruction_data {
struct rt_variables *vars;
struct radv_rt_shader_info *out_info;
};
static bool
radv_lower_rt_instruction(nir_builder *b, nir_instr *instr, void *_data)
{
if (instr->type == nir_instr_type_jump) {
nir_jump_instr *jump = nir_instr_as_jump(instr);
if (jump->type == nir_jump_halt) {
jump->type = nir_jump_return;
return true;
}
return false;
} else if (instr->type != nir_instr_type_intrinsic) {
return false;
}
nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
struct radv_lower_rt_instruction_data *data = _data;
struct rt_variables *vars = data->vars;
b->cursor = nir_before_instr(&intr->instr);
nir_def *ret = NULL;
switch (intr->intrinsic) {
case nir_intrinsic_rt_execute_callable: {
uint32_t size = align(nir_intrinsic_stack_size(intr), 16);
nir_def *ret_ptr = nir_load_resume_shader_address_amd(b, nir_intrinsic_call_idx(intr));
ret_ptr = nir_ior_imm(b, ret_ptr, radv_get_rt_priority(b->shader->info.stage));
nir_store_var(b, vars->stack_ptr, nir_iadd_imm_nuw(b, nir_load_var(b, vars->stack_ptr), size), 1);
nir_store_scratch(b, ret_ptr, nir_load_var(b, vars->stack_ptr), .align_mul = 16);
nir_store_var(b, vars->stack_ptr, nir_iadd_imm_nuw(b, nir_load_var(b, vars->stack_ptr), 16), 1);
struct radv_nir_sbt_data sbt_data =
radv_nir_load_sbt_entry(b, intr->src[0].ssa, SBT_CALLABLE, SBT_RECURSIVE_PTR);
nir_store_var(b, vars->shader_addr, sbt_data.shader_addr, 0x1);
nir_store_var(b, vars->shader_record_ptr, sbt_data.shader_record_ptr, 0x1);
nir_store_var(b, vars->arg, nir_iadd_imm(b, intr->src[1].ssa, -size - 16), 1);
vars->stack_size = MAX2(vars->stack_size, size + 16);
break;
}
case nir_intrinsic_rt_trace_ray: {
uint32_t size = align(nir_intrinsic_stack_size(intr), 16);
nir_def *ret_ptr = nir_load_resume_shader_address_amd(b, nir_intrinsic_call_idx(intr));
ret_ptr = nir_ior_imm(b, ret_ptr, radv_get_rt_priority(b->shader->info.stage));
nir_store_var(b, vars->stack_ptr, nir_iadd_imm_nuw(b, nir_load_var(b, vars->stack_ptr), size), 1);
nir_store_scratch(b, ret_ptr, nir_load_var(b, vars->stack_ptr), .align_mul = 16);
nir_store_var(b, vars->stack_ptr, nir_iadd_imm_nuw(b, nir_load_var(b, vars->stack_ptr), 16), 1);
nir_store_var(b, vars->shader_addr, nir_load_var(b, vars->traversal_addr), 1);
nir_store_var(b, vars->arg, nir_iadd_imm(b, intr->src[10].ssa, -size - 16), 1);
vars->stack_size = MAX2(vars->stack_size, size + 16);
/* Per the SPIR-V extension spec we have to ignore some bits for some arguments. */
nir_store_var(b, vars->accel_struct, intr->src[0].ssa, 0x1);
nir_store_var(b, vars->cull_mask_and_flags, nir_ior(b, nir_ishl_imm(b, intr->src[2].ssa, 24), intr->src[1].ssa),
0x1);
nir_store_var(b, vars->sbt_offset, nir_iand_imm(b, intr->src[3].ssa, 0xf), 0x1);
nir_store_var(b, vars->sbt_stride, nir_iand_imm(b, intr->src[4].ssa, 0xf), 0x1);
nir_store_var(b, vars->miss_index, nir_iand_imm(b, intr->src[5].ssa, 0xffff), 0x1);
nir_store_var(b, vars->origin, intr->src[6].ssa, 0x7);
nir_store_var(b, vars->tmin, intr->src[7].ssa, 0x1);
nir_store_var(b, vars->direction, intr->src[8].ssa, 0x7);
nir_store_var(b, vars->tmax, intr->src[9].ssa, 0x1);
break;
}
case nir_intrinsic_rt_resume: {
uint32_t size = align(nir_intrinsic_stack_size(intr), 16);
nir_store_var(b, vars->stack_ptr, nir_iadd_imm(b, nir_load_var(b, vars->stack_ptr), -size), 1);
break;
}
case nir_intrinsic_rt_return_amd: {
if (b->shader->info.stage == MESA_SHADER_RAYGEN) {
nir_terminate(b);
break;
}
insert_rt_return(b, vars);
break;
}
case nir_intrinsic_load_scratch: {
nir_src_rewrite(&intr->src[0], nir_iadd_nuw(b, nir_load_var(b, vars->stack_ptr), intr->src[0].ssa));
return true;
}
case nir_intrinsic_store_scratch: {
nir_src_rewrite(&intr->src[1], nir_iadd_nuw(b, nir_load_var(b, vars->stack_ptr), intr->src[1].ssa));
return true;
}
case nir_intrinsic_load_rt_arg_scratch_offset_amd: {
ret = nir_load_var(b, vars->arg);
break;
}
case nir_intrinsic_load_shader_record_ptr: {
ret = nir_load_var(b, vars->shader_record_ptr);
break;
}
case nir_intrinsic_load_ray_launch_size: {
if (data->out_info)
data->out_info->uses_launch_size = true;
ret = nir_vec3(b, nir_load_var(b, vars->launch_sizes[0]), nir_load_var(b, vars->launch_sizes[1]),
nir_load_var(b, vars->launch_sizes[2]));
break;
};
case nir_intrinsic_load_ray_launch_id: {
if (data->out_info)
data->out_info->uses_launch_id = true;
ret = nir_vec3(b, nir_load_var(b, vars->launch_ids[0]), nir_load_var(b, vars->launch_ids[1]),
nir_load_var(b, vars->launch_ids[2]));
break;
}
case nir_intrinsic_load_ray_t_min: {
ret = nir_load_var(b, vars->tmin);
break;
}
case nir_intrinsic_load_ray_t_max: {
ret = nir_load_var(b, vars->tmax);
break;
}
case nir_intrinsic_load_ray_world_origin: {
ret = nir_load_var(b, vars->origin);
break;
}
case nir_intrinsic_load_ray_world_direction: {
ret = nir_load_var(b, vars->direction);
break;
}
case nir_intrinsic_load_ray_instance_custom_index: {
ret = radv_load_custom_instance(vars->device, b, nir_load_var(b, vars->instance_addr));
break;
}
case nir_intrinsic_load_primitive_id: {
ret = nir_load_var(b, vars->primitive_id);
break;
}
case nir_intrinsic_load_ray_geometry_index: {
ret = nir_load_var(b, vars->geometry_id_and_flags);
ret = nir_iand_imm(b, ret, 0xFFFFFFF);
break;
}
case nir_intrinsic_load_instance_id: {
ret = radv_load_instance_id(vars->device, b, nir_load_var(b, vars->instance_addr));
break;
}
case nir_intrinsic_load_ray_flags: {
ret = nir_iand_imm(b, nir_load_var(b, vars->cull_mask_and_flags), 0xFFFFFF);
break;
}
case nir_intrinsic_load_ray_hit_kind: {
ret = nir_load_var(b, vars->hit_kind);
break;
}
case nir_intrinsic_load_ray_world_to_object: {
unsigned c = nir_intrinsic_column(intr);
nir_def *instance_node_addr = nir_load_var(b, vars->instance_addr);
nir_def *wto_matrix[3];
radv_load_wto_matrix(vars->device, b, instance_node_addr, wto_matrix);
nir_def *vals[3];
for (unsigned i = 0; i < 3; ++i)
vals[i] = nir_channel(b, wto_matrix[i], c);
ret = nir_vec(b, vals, 3);
break;
}
case nir_intrinsic_load_ray_object_to_world: {
unsigned c = nir_intrinsic_column(intr);
nir_def *otw_matrix[3];
radv_load_otw_matrix(vars->device, b, nir_load_var(b, vars->instance_addr), otw_matrix);
ret = nir_vec3(b, nir_channel(b, otw_matrix[0], c), nir_channel(b, otw_matrix[1], c),
nir_channel(b, otw_matrix[2], c));
break;
}
case nir_intrinsic_load_ray_object_origin: {
nir_def *wto_matrix[3];
radv_load_wto_matrix(vars->device, b, nir_load_var(b, vars->instance_addr), wto_matrix);
ret = nir_build_vec3_mat_mult(b, nir_load_var(b, vars->origin), wto_matrix, true);
break;
}
case nir_intrinsic_load_ray_object_direction: {
nir_def *wto_matrix[3];
radv_load_wto_matrix(vars->device, b, nir_load_var(b, vars->instance_addr), wto_matrix);
ret = nir_build_vec3_mat_mult(b, nir_load_var(b, vars->direction), wto_matrix, false);
break;
}
case nir_intrinsic_load_cull_mask: {
ret = nir_ushr_imm(b, nir_load_var(b, vars->cull_mask_and_flags), 24);
break;
}
case nir_intrinsic_load_sbt_offset_amd: {
ret = nir_load_var(b, vars->sbt_offset);
break;
}
case nir_intrinsic_load_sbt_stride_amd: {
ret = nir_load_var(b, vars->sbt_stride);
break;
}
case nir_intrinsic_load_accel_struct_amd: {
ret = nir_load_var(b, vars->accel_struct);
break;
}
case nir_intrinsic_load_cull_mask_and_flags_amd: {
ret = nir_load_var(b, vars->cull_mask_and_flags);
break;
}
case nir_intrinsic_execute_closest_hit_amd: {
nir_store_var(b, vars->tmax, intr->src[1].ssa, 0x1);
nir_store_var(b, vars->primitive_addr, intr->src[2].ssa, 0x1);
nir_store_var(b, vars->primitive_id, intr->src[3].ssa, 0x1);
nir_store_var(b, vars->instance_addr, intr->src[4].ssa, 0x1);
nir_store_var(b, vars->geometry_id_and_flags, intr->src[5].ssa, 0x1);
nir_store_var(b, vars->hit_kind, intr->src[6].ssa, 0x1);
struct radv_nir_sbt_data sbt_data =
radv_nir_load_sbt_entry(b, intr->src[0].ssa, SBT_HIT, SBT_RECURSIVE_PTR);
nir_store_var(b, vars->shader_addr, sbt_data.shader_addr, 0x1);
nir_store_var(b, vars->shader_record_ptr, sbt_data.shader_record_ptr, 0x1);
nir_def *should_return =
nir_test_mask(b, nir_load_var(b, vars->cull_mask_and_flags), SpvRayFlagsSkipClosestHitShaderKHRMask);
if (!(vars->flags & VK_PIPELINE_CREATE_2_RAY_TRACING_NO_NULL_CLOSEST_HIT_SHADERS_BIT_KHR)) {
should_return = nir_ior(b, should_return, nir_ieq_imm(b, nir_load_var(b, vars->shader_addr), 0));
}
/* should_return is set if we had a hit but we won't be calling the closest hit
* shader and hence need to return immediately to the calling shader. */
nir_push_if(b, should_return);
insert_rt_return(b, vars);
nir_pop_if(b, NULL);
break;
}
case nir_intrinsic_execute_miss_amd: {
nir_store_var(b, vars->tmax, intr->src[0].ssa, 0x1);
nir_def *undef = nir_undef(b, 1, 32);
nir_store_var(b, vars->primitive_id, undef, 0x1);
nir_store_var(b, vars->instance_addr, nir_undef(b, 1, 64), 0x1);
nir_store_var(b, vars->geometry_id_and_flags, undef, 0x1);
nir_store_var(b, vars->hit_kind, undef, 0x1);
nir_def *miss_index = nir_load_var(b, vars->miss_index);
struct radv_nir_sbt_data sbt_data =
radv_nir_load_sbt_entry(b, miss_index, SBT_MISS, SBT_RECURSIVE_PTR);
nir_store_var(b, vars->shader_addr, sbt_data.shader_addr, 0x1);
nir_store_var(b, vars->shader_record_ptr, sbt_data.shader_record_ptr, 0x1);
if (!(vars->flags & VK_PIPELINE_CREATE_2_RAY_TRACING_NO_NULL_MISS_SHADERS_BIT_KHR)) {
/* In case of a NULL miss shader, do nothing and just return. */
nir_push_if(b, nir_ieq_imm(b, nir_load_var(b, vars->shader_addr), 0));
insert_rt_return(b, vars);
nir_pop_if(b, NULL);
}
break;
}
case nir_intrinsic_load_ray_triangle_vertex_positions: {
nir_def *primitive_addr = nir_load_var(b, vars->primitive_addr);
ret = radv_load_vertex_position(vars->device, b, primitive_addr, nir_intrinsic_column(intr));
break;
}
default:
return false;
}
if (ret)
nir_def_rewrite_uses(&intr->def, ret);
nir_instr_remove(&intr->instr);
return true;
}
/* This lowers all the RT instructions that we do not want to pass on to the combined shader and
* that we can implement using the variables from the shader we are going to inline into. */
static bool
lower_rt_instructions(nir_shader *shader, struct rt_variables *vars, struct radv_rt_shader_info *out_info)
{
struct radv_lower_rt_instruction_data data = {
.vars = vars,
.out_info = out_info,
};
return nir_shader_instructions_pass(shader, radv_lower_rt_instruction, nir_metadata_none, &data);
}
void
radv_nir_lower_rt_io_cps(nir_shader *nir)
{
NIR_PASS(_, nir, nir_lower_vars_to_explicit_types, nir_var_function_temp | nir_var_shader_call_data,
glsl_get_natural_size_align_bytes);
NIR_PASS(_, nir, radv_nir_lower_rt_derefs);
NIR_PASS(_, nir, nir_lower_explicit_io, nir_var_function_temp, nir_address_format_32bit_offset);
}
/** Select the next shader based on priorities:
*
* Detect the priority of the shader stage by the lowest bits in the address (low to high):
* - Raygen - idx 0
* - Traversal - idx 1
* - Closest Hit / Miss - idx 2
* - Callable - idx 3
*
*
* This gives us the following priorities:
* Raygen : Callable > > Traversal > Raygen
* Traversal : > Chit / Miss > > Raygen
* CHit / Miss : Callable > Chit / Miss > Traversal > Raygen
* Callable : Callable > Chit / Miss > > Raygen
*/
static nir_def *
select_next_shader(nir_builder *b, nir_def *shader_addr, unsigned wave_size)
{
mesa_shader_stage stage = b->shader->info.stage;
nir_def *prio = nir_iand_imm(b, shader_addr, radv_rt_priority_mask);
nir_def *ballot = nir_ballot(b, 1, wave_size, nir_imm_bool(b, true));
nir_def *ballot_traversal = nir_ballot(b, 1, wave_size, nir_ieq_imm(b, prio, radv_rt_priority_traversal));
nir_def *ballot_hit_miss = nir_ballot(b, 1, wave_size, nir_ieq_imm(b, prio, radv_rt_priority_hit_miss));
nir_def *ballot_callable = nir_ballot(b, 1, wave_size, nir_ieq_imm(b, prio, radv_rt_priority_callable));
if (stage != MESA_SHADER_CALLABLE && stage != MESA_SHADER_INTERSECTION)
ballot = nir_bcsel(b, nir_ine_imm(b, ballot_traversal, 0), ballot_traversal, ballot);
if (stage != MESA_SHADER_RAYGEN)
ballot = nir_bcsel(b, nir_ine_imm(b, ballot_hit_miss, 0), ballot_hit_miss, ballot);
if (stage != MESA_SHADER_INTERSECTION)
ballot = nir_bcsel(b, nir_ine_imm(b, ballot_callable, 0), ballot_callable, ballot);
nir_def *lsb = nir_find_lsb(b, ballot);
nir_def *next = nir_read_invocation(b, shader_addr, lsb);
return nir_iand_imm(b, next, ~radv_rt_priority_mask);
}
static void
radv_store_arg(nir_builder *b, const struct radv_shader_args *args, const struct radv_ray_tracing_stage_info *info,
struct ac_arg arg, nir_def *value)
{
/* Do not pass unused data to the next stage. */
if (!info || !BITSET_TEST(info->unused_args, arg.arg_index))
ac_nir_store_arg(b, &args->ac, arg, value);
}
void
radv_nir_lower_rt_abi_cps(nir_shader *shader, const struct radv_shader_args *args, const struct radv_shader_info *info,
uint32_t *stack_size, bool resume_shader, struct radv_device *device,
struct radv_ray_tracing_pipeline *pipeline, bool has_position_fetch,
const struct radv_ray_tracing_stage_info *traversal_info)
{
const struct radv_physical_device *pdev = radv_device_physical(device);
nir_function_impl *impl = nir_shader_get_entrypoint(shader);
struct rt_variables vars = create_rt_variables(shader, device, pipeline->base.base.create_flags);
struct radv_rt_shader_info rt_info = {0};
lower_rt_instructions(shader, &vars, &rt_info);
if (stack_size) {
vars.stack_size = MAX2(vars.stack_size, shader->scratch_size);
*stack_size = MAX2(*stack_size, vars.stack_size);
}
shader->scratch_size = 0;
/* This can't use NIR_PASS because NIR_DEBUG=serialize,clone invalidates pointers. */
nir_lower_returns(shader);
nir_cf_list list;
nir_cf_extract(&list, nir_before_impl(impl), nir_after_impl(impl));
/* initialize variables */
nir_builder b = nir_builder_at(nir_before_impl(impl));
nir_def *descriptors = ac_nir_load_arg(&b, &args->ac, args->descriptors[0]);
nir_def *push_constants = ac_nir_load_arg(&b, &args->ac, args->ac.push_constants);
nir_def *dynamic_descriptors = ac_nir_load_arg(&b, &args->ac, args->ac.dynamic_descriptors);
nir_def *sbt_descriptors = ac_nir_load_arg(&b, &args->ac, args->ac.rt.sbt_descriptors);
nir_def *launch_sizes[3];
for (uint32_t i = 0; i < ARRAY_SIZE(launch_sizes); i++) {
launch_sizes[i] = ac_nir_load_arg(&b, &args->ac, args->ac.rt.launch_sizes[i]);
nir_store_var(&b, vars.launch_sizes[i], launch_sizes[i], 1);
}
nir_def *scratch_offset = NULL;
if (args->ac.scratch_offset.used)
scratch_offset = ac_nir_load_arg(&b, &args->ac, args->ac.scratch_offset);
nir_def *ring_offsets = NULL;
if (args->ac.ring_offsets.used)
ring_offsets = ac_nir_load_arg(&b, &args->ac, args->ac.ring_offsets);
nir_def *launch_ids[3];
for (uint32_t i = 0; i < ARRAY_SIZE(launch_ids); i++) {
launch_ids[i] = ac_nir_load_arg(&b, &args->ac, args->ac.rt.launch_ids[i]);
nir_store_var(&b, vars.launch_ids[i], launch_ids[i], 1);
}
nir_def *traversal_addr = ac_nir_load_arg(&b, &args->ac, args->ac.rt.traversal_shader_addr);
nir_store_var(&b, vars.traversal_addr,
nir_pack_64_2x32_split(&b, traversal_addr, nir_imm_int(&b, pdev->info.address32_hi)), 1);
nir_def *shader_addr = ac_nir_load_arg(&b, &args->ac, args->ac.rt.shader_addr);
shader_addr = nir_pack_64_2x32(&b, shader_addr);
nir_store_var(&b, vars.shader_addr, shader_addr, 1);
nir_store_var(&b, vars.stack_ptr, ac_nir_load_arg(&b, &args->ac, args->ac.rt.dynamic_callable_stack_base), 1);
nir_def *record_ptr = ac_nir_load_arg(&b, &args->ac, args->ac.rt.shader_record);
nir_store_var(&b, vars.shader_record_ptr, nir_pack_64_2x32(&b, record_ptr), 1);
nir_store_var(&b, vars.arg, ac_nir_load_arg(&b, &args->ac, args->ac.rt.payload_offset), 1);
nir_def *accel_struct = ac_nir_load_arg(&b, &args->ac, args->ac.rt.accel_struct);
nir_store_var(&b, vars.accel_struct, nir_pack_64_2x32(&b, accel_struct), 1);
nir_store_var(&b, vars.cull_mask_and_flags, ac_nir_load_arg(&b, &args->ac, args->ac.rt.cull_mask_and_flags), 1);
nir_store_var(&b, vars.sbt_offset, ac_nir_load_arg(&b, &args->ac, args->ac.rt.sbt_offset), 1);
nir_store_var(&b, vars.sbt_stride, ac_nir_load_arg(&b, &args->ac, args->ac.rt.sbt_stride), 1);
nir_store_var(&b, vars.origin, ac_nir_load_arg(&b, &args->ac, args->ac.rt.ray_origin), 0x7);
nir_store_var(&b, vars.tmin, ac_nir_load_arg(&b, &args->ac, args->ac.rt.ray_tmin), 1);
nir_store_var(&b, vars.direction, ac_nir_load_arg(&b, &args->ac, args->ac.rt.ray_direction), 0x7);
nir_store_var(&b, vars.tmax, ac_nir_load_arg(&b, &args->ac, args->ac.rt.ray_tmax), 1);
if (traversal_info && traversal_info->miss_index.state == RADV_RT_CONST_ARG_STATE_VALID)
nir_store_var(&b, vars.miss_index, nir_imm_int(&b, traversal_info->miss_index.value), 0x1);
else
nir_store_var(&b, vars.miss_index, ac_nir_load_arg(&b, &args->ac, args->ac.rt.miss_index), 0x1);
nir_def *primitive_addr = ac_nir_load_arg(&b, &args->ac, args->ac.rt.primitive_addr);
nir_store_var(&b, vars.primitive_addr, nir_pack_64_2x32(&b, primitive_addr), 1);
nir_store_var(&b, vars.primitive_id, ac_nir_load_arg(&b, &args->ac, args->ac.rt.primitive_id), 1);
nir_def *instance_addr = ac_nir_load_arg(&b, &args->ac, args->ac.rt.instance_addr);
nir_store_var(&b, vars.instance_addr, nir_pack_64_2x32(&b, instance_addr), 1);
nir_store_var(&b, vars.geometry_id_and_flags, ac_nir_load_arg(&b, &args->ac, args->ac.rt.geometry_id_and_flags), 1);
nir_store_var(&b, vars.hit_kind, ac_nir_load_arg(&b, &args->ac, args->ac.rt.hit_kind), 1);
/* guard the shader, so that only the correct invocations execute it */
nir_if *shader_guard = NULL;
if (shader->info.stage != MESA_SHADER_RAYGEN || resume_shader) {
nir_def *uniform_shader_addr = ac_nir_load_arg(&b, &args->ac, args->ac.rt.uniform_shader_addr);
uniform_shader_addr = nir_pack_64_2x32(&b, uniform_shader_addr);
uniform_shader_addr = nir_ior_imm(&b, uniform_shader_addr, radv_get_rt_priority(shader->info.stage));
shader_guard = nir_push_if(&b, nir_ieq(&b, uniform_shader_addr, shader_addr));
shader_guard->control = nir_selection_control_divergent_always_taken;
}
nir_cf_reinsert(&list, b.cursor);
if (shader_guard)
nir_pop_if(&b, shader_guard);
b.cursor = nir_after_impl(impl);
/* select next shader */
shader_addr = nir_load_var(&b, vars.shader_addr);
nir_def *next = select_next_shader(&b, shader_addr, info->wave_size);
ac_nir_store_arg(&b, &args->ac, args->ac.rt.uniform_shader_addr, next);
ac_nir_store_arg(&b, &args->ac, args->descriptors[0], descriptors);
ac_nir_store_arg(&b, &args->ac, args->ac.push_constants, push_constants);
ac_nir_store_arg(&b, &args->ac, args->ac.dynamic_descriptors, dynamic_descriptors);
ac_nir_store_arg(&b, &args->ac, args->ac.rt.sbt_descriptors, sbt_descriptors);
ac_nir_store_arg(&b, &args->ac, args->ac.rt.traversal_shader_addr, traversal_addr);
for (uint32_t i = 0; i < ARRAY_SIZE(launch_sizes); i++) {
if (rt_info.uses_launch_size)
ac_nir_store_arg(&b, &args->ac, args->ac.rt.launch_sizes[i], launch_sizes[i]);
else
radv_store_arg(&b, args, traversal_info, args->ac.rt.launch_sizes[i], launch_sizes[i]);
}
if (scratch_offset)
ac_nir_store_arg(&b, &args->ac, args->ac.scratch_offset, scratch_offset);
if (ring_offsets)
ac_nir_store_arg(&b, &args->ac, args->ac.ring_offsets, ring_offsets);
for (uint32_t i = 0; i < ARRAY_SIZE(launch_ids); i++) {
if (rt_info.uses_launch_id)
ac_nir_store_arg(&b, &args->ac, args->ac.rt.launch_ids[i], launch_ids[i]);
else
radv_store_arg(&b, args, traversal_info, args->ac.rt.launch_ids[i], launch_ids[i]);
}
/* store back all variables to registers */
ac_nir_store_arg(&b, &args->ac, args->ac.rt.dynamic_callable_stack_base, nir_load_var(&b, vars.stack_ptr));
ac_nir_store_arg(&b, &args->ac, args->ac.rt.shader_addr, shader_addr);
radv_store_arg(&b, args, traversal_info, args->ac.rt.shader_record, nir_load_var(&b, vars.shader_record_ptr));
radv_store_arg(&b, args, traversal_info, args->ac.rt.payload_offset, nir_load_var(&b, vars.arg));
radv_store_arg(&b, args, traversal_info, args->ac.rt.accel_struct, nir_load_var(&b, vars.accel_struct));
radv_store_arg(&b, args, traversal_info, args->ac.rt.cull_mask_and_flags,
nir_load_var(&b, vars.cull_mask_and_flags));
radv_store_arg(&b, args, traversal_info, args->ac.rt.sbt_offset, nir_load_var(&b, vars.sbt_offset));
radv_store_arg(&b, args, traversal_info, args->ac.rt.sbt_stride, nir_load_var(&b, vars.sbt_stride));
radv_store_arg(&b, args, traversal_info, args->ac.rt.miss_index, nir_load_var(&b, vars.miss_index));
radv_store_arg(&b, args, traversal_info, args->ac.rt.ray_origin, nir_load_var(&b, vars.origin));
radv_store_arg(&b, args, traversal_info, args->ac.rt.ray_tmin, nir_load_var(&b, vars.tmin));
radv_store_arg(&b, args, traversal_info, args->ac.rt.ray_direction, nir_load_var(&b, vars.direction));
radv_store_arg(&b, args, traversal_info, args->ac.rt.ray_tmax, nir_load_var(&b, vars.tmax));
if (has_position_fetch)
radv_store_arg(&b, args, traversal_info, args->ac.rt.primitive_addr, nir_load_var(&b, vars.primitive_addr));
radv_store_arg(&b, args, traversal_info, args->ac.rt.primitive_id, nir_load_var(&b, vars.primitive_id));
radv_store_arg(&b, args, traversal_info, args->ac.rt.instance_addr, nir_load_var(&b, vars.instance_addr));
radv_store_arg(&b, args, traversal_info, args->ac.rt.geometry_id_and_flags,
nir_load_var(&b, vars.geometry_id_and_flags));
radv_store_arg(&b, args, traversal_info, args->ac.rt.hit_kind, nir_load_var(&b, vars.hit_kind));
nir_progress(true, impl, nir_metadata_none);
/* cleanup passes */
NIR_PASS(_, shader, nir_lower_global_vars_to_local);
NIR_PASS(_, shader, nir_lower_vars_to_ssa);
if (shader->info.stage == MESA_SHADER_CLOSEST_HIT || shader->info.stage == MESA_SHADER_INTERSECTION)
NIR_PASS(_, shader, radv_nir_lower_hit_attribs, NULL, info->wave_size);
}

22
src/amd/vulkan/nir/radv_nir_rt_stage_cps.h Normal file
View file

@ -0,0 +1,22 @@
/*
* Copyright © 2025 Valve Corporation
*
* SPDX-License-Identifier: MIT
*/
/* This file contains the public interface for all RT pipeline stage lowering. */
#ifndef RADV_NIR_RT_STAGE_CPS_H
#define RADV_NIR_RT_STAGE_CPS_H
#include "radv_pipeline_rt.h"
void radv_gather_unused_args(struct radv_ray_tracing_stage_info *info, nir_shader *nir);
void radv_nir_lower_rt_abi_cps(nir_shader *shader, const struct radv_shader_args *args,
const struct radv_shader_info *info, uint32_t *stack_size, bool resume_shader,
struct radv_device *device, struct radv_ray_tracing_pipeline *pipeline,
bool has_position_fetch, const struct radv_ray_tracing_stage_info *traversal_info);
void radv_nir_lower_rt_io_cps(nir_shader *shader);
#endif // RADV_NIR_RT_STAGE_CPS_H

487
src/amd/vulkan/nir/radv_nir_rt_stage_monolithic.c Normal file
View file

@ -0,0 +1,487 @@
/*
* Copyright © 2025 Valve Corporation
* Copyright © 2021 Google
*
* SPDX-License-Identifier: MIT
*/
#include "nir/radv_nir_rt_common.h"
#include "nir/radv_nir_rt_stage_common.h"
#include "nir/radv_nir_rt_stage_monolithic.h"
#include "nir_builder.h"
#include "radv_device.h"
#include "radv_physical_device.h"
struct chit_miss_inlining_params {
struct radv_device *device;
struct radv_nir_rt_traversal_params *trav_params;
struct radv_nir_rt_traversal_result *trav_result;
struct radv_nir_sbt_data *sbt;
unsigned payload_offset;
};
struct chit_miss_inlining_vars {
struct radv_device *device;
nir_variable *shader_record_ptr;
nir_variable *origin;
nir_variable *direction;
nir_variable *tmin;
nir_variable *tmax;
nir_variable *primitive_addr;
nir_variable *primitive_id;
nir_variable *geometry_id_and_flags;
nir_variable *cull_mask_and_flags;
nir_variable *instance_addr;
nir_variable *hit_kind;
};
static void
init_chit_miss_inlining_vars(nir_shader *shader, struct chit_miss_inlining_vars *vars)
{
const struct glsl_type *vec3_type = glsl_vector_type(GLSL_TYPE_FLOAT, 3);
vars->shader_record_ptr =
nir_variable_create(shader, nir_var_shader_temp, glsl_uint64_t_type(), "shader_record_ptr");
vars->origin = nir_variable_create(shader, nir_var_shader_temp, vec3_type, "origin");
vars->direction = nir_variable_create(shader, nir_var_shader_temp, vec3_type, "direction");
vars->tmin = nir_variable_create(shader, nir_var_shader_temp, glsl_float_type(), "tmin");
vars->tmax = nir_variable_create(shader, nir_var_shader_temp, glsl_float_type(), "tmax");
vars->primitive_addr = nir_variable_create(shader, nir_var_shader_temp, glsl_uint64_t_type(), "primitive_addr");
vars->primitive_id = nir_variable_create(shader, nir_var_shader_temp, glsl_uint_type(), "primitive_id");
vars->geometry_id_and_flags =
nir_variable_create(shader, nir_var_shader_temp, glsl_uint_type(), "geometry_id_and_flags");
vars->cull_mask_and_flags =
nir_variable_create(shader, nir_var_shader_temp, glsl_uint_type(), "cull_mask_and_flags");
vars->instance_addr = nir_variable_create(shader, nir_var_shader_temp, glsl_uint64_t_type(), "instance_addr");
vars->hit_kind = nir_variable_create(shader, nir_var_shader_temp, glsl_uint_type(), "hit_kind");
}
static void
setup_chit_miss_inlining(struct chit_miss_inlining_vars *vars, const struct chit_miss_inlining_params *params,
nir_builder *b, nir_shader *chit, struct hash_table *var_remap)
{
nir_shader *inline_target = b->shader;
struct chit_miss_inlining_vars dst_vars;
init_chit_miss_inlining_vars(inline_target, &dst_vars);
init_chit_miss_inlining_vars(chit, vars);
dst_vars.tmax = params->trav_result->tmax;
dst_vars.primitive_addr = params->trav_result->primitive_addr;
dst_vars.primitive_id = params->trav_result->primitive_id;
dst_vars.geometry_id_and_flags = params->trav_result->geometry_id_and_flags;
dst_vars.instance_addr = params->trav_result->instance_addr;
dst_vars.hit_kind = params->trav_result->hit_kind;
nir_store_var(b, dst_vars.shader_record_ptr, params->sbt->shader_record_ptr, 0x1);
nir_store_var(b, dst_vars.origin, params->trav_params->origin, 0x7);
nir_store_var(b, dst_vars.direction, params->trav_params->direction, 0x7);
nir_store_var(b, dst_vars.tmin, params->trav_params->tmin, 0x1);
nir_store_var(b, dst_vars.cull_mask_and_flags, params->trav_params->cull_mask_and_flags, 0x1);
_mesa_hash_table_insert(var_remap, vars->shader_record_ptr, dst_vars.shader_record_ptr);
_mesa_hash_table_insert(var_remap, vars->origin, dst_vars.origin);
_mesa_hash_table_insert(var_remap, vars->direction, dst_vars.direction);
_mesa_hash_table_insert(var_remap, vars->tmin, dst_vars.tmin);
_mesa_hash_table_insert(var_remap, vars->tmax, dst_vars.tmax);
_mesa_hash_table_insert(var_remap, vars->primitive_addr, dst_vars.primitive_addr);
_mesa_hash_table_insert(var_remap, vars->primitive_id, dst_vars.primitive_id);
_mesa_hash_table_insert(var_remap, vars->geometry_id_and_flags, dst_vars.geometry_id_and_flags);
_mesa_hash_table_insert(var_remap, vars->cull_mask_and_flags, dst_vars.cull_mask_and_flags);
_mesa_hash_table_insert(var_remap, vars->instance_addr, dst_vars.instance_addr);
_mesa_hash_table_insert(var_remap, vars->hit_kind, dst_vars.hit_kind);
}
static bool
lower_rt_instruction_chit_miss(nir_builder *b, nir_intrinsic_instr *intr, void *_vars)
{
struct chit_miss_inlining_vars *vars = _vars;
b->cursor = nir_after_instr(&intr->instr);
nir_def *ret = NULL;
switch (intr->intrinsic) {
case nir_intrinsic_load_ray_world_origin:
ret = nir_load_var(b, vars->origin);
break;
case nir_intrinsic_load_ray_world_direction:
ret = nir_load_var(b, vars->direction);
break;
case nir_intrinsic_load_shader_record_ptr:
ret = nir_load_var(b, vars->shader_record_ptr);
break;
case nir_intrinsic_load_ray_t_max:
ret = nir_load_var(b, vars->tmax);
break;
case nir_intrinsic_load_ray_t_min:
ret = nir_load_var(b, vars->tmin);
break;
case nir_intrinsic_load_ray_instance_custom_index:
ret = radv_load_custom_instance(vars->device, b, nir_load_var(b, vars->instance_addr));
break;
case nir_intrinsic_load_primitive_id:
ret = nir_load_var(b, vars->primitive_id);
break;
case nir_intrinsic_load_instance_id:
ret = radv_load_instance_id(vars->device, b, nir_load_var(b, vars->instance_addr));
break;
case nir_intrinsic_load_ray_hit_kind:
ret = nir_load_var(b, vars->hit_kind);
break;
case nir_intrinsic_load_ray_flags:
ret = nir_iand_imm(b, nir_load_var(b, vars->cull_mask_and_flags), 0xFFFFFF);
break;
case nir_intrinsic_load_cull_mask:
ret = nir_ushr_imm(b, nir_load_var(b, vars->cull_mask_and_flags), 24);
break;
case nir_intrinsic_load_ray_geometry_index: {
ret = nir_load_var(b, vars->geometry_id_and_flags);
ret = nir_iand_imm(b, ret, 0xFFFFFFF);
break;
}
case nir_intrinsic_load_ray_world_to_object: {
unsigned c = nir_intrinsic_column(intr);
nir_def *instance_node_addr = nir_load_var(b, vars->instance_addr);
nir_def *wto_matrix[3];
radv_load_wto_matrix(vars->device, b, instance_node_addr, wto_matrix);
nir_def *vals[3];
for (unsigned i = 0; i < 3; ++i)
vals[i] = nir_channel(b, wto_matrix[i], c);
ret = nir_vec(b, vals, 3);
break;
}
case nir_intrinsic_load_ray_object_to_world: {
unsigned c = nir_intrinsic_column(intr);
nir_def *otw_matrix[3];
radv_load_otw_matrix(vars->device, b, nir_load_var(b, vars->instance_addr), otw_matrix);
ret = nir_vec3(b, nir_channel(b, otw_matrix[0], c), nir_channel(b, otw_matrix[1], c),
nir_channel(b, otw_matrix[2], c));
break;
}
case nir_intrinsic_load_ray_object_origin: {
nir_def *wto_matrix[3];
radv_load_wto_matrix(vars->device, b, nir_load_var(b, vars->instance_addr), wto_matrix);
ret = nir_build_vec3_mat_mult(b, nir_load_var(b, vars->origin), wto_matrix, true);
break;
}
case nir_intrinsic_load_ray_object_direction: {
nir_def *wto_matrix[3];
radv_load_wto_matrix(vars->device, b, nir_load_var(b, vars->instance_addr), wto_matrix);
ret = nir_build_vec3_mat_mult(b, nir_load_var(b, vars->direction), wto_matrix, false);
break;
}
case nir_intrinsic_load_ray_triangle_vertex_positions: {
nir_def *primitive_addr = nir_load_var(b, vars->primitive_addr);
ret = radv_load_vertex_position(vars->device, b, primitive_addr, nir_intrinsic_column(intr));
break;
}
default:
return false;
}
nir_def_replace(&intr->def, ret);
return true;
}
static void
radv_ray_tracing_group_chit_info(struct radv_ray_tracing_group *group, uint32_t *shader_index, uint32_t *handle_index,
struct radv_rt_case_data *data)
{
if (group->type != VK_RAY_TRACING_SHADER_GROUP_TYPE_GENERAL_KHR) {
*shader_index = group->recursive_shader;
*handle_index = group->handle.closest_hit_index;
}
}
static void
radv_ray_tracing_group_miss_info(struct radv_ray_tracing_group *group, uint32_t *shader_index, uint32_t *handle_index,
struct radv_rt_case_data *data)
{
if (group->type == VK_RAY_TRACING_SHADER_GROUP_TYPE_GENERAL_KHR) {
if (data->pipeline->stages[group->recursive_shader].stage != MESA_SHADER_MISS)
return;
*shader_index = group->recursive_shader;
*handle_index = group->handle.general_index;
}
}
static void
preprocess_shader_cb_monolithic(nir_shader *nir, void *_data)
{
uint32_t *payload_offset = _data;
NIR_PASS(_, nir, radv_nir_lower_ray_payload_derefs, *payload_offset);
}
void
radv_nir_lower_rt_io_monolithic(nir_shader *nir)
{
uint32_t raygen_payload_offset = 0;
preprocess_shader_cb_monolithic(nir, &raygen_payload_offset);
}
struct rt_variables {
struct radv_device *device;
const VkPipelineCreateFlags2 flags;
uint32_t payload_offset;
unsigned stack_size;
nir_def *launch_sizes[3];
nir_def *launch_ids[3];
nir_def *shader_record_ptr;
nir_variable *stack_ptr;
};
static void
radv_build_recursive_case(nir_builder *b, nir_def *idx, struct radv_ray_tracing_group *group,
struct radv_rt_case_data *data)
{
nir_shader *shader =
radv_pipeline_cache_handle_to_nir(data->device, data->pipeline->stages[group->recursive_shader].nir);
assert(shader);
struct chit_miss_inlining_params *params = data->param_data;
struct chit_miss_inlining_vars vars = {
.device = params->device,
};
struct hash_table *var_remap = _mesa_pointer_hash_table_create(NULL);
setup_chit_miss_inlining(&vars, params, b, shader, var_remap);
nir_opt_dead_cf(shader);
preprocess_shader_cb_monolithic(shader, &params->payload_offset);
nir_shader_intrinsics_pass(shader, lower_rt_instruction_chit_miss, nir_metadata_control_flow, &vars);
nir_lower_returns(shader);
nir_opt_dce(shader);
radv_nir_inline_constants(b->shader, shader);
nir_push_if(b, nir_ieq_imm(b, idx, group->handle.general_index));
nir_inline_function_impl(b, nir_shader_get_entrypoint(shader), NULL, var_remap);
nir_pop_if(b, NULL);
ralloc_free(shader);
}
struct lower_rt_instruction_monolithic_state {
struct radv_device *device;
struct radv_ray_tracing_pipeline *pipeline;
const VkRayTracingPipelineCreateInfoKHR *pCreateInfo;
struct rt_variables *vars;
};
static bool
lower_rt_call_monolithic(nir_builder *b, nir_intrinsic_instr *intr, void *data)
{
b->cursor = nir_after_instr(&intr->instr);
struct lower_rt_instruction_monolithic_state *state = data;
const struct radv_physical_device *pdev = radv_device_physical(state->device);
struct rt_variables *vars = state->vars;
switch (intr->intrinsic) {
case nir_intrinsic_execute_callable:
/* It's allowed to place OpExecuteCallableKHR in a SPIR-V, even if the RT pipeline doesn't contain
* any callable shaders. However, it's impossible to execute the instruction in a valid way, so just remove any
* nir_intrinsic_execute_callable we encounter.
*/
nir_instr_remove(&intr->instr);
return true;
case nir_intrinsic_trace_ray: {
vars->payload_offset = nir_src_as_uint(intr->src[10]);
nir_src cull_mask = intr->src[2];
bool ignore_cull_mask = nir_src_is_const(cull_mask) && (nir_src_as_uint(cull_mask) & 0xFF) == 0xFF;
/* Per the SPIR-V extension spec we have to ignore some bits for some arguments. */
struct radv_nir_rt_traversal_params params = {
.accel_struct = intr->src[0].ssa,
.cull_mask_and_flags = nir_ior(b, nir_ishl_imm(b, cull_mask.ssa, 24), intr->src[1].ssa),
.sbt_offset = nir_iand_imm(b, intr->src[3].ssa, 0xf),
.sbt_stride = nir_iand_imm(b, intr->src[4].ssa, 0xf),
.miss_index = nir_iand_imm(b, intr->src[5].ssa, 0xffff),
.origin = intr->src[6].ssa,
.tmin = intr->src[7].ssa,
.direction = intr->src[8].ssa,
.tmax = intr->src[9].ssa,
.ignore_cull_mask = ignore_cull_mask,
.preprocess_ahit_isec = preprocess_shader_cb_monolithic,
.cb_data = &vars->payload_offset,
};
nir_def *stack_ptr = nir_load_var(b, vars->stack_ptr);
nir_store_var(b, vars->stack_ptr, nir_iadd_imm(b, stack_ptr, b->shader->scratch_size), 0x1);
struct radv_nir_rt_traversal_result result =
radv_build_traversal(state->device, state->pipeline, b, &params, NULL);
nir_store_var(b, vars->stack_ptr, stack_ptr, 0x1);
struct chit_miss_inlining_params inline_params = {
.device = state->device,
.trav_params = &params,
.trav_result = &result,
.payload_offset = vars->payload_offset,
};
struct radv_rt_case_data case_data = {
.device = state->device,
.pipeline = state->pipeline,
.param_data = &inline_params,
};
nir_push_if(b, nir_load_var(b, result.hit));
{
struct radv_nir_sbt_data hit_sbt =
radv_nir_load_sbt_entry(b, nir_load_var(b, result.sbt_index), SBT_HIT, SBT_CLOSEST_HIT_IDX);
inline_params.sbt = &hit_sbt;
nir_def *should_return = nir_test_mask(b, params.cull_mask_and_flags, SpvRayFlagsSkipClosestHitShaderKHRMask);
/* should_return is set if we had a hit but we won't be calling the closest hit
* shader and hence need to return immediately to the calling shader. */
nir_push_if(b, nir_inot(b, should_return));
radv_visit_inlined_shaders(b, hit_sbt.shader_addr,
!(state->pipeline->base.base.create_flags &
VK_PIPELINE_CREATE_2_RAY_TRACING_NO_NULL_CLOSEST_HIT_SHADERS_BIT_KHR),
&case_data, radv_ray_tracing_group_chit_info, radv_build_recursive_case);
nir_pop_if(b, NULL);
}
nir_push_else(b, NULL);
{
struct radv_nir_sbt_data miss_sbt = radv_nir_load_sbt_entry(b, params.miss_index, SBT_MISS, SBT_GENERAL_IDX);
inline_params.sbt = &miss_sbt;
radv_visit_inlined_shaders(b, miss_sbt.shader_addr,
!(vars->flags & VK_PIPELINE_CREATE_2_RAY_TRACING_NO_NULL_MISS_SHADERS_BIT_KHR),
&case_data, radv_ray_tracing_group_miss_info, radv_build_recursive_case);
}
nir_pop_if(b, NULL);
b->shader->info.shared_size =
MAX2(b->shader->info.shared_size, pdev->rt_wave_size * MAX_STACK_ENTRY_COUNT * sizeof(uint32_t));
nir_instr_remove(&intr->instr);
return true;
}
default:
return false;
}
}
static bool
lower_rt_instruction_monolithic(nir_builder *b, nir_intrinsic_instr *intr, void *data)
{
b->cursor = nir_before_instr(&intr->instr);
struct lower_rt_instruction_monolithic_state *state = data;
struct rt_variables *vars = state->vars;
switch (intr->intrinsic) {
case nir_intrinsic_load_shader_record_ptr: {
nir_def_replace(&intr->def, vars->shader_record_ptr);
return true;
}
case nir_intrinsic_load_ray_launch_size: {
nir_def_replace(&intr->def, nir_vec(b, vars->launch_sizes, 3));
return true;
};
case nir_intrinsic_load_ray_launch_id: {
nir_def_replace(&intr->def, nir_vec(b, vars->launch_ids, 3));
return true;
}
case nir_intrinsic_load_scratch: {
nir_src_rewrite(&intr->src[0], nir_iadd_nuw(b, nir_load_var(b, vars->stack_ptr), intr->src[0].ssa));
return true;
}
case nir_intrinsic_store_scratch: {
nir_src_rewrite(&intr->src[1], nir_iadd_nuw(b, nir_load_var(b, vars->stack_ptr), intr->src[1].ssa));
return true;
}
default:
return false;
}
}
static bool
radv_count_hit_attrib_slots(nir_builder *b, nir_intrinsic_instr *instr, void *data)
{
uint32_t *count = data;
if (instr->intrinsic == nir_intrinsic_load_hit_attrib_amd || instr->intrinsic == nir_intrinsic_store_hit_attrib_amd)
*count = MAX2(*count, nir_intrinsic_base(instr) + 1);
return false;
}
void
radv_nir_lower_rt_abi_monolithic(nir_shader *shader, const struct radv_shader_args *args, uint32_t *stack_size,
struct radv_device *device, struct radv_ray_tracing_pipeline *pipeline)
{
nir_function_impl *impl = nir_shader_get_entrypoint(shader);
nir_builder b = nir_builder_at(nir_before_impl(impl));
struct rt_variables vars = {
.device = device,
.flags = pipeline->base.base.create_flags,
};
vars.stack_ptr = nir_variable_create(shader, nir_var_shader_temp, glsl_uint_type(), "stack_ptr");
for (uint32_t i = 0; i < ARRAY_SIZE(vars.launch_sizes); i++)
vars.launch_sizes[i] = ac_nir_load_arg(&b, &args->ac, args->ac.rt.launch_sizes[i]);
for (uint32_t i = 0; i < ARRAY_SIZE(vars.launch_sizes); i++) {
vars.launch_ids[i] = ac_nir_load_arg(&b, &args->ac, args->ac.rt.launch_ids[i]);
}
nir_def *record_ptr = ac_nir_load_arg(&b, &args->ac, args->ac.rt.shader_record);
vars.shader_record_ptr = nir_pack_64_2x32(&b, record_ptr);
nir_def *stack_ptr = ac_nir_load_arg(&b, &args->ac, args->ac.rt.dynamic_callable_stack_base);
nir_store_var(&b, vars.stack_ptr, stack_ptr, 0x1);
struct lower_rt_instruction_monolithic_state state = {
.device = device,
.pipeline = pipeline,
.vars = &vars,
};
nir_shader_intrinsics_pass(shader, lower_rt_call_monolithic, nir_metadata_none, &state);
nir_shader_intrinsics_pass(shader, lower_rt_instruction_monolithic, nir_metadata_none, &state);
nir_index_ssa_defs(impl);
uint32_t hit_attrib_count = 0;
nir_shader_intrinsics_pass(shader, radv_count_hit_attrib_slots, nir_metadata_all, &hit_attrib_count);
/* Register storage for hit attributes */
STACK_ARRAY(nir_variable *, hit_attribs, hit_attrib_count);
for (uint32_t i = 0; i < hit_attrib_count; i++)
hit_attribs[i] = nir_local_variable_create(impl, glsl_uint_type(), "ahit_attrib");
radv_nir_lower_hit_attribs(shader, hit_attribs, 0);
vars.stack_size = MAX2(vars.stack_size, shader->scratch_size);
*stack_size = MAX2(*stack_size, vars.stack_size);
shader->scratch_size = 0;
nir_progress(true, impl, nir_metadata_none);
/* cleanup passes */
NIR_PASS(_, shader, nir_lower_returns);
NIR_PASS(_, shader, nir_lower_global_vars_to_local);
NIR_PASS(_, shader, nir_lower_vars_to_ssa);
STACK_ARRAY_FINISH(hit_attribs);
}

18
src/amd/vulkan/nir/radv_nir_rt_stage_monolithic.h Normal file
View file

@ -0,0 +1,18 @@
/*
* Copyright © 2025 Valve Corporation
*
* SPDX-License-Identifier: MIT
*/
/* This file contains the public interface for all RT pipeline stage lowering. */
#ifndef RADV_NIR_RT_STAGE_MONOLITHIC_H
#define RADV_NIR_RT_STAGE_MONOLITHIC_H
#include "radv_pipeline_rt.h"
void radv_nir_lower_rt_abi_monolithic(nir_shader *shader, const struct radv_shader_args *args, uint32_t *stack_size,
struct radv_device *device, struct radv_ray_tracing_pipeline *pipeline);
void radv_nir_lower_rt_io_monolithic(nir_shader *shader);
#endif // RADV_NIR_RT_STAGE_MONOLITHIC_H

1228
src/amd/vulkan/nir/radv_nir_rt_traversal_shader.c Normal file
View file

File diff suppressed because it is too large Load diff

15
src/amd/vulkan/nir/radv_nir_rt_traversal_shader.h Normal file
View file

@ -0,0 +1,15 @@
/*
* Copyright © 2025 Valve Corporation
*
* SPDX-License-Identifier: MIT
*/
#ifndef RADV_NIR_RT_TRAVERSAL_SHADER_H
#define RADV_NIR_RT_TRAVERSAL_SHADER_H
#include "radv_pipeline_rt.h"
nir_shader *radv_build_traversal_shader(struct radv_device *device, struct radv_ray_tracing_pipeline *pipeline,
struct radv_ray_tracing_stage_info *info);
#endif // RADV_NIR_RT_TRAVERSAL_SHADER_H

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

@ -12,6 +12,9 @@
#include "vk_shader_module.h"
#include "nir/radv_nir.h"
#include "nir/radv_nir_rt_stage_cps.h"
#include "nir/radv_nir_rt_stage_monolithic.h"
#include "nir/radv_nir_rt_traversal_shader.h"
#include "ac_nir.h"
#include "radv_debug.h"
#include "radv_descriptor_set.h"
@ -360,9 +363,8 @@ move_rt_instructions(nir_shader *shader)
static VkResult
radv_rt_nir_to_asm(struct radv_device *device, struct vk_pipeline_cache *cache,
const VkRayTracingPipelineCreateInfoKHR *pCreateInfo, struct radv_ray_tracing_pipeline *pipeline,
bool monolithic, struct radv_shader_stage *stage, uint32_t *stack_size,
struct radv_ray_tracing_stage_info *stage_info,
struct radv_ray_tracing_pipeline *pipeline, bool monolithic, struct radv_shader_stage *stage,
uint32_t *stack_size, struct radv_ray_tracing_stage_info *stage_info,
const struct radv_ray_tracing_stage_info *traversal_stage_info,
struct radv_serialized_shader_arena_block *replay_block, bool skip_shaders_cache,
bool has_position_fetch, struct radv_shader **out_shader)
@ -374,7 +376,10 @@ radv_rt_nir_to_asm(struct radv_device *device, struct vk_pipeline_cache *cache,
bool keep_executable_info = radv_pipeline_capture_shaders(device, pipeline->base.base.create_flags);
bool keep_statistic_info = radv_pipeline_capture_shader_stats(device, pipeline->base.base.create_flags);
radv_nir_lower_rt_io(stage->nir, monolithic, 0);
if (monolithic)
radv_nir_lower_rt_io_monolithic(stage->nir);
else
radv_nir_lower_rt_io_cps(stage->nir);
/* Gather shader info. */
nir_shader_gather_info(stage->nir, nir_shader_get_entrypoint(stage->nir));
@ -427,8 +432,13 @@ radv_rt_nir_to_asm(struct radv_device *device, struct vk_pipeline_cache *cache,
for (uint32_t i = 0; i < num_shaders; i++) {
struct radv_shader_stage temp_stage = *stage;
temp_stage.nir = shaders[i];
radv_nir_lower_rt_abi(temp_stage.nir, pCreateInfo, &temp_stage.args, &stage->info, stack_size, i > 0, device,
pipeline, monolithic, has_position_fetch, traversal_stage_info);
if (monolithic) {
assert(num_shaders == 1);
radv_nir_lower_rt_abi_monolithic(temp_stage.nir, &temp_stage.args, stack_size, device, pipeline);
} else {
radv_nir_lower_rt_abi_cps(temp_stage.nir, &temp_stage.args, &stage->info, stack_size, i > 0, device, pipeline,
has_position_fetch, traversal_stage_info);
}
/* Info might be out-of-date after inlining in radv_nir_lower_rt_abi(). */
nir_shader_gather_info(temp_stage.nir, nir_shader_get_entrypoint(temp_stage.nir));
@ -705,8 +715,8 @@ radv_rt_compile_shaders(struct radv_device *device, struct vk_pipeline_cache *ca
bool monolithic_raygen = monolithic && stage->stage == MESA_SHADER_RAYGEN;
result = radv_rt_nir_to_asm(device, cache, pCreateInfo, pipeline, monolithic_raygen, stage, &stack_size,
&rt_stages[idx].info, NULL, replay_block, skip_shaders_cache, has_position_fetch,
result =
radv_rt_nir_to_asm(device, cache, pipeline, monolithic_raygen, stage, &stack_size, &rt_stages[idx].info, NULL, replay_block, skip_shaders_cache, has_position_fetch,
&rt_stages[idx].shader);
if (result != VK_SUCCESS)
goto cleanup;
@ -758,15 +768,15 @@ radv_rt_compile_shaders(struct radv_device *device, struct vk_pipeline_cache *ca
}
/* create traversal shader */
nir_shader *traversal_nir = radv_build_traversal_shader(device, pipeline, pCreateInfo, &traversal_info);
nir_shader *traversal_nir = radv_build_traversal_shader(device, pipeline, &traversal_info);
struct radv_shader_stage traversal_stage = {
.stage = MESA_SHADER_INTERSECTION,
.nir = traversal_nir,
.key = stage_keys[MESA_SHADER_INTERSECTION],
};
radv_shader_layout_init(pipeline_layout, MESA_SHADER_INTERSECTION, &traversal_stage.layout);
result = radv_rt_nir_to_asm(device, cache, pCreateInfo, pipeline, false, &traversal_stage, NULL, NULL,
&traversal_info, NULL, skip_shaders_cache, has_position_fetch,
result = radv_rt_nir_to_asm(device, cache, pipeline, false, &traversal_stage, NULL, NULL, &traversal_info, NULL,
skip_shaders_cache, has_position_fetch,
&pipeline->base.base.shaders[MESA_SHADER_INTERSECTION]);
ralloc_free(traversal_nir);

16
src/amd/vulkan/radv_shader.h
View file

@ -492,17 +492,6 @@ void radv_optimize_nir_algebraic_late(nir_shader *shader);
void radv_optimize_nir_algebraic(nir_shader *shader, bool opt_offsets, bool opt_mqsad,
enum amd_gfx_level gfx_level);
void radv_nir_lower_rt_io(nir_shader *shader, bool monolithic, uint32_t payload_offset);
struct radv_ray_tracing_stage_info;
void radv_nir_lower_rt_abi(nir_shader *shader, const VkRayTracingPipelineCreateInfoKHR *pCreateInfo,
const struct radv_shader_args *args, const struct radv_shader_info *info,
uint32_t *stack_size, bool resume_shader, struct radv_device *device,
struct radv_ray_tracing_pipeline *pipeline, bool monolithic, bool has_position_fetch,
const struct radv_ray_tracing_stage_info *traversal_info);
void radv_gather_unused_args(struct radv_ray_tracing_stage_info *info, nir_shader *nir);
struct radv_shader_stage;
@ -658,11 +647,6 @@ bool radv_consider_culling(const struct radv_physical_device *pdev, struct nir_s
void radv_get_nir_options(struct radv_physical_device *pdev);
struct radv_ray_tracing_stage_info;
nir_shader *radv_build_traversal_shader(struct radv_device *device, struct radv_ray_tracing_pipeline *pipeline,
const VkRayTracingPipelineCreateInfoKHR *pCreateInfo,
struct radv_ray_tracing_stage_info *info);
enum radv_rt_priority {
radv_rt_priority_raygen = 0,