pub struct Executable { /* private fields */ }Expand description
A program compiled once into a flat form: typed operations on four-lane word registers, static pointer arithmetic
folded, bitcasts aliased, dead operations removed. Running it is Runner’s.
Implementations§
Source§impl Executable
impl Executable
Sourcepub fn bind<'a>(
&self,
constants: impl Fn(ResourceId) -> Option<&'a [[u32; 4]]>,
buffers: impl Fn(ResourceId) -> Option<&'a [u32]>,
) -> Result<Bound<'a>, String>
pub fn bind<'a>( &self, constants: impl Fn(ResourceId) -> Option<&'a [[u32; 4]]>, buffers: impl Fn(ResourceId) -> Option<&'a [u32]>, ) -> Result<Bound<'a>, String>
Bind the buffers the program reads, as constants and buffers find them by resource. A constant buffer the
program reads and the caller does not give is an error; a buffer not given is one the backend answers reads of
(Backend::read_buffer). Nothing is copied: the caller’s buffers are read in place.
Sourcepub fn bind_elements<'a>(
&self,
constants: impl Fn(ResourceId) -> Option<&'a [[u32; 4]]>,
buffers: impl Fn(ResourceId) -> Option<&'a [u32]>,
elements: impl Fn(ResourceId, u32) -> Option<&'a [u32]>,
) -> Result<Bound<'a>, String>
pub fn bind_elements<'a>( &self, constants: impl Fn(ResourceId) -> Option<&'a [[u32; 4]]>, buffers: impl Fn(ResourceId) -> Option<&'a [u32]>, elements: impl Fn(ResourceId, u32) -> Option<&'a [u32]>, ) -> Result<Bound<'a>, String>
Executable::bind for a program that also reads arrays of buffers: elements finds the words of element
i of the array resource id. An element not given errors when a run reads it.
Sourcepub fn bind_resource_arrays<'a>(
&self,
constants: impl Fn(ResourceId) -> Option<&'a [[u32; 4]]>,
buffers: impl Fn(ResourceId) -> Option<&'a [u32]>,
elements: impl Fn(ResourceId) -> BTreeMap<u32, &'a [u32]>,
constant_elements: impl Fn(ResourceId) -> BTreeMap<u32, &'a [[u32; 4]]>,
) -> Result<Bound<'a>, String>
pub fn bind_resource_arrays<'a>( &self, constants: impl Fn(ResourceId) -> Option<&'a [[u32; 4]]>, buffers: impl Fn(ResourceId) -> Option<&'a [u32]>, elements: impl Fn(ResourceId) -> BTreeMap<u32, &'a [u32]>, constant_elements: impl Fn(ResourceId) -> BTreeMap<u32, &'a [[u32; 4]]>, ) -> Result<Bound<'a>, String>
Bind descriptor arrays by their supplied elements, including arrays of no stated size. Keys are zero-based descriptor indices, independent of D3D register numbers. Missing elements error only when read; a key past a bounded declaration’s end errors at binding. The maps are retained by the binding; the buffers themselves are borrowed without copying.
Sourcepub fn bind_resources<'a>(
&self,
resources: &'a Resources,
) -> Result<Bound<'a>, String>
pub fn bind_resources<'a>( &self, resources: &'a Resources, ) -> Result<Bound<'a>, String>
Executable::bind over a Resources.
Source§impl Executable
impl Executable
Sourcepub fn workgroup_size(&self) -> Option<[u32; 3]>
pub fn workgroup_size(&self) -> Option<[u32; 3]>
The thread group of a compute program (numthreads), or None for any other.
Sourcepub fn dispatch(
&self,
bound: &Bound<'_>,
backend: &mut dyn Backend,
buffers: &mut BTreeMap<ResourceId, Vec<u32>>,
groups: [u32; 3],
wave: u32,
) -> Result<Dispatched, String>
pub fn dispatch( &self, bound: &Bound<'_>, backend: &mut dyn Backend, buffers: &mut BTreeMap<ResourceId, Vec<u32>>, groups: [u32; 3], wave: u32, ) -> Result<Dispatched, String>
Run groups thread groups of a compute program. Each group’s invocations run in lockstep (a barrier is the
point where every invocation still in the program has arrived), thread-group memory starts at zero in every
group, and the groups run one after another in x, y, z order. buffers holds the words of every read-write
buffer by resource: writes and atomics change them and later reads see the changes, as they do between
groups. A buffer not given is one the backend answers reads of and takes writes to (Backend::read_buffer,
Backend::store_buffer). wave is the wave width the wave operations and lane built-ins report; it must
divide the thread group’s invocations into whole waves only for programs that care.
Sourcepub fn dispatch_elements(
&self,
bound: &Bound<'_>,
backend: &mut dyn Backend,
buffers: &mut BTreeMap<ResourceId, Vec<u32>>,
elements: &mut BTreeMap<(ResourceId, u32), Vec<u32>>,
groups: [u32; 3],
wave: u32,
) -> Result<Dispatched, String>
pub fn dispatch_elements( &self, bound: &Bound<'_>, backend: &mut dyn Backend, buffers: &mut BTreeMap<ResourceId, Vec<u32>>, elements: &mut BTreeMap<(ResourceId, u32), Vec<u32>>, groups: [u32; 3], wave: u32, ) -> Result<Dispatched, String>
Executable::dispatch for a program that also writes elements of arrays of buffers, elements holding
their words by resource and element.
Source§impl Executable
impl Executable
Sourcepub fn new(program: &Program) -> Result<Self, Error>
pub fn new(program: &Program) -> Result<Self, Error>
Check program (its module and its interface) and compile it. A construct the executor does not run is
refused here, by name.
Sourcepub fn function(program: &Program, name: &str) -> Result<Self, Error>
pub fn function(program: &Program, name: &str) -> Result<Self, Error>
Compile one function of a program made of functions (a hull program’s main and patch_constant, a domain
program’s or a position-less vertex program’s main; see iridia::stage_functions) to run for one
invocation. The registers it takes are its arguments in order, an array argument (the control points) as one
register per element; the members of its result are the members of iridia::FunctionSig::result.
Sourcepub fn functions(
program: &Program,
) -> Result<Option<Vec<(&'static str, Self)>>, Error>
pub fn functions( program: &Program, ) -> Result<Option<Vec<(&'static str, Self)>>, Error>
Every function of a program made of functions (iridia::stage_functions: a hull, domain, geometry or ray tracing
program, a vertex program that feeds a tessellator), each compiled by Executable::function under its name;
None for a program that is an entry point, which Executable::new takes.
Sourcepub fn sampled(&self) -> (BTreeSet<ResourceId>, BTreeSet<ResourceId>)
pub fn sampled(&self) -> (BTreeSet<ResourceId>, BTreeSet<ResourceId>)
The textures and the samplers the program’s compiled code samples, loads from or asks the size of (an operation nothing reads is not in it): the resources a caller must bind for a run that may reach any path.
Sourcepub fn needs_quad(&self) -> bool
pub fn needs_quad(&self) -> bool
Whether the program needs the other lanes of its 2 x 2 quad: it takes a screen derivative, or samples at an
implicit level of detail (sample, sample_b), whose footprint is the quad’s. A program that does not runs
the same one invocation at a time as in a quad.