VkPipelineShaderStageCreateInfo stage = {
.sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
.pNext = NULL,
- .stage = VK_SHADER_STAGE_COMPUTE_BIT,
+ .stage = VK_SHADER_STAGE_RAYGEN_BIT_KHR,
.module = vk_shader_module_to_handle(&module),
.pName = "main",
};
goto shader_fail;
}
- radv_compute_pipeline_init(&rt_pipeline->base, pipeline_layout);
-
rt_pipeline->stack_size = compute_rt_stack_size(pCreateInfo, rt_pipeline->groups);
rt_pipeline->base.base.shaders[MESA_SHADER_COMPUTE] = radv_create_rt_prolog(device);
- *pPipeline = radv_pipeline_to_handle(&rt_pipeline->base.base);
+ combine_config(&rt_pipeline->base.base.shaders[MESA_SHADER_COMPUTE]->config,
+ &rt_pipeline->base.base.shaders[MESA_SHADER_RAYGEN]->config);
+
+ postprocess_rt_config(&rt_pipeline->base.base.shaders[MESA_SHADER_COMPUTE]->config,
+ device->physical_device->rt_wave_size);
+ radv_compute_pipeline_init(&rt_pipeline->base, pipeline_layout);
+
+ *pPipeline = radv_pipeline_to_handle(&rt_pipeline->base.base);
shader_fail:
ralloc_free(shader);
pipeline_fail:
/* global address of the SBT entry used for the shader */
nir_variable *shader_record_ptr;
- nir_variable *launch_size;
- nir_variable *launch_id;
-
/* trace_ray arguments */
nir_variable *accel_struct;
nir_variable *cull_mask_and_flags;
vars.shader_record_ptr =
nir_variable_create(shader, nir_var_shader_temp, glsl_uint64_t_type(), "shader_record_ptr");
- const struct glsl_type *uvec3_type = glsl_vector_type(GLSL_TYPE_UINT, 3);
- vars.launch_size = nir_variable_create(shader, nir_var_shader_temp, uvec3_type, "launch_size");
- vars.launch_id = nir_variable_create(shader, nir_var_shader_temp, uvec3_type, "launch_id");
-
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");
_mesa_hash_table_insert(var_remap, src->arg, dst->arg);
_mesa_hash_table_insert(var_remap, src->stack_ptr, dst->stack_ptr);
_mesa_hash_table_insert(var_remap, src->shader_record_ptr, dst->shader_record_ptr);
- _mesa_hash_table_insert(var_remap, src->launch_size, dst->launch_size);
- _mesa_hash_table_insert(var_remap, src->launch_id, dst->launch_id);
_mesa_hash_table_insert(var_remap, src->accel_struct, dst->accel_struct);
_mesa_hash_table_insert(var_remap, src->cull_mask_and_flags, dst->cull_mask_and_flags);
ret = nir_load_var(&b_shader, vars->shader_record_ptr);
break;
}
- case nir_intrinsic_load_ray_launch_id: {
- ret = nir_load_var(&b_shader, vars->launch_id);
- break;
- }
- case nir_intrinsic_load_ray_launch_size: {
- ret = nir_load_var(&b_shader, vars->launch_size);
- break;
- }
case nir_intrinsic_load_ray_t_min: {
ret = nir_load_var(&b_shader, vars->tmin);
break;
nir_store_var(&b, vars.tmax, nir_load_ray_t_max(&b), 0x1);
nir_store_var(&b, vars.arg, nir_load_rt_arg_scratch_offset_amd(&b), 0x1);
nir_store_var(&b, vars.stack_ptr, nir_imm_int(&b, 0), 0x1);
- nir_store_var(&b, vars.launch_size, nir_load_ray_launch_size(&b), 0x7);
- nir_store_var(&b, vars.launch_id, nir_load_ray_launch_id(&b), 0x7);
struct rt_traversal_vars trav_vars = init_traversal_vars(&b);
create_rt_shader(struct radv_device *device, const VkRayTracingPipelineCreateInfoKHR *pCreateInfo,
struct radv_ray_tracing_module *groups, const struct radv_pipeline_key *key)
{
- nir_builder b = radv_meta_init_shader(device, MESA_SHADER_COMPUTE, "rt_combined");
+ nir_builder b = radv_meta_init_shader(device, MESA_SHADER_RAYGEN, "rt_combined");
b.shader->info.internal = false;
b.shader->info.workgroup_size[0] = 8;
b.shader->info.workgroup_size[1] = device->physical_device->rt_wave_size == 64 ? 8 : 4;
load_sbt_entry(&b, &vars, nir_imm_int(&b, 0), SBT_RAYGEN, SBT_GENERAL_IDX);
nir_store_var(&b, vars.stack_ptr, nir_load_rt_dynamic_callable_stack_base_amd(&b), 0x1);
- nir_store_var(&b, vars.launch_id, nir_load_global_invocation_id(&b, 32), 0x7);
- nir_ssa_def *launch_size_addr = nir_load_ray_launch_size_addr_amd(&b);
- nir_ssa_def *xy = nir_build_load_smem_amd(&b, 2, launch_size_addr, nir_imm_int(&b, 0));
- nir_ssa_def *z = nir_build_load_smem_amd(&b, 1, launch_size_addr, nir_imm_int(&b, 8));
- nir_ssa_def *xyz[3] = {
- nir_channel(&b, xy, 0),
- nir_channel(&b, xy, 1),
- z,
- };
- nir_store_var(&b, vars.launch_size, nir_vec(&b, xyz, 3), 0x7);
-
nir_loop *loop = nir_push_loop(&b);
nir_ssa_def *idx = nir_load_var(&b, vars.idx);
case MESA_SHADER_TASK:
declare_global_input_sgprs(info, &user_sgpr_info, args);
- if (info->cs.is_rt_shader) {
- ac_add_arg(&args->ac, AC_ARG_SGPR, 2, AC_ARG_CONST_PTR, &args->ac.sbt_descriptors);
- }
-
if (info->cs.uses_grid_size) {
if (args->load_grid_size_from_user_sgpr)
ac_add_arg(&args->ac, AC_ARG_SGPR, 3, AC_ARG_INT, &args->ac.num_work_groups);
ac_add_arg(&args->ac, AC_ARG_SGPR, 2, AC_ARG_CONST_PTR, &args->ac.num_work_groups);
}
- if (info->cs.uses_ray_launch_size) {
+ if (info->cs.is_rt_shader) {
+ ac_add_arg(&args->ac, AC_ARG_SGPR, 2, AC_ARG_CONST_DESC_PTR, &args->ac.sbt_descriptors);
ac_add_arg(&args->ac, AC_ARG_SGPR, 2, AC_ARG_CONST_PTR, &args->ac.ray_launch_size_addr);
- }
-
- if (info->cs.uses_dynamic_rt_callable_stack) {
- ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT,
- &args->ac.rt_dynamic_callable_stack_base);
ac_add_arg(&args->ac, AC_ARG_SGPR, 2, AC_ARG_CONST_PTR,
&args->ac.rt_traversal_shader_addr);
+ ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT,
+ &args->ac.rt_dynamic_callable_stack_base);
}
if (info->vs.needs_draw_id) {
if (args->ac.ray_launch_size_addr.used) {
set_loc_shader_ptr(args, AC_UD_CS_RAY_LAUNCH_SIZE_ADDR, &user_sgpr_idx);
}
- if (args->ac.rt_dynamic_callable_stack_base.used) {
- set_loc_shader(args, AC_UD_CS_RAY_DYNAMIC_CALLABLE_STACK_BASE, &user_sgpr_idx, 1);
- }
if (args->ac.rt_traversal_shader_addr.used) {
set_loc_shader_ptr(args, AC_UD_CS_TRAVERSAL_SHADER_ADDR, &user_sgpr_idx);
}
+ if (args->ac.rt_dynamic_callable_stack_base.used) {
+ set_loc_shader(args, AC_UD_CS_RAY_DYNAMIC_CALLABLE_STACK_BASE, &user_sgpr_idx, 1);
+ }
if (args->ac.draw_id.used) {
set_loc_shader(args, AC_UD_CS_TASK_DRAW_ID, &user_sgpr_idx, 1);
}