archive-gh-me/wgpu
1
0
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2025-12-08 21:26:17 +00:00
Code Issues Projects Releases Wiki Activity

[spirv] Make ray queries safer (#8390)

Browse Source
This commit is contained in:
Vecvec 2025-11-13 18:33:32 +13:00 committed by GitHub
parent a05c70cef7
commit efca3f5066
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
24 changed files with 3755 additions and 585 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
CHANGELOG.md
View File

@ -125,6 +125,10 @@ By @SupaMaggie70Incorporated in [#8206](https://github.com/gfx-rs/wgpu/pull/8206
- `util::StagingBelt` now takes a `Device` when it is created instead of when it is used. By @kpreid in [#8462](https://github.com/gfx-rs/wgpu/pull/8462).
- `wgpu_hal::vulkan::Device::texture_from_raw` now takes an `external_memory` argument. By @s-ol in [#8512](https://github.com/gfx-rs/wgpu/pull/8512)
#### Naga
- Prevent UB with invalid ray query calls on spirv. By @Vecvec in [#8390](https://github.com/gfx-rs/wgpu/pull/8390).
### Bug Fixes
#### naga

3
naga-test/src/lib.rs
View File

@ -114,6 +114,7 @@ pub struct SpirvOutParameters {
pub separate_entry_points: bool,
#[serde(deserialize_with = "deserialize_binding_map")]
pub binding_map: naga::back::spv::BindingMap,
pub ray_query_initialization_tracking: bool,
pub use_storage_input_output_16: bool,
}
impl Default for SpirvOutParameters {
@ -126,6 +127,7 @@ impl Default for SpirvOutParameters {
force_point_size: false,
clamp_frag_depth: false,
separate_entry_points: false,
ray_query_initialization_tracking: true,
use_storage_input_output_16: true,
binding_map: naga::back::spv::BindingMap::default(),
}
@ -159,6 +161,7 @@ impl SpirvOutParameters {
binding_map: self.binding_map.clone(),
zero_initialize_workgroup_memory: spv::ZeroInitializeWorkgroupMemoryMode::Polyfill,
force_loop_bounding: true,
ray_query_initialization_tracking: true,
debug_info,
use_storage_input_output_16: self.use_storage_input_output_16,
}

60
naga/src/back/spv/block.rs
View File

@ -203,7 +203,7 @@ impl Writer {
));
let clamp_id = self.id_gen.next();
body.push(Instruction::ext_inst(
body.push(Instruction::ext_inst_gl_op(
self.gl450_ext_inst_id,
spirv::GLOp::FClamp,
float_type_id,
@ -1026,7 +1026,7 @@ impl BlockContext<'_> {
};
let max_id = self.gen_id();
block.body.push(Instruction::ext_inst(
block.body.push(Instruction::ext_inst_gl_op(
self.writer.gl450_ext_inst_id,
max_op,
result_type_id,
@ -1034,7 +1034,7 @@ impl BlockContext<'_> {
&[arg0_id, arg1_id],
));
MathOp::Custom(Instruction::ext_inst(
MathOp::Custom(Instruction::ext_inst_gl_op(
self.writer.gl450_ext_inst_id,
min_op,
result_type_id,
@ -1068,7 +1068,7 @@ impl BlockContext<'_> {
arg2_id = self.writer.get_constant_composite(ty, &self.temp_list);
}
MathOp::Custom(Instruction::ext_inst(
MathOp::Custom(Instruction::ext_inst_gl_op(
self.writer.gl450_ext_inst_id,
spirv::GLOp::FClamp,
result_type_id,
@ -1282,7 +1282,7 @@ impl BlockContext<'_> {
&self.temp_list,
));
MathOp::Custom(Instruction::ext_inst(
MathOp::Custom(Instruction::ext_inst_gl_op(
self.writer.gl450_ext_inst_id,
spirv::GLOp::FMix,
result_type_id,
@ -1339,7 +1339,7 @@ impl BlockContext<'_> {
};
let lsb_id = self.gen_id();
block.body.push(Instruction::ext_inst(
block.body.push(Instruction::ext_inst_gl_op(
self.writer.gl450_ext_inst_id,
spirv::GLOp::FindILsb,
result_type_id,
@ -1347,7 +1347,7 @@ impl BlockContext<'_> {
&[arg0_id],
));
MathOp::Custom(Instruction::ext_inst(
MathOp::Custom(Instruction::ext_inst_gl_op(
self.writer.gl450_ext_inst_id,
spirv::GLOp::UMin,
result_type_id,
@ -1388,7 +1388,7 @@ impl BlockContext<'_> {
};
let msb_id = self.gen_id();
block.body.push(Instruction::ext_inst(
block.body.push(Instruction::ext_inst_gl_op(
self.writer.gl450_ext_inst_id,
if width != 4 {
spirv::GLOp::FindILsb
@ -1445,7 +1445,7 @@ impl BlockContext<'_> {
// o = min(offset, w)
let offset_id = self.gen_id();
block.body.push(Instruction::ext_inst(
block.body.push(Instruction::ext_inst_gl_op(
self.writer.gl450_ext_inst_id,
spirv::GLOp::UMin,
u32_type,
@ -1465,7 +1465,7 @@ impl BlockContext<'_> {
// c = min(count, tmp)
let count_id = self.gen_id();
block.body.push(Instruction::ext_inst(
block.body.push(Instruction::ext_inst_gl_op(
self.writer.gl450_ext_inst_id,
spirv::GLOp::UMin,
u32_type,
@ -1495,7 +1495,7 @@ impl BlockContext<'_> {
// o = min(offset, w)
let offset_id = self.gen_id();
block.body.push(Instruction::ext_inst(
block.body.push(Instruction::ext_inst_gl_op(
self.writer.gl450_ext_inst_id,
spirv::GLOp::UMin,
u32_type,
@ -1515,7 +1515,7 @@ impl BlockContext<'_> {
// c = min(count, tmp)
let count_id = self.gen_id();
block.body.push(Instruction::ext_inst(
block.body.push(Instruction::ext_inst_gl_op(
self.writer.gl450_ext_inst_id,
spirv::GLOp::UMin,
u32_type,
@ -1610,7 +1610,7 @@ impl BlockContext<'_> {
};
block.body.push(match math_op {
MathOp::Ext(op) => Instruction::ext_inst(
MathOp::Ext(op) => Instruction::ext_inst_gl_op(
self.writer.gl450_ext_inst_id,
op,
result_type_id,
@ -1621,7 +1621,27 @@ impl BlockContext<'_> {
});
id
}
crate::Expression::LocalVariable(variable) => self.function.variables[&variable].id,
crate::Expression::LocalVariable(variable) => {
if let Some(rq_tracker) = self
.function
.ray_query_initialization_tracker_variables
.get(&variable)
{
self.ray_query_tracker_expr.insert(
expr_handle,
super::RayQueryTrackers {
initialized_tracker: rq_tracker.id,
t_max_tracker: self
.function
.ray_query_t_max_tracker_variables
.get(&variable)
.expect("Both trackers are set at the same time.")
.id,
},
);
}
self.function.variables[&variable].id
}
crate::Expression::Load { pointer } => {
self.write_checked_load(pointer, block, AccessTypeAdjustment::None, result_type_id)?
}
@ -1772,6 +1792,10 @@ impl BlockContext<'_> {
crate::Expression::ArrayLength(expr) => self.write_runtime_array_length(expr, block)?,
crate::Expression::RayQueryGetIntersection { query, committed } => {
let query_id = self.cached[query];
let init_tracker_id = *self
.ray_query_tracker_expr
.get(&query)
.expect("not a cached ray query");
let func_id = self
.writer
.write_ray_query_get_intersection_function(committed, self.ir_module);
@ -1782,7 +1806,7 @@ impl BlockContext<'_> {
intersection_type_id,
id,
func_id,
&[query_id],
&[query_id, init_tracker_id.initialized_tracker],
));
id
}
@ -2008,7 +2032,7 @@ impl BlockContext<'_> {
let max_const_id = maybe_splat_const(self.writer, max_const_id);
let clamp_id = self.gen_id();
block.body.push(Instruction::ext_inst(
block.body.push(Instruction::ext_inst_gl_op(
self.writer.gl450_ext_inst_id,
spirv::GLOp::FClamp,
expr_type_id,
@ -2671,7 +2695,7 @@ impl BlockContext<'_> {
});
let clamp_id = self.gen_id();
block.body.push(Instruction::ext_inst(
block.body.push(Instruction::ext_inst_gl_op(
self.writer.gl450_ext_inst_id,
clamp_op,
wide_vector_type_id,
@ -2765,7 +2789,7 @@ impl BlockContext<'_> {
let [min, max] = [min, max].map(|lit| self.writer.get_constant_scalar(lit));
let clamp_id = self.gen_id();
block.body.push(Instruction::ext_inst(
block.body.push(Instruction::ext_inst_gl_op(
self.writer.gl450_ext_inst_id,
clamp_op,
result_type_id,

6
naga/src/back/spv/image.rs
View File

@ -446,7 +446,7 @@ impl BlockContext<'_> {
// and negative values in a single instruction: negative values of
// `input_id` get treated as very large positive values.
let restricted_id = self.gen_id();
block.body.push(Instruction::ext_inst(
block.body.push(Instruction::ext_inst_gl_op(
self.writer.gl450_ext_inst_id,
spirv::GLOp::UMin,
type_id,
@ -580,7 +580,7 @@ impl BlockContext<'_> {
// and negative values in a single instruction: negative values of
// `coordinates` get treated as very large positive values.
let restricted_coordinates_id = self.gen_id();
block.body.push(Instruction::ext_inst(
block.body.push(Instruction::ext_inst_gl_op(
self.writer.gl450_ext_inst_id,
spirv::GLOp::UMin,
coordinates.type_id,
@ -923,7 +923,7 @@ impl BlockContext<'_> {
// Clamp the coords to the calculated margins
let clamped_coords_id = self.gen_id();
block.body.push(Instruction::ext_inst(
block.body.push(Instruction::ext_inst_gl_op(
self.writer.gl450_ext_inst_id,
spirv::GLOp::NClamp,
vec2f_type_id,

2
naga/src/back/spv/index.rs
View File

@ -366,7 +366,7 @@ impl BlockContext<'_> {
// One or the other of the index or length is dynamic, so emit code for
// BoundsCheckPolicy::Restrict.
let restricted_index_id = self.gen_id();
block.body.push(Instruction::ext_inst(
block.body.push(Instruction::ext_inst_gl_op(
self.writer.gl450_ext_inst_id,
spirv::GLOp::UMin,
self.writer.get_u32_type_id(),

22
naga/src/back/spv/instructions.rs
View File

@ -156,18 +156,28 @@ impl super::Instruction {
instruction
}
pub(super) fn ext_inst(
pub(super) fn ext_inst_gl_op(
set_id: Word,
op: spirv::GLOp,
result_type_id: Word,
id: Word,
operands: &[Word],
) -> Self {
Self::ext_inst(set_id, op as u32, result_type_id, id, operands)
}
pub(super) fn ext_inst(
set_id: Word,
op: u32,
result_type_id: Word,
id: Word,
operands: &[Word],
) -> Self {
let mut instruction = Self::new(Op::ExtInst);
instruction.set_type(result_type_id);
instruction.set_result(id);
instruction.add_operand(set_id);
instruction.add_operand(op as u32);
instruction.add_operand(op);
for operand in operands {
instruction.add_operand(*operand)
}
@ -824,6 +834,14 @@ impl super::Instruction {
instruction
}
pub(super) fn ray_query_get_t_min(result_type_id: Word, id: Word, query: Word) -> Self {
let mut instruction = Self::new(Op::RayQueryGetRayTMinKHR);
instruction.set_type(result_type_id);
instruction.set_result(id);
instruction.add_operand(query);
instruction
}
//
// Conversion Instructions
//

64
naga/src/back/spv/mod.rs
View File

@ -151,6 +151,12 @@ struct Function {
signature: Option<Instruction>,
parameters: Vec<FunctionArgument>,
variables: crate::FastHashMap<Handle<crate::LocalVariable>, LocalVariable>,
/// Map from a local variable that is a ray query to its u32 tracker.
ray_query_initialization_tracker_variables:
crate::FastHashMap<Handle<crate::LocalVariable>, LocalVariable>,
/// Map from a local variable that is a ray query to its tracker for the t max.
ray_query_t_max_tracker_variables:
crate::FastHashMap<Handle<crate::LocalVariable>, LocalVariable>,
/// List of local variables used as a counters to ensure that all loops are bounded.
force_loop_bounding_vars: Vec<LocalVariable>,
@ -445,6 +451,16 @@ struct LookupFunctionType {
return_type_id: Word,
}
#[derive(Debug, PartialEq, Clone, Hash, Eq)]
enum LookupRayQueryFunction {
Initialize,
Proceed,
GenerateIntersection,
ConfirmIntersection,
GetVertexPositions { committed: bool },
GetIntersection { committed: bool },
}
#[derive(Debug)]
enum Dimension {
Scalar,
@ -685,6 +701,21 @@ struct BlockContext<'w> {
expression_constness: ExpressionConstnessTracker,
force_loop_bounding: bool,
/// Hash from an expression whose type is a ray query / pointer to a ray query to its tracker.
/// Note: this is sparse, so can't be a handle vec
ray_query_tracker_expr: crate::FastHashMap<Handle<crate::Expression>, RayQueryTrackers>,
}
#[derive(Clone, Copy)]
struct RayQueryTrackers {
// Initialization tracker
initialized_tracker: Word,
// Tracks the t max from ray query initialize.
// Unlike HLSL, spir-v's equivalent getter for the current committed t has UB (instead of just
// returning t_max) if there was no previous hit (though in some places it treats the behaviour as
// defined), therefore we must track the tmax inputted into ray query initialize.
t_max_tracker: Word,
}
impl BlockContext<'_> {
@ -741,6 +772,7 @@ pub struct Writer {
/// The set of spirv extensions used.
extensions_used: crate::FastIndexSet<&'static str>,
debug_strings: Vec<Instruction>,
debugs: Vec<Instruction>,
annotations: Vec<Instruction>,
flags: WriterFlags,
@ -773,12 +805,15 @@ pub struct Writer {
// Just a temporary list of SPIR-V ids
temp_list: Vec<Word>,
ray_get_committed_intersection_function: Option<Word>,
ray_get_candidate_intersection_function: Option<Word>,
ray_query_functions: crate::FastHashMap<LookupRayQueryFunction, Word>,
/// F16 I/O polyfill manager for handling `f16` input/output variables
/// when `StorageInputOutput16` capability is not available.
io_f16_polyfills: f16_polyfill::F16IoPolyfill,
/// Non semantic debug printf extension `OpExtInstImport`
debug_printf: Option<Word>,
pub(crate) ray_query_initialization_tracking: bool,
}
bitflags::bitflags! {
@ -810,6 +845,26 @@ bitflags::bitflags! {
///
/// [`BuiltIn::FragDepth`]: crate::BuiltIn::FragDepth
const CLAMP_FRAG_DEPTH = 0x10;
/// Instead of silently failing if the arguments to generate a ray query are
/// invalid, uses debug printf extension to print to the command line
///
/// Note: VK_KHR_shader_non_semantic_info must be enabled. This will have no
/// effect if `options.ray_query_initialization_tracking` is set to false.
const PRINT_ON_RAY_QUERY_INITIALIZATION_FAIL = 0x20;
}
}
bitflags::bitflags! {
/// How far through a ray query are we
#[derive(Clone, Copy, Debug, Eq, PartialEq)]
pub(super) struct RayQueryPoint: u32 {
/// Ray query has been successfully initialized.
const INITIALIZED = 1 << 0;
/// Proceed has been called on ray query.
const PROCEED = 1 << 1;
/// Proceed has returned false (have finished traversal).
const FINISHED_TRAVERSAL = 1 << 2;
}
}
@ -867,6 +922,10 @@ pub struct Options<'a> {
/// to think the number of iterations is bounded.
pub force_loop_bounding: bool,
/// if set, ray queries will get a variable to track their state to prevent
/// misuse.
pub ray_query_initialization_tracking: bool,
/// Whether to use the `StorageInputOutput16` capability for `f16` shader I/O.
/// When false, `f16` I/O is polyfilled using `f32` types with conversions.
pub use_storage_input_output_16: bool,
@ -891,6 +950,7 @@ impl Default for Options<'_> {
bounds_check_policies: BoundsCheckPolicies::default(),
zero_initialize_workgroup_memory: ZeroInitializeWorkgroupMemoryMode::Polyfill,
force_loop_bounding: true,
ray_query_initialization_tracking: true,
use_storage_input_output_16: true,
debug_info: None,
}

1607
naga/src/back/spv/ray.rs
View File

File diff suppressed because it is too large Load Diff

109
naga/src/back/spv/writer.rs
View File

@ -35,6 +35,12 @@ impl Function {
for local_var in self.variables.values() {
local_var.instruction.to_words(sink);
}
for local_var in self.ray_query_initialization_tracker_variables.values() {
local_var.instruction.to_words(sink);
}
for local_var in self.ray_query_t_max_tracker_variables.values() {
local_var.instruction.to_words(sink);
}
for local_var in self.force_loop_bounding_vars.iter() {
local_var.instruction.to_words(sink);
}
@ -71,12 +77,14 @@ impl Writer {
capabilities_available: options.capabilities.clone(),
capabilities_used,
extensions_used: crate::FastIndexSet::default(),
debug_strings: vec![],
debugs: vec![],
annotations: vec![],
flags: options.flags,
bounds_check_policies: options.bounds_check_policies,
zero_initialize_workgroup_memory: options.zero_initialize_workgroup_memory,
force_loop_bounding: options.force_loop_bounding,
ray_query_initialization_tracking: options.ray_query_initialization_tracking,
use_storage_input_output_16: options.use_storage_input_output_16,
void_type,
lookup_type: crate::FastHashMap::default(),
@ -91,11 +99,11 @@ impl Writer {
saved_cached: CachedExpressions::default(),
gl450_ext_inst_id,
temp_list: Vec::new(),
ray_get_committed_intersection_function: None,
ray_get_candidate_intersection_function: None,
ray_query_functions: crate::FastHashMap::default(),
io_f16_polyfills: super::f16_polyfill::F16IoPolyfill::new(
options.use_storage_input_output_16,
),
debug_printf: None,
})
}
@ -147,6 +155,7 @@ impl Writer {
bounds_check_policies: self.bounds_check_policies,
zero_initialize_workgroup_memory: self.zero_initialize_workgroup_memory,
force_loop_bounding: self.force_loop_bounding,
ray_query_initialization_tracking: self.ray_query_initialization_tracking,
use_storage_input_output_16: self.use_storage_input_output_16,
capabilities_available: take(&mut self.capabilities_available),
fake_missing_bindings: self.fake_missing_bindings,
@ -162,6 +171,7 @@ impl Writer {
extensions_used: take(&mut self.extensions_used).recycle(),
physical_layout: self.physical_layout.clone().recycle(),
logical_layout: take(&mut self.logical_layout).recycle(),
debug_strings: take(&mut self.debug_strings).recycle(),
debugs: take(&mut self.debugs).recycle(),
annotations: take(&mut self.annotations).recycle(),
lookup_type: take(&mut self.lookup_type).recycle(),
@ -173,9 +183,9 @@ impl Writer {
global_variables: take(&mut self.global_variables).recycle(),
saved_cached: take(&mut self.saved_cached).recycle(),
temp_list: take(&mut self.temp_list).recycle(),
ray_get_candidate_intersection_function: None,
ray_get_committed_intersection_function: None,
ray_query_functions: take(&mut self.ray_query_functions).recycle(),
io_f16_polyfills: take(&mut self.io_f16_polyfills).recycle(),
debug_printf: None,
};
*self = fresh;
@ -1022,6 +1032,7 @@ impl Writer {
expression_constness: super::ExpressionConstnessTracker::from_arena(
&ir_function.expressions,
),
ray_query_tracker_expr: crate::FastHashMap::default(),
};
// fill up the pre-emitted and const expressions
@ -1063,6 +1074,58 @@ impl Writer {
.function
.variables
.insert(handle, LocalVariable { id, instruction });
if let crate::TypeInner::RayQuery { .. } = ir_module.types[variable.ty].inner {
// Don't refactor this into a struct: Although spirv itself allows opaque types in structs,
// the vulkan environment for spirv does not. Putting ray queries into structs can cause
// confusing bugs.
let u32_type_id = context.writer.get_u32_type_id();
let ptr_u32_type_id = context
.writer
.get_pointer_type_id(u32_type_id, spirv::StorageClass::Function);
let tracker_id = context.gen_id();
let tracker_init_id = context
.writer
.get_constant_scalar(crate::Literal::U32(super::RayQueryPoint::empty().bits()));
let tracker_instruction = Instruction::variable(
ptr_u32_type_id,
tracker_id,
spirv::StorageClass::Function,
Some(tracker_init_id),
);
context
.function
.ray_query_initialization_tracker_variables
.insert(
handle,
LocalVariable {
id: tracker_id,
instruction: tracker_instruction,
},
);
let f32_type_id = context.writer.get_f32_type_id();
let ptr_f32_type_id = context
.writer
.get_pointer_type_id(f32_type_id, spirv::StorageClass::Function);
let t_max_tracker_id = context.gen_id();
let t_max_tracker_init_id =
context.writer.get_constant_scalar(crate::Literal::F32(0.0));
let t_max_tracker_instruction = Instruction::variable(
ptr_f32_type_id,
t_max_tracker_id,
spirv::StorageClass::Function,
Some(t_max_tracker_init_id),
);
context.function.ray_query_t_max_tracker_variables.insert(
handle,
LocalVariable {
id: t_max_tracker_id,
instruction: t_max_tracker_instruction,
},
);
}
}
for (handle, expr) in ir_function.expressions.iter() {
@ -2655,6 +2718,10 @@ impl Writer {
Instruction::memory_model(addressing_model, memory_model)
.to_words(&mut self.logical_layout.memory_model);
for debug_string in self.debug_strings.iter() {
debug_string.to_words(&mut self.logical_layout.debugs);
}
if self.flags.contains(WriterFlags::DEBUG) {
for debug in self.debugs.iter() {
debug.to_words(&mut self.logical_layout.debugs);
@ -2714,6 +2781,40 @@ impl Writer {
pub(super) fn needs_f16_polyfill(&self, ty_inner: &crate::TypeInner) -> bool {
self.io_f16_polyfills.needs_polyfill(ty_inner)
}
pub(super) fn write_debug_printf(
&mut self,
block: &mut Block,
string: &str,
format_params: &[Word],
) {
if self.debug_printf.is_none() {
self.use_extension("SPV_KHR_non_semantic_info");
let import_id = self.id_gen.next();
Instruction::ext_inst_import(import_id, "NonSemantic.DebugPrintf")
.to_words(&mut self.logical_layout.ext_inst_imports);
self.debug_printf = Some(import_id)
}
let import_id = self.debug_printf.unwrap();
let string_id = self.id_gen.next();
self.debug_strings
.push(Instruction::string(string, string_id));
let mut operands = Vec::with_capacity(1 + format_params.len());
operands.push(string_id);
operands.extend(format_params.iter());
let print_id = self.id_gen.next();
block.body.push(Instruction::ext_inst(
import_id,
1,
self.void_type,
print_id,
&operands,
));
}
}
#[test]

19
naga/tests/in/wgsl/ray-query-no-init-tracking.toml Normal file
View File

@ -0,0 +1,19 @@
god_mode = true
targets = "SPIRV | METAL | HLSL"
[msl]
fake_missing_bindings = true
lang_version = [2, 4]
spirv_cross_compatibility = false
zero_initialize_workgroup_memory = false
[hlsl]
shader_model = "V6_5"
fake_missing_bindings = true
zero_initialize_workgroup_memory = true
# Not yet implemented
# ray_query_initialization_tracking = false
[spv]
version = [1, 4]
ray_query_initialization_tracking = false

97
naga/tests/in/wgsl/ray-query-no-init-tracking.wgsl Normal file
View File

@ -0,0 +1,97 @@
/*
let RAY_FLAG_NONE = 0x00u;
let RAY_FLAG_FORCE_OPAQUE = 0x01u;
let RAY_FLAG_FORCE_NO_OPAQUE = 0x02u;
let RAY_FLAG_TERMINATE_ON_FIRST_HIT = 0x04u;
let RAY_FLAG_SKIP_CLOSEST_HIT_SHADER = 0x08u;
let RAY_FLAG_CULL_BACK_FACING = 0x10u;
let RAY_FLAG_CULL_FRONT_FACING = 0x20u;
let RAY_FLAG_CULL_OPAQUE = 0x40u;
let RAY_FLAG_CULL_NO_OPAQUE = 0x80u;
let RAY_FLAG_SKIP_TRIANGLES = 0x100u;
let RAY_FLAG_SKIP_AABBS = 0x200u;
let RAY_QUERY_INTERSECTION_NONE = 0u;
let RAY_QUERY_INTERSECTION_TRIANGLE = 1u;
let RAY_QUERY_INTERSECTION_GENERATED = 2u;
let RAY_QUERY_INTERSECTION_AABB = 3u;
struct RayDesc {
flags: u32,
cull_mask: u32,
t_min: f32,
t_max: f32,
origin: vec3<f32>,
dir: vec3<f32>,
}
struct RayIntersection {
kind: u32,
t: f32,
instance_custom_data: u32,
instance_index: u32,
sbt_record_offset: u32,
geometry_index: u32,
primitive_index: u32,
barycentrics: vec2<f32>,
front_face: bool,
object_to_world: mat4x3<f32>,
world_to_object: mat4x3<f32>,
}
*/
fn query_loop(pos: vec3<f32>, dir: vec3<f32>, acs: acceleration_structure) -> RayIntersection {
var rq: ray_query;
rayQueryInitialize(&rq, acs, RayDesc(RAY_FLAG_TERMINATE_ON_FIRST_HIT, 0xFFu, 0.1, 100.0, pos, dir));
while (rayQueryProceed(&rq)) {}
return rayQueryGetCommittedIntersection(&rq);
}
@group(0) @binding(0)
var acc_struct: acceleration_structure;
struct Output {
visible: u32,
normal: vec3<f32>,
}
@group(0) @binding(1)
var<storage, read_write> output: Output;
fn get_torus_normal(world_point: vec3<f32>, intersection: RayIntersection) -> vec3<f32> {
let local_point = intersection.world_to_object * vec4<f32>(world_point, 1.0);
let point_on_guiding_line = normalize(local_point.xy) * 2.4;
let world_point_on_guiding_line = intersection.object_to_world * vec4<f32>(point_on_guiding_line, 0.0, 1.0);
return normalize(world_point - world_point_on_guiding_line);
}
@compute @workgroup_size(1)
fn main() {
let pos = vec3<f32>(0.0);
let dir = vec3<f32>(0.0, 1.0, 0.0);
let intersection = query_loop(pos, dir, acc_struct);
output.visible = u32(intersection.kind == RAY_QUERY_INTERSECTION_NONE);
output.normal = get_torus_normal(dir * intersection.t, intersection);
}
@compute @workgroup_size(1)
fn main_candidate() {
let pos = vec3<f32>(0.0);
let dir = vec3<f32>(0.0, 1.0, 0.0);
var rq: ray_query;
rayQueryInitialize(&rq, acc_struct, RayDesc(RAY_FLAG_TERMINATE_ON_FIRST_HIT, 0xFFu, 0.1, 100.0, pos, dir));
let intersection = rayQueryGetCandidateIntersection(&rq);
if (intersection.kind == RAY_QUERY_INTERSECTION_AABB) {
rayQueryGenerateIntersection(&rq, 10.0);
} else if (intersection.kind == RAY_QUERY_INTERSECTION_TRIANGLE) {
rayQueryConfirmIntersection(&rq);
} else {
rayQueryTerminate(&rq);
}
}

165
naga/tests/out/hlsl/wgsl-ray-query-no-init-tracking.hlsl Normal file
View File

@ -0,0 +1,165 @@
struct RayIntersection {
uint kind;
float t;
uint instance_custom_data;
uint instance_index;
uint sbt_record_offset;
uint geometry_index;
uint primitive_index;
float2 barycentrics;
bool front_face;
int _pad9_0;
int _pad9_1;
row_major float4x3 object_to_world;
int _pad10_0;
row_major float4x3 world_to_object;
int _end_pad_0;
};
struct RayDesc_ {
uint flags;
uint cull_mask;
float tmin;
float tmax;
float3 origin;
int _pad5_0;
float3 dir;
int _end_pad_0;
};
struct Output {
uint visible;
int _pad1_0;
int _pad1_1;
int _pad1_2;
float3 normal;
int _end_pad_0;
};
RayDesc RayDescFromRayDesc_(RayDesc_ arg0) {
RayDesc ret = (RayDesc)0;
ret.Origin = arg0.origin;
ret.TMin = arg0.tmin;
ret.Direction = arg0.dir;
ret.TMax = arg0.tmax;
return ret;
}
RaytracingAccelerationStructure acc_struct : register(t0);
RWByteAddressBuffer output : register(u1);
RayDesc_ ConstructRayDesc_(uint arg0, uint arg1, float arg2, float arg3, float3 arg4, float3 arg5) {
RayDesc_ ret = (RayDesc_)0;
ret.flags = arg0;
ret.cull_mask = arg1;
ret.tmin = arg2;
ret.tmax = arg3;
ret.origin = arg4;
ret.dir = arg5;
return ret;
}
RayIntersection GetCommittedIntersection(RayQuery<RAY_FLAG_NONE> rq) {
RayIntersection ret = (RayIntersection)0;
ret.kind = rq.CommittedStatus();
if( rq.CommittedStatus() == COMMITTED_NOTHING) {} else {
ret.t = rq.CommittedRayT();
ret.instance_custom_data = rq.CommittedInstanceID();
ret.instance_index = rq.CommittedInstanceIndex();
ret.sbt_record_offset = rq.CommittedInstanceContributionToHitGroupIndex();
ret.geometry_index = rq.CommittedGeometryIndex();
ret.primitive_index = rq.CommittedPrimitiveIndex();
if( rq.CommittedStatus() == COMMITTED_TRIANGLE_HIT ) {
ret.barycentrics = rq.CommittedTriangleBarycentrics();
ret.front_face = rq.CommittedTriangleFrontFace();
}
ret.object_to_world = rq.CommittedObjectToWorld4x3();
ret.world_to_object = rq.CommittedWorldToObject4x3();
}
return ret;
}
RayIntersection query_loop(float3 pos, float3 dir, RaytracingAccelerationStructure acs)
{
RayQuery<RAY_FLAG_NONE> rq_1;
rq_1.TraceRayInline(acs, ConstructRayDesc_(4u, 255u, 0.1, 100.0, pos, dir).flags, ConstructRayDesc_(4u, 255u, 0.1, 100.0, pos, dir).cull_mask, RayDescFromRayDesc_(ConstructRayDesc_(4u, 255u, 0.1, 100.0, pos, dir)));
uint2 loop_bound = uint2(4294967295u, 4294967295u);
while(true) {
if (all(loop_bound == uint2(0u, 0u))) { break; }
loop_bound -= uint2(loop_bound.y == 0u, 1u);
const bool _e9 = rq_1.Proceed();
if (_e9) {
} else {
break;
}
{
}
}
const RayIntersection rayintersection = GetCommittedIntersection(rq_1);
return rayintersection;
}
float3 get_torus_normal(float3 world_point, RayIntersection intersection)
{
float3 local_point = mul(float4(world_point, 1.0), intersection.world_to_object);
float2 point_on_guiding_line = (normalize(local_point.xy) * 2.4);
float3 world_point_on_guiding_line = mul(float4(point_on_guiding_line, 0.0, 1.0), intersection.object_to_world);
return normalize((world_point - world_point_on_guiding_line));
}
[numthreads(1, 1, 1)]
void main()
{
float3 pos_1 = (0.0).xxx;
float3 dir_1 = float3(0.0, 1.0, 0.0);
const RayIntersection _e7 = query_loop(pos_1, dir_1, acc_struct);
output.Store(0, asuint(uint((_e7.kind == 0u))));
const float3 _e18 = get_torus_normal((dir_1 * _e7.t), _e7);
output.Store3(16, asuint(_e18));
return;
}
RayIntersection GetCandidateIntersection(RayQuery<RAY_FLAG_NONE> rq) {
RayIntersection ret = (RayIntersection)0;
CANDIDATE_TYPE kind = rq.CandidateType();
if (kind == CANDIDATE_NON_OPAQUE_TRIANGLE) {
ret.kind = 1;
ret.t = rq.CandidateTriangleRayT();
ret.barycentrics = rq.CandidateTriangleBarycentrics();
ret.front_face = rq.CandidateTriangleFrontFace();
} else {
ret.kind = 3;
}
ret.instance_custom_data = rq.CandidateInstanceID();
ret.instance_index = rq.CandidateInstanceIndex();
ret.sbt_record_offset = rq.CandidateInstanceContributionToHitGroupIndex();
ret.geometry_index = rq.CandidateGeometryIndex();
ret.primitive_index = rq.CandidatePrimitiveIndex();
ret.object_to_world = rq.CandidateObjectToWorld4x3();
ret.world_to_object = rq.CandidateWorldToObject4x3();
return ret;
}
[numthreads(1, 1, 1)]
void main_candidate()
{
RayQuery<RAY_FLAG_NONE> rq;
float3 pos_2 = (0.0).xxx;
float3 dir_2 = float3(0.0, 1.0, 0.0);
rq.TraceRayInline(acc_struct, ConstructRayDesc_(4u, 255u, 0.1, 100.0, pos_2, dir_2).flags, ConstructRayDesc_(4u, 255u, 0.1, 100.0, pos_2, dir_2).cull_mask, RayDescFromRayDesc_(ConstructRayDesc_(4u, 255u, 0.1, 100.0, pos_2, dir_2)));
RayIntersection intersection_1 = GetCandidateIntersection(rq);
if ((intersection_1.kind == 3u)) {
rq.CommitProceduralPrimitiveHit(10.0);
return;
} else {
if ((intersection_1.kind == 1u)) {
rq.CommitNonOpaqueTriangleHit();
return;
} else {
rq.Abort();
return;
}
}
}

16
naga/tests/out/hlsl/wgsl-ray-query-no-init-tracking.ron Normal file
View File

@ -0,0 +1,16 @@
(
vertex:[
],
fragment:[
],
compute:[
(
entry_point:"main",
target_profile:"cs_6_5",
),
(
entry_point:"main_candidate",
target_profile:"cs_6_5",
),
],
)

116
naga/tests/out/msl/wgsl-ray-query-no-init-tracking.msl Normal file
View File

@ -0,0 +1,116 @@
// language: metal2.4
#include <metal_stdlib>
#include <simd/simd.h>
using metal::uint;
struct _RayQuery {
metal::raytracing::intersector<metal::raytracing::instancing, metal::raytracing::triangle_data, metal::raytracing::world_space_data> intersector;
metal::raytracing::intersector<metal::raytracing::instancing, metal::raytracing::triangle_data, metal::raytracing::world_space_data>::result_type intersection;
bool ready = false;
};
constexpr metal::uint _map_intersection_type(const metal::raytracing::intersection_type ty) {
return ty==metal::raytracing::intersection_type::triangle ? 1 :
ty==metal::raytracing::intersection_type::bounding_box ? 4 : 0;
}
struct RayIntersection {
uint kind;
float t;
uint instance_custom_data;
uint instance_index;
uint sbt_record_offset;
uint geometry_index;
uint primitive_index;
metal::float2 barycentrics;
bool front_face;
char _pad9[11];
metal::float4x3 object_to_world;
metal::float4x3 world_to_object;
};
struct RayDesc {
uint flags;
uint cull_mask;
float tmin;
float tmax;
metal::float3 origin;
metal::float3 dir;
};
struct Output {
uint visible;
char _pad1[12];
metal::float3 normal;
};
RayIntersection query_loop(
metal::float3 pos,
metal::float3 dir,
metal::raytracing::instance_acceleration_structure acs
) {
_RayQuery rq_1 = {};
RayDesc _e8 = RayDesc {4u, 255u, 0.1, 100.0, pos, dir};
rq_1.intersector.assume_geometry_type(metal::raytracing::geometry_type::triangle);
rq_1.intersector.set_opacity_cull_mode((_e8.flags & 64) != 0 ? metal::raytracing::opacity_cull_mode::opaque : (_e8.flags & 128) != 0 ? metal::raytracing::opacity_cull_mode::non_opaque : metal::raytracing::opacity_cull_mode::none);
rq_1.intersector.force_opacity((_e8.flags & 1) != 0 ? metal::raytracing::forced_opacity::opaque : (_e8.flags & 2) != 0 ? metal::raytracing::forced_opacity::non_opaque : metal::raytracing::forced_opacity::none);
rq_1.intersector.accept_any_intersection((_e8.flags & 4) != 0);
rq_1.intersection = rq_1.intersector.intersect(metal::raytracing::ray(_e8.origin, _e8.dir, _e8.tmin, _e8.tmax), acs, _e8.cull_mask); rq_1.ready = true;
uint2 loop_bound = uint2(4294967295u);
while(true) {
if (metal::all(loop_bound == uint2(0u))) { break; }
loop_bound -= uint2(loop_bound.y == 0u, 1u);
bool _e9 = rq_1.ready;
if (_e9) {
} else {
break;
}
}
return RayIntersection {_map_intersection_type(rq_1.intersection.type), rq_1.intersection.distance, rq_1.intersection.user_instance_id, rq_1.intersection.instance_id, {}, rq_1.intersection.geometry_id, rq_1.intersection.primitive_id, rq_1.intersection.triangle_barycentric_coord, rq_1.intersection.triangle_front_facing, {}, rq_1.intersection.object_to_world_transform, rq_1.intersection.world_to_object_transform};
}
metal::float3 get_torus_normal(
metal::float3 world_point,
RayIntersection intersection
) {
metal::float3 local_point = intersection.world_to_object * metal::float4(world_point, 1.0);
metal::float2 point_on_guiding_line = metal::normalize(local_point.xy) * 2.4;
metal::float3 world_point_on_guiding_line = intersection.object_to_world * metal::float4(point_on_guiding_line, 0.0, 1.0);
return metal::normalize(world_point - world_point_on_guiding_line);
}
kernel void main_(
metal::raytracing::instance_acceleration_structure acc_struct [[user(fake0)]]
, device Output& output [[user(fake0)]]
) {
metal::float3 pos_1 = metal::float3(0.0);
metal::float3 dir_1 = metal::float3(0.0, 1.0, 0.0);
RayIntersection _e7 = query_loop(pos_1, dir_1, acc_struct);
output.visible = static_cast<uint>(_e7.kind == 0u);
metal::float3 _e18 = get_torus_normal(dir_1 * _e7.t, _e7);
output.normal = _e18;
return;
}
kernel void main_candidate(
metal::raytracing::instance_acceleration_structure acc_struct [[user(fake0)]]
) {
_RayQuery rq = {};
metal::float3 pos_2 = metal::float3(0.0);
metal::float3 dir_2 = metal::float3(0.0, 1.0, 0.0);
RayDesc _e12 = RayDesc {4u, 255u, 0.1, 100.0, pos_2, dir_2};
rq.intersector.assume_geometry_type(metal::raytracing::geometry_type::triangle);
rq.intersector.set_opacity_cull_mode((_e12.flags & 64) != 0 ? metal::raytracing::opacity_cull_mode::opaque : (_e12.flags & 128) != 0 ? metal::raytracing::opacity_cull_mode::non_opaque : metal::raytracing::opacity_cull_mode::none);
rq.intersector.force_opacity((_e12.flags & 1) != 0 ? metal::raytracing::forced_opacity::opaque : (_e12.flags & 2) != 0 ? metal::raytracing::forced_opacity::non_opaque : metal::raytracing::forced_opacity::none);
rq.intersector.accept_any_intersection((_e12.flags & 4) != 0);
rq.intersection = rq.intersector.intersect(metal::raytracing::ray(_e12.origin, _e12.dir, _e12.tmin, _e12.tmax), acc_struct, _e12.cull_mask); rq.ready = true;
RayIntersection intersection_1 = RayIntersection {_map_intersection_type(rq.intersection.type), rq.intersection.distance, rq.intersection.user_instance_id, rq.intersection.instance_id, {}, rq.intersection.geometry_id, rq.intersection.primitive_id, rq.intersection.triangle_barycentric_coord, rq.intersection.triangle_front_facing, {}, rq.intersection.object_to_world_transform, rq.intersection.world_to_object_transform};
if (intersection_1.kind == 3u) {
return;
} else {
if (intersection_1.kind == 1u) {
return;
} else {
rq.ready = false;
return;
}
}
}

369
naga/tests/out/spv/wgsl-aliased-ray-query.spvasm
View File

@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.4
; Generator: rspirv
; Bound: 102
; Bound: 244
OpCapability Shader
OpCapability RayQueryKHR
OpExtension "SPV_KHR_ray_query"
@ -59,109 +59,284 @@ OpDecorate %13 Binding 0
%29 = OpConstant %5 10
%30 = OpConstant %7 1
%32 = OpTypePointer Function %3
%40 = OpTypePointer Function %12
%41 = OpTypePointer Function %7
%42 = OpTypePointer Function %11
%43 = OpTypePointer Function %9
%44 = OpTypePointer Function %10
%45 = OpTypePointer Function %5
%46 = OpTypeFunction %12 %32
%48 = OpConstantNull %12
%52 = OpConstant %7 0
%67 = OpConstant %7 2
%71 = OpConstant %7 5
%73 = OpConstant %7 6
%75 = OpConstant %7 9
%77 = OpConstant %7 10
%86 = OpConstant %7 7
%88 = OpConstant %7 8
%47 = OpFunction %12 None %46
%49 = OpFunctionParameter %32
%50 = OpLabel
%51 = OpVariable %40 Function %48
%53 = OpRayQueryGetIntersectionTypeKHR %7 %49 %52
%54 = OpIEqual %10 %53 %52
%55 = OpSelect %7 %54 %30 %28
%56 = OpAccessChain %41 %51 %52
OpStore %56 %55
%57 = OpINotEqual %10 %55 %52
OpSelectionMerge %59 None
OpBranchConditional %57 %58 %59
%58 = OpLabel
%60 = OpRayQueryGetIntersectionInstanceCustomIndexKHR %7 %49 %52
%61 = OpRayQueryGetIntersectionInstanceIdKHR %7 %49 %52
%62 = OpRayQueryGetIntersectionInstanceShaderBindingTableRecordOffsetKHR %7 %49 %52
%63 = OpRayQueryGetIntersectionGeometryIndexKHR %7 %49 %52
%64 = OpRayQueryGetIntersectionPrimitiveIndexKHR %7 %49 %52
%65 = OpRayQueryGetIntersectionObjectToWorldKHR %11 %49 %52
%66 = OpRayQueryGetIntersectionWorldToObjectKHR %11 %49 %52
%68 = OpAccessChain %41 %51 %67
OpStore %68 %60
%69 = OpAccessChain %41 %51 %28
OpStore %69 %61
%70 = OpAccessChain %41 %51 %23
OpStore %70 %62
%72 = OpAccessChain %41 %51 %71
OpStore %72 %63
%74 = OpAccessChain %41 %51 %73
OpStore %74 %64
%76 = OpAccessChain %42 %51 %75
OpStore %76 %65
%78 = OpAccessChain %42 %51 %77
OpStore %78 %66
%79 = OpIEqual %10 %55 %30
OpSelectionMerge %81 None
OpBranchConditional %57 %80 %81
%80 = OpLabel
%82 = OpRayQueryGetIntersectionTKHR %5 %49 %52
%83 = OpAccessChain %45 %51 %30
OpStore %83 %82
%84 = OpRayQueryGetIntersectionBarycentricsKHR %9 %49 %52
%85 = OpRayQueryGetIntersectionFrontFaceKHR %10 %49 %52
%87 = OpAccessChain %43 %51 %86
OpStore %87 %84
%89 = OpAccessChain %44 %51 %88
OpStore %89 %85
OpBranch %81
%81 = OpLabel
OpBranch %59
%59 = OpLabel
%90 = OpLoad %12 %51
OpReturnValue %90
%33 = OpTypePointer Function %7
%35 = OpConstant %7 0
%36 = OpTypePointer Function %5
%39 = OpTypeVector %10 3
%40 = OpTypeFunction %2 %32 %4 %8 %33 %36
%68 = OpConstant %7 256
%71 = OpConstant %7 512
%76 = OpConstant %7 16
%79 = OpConstant %7 32
%90 = OpConstant %7 2
%93 = OpConstant %7 64
%96 = OpConstant %7 128
%121 = OpTypePointer Function %12
%122 = OpTypePointer Function %11
%123 = OpTypePointer Function %9
%124 = OpTypePointer Function %10
%125 = OpTypeFunction %12 %32 %33
%130 = OpConstantNull %12
%158 = OpConstant %7 5
%160 = OpConstant %7 6
%162 = OpConstant %7 9
%164 = OpConstant %7 10
%173 = OpConstant %7 7
%175 = OpConstant %7 8
%184 = OpTypeFunction %2 %32 %33 %5 %36
%225 = OpTypeFunction %2 %32 %33
%41 = OpFunction %2 None %40
%42 = OpFunctionParameter %32
%43 = OpFunctionParameter %4
%44 = OpFunctionParameter %8
%45 = OpFunctionParameter %33
%46 = OpFunctionParameter %36
%47 = OpLabel
%48 = OpCompositeExtract %7 %44 0
%49 = OpCompositeExtract %7 %44 1
%50 = OpCompositeExtract %5 %44 2
%51 = OpCompositeExtract %5 %44 3
OpStore %46 %51
%52 = OpCompositeExtract %6 %44 4
%53 = OpCompositeExtract %6 %44 5
%54 = OpFOrdLessThanEqual %10 %50 %51
%55 = OpFOrdGreaterThanEqual %10 %50 %19
%56 = OpIsInf %39 %52
%57 = OpAny %10 %56
%58 = OpIsNan %39 %52
%59 = OpAny %10 %58
%60 = OpLogicalOr %10 %59 %57
%61 = OpLogicalNot %10 %60
%62 = OpIsInf %39 %53
%63 = OpAny %10 %62
%64 = OpIsNan %39 %53
%65 = OpAny %10 %64
%66 = OpLogicalOr %10 %65 %63
%67 = OpLogicalNot %10 %66
%69 = OpBitwiseAnd %7 %48 %68
%70 = OpINotEqual %10 %69 %35
%72 = OpBitwiseAnd %7 %48 %71
%73 = OpINotEqual %10 %72 %35
%74 = OpLogicalAnd %10 %73 %70
%75 = OpLogicalNot %10 %74
%77 = OpBitwiseAnd %7 %48 %76
%78 = OpINotEqual %10 %77 %35
%80 = OpBitwiseAnd %7 %48 %79
%81 = OpINotEqual %10 %80 %35
%82 = OpLogicalAnd %10 %81 %70
%83 = OpLogicalAnd %10 %81 %78
%84 = OpLogicalAnd %10 %78 %70
%85 = OpLogicalOr %10 %84 %82
%86 = OpLogicalOr %10 %85 %83
%87 = OpLogicalNot %10 %86
%88 = OpBitwiseAnd %7 %48 %30
%89 = OpINotEqual %10 %88 %35
%91 = OpBitwiseAnd %7 %48 %90
%92 = OpINotEqual %10 %91 %35
%94 = OpBitwiseAnd %7 %48 %93
%95 = OpINotEqual %10 %94 %35
%97 = OpBitwiseAnd %7 %48 %96
%98 = OpINotEqual %10 %97 %35
%99 = OpLogicalAnd %10 %98 %89
%100 = OpLogicalAnd %10 %98 %92
%101 = OpLogicalAnd %10 %98 %95
%102 = OpLogicalAnd %10 %95 %89
%103 = OpLogicalAnd %10 %95 %92
%104 = OpLogicalAnd %10 %92 %89
%105 = OpLogicalOr %10 %104 %99
%106 = OpLogicalOr %10 %105 %100
%107 = OpLogicalOr %10 %106 %101
%108 = OpLogicalOr %10 %107 %102
%109 = OpLogicalOr %10 %108 %103
%110 = OpLogicalNot %10 %109
%111 = OpLogicalAnd %10 %110 %54
%112 = OpLogicalAnd %10 %111 %55
%113 = OpLogicalAnd %10 %112 %61
%114 = OpLogicalAnd %10 %113 %67
%115 = OpLogicalAnd %10 %114 %75
%116 = OpLogicalAnd %10 %115 %87
OpSelectionMerge %117 None
OpBranchConditional %116 %119 %118
%119 = OpLabel
OpRayQueryInitializeKHR %42 %43 %48 %49 %52 %50 %53 %51
OpStore %45 %30
OpBranch %117
%118 = OpLabel
OpBranch %117
%117 = OpLabel
OpReturn
OpFunctionEnd
%126 = OpFunction %12 None %125
%127 = OpFunctionParameter %32
%128 = OpFunctionParameter %33
%129 = OpLabel
%131 = OpVariable %121 Function %130
%132 = OpLoad %7 %128
%133 = OpBitwiseAnd %7 %132 %90
%134 = OpINotEqual %10 %133 %35
%135 = OpBitwiseAnd %7 %132 %23
%136 = OpINotEqual %10 %135 %35
%137 = OpLogicalNot %10 %136
%138 = OpLogicalAnd %10 %137 %134
OpSelectionMerge %140 None
OpBranchConditional %138 %139 %140
%139 = OpLabel
%141 = OpRayQueryGetIntersectionTypeKHR %7 %127 %35
%142 = OpIEqual %10 %141 %35
%143 = OpSelect %7 %142 %30 %28
%144 = OpAccessChain %33 %131 %35
OpStore %144 %143
%145 = OpINotEqual %10 %143 %35
OpSelectionMerge %147 None
OpBranchConditional %145 %146 %147
%146 = OpLabel
%148 = OpRayQueryGetIntersectionInstanceCustomIndexKHR %7 %127 %35
%149 = OpRayQueryGetIntersectionInstanceIdKHR %7 %127 %35
%150 = OpRayQueryGetIntersectionInstanceShaderBindingTableRecordOffsetKHR %7 %127 %35
%151 = OpRayQueryGetIntersectionGeometryIndexKHR %7 %127 %35
%152 = OpRayQueryGetIntersectionPrimitiveIndexKHR %7 %127 %35
%153 = OpRayQueryGetIntersectionObjectToWorldKHR %11 %127 %35
%154 = OpRayQueryGetIntersectionWorldToObjectKHR %11 %127 %35
%155 = OpAccessChain %33 %131 %90
OpStore %155 %148
%156 = OpAccessChain %33 %131 %28
OpStore %156 %149
%157 = OpAccessChain %33 %131 %23
OpStore %157 %150
%159 = OpAccessChain %33 %131 %158
OpStore %159 %151
%161 = OpAccessChain %33 %131 %160
OpStore %161 %152
%163 = OpAccessChain %122 %131 %162
OpStore %163 %153
%165 = OpAccessChain %122 %131 %164
OpStore %165 %154
%166 = OpIEqual %10 %143 %30
OpSelectionMerge %168 None
OpBranchConditional %145 %167 %168
%167 = OpLabel
%169 = OpRayQueryGetIntersectionTKHR %5 %127 %35
%170 = OpAccessChain %36 %131 %30
OpStore %170 %169
%171 = OpRayQueryGetIntersectionBarycentricsKHR %9 %127 %35
%172 = OpRayQueryGetIntersectionFrontFaceKHR %10 %127 %35
%174 = OpAccessChain %123 %131 %173
OpStore %174 %171
%176 = OpAccessChain %124 %131 %175
OpStore %176 %172
OpBranch %168
%168 = OpLabel
OpBranch %147
%147 = OpLabel
OpBranch %140
%140 = OpLabel
%177 = OpLoad %12 %131
OpReturnValue %177
OpFunctionEnd
%185 = OpFunction %2 None %184
%186 = OpFunctionParameter %32
%187 = OpFunctionParameter %33
%188 = OpFunctionParameter %5
%189 = OpFunctionParameter %36
%190 = OpLabel
%191 = OpVariable %36 Function
%192 = OpVariable %36 Function
%195 = OpLoad %7 %187
%196 = OpBitwiseAnd %7 %195 %90
%197 = OpINotEqual %10 %196 %35
%198 = OpBitwiseAnd %7 %195 %23
%199 = OpINotEqual %10 %198 %35
%200 = OpLogicalNot %10 %199
%201 = OpLogicalAnd %10 %200 %197
OpSelectionMerge %194 None
OpBranchConditional %201 %193 %194
%193 = OpLabel
%202 = OpRayQueryGetIntersectionTypeKHR %7 %186 %35
%203 = OpIEqual %10 %202 %30
%204 = OpRayQueryGetRayTMinKHR %5 %186
%205 = OpRayQueryGetIntersectionTypeKHR %7 %186 %30
%206 = OpIEqual %10 %205 %35
OpSelectionMerge %207 None
OpBranchConditional %206 %208 %209
%208 = OpLabel
%210 = OpLoad %5 %189
OpStore %192 %210
OpBranch %207
%209 = OpLabel
%211 = OpRayQueryGetIntersectionTKHR %5 %186 %35
OpStore %192 %211
OpBranch %207
%207 = OpLabel
%212 = OpFOrdGreaterThanEqual %10 %188 %204
%213 = OpLoad %5 %192
%214 = OpFOrdLessThanEqual %10 %188 %213
%215 = OpLogicalAnd %10 %212 %214
%216 = OpLogicalAnd %10 %215 %203
OpSelectionMerge %218 None
OpBranchConditional %216 %217 %218
%217 = OpLabel
OpRayQueryGenerateIntersectionKHR %186 %188
OpBranch %218
%218 = OpLabel
OpBranch %194
%194 = OpLabel
OpReturn
OpFunctionEnd
%226 = OpFunction %2 None %225
%227 = OpFunctionParameter %32
%228 = OpFunctionParameter %33
%229 = OpLabel
%232 = OpLoad %7 %228
%233 = OpBitwiseAnd %7 %232 %90
%234 = OpINotEqual %10 %233 %35
%235 = OpBitwiseAnd %7 %232 %23
%236 = OpINotEqual %10 %235 %35
%237 = OpLogicalNot %10 %236
%238 = OpLogicalAnd %10 %237 %234
OpSelectionMerge %231 None
OpBranchConditional %238 %230 %231
%230 = OpLabel
%239 = OpRayQueryGetIntersectionTypeKHR %7 %227 %35
%240 = OpIEqual %10 %239 %35
OpSelectionMerge %242 None
OpBranchConditional %240 %241 %242
%241 = OpLabel
OpRayQueryConfirmIntersectionKHR %227
OpBranch %242
%242 = OpLabel
OpBranch %231
%231 = OpLabel
OpReturn
OpFunctionEnd
%16 = OpFunction %2 None %17
%15 = OpLabel
%31 = OpVariable %32 Function
%34 = OpVariable %33 Function %35
%37 = OpVariable %36 Function %19
%18 = OpLoad %4 %13
OpBranch %33
%33 = OpLabel
%34 = OpCompositeExtract %7 %27 0
%35 = OpCompositeExtract %7 %27 1
%36 = OpCompositeExtract %5 %27 2
%37 = OpCompositeExtract %5 %27 3
%38 = OpCompositeExtract %6 %27 4
%39 = OpCompositeExtract %6 %27 5
OpRayQueryInitializeKHR %31 %18 %34 %35 %38 %36 %39 %37
%91 = OpFunctionCall %12 %47 %31
%92 = OpCompositeExtract %7 %91 0
%93 = OpIEqual %10 %92 %28
OpSelectionMerge %94 None
OpBranchConditional %93 %95 %96
%95 = OpLabel
OpRayQueryGenerateIntersectionKHR %31 %29
OpBranch %38
%38 = OpLabel
%120 = OpFunctionCall %2 %41 %31 %18 %27 %34 %37
%178 = OpFunctionCall %12 %126 %31 %34
%179 = OpCompositeExtract %7 %178 0
%180 = OpIEqual %10 %179 %28
OpSelectionMerge %181 None
OpBranchConditional %180 %182 %183
%182 = OpLabel
%219 = OpFunctionCall %2 %185 %31 %34 %29 %37
OpReturn
%96 = OpLabel
%97 = OpCompositeExtract %7 %91 0
%98 = OpIEqual %10 %97 %30
OpSelectionMerge %99 None
OpBranchConditional %98 %100 %101
%100 = OpLabel
OpRayQueryConfirmIntersectionKHR %31
%183 = OpLabel
%220 = OpCompositeExtract %7 %178 0
%221 = OpIEqual %10 %220 %30
OpSelectionMerge %222 None
OpBranchConditional %221 %223 %224
%223 = OpLabel
%243 = OpFunctionCall %2 %226 %31 %34
OpReturn
%101 = OpLabel
%224 = OpLabel
OpReturn
%99 = OpLabel
OpBranch %94
%94 = OpLabel
%222 = OpLabel
OpBranch %181
%181 = OpLabel
OpReturn
OpFunctionEnd

220
naga/tests/out/spv/wgsl-overrides-ray-query.main.spvasm
View File

@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.4
; Generator: rspirv
; Bound: 65
; Bound: 164
OpCapability Shader
OpCapability RayQueryKHR
OpExtension "SPV_KHR_ray_query"
@ -40,61 +40,175 @@ OpDecorate %10 Binding 0
%25 = OpConstantComposite %7 %22 %23 %24
%26 = OpConstantComposite %8 %16 %17 %18 %19 %21 %25
%28 = OpTypePointer Function %5
%40 = OpTypeVector %6 2
%41 = OpTypePointer Function %40
%42 = OpTypeBool
%43 = OpTypeVector %42 2
%44 = OpConstant %6 0
%45 = OpConstantComposite %40 %44 %44
%46 = OpConstant %6 1
%47 = OpConstant %6 4294967295
%48 = OpConstantComposite %40 %47 %47
%29 = OpTypePointer Function %6
%31 = OpConstant %6 0
%32 = OpTypePointer Function %3
%34 = OpConstant %3 0
%36 = OpTypeBool
%37 = OpTypeVector %36 3
%38 = OpTypeFunction %2 %28 %4 %8 %29 %32
%66 = OpConstant %6 256
%69 = OpConstant %6 512
%74 = OpConstant %6 16
%77 = OpConstant %6 32
%86 = OpConstant %6 1
%89 = OpConstant %6 2
%92 = OpConstant %6 64
%95 = OpConstant %6 128
%124 = OpTypeVector %6 2
%125 = OpTypePointer Function %124
%126 = OpTypeVector %36 2
%127 = OpConstantComposite %124 %31 %31
%128 = OpConstant %6 4294967295
%129 = OpConstantComposite %124 %128 %128
%142 = OpTypePointer Function %36
%143 = OpTypeFunction %36 %28 %29
%149 = OpConstantFalse %36
%156 = OpConstant %6 6
%39 = OpFunction %2 None %38
%40 = OpFunctionParameter %28
%41 = OpFunctionParameter %4
%42 = OpFunctionParameter %8
%43 = OpFunctionParameter %29
%44 = OpFunctionParameter %32
%45 = OpLabel
%46 = OpCompositeExtract %6 %42 0
%47 = OpCompositeExtract %6 %42 1
%48 = OpCompositeExtract %3 %42 2
%49 = OpCompositeExtract %3 %42 3
OpStore %44 %49
%50 = OpCompositeExtract %7 %42 4
%51 = OpCompositeExtract %7 %42 5
%52 = OpFOrdLessThanEqual %36 %48 %49
%53 = OpFOrdGreaterThanEqual %36 %48 %34
%54 = OpIsInf %37 %50
%55 = OpAny %36 %54
%56 = OpIsNan %37 %50
%57 = OpAny %36 %56
%58 = OpLogicalOr %36 %57 %55
%59 = OpLogicalNot %36 %58
%60 = OpIsInf %37 %51
%61 = OpAny %36 %60
%62 = OpIsNan %37 %51
%63 = OpAny %36 %62
%64 = OpLogicalOr %36 %63 %61
%65 = OpLogicalNot %36 %64
%67 = OpBitwiseAnd %6 %46 %66
%68 = OpINotEqual %36 %67 %31
%70 = OpBitwiseAnd %6 %46 %69
%71 = OpINotEqual %36 %70 %31
%72 = OpLogicalAnd %36 %71 %68
%73 = OpLogicalNot %36 %72
%75 = OpBitwiseAnd %6 %46 %74
%76 = OpINotEqual %36 %75 %31
%78 = OpBitwiseAnd %6 %46 %77
%79 = OpINotEqual %36 %78 %31
%80 = OpLogicalAnd %36 %79 %68
%81 = OpLogicalAnd %36 %79 %76
%82 = OpLogicalAnd %36 %76 %68
%83 = OpLogicalOr %36 %82 %80
%84 = OpLogicalOr %36 %83 %81
%85 = OpLogicalNot %36 %84
%87 = OpBitwiseAnd %6 %46 %86
%88 = OpINotEqual %36 %87 %31
%90 = OpBitwiseAnd %6 %46 %89
%91 = OpINotEqual %36 %90 %31
%93 = OpBitwiseAnd %6 %46 %92
%94 = OpINotEqual %36 %93 %31
%96 = OpBitwiseAnd %6 %46 %95
%97 = OpINotEqual %36 %96 %31
%98 = OpLogicalAnd %36 %97 %88
%99 = OpLogicalAnd %36 %97 %91
%100 = OpLogicalAnd %36 %97 %94
%101 = OpLogicalAnd %36 %94 %88
%102 = OpLogicalAnd %36 %94 %91
%103 = OpLogicalAnd %36 %91 %88
%104 = OpLogicalOr %36 %103 %98
%105 = OpLogicalOr %36 %104 %99
%106 = OpLogicalOr %36 %105 %100
%107 = OpLogicalOr %36 %106 %101
%108 = OpLogicalOr %36 %107 %102
%109 = OpLogicalNot %36 %108
%110 = OpLogicalAnd %36 %109 %52
%111 = OpLogicalAnd %36 %110 %53
%112 = OpLogicalAnd %36 %111 %59
%113 = OpLogicalAnd %36 %112 %65
%114 = OpLogicalAnd %36 %113 %73
%115 = OpLogicalAnd %36 %114 %85
OpSelectionMerge %116 None
OpBranchConditional %115 %118 %117
%118 = OpLabel
OpRayQueryInitializeKHR %40 %41 %46 %47 %50 %48 %51 %49
OpStore %43 %86
OpBranch %116
%117 = OpLabel
OpBranch %116
%116 = OpLabel
OpReturn
OpFunctionEnd
%144 = OpFunction %36 None %143
%145 = OpFunctionParameter %28
%146 = OpFunctionParameter %29
%147 = OpLabel
%148 = OpVariable %142 Function %149
%150 = OpLoad %6 %146
%153 = OpBitwiseAnd %6 %150 %86
%154 = OpINotEqual %36 %153 %31
OpSelectionMerge %151 None
OpBranchConditional %154 %152 %151
%152 = OpLabel
%155 = OpRayQueryProceedKHR %36 %145
OpStore %148 %155
%157 = OpSelect %6 %155 %89 %156
%158 = OpBitwiseOr %6 %150 %157
OpStore %146 %158
OpBranch %151
%151 = OpLabel
%159 = OpLoad %36 %148
OpReturnValue %159
OpFunctionEnd
%13 = OpFunction %2 None %14
%12 = OpLabel
%27 = OpVariable %28 Function
%49 = OpVariable %41 Function %48
%30 = OpVariable %29 Function %31
%33 = OpVariable %32 Function %34
%130 = OpVariable %125 Function %129
%15 = OpLoad %4 %10
OpBranch %29
%29 = OpLabel
%30 = OpCompositeExtract %6 %26 0
%31 = OpCompositeExtract %6 %26 1
%32 = OpCompositeExtract %3 %26 2
%33 = OpCompositeExtract %3 %26 3
%34 = OpCompositeExtract %7 %26 4
%35 = OpCompositeExtract %7 %26 5
OpRayQueryInitializeKHR %27 %15 %30 %31 %34 %32 %35 %33
OpBranch %36
%36 = OpLabel
OpLoopMerge %37 %39 None
OpBranch %50
%50 = OpLabel
%51 = OpLoad %40 %49
%52 = OpIEqual %43 %45 %51
%53 = OpAll %42 %52
OpSelectionMerge %54 None
OpBranchConditional %53 %37 %54
%54 = OpLabel
%55 = OpCompositeExtract %6 %51 1
%56 = OpIEqual %42 %55 %44
%57 = OpSelect %6 %56 %46 %44
%58 = OpCompositeConstruct %40 %57 %46
%59 = OpISub %40 %51 %58
OpStore %49 %59
OpBranch %38
%38 = OpLabel
%60 = OpRayQueryProceedKHR %42 %27
OpSelectionMerge %61 None
OpBranchConditional %60 %61 %62
%62 = OpLabel
OpBranch %37
%61 = OpLabel
OpBranch %63
%63 = OpLabel
OpBranch %64
%64 = OpLabel
OpBranch %39
%39 = OpLabel
OpBranch %36
%37 = OpLabel
OpBranch %35
%35 = OpLabel
%119 = OpFunctionCall %2 %39 %27 %15 %26 %30 %33
OpBranch %120
%120 = OpLabel
OpLoopMerge %121 %123 None
OpBranch %131
%131 = OpLabel
%132 = OpLoad %124 %130
%133 = OpIEqual %126 %127 %132
%134 = OpAll %36 %133
OpSelectionMerge %135 None
OpBranchConditional %134 %121 %135
%135 = OpLabel
%136 = OpCompositeExtract %6 %132 1
%137 = OpIEqual %36 %136 %31
%138 = OpSelect %6 %137 %86 %31
%139 = OpCompositeConstruct %124 %138 %86
%140 = OpISub %124 %132 %139
OpStore %130 %140
OpBranch %122
%122 = OpLabel
%141 = OpFunctionCall %36 %144 %27 %30
OpSelectionMerge %160 None
OpBranchConditional %141 %160 %161
%161 = OpLabel
OpBranch %121
%160 = OpLabel
OpBranch %162
%162 = OpLabel
OpBranch %163
%163 = OpLabel
OpBranch %123
%123 = OpLabel
OpBranch %120
%121 = OpLabel
OpReturn
OpFunctionEnd

542
naga/tests/out/spv/wgsl-ray-query-no-init-tracking.spvasm Normal file
View File

@ -0,0 +1,542 @@
; SPIR-V
; Version: 1.4
; Generator: rspirv
; Bound: 382
OpCapability Shader
OpCapability RayQueryKHR
OpExtension "SPV_KHR_ray_query"
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %242 "main" %15 %17
OpEntryPoint GLCompute %262 "main_candidate" %15
OpExecutionMode %242 LocalSize 1 1 1
OpExecutionMode %262 LocalSize 1 1 1
OpMemberDecorate %10 0 Offset 0
OpMemberDecorate %10 1 Offset 4
OpMemberDecorate %10 2 Offset 8
OpMemberDecorate %10 3 Offset 12
OpMemberDecorate %10 4 Offset 16
OpMemberDecorate %10 5 Offset 20
OpMemberDecorate %10 6 Offset 24
OpMemberDecorate %10 7 Offset 28
OpMemberDecorate %10 8 Offset 36
OpMemberDecorate %10 9 Offset 48
OpMemberDecorate %10 9 ColMajor
OpMemberDecorate %10 9 MatrixStride 16
OpMemberDecorate %10 10 Offset 112
OpMemberDecorate %10 10 ColMajor
OpMemberDecorate %10 10 MatrixStride 16
OpMemberDecorate %12 0 Offset 0
OpMemberDecorate %12 1 Offset 4
OpMemberDecorate %12 2 Offset 8
OpMemberDecorate %12 3 Offset 12
OpMemberDecorate %12 4 Offset 16
OpMemberDecorate %12 5 Offset 32
OpMemberDecorate %13 0 Offset 0
OpMemberDecorate %13 1 Offset 16
OpDecorate %15 DescriptorSet 0
OpDecorate %15 Binding 0
OpDecorate %17 DescriptorSet 0
OpDecorate %17 Binding 1
OpDecorate %18 Block
OpMemberDecorate %18 0 Offset 0
%2 = OpTypeVoid
%3 = OpTypeFloat 32
%4 = OpTypeVector %3 3
%5 = OpTypeAccelerationStructureKHR
%6 = OpTypeInt 32 0
%7 = OpTypeVector %3 2
%8 = OpTypeBool
%9 = OpTypeMatrix %4 4
%10 = OpTypeStruct %6 %3 %6 %6 %6 %6 %6 %7 %8 %9 %9
%11 = OpTypeRayQueryKHR
%12 = OpTypeStruct %6 %6 %3 %3 %4 %4
%13 = OpTypeStruct %6 %4
%14 = OpTypeVector %3 4
%16 = OpTypePointer UniformConstant %5
%15 = OpVariable %16 UniformConstant
%18 = OpTypeStruct %13
%19 = OpTypePointer StorageBuffer %18
%17 = OpVariable %19 StorageBuffer
%26 = OpTypeFunction %10 %4 %4 %16
%27 = OpConstant %6 4
%28 = OpConstant %6 255
%29 = OpConstant %3 0.1
%30 = OpConstant %3 100
%32 = OpTypePointer Function %11
%33 = OpTypePointer Function %6
%35 = OpConstant %6 0
%36 = OpTypePointer Function %3
%38 = OpConstant %3 0
%41 = OpTypeVector %8 3
%42 = OpTypeFunction %2 %32 %5 %12 %33 %36
%70 = OpConstant %6 256
%73 = OpConstant %6 512
%78 = OpConstant %6 16
%81 = OpConstant %6 32
%90 = OpConstant %6 1
%93 = OpConstant %6 2
%96 = OpConstant %6 64
%99 = OpConstant %6 128
%128 = OpTypeVector %6 2
%129 = OpTypePointer Function %128
%130 = OpTypeVector %8 2
%131 = OpConstantComposite %128 %35 %35
%132 = OpConstant %6 4294967295
%133 = OpConstantComposite %128 %132 %132
%146 = OpTypePointer Function %8
%147 = OpTypeFunction %8 %32 %33
%153 = OpConstantFalse %8
%160 = OpConstant %6 6
%168 = OpTypePointer Function %10
%169 = OpTypePointer Function %9
%170 = OpTypePointer Function %7
%171 = OpTypeFunction %10 %32 %33
%176 = OpConstantNull %10
%199 = OpConstant %6 3
%202 = OpConstant %6 5
%205 = OpConstant %6 9
%207 = OpConstant %6 10
%216 = OpConstant %6 7
%218 = OpConstant %6 8
%226 = OpTypeFunction %4 %4 %10
%227 = OpConstant %3 1
%228 = OpConstant %3 2.4
%243 = OpTypeFunction %2
%245 = OpTypePointer StorageBuffer %13
%247 = OpConstantComposite %4 %38 %38 %38
%248 = OpConstantComposite %4 %38 %227 %38
%251 = OpTypePointer StorageBuffer %6
%256 = OpTypePointer StorageBuffer %4
%264 = OpConstantComposite %12 %27 %28 %29 %30 %247 %248
%265 = OpConstant %3 10
%322 = OpTypeFunction %2 %32 %33 %3 %36
%363 = OpTypeFunction %2 %32 %33
%43 = OpFunction %2 None %42
%44 = OpFunctionParameter %32
%45 = OpFunctionParameter %5
%46 = OpFunctionParameter %12
%47 = OpFunctionParameter %33
%48 = OpFunctionParameter %36
%49 = OpLabel
%50 = OpCompositeExtract %6 %46 0
%51 = OpCompositeExtract %6 %46 1
%52 = OpCompositeExtract %3 %46 2
%53 = OpCompositeExtract %3 %46 3
OpStore %48 %53
%54 = OpCompositeExtract %4 %46 4
%55 = OpCompositeExtract %4 %46 5
%56 = OpFOrdLessThanEqual %8 %52 %53
%57 = OpFOrdGreaterThanEqual %8 %52 %38
%58 = OpIsInf %41 %54
%59 = OpAny %8 %58
%60 = OpIsNan %41 %54
%61 = OpAny %8 %60
%62 = OpLogicalOr %8 %61 %59
%63 = OpLogicalNot %8 %62
%64 = OpIsInf %41 %55
%65 = OpAny %8 %64
%66 = OpIsNan %41 %55
%67 = OpAny %8 %66
%68 = OpLogicalOr %8 %67 %65
%69 = OpLogicalNot %8 %68
%71 = OpBitwiseAnd %6 %50 %70
%72 = OpINotEqual %8 %71 %35
%74 = OpBitwiseAnd %6 %50 %73
%75 = OpINotEqual %8 %74 %35
%76 = OpLogicalAnd %8 %75 %72
%77 = OpLogicalNot %8 %76
%79 = OpBitwiseAnd %6 %50 %78
%80 = OpINotEqual %8 %79 %35
%82 = OpBitwiseAnd %6 %50 %81
%83 = OpINotEqual %8 %82 %35
%84 = OpLogicalAnd %8 %83 %72
%85 = OpLogicalAnd %8 %83 %80
%86 = OpLogicalAnd %8 %80 %72
%87 = OpLogicalOr %8 %86 %84
%88 = OpLogicalOr %8 %87 %85
%89 = OpLogicalNot %8 %88
%91 = OpBitwiseAnd %6 %50 %90
%92 = OpINotEqual %8 %91 %35
%94 = OpBitwiseAnd %6 %50 %93
%95 = OpINotEqual %8 %94 %35
%97 = OpBitwiseAnd %6 %50 %96
%98 = OpINotEqual %8 %97 %35
%100 = OpBitwiseAnd %6 %50 %99
%101 = OpINotEqual %8 %100 %35
%102 = OpLogicalAnd %8 %101 %92
%103 = OpLogicalAnd %8 %101 %95
%104 = OpLogicalAnd %8 %101 %98
%105 = OpLogicalAnd %8 %98 %92
%106 = OpLogicalAnd %8 %98 %95
%107 = OpLogicalAnd %8 %95 %92
%108 = OpLogicalOr %8 %107 %102
%109 = OpLogicalOr %8 %108 %103
%110 = OpLogicalOr %8 %109 %104
%111 = OpLogicalOr %8 %110 %105
%112 = OpLogicalOr %8 %111 %106
%113 = OpLogicalNot %8 %112
%114 = OpLogicalAnd %8 %113 %56
%115 = OpLogicalAnd %8 %114 %57
%116 = OpLogicalAnd %8 %115 %63
%117 = OpLogicalAnd %8 %116 %69
%118 = OpLogicalAnd %8 %117 %77
%119 = OpLogicalAnd %8 %118 %89
OpSelectionMerge %120 None
OpBranchConditional %119 %122 %121
%122 = OpLabel
OpRayQueryInitializeKHR %44 %45 %50 %51 %54 %52 %55 %53
OpStore %47 %90
OpBranch %120
%121 = OpLabel
OpBranch %120
%120 = OpLabel
OpReturn
OpFunctionEnd
%148 = OpFunction %8 None %147
%149 = OpFunctionParameter %32
%150 = OpFunctionParameter %33
%151 = OpLabel
%152 = OpVariable %146 Function %153
%154 = OpLoad %6 %150
%157 = OpBitwiseAnd %6 %154 %90
%158 = OpINotEqual %8 %157 %35
OpSelectionMerge %155 None
OpBranchConditional %158 %156 %155
%156 = OpLabel
%159 = OpRayQueryProceedKHR %8 %149
OpStore %152 %159
%161 = OpSelect %6 %159 %93 %160
%162 = OpBitwiseOr %6 %154 %161
OpStore %150 %162
OpBranch %155
%155 = OpLabel
%163 = OpLoad %8 %152
OpReturnValue %163
OpFunctionEnd
%172 = OpFunction %10 None %171
%173 = OpFunctionParameter %32
%174 = OpFunctionParameter %33
%175 = OpLabel
%177 = OpVariable %168 Function %176
%178 = OpLoad %6 %174
%179 = OpBitwiseAnd %6 %178 %93
%180 = OpINotEqual %8 %179 %35
%181 = OpBitwiseAnd %6 %178 %27
%182 = OpINotEqual %8 %181 %35
%183 = OpLogicalAnd %8 %182 %180
OpSelectionMerge %185 None
OpBranchConditional %183 %184 %185
%184 = OpLabel
%186 = OpRayQueryGetIntersectionTypeKHR %6 %173 %90
%187 = OpAccessChain %33 %177 %35
OpStore %187 %186
%188 = OpINotEqual %8 %186 %35
OpSelectionMerge %190 None
OpBranchConditional %188 %189 %190
%189 = OpLabel
%191 = OpRayQueryGetIntersectionInstanceCustomIndexKHR %6 %173 %90
%192 = OpRayQueryGetIntersectionInstanceIdKHR %6 %173 %90
%193 = OpRayQueryGetIntersectionInstanceShaderBindingTableRecordOffsetKHR %6 %173 %90
%194 = OpRayQueryGetIntersectionGeometryIndexKHR %6 %173 %90
%195 = OpRayQueryGetIntersectionPrimitiveIndexKHR %6 %173 %90
%196 = OpRayQueryGetIntersectionObjectToWorldKHR %9 %173 %90
%197 = OpRayQueryGetIntersectionWorldToObjectKHR %9 %173 %90
%198 = OpAccessChain %33 %177 %93
OpStore %198 %191
%200 = OpAccessChain %33 %177 %199
OpStore %200 %192
%201 = OpAccessChain %33 %177 %27
OpStore %201 %193
%203 = OpAccessChain %33 %177 %202
OpStore %203 %194
%204 = OpAccessChain %33 %177 %160
OpStore %204 %195
%206 = OpAccessChain %169 %177 %205
OpStore %206 %196
%208 = OpAccessChain %169 %177 %207
OpStore %208 %197
%209 = OpIEqual %8 %186 %90
%212 = OpRayQueryGetIntersectionTKHR %3 %173 %90
%213 = OpAccessChain %36 %177 %90
OpStore %213 %212
OpSelectionMerge %211 None
OpBranchConditional %188 %210 %211
%210 = OpLabel
%214 = OpRayQueryGetIntersectionBarycentricsKHR %7 %173 %90
%215 = OpRayQueryGetIntersectionFrontFaceKHR %8 %173 %90
%217 = OpAccessChain %170 %177 %216
OpStore %217 %214
%219 = OpAccessChain %146 %177 %218
OpStore %219 %215
OpBranch %211
%211 = OpLabel
OpBranch %190
%190 = OpLabel
OpBranch %185
%185 = OpLabel
%220 = OpLoad %10 %177
OpReturnValue %220
OpFunctionEnd
%25 = OpFunction %10 None %26
%21 = OpFunctionParameter %4
%22 = OpFunctionParameter %4
%23 = OpFunctionParameter %16
%20 = OpLabel
%31 = OpVariable %32 Function
%34 = OpVariable %33 Function %35
%37 = OpVariable %36 Function %38
%134 = OpVariable %129 Function %133
%24 = OpLoad %5 %23
OpBranch %39
%39 = OpLabel
%40 = OpCompositeConstruct %12 %27 %28 %29 %30 %21 %22
%123 = OpFunctionCall %2 %43 %31 %24 %40 %34 %37
OpBranch %124
%124 = OpLabel
OpLoopMerge %125 %127 None
OpBranch %135
%135 = OpLabel
%136 = OpLoad %128 %134
%137 = OpIEqual %130 %131 %136
%138 = OpAll %8 %137
OpSelectionMerge %139 None
OpBranchConditional %138 %125 %139
%139 = OpLabel
%140 = OpCompositeExtract %6 %136 1
%141 = OpIEqual %8 %140 %35
%142 = OpSelect %6 %141 %90 %35
%143 = OpCompositeConstruct %128 %142 %90
%144 = OpISub %128 %136 %143
OpStore %134 %144
OpBranch %126
%126 = OpLabel
%145 = OpFunctionCall %8 %148 %31 %34
OpSelectionMerge %164 None
OpBranchConditional %145 %164 %165
%165 = OpLabel
OpBranch %125
%164 = OpLabel
OpBranch %166
%166 = OpLabel
OpBranch %167
%167 = OpLabel
OpBranch %127
%127 = OpLabel
OpBranch %124
%125 = OpLabel
%221 = OpFunctionCall %10 %172 %31 %34
OpReturnValue %221
OpFunctionEnd
%225 = OpFunction %4 None %226
%223 = OpFunctionParameter %4
%224 = OpFunctionParameter %10
%222 = OpLabel
OpBranch %229
%229 = OpLabel
%230 = OpCompositeExtract %9 %224 10
%231 = OpCompositeConstruct %14 %223 %227
%232 = OpMatrixTimesVector %4 %230 %231
%233 = OpVectorShuffle %7 %232 %232 0 1
%234 = OpExtInst %7 %1 Normalize %233
%235 = OpVectorTimesScalar %7 %234 %228
%236 = OpCompositeExtract %9 %224 9
%237 = OpCompositeConstruct %14 %235 %38 %227
%238 = OpMatrixTimesVector %4 %236 %237
%239 = OpFSub %4 %223 %238
%240 = OpExtInst %4 %1 Normalize %239
OpReturnValue %240
OpFunctionEnd
%242 = OpFunction %2 None %243
%241 = OpLabel
%244 = OpLoad %5 %15
%246 = OpAccessChain %245 %17 %35
OpBranch %249
%249 = OpLabel
%250 = OpFunctionCall %10 %25 %247 %248 %15
%252 = OpCompositeExtract %6 %250 0
%253 = OpIEqual %8 %252 %35
%254 = OpSelect %6 %253 %90 %35
%255 = OpAccessChain %251 %246 %35
OpStore %255 %254
%257 = OpCompositeExtract %3 %250 1
%258 = OpVectorTimesScalar %4 %248 %257
%259 = OpFunctionCall %4 %225 %258 %250
%260 = OpAccessChain %256 %246 %90
OpStore %260 %259
OpReturn
OpFunctionEnd
%271 = OpFunction %10 None %171
%272 = OpFunctionParameter %32
%273 = OpFunctionParameter %33
%274 = OpLabel
%275 = OpVariable %168 Function %176
%276 = OpLoad %6 %273
%277 = OpBitwiseAnd %6 %276 %93
%278 = OpINotEqual %8 %277 %35
%279 = OpBitwiseAnd %6 %276 %27
%280 = OpINotEqual %8 %279 %35
%281 = OpLogicalNot %8 %280
%282 = OpLogicalAnd %8 %281 %278
OpSelectionMerge %284 None
OpBranchConditional %282 %283 %284
%283 = OpLabel
%285 = OpRayQueryGetIntersectionTypeKHR %6 %272 %35
%286 = OpIEqual %8 %285 %35
%287 = OpSelect %6 %286 %90 %199
%288 = OpAccessChain %33 %275 %35
OpStore %288 %287
%289 = OpINotEqual %8 %287 %35
OpSelectionMerge %291 None
OpBranchConditional %289 %290 %291
%290 = OpLabel
%292 = OpRayQueryGetIntersectionInstanceCustomIndexKHR %6 %272 %35
%293 = OpRayQueryGetIntersectionInstanceIdKHR %6 %272 %35
%294 = OpRayQueryGetIntersectionInstanceShaderBindingTableRecordOffsetKHR %6 %272 %35
%295 = OpRayQueryGetIntersectionGeometryIndexKHR %6 %272 %35
%296 = OpRayQueryGetIntersectionPrimitiveIndexKHR %6 %272 %35
%297 = OpRayQueryGetIntersectionObjectToWorldKHR %9 %272 %35
%298 = OpRayQueryGetIntersectionWorldToObjectKHR %9 %272 %35
%299 = OpAccessChain %33 %275 %93
OpStore %299 %292
%300 = OpAccessChain %33 %275 %199
OpStore %300 %293
%301 = OpAccessChain %33 %275 %27
OpStore %301 %294
%302 = OpAccessChain %33 %275 %202
OpStore %302 %295
%303 = OpAccessChain %33 %275 %160
OpStore %303 %296
%304 = OpAccessChain %169 %275 %205
OpStore %304 %297
%305 = OpAccessChain %169 %275 %207
OpStore %305 %298
%306 = OpIEqual %8 %287 %90
OpSelectionMerge %308 None
OpBranchConditional %289 %307 %308
%307 = OpLabel
%309 = OpRayQueryGetIntersectionTKHR %3 %272 %35
%310 = OpAccessChain %36 %275 %90
OpStore %310 %309
%311 = OpRayQueryGetIntersectionBarycentricsKHR %7 %272 %35
%312 = OpRayQueryGetIntersectionFrontFaceKHR %8 %272 %35
%313 = OpAccessChain %170 %275 %216
OpStore %313 %311
%314 = OpAccessChain %146 %275 %218
OpStore %314 %312
OpBranch %308
%308 = OpLabel
OpBranch %291
%291 = OpLabel
OpBranch %284
%284 = OpLabel
%315 = OpLoad %10 %275
OpReturnValue %315
OpFunctionEnd
%323 = OpFunction %2 None %322
%324 = OpFunctionParameter %32
%325 = OpFunctionParameter %33
%326 = OpFunctionParameter %3
%327 = OpFunctionParameter %36
%328 = OpLabel
%329 = OpVariable %36 Function
%330 = OpVariable %36 Function
%333 = OpLoad %6 %325
%334 = OpBitwiseAnd %6 %333 %93
%335 = OpINotEqual %8 %334 %35
%336 = OpBitwiseAnd %6 %333 %27
%337 = OpINotEqual %8 %336 %35
%338 = OpLogicalNot %8 %337
%339 = OpLogicalAnd %8 %338 %335
OpSelectionMerge %332 None
OpBranchConditional %339 %331 %332
%331 = OpLabel
%340 = OpRayQueryGetIntersectionTypeKHR %6 %324 %35
%341 = OpIEqual %8 %340 %90
%342 = OpRayQueryGetRayTMinKHR %3 %324
%343 = OpRayQueryGetIntersectionTypeKHR %6 %324 %90
%344 = OpIEqual %8 %343 %35
OpSelectionMerge %345 None
OpBranchConditional %344 %346 %347
%346 = OpLabel
%348 = OpLoad %3 %327
OpStore %330 %348
OpBranch %345
%347 = OpLabel
%349 = OpRayQueryGetIntersectionTKHR %3 %324 %35
OpStore %330 %349
OpBranch %345
%345 = OpLabel
%350 = OpFOrdGreaterThanEqual %8 %326 %342
%351 = OpLoad %3 %330
%352 = OpFOrdLessThanEqual %8 %326 %351
%353 = OpLogicalAnd %8 %350 %352
%354 = OpLogicalAnd %8 %353 %341
OpSelectionMerge %356 None
OpBranchConditional %354 %355 %356
%355 = OpLabel
OpRayQueryGenerateIntersectionKHR %324 %326
OpBranch %356
%356 = OpLabel
OpBranch %332
%332 = OpLabel
OpReturn
OpFunctionEnd
%364 = OpFunction %2 None %363
%365 = OpFunctionParameter %32
%366 = OpFunctionParameter %33
%367 = OpLabel
%370 = OpLoad %6 %366
%371 = OpBitwiseAnd %6 %370 %93
%372 = OpINotEqual %8 %371 %35
%373 = OpBitwiseAnd %6 %370 %27
%374 = OpINotEqual %8 %373 %35
%375 = OpLogicalNot %8 %374
%376 = OpLogicalAnd %8 %375 %372
OpSelectionMerge %369 None
OpBranchConditional %376 %368 %369
%368 = OpLabel
%377 = OpRayQueryGetIntersectionTypeKHR %6 %365 %35
%378 = OpIEqual %8 %377 %35
OpSelectionMerge %380 None
OpBranchConditional %378 %379 %380
%379 = OpLabel
OpRayQueryConfirmIntersectionKHR %365
OpBranch %380
%380 = OpLabel
OpBranch %369
%369 = OpLabel
OpReturn
OpFunctionEnd
%262 = OpFunction %2 None %243
%261 = OpLabel
%266 = OpVariable %32 Function
%267 = OpVariable %33 Function %35
%268 = OpVariable %36 Function %38
%263 = OpLoad %5 %15
OpBranch %269
%269 = OpLabel
%270 = OpFunctionCall %2 %43 %266 %263 %264 %267 %268
%316 = OpFunctionCall %10 %271 %266 %267
%317 = OpCompositeExtract %6 %316 0
%318 = OpIEqual %8 %317 %199
OpSelectionMerge %319 None
OpBranchConditional %318 %320 %321
%320 = OpLabel
%357 = OpFunctionCall %2 %323 %266 %267 %265 %268
OpReturn
%321 = OpLabel
%358 = OpCompositeExtract %6 %316 0
%359 = OpIEqual %8 %358 %90
OpSelectionMerge %360 None
OpBranchConditional %359 %361 %362
%361 = OpLabel
%381 = OpFunctionCall %2 %364 %266 %267
OpReturn
%362 = OpLabel
OpReturn
%360 = OpLabel
OpBranch %319
%319 = OpLabel
OpReturn
OpFunctionEnd

716
naga/tests/out/spv/wgsl-ray-query.spvasm
View File

@ -1,16 +1,16 @@
; SPIR-V
; Version: 1.4
; Generator: rspirv
; Bound: 218
; Bound: 382
OpCapability Shader
OpCapability RayQueryKHR
OpExtension "SPV_KHR_ray_query"
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %140 "main" %15 %17
OpEntryPoint GLCompute %160 "main_candidate" %15
OpExecutionMode %140 LocalSize 1 1 1
OpExecutionMode %160 LocalSize 1 1 1
OpEntryPoint GLCompute %242 "main" %15 %17
OpEntryPoint GLCompute %262 "main_candidate" %15
OpExecutionMode %242 LocalSize 1 1 1
OpExecutionMode %262 LocalSize 1 1 1
OpMemberDecorate %10 0 Offset 0
OpMemberDecorate %10 1 Offset 4
OpMemberDecorate %10 2 Offset 8
@ -64,93 +64,219 @@ OpMemberDecorate %18 0 Offset 0
%29 = OpConstant %3 0.1
%30 = OpConstant %3 100
%32 = OpTypePointer Function %11
%45 = OpTypeVector %6 2
%46 = OpTypePointer Function %45
%47 = OpTypeVector %8 2
%48 = OpConstant %6 0
%49 = OpConstantComposite %45 %48 %48
%50 = OpConstant %6 1
%51 = OpConstant %6 4294967295
%52 = OpConstantComposite %45 %51 %51
%69 = OpTypePointer Function %10
%70 = OpTypePointer Function %6
%71 = OpTypePointer Function %9
%72 = OpTypePointer Function %7
%73 = OpTypePointer Function %8
%74 = OpTypePointer Function %3
%75 = OpTypeFunction %10 %32
%77 = OpConstantNull %10
%33 = OpTypePointer Function %6
%35 = OpConstant %6 0
%36 = OpTypePointer Function %3
%38 = OpConstant %3 0
%41 = OpTypeVector %8 3
%42 = OpTypeFunction %2 %32 %5 %12 %33 %36
%70 = OpConstant %6 256
%73 = OpConstant %6 512
%78 = OpConstant %6 16
%81 = OpConstant %6 32
%90 = OpConstant %6 1
%93 = OpConstant %6 2
%95 = OpConstant %6 3
%98 = OpConstant %6 5
%100 = OpConstant %6 6
%102 = OpConstant %6 9
%104 = OpConstant %6 10
%113 = OpConstant %6 7
%115 = OpConstant %6 8
%123 = OpTypeFunction %4 %4 %10
%124 = OpConstant %3 1
%125 = OpConstant %3 2.4
%126 = OpConstant %3 0
%141 = OpTypeFunction %2
%143 = OpTypePointer StorageBuffer %13
%145 = OpConstantComposite %4 %126 %126 %126
%146 = OpConstantComposite %4 %126 %124 %126
%149 = OpTypePointer StorageBuffer %6
%154 = OpTypePointer StorageBuffer %4
%162 = OpConstantComposite %12 %27 %28 %29 %30 %145 %146
%163 = OpConstant %3 10
%76 = OpFunction %10 None %75
%78 = OpFunctionParameter %32
%79 = OpLabel
%80 = OpVariable %69 Function %77
%81 = OpRayQueryGetIntersectionTypeKHR %6 %78 %50
%82 = OpAccessChain %70 %80 %48
OpStore %82 %81
%83 = OpINotEqual %8 %81 %48
OpSelectionMerge %85 None
OpBranchConditional %83 %84 %85
%84 = OpLabel
%86 = OpRayQueryGetIntersectionInstanceCustomIndexKHR %6 %78 %50
%87 = OpRayQueryGetIntersectionInstanceIdKHR %6 %78 %50
%88 = OpRayQueryGetIntersectionInstanceShaderBindingTableRecordOffsetKHR %6 %78 %50
%89 = OpRayQueryGetIntersectionGeometryIndexKHR %6 %78 %50
%90 = OpRayQueryGetIntersectionPrimitiveIndexKHR %6 %78 %50
%91 = OpRayQueryGetIntersectionObjectToWorldKHR %9 %78 %50
%92 = OpRayQueryGetIntersectionWorldToObjectKHR %9 %78 %50
%94 = OpAccessChain %70 %80 %93
OpStore %94 %86
%96 = OpAccessChain %70 %80 %95
OpStore %96 %87
%97 = OpAccessChain %70 %80 %27
OpStore %97 %88
%99 = OpAccessChain %70 %80 %98
OpStore %99 %89
%101 = OpAccessChain %70 %80 %100
OpStore %101 %90
%103 = OpAccessChain %71 %80 %102
OpStore %103 %91
%105 = OpAccessChain %71 %80 %104
OpStore %105 %92
%106 = OpIEqual %8 %81 %50
%109 = OpRayQueryGetIntersectionTKHR %3 %78 %50
%110 = OpAccessChain %74 %80 %50
OpStore %110 %109
OpSelectionMerge %108 None
OpBranchConditional %83 %107 %108
%107 = OpLabel
%111 = OpRayQueryGetIntersectionBarycentricsKHR %7 %78 %50
%112 = OpRayQueryGetIntersectionFrontFaceKHR %8 %78 %50
%114 = OpAccessChain %72 %80 %113
OpStore %114 %111
%116 = OpAccessChain %73 %80 %115
OpStore %116 %112
OpBranch %108
%108 = OpLabel
OpBranch %85
%85 = OpLabel
%117 = OpLoad %10 %80
OpReturnValue %117
%96 = OpConstant %6 64
%99 = OpConstant %6 128
%128 = OpTypeVector %6 2
%129 = OpTypePointer Function %128
%130 = OpTypeVector %8 2
%131 = OpConstantComposite %128 %35 %35
%132 = OpConstant %6 4294967295
%133 = OpConstantComposite %128 %132 %132
%146 = OpTypePointer Function %8
%147 = OpTypeFunction %8 %32 %33
%153 = OpConstantFalse %8
%160 = OpConstant %6 6
%168 = OpTypePointer Function %10
%169 = OpTypePointer Function %9
%170 = OpTypePointer Function %7
%171 = OpTypeFunction %10 %32 %33
%176 = OpConstantNull %10
%199 = OpConstant %6 3
%202 = OpConstant %6 5
%205 = OpConstant %6 9
%207 = OpConstant %6 10
%216 = OpConstant %6 7
%218 = OpConstant %6 8
%226 = OpTypeFunction %4 %4 %10
%227 = OpConstant %3 1
%228 = OpConstant %3 2.4
%243 = OpTypeFunction %2
%245 = OpTypePointer StorageBuffer %13
%247 = OpConstantComposite %4 %38 %38 %38
%248 = OpConstantComposite %4 %38 %227 %38
%251 = OpTypePointer StorageBuffer %6
%256 = OpTypePointer StorageBuffer %4
%264 = OpConstantComposite %12 %27 %28 %29 %30 %247 %248
%265 = OpConstant %3 10
%322 = OpTypeFunction %2 %32 %33 %3 %36
%363 = OpTypeFunction %2 %32 %33
%43 = OpFunction %2 None %42
%44 = OpFunctionParameter %32
%45 = OpFunctionParameter %5
%46 = OpFunctionParameter %12
%47 = OpFunctionParameter %33
%48 = OpFunctionParameter %36
%49 = OpLabel
%50 = OpCompositeExtract %6 %46 0
%51 = OpCompositeExtract %6 %46 1
%52 = OpCompositeExtract %3 %46 2
%53 = OpCompositeExtract %3 %46 3
OpStore %48 %53
%54 = OpCompositeExtract %4 %46 4
%55 = OpCompositeExtract %4 %46 5
%56 = OpFOrdLessThanEqual %8 %52 %53
%57 = OpFOrdGreaterThanEqual %8 %52 %38
%58 = OpIsInf %41 %54
%59 = OpAny %8 %58
%60 = OpIsNan %41 %54
%61 = OpAny %8 %60
%62 = OpLogicalOr %8 %61 %59
%63 = OpLogicalNot %8 %62
%64 = OpIsInf %41 %55
%65 = OpAny %8 %64
%66 = OpIsNan %41 %55
%67 = OpAny %8 %66
%68 = OpLogicalOr %8 %67 %65
%69 = OpLogicalNot %8 %68
%71 = OpBitwiseAnd %6 %50 %70
%72 = OpINotEqual %8 %71 %35
%74 = OpBitwiseAnd %6 %50 %73
%75 = OpINotEqual %8 %74 %35
%76 = OpLogicalAnd %8 %75 %72
%77 = OpLogicalNot %8 %76
%79 = OpBitwiseAnd %6 %50 %78
%80 = OpINotEqual %8 %79 %35
%82 = OpBitwiseAnd %6 %50 %81
%83 = OpINotEqual %8 %82 %35
%84 = OpLogicalAnd %8 %83 %72
%85 = OpLogicalAnd %8 %83 %80
%86 = OpLogicalAnd %8 %80 %72
%87 = OpLogicalOr %8 %86 %84
%88 = OpLogicalOr %8 %87 %85
%89 = OpLogicalNot %8 %88
%91 = OpBitwiseAnd %6 %50 %90
%92 = OpINotEqual %8 %91 %35
%94 = OpBitwiseAnd %6 %50 %93
%95 = OpINotEqual %8 %94 %35
%97 = OpBitwiseAnd %6 %50 %96
%98 = OpINotEqual %8 %97 %35
%100 = OpBitwiseAnd %6 %50 %99
%101 = OpINotEqual %8 %100 %35
%102 = OpLogicalAnd %8 %101 %92
%103 = OpLogicalAnd %8 %101 %95
%104 = OpLogicalAnd %8 %101 %98
%105 = OpLogicalAnd %8 %98 %92
%106 = OpLogicalAnd %8 %98 %95
%107 = OpLogicalAnd %8 %95 %92
%108 = OpLogicalOr %8 %107 %102
%109 = OpLogicalOr %8 %108 %103
%110 = OpLogicalOr %8 %109 %104
%111 = OpLogicalOr %8 %110 %105
%112 = OpLogicalOr %8 %111 %106
%113 = OpLogicalNot %8 %112
%114 = OpLogicalAnd %8 %113 %56
%115 = OpLogicalAnd %8 %114 %57
%116 = OpLogicalAnd %8 %115 %63
%117 = OpLogicalAnd %8 %116 %69
%118 = OpLogicalAnd %8 %117 %77
%119 = OpLogicalAnd %8 %118 %89
OpSelectionMerge %120 None
OpBranchConditional %119 %122 %121
%122 = OpLabel
OpRayQueryInitializeKHR %44 %45 %50 %51 %54 %52 %55 %53
OpStore %47 %90
OpBranch %120
%121 = OpLabel
OpBranch %120
%120 = OpLabel
OpReturn
OpFunctionEnd
%148 = OpFunction %8 None %147
%149 = OpFunctionParameter %32
%150 = OpFunctionParameter %33
%151 = OpLabel
%152 = OpVariable %146 Function %153
%154 = OpLoad %6 %150
%157 = OpBitwiseAnd %6 %154 %90
%158 = OpINotEqual %8 %157 %35
OpSelectionMerge %155 None
OpBranchConditional %158 %156 %155
%156 = OpLabel
%159 = OpRayQueryProceedKHR %8 %149
OpStore %152 %159
%161 = OpSelect %6 %159 %93 %160
%162 = OpBitwiseOr %6 %154 %161
OpStore %150 %162
OpBranch %155
%155 = OpLabel
%163 = OpLoad %8 %152
OpReturnValue %163
OpFunctionEnd
%172 = OpFunction %10 None %171
%173 = OpFunctionParameter %32
%174 = OpFunctionParameter %33
%175 = OpLabel
%177 = OpVariable %168 Function %176
%178 = OpLoad %6 %174
%179 = OpBitwiseAnd %6 %178 %93
%180 = OpINotEqual %8 %179 %35
%181 = OpBitwiseAnd %6 %178 %27
%182 = OpINotEqual %8 %181 %35
%183 = OpLogicalAnd %8 %182 %180
OpSelectionMerge %185 None
OpBranchConditional %183 %184 %185
%184 = OpLabel
%186 = OpRayQueryGetIntersectionTypeKHR %6 %173 %90
%187 = OpAccessChain %33 %177 %35
OpStore %187 %186
%188 = OpINotEqual %8 %186 %35
OpSelectionMerge %190 None
OpBranchConditional %188 %189 %190
%189 = OpLabel
%191 = OpRayQueryGetIntersectionInstanceCustomIndexKHR %6 %173 %90
%192 = OpRayQueryGetIntersectionInstanceIdKHR %6 %173 %90
%193 = OpRayQueryGetIntersectionInstanceShaderBindingTableRecordOffsetKHR %6 %173 %90
%194 = OpRayQueryGetIntersectionGeometryIndexKHR %6 %173 %90
%195 = OpRayQueryGetIntersectionPrimitiveIndexKHR %6 %173 %90
%196 = OpRayQueryGetIntersectionObjectToWorldKHR %9 %173 %90
%197 = OpRayQueryGetIntersectionWorldToObjectKHR %9 %173 %90
%198 = OpAccessChain %33 %177 %93
OpStore %198 %191
%200 = OpAccessChain %33 %177 %199
OpStore %200 %192
%201 = OpAccessChain %33 %177 %27
OpStore %201 %193
%203 = OpAccessChain %33 %177 %202
OpStore %203 %194
%204 = OpAccessChain %33 %177 %160
OpStore %204 %195
%206 = OpAccessChain %169 %177 %205
OpStore %206 %196
%208 = OpAccessChain %169 %177 %207
OpStore %208 %197
%209 = OpIEqual %8 %186 %90
%212 = OpRayQueryGetIntersectionTKHR %3 %173 %90
%213 = OpAccessChain %36 %177 %90
OpStore %213 %212
OpSelectionMerge %211 None
OpBranchConditional %188 %210 %211
%210 = OpLabel
%214 = OpRayQueryGetIntersectionBarycentricsKHR %7 %173 %90
%215 = OpRayQueryGetIntersectionFrontFaceKHR %8 %173 %90
%217 = OpAccessChain %170 %177 %216
OpStore %217 %214
%219 = OpAccessChain %146 %177 %218
OpStore %219 %215
OpBranch %211
%211 = OpLabel
OpBranch %190
%190 = OpLabel
OpBranch %185
%185 = OpLabel
%220 = OpLoad %10 %177
OpReturnValue %220
OpFunctionEnd
%25 = OpFunction %10 None %26
%21 = OpFunctionParameter %4
@ -158,179 +284,259 @@ OpFunctionEnd
%23 = OpFunctionParameter %16
%20 = OpLabel
%31 = OpVariable %32 Function
%53 = OpVariable %46 Function %52
%34 = OpVariable %33 Function %35
%37 = OpVariable %36 Function %38
%134 = OpVariable %129 Function %133
%24 = OpLoad %5 %23
OpBranch %33
%33 = OpLabel
%34 = OpCompositeConstruct %12 %27 %28 %29 %30 %21 %22
%35 = OpCompositeExtract %6 %34 0
%36 = OpCompositeExtract %6 %34 1
%37 = OpCompositeExtract %3 %34 2
%38 = OpCompositeExtract %3 %34 3
%39 = OpCompositeExtract %4 %34 4
%40 = OpCompositeExtract %4 %34 5
OpRayQueryInitializeKHR %31 %24 %35 %36 %39 %37 %40 %38
OpBranch %41
%41 = OpLabel
OpLoopMerge %42 %44 None
OpBranch %54
%54 = OpLabel
%55 = OpLoad %45 %53
%56 = OpIEqual %47 %49 %55
%57 = OpAll %8 %56
OpSelectionMerge %58 None
OpBranchConditional %57 %42 %58
%58 = OpLabel
%59 = OpCompositeExtract %6 %55 1
%60 = OpIEqual %8 %59 %48
%61 = OpSelect %6 %60 %50 %48
%62 = OpCompositeConstruct %45 %61 %50
%63 = OpISub %45 %55 %62
OpStore %53 %63
OpBranch %43
%43 = OpLabel
%64 = OpRayQueryProceedKHR %8 %31
OpSelectionMerge %65 None
OpBranchConditional %64 %65 %66
%66 = OpLabel
OpBranch %42
%65 = OpLabel
OpBranch %67
%67 = OpLabel
OpBranch %68
%68 = OpLabel
OpBranch %44
%44 = OpLabel
OpBranch %41
%42 = OpLabel
%118 = OpFunctionCall %10 %76 %31
OpReturnValue %118
OpFunctionEnd
%122 = OpFunction %4 None %123
%120 = OpFunctionParameter %4
%121 = OpFunctionParameter %10
%119 = OpLabel
OpBranch %39
%39 = OpLabel
%40 = OpCompositeConstruct %12 %27 %28 %29 %30 %21 %22
%123 = OpFunctionCall %2 %43 %31 %24 %40 %34 %37
OpBranch %124
%124 = OpLabel
OpLoopMerge %125 %127 None
OpBranch %135
%135 = OpLabel
%136 = OpLoad %128 %134
%137 = OpIEqual %130 %131 %136
%138 = OpAll %8 %137
OpSelectionMerge %139 None
OpBranchConditional %138 %125 %139
%139 = OpLabel
%140 = OpCompositeExtract %6 %136 1
%141 = OpIEqual %8 %140 %35
%142 = OpSelect %6 %141 %90 %35
%143 = OpCompositeConstruct %128 %142 %90
%144 = OpISub %128 %136 %143
OpStore %134 %144
OpBranch %126
%126 = OpLabel
%145 = OpFunctionCall %8 %148 %31 %34
OpSelectionMerge %164 None
OpBranchConditional %145 %164 %165
%165 = OpLabel
OpBranch %125
%164 = OpLabel
OpBranch %166
%166 = OpLabel
OpBranch %167
%167 = OpLabel
OpBranch %127
%127 = OpLabel
%128 = OpCompositeExtract %9 %121 10
%129 = OpCompositeConstruct %14 %120 %124
%130 = OpMatrixTimesVector %4 %128 %129
%131 = OpVectorShuffle %7 %130 %130 0 1
%132 = OpExtInst %7 %1 Normalize %131
%133 = OpVectorTimesScalar %7 %132 %125
%134 = OpCompositeExtract %9 %121 9
%135 = OpCompositeConstruct %14 %133 %126 %124
%136 = OpMatrixTimesVector %4 %134 %135
%137 = OpFSub %4 %120 %136
%138 = OpExtInst %4 %1 Normalize %137
OpReturnValue %138
OpBranch %124
%125 = OpLabel
%221 = OpFunctionCall %10 %172 %31 %34
OpReturnValue %221
OpFunctionEnd
%140 = OpFunction %2 None %141
%139 = OpLabel
%142 = OpLoad %5 %15
%144 = OpAccessChain %143 %17 %48
OpBranch %147
%147 = OpLabel
%148 = OpFunctionCall %10 %25 %145 %146 %15
%150 = OpCompositeExtract %6 %148 0
%151 = OpIEqual %8 %150 %48
%152 = OpSelect %6 %151 %50 %48
%153 = OpAccessChain %149 %144 %48
OpStore %153 %152
%155 = OpCompositeExtract %3 %148 1
%156 = OpVectorTimesScalar %4 %146 %155
%157 = OpFunctionCall %4 %122 %156 %148
%158 = OpAccessChain %154 %144 %50
OpStore %158 %157
%225 = OpFunction %4 None %226
%223 = OpFunctionParameter %4
%224 = OpFunctionParameter %10
%222 = OpLabel
OpBranch %229
%229 = OpLabel
%230 = OpCompositeExtract %9 %224 10
%231 = OpCompositeConstruct %14 %223 %227
%232 = OpMatrixTimesVector %4 %230 %231
%233 = OpVectorShuffle %7 %232 %232 0 1
%234 = OpExtInst %7 %1 Normalize %233
%235 = OpVectorTimesScalar %7 %234 %228
%236 = OpCompositeExtract %9 %224 9
%237 = OpCompositeConstruct %14 %235 %38 %227
%238 = OpMatrixTimesVector %4 %236 %237
%239 = OpFSub %4 %223 %238
%240 = OpExtInst %4 %1 Normalize %239
OpReturnValue %240
OpFunctionEnd
%242 = OpFunction %2 None %243
%241 = OpLabel
%244 = OpLoad %5 %15
%246 = OpAccessChain %245 %17 %35
OpBranch %249
%249 = OpLabel
%250 = OpFunctionCall %10 %25 %247 %248 %15
%252 = OpCompositeExtract %6 %250 0
%253 = OpIEqual %8 %252 %35
%254 = OpSelect %6 %253 %90 %35
%255 = OpAccessChain %251 %246 %35
OpStore %255 %254
%257 = OpCompositeExtract %3 %250 1
%258 = OpVectorTimesScalar %4 %248 %257
%259 = OpFunctionCall %4 %225 %258 %250
%260 = OpAccessChain %256 %246 %90
OpStore %260 %259
OpReturn
OpFunctionEnd
%172 = OpFunction %10 None %75
%173 = OpFunctionParameter %32
%174 = OpLabel
%175 = OpVariable %69 Function %77
%176 = OpRayQueryGetIntersectionTypeKHR %6 %173 %48
%177 = OpIEqual %8 %176 %48
%178 = OpSelect %6 %177 %50 %95
%179 = OpAccessChain %70 %175 %48
OpStore %179 %178
%180 = OpINotEqual %8 %178 %48
OpSelectionMerge %182 None
OpBranchConditional %180 %181 %182
%181 = OpLabel
%183 = OpRayQueryGetIntersectionInstanceCustomIndexKHR %6 %173 %48
%184 = OpRayQueryGetIntersectionInstanceIdKHR %6 %173 %48
%185 = OpRayQueryGetIntersectionInstanceShaderBindingTableRecordOffsetKHR %6 %173 %48
%186 = OpRayQueryGetIntersectionGeometryIndexKHR %6 %173 %48
%187 = OpRayQueryGetIntersectionPrimitiveIndexKHR %6 %173 %48
%188 = OpRayQueryGetIntersectionObjectToWorldKHR %9 %173 %48
%189 = OpRayQueryGetIntersectionWorldToObjectKHR %9 %173 %48
%190 = OpAccessChain %70 %175 %93
OpStore %190 %183
%191 = OpAccessChain %70 %175 %95
OpStore %191 %184
%192 = OpAccessChain %70 %175 %27
OpStore %192 %185
%193 = OpAccessChain %70 %175 %98
OpStore %193 %186
%194 = OpAccessChain %70 %175 %100
OpStore %194 %187
%195 = OpAccessChain %71 %175 %102
OpStore %195 %188
%196 = OpAccessChain %71 %175 %104
OpStore %196 %189
%197 = OpIEqual %8 %178 %50
OpSelectionMerge %199 None
OpBranchConditional %180 %198 %199
%198 = OpLabel
%200 = OpRayQueryGetIntersectionTKHR %3 %173 %48
%201 = OpAccessChain %74 %175 %50
OpStore %201 %200
%202 = OpRayQueryGetIntersectionBarycentricsKHR %7 %173 %48
%203 = OpRayQueryGetIntersectionFrontFaceKHR %8 %173 %48
%204 = OpAccessChain %72 %175 %113
OpStore %204 %202
%205 = OpAccessChain %73 %175 %115
OpStore %205 %203
OpBranch %199
%199 = OpLabel
OpBranch %182
%182 = OpLabel
%206 = OpLoad %10 %175
OpReturnValue %206
%271 = OpFunction %10 None %171
%272 = OpFunctionParameter %32
%273 = OpFunctionParameter %33
%274 = OpLabel
%275 = OpVariable %168 Function %176
%276 = OpLoad %6 %273
%277 = OpBitwiseAnd %6 %276 %93
%278 = OpINotEqual %8 %277 %35
%279 = OpBitwiseAnd %6 %276 %27
%280 = OpINotEqual %8 %279 %35
%281 = OpLogicalNot %8 %280
%282 = OpLogicalAnd %8 %281 %278
OpSelectionMerge %284 None
OpBranchConditional %282 %283 %284
%283 = OpLabel
%285 = OpRayQueryGetIntersectionTypeKHR %6 %272 %35
%286 = OpIEqual %8 %285 %35
%287 = OpSelect %6 %286 %90 %199
%288 = OpAccessChain %33 %275 %35
OpStore %288 %287
%289 = OpINotEqual %8 %287 %35
OpSelectionMerge %291 None
OpBranchConditional %289 %290 %291
%290 = OpLabel
%292 = OpRayQueryGetIntersectionInstanceCustomIndexKHR %6 %272 %35
%293 = OpRayQueryGetIntersectionInstanceIdKHR %6 %272 %35
%294 = OpRayQueryGetIntersectionInstanceShaderBindingTableRecordOffsetKHR %6 %272 %35
%295 = OpRayQueryGetIntersectionGeometryIndexKHR %6 %272 %35
%296 = OpRayQueryGetIntersectionPrimitiveIndexKHR %6 %272 %35
%297 = OpRayQueryGetIntersectionObjectToWorldKHR %9 %272 %35
%298 = OpRayQueryGetIntersectionWorldToObjectKHR %9 %272 %35
%299 = OpAccessChain %33 %275 %93
OpStore %299 %292
%300 = OpAccessChain %33 %275 %199
OpStore %300 %293
%301 = OpAccessChain %33 %275 %27
OpStore %301 %294
%302 = OpAccessChain %33 %275 %202
OpStore %302 %295
%303 = OpAccessChain %33 %275 %160
OpStore %303 %296
%304 = OpAccessChain %169 %275 %205
OpStore %304 %297
%305 = OpAccessChain %169 %275 %207
OpStore %305 %298
%306 = OpIEqual %8 %287 %90
OpSelectionMerge %308 None
OpBranchConditional %289 %307 %308
%307 = OpLabel
%309 = OpRayQueryGetIntersectionTKHR %3 %272 %35
%310 = OpAccessChain %36 %275 %90
OpStore %310 %309
%311 = OpRayQueryGetIntersectionBarycentricsKHR %7 %272 %35
%312 = OpRayQueryGetIntersectionFrontFaceKHR %8 %272 %35
%313 = OpAccessChain %170 %275 %216
OpStore %313 %311
%314 = OpAccessChain %146 %275 %218
OpStore %314 %312
OpBranch %308
%308 = OpLabel
OpBranch %291
%291 = OpLabel
OpBranch %284
%284 = OpLabel
%315 = OpLoad %10 %275
OpReturnValue %315
OpFunctionEnd
%160 = OpFunction %2 None %141
%159 = OpLabel
%164 = OpVariable %32 Function
%161 = OpLoad %5 %15
OpBranch %165
%165 = OpLabel
%166 = OpCompositeExtract %6 %162 0
%167 = OpCompositeExtract %6 %162 1
%168 = OpCompositeExtract %3 %162 2
%169 = OpCompositeExtract %3 %162 3
%170 = OpCompositeExtract %4 %162 4
%171 = OpCompositeExtract %4 %162 5
OpRayQueryInitializeKHR %164 %161 %166 %167 %170 %168 %171 %169
%207 = OpFunctionCall %10 %172 %164
%208 = OpCompositeExtract %6 %207 0
%209 = OpIEqual %8 %208 %95
OpSelectionMerge %210 None
OpBranchConditional %209 %211 %212
%211 = OpLabel
OpRayQueryGenerateIntersectionKHR %164 %163
%323 = OpFunction %2 None %322
%324 = OpFunctionParameter %32
%325 = OpFunctionParameter %33
%326 = OpFunctionParameter %3
%327 = OpFunctionParameter %36
%328 = OpLabel
%329 = OpVariable %36 Function
%330 = OpVariable %36 Function
%333 = OpLoad %6 %325
%334 = OpBitwiseAnd %6 %333 %93
%335 = OpINotEqual %8 %334 %35
%336 = OpBitwiseAnd %6 %333 %27
%337 = OpINotEqual %8 %336 %35
%338 = OpLogicalNot %8 %337
%339 = OpLogicalAnd %8 %338 %335
OpSelectionMerge %332 None
OpBranchConditional %339 %331 %332
%331 = OpLabel
%340 = OpRayQueryGetIntersectionTypeKHR %6 %324 %35
%341 = OpIEqual %8 %340 %90
%342 = OpRayQueryGetRayTMinKHR %3 %324
%343 = OpRayQueryGetIntersectionTypeKHR %6 %324 %90
%344 = OpIEqual %8 %343 %35
OpSelectionMerge %345 None
OpBranchConditional %344 %346 %347
%346 = OpLabel
%348 = OpLoad %3 %327
OpStore %330 %348
OpBranch %345
%347 = OpLabel
%349 = OpRayQueryGetIntersectionTKHR %3 %324 %35
OpStore %330 %349
OpBranch %345
%345 = OpLabel
%350 = OpFOrdGreaterThanEqual %8 %326 %342
%351 = OpLoad %3 %330
%352 = OpFOrdLessThanEqual %8 %326 %351
%353 = OpLogicalAnd %8 %350 %352
%354 = OpLogicalAnd %8 %353 %341
OpSelectionMerge %356 None
OpBranchConditional %354 %355 %356
%355 = OpLabel
OpRayQueryGenerateIntersectionKHR %324 %326
OpBranch %356
%356 = OpLabel
OpBranch %332
%332 = OpLabel
OpReturn
%212 = OpLabel
%213 = OpCompositeExtract %6 %207 0
%214 = OpIEqual %8 %213 %50
OpSelectionMerge %215 None
OpBranchConditional %214 %216 %217
%216 = OpLabel
OpRayQueryConfirmIntersectionKHR %164
OpFunctionEnd
%364 = OpFunction %2 None %363
%365 = OpFunctionParameter %32
%366 = OpFunctionParameter %33
%367 = OpLabel
%370 = OpLoad %6 %366
%371 = OpBitwiseAnd %6 %370 %93
%372 = OpINotEqual %8 %371 %35
%373 = OpBitwiseAnd %6 %370 %27
%374 = OpINotEqual %8 %373 %35
%375 = OpLogicalNot %8 %374
%376 = OpLogicalAnd %8 %375 %372
OpSelectionMerge %369 None
OpBranchConditional %376 %368 %369
%368 = OpLabel
%377 = OpRayQueryGetIntersectionTypeKHR %6 %365 %35
%378 = OpIEqual %8 %377 %35
OpSelectionMerge %380 None
OpBranchConditional %378 %379 %380
%379 = OpLabel
OpRayQueryConfirmIntersectionKHR %365
OpBranch %380
%380 = OpLabel
OpBranch %369
%369 = OpLabel
OpReturn
%217 = OpLabel
OpFunctionEnd
%262 = OpFunction %2 None %243
%261 = OpLabel
%266 = OpVariable %32 Function
%267 = OpVariable %33 Function %35
%268 = OpVariable %36 Function %38
%263 = OpLoad %5 %15
OpBranch %269
%269 = OpLabel
%270 = OpFunctionCall %2 %43 %266 %263 %264 %267 %268
%316 = OpFunctionCall %10 %271 %266 %267
%317 = OpCompositeExtract %6 %316 0
%318 = OpIEqual %8 %317 %199
OpSelectionMerge %319 None
OpBranchConditional %318 %320 %321
%320 = OpLabel
%357 = OpFunctionCall %2 %323 %266 %267 %265 %268
OpReturn
%215 = OpLabel
OpBranch %210
%210 = OpLabel
%321 = OpLabel
%358 = OpCompositeExtract %6 %316 0
%359 = OpIEqual %8 %358 %90
OpSelectionMerge %360 None
OpBranchConditional %359 %361 %362
%361 = OpLabel
%381 = OpFunctionCall %2 %364 %266 %267
OpReturn
%362 = OpLabel
OpReturn
%360 = OpLabel
OpBranch %319
%319 = OpLabel
OpReturn
OpFunctionEnd

102
tests/tests/wgpu-gpu/ray_tracing/shader.rs
View File

@ -1,17 +1,19 @@
use crate::ray_tracing::{acceleration_structure_limits, AsBuildContext};
use wgpu::util::{BufferInitDescriptor, DeviceExt};
use wgpu::{
include_wgsl, BindGroupDescriptor, BindGroupEntry, BindingResource, BufferDescriptor,
CommandEncoderDescriptor, ComputePassDescriptor, ComputePipelineDescriptor,
include_wgsl, Backends, BindGroupDescriptor, BindGroupEntry, BindingResource, BufferDescriptor,
CommandEncoderDescriptor, ComputePassDescriptor, ComputePipelineDescriptor, InstanceFlags,
};
use wgpu::{AccelerationStructureFlags, BufferUsages};
use wgpu_macros::gpu_test;
use wgpu_test::GpuTestInitializer;
use wgpu_test::{FailureCase, GpuTestInitializer};
use wgpu_test::{GpuTestConfiguration, TestParameters, TestingContext};
const STRUCT_SIZE: wgpu::BufferAddress = 176;
pub fn all_tests(tests: &mut Vec<GpuTestInitializer>) {
tests.push(ACCESS_ALL_STRUCT_MEMBERS);
tests.push(PREVENT_INVALID_RAY_QUERY_CALLS);
}
#[gpu_test]
@ -103,3 +105,97 @@ fn access_all_struct_members(ctx: TestingContext) {
ctx.queue.submit([encoder_compute.finish()]);
}
#[gpu_test]
static PREVENT_INVALID_RAY_QUERY_CALLS: GpuTestConfiguration = GpuTestConfiguration::new()
.parameters(
TestParameters::default()
.test_features_limits()
.limits(acceleration_structure_limits())
.features(wgpu::Features::EXPERIMENTAL_RAY_QUERY)
// Otherwise, mistakes in the generated code won't be caught.
.instance_flags(InstanceFlags::GPU_BASED_VALIDATION)
// not yet implemented in directx12
.skip(FailureCase::backend(Backends::DX12 | Backends::METAL)),
)
.run_sync(prevent_invalid_ray_query_calls);
fn prevent_invalid_ray_query_calls(ctx: TestingContext) {
let invalid_values_buffer = ctx.device.create_buffer_init(&BufferInitDescriptor {
label: Some("invalid values buffer"),
contents: bytemuck::cast_slice(&[f32::NAN, f32::INFINITY]),
usage: BufferUsages::STORAGE,
});
//
// Create a clean `AsBuildContext`
//
let as_ctx = AsBuildContext::new(
&ctx,
AccelerationStructureFlags::empty(),
AccelerationStructureFlags::empty(),
);
let mut encoder_build = ctx
.device
.create_command_encoder(&CommandEncoderDescriptor {
label: Some("Build"),
});
encoder_build.build_acceleration_structures([&as_ctx.blas_build_entry()], [&as_ctx.tlas]);
ctx.queue.submit([encoder_build.finish()]);
//
// Create shader
//
let shader = ctx
.device
.create_shader_module(include_wgsl!("shader.wgsl"));
let compute_pipeline = ctx
.device
.create_compute_pipeline(&ComputePipelineDescriptor {
label: None,
layout: None,
module: &shader,
entry_point: Some("invalid_usages"),
compilation_options: Default::default(),
cache: None,
});
let bind_group = ctx.device.create_bind_group(&BindGroupDescriptor {
label: None,
layout: &compute_pipeline.get_bind_group_layout(0),
entries: &[
BindGroupEntry {
binding: 0,
resource: BindingResource::AccelerationStructure(&as_ctx.tlas),
},
BindGroupEntry {
binding: 1,
resource: BindingResource::Buffer(invalid_values_buffer.as_entire_buffer_binding()),
},
],
});
//
// Submit once to check for no issues
//
let mut encoder_compute = ctx
.device
.create_command_encoder(&CommandEncoderDescriptor::default());
{
let mut pass = encoder_compute.begin_compute_pass(&ComputePassDescriptor {
label: None,
timestamp_writes: None,
});
pass.set_pipeline(&compute_pipeline);
pass.set_bind_group(0, Some(&bind_group), &[]);
pass.dispatch_workgroups(1, 1, 1)
}
ctx.queue.submit([encoder_compute.finish()]);
}

74
tests/tests/wgpu-gpu/ray_tracing/shader.wgsl
View File

@ -48,4 +48,78 @@ fn all_of_struct() {
intersection.world_to_object,
intersection.object_to_world,
);
}
struct MaybeInvalidValues {
nan: f32,
inf: f32,
}
@group(0) @binding(1)
var<storage> invalid_values: MaybeInvalidValues;
@workgroup_size(1)
@compute
fn invalid_usages() {
{
var rq: ray_query;
// no initialize
rayQueryProceed(&rq);
let intersection = rayQueryGetCommittedIntersection(&rq);
}
{
var rq: ray_query;
rayQueryInitialize(&rq, acc_struct, RayDesc(0u, 0xFFu, 0.001, 100000.0, vec3f(0.0, 0.0, 0.0), vec3f(0.0, 0.0, 1.0)));
// no proceed
let intersection = rayQueryGetCommittedIntersection(&rq);
}
{
var rq: ray_query;
rayQueryInitialize(&rq, acc_struct, RayDesc(0u, 0xFFu, 0.001, 100000.0, vec3f(0.0, 0.0, 0.0), vec3f(0.0, 0.0, 1.0)));
rayQueryProceed(&rq);
// The acceleration structure has been set up to not generate an intersections, meaning it will be a committed intersection, not candidate.
let intersection = rayQueryGetCandidateIntersection(&rq);
}
{
var rq: ray_query;
// NaN in origin
rayQueryInitialize(&rq, acc_struct, RayDesc(0u, 0xFFu, 0.001, 100000.0, vec3f(0.0, invalid_values.nan, 0.0), vec3f(0.0, 0.0, 1.0)));
rayQueryProceed(&rq);
let intersection = rayQueryGetCommittedIntersection(&rq);
}
{
var rq: ray_query;
// Inf in origin
rayQueryInitialize(&rq, acc_struct, RayDesc(0u, 0xFFu, 0.001, 100000.0, vec3f(0.0, invalid_values.inf, 0.0), vec3f(0.0, 0.0, 1.0)));
rayQueryProceed(&rq);
let intersection = rayQueryGetCommittedIntersection(&rq);
}
{
var rq: ray_query;
// NaN in direction
rayQueryInitialize(&rq, acc_struct, RayDesc(0u, 0xFFu, 0.001, 100000.0, vec3f(0.0, 0.0, 0.0), vec3f(0.0, invalid_values.nan, 1.0)));
rayQueryProceed(&rq);
let intersection = rayQueryGetCommittedIntersection(&rq);
}
{
var rq: ray_query;
// Inf in direction
rayQueryInitialize(&rq, acc_struct, RayDesc(0u, 0xFFu, 0.001, 100000.0, vec3f(0.0, 0.0, 0.0), vec3f(0.0, invalid_values.inf, 1.0)));
rayQueryProceed(&rq);
let intersection = rayQueryGetCommittedIntersection(&rq);
}
{
var rq: ray_query;
// t_min greater than t_max
rayQueryInitialize(&rq, acc_struct, RayDesc(0u, 0xFFu, 100000.0, 0.1, vec3f(0.0, 0.0, 0.0), vec3f(0.0, 0.0, 1.0)));
rayQueryProceed(&rq);
let intersection = rayQueryGetCommittedIntersection(&rq);
}
{
var rq: ray_query;
// t_min less than 0
rayQueryInitialize(&rq, acc_struct, RayDesc(0u, 0xFFu, -0.001, 100000.0, vec3f(0.0, 0.0, 0.0), vec3f(0.0, 0.0, 1.0)));
rayQueryProceed(&rq);
let intersection = rayQueryGetCommittedIntersection(&rq);
}
}

1
wgpu-hal/src/dx12/device.rs
View File

@ -290,6 +290,7 @@ impl super::Device {
|| stage.module.runtime_checks.bounds_checks != layout.naga_options.restrict_indexing
|| stage.module.runtime_checks.force_loop_bounding
!= layout.naga_options.force_loop_bounding;
// Note: ray query initialization tracking not yet implemented
let mut temp_options;
let naga_options = if needs_temp_options {
temp_options = layout.naga_options.clone();

7
wgpu-hal/src/vulkan/adapter.rs
View File

@ -2219,6 +2219,12 @@ impl super::Adapter {
// But this requires cloning the `spv::Options` struct, which has heap allocations.
true, // could check `super::Workarounds::SEPARATE_ENTRY_POINTS`
);
flags.set(
spv::WriterFlags::PRINT_ON_RAY_QUERY_INITIALIZATION_FAIL,
self.instance.flags.contains(wgt::InstanceFlags::DEBUG)
&& (self.instance.instance_api_version >= vk::API_VERSION_1_3
|| enabled_extensions.contains(&khr::shader_non_semantic_info::NAME)),
);
if features.contains(wgt::Features::EXPERIMENTAL_RAY_QUERY) {
capabilities.push(spv::Capability::RayQueryKHR);
}
@ -2277,6 +2283,7 @@ impl super::Adapter {
spv::ZeroInitializeWorkgroupMemoryMode::Polyfill
},
force_loop_bounding: true,
ray_query_initialization_tracking: true,
use_storage_input_output_16: features.contains(wgt::Features::SHADER_F16)
&& self.phd_features.supports_storage_input_output_16(),
fake_missing_bindings: false,

4
wgpu-hal/src/vulkan/device.rs
View File

@ -762,6 +762,7 @@ impl super::Device {
};
let needs_temp_options = !runtime_checks.bounds_checks
|| !runtime_checks.force_loop_bounding
|| !runtime_checks.ray_query_initialization_tracking
|| !binding_map.is_empty()
|| naga_shader.debug_source.is_some()
|| !stage.zero_initialize_workgroup_memory;
@ -779,6 +780,9 @@ impl super::Device {
if !runtime_checks.force_loop_bounding {
temp_options.force_loop_bounding = false;
}
if !runtime_checks.ray_query_initialization_tracking {
temp_options.ray_query_initialization_tracking = false;
}
if !binding_map.is_empty() {
temp_options.binding_map = binding_map.clone();
}

15
wgpu-types/src/lib.rs
View File

@ -7977,6 +7977,20 @@ pub struct ShaderRuntimeChecks {
/// conclusions about other safety-critical code paths. This option SHOULD NOT be disabled
/// when running untrusted code.
pub force_loop_bounding: bool,
/// If false, the caller **MUST** ensure that in all passed shaders every function operating
/// on a ray query must obey these rules (functions using wgsl naming)
/// - `rayQueryInitialize` must have called before `rayQueryProceed`
/// - `rayQueryProceed` must have been called, returned true and have hit an AABB before
/// `rayQueryGenerateIntersection` is called
/// - `rayQueryProceed` must have been called, returned true and have hit a triangle before
/// `rayQueryConfirmIntersection` is called
/// - `rayQueryProceed` must have been called and have returned true before `rayQueryTerminate`,
/// `getCandidateHitVertexPositions` or `rayQueryGetCandidateIntersection` is called
/// - `rayQueryProceed` must have been called and have returned false before `rayQueryGetCommittedIntersection`
/// or `getCommittedHitVertexPositions` are called
///
/// It is the aim that these cases will not cause UB if this is set to true, but currently this will still happen on DX12 and Metal.
pub ray_query_initialization_tracking: bool,
}
impl ShaderRuntimeChecks {
@ -8009,6 +8023,7 @@ impl ShaderRuntimeChecks {
Self {
bounds_checks: all_checks,
force_loop_bounding: all_checks,
ray_query_initialization_tracking: all_checks,
}
}
}