mirror of
https://github.com/gfx-rs/wgpu.git
synced 2025-02-20 02:42:43 +00:00
glsl-out: Implement bounds checks for ImageLoad (#1889)
* glsl-out: Implement bounds checks for `ImageLoad` * Enable image bounds check snapshot tests for GLSL. In addition to the snapshot.rs changes, this entails adding an entry point function to `bounds-check-image-restrict.wgsl` and `bounds-check-image-rzsw.wgsl`, including appropriate data in the param.ron files. * Apply comments Snapshot test changes: Co-authored-by: Jim Blandy <jimb@red-bean.com>
This commit is contained in:
João Capucho
GitHub
parent
c7e245c083
commit
0aa7681165
@ -455,6 +455,7 @@ fn run() -> Result<(), Box<dyn std::error::Error>> {
|
||||
))?,
|
||||
¶ms.glsl,
|
||||
&pipeline_options,
|
||||
params.bounds_check_policies,
|
||||
)
|
||||
.unwrap_pretty();
|
||||
writer.write()?;
|
||||
|
||||
@ -38,6 +38,10 @@ bitflags::bitflags! {
|
||||
const FMA = 1 << 18;
|
||||
/// Texture samples query
|
||||
const TEXTURE_SAMPLES = 1 << 19;
|
||||
/// Texture levels query
|
||||
const TEXTURE_LEVELS = 1 << 20;
|
||||
/// Image size query
|
||||
const IMAGE_SIZE = 1 << 21;
|
||||
}
|
||||
}
|
||||
|
||||
@ -104,9 +108,11 @@ impl FeaturesManager {
|
||||
check_feature!(DYNAMIC_ARRAY_SIZE, 430, 310);
|
||||
check_feature!(MULTI_VIEW, 140, 310);
|
||||
// Only available on glsl core, this means that opengl es can't query the number
|
||||
// of samples in a image and neither do bound checks on the sample argument
|
||||
// of texelFecth
|
||||
// of samples nor levels in a image and neither do bound checks on the sample nor
|
||||
// the level argument of texelFecth
|
||||
check_feature!(TEXTURE_SAMPLES, 150);
|
||||
check_feature!(TEXTURE_LEVELS, 130);
|
||||
check_feature!(IMAGE_SIZE, 430, 310);
|
||||
|
||||
// Return an error if there are missing features
|
||||
if missing.is_empty() {
|
||||
@ -223,6 +229,11 @@ impl FeaturesManager {
|
||||
)?;
|
||||
}
|
||||
|
||||
if self.0.contains(Features::TEXTURE_LEVELS) && version < Version::Desktop(430) {
|
||||
// https://www.khronos.org/registry/OpenGL/extensions/ARB/ARB_texture_query_levels.txt
|
||||
writeln!(out, "#extension GL_ARB_texture_query_levels : require")?;
|
||||
}
|
||||
|
||||
Ok(())
|
||||
}
|
||||
}
|
||||
@ -376,27 +387,75 @@ impl<'a, W> Writer<'a, W> {
|
||||
}
|
||||
}
|
||||
|
||||
// Loop trough all expressions in both functions and entry points
|
||||
// We will need to pass some of the members to a closure, so we need
|
||||
// to separate them otherwise the borrow checker will complain, this
|
||||
// shouldn't be needed in rust 2021
|
||||
let &mut Self {
|
||||
module,
|
||||
info,
|
||||
ref mut features,
|
||||
entry_point,
|
||||
entry_point_idx,
|
||||
ref policies,
|
||||
..
|
||||
} = self;
|
||||
|
||||
// Loop trough all expressions in both functions and the entry point
|
||||
// to check for needed features
|
||||
for (_, expr) in self
|
||||
.module
|
||||
for (expressions, info) in module
|
||||
.functions
|
||||
.iter()
|
||||
.flat_map(|(_, f)| f.expressions.iter())
|
||||
.chain(self.entry_point.function.expressions.iter())
|
||||
.map(|(h, f)| (&f.expressions, &info[h]))
|
||||
.chain(std::iter::once((
|
||||
&entry_point.function.expressions,
|
||||
info.get_entry_point(entry_point_idx as usize),
|
||||
)))
|
||||
{
|
||||
match *expr {
|
||||
for (_, expr) in expressions.iter() {
|
||||
match *expr {
|
||||
// Check for fused multiply add use
|
||||
Expression::Math { fun, .. } if fun == MathFunction::Fma => {
|
||||
self.features.request(Features::FMA)
|
||||
features.request(Features::FMA)
|
||||
}
|
||||
// Check for samples query
|
||||
// Check for queries that neeed aditonal features
|
||||
Expression::ImageQuery {
|
||||
query: crate::ImageQuery::NumSamples,
|
||||
image,
|
||||
query,
|
||||
..
|
||||
} => self.features.request(Features::TEXTURE_SAMPLES),
|
||||
} => match query {
|
||||
// Storage images use `imageSize` which is only available
|
||||
// in glsl > 420
|
||||
//
|
||||
// layers queries are also implemented as size queries
|
||||
crate::ImageQuery::Size { .. } | crate::ImageQuery::NumLayers => {
|
||||
if let TypeInner::Image {
|
||||
class: crate::ImageClass::Storage { .. }, ..
|
||||
} = *info[image].ty.inner_with(&module.types) {
|
||||
features.request(Features::IMAGE_SIZE)
|
||||
}
|
||||
},
|
||||
crate::ImageQuery::NumLevels => features.request(Features::TEXTURE_LEVELS),
|
||||
crate::ImageQuery::NumSamples => features.request(Features::TEXTURE_SAMPLES),
|
||||
}
|
||||
,
|
||||
// Check for image loads that needs bound checking on the sample
|
||||
// or level argument since this requires a feature
|
||||
Expression::ImageLoad {
|
||||
sample, level, ..
|
||||
} => {
|
||||
if policies.image != crate::proc::BoundsCheckPolicy::Unchecked {
|
||||
if sample.is_some() {
|
||||
features.request(Features::TEXTURE_SAMPLES)
|
||||
}
|
||||
|
||||
if level.is_some() {
|
||||
features.request(Features::TEXTURE_LEVELS)
|
||||
}
|
||||
}
|
||||
}
|
||||
_ => {}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
self.features.check_availability(self.options.version)
|
||||
|
||||
@ -12,7 +12,6 @@ to output a [`Module`](crate::Module) into glsl
|
||||
- 420
|
||||
- 430
|
||||
- 450
|
||||
- 460
|
||||
|
||||
### ES
|
||||
- 300
|
||||
@ -69,6 +68,10 @@ pub const SUPPORTED_CORE_VERSIONS: &[u16] = &[330, 400, 410, 420, 430, 440, 450]
|
||||
/// List of supported `es` GLSL versions.
|
||||
pub const SUPPORTED_ES_VERSIONS: &[u16] = &[300, 310, 320];
|
||||
|
||||
/// The suffix of the variable that will hold the calculated clamped level
|
||||
/// of detail for bounds checking in `ImageLoad`
|
||||
const CLAMPED_LOD_SUFFIX: &str = "_clamped_lod";
|
||||
|
||||
/// Mapping between resources and bindings.
|
||||
pub type BindingMap = std::collections::BTreeMap<crate::ResourceBinding, u8>;
|
||||
|
||||
@ -375,6 +378,8 @@ pub struct Writer<'a, W> {
|
||||
out: W,
|
||||
/// User defined configuration to be used.
|
||||
options: &'a Options,
|
||||
/// The bound checking policies to be used
|
||||
policies: proc::BoundsCheckPolicies,
|
||||
|
||||
// Internal State
|
||||
/// Features manager used to store all the needed features and write them.
|
||||
@ -410,6 +415,7 @@ impl<'a, W: Write> Writer<'a, W> {
|
||||
info: &'a valid::ModuleInfo,
|
||||
options: &'a Options,
|
||||
pipeline_options: &'a PipelineOptions,
|
||||
policies: proc::BoundsCheckPolicies,
|
||||
) -> Result<Self, Error> {
|
||||
// Check if the requested version is supported
|
||||
if !options.version.is_supported() {
|
||||
@ -437,6 +443,8 @@ impl<'a, W: Write> Writer<'a, W> {
|
||||
info,
|
||||
out,
|
||||
options,
|
||||
policies,
|
||||
|
||||
namer,
|
||||
features: FeaturesManager::new(),
|
||||
names,
|
||||
@ -1635,6 +1643,27 @@ impl<'a, W: Write> Writer<'a, W> {
|
||||
None
|
||||
};
|
||||
|
||||
// If we are going to write an `ImageLoad` next and the target image
|
||||
// is sampled and we are using the `Restrict` policy for bounds
|
||||
// checking images we need to write a local holding the clamped lod.
|
||||
if let crate::Expression::ImageLoad {
|
||||
image,
|
||||
level: Some(level_expr),
|
||||
..
|
||||
} = ctx.expressions[handle]
|
||||
{
|
||||
if let TypeInner::Image {
|
||||
class: crate::ImageClass::Sampled { .. },
|
||||
..
|
||||
} = *ctx.info[image].ty.inner_with(&self.module.types)
|
||||
{
|
||||
if let proc::BoundsCheckPolicy::Restrict = self.policies.image {
|
||||
write!(self.out, "{}", level)?;
|
||||
self.write_clamped_lod(ctx, handle, image, level_expr)?
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
if let Some(name) = expr_name {
|
||||
write!(self.out, "{}", level)?;
|
||||
self.write_named_expr(handle, name, ctx)?;
|
||||
@ -1933,19 +1962,7 @@ impl<'a, W: Write> Writer<'a, W> {
|
||||
value,
|
||||
} => {
|
||||
write!(self.out, "{}", level)?;
|
||||
// This will only panic if the module is invalid
|
||||
let dim = match *ctx.info[image].ty.inner_with(&self.module.types) {
|
||||
TypeInner::Image { dim, .. } => dim,
|
||||
_ => unreachable!(),
|
||||
};
|
||||
|
||||
write!(self.out, "imageStore(")?;
|
||||
self.write_expr(image, ctx)?;
|
||||
write!(self.out, ", ")?;
|
||||
self.write_texture_coordinates(coordinate, array_index, dim, ctx)?;
|
||||
write!(self.out, ", ")?;
|
||||
self.write_expr(value, ctx)?;
|
||||
writeln!(self.out, ");")?;
|
||||
self.write_image_store(ctx, image, coordinate, array_index, value)?
|
||||
}
|
||||
// A `Call` is written `name(arguments)` where `arguments` is a comma separated expressions list
|
||||
Statement::Call {
|
||||
@ -2320,51 +2337,13 @@ impl<'a, W: Write> Writer<'a, W> {
|
||||
// End the function
|
||||
write!(self.out, ")")?
|
||||
}
|
||||
// `ImageLoad` is also a bit complicated.
|
||||
// There are two functions one for sampled
|
||||
// images another for storage images, the former uses `texelFetch` and the latter uses
|
||||
// `imageLoad`.
|
||||
// Furthermore we have `index` which is always `Some` for sampled images
|
||||
// and `None` for storage images, so we end up with two functions:
|
||||
// `texelFetch(image, coordinate, index)` - for sampled images
|
||||
// `imageLoad(image, coordinate)` - for storage images
|
||||
Expression::ImageLoad {
|
||||
image,
|
||||
coordinate,
|
||||
array_index,
|
||||
sample,
|
||||
level,
|
||||
} => {
|
||||
// This will only panic if the module is invalid
|
||||
let (dim, class) = match *ctx.info[image].ty.inner_with(&self.module.types) {
|
||||
TypeInner::Image {
|
||||
dim,
|
||||
arrayed: _,
|
||||
class,
|
||||
} => (dim, class),
|
||||
_ => unreachable!(),
|
||||
};
|
||||
|
||||
let fun_name = match class {
|
||||
crate::ImageClass::Sampled { .. } => "texelFetch",
|
||||
crate::ImageClass::Storage { .. } => "imageLoad",
|
||||
// TODO: Is there even a function for this?
|
||||
crate::ImageClass::Depth { multi: _ } => {
|
||||
return Err(Error::Custom("TODO: depth sample loads".to_string()))
|
||||
}
|
||||
};
|
||||
|
||||
write!(self.out, "{}(", fun_name)?;
|
||||
self.write_expr(image, ctx)?;
|
||||
write!(self.out, ", ")?;
|
||||
self.write_texture_coordinates(coordinate, array_index, dim, ctx)?;
|
||||
|
||||
if let Some(sample_or_level) = sample.or(level) {
|
||||
write!(self.out, ", ")?;
|
||||
self.write_expr(sample_or_level, ctx)?;
|
||||
}
|
||||
write!(self.out, ")")?;
|
||||
}
|
||||
} => self.write_image_load(expr, ctx, image, coordinate, array_index, sample, level)?,
|
||||
// Query translates into one of the:
|
||||
// - textureSize/imageSize
|
||||
// - textureQueryLevels
|
||||
@ -2961,25 +2940,71 @@ impl<'a, W: Write> Writer<'a, W> {
|
||||
Ok(())
|
||||
}
|
||||
|
||||
fn write_texture_coordinates(
|
||||
/// Helper function to write the local holding the clamped lod
|
||||
fn write_clamped_lod(
|
||||
&mut self,
|
||||
ctx: &back::FunctionCtx,
|
||||
expr: Handle<crate::Expression>,
|
||||
image: Handle<crate::Expression>,
|
||||
level_expr: Handle<crate::Expression>,
|
||||
) -> Result<(), Error> {
|
||||
// Define our local and start a call to `clamp`
|
||||
write!(
|
||||
self.out,
|
||||
"int {}{}{} = clamp(",
|
||||
back::BAKE_PREFIX,
|
||||
expr.index(),
|
||||
CLAMPED_LOD_SUFFIX
|
||||
)?;
|
||||
// Write the lod that will be clamped
|
||||
self.write_expr(level_expr, ctx)?;
|
||||
// Set the min value to 0 and start a call to `textureQueryLevels` to get
|
||||
// the maximum value
|
||||
write!(self.out, ", 0, textureQueryLevels(")?;
|
||||
// Write the target image as an argument to `textureQueryLevels`
|
||||
self.write_expr(image, ctx)?;
|
||||
// Close the call to `textureQueryLevels` subtract 1 from it since
|
||||
// the lod argument is 0 based, close the `clamp` call and end the
|
||||
// local declaration statement.
|
||||
writeln!(self.out, ") - 1);")?;
|
||||
|
||||
Ok(())
|
||||
}
|
||||
|
||||
// Helper method used to retrieve how many elements a coordinate vector
|
||||
// for the images operations need.
|
||||
fn get_coordinate_vector_size(&self, dim: crate::ImageDimension, arrayed: bool) -> u8 {
|
||||
// openGL es doesn't have 1D images so we need workaround it
|
||||
let tex_1d_hack = dim == crate::ImageDimension::D1 && self.options.version.is_es();
|
||||
// Get how many components the coordinate vector needs for the dimensions only
|
||||
let tex_coord_size = match dim {
|
||||
crate::ImageDimension::D1 => 1,
|
||||
crate::ImageDimension::D2 => 2,
|
||||
crate::ImageDimension::D3 => 3,
|
||||
crate::ImageDimension::Cube => 2,
|
||||
};
|
||||
// Calculate the true size of the coordinate vector by adding 1 for arrayed images
|
||||
// and another 1 if we need to workaround 1D images by making them 2D
|
||||
tex_coord_size + tex_1d_hack as u8 + arrayed as u8
|
||||
}
|
||||
|
||||
/// Helper method to write the coordinate vector for image operations
|
||||
fn write_texture_coord(
|
||||
&mut self,
|
||||
ctx: &back::FunctionCtx,
|
||||
vector_size: u8,
|
||||
coordinate: Handle<crate::Expression>,
|
||||
array_index: Option<Handle<crate::Expression>>,
|
||||
dim: crate::ImageDimension,
|
||||
ctx: &back::FunctionCtx,
|
||||
// Emulate 1D images as 2D for profiles that don't support it (glsl es)
|
||||
tex_1d_hack: bool,
|
||||
) -> Result<(), Error> {
|
||||
use crate::ImageDimension as IDim;
|
||||
|
||||
let tex_1d_hack = dim == IDim::D1 && self.options.version.is_es();
|
||||
match array_index {
|
||||
// If the image needs an array indice we need to add it to the end of our
|
||||
// coordinate vector, to do so we will use the `ivec(ivec, scalar)`
|
||||
// constructor notation (NOTE: the inner `ivec` can also be a scalar, this
|
||||
// is important for 1D arrayed images).
|
||||
Some(layer_expr) => {
|
||||
let tex_coord_size = match dim {
|
||||
IDim::D1 => 2,
|
||||
IDim::D2 => 3,
|
||||
IDim::D3 => 4,
|
||||
IDim::Cube => 4,
|
||||
};
|
||||
write!(self.out, "ivec{}(", tex_coord_size + tex_1d_hack as u8)?;
|
||||
write!(self.out, "ivec{}(", vector_size)?;
|
||||
self.write_expr(coordinate, ctx)?;
|
||||
write!(self.out, ", ")?;
|
||||
// If we are replacing sampler1D with sampler2D we also need
|
||||
@ -2990,16 +3015,326 @@ impl<'a, W: Write> Writer<'a, W> {
|
||||
self.write_expr(layer_expr, ctx)?;
|
||||
write!(self.out, ")")?;
|
||||
}
|
||||
// Otherwise write just the expression (and the 1D hack if needed)
|
||||
None => {
|
||||
if tex_1d_hack {
|
||||
write!(self.out, "ivec2(")?;
|
||||
}
|
||||
self.write_expr(coordinate, ctx)?;
|
||||
if tex_1d_hack {
|
||||
write!(self.out, ", 0.0)")?;
|
||||
write!(self.out, ", 0)")?;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
Ok(())
|
||||
}
|
||||
|
||||
/// Helper method to write the `ImageStore` statement
|
||||
fn write_image_store(
|
||||
&mut self,
|
||||
ctx: &back::FunctionCtx,
|
||||
image: Handle<crate::Expression>,
|
||||
coordinate: Handle<crate::Expression>,
|
||||
array_index: Option<Handle<crate::Expression>>,
|
||||
value: Handle<crate::Expression>,
|
||||
) -> Result<(), Error> {
|
||||
use crate::ImageDimension as IDim;
|
||||
|
||||
// NOTE: openGL requires that `imageStore`s have no effets when the texel is invalid
|
||||
// so we don't need to generate bounds checks (OpenGL 4.2 Core §3.9.20)
|
||||
|
||||
// This will only panic if the module is invalid
|
||||
let dim = match *ctx.info[image].ty.inner_with(&self.module.types) {
|
||||
TypeInner::Image { dim, .. } => dim,
|
||||
_ => unreachable!(),
|
||||
};
|
||||
|
||||
// Begin our call to `imageStore`
|
||||
write!(self.out, "imageStore(")?;
|
||||
self.write_expr(image, ctx)?;
|
||||
// Separate the image argument from the coordinates
|
||||
write!(self.out, ", ")?;
|
||||
|
||||
// openGL es doesn't have 1D images so we need workaround it
|
||||
let tex_1d_hack = dim == IDim::D1 && self.options.version.is_es();
|
||||
// Write the coordinate vector
|
||||
self.write_texture_coord(
|
||||
ctx,
|
||||
// Get the size of the coordinate vector
|
||||
self.get_coordinate_vector_size(dim, array_index.is_some()),
|
||||
coordinate,
|
||||
array_index,
|
||||
tex_1d_hack,
|
||||
)?;
|
||||
|
||||
// Separate the coordinate from the value to write and write the expression
|
||||
// of the value to write.
|
||||
write!(self.out, ", ")?;
|
||||
self.write_expr(value, ctx)?;
|
||||
// End the call to `imageStore` and the statement.
|
||||
writeln!(self.out, ");")?;
|
||||
|
||||
Ok(())
|
||||
}
|
||||
|
||||
/// Helper method for writing an `ImageLoad` expression.
|
||||
#[allow(clippy::too_many_arguments)]
|
||||
fn write_image_load(
|
||||
&mut self,
|
||||
handle: Handle<crate::Expression>,
|
||||
ctx: &back::FunctionCtx,
|
||||
image: Handle<crate::Expression>,
|
||||
coordinate: Handle<crate::Expression>,
|
||||
array_index: Option<Handle<crate::Expression>>,
|
||||
sample: Option<Handle<crate::Expression>>,
|
||||
level: Option<Handle<crate::Expression>>,
|
||||
) -> Result<(), Error> {
|
||||
use crate::ImageDimension as IDim;
|
||||
|
||||
// `ImageLoad` is a bit complicated.
|
||||
// There are two functions one for sampled
|
||||
// images another for storage images, the former uses `texelFetch` and the
|
||||
// latter uses `imageLoad`.
|
||||
//
|
||||
// Furthermore we have `level` which is always `Some` for sampled images
|
||||
// and `None` for storage images, so we end up with two functions:
|
||||
// - `texelFetch(image, coordinate, level)` for sampled images
|
||||
// - `imageLoad(image, coordinate)` for storage images
|
||||
//
|
||||
// Finally we also have to consider bounds checking, for storage images
|
||||
// this is easy since openGL requires that invalid texels always return
|
||||
// 0, for sampled images we need to either verify that all arguments are
|
||||
// in bounds (`ReadZeroSkipWrite`) or make them a valid texel (`Restrict`).
|
||||
|
||||
// This will only panic if the module is invalid
|
||||
let (dim, class) = match *ctx.info[image].ty.inner_with(&self.module.types) {
|
||||
TypeInner::Image {
|
||||
dim,
|
||||
arrayed: _,
|
||||
class,
|
||||
} => (dim, class),
|
||||
_ => unreachable!(),
|
||||
};
|
||||
|
||||
// Get the name of the function to be used for the load operation
|
||||
// and the policy to be used with it.
|
||||
let (fun_name, policy) = match class {
|
||||
// Sampled images inherit the policy from the user passed policies
|
||||
crate::ImageClass::Sampled { .. } => ("texelFetch", self.policies.image),
|
||||
crate::ImageClass::Storage { .. } => {
|
||||
// OpenGL 4.2 Core §3.9.20 defines that out of bounds texels in `imageLoad`s
|
||||
// always return zero values so we don't need to generate bounds checks
|
||||
("imageLoad", proc::BoundsCheckPolicy::Unchecked)
|
||||
}
|
||||
// TODO: Is there even a function for this?
|
||||
crate::ImageClass::Depth { multi: _ } => {
|
||||
return Err(Error::Custom(
|
||||
"WGSL `textureLoad` from depth textures is not supported in GLSL".to_string(),
|
||||
))
|
||||
}
|
||||
};
|
||||
|
||||
// openGL es doesn't have 1D images so we need workaround it
|
||||
let tex_1d_hack = dim == IDim::D1 && self.options.version.is_es();
|
||||
// Get the size of the coordinate vector
|
||||
let vector_size = self.get_coordinate_vector_size(dim, array_index.is_some());
|
||||
|
||||
if let proc::BoundsCheckPolicy::ReadZeroSkipWrite = policy {
|
||||
// To write the bounds checks for `ReadZeroSkipWrite` we will use a
|
||||
// ternary operator since we are in the middle of an expression and
|
||||
// need to return a value.
|
||||
//
|
||||
// NOTE: glsl does short circuit when evaluating logical
|
||||
// expressions so we can be sure that after we test a
|
||||
// condition it will be true for the next ones
|
||||
|
||||
// Write parantheses around the ternary operator to prevent problems with
|
||||
// expressions emitted before or after it having more precedence
|
||||
write!(self.out, "(",)?;
|
||||
|
||||
// The lod check needs to precede the size check since we need
|
||||
// to use the lod to get the size of the image at that level.
|
||||
if let Some(level_expr) = level {
|
||||
self.write_expr(level_expr, ctx)?;
|
||||
write!(self.out, " < textureQueryLevels(",)?;
|
||||
self.write_expr(image, ctx)?;
|
||||
// Chain the next check
|
||||
write!(self.out, ") && ")?;
|
||||
}
|
||||
|
||||
// Check that the sample arguments doesn't exceed the number of samples
|
||||
if let Some(sample_expr) = sample {
|
||||
self.write_expr(sample_expr, ctx)?;
|
||||
write!(self.out, " < textureSamples(",)?;
|
||||
self.write_expr(image, ctx)?;
|
||||
// Chain the next check
|
||||
write!(self.out, ") && ")?;
|
||||
}
|
||||
|
||||
// We now need to write the size checks for the coordinates and array index
|
||||
// first we write the comparation function in case the image is 1D non arrayed
|
||||
// (and no 1D to 2D hack was needed) we are comparing scalars so the less than
|
||||
// operator will suffice, but otherwise we'll be comparing two vectors so we'll
|
||||
// need to use the `lessThan` function but it returns a vector of booleans (one
|
||||
// for each comparison) so we need to fold it all in one scalar boolean, since
|
||||
// we want all comparisons to pass we use the `all` function which will only
|
||||
// return `true` if all the elements of the boolean vector are also `true`.
|
||||
//
|
||||
// So we'll end with one of the following forms
|
||||
// - `coord < textureSize(image, lod)` for 1D images
|
||||
// - `all(lessThan(coord, textureSize(image, lod)))` for normal images
|
||||
// - `all(lessThan(ivec(coord, array_index), textureSize(image, lod)))`
|
||||
// for arrayed images
|
||||
// - `all(lessThan(coord, textureSize(image)))` for multi sampled images
|
||||
|
||||
if vector_size != 1 {
|
||||
write!(self.out, "all(lessThan(")?;
|
||||
}
|
||||
|
||||
// Write the coordinate vector
|
||||
self.write_texture_coord(ctx, vector_size, coordinate, array_index, tex_1d_hack)?;
|
||||
|
||||
if vector_size != 1 {
|
||||
// If we used the `lessThan` function we need to separate the
|
||||
// coordinates from the image size.
|
||||
write!(self.out, ", ")?;
|
||||
} else {
|
||||
// If we didn't use it (ie. 1D images) we perform the comparsion
|
||||
// using the less than operator.
|
||||
write!(self.out, " < ")?;
|
||||
}
|
||||
|
||||
// Call `textureSize` to get our image size
|
||||
write!(self.out, "textureSize(")?;
|
||||
self.write_expr(image, ctx)?;
|
||||
// `textureSize` uses the lod as a second argument for mipmapped images
|
||||
if let Some(level_expr) = level {
|
||||
// Separate the image from the lod
|
||||
write!(self.out, ", ")?;
|
||||
self.write_expr(level_expr, ctx)?;
|
||||
}
|
||||
// Close the `textureSize` call
|
||||
write!(self.out, ")")?;
|
||||
|
||||
if vector_size != 1 {
|
||||
// Close the `all` and `lessThan` calls
|
||||
write!(self.out, "))")?;
|
||||
}
|
||||
|
||||
// Finally end the condition part of the ternary operator
|
||||
write!(self.out, " ? ")?;
|
||||
}
|
||||
|
||||
// Begin the call to the function used to load the texel
|
||||
write!(self.out, "{}(", fun_name)?;
|
||||
self.write_expr(image, ctx)?;
|
||||
write!(self.out, ", ")?;
|
||||
|
||||
// If we are using `Restrict` bounds checking we need to pass valid texel
|
||||
// coordinates, to do so we use the `clamp` function to get a value between
|
||||
// 0 and the image size - 1 (indexing begins at 0)
|
||||
if let proc::BoundsCheckPolicy::Restrict = policy {
|
||||
write!(self.out, "clamp(")?;
|
||||
}
|
||||
|
||||
// Write the coordinate vector
|
||||
self.write_texture_coord(ctx, vector_size, coordinate, array_index, tex_1d_hack)?;
|
||||
|
||||
// If we are using `Restrict` bounds checking we need to write the rest of the
|
||||
// clamp we initiated before writing the coordinates.
|
||||
if let proc::BoundsCheckPolicy::Restrict = policy {
|
||||
// Write the min value 0
|
||||
if vector_size == 1 {
|
||||
write!(self.out, ", 0")?;
|
||||
} else {
|
||||
write!(self.out, ", ivec{}(0)", vector_size)?;
|
||||
}
|
||||
// Start the `textureSize` call to use as the max value.
|
||||
write!(self.out, ", textureSize(")?;
|
||||
self.write_expr(image, ctx)?;
|
||||
// If the image is mipmapped we need to add the lod argument to the
|
||||
// `textureSize` call, but this needs to be the clamped lod, this should
|
||||
// have been generated earlier and put in a local.
|
||||
if class.is_mipmapped() {
|
||||
write!(
|
||||
self.out,
|
||||
", {}{}{}",
|
||||
back::BAKE_PREFIX,
|
||||
handle.index(),
|
||||
CLAMPED_LOD_SUFFIX
|
||||
)?;
|
||||
}
|
||||
// Close the `textureSize` call
|
||||
write!(self.out, ")")?;
|
||||
|
||||
// Subtract 1 from the `textureSize` call since the coordinates are zero based.
|
||||
if vector_size == 1 {
|
||||
write!(self.out, " - 1")?;
|
||||
} else {
|
||||
write!(self.out, " - ivec{}(1)", vector_size)?;
|
||||
}
|
||||
|
||||
// Close the `clamp` call
|
||||
write!(self.out, ")")?;
|
||||
|
||||
// Add the clamped lod (if present) as the second argument to the
|
||||
// image load function.
|
||||
if level.is_some() {
|
||||
write!(
|
||||
self.out,
|
||||
", {}{}{}",
|
||||
back::BAKE_PREFIX,
|
||||
handle.index(),
|
||||
CLAMPED_LOD_SUFFIX
|
||||
)?;
|
||||
}
|
||||
|
||||
// If a sample argument is needed we need to clamp it between 0 and
|
||||
// the number of samples the image has.
|
||||
if let Some(sample_expr) = sample {
|
||||
write!(self.out, ", clamp(")?;
|
||||
self.write_expr(sample_expr, ctx)?;
|
||||
// Set the min value to 0 and start the call to `textureSamples`
|
||||
write!(self.out, ", 0, textureSamples(")?;
|
||||
self.write_expr(image, ctx)?;
|
||||
// Close the `textureSamples` call, subtract 1 from it since the sample
|
||||
// argument is zero based, and close the `clamp` call
|
||||
writeln!(self.out, ") - 1)")?;
|
||||
}
|
||||
} else if let Some(sample_or_level) = sample.or(level) {
|
||||
// If no bounds checking is need just add the sample or level argument
|
||||
// after the coordinates
|
||||
write!(self.out, ", ")?;
|
||||
self.write_expr(sample_or_level, ctx)?;
|
||||
}
|
||||
|
||||
// Close the image load function.
|
||||
write!(self.out, ")")?;
|
||||
|
||||
// If we were using the `ReadZeroSkipWrite` policy we need to end the first branch
|
||||
// (which is taken if the condition is `true`) with a colon (`:`) and write the
|
||||
// second branch which is just a 0 value.
|
||||
if let proc::BoundsCheckPolicy::ReadZeroSkipWrite = policy {
|
||||
// Get the kind of the output value.
|
||||
let kind = match class {
|
||||
// Only sampled images can reach here since storage images
|
||||
// don't need bounds checks and depth images aren't implmented
|
||||
crate::ImageClass::Sampled { kind, .. } => kind,
|
||||
_ => unreachable!(),
|
||||
};
|
||||
|
||||
// End the first branch
|
||||
write!(self.out, " : ")?;
|
||||
// Write the 0 value
|
||||
write!(self.out, "{}vec4(", glsl_scalar(kind, 4)?.prefix,)?;
|
||||
self.write_zero_init_scalar(kind)?;
|
||||
// Close the zero value constructor
|
||||
write!(self.out, ")")?;
|
||||
// Close the parantheses surrounding our ternary
|
||||
write!(self.out, ")")?;
|
||||
}
|
||||
|
||||
Ok(())
|
||||
}
|
||||
|
||||
|
||||
@ -6,4 +6,9 @@
|
||||
version: (1, 1),
|
||||
debug: true,
|
||||
),
|
||||
glsl: (
|
||||
version: Desktop(430),
|
||||
writer_flags: (bits: 0),
|
||||
binding_map: { },
|
||||
),
|
||||
)
|
||||
|
||||
@ -5,79 +5,100 @@ fn test_textureLoad_1d(coords: i32, level: i32) -> vec4<f32> {
|
||||
return textureLoad(image_1d, coords, level);
|
||||
}
|
||||
|
||||
@group(0) @binding(0)
|
||||
@group(0) @binding(1)
|
||||
var image_2d: texture_2d<f32>;
|
||||
|
||||
fn test_textureLoad_2d(coords: vec2<i32>, level: i32) -> vec4<f32> {
|
||||
return textureLoad(image_2d, coords, level);
|
||||
}
|
||||
|
||||
@group(0) @binding(0)
|
||||
@group(0) @binding(2)
|
||||
var image_2d_array: texture_2d_array<f32>;
|
||||
|
||||
fn test_textureLoad_2d_array(coords: vec2<i32>, index: i32, level: i32) -> vec4<f32> {
|
||||
return textureLoad(image_2d_array, coords, index, level);
|
||||
}
|
||||
|
||||
@group(0) @binding(0)
|
||||
@group(0) @binding(3)
|
||||
var image_3d: texture_3d<f32>;
|
||||
|
||||
fn test_textureLoad_3d(coords: vec3<i32>, level: i32) -> vec4<f32> {
|
||||
return textureLoad(image_3d, coords, level);
|
||||
}
|
||||
|
||||
@group(0) @binding(0)
|
||||
@group(0) @binding(4)
|
||||
var image_multisampled_2d: texture_multisampled_2d<f32>;
|
||||
|
||||
fn test_textureLoad_multisampled_2d(coords: vec2<i32>, _sample: i32) -> vec4<f32> {
|
||||
return textureLoad(image_multisampled_2d, coords, _sample);
|
||||
}
|
||||
|
||||
@group(0) @binding(0)
|
||||
@group(0) @binding(5)
|
||||
var image_depth_2d: texture_depth_2d;
|
||||
|
||||
fn test_textureLoad_depth_2d(coords: vec2<i32>, level: i32) -> f32 {
|
||||
return textureLoad(image_depth_2d, coords, level);
|
||||
}
|
||||
|
||||
@group(0) @binding(0)
|
||||
@group(0) @binding(6)
|
||||
var image_depth_2d_array: texture_depth_2d_array;
|
||||
|
||||
fn test_textureLoad_depth_2d_array(coords: vec2<i32>, index: i32, level: i32) -> f32 {
|
||||
return textureLoad(image_depth_2d_array, coords, index, level);
|
||||
}
|
||||
|
||||
@group(0) @binding(0)
|
||||
@group(0) @binding(7)
|
||||
var image_depth_multisampled_2d: texture_depth_multisampled_2d;
|
||||
|
||||
fn test_textureLoad_depth_multisampled_2d(coords: vec2<i32>, _sample: i32) -> f32 {
|
||||
return textureLoad(image_depth_multisampled_2d, coords, _sample);
|
||||
}
|
||||
|
||||
@group(0) @binding(0)
|
||||
@group(0) @binding(8)
|
||||
var image_storage_1d: texture_storage_1d<rgba8unorm, write>;
|
||||
|
||||
fn test_textureStore_1d(coords: i32, value: vec4<f32>) {
|
||||
textureStore(image_storage_1d, coords, value);
|
||||
}
|
||||
|
||||
@group(0) @binding(0)
|
||||
@group(0) @binding(9)
|
||||
var image_storage_2d: texture_storage_2d<rgba8unorm, write>;
|
||||
|
||||
fn test_textureStore_2d(coords: vec2<i32>, value: vec4<f32>) {
|
||||
textureStore(image_storage_2d, coords, value);
|
||||
}
|
||||
|
||||
@group(0) @binding(0)
|
||||
@group(0) @binding(10)
|
||||
var image_storage_2d_array: texture_storage_2d_array<rgba8unorm, write>;
|
||||
|
||||
fn test_textureStore_2d_array(coords: vec2<i32>, array_index: i32, value: vec4<f32>) {
|
||||
textureStore(image_storage_2d_array, coords, array_index, value);
|
||||
}
|
||||
|
||||
@group(0) @binding(0)
|
||||
@group(0) @binding(11)
|
||||
var image_storage_3d: texture_storage_3d<rgba8unorm, write>;
|
||||
|
||||
fn test_textureStore_3d(coords: vec3<i32>, value: vec4<f32>) {
|
||||
textureStore(image_storage_3d, coords, value);
|
||||
}
|
||||
|
||||
// GLSL output requires that we identify an entry point, so
|
||||
// that it can tell what "in" and "out" globals to write.
|
||||
@fragment
|
||||
fn fragment_shader() -> @location(0) vec4<f32> {
|
||||
test_textureLoad_1d(0, 0);
|
||||
test_textureLoad_2d(vec2<i32>(), 0);
|
||||
test_textureLoad_2d_array(vec2<i32>(), 0, 0);
|
||||
test_textureLoad_3d(vec3<i32>(), 0);
|
||||
test_textureLoad_multisampled_2d(vec2<i32>(), 0);
|
||||
// Not yet implemented for GLSL:
|
||||
// test_textureLoad_depth_2d(vec2<i32>(), 0);
|
||||
// test_textureLoad_depth_2d_array(vec2<i32>(), 0, 0);
|
||||
// test_textureLoad_depth_multisampled_2d(vec2<i32>(), 0);
|
||||
test_textureStore_1d(0, vec4<f32>());
|
||||
test_textureStore_2d(vec2<i32>(), vec4<f32>());
|
||||
test_textureStore_2d_array(vec2<i32>(), 0, vec4<f32>());
|
||||
test_textureStore_3d(vec3<i32>(), vec4<f32>());
|
||||
|
||||
return vec4<f32>(0.,0.,0.,0.);
|
||||
}
|
||||
|
||||
@ -6,4 +6,9 @@
|
||||
version: (1, 1),
|
||||
debug: true,
|
||||
),
|
||||
glsl: (
|
||||
version: Desktop(430),
|
||||
writer_flags: (bits: 0),
|
||||
binding_map: { },
|
||||
),
|
||||
)
|
||||
|
||||
@ -5,79 +5,100 @@ fn test_textureLoad_1d(coords: i32, level: i32) -> vec4<f32> {
|
||||
return textureLoad(image_1d, coords, level);
|
||||
}
|
||||
|
||||
@group(0) @binding(0)
|
||||
@group(0) @binding(1)
|
||||
var image_2d: texture_2d<f32>;
|
||||
|
||||
fn test_textureLoad_2d(coords: vec2<i32>, level: i32) -> vec4<f32> {
|
||||
return textureLoad(image_2d, coords, level);
|
||||
}
|
||||
|
||||
@group(0) @binding(0)
|
||||
@group(0) @binding(2)
|
||||
var image_2d_array: texture_2d_array<f32>;
|
||||
|
||||
fn test_textureLoad_2d_array(coords: vec2<i32>, index: i32, level: i32) -> vec4<f32> {
|
||||
return textureLoad(image_2d_array, coords, index, level);
|
||||
}
|
||||
|
||||
@group(0) @binding(0)
|
||||
@group(0) @binding(3)
|
||||
var image_3d: texture_3d<f32>;
|
||||
|
||||
fn test_textureLoad_3d(coords: vec3<i32>, level: i32) -> vec4<f32> {
|
||||
return textureLoad(image_3d, coords, level);
|
||||
}
|
||||
|
||||
@group(0) @binding(0)
|
||||
@group(0) @binding(4)
|
||||
var image_multisampled_2d: texture_multisampled_2d<f32>;
|
||||
|
||||
fn test_textureLoad_multisampled_2d(coords: vec2<i32>, _sample: i32) -> vec4<f32> {
|
||||
return textureLoad(image_multisampled_2d, coords, _sample);
|
||||
}
|
||||
|
||||
@group(0) @binding(0)
|
||||
@group(0) @binding(5)
|
||||
var image_depth_2d: texture_depth_2d;
|
||||
|
||||
fn test_textureLoad_depth_2d(coords: vec2<i32>, level: i32) -> f32 {
|
||||
return textureLoad(image_depth_2d, coords, level);
|
||||
}
|
||||
|
||||
@group(0) @binding(0)
|
||||
@group(0) @binding(6)
|
||||
var image_depth_2d_array: texture_depth_2d_array;
|
||||
|
||||
fn test_textureLoad_depth_2d_array(coords: vec2<i32>, index: i32, level: i32) -> f32 {
|
||||
return textureLoad(image_depth_2d_array, coords, index, level);
|
||||
}
|
||||
|
||||
@group(0) @binding(0)
|
||||
@group(0) @binding(7)
|
||||
var image_depth_multisampled_2d: texture_depth_multisampled_2d;
|
||||
|
||||
fn test_textureLoad_depth_multisampled_2d(coords: vec2<i32>, _sample: i32) -> f32 {
|
||||
return textureLoad(image_depth_multisampled_2d, coords, _sample);
|
||||
}
|
||||
|
||||
@group(0) @binding(0)
|
||||
@group(0) @binding(8)
|
||||
var image_storage_1d: texture_storage_1d<rgba8unorm, write>;
|
||||
|
||||
fn test_textureStore_1d(coords: i32, value: vec4<f32>) {
|
||||
textureStore(image_storage_1d, coords, value);
|
||||
}
|
||||
|
||||
@group(0) @binding(0)
|
||||
@group(0) @binding(9)
|
||||
var image_storage_2d: texture_storage_2d<rgba8unorm, write>;
|
||||
|
||||
fn test_textureStore_2d(coords: vec2<i32>, value: vec4<f32>) {
|
||||
textureStore(image_storage_2d, coords, value);
|
||||
}
|
||||
|
||||
@group(0) @binding(0)
|
||||
@group(0) @binding(10)
|
||||
var image_storage_2d_array: texture_storage_2d_array<rgba8unorm, write>;
|
||||
|
||||
fn test_textureStore_2d_array(coords: vec2<i32>, array_index: i32, value: vec4<f32>) {
|
||||
textureStore(image_storage_2d_array, coords, array_index, value);
|
||||
}
|
||||
|
||||
@group(0) @binding(0)
|
||||
@group(0) @binding(11)
|
||||
var image_storage_3d: texture_storage_3d<rgba8unorm, write>;
|
||||
|
||||
fn test_textureStore_3d(coords: vec3<i32>, value: vec4<f32>) {
|
||||
textureStore(image_storage_3d, coords, value);
|
||||
}
|
||||
|
||||
// GLSL output requires that we identify an entry point, so
|
||||
// that it can tell what "in" and "out" globals to write.
|
||||
@fragment
|
||||
fn fragment_shader() -> @location(0) vec4<f32> {
|
||||
test_textureLoad_1d(0, 0);
|
||||
test_textureLoad_2d(vec2<i32>(), 0);
|
||||
test_textureLoad_2d_array(vec2<i32>(), 0, 0);
|
||||
test_textureLoad_3d(vec3<i32>(), 0);
|
||||
test_textureLoad_multisampled_2d(vec2<i32>(), 0);
|
||||
// Not yet implemented for GLSL:
|
||||
// test_textureLoad_depth_2d(vec2<i32>(), 0);
|
||||
// test_textureLoad_depth_2d_array(vec2<i32>(), 0, 0);
|
||||
// test_textureLoad_depth_multisampled_2d(vec2<i32>(), 0);
|
||||
test_textureStore_1d(0, vec4<f32>());
|
||||
test_textureStore_2d(vec2<i32>(), vec4<f32>());
|
||||
test_textureStore_2d_array(vec2<i32>(), 0, vec4<f32>());
|
||||
test_textureStore_3d(vec3<i32>(), vec4<f32>());
|
||||
|
||||
return vec4<f32>(0.,0.,0.,0.);
|
||||
}
|
||||
|
||||
@ -0,0 +1,86 @@
|
||||
#version 430 core
|
||||
#extension GL_ARB_shader_texture_image_samples : require
|
||||
uniform highp sampler1D _group_0_binding_0_fs;
|
||||
|
||||
uniform highp sampler2D _group_0_binding_1_fs;
|
||||
|
||||
uniform highp sampler2DArray _group_0_binding_2_fs;
|
||||
|
||||
uniform highp sampler3D _group_0_binding_3_fs;
|
||||
|
||||
uniform highp sampler2DMS _group_0_binding_4_fs;
|
||||
|
||||
layout(rgba8) writeonly uniform highp image1D _group_0_binding_8_fs;
|
||||
|
||||
layout(rgba8) writeonly uniform highp image2D _group_0_binding_9_fs;
|
||||
|
||||
layout(rgba8) writeonly uniform highp image2DArray _group_0_binding_10_fs;
|
||||
|
||||
layout(rgba8) writeonly uniform highp image3D _group_0_binding_11_fs;
|
||||
|
||||
layout(location = 0) out vec4 _fs2p_location0;
|
||||
|
||||
vec4 test_textureLoad_1d(int coords, int level) {
|
||||
int _e3_clamped_lod = clamp(level, 0, textureQueryLevels(_group_0_binding_0_fs) - 1);
|
||||
vec4 _e3 = texelFetch(_group_0_binding_0_fs, clamp(coords, 0, textureSize(_group_0_binding_0_fs, _e3_clamped_lod) - 1), _e3_clamped_lod);
|
||||
return _e3;
|
||||
}
|
||||
|
||||
vec4 test_textureLoad_2d(ivec2 coords_1, int level_1) {
|
||||
int _e4_clamped_lod = clamp(level_1, 0, textureQueryLevels(_group_0_binding_1_fs) - 1);
|
||||
vec4 _e4 = texelFetch(_group_0_binding_1_fs, clamp(coords_1, ivec2(0), textureSize(_group_0_binding_1_fs, _e4_clamped_lod) - ivec2(1)), _e4_clamped_lod);
|
||||
return _e4;
|
||||
}
|
||||
|
||||
vec4 test_textureLoad_2d_array(ivec2 coords_2, int index, int level_2) {
|
||||
int _e6_clamped_lod = clamp(level_2, 0, textureQueryLevels(_group_0_binding_2_fs) - 1);
|
||||
vec4 _e6 = texelFetch(_group_0_binding_2_fs, clamp(ivec3(coords_2, index), ivec3(0), textureSize(_group_0_binding_2_fs, _e6_clamped_lod) - ivec3(1)), _e6_clamped_lod);
|
||||
return _e6;
|
||||
}
|
||||
|
||||
vec4 test_textureLoad_3d(ivec3 coords_3, int level_3) {
|
||||
int _e6_clamped_lod = clamp(level_3, 0, textureQueryLevels(_group_0_binding_3_fs) - 1);
|
||||
vec4 _e6 = texelFetch(_group_0_binding_3_fs, clamp(coords_3, ivec3(0), textureSize(_group_0_binding_3_fs, _e6_clamped_lod) - ivec3(1)), _e6_clamped_lod);
|
||||
return _e6;
|
||||
}
|
||||
|
||||
vec4 test_textureLoad_multisampled_2d(ivec2 coords_4, int _sample) {
|
||||
vec4 _e7 = texelFetch(_group_0_binding_4_fs, clamp(coords_4, ivec2(0), textureSize(_group_0_binding_4_fs) - ivec2(1)), clamp(_sample, 0, textureSamples(_group_0_binding_4_fs) - 1)
|
||||
);
|
||||
return _e7;
|
||||
}
|
||||
|
||||
void test_textureStore_1d(int coords_8, vec4 value) {
|
||||
imageStore(_group_0_binding_8_fs, coords_8, value);
|
||||
return;
|
||||
}
|
||||
|
||||
void test_textureStore_2d(ivec2 coords_9, vec4 value_1) {
|
||||
imageStore(_group_0_binding_9_fs, coords_9, value_1);
|
||||
return;
|
||||
}
|
||||
|
||||
void test_textureStore_2d_array(ivec2 coords_10, int array_index, vec4 value_2) {
|
||||
imageStore(_group_0_binding_10_fs, ivec3(coords_10, array_index), value_2);
|
||||
return;
|
||||
}
|
||||
|
||||
void test_textureStore_3d(ivec3 coords_11, vec4 value_3) {
|
||||
imageStore(_group_0_binding_11_fs, coords_11, value_3);
|
||||
return;
|
||||
}
|
||||
|
||||
void main() {
|
||||
vec4 _e14 = test_textureLoad_1d(0, 0);
|
||||
vec4 _e17 = test_textureLoad_2d(ivec2(0, 0), 0);
|
||||
vec4 _e21 = test_textureLoad_2d_array(ivec2(0, 0), 0, 0);
|
||||
vec4 _e24 = test_textureLoad_3d(ivec3(0, 0, 0), 0);
|
||||
vec4 _e27 = test_textureLoad_multisampled_2d(ivec2(0, 0), 0);
|
||||
test_textureStore_1d(0, vec4(0.0, 0.0, 0.0, 0.0));
|
||||
test_textureStore_2d(ivec2(0, 0), vec4(0.0, 0.0, 0.0, 0.0));
|
||||
test_textureStore_2d_array(ivec2(0, 0), 0, vec4(0.0, 0.0, 0.0, 0.0));
|
||||
test_textureStore_3d(ivec3(0, 0, 0), vec4(0.0, 0.0, 0.0, 0.0));
|
||||
_fs2p_location0 = vec4(0.0, 0.0, 0.0, 0.0);
|
||||
return;
|
||||
}
|
||||
|
||||
@ -0,0 +1,81 @@
|
||||
#version 430 core
|
||||
#extension GL_ARB_shader_texture_image_samples : require
|
||||
uniform highp sampler1D _group_0_binding_0_fs;
|
||||
|
||||
uniform highp sampler2D _group_0_binding_1_fs;
|
||||
|
||||
uniform highp sampler2DArray _group_0_binding_2_fs;
|
||||
|
||||
uniform highp sampler3D _group_0_binding_3_fs;
|
||||
|
||||
uniform highp sampler2DMS _group_0_binding_4_fs;
|
||||
|
||||
layout(rgba8) writeonly uniform highp image1D _group_0_binding_8_fs;
|
||||
|
||||
layout(rgba8) writeonly uniform highp image2D _group_0_binding_9_fs;
|
||||
|
||||
layout(rgba8) writeonly uniform highp image2DArray _group_0_binding_10_fs;
|
||||
|
||||
layout(rgba8) writeonly uniform highp image3D _group_0_binding_11_fs;
|
||||
|
||||
layout(location = 0) out vec4 _fs2p_location0;
|
||||
|
||||
vec4 test_textureLoad_1d(int coords, int level) {
|
||||
vec4 _e3 = (level < textureQueryLevels(_group_0_binding_0_fs) && coords < textureSize(_group_0_binding_0_fs, level) ? texelFetch(_group_0_binding_0_fs, coords, level) : vec4(0.0));
|
||||
return _e3;
|
||||
}
|
||||
|
||||
vec4 test_textureLoad_2d(ivec2 coords_1, int level_1) {
|
||||
vec4 _e4 = (level_1 < textureQueryLevels(_group_0_binding_1_fs) && all(lessThan(coords_1, textureSize(_group_0_binding_1_fs, level_1))) ? texelFetch(_group_0_binding_1_fs, coords_1, level_1) : vec4(0.0));
|
||||
return _e4;
|
||||
}
|
||||
|
||||
vec4 test_textureLoad_2d_array(ivec2 coords_2, int index, int level_2) {
|
||||
vec4 _e6 = (level_2 < textureQueryLevels(_group_0_binding_2_fs) && all(lessThan(ivec3(coords_2, index), textureSize(_group_0_binding_2_fs, level_2))) ? texelFetch(_group_0_binding_2_fs, ivec3(coords_2, index), level_2) : vec4(0.0));
|
||||
return _e6;
|
||||
}
|
||||
|
||||
vec4 test_textureLoad_3d(ivec3 coords_3, int level_3) {
|
||||
vec4 _e6 = (level_3 < textureQueryLevels(_group_0_binding_3_fs) && all(lessThan(coords_3, textureSize(_group_0_binding_3_fs, level_3))) ? texelFetch(_group_0_binding_3_fs, coords_3, level_3) : vec4(0.0));
|
||||
return _e6;
|
||||
}
|
||||
|
||||
vec4 test_textureLoad_multisampled_2d(ivec2 coords_4, int _sample) {
|
||||
vec4 _e7 = (_sample < textureSamples(_group_0_binding_4_fs) && all(lessThan(coords_4, textureSize(_group_0_binding_4_fs))) ? texelFetch(_group_0_binding_4_fs, coords_4, _sample) : vec4(0.0));
|
||||
return _e7;
|
||||
}
|
||||
|
||||
void test_textureStore_1d(int coords_8, vec4 value) {
|
||||
imageStore(_group_0_binding_8_fs, coords_8, value);
|
||||
return;
|
||||
}
|
||||
|
||||
void test_textureStore_2d(ivec2 coords_9, vec4 value_1) {
|
||||
imageStore(_group_0_binding_9_fs, coords_9, value_1);
|
||||
return;
|
||||
}
|
||||
|
||||
void test_textureStore_2d_array(ivec2 coords_10, int array_index, vec4 value_2) {
|
||||
imageStore(_group_0_binding_10_fs, ivec3(coords_10, array_index), value_2);
|
||||
return;
|
||||
}
|
||||
|
||||
void test_textureStore_3d(ivec3 coords_11, vec4 value_3) {
|
||||
imageStore(_group_0_binding_11_fs, coords_11, value_3);
|
||||
return;
|
||||
}
|
||||
|
||||
void main() {
|
||||
vec4 _e14 = test_textureLoad_1d(0, 0);
|
||||
vec4 _e17 = test_textureLoad_2d(ivec2(0, 0), 0);
|
||||
vec4 _e21 = test_textureLoad_2d_array(ivec2(0, 0), 0, 0);
|
||||
vec4 _e24 = test_textureLoad_3d(ivec3(0, 0, 0), 0);
|
||||
vec4 _e27 = test_textureLoad_multisampled_2d(ivec2(0, 0), 0);
|
||||
test_textureStore_1d(0, vec4(0.0, 0.0, 0.0, 0.0));
|
||||
test_textureStore_2d(ivec2(0, 0), vec4(0.0, 0.0, 0.0, 0.0));
|
||||
test_textureStore_2d_array(ivec2(0, 0), 0, vec4(0.0, 0.0, 0.0, 0.0));
|
||||
test_textureStore_3d(ivec3(0, 0, 0), vec4(0.0, 0.0, 0.0, 0.0));
|
||||
_fs2p_location0 = vec4(0.0, 0.0, 0.0, 0.0);
|
||||
return;
|
||||
}
|
||||
|
||||
@ -27,8 +27,8 @@ void main() {
|
||||
uvec4 value2_ = texelFetch(_group_0_binding_3_cs, itc, int(local_id.z));
|
||||
uvec4 value4_ = imageLoad(_group_0_binding_1_cs, itc);
|
||||
uvec4 value5_ = texelFetch(_group_0_binding_5_cs, ivec3(itc, int(local_id.z)), (int(local_id.z) + 1));
|
||||
uvec4 value6_ = texelFetch(_group_0_binding_7_cs, ivec2(int(local_id.x), 0.0), int(local_id.z));
|
||||
imageStore(_group_0_binding_2_cs, ivec2(itc.x, 0.0), ((((value1_ + value2_) + value4_) + value5_) + value6_));
|
||||
uvec4 value6_ = texelFetch(_group_0_binding_7_cs, ivec2(int(local_id.x), 0), int(local_id.z));
|
||||
imageStore(_group_0_binding_2_cs, ivec2(itc.x, 0), ((((value1_ + value2_) + value4_) + value5_) + value6_));
|
||||
return;
|
||||
}
|
||||
|
||||
|
||||
@ -4,6 +4,9 @@
|
||||
|
||||
using metal::uint;
|
||||
|
||||
constant metal::int2 const_type_4_ = {0, 0};
|
||||
constant metal::int3 const_type_7_ = {0, 0, 0};
|
||||
constant metal::float4 const_type_2_ = {0.0, 0.0, 0.0, 0.0};
|
||||
|
||||
metal::float4 test_textureLoad_1d(
|
||||
int coords,
|
||||
@ -120,3 +123,29 @@ void test_textureStore_3d(
|
||||
image_storage_3d.write(value_3, metal::min(metal::uint3(coords_11), metal::uint3(image_storage_3d.get_width(), image_storage_3d.get_height(), image_storage_3d.get_depth()) - 1));
|
||||
return;
|
||||
}
|
||||
|
||||
struct fragment_shaderOutput {
|
||||
metal::float4 member [[color(0)]];
|
||||
};
|
||||
fragment fragment_shaderOutput fragment_shader(
|
||||
metal::texture1d<float, metal::access::sample> image_1d [[user(fake0)]]
|
||||
, metal::texture2d<float, metal::access::sample> image_2d [[user(fake0)]]
|
||||
, metal::texture2d_array<float, metal::access::sample> image_2d_array [[user(fake0)]]
|
||||
, metal::texture3d<float, metal::access::sample> image_3d [[user(fake0)]]
|
||||
, metal::texture2d_ms<float, metal::access::read> image_multisampled_2d [[user(fake0)]]
|
||||
, metal::texture1d<float, metal::access::write> image_storage_1d [[user(fake0)]]
|
||||
, metal::texture2d<float, metal::access::write> image_storage_2d [[user(fake0)]]
|
||||
, metal::texture2d_array<float, metal::access::write> image_storage_2d_array [[user(fake0)]]
|
||||
, metal::texture3d<float, metal::access::write> image_storage_3d [[user(fake0)]]
|
||||
) {
|
||||
metal::float4 _e14 = test_textureLoad_1d(0, 0, image_1d);
|
||||
metal::float4 _e17 = test_textureLoad_2d(const_type_4_, 0, image_2d);
|
||||
metal::float4 _e21 = test_textureLoad_2d_array(const_type_4_, 0, 0, image_2d_array);
|
||||
metal::float4 _e24 = test_textureLoad_3d(const_type_7_, 0, image_3d);
|
||||
metal::float4 _e27 = test_textureLoad_multisampled_2d(const_type_4_, 0, image_multisampled_2d);
|
||||
test_textureStore_1d(0, const_type_2_, image_storage_1d);
|
||||
test_textureStore_2d(const_type_4_, const_type_2_, image_storage_2d);
|
||||
test_textureStore_2d_array(const_type_4_, 0, const_type_2_, image_storage_2d_array);
|
||||
test_textureStore_3d(const_type_7_, const_type_2_, image_storage_3d);
|
||||
return fragment_shaderOutput { metal::float4(0.0, 0.0, 0.0, 0.0) };
|
||||
}
|
||||
|
||||
@ -10,6 +10,9 @@ struct DefaultConstructible {
|
||||
return T {};
|
||||
}
|
||||
};
|
||||
constant metal::int2 const_type_4_ = {0, 0};
|
||||
constant metal::int3 const_type_7_ = {0, 0, 0};
|
||||
constant metal::float4 const_type_2_ = {0.0, 0.0, 0.0, 0.0};
|
||||
|
||||
metal::float4 test_textureLoad_1d(
|
||||
int coords,
|
||||
@ -129,3 +132,29 @@ void test_textureStore_3d(
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
||||
struct fragment_shaderOutput {
|
||||
metal::float4 member [[color(0)]];
|
||||
};
|
||||
fragment fragment_shaderOutput fragment_shader(
|
||||
metal::texture1d<float, metal::access::sample> image_1d [[user(fake0)]]
|
||||
, metal::texture2d<float, metal::access::sample> image_2d [[user(fake0)]]
|
||||
, metal::texture2d_array<float, metal::access::sample> image_2d_array [[user(fake0)]]
|
||||
, metal::texture3d<float, metal::access::sample> image_3d [[user(fake0)]]
|
||||
, metal::texture2d_ms<float, metal::access::read> image_multisampled_2d [[user(fake0)]]
|
||||
, metal::texture1d<float, metal::access::write> image_storage_1d [[user(fake0)]]
|
||||
, metal::texture2d<float, metal::access::write> image_storage_2d [[user(fake0)]]
|
||||
, metal::texture2d_array<float, metal::access::write> image_storage_2d_array [[user(fake0)]]
|
||||
, metal::texture3d<float, metal::access::write> image_storage_3d [[user(fake0)]]
|
||||
) {
|
||||
metal::float4 _e14 = test_textureLoad_1d(0, 0, image_1d);
|
||||
metal::float4 _e17 = test_textureLoad_2d(const_type_4_, 0, image_2d);
|
||||
metal::float4 _e21 = test_textureLoad_2d_array(const_type_4_, 0, 0, image_2d_array);
|
||||
metal::float4 _e24 = test_textureLoad_3d(const_type_7_, 0, image_3d);
|
||||
metal::float4 _e27 = test_textureLoad_multisampled_2d(const_type_4_, 0, image_multisampled_2d);
|
||||
test_textureStore_1d(0, const_type_2_, image_storage_1d);
|
||||
test_textureStore_2d(const_type_4_, const_type_2_, image_storage_2d);
|
||||
test_textureStore_2d_array(const_type_4_, 0, const_type_2_, image_storage_2d_array);
|
||||
test_textureStore_3d(const_type_7_, const_type_2_, image_storage_3d);
|
||||
return fragment_shaderOutput { metal::float4(0.0, 0.0, 0.0, 0.0) };
|
||||
}
|
||||
|
||||
@ -1,343 +1,380 @@
|
||||
; SPIR-V
|
||||
; Version: 1.1
|
||||
; Generator: rspirv
|
||||
; Bound: 214
|
||||
; Bound: 244
|
||||
OpCapability ImageQuery
|
||||
OpCapability Image1D
|
||||
OpCapability Shader
|
||||
OpCapability Sampled1D
|
||||
OpCapability Linkage
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint Fragment %222 "fragment_shader" %220
|
||||
OpExecutionMode %222 OriginUpperLeft
|
||||
OpSource GLSL 450
|
||||
OpName %20 "image_1d"
|
||||
OpName %22 "image_2d"
|
||||
OpName %24 "image_2d_array"
|
||||
OpName %26 "image_3d"
|
||||
OpName %28 "image_multisampled_2d"
|
||||
OpName %30 "image_depth_2d"
|
||||
OpName %32 "image_depth_2d_array"
|
||||
OpName %34 "image_depth_multisampled_2d"
|
||||
OpName %36 "image_storage_1d"
|
||||
OpName %38 "image_storage_2d"
|
||||
OpName %40 "image_storage_2d_array"
|
||||
OpName %42 "image_storage_3d"
|
||||
OpName %45 "coords"
|
||||
OpName %46 "level"
|
||||
OpName %47 "test_textureLoad_1d"
|
||||
OpName %60 "coords"
|
||||
OpName %61 "level"
|
||||
OpName %62 "test_textureLoad_2d"
|
||||
OpName %75 "coords"
|
||||
OpName %76 "index"
|
||||
OpName %77 "level"
|
||||
OpName %78 "test_textureLoad_2d_array"
|
||||
OpName %92 "coords"
|
||||
OpName %93 "level"
|
||||
OpName %94 "test_textureLoad_3d"
|
||||
OpName %107 "coords"
|
||||
OpName %108 "_sample"
|
||||
OpName %109 "test_textureLoad_multisampled_2d"
|
||||
OpName %121 "coords"
|
||||
OpName %122 "level"
|
||||
OpName %123 "test_textureLoad_depth_2d"
|
||||
OpName %137 "coords"
|
||||
OpName %138 "index"
|
||||
OpName %139 "level"
|
||||
OpName %140 "test_textureLoad_depth_2d_array"
|
||||
OpName %155 "coords"
|
||||
OpName %156 "_sample"
|
||||
OpName %157 "test_textureLoad_depth_multisampled_2d"
|
||||
OpName %170 "coords"
|
||||
OpName %171 "value"
|
||||
OpName %172 "test_textureStore_1d"
|
||||
OpName %180 "coords"
|
||||
OpName %181 "value"
|
||||
OpName %182 "test_textureStore_2d"
|
||||
OpName %191 "coords"
|
||||
OpName %192 "array_index"
|
||||
OpName %193 "value"
|
||||
OpName %194 "test_textureStore_2d_array"
|
||||
OpName %204 "coords"
|
||||
OpName %205 "value"
|
||||
OpName %206 "test_textureStore_3d"
|
||||
OpDecorate %20 DescriptorSet 0
|
||||
OpDecorate %20 Binding 0
|
||||
OpDecorate %22 DescriptorSet 0
|
||||
OpDecorate %22 Binding 0
|
||||
OpDecorate %24 DescriptorSet 0
|
||||
OpDecorate %24 Binding 0
|
||||
OpDecorate %26 DescriptorSet 0
|
||||
OpDecorate %26 Binding 0
|
||||
OpDecorate %28 DescriptorSet 0
|
||||
OpDecorate %28 Binding 0
|
||||
OpDecorate %30 DescriptorSet 0
|
||||
OpDecorate %30 Binding 0
|
||||
OpDecorate %32 DescriptorSet 0
|
||||
OpDecorate %32 Binding 0
|
||||
OpDecorate %34 DescriptorSet 0
|
||||
OpDecorate %34 Binding 0
|
||||
OpDecorate %36 NonReadable
|
||||
OpDecorate %36 DescriptorSet 0
|
||||
OpDecorate %36 Binding 0
|
||||
OpDecorate %38 NonReadable
|
||||
OpDecorate %38 DescriptorSet 0
|
||||
OpDecorate %38 Binding 0
|
||||
OpDecorate %40 NonReadable
|
||||
OpDecorate %40 DescriptorSet 0
|
||||
OpDecorate %40 Binding 0
|
||||
OpDecorate %42 NonReadable
|
||||
OpDecorate %42 DescriptorSet 0
|
||||
OpDecorate %42 Binding 0
|
||||
OpName %25 "image_1d"
|
||||
OpName %27 "image_2d"
|
||||
OpName %29 "image_2d_array"
|
||||
OpName %31 "image_3d"
|
||||
OpName %33 "image_multisampled_2d"
|
||||
OpName %35 "image_depth_2d"
|
||||
OpName %37 "image_depth_2d_array"
|
||||
OpName %39 "image_depth_multisampled_2d"
|
||||
OpName %41 "image_storage_1d"
|
||||
OpName %43 "image_storage_2d"
|
||||
OpName %45 "image_storage_2d_array"
|
||||
OpName %47 "image_storage_3d"
|
||||
OpName %50 "coords"
|
||||
OpName %51 "level"
|
||||
OpName %52 "test_textureLoad_1d"
|
||||
OpName %65 "coords"
|
||||
OpName %66 "level"
|
||||
OpName %67 "test_textureLoad_2d"
|
||||
OpName %80 "coords"
|
||||
OpName %81 "index"
|
||||
OpName %82 "level"
|
||||
OpName %83 "test_textureLoad_2d_array"
|
||||
OpName %97 "coords"
|
||||
OpName %98 "level"
|
||||
OpName %99 "test_textureLoad_3d"
|
||||
OpName %112 "coords"
|
||||
OpName %113 "_sample"
|
||||
OpName %114 "test_textureLoad_multisampled_2d"
|
||||
OpName %126 "coords"
|
||||
OpName %127 "level"
|
||||
OpName %128 "test_textureLoad_depth_2d"
|
||||
OpName %142 "coords"
|
||||
OpName %143 "index"
|
||||
OpName %144 "level"
|
||||
OpName %145 "test_textureLoad_depth_2d_array"
|
||||
OpName %160 "coords"
|
||||
OpName %161 "_sample"
|
||||
OpName %162 "test_textureLoad_depth_multisampled_2d"
|
||||
OpName %175 "coords"
|
||||
OpName %176 "value"
|
||||
OpName %177 "test_textureStore_1d"
|
||||
OpName %185 "coords"
|
||||
OpName %186 "value"
|
||||
OpName %187 "test_textureStore_2d"
|
||||
OpName %196 "coords"
|
||||
OpName %197 "array_index"
|
||||
OpName %198 "value"
|
||||
OpName %199 "test_textureStore_2d_array"
|
||||
OpName %209 "coords"
|
||||
OpName %210 "value"
|
||||
OpName %211 "test_textureStore_3d"
|
||||
OpName %222 "fragment_shader"
|
||||
OpDecorate %25 DescriptorSet 0
|
||||
OpDecorate %25 Binding 0
|
||||
OpDecorate %27 DescriptorSet 0
|
||||
OpDecorate %27 Binding 1
|
||||
OpDecorate %29 DescriptorSet 0
|
||||
OpDecorate %29 Binding 2
|
||||
OpDecorate %31 DescriptorSet 0
|
||||
OpDecorate %31 Binding 3
|
||||
OpDecorate %33 DescriptorSet 0
|
||||
OpDecorate %33 Binding 4
|
||||
OpDecorate %35 DescriptorSet 0
|
||||
OpDecorate %35 Binding 5
|
||||
OpDecorate %37 DescriptorSet 0
|
||||
OpDecorate %37 Binding 6
|
||||
OpDecorate %39 DescriptorSet 0
|
||||
OpDecorate %39 Binding 7
|
||||
OpDecorate %41 NonReadable
|
||||
OpDecorate %41 DescriptorSet 0
|
||||
OpDecorate %41 Binding 8
|
||||
OpDecorate %43 NonReadable
|
||||
OpDecorate %43 DescriptorSet 0
|
||||
OpDecorate %43 Binding 9
|
||||
OpDecorate %45 NonReadable
|
||||
OpDecorate %45 DescriptorSet 0
|
||||
OpDecorate %45 Binding 10
|
||||
OpDecorate %47 NonReadable
|
||||
OpDecorate %47 DescriptorSet 0
|
||||
OpDecorate %47 Binding 11
|
||||
OpDecorate %220 Location 0
|
||||
%2 = OpTypeVoid
|
||||
%4 = OpTypeFloat 32
|
||||
%3 = OpTypeImage %4 1D 0 0 0 1 Unknown
|
||||
%5 = OpTypeInt 32 1
|
||||
%6 = OpTypeVector %4 4
|
||||
%7 = OpTypeImage %4 2D 0 0 0 1 Unknown
|
||||
%8 = OpTypeVector %5 2
|
||||
%9 = OpTypeImage %4 2D 0 1 0 1 Unknown
|
||||
%10 = OpTypeImage %4 3D 0 0 0 1 Unknown
|
||||
%11 = OpTypeVector %5 3
|
||||
%12 = OpTypeImage %4 2D 0 0 1 1 Unknown
|
||||
%13 = OpTypeImage %4 2D 1 0 0 1 Unknown
|
||||
%14 = OpTypeImage %4 2D 1 1 0 1 Unknown
|
||||
%15 = OpTypeImage %4 2D 1 0 1 1 Unknown
|
||||
%16 = OpTypeImage %4 1D 0 0 0 2 Rgba8
|
||||
%17 = OpTypeImage %4 2D 0 0 0 2 Rgba8
|
||||
%18 = OpTypeImage %4 2D 0 1 0 2 Rgba8
|
||||
%19 = OpTypeImage %4 3D 0 0 0 2 Rgba8
|
||||
%21 = OpTypePointer UniformConstant %3
|
||||
%20 = OpVariable %21 UniformConstant
|
||||
%23 = OpTypePointer UniformConstant %7
|
||||
%22 = OpVariable %23 UniformConstant
|
||||
%25 = OpTypePointer UniformConstant %9
|
||||
%24 = OpVariable %25 UniformConstant
|
||||
%27 = OpTypePointer UniformConstant %10
|
||||
%26 = OpVariable %27 UniformConstant
|
||||
%29 = OpTypePointer UniformConstant %12
|
||||
%28 = OpVariable %29 UniformConstant
|
||||
%31 = OpTypePointer UniformConstant %13
|
||||
%30 = OpVariable %31 UniformConstant
|
||||
%33 = OpTypePointer UniformConstant %14
|
||||
%32 = OpVariable %33 UniformConstant
|
||||
%35 = OpTypePointer UniformConstant %15
|
||||
%34 = OpVariable %35 UniformConstant
|
||||
%37 = OpTypePointer UniformConstant %16
|
||||
%36 = OpVariable %37 UniformConstant
|
||||
%39 = OpTypePointer UniformConstant %17
|
||||
%38 = OpVariable %39 UniformConstant
|
||||
%41 = OpTypePointer UniformConstant %18
|
||||
%40 = OpVariable %41 UniformConstant
|
||||
%43 = OpTypePointer UniformConstant %19
|
||||
%42 = OpVariable %43 UniformConstant
|
||||
%48 = OpTypeFunction %6 %5 %5
|
||||
%52 = OpConstant %5 1
|
||||
%63 = OpTypeFunction %6 %8 %5
|
||||
%70 = OpConstantComposite %8 %52 %52
|
||||
%79 = OpTypeFunction %6 %8 %5 %5
|
||||
%87 = OpConstantComposite %11 %52 %52 %52
|
||||
%95 = OpTypeFunction %6 %11 %5
|
||||
%102 = OpConstantComposite %11 %52 %52 %52
|
||||
%116 = OpConstantComposite %8 %52 %52
|
||||
%124 = OpTypeFunction %4 %8 %5
|
||||
%131 = OpConstantComposite %8 %52 %52
|
||||
%141 = OpTypeFunction %4 %8 %5 %5
|
||||
%149 = OpConstantComposite %11 %52 %52 %52
|
||||
%164 = OpConstantComposite %8 %52 %52
|
||||
%173 = OpTypeFunction %2 %5 %6
|
||||
%183 = OpTypeFunction %2 %8 %6
|
||||
%187 = OpConstantComposite %8 %52 %52
|
||||
%195 = OpTypeFunction %2 %8 %5 %6
|
||||
%200 = OpConstantComposite %11 %52 %52 %52
|
||||
%207 = OpTypeFunction %2 %11 %6
|
||||
%211 = OpConstantComposite %11 %52 %52 %52
|
||||
%47 = OpFunction %6 None %48
|
||||
%45 = OpFunctionParameter %5
|
||||
%46 = OpFunctionParameter %5
|
||||
%44 = OpLabel
|
||||
%49 = OpLoad %3 %20
|
||||
OpBranch %50
|
||||
%50 = OpLabel
|
||||
%51 = OpImageQueryLevels %5 %49
|
||||
%53 = OpISub %5 %51 %52
|
||||
%54 = OpExtInst %5 %1 UMin %46 %53
|
||||
%55 = OpImageQuerySizeLod %5 %49 %54
|
||||
%56 = OpISub %5 %55 %52
|
||||
%57 = OpExtInst %5 %1 UMin %45 %56
|
||||
%58 = OpImageFetch %6 %49 %57 Lod %54
|
||||
OpReturnValue %58
|
||||
%4 = OpTypeInt 32 1
|
||||
%3 = OpConstant %4 0
|
||||
%6 = OpTypeFloat 32
|
||||
%5 = OpConstant %6 0.0
|
||||
%7 = OpTypeImage %6 1D 0 0 0 1 Unknown
|
||||
%8 = OpTypeVector %6 4
|
||||
%9 = OpTypeImage %6 2D 0 0 0 1 Unknown
|
||||
%10 = OpTypeVector %4 2
|
||||
%11 = OpTypeImage %6 2D 0 1 0 1 Unknown
|
||||
%12 = OpTypeImage %6 3D 0 0 0 1 Unknown
|
||||
%13 = OpTypeVector %4 3
|
||||
%14 = OpTypeImage %6 2D 0 0 1 1 Unknown
|
||||
%15 = OpTypeImage %6 2D 1 0 0 1 Unknown
|
||||
%16 = OpTypeImage %6 2D 1 1 0 1 Unknown
|
||||
%17 = OpTypeImage %6 2D 1 0 1 1 Unknown
|
||||
%18 = OpTypeImage %6 1D 0 0 0 2 Rgba8
|
||||
%19 = OpTypeImage %6 2D 0 0 0 2 Rgba8
|
||||
%20 = OpTypeImage %6 2D 0 1 0 2 Rgba8
|
||||
%21 = OpTypeImage %6 3D 0 0 0 2 Rgba8
|
||||
%22 = OpConstantComposite %10 %3 %3
|
||||
%23 = OpConstantComposite %13 %3 %3 %3
|
||||
%24 = OpConstantComposite %8 %5 %5 %5 %5
|
||||
%26 = OpTypePointer UniformConstant %7
|
||||
%25 = OpVariable %26 UniformConstant
|
||||
%28 = OpTypePointer UniformConstant %9
|
||||
%27 = OpVariable %28 UniformConstant
|
||||
%30 = OpTypePointer UniformConstant %11
|
||||
%29 = OpVariable %30 UniformConstant
|
||||
%32 = OpTypePointer UniformConstant %12
|
||||
%31 = OpVariable %32 UniformConstant
|
||||
%34 = OpTypePointer UniformConstant %14
|
||||
%33 = OpVariable %34 UniformConstant
|
||||
%36 = OpTypePointer UniformConstant %15
|
||||
%35 = OpVariable %36 UniformConstant
|
||||
%38 = OpTypePointer UniformConstant %16
|
||||
%37 = OpVariable %38 UniformConstant
|
||||
%40 = OpTypePointer UniformConstant %17
|
||||
%39 = OpVariable %40 UniformConstant
|
||||
%42 = OpTypePointer UniformConstant %18
|
||||
%41 = OpVariable %42 UniformConstant
|
||||
%44 = OpTypePointer UniformConstant %19
|
||||
%43 = OpVariable %44 UniformConstant
|
||||
%46 = OpTypePointer UniformConstant %20
|
||||
%45 = OpVariable %46 UniformConstant
|
||||
%48 = OpTypePointer UniformConstant %21
|
||||
%47 = OpVariable %48 UniformConstant
|
||||
%53 = OpTypeFunction %8 %4 %4
|
||||
%57 = OpConstant %4 1
|
||||
%68 = OpTypeFunction %8 %10 %4
|
||||
%75 = OpConstantComposite %10 %57 %57
|
||||
%84 = OpTypeFunction %8 %10 %4 %4
|
||||
%92 = OpConstantComposite %13 %57 %57 %57
|
||||
%100 = OpTypeFunction %8 %13 %4
|
||||
%107 = OpConstantComposite %13 %57 %57 %57
|
||||
%121 = OpConstantComposite %10 %57 %57
|
||||
%129 = OpTypeFunction %6 %10 %4
|
||||
%136 = OpConstantComposite %10 %57 %57
|
||||
%146 = OpTypeFunction %6 %10 %4 %4
|
||||
%154 = OpConstantComposite %13 %57 %57 %57
|
||||
%169 = OpConstantComposite %10 %57 %57
|
||||
%178 = OpTypeFunction %2 %4 %8
|
||||
%188 = OpTypeFunction %2 %10 %8
|
||||
%192 = OpConstantComposite %10 %57 %57
|
||||
%200 = OpTypeFunction %2 %10 %4 %8
|
||||
%205 = OpConstantComposite %13 %57 %57 %57
|
||||
%212 = OpTypeFunction %2 %13 %8
|
||||
%216 = OpConstantComposite %13 %57 %57 %57
|
||||
%221 = OpTypePointer Output %8
|
||||
%220 = OpVariable %221 Output
|
||||
%223 = OpTypeFunction %2
|
||||
%52 = OpFunction %8 None %53
|
||||
%50 = OpFunctionParameter %4
|
||||
%51 = OpFunctionParameter %4
|
||||
%49 = OpLabel
|
||||
%54 = OpLoad %7 %25
|
||||
OpBranch %55
|
||||
%55 = OpLabel
|
||||
%56 = OpImageQueryLevels %4 %54
|
||||
%58 = OpISub %4 %56 %57
|
||||
%59 = OpExtInst %4 %1 UMin %51 %58
|
||||
%60 = OpImageQuerySizeLod %4 %54 %59
|
||||
%61 = OpISub %4 %60 %57
|
||||
%62 = OpExtInst %4 %1 UMin %50 %61
|
||||
%63 = OpImageFetch %8 %54 %62 Lod %59
|
||||
OpReturnValue %63
|
||||
OpFunctionEnd
|
||||
%62 = OpFunction %6 None %63
|
||||
%60 = OpFunctionParameter %8
|
||||
%61 = OpFunctionParameter %5
|
||||
%59 = OpLabel
|
||||
%64 = OpLoad %7 %22
|
||||
OpBranch %65
|
||||
%65 = OpLabel
|
||||
%66 = OpImageQueryLevels %5 %64
|
||||
%67 = OpISub %5 %66 %52
|
||||
%68 = OpExtInst %5 %1 UMin %61 %67
|
||||
%69 = OpImageQuerySizeLod %8 %64 %68
|
||||
%71 = OpISub %8 %69 %70
|
||||
%72 = OpExtInst %8 %1 UMin %60 %71
|
||||
%73 = OpImageFetch %6 %64 %72 Lod %68
|
||||
OpReturnValue %73
|
||||
%67 = OpFunction %8 None %68
|
||||
%65 = OpFunctionParameter %10
|
||||
%66 = OpFunctionParameter %4
|
||||
%64 = OpLabel
|
||||
%69 = OpLoad %9 %27
|
||||
OpBranch %70
|
||||
%70 = OpLabel
|
||||
%71 = OpImageQueryLevels %4 %69
|
||||
%72 = OpISub %4 %71 %57
|
||||
%73 = OpExtInst %4 %1 UMin %66 %72
|
||||
%74 = OpImageQuerySizeLod %10 %69 %73
|
||||
%76 = OpISub %10 %74 %75
|
||||
%77 = OpExtInst %10 %1 UMin %65 %76
|
||||
%78 = OpImageFetch %8 %69 %77 Lod %73
|
||||
OpReturnValue %78
|
||||
OpFunctionEnd
|
||||
%78 = OpFunction %6 None %79
|
||||
%75 = OpFunctionParameter %8
|
||||
%76 = OpFunctionParameter %5
|
||||
%77 = OpFunctionParameter %5
|
||||
%74 = OpLabel
|
||||
%80 = OpLoad %9 %24
|
||||
OpBranch %81
|
||||
%81 = OpLabel
|
||||
%82 = OpCompositeConstruct %11 %75 %76
|
||||
%83 = OpImageQueryLevels %5 %80
|
||||
%84 = OpISub %5 %83 %52
|
||||
%85 = OpExtInst %5 %1 UMin %77 %84
|
||||
%86 = OpImageQuerySizeLod %11 %80 %85
|
||||
%88 = OpISub %11 %86 %87
|
||||
%89 = OpExtInst %11 %1 UMin %82 %88
|
||||
%90 = OpImageFetch %6 %80 %89 Lod %85
|
||||
OpReturnValue %90
|
||||
%83 = OpFunction %8 None %84
|
||||
%80 = OpFunctionParameter %10
|
||||
%81 = OpFunctionParameter %4
|
||||
%82 = OpFunctionParameter %4
|
||||
%79 = OpLabel
|
||||
%85 = OpLoad %11 %29
|
||||
OpBranch %86
|
||||
%86 = OpLabel
|
||||
%87 = OpCompositeConstruct %13 %80 %81
|
||||
%88 = OpImageQueryLevels %4 %85
|
||||
%89 = OpISub %4 %88 %57
|
||||
%90 = OpExtInst %4 %1 UMin %82 %89
|
||||
%91 = OpImageQuerySizeLod %13 %85 %90
|
||||
%93 = OpISub %13 %91 %92
|
||||
%94 = OpExtInst %13 %1 UMin %87 %93
|
||||
%95 = OpImageFetch %8 %85 %94 Lod %90
|
||||
OpReturnValue %95
|
||||
OpFunctionEnd
|
||||
%94 = OpFunction %6 None %95
|
||||
%92 = OpFunctionParameter %11
|
||||
%93 = OpFunctionParameter %5
|
||||
%91 = OpLabel
|
||||
%96 = OpLoad %10 %26
|
||||
OpBranch %97
|
||||
%97 = OpLabel
|
||||
%98 = OpImageQueryLevels %5 %96
|
||||
%99 = OpISub %5 %98 %52
|
||||
%100 = OpExtInst %5 %1 UMin %93 %99
|
||||
%101 = OpImageQuerySizeLod %11 %96 %100
|
||||
%103 = OpISub %11 %101 %102
|
||||
%104 = OpExtInst %11 %1 UMin %92 %103
|
||||
%105 = OpImageFetch %6 %96 %104 Lod %100
|
||||
OpReturnValue %105
|
||||
%99 = OpFunction %8 None %100
|
||||
%97 = OpFunctionParameter %13
|
||||
%98 = OpFunctionParameter %4
|
||||
%96 = OpLabel
|
||||
%101 = OpLoad %12 %31
|
||||
OpBranch %102
|
||||
%102 = OpLabel
|
||||
%103 = OpImageQueryLevels %4 %101
|
||||
%104 = OpISub %4 %103 %57
|
||||
%105 = OpExtInst %4 %1 UMin %98 %104
|
||||
%106 = OpImageQuerySizeLod %13 %101 %105
|
||||
%108 = OpISub %13 %106 %107
|
||||
%109 = OpExtInst %13 %1 UMin %97 %108
|
||||
%110 = OpImageFetch %8 %101 %109 Lod %105
|
||||
OpReturnValue %110
|
||||
OpFunctionEnd
|
||||
%109 = OpFunction %6 None %63
|
||||
%107 = OpFunctionParameter %8
|
||||
%108 = OpFunctionParameter %5
|
||||
%106 = OpLabel
|
||||
%110 = OpLoad %12 %28
|
||||
OpBranch %111
|
||||
%114 = OpFunction %8 None %68
|
||||
%112 = OpFunctionParameter %10
|
||||
%113 = OpFunctionParameter %4
|
||||
%111 = OpLabel
|
||||
%112 = OpImageQuerySamples %5 %110
|
||||
%113 = OpISub %5 %112 %52
|
||||
%114 = OpExtInst %5 %1 UMin %108 %113
|
||||
%115 = OpImageQuerySize %8 %110
|
||||
%117 = OpISub %8 %115 %116
|
||||
%118 = OpExtInst %8 %1 UMin %107 %117
|
||||
%119 = OpImageFetch %6 %110 %118 Sample %114
|
||||
OpReturnValue %119
|
||||
%115 = OpLoad %14 %33
|
||||
OpBranch %116
|
||||
%116 = OpLabel
|
||||
%117 = OpImageQuerySamples %4 %115
|
||||
%118 = OpISub %4 %117 %57
|
||||
%119 = OpExtInst %4 %1 UMin %113 %118
|
||||
%120 = OpImageQuerySize %10 %115
|
||||
%122 = OpISub %10 %120 %121
|
||||
%123 = OpExtInst %10 %1 UMin %112 %122
|
||||
%124 = OpImageFetch %8 %115 %123 Sample %119
|
||||
OpReturnValue %124
|
||||
OpFunctionEnd
|
||||
%123 = OpFunction %4 None %124
|
||||
%121 = OpFunctionParameter %8
|
||||
%122 = OpFunctionParameter %5
|
||||
%120 = OpLabel
|
||||
%125 = OpLoad %13 %30
|
||||
OpBranch %126
|
||||
%126 = OpLabel
|
||||
%127 = OpImageQueryLevels %5 %125
|
||||
%128 = OpISub %5 %127 %52
|
||||
%129 = OpExtInst %5 %1 UMin %122 %128
|
||||
%130 = OpImageQuerySizeLod %8 %125 %129
|
||||
%132 = OpISub %8 %130 %131
|
||||
%133 = OpExtInst %8 %1 UMin %121 %132
|
||||
%134 = OpImageFetch %6 %125 %133 Lod %129
|
||||
%135 = OpCompositeExtract %4 %134 0
|
||||
OpReturnValue %135
|
||||
%128 = OpFunction %6 None %129
|
||||
%126 = OpFunctionParameter %10
|
||||
%127 = OpFunctionParameter %4
|
||||
%125 = OpLabel
|
||||
%130 = OpLoad %15 %35
|
||||
OpBranch %131
|
||||
%131 = OpLabel
|
||||
%132 = OpImageQueryLevels %4 %130
|
||||
%133 = OpISub %4 %132 %57
|
||||
%134 = OpExtInst %4 %1 UMin %127 %133
|
||||
%135 = OpImageQuerySizeLod %10 %130 %134
|
||||
%137 = OpISub %10 %135 %136
|
||||
%138 = OpExtInst %10 %1 UMin %126 %137
|
||||
%139 = OpImageFetch %8 %130 %138 Lod %134
|
||||
%140 = OpCompositeExtract %6 %139 0
|
||||
OpReturnValue %140
|
||||
OpFunctionEnd
|
||||
%140 = OpFunction %4 None %141
|
||||
%137 = OpFunctionParameter %8
|
||||
%138 = OpFunctionParameter %5
|
||||
%139 = OpFunctionParameter %5
|
||||
%136 = OpLabel
|
||||
%142 = OpLoad %14 %32
|
||||
OpBranch %143
|
||||
%143 = OpLabel
|
||||
%144 = OpCompositeConstruct %11 %137 %138
|
||||
%145 = OpImageQueryLevels %5 %142
|
||||
%146 = OpISub %5 %145 %52
|
||||
%147 = OpExtInst %5 %1 UMin %139 %146
|
||||
%148 = OpImageQuerySizeLod %11 %142 %147
|
||||
%150 = OpISub %11 %148 %149
|
||||
%151 = OpExtInst %11 %1 UMin %144 %150
|
||||
%152 = OpImageFetch %6 %142 %151 Lod %147
|
||||
%153 = OpCompositeExtract %4 %152 0
|
||||
OpReturnValue %153
|
||||
%145 = OpFunction %6 None %146
|
||||
%142 = OpFunctionParameter %10
|
||||
%143 = OpFunctionParameter %4
|
||||
%144 = OpFunctionParameter %4
|
||||
%141 = OpLabel
|
||||
%147 = OpLoad %16 %37
|
||||
OpBranch %148
|
||||
%148 = OpLabel
|
||||
%149 = OpCompositeConstruct %13 %142 %143
|
||||
%150 = OpImageQueryLevels %4 %147
|
||||
%151 = OpISub %4 %150 %57
|
||||
%152 = OpExtInst %4 %1 UMin %144 %151
|
||||
%153 = OpImageQuerySizeLod %13 %147 %152
|
||||
%155 = OpISub %13 %153 %154
|
||||
%156 = OpExtInst %13 %1 UMin %149 %155
|
||||
%157 = OpImageFetch %8 %147 %156 Lod %152
|
||||
%158 = OpCompositeExtract %6 %157 0
|
||||
OpReturnValue %158
|
||||
OpFunctionEnd
|
||||
%157 = OpFunction %4 None %124
|
||||
%155 = OpFunctionParameter %8
|
||||
%156 = OpFunctionParameter %5
|
||||
%154 = OpLabel
|
||||
%158 = OpLoad %15 %34
|
||||
OpBranch %159
|
||||
%162 = OpFunction %6 None %129
|
||||
%160 = OpFunctionParameter %10
|
||||
%161 = OpFunctionParameter %4
|
||||
%159 = OpLabel
|
||||
%160 = OpImageQuerySamples %5 %158
|
||||
%161 = OpISub %5 %160 %52
|
||||
%162 = OpExtInst %5 %1 UMin %156 %161
|
||||
%163 = OpImageQuerySize %8 %158
|
||||
%165 = OpISub %8 %163 %164
|
||||
%166 = OpExtInst %8 %1 UMin %155 %165
|
||||
%167 = OpImageFetch %6 %158 %166 Sample %162
|
||||
%168 = OpCompositeExtract %4 %167 0
|
||||
OpReturnValue %168
|
||||
%163 = OpLoad %17 %39
|
||||
OpBranch %164
|
||||
%164 = OpLabel
|
||||
%165 = OpImageQuerySamples %4 %163
|
||||
%166 = OpISub %4 %165 %57
|
||||
%167 = OpExtInst %4 %1 UMin %161 %166
|
||||
%168 = OpImageQuerySize %10 %163
|
||||
%170 = OpISub %10 %168 %169
|
||||
%171 = OpExtInst %10 %1 UMin %160 %170
|
||||
%172 = OpImageFetch %8 %163 %171 Sample %167
|
||||
%173 = OpCompositeExtract %6 %172 0
|
||||
OpReturnValue %173
|
||||
OpFunctionEnd
|
||||
%172 = OpFunction %2 None %173
|
||||
%170 = OpFunctionParameter %5
|
||||
%171 = OpFunctionParameter %6
|
||||
%169 = OpLabel
|
||||
%174 = OpLoad %16 %36
|
||||
OpBranch %175
|
||||
%175 = OpLabel
|
||||
%176 = OpImageQuerySize %5 %174
|
||||
%177 = OpISub %5 %176 %52
|
||||
%178 = OpExtInst %5 %1 UMin %170 %177
|
||||
OpImageWrite %174 %178 %171
|
||||
%177 = OpFunction %2 None %178
|
||||
%175 = OpFunctionParameter %4
|
||||
%176 = OpFunctionParameter %8
|
||||
%174 = OpLabel
|
||||
%179 = OpLoad %18 %41
|
||||
OpBranch %180
|
||||
%180 = OpLabel
|
||||
%181 = OpImageQuerySize %4 %179
|
||||
%182 = OpISub %4 %181 %57
|
||||
%183 = OpExtInst %4 %1 UMin %175 %182
|
||||
OpImageWrite %179 %183 %176
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%182 = OpFunction %2 None %183
|
||||
%180 = OpFunctionParameter %8
|
||||
%181 = OpFunctionParameter %6
|
||||
%179 = OpLabel
|
||||
%184 = OpLoad %17 %38
|
||||
OpBranch %185
|
||||
%185 = OpLabel
|
||||
%186 = OpImageQuerySize %8 %184
|
||||
%188 = OpISub %8 %186 %187
|
||||
%189 = OpExtInst %8 %1 UMin %180 %188
|
||||
OpImageWrite %184 %189 %181
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%194 = OpFunction %2 None %195
|
||||
%191 = OpFunctionParameter %8
|
||||
%192 = OpFunctionParameter %5
|
||||
%193 = OpFunctionParameter %6
|
||||
%187 = OpFunction %2 None %188
|
||||
%185 = OpFunctionParameter %10
|
||||
%186 = OpFunctionParameter %8
|
||||
%184 = OpLabel
|
||||
%189 = OpLoad %19 %43
|
||||
OpBranch %190
|
||||
%190 = OpLabel
|
||||
%196 = OpLoad %18 %40
|
||||
OpBranch %197
|
||||
%197 = OpLabel
|
||||
%198 = OpCompositeConstruct %11 %191 %192
|
||||
%199 = OpImageQuerySize %11 %196
|
||||
%201 = OpISub %11 %199 %200
|
||||
%202 = OpExtInst %11 %1 UMin %198 %201
|
||||
OpImageWrite %196 %202 %193
|
||||
%191 = OpImageQuerySize %10 %189
|
||||
%193 = OpISub %10 %191 %192
|
||||
%194 = OpExtInst %10 %1 UMin %185 %193
|
||||
OpImageWrite %189 %194 %186
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%206 = OpFunction %2 None %207
|
||||
%204 = OpFunctionParameter %11
|
||||
%205 = OpFunctionParameter %6
|
||||
%203 = OpLabel
|
||||
%208 = OpLoad %19 %42
|
||||
OpBranch %209
|
||||
%209 = OpLabel
|
||||
%210 = OpImageQuerySize %11 %208
|
||||
%212 = OpISub %11 %210 %211
|
||||
%213 = OpExtInst %11 %1 UMin %204 %212
|
||||
OpImageWrite %208 %213 %205
|
||||
%199 = OpFunction %2 None %200
|
||||
%196 = OpFunctionParameter %10
|
||||
%197 = OpFunctionParameter %4
|
||||
%198 = OpFunctionParameter %8
|
||||
%195 = OpLabel
|
||||
%201 = OpLoad %20 %45
|
||||
OpBranch %202
|
||||
%202 = OpLabel
|
||||
%203 = OpCompositeConstruct %13 %196 %197
|
||||
%204 = OpImageQuerySize %13 %201
|
||||
%206 = OpISub %13 %204 %205
|
||||
%207 = OpExtInst %13 %1 UMin %203 %206
|
||||
OpImageWrite %201 %207 %198
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%211 = OpFunction %2 None %212
|
||||
%209 = OpFunctionParameter %13
|
||||
%210 = OpFunctionParameter %8
|
||||
%208 = OpLabel
|
||||
%213 = OpLoad %21 %47
|
||||
OpBranch %214
|
||||
%214 = OpLabel
|
||||
%215 = OpImageQuerySize %13 %213
|
||||
%217 = OpISub %13 %215 %216
|
||||
%218 = OpExtInst %13 %1 UMin %209 %217
|
||||
OpImageWrite %213 %218 %210
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%222 = OpFunction %2 None %223
|
||||
%219 = OpLabel
|
||||
%224 = OpLoad %7 %25
|
||||
%225 = OpLoad %9 %27
|
||||
%226 = OpLoad %11 %29
|
||||
%227 = OpLoad %12 %31
|
||||
%228 = OpLoad %14 %33
|
||||
%229 = OpLoad %18 %41
|
||||
%230 = OpLoad %19 %43
|
||||
%231 = OpLoad %20 %45
|
||||
%232 = OpLoad %21 %47
|
||||
OpBranch %233
|
||||
%233 = OpLabel
|
||||
%234 = OpFunctionCall %8 %52 %3 %3
|
||||
%235 = OpFunctionCall %8 %67 %22 %3
|
||||
%236 = OpFunctionCall %8 %83 %22 %3 %3
|
||||
%237 = OpFunctionCall %8 %99 %23 %3
|
||||
%238 = OpFunctionCall %8 %114 %22 %3
|
||||
%239 = OpFunctionCall %2 %177 %3 %24
|
||||
%240 = OpFunctionCall %2 %187 %22 %24
|
||||
%241 = OpFunctionCall %2 %199 %22 %3 %24
|
||||
%242 = OpFunctionCall %2 %211 %23 %24
|
||||
%243 = OpCompositeConstruct %8 %5 %5 %5 %5
|
||||
OpStore %220 %243
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@ -1,417 +1,454 @@
|
||||
; SPIR-V
|
||||
; Version: 1.1
|
||||
; Generator: rspirv
|
||||
; Bound: 244
|
||||
; Bound: 274
|
||||
OpCapability ImageQuery
|
||||
OpCapability Image1D
|
||||
OpCapability Shader
|
||||
OpCapability Sampled1D
|
||||
OpCapability Linkage
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint Fragment %252 "fragment_shader" %250
|
||||
OpExecutionMode %252 OriginUpperLeft
|
||||
OpSource GLSL 450
|
||||
OpName %20 "image_1d"
|
||||
OpName %22 "image_2d"
|
||||
OpName %24 "image_2d_array"
|
||||
OpName %26 "image_3d"
|
||||
OpName %28 "image_multisampled_2d"
|
||||
OpName %30 "image_depth_2d"
|
||||
OpName %32 "image_depth_2d_array"
|
||||
OpName %34 "image_depth_multisampled_2d"
|
||||
OpName %36 "image_storage_1d"
|
||||
OpName %38 "image_storage_2d"
|
||||
OpName %40 "image_storage_2d_array"
|
||||
OpName %42 "image_storage_3d"
|
||||
OpName %45 "coords"
|
||||
OpName %46 "level"
|
||||
OpName %47 "test_textureLoad_1d"
|
||||
OpName %63 "coords"
|
||||
OpName %64 "level"
|
||||
OpName %65 "test_textureLoad_2d"
|
||||
OpName %82 "coords"
|
||||
OpName %83 "index"
|
||||
OpName %84 "level"
|
||||
OpName %85 "test_textureLoad_2d_array"
|
||||
OpName %103 "coords"
|
||||
OpName %104 "level"
|
||||
OpName %105 "test_textureLoad_3d"
|
||||
OpName %121 "coords"
|
||||
OpName %122 "_sample"
|
||||
OpName %123 "test_textureLoad_multisampled_2d"
|
||||
OpName %138 "coords"
|
||||
OpName %139 "level"
|
||||
OpName %140 "test_textureLoad_depth_2d"
|
||||
OpName %157 "coords"
|
||||
OpName %158 "index"
|
||||
OpName %159 "level"
|
||||
OpName %160 "test_textureLoad_depth_2d_array"
|
||||
OpName %178 "coords"
|
||||
OpName %179 "_sample"
|
||||
OpName %180 "test_textureLoad_depth_multisampled_2d"
|
||||
OpName %196 "coords"
|
||||
OpName %197 "value"
|
||||
OpName %198 "test_textureStore_1d"
|
||||
OpName %207 "coords"
|
||||
OpName %208 "value"
|
||||
OpName %209 "test_textureStore_2d"
|
||||
OpName %219 "coords"
|
||||
OpName %220 "array_index"
|
||||
OpName %221 "value"
|
||||
OpName %222 "test_textureStore_2d_array"
|
||||
OpName %233 "coords"
|
||||
OpName %234 "value"
|
||||
OpName %235 "test_textureStore_3d"
|
||||
OpDecorate %20 DescriptorSet 0
|
||||
OpDecorate %20 Binding 0
|
||||
OpDecorate %22 DescriptorSet 0
|
||||
OpDecorate %22 Binding 0
|
||||
OpDecorate %24 DescriptorSet 0
|
||||
OpDecorate %24 Binding 0
|
||||
OpDecorate %26 DescriptorSet 0
|
||||
OpDecorate %26 Binding 0
|
||||
OpDecorate %28 DescriptorSet 0
|
||||
OpDecorate %28 Binding 0
|
||||
OpDecorate %30 DescriptorSet 0
|
||||
OpDecorate %30 Binding 0
|
||||
OpDecorate %32 DescriptorSet 0
|
||||
OpDecorate %32 Binding 0
|
||||
OpDecorate %34 DescriptorSet 0
|
||||
OpDecorate %34 Binding 0
|
||||
OpDecorate %36 NonReadable
|
||||
OpDecorate %36 DescriptorSet 0
|
||||
OpDecorate %36 Binding 0
|
||||
OpDecorate %38 NonReadable
|
||||
OpDecorate %38 DescriptorSet 0
|
||||
OpDecorate %38 Binding 0
|
||||
OpDecorate %40 NonReadable
|
||||
OpDecorate %40 DescriptorSet 0
|
||||
OpDecorate %40 Binding 0
|
||||
OpDecorate %42 NonReadable
|
||||
OpDecorate %42 DescriptorSet 0
|
||||
OpDecorate %42 Binding 0
|
||||
OpName %25 "image_1d"
|
||||
OpName %27 "image_2d"
|
||||
OpName %29 "image_2d_array"
|
||||
OpName %31 "image_3d"
|
||||
OpName %33 "image_multisampled_2d"
|
||||
OpName %35 "image_depth_2d"
|
||||
OpName %37 "image_depth_2d_array"
|
||||
OpName %39 "image_depth_multisampled_2d"
|
||||
OpName %41 "image_storage_1d"
|
||||
OpName %43 "image_storage_2d"
|
||||
OpName %45 "image_storage_2d_array"
|
||||
OpName %47 "image_storage_3d"
|
||||
OpName %50 "coords"
|
||||
OpName %51 "level"
|
||||
OpName %52 "test_textureLoad_1d"
|
||||
OpName %68 "coords"
|
||||
OpName %69 "level"
|
||||
OpName %70 "test_textureLoad_2d"
|
||||
OpName %87 "coords"
|
||||
OpName %88 "index"
|
||||
OpName %89 "level"
|
||||
OpName %90 "test_textureLoad_2d_array"
|
||||
OpName %108 "coords"
|
||||
OpName %109 "level"
|
||||
OpName %110 "test_textureLoad_3d"
|
||||
OpName %126 "coords"
|
||||
OpName %127 "_sample"
|
||||
OpName %128 "test_textureLoad_multisampled_2d"
|
||||
OpName %143 "coords"
|
||||
OpName %144 "level"
|
||||
OpName %145 "test_textureLoad_depth_2d"
|
||||
OpName %162 "coords"
|
||||
OpName %163 "index"
|
||||
OpName %164 "level"
|
||||
OpName %165 "test_textureLoad_depth_2d_array"
|
||||
OpName %183 "coords"
|
||||
OpName %184 "_sample"
|
||||
OpName %185 "test_textureLoad_depth_multisampled_2d"
|
||||
OpName %201 "coords"
|
||||
OpName %202 "value"
|
||||
OpName %203 "test_textureStore_1d"
|
||||
OpName %212 "coords"
|
||||
OpName %213 "value"
|
||||
OpName %214 "test_textureStore_2d"
|
||||
OpName %224 "coords"
|
||||
OpName %225 "array_index"
|
||||
OpName %226 "value"
|
||||
OpName %227 "test_textureStore_2d_array"
|
||||
OpName %238 "coords"
|
||||
OpName %239 "value"
|
||||
OpName %240 "test_textureStore_3d"
|
||||
OpName %252 "fragment_shader"
|
||||
OpDecorate %25 DescriptorSet 0
|
||||
OpDecorate %25 Binding 0
|
||||
OpDecorate %27 DescriptorSet 0
|
||||
OpDecorate %27 Binding 1
|
||||
OpDecorate %29 DescriptorSet 0
|
||||
OpDecorate %29 Binding 2
|
||||
OpDecorate %31 DescriptorSet 0
|
||||
OpDecorate %31 Binding 3
|
||||
OpDecorate %33 DescriptorSet 0
|
||||
OpDecorate %33 Binding 4
|
||||
OpDecorate %35 DescriptorSet 0
|
||||
OpDecorate %35 Binding 5
|
||||
OpDecorate %37 DescriptorSet 0
|
||||
OpDecorate %37 Binding 6
|
||||
OpDecorate %39 DescriptorSet 0
|
||||
OpDecorate %39 Binding 7
|
||||
OpDecorate %41 NonReadable
|
||||
OpDecorate %41 DescriptorSet 0
|
||||
OpDecorate %41 Binding 8
|
||||
OpDecorate %43 NonReadable
|
||||
OpDecorate %43 DescriptorSet 0
|
||||
OpDecorate %43 Binding 9
|
||||
OpDecorate %45 NonReadable
|
||||
OpDecorate %45 DescriptorSet 0
|
||||
OpDecorate %45 Binding 10
|
||||
OpDecorate %47 NonReadable
|
||||
OpDecorate %47 DescriptorSet 0
|
||||
OpDecorate %47 Binding 11
|
||||
OpDecorate %250 Location 0
|
||||
%2 = OpTypeVoid
|
||||
%4 = OpTypeFloat 32
|
||||
%3 = OpTypeImage %4 1D 0 0 0 1 Unknown
|
||||
%5 = OpTypeInt 32 1
|
||||
%6 = OpTypeVector %4 4
|
||||
%7 = OpTypeImage %4 2D 0 0 0 1 Unknown
|
||||
%8 = OpTypeVector %5 2
|
||||
%9 = OpTypeImage %4 2D 0 1 0 1 Unknown
|
||||
%10 = OpTypeImage %4 3D 0 0 0 1 Unknown
|
||||
%11 = OpTypeVector %5 3
|
||||
%12 = OpTypeImage %4 2D 0 0 1 1 Unknown
|
||||
%13 = OpTypeImage %4 2D 1 0 0 1 Unknown
|
||||
%14 = OpTypeImage %4 2D 1 1 0 1 Unknown
|
||||
%15 = OpTypeImage %4 2D 1 0 1 1 Unknown
|
||||
%16 = OpTypeImage %4 1D 0 0 0 2 Rgba8
|
||||
%17 = OpTypeImage %4 2D 0 0 0 2 Rgba8
|
||||
%18 = OpTypeImage %4 2D 0 1 0 2 Rgba8
|
||||
%19 = OpTypeImage %4 3D 0 0 0 2 Rgba8
|
||||
%21 = OpTypePointer UniformConstant %3
|
||||
%20 = OpVariable %21 UniformConstant
|
||||
%23 = OpTypePointer UniformConstant %7
|
||||
%22 = OpVariable %23 UniformConstant
|
||||
%25 = OpTypePointer UniformConstant %9
|
||||
%24 = OpVariable %25 UniformConstant
|
||||
%27 = OpTypePointer UniformConstant %10
|
||||
%26 = OpVariable %27 UniformConstant
|
||||
%29 = OpTypePointer UniformConstant %12
|
||||
%28 = OpVariable %29 UniformConstant
|
||||
%31 = OpTypePointer UniformConstant %13
|
||||
%30 = OpVariable %31 UniformConstant
|
||||
%33 = OpTypePointer UniformConstant %14
|
||||
%32 = OpVariable %33 UniformConstant
|
||||
%35 = OpTypePointer UniformConstant %15
|
||||
%34 = OpVariable %35 UniformConstant
|
||||
%37 = OpTypePointer UniformConstant %16
|
||||
%36 = OpVariable %37 UniformConstant
|
||||
%39 = OpTypePointer UniformConstant %17
|
||||
%38 = OpVariable %39 UniformConstant
|
||||
%41 = OpTypePointer UniformConstant %18
|
||||
%40 = OpVariable %41 UniformConstant
|
||||
%43 = OpTypePointer UniformConstant %19
|
||||
%42 = OpVariable %43 UniformConstant
|
||||
%48 = OpTypeFunction %6 %5 %5
|
||||
%51 = OpTypeBool
|
||||
%52 = OpConstantNull %6
|
||||
%66 = OpTypeFunction %6 %8 %5
|
||||
%69 = OpConstantNull %6
|
||||
%75 = OpTypeVector %51 2
|
||||
%86 = OpTypeFunction %6 %8 %5 %5
|
||||
%90 = OpConstantNull %6
|
||||
%96 = OpTypeVector %51 3
|
||||
%106 = OpTypeFunction %6 %11 %5
|
||||
%109 = OpConstantNull %6
|
||||
%126 = OpConstantNull %6
|
||||
%141 = OpTypeFunction %4 %8 %5
|
||||
%144 = OpConstantNull %6
|
||||
%161 = OpTypeFunction %4 %8 %5 %5
|
||||
%165 = OpConstantNull %6
|
||||
%183 = OpConstantNull %6
|
||||
%199 = OpTypeFunction %2 %5 %6
|
||||
%210 = OpTypeFunction %2 %8 %6
|
||||
%223 = OpTypeFunction %2 %8 %5 %6
|
||||
%236 = OpTypeFunction %2 %11 %6
|
||||
%47 = OpFunction %6 None %48
|
||||
%45 = OpFunctionParameter %5
|
||||
%46 = OpFunctionParameter %5
|
||||
%44 = OpLabel
|
||||
%49 = OpLoad %3 %20
|
||||
OpBranch %50
|
||||
%50 = OpLabel
|
||||
%53 = OpImageQueryLevels %5 %49
|
||||
%54 = OpULessThan %51 %46 %53
|
||||
OpSelectionMerge %55 None
|
||||
OpBranchConditional %54 %56 %55
|
||||
%56 = OpLabel
|
||||
%57 = OpImageQuerySizeLod %5 %49 %46
|
||||
%58 = OpULessThan %51 %45 %57
|
||||
OpBranchConditional %58 %59 %55
|
||||
%59 = OpLabel
|
||||
%60 = OpImageFetch %6 %49 %45 Lod %46
|
||||
%4 = OpTypeInt 32 1
|
||||
%3 = OpConstant %4 0
|
||||
%6 = OpTypeFloat 32
|
||||
%5 = OpConstant %6 0.0
|
||||
%7 = OpTypeImage %6 1D 0 0 0 1 Unknown
|
||||
%8 = OpTypeVector %6 4
|
||||
%9 = OpTypeImage %6 2D 0 0 0 1 Unknown
|
||||
%10 = OpTypeVector %4 2
|
||||
%11 = OpTypeImage %6 2D 0 1 0 1 Unknown
|
||||
%12 = OpTypeImage %6 3D 0 0 0 1 Unknown
|
||||
%13 = OpTypeVector %4 3
|
||||
%14 = OpTypeImage %6 2D 0 0 1 1 Unknown
|
||||
%15 = OpTypeImage %6 2D 1 0 0 1 Unknown
|
||||
%16 = OpTypeImage %6 2D 1 1 0 1 Unknown
|
||||
%17 = OpTypeImage %6 2D 1 0 1 1 Unknown
|
||||
%18 = OpTypeImage %6 1D 0 0 0 2 Rgba8
|
||||
%19 = OpTypeImage %6 2D 0 0 0 2 Rgba8
|
||||
%20 = OpTypeImage %6 2D 0 1 0 2 Rgba8
|
||||
%21 = OpTypeImage %6 3D 0 0 0 2 Rgba8
|
||||
%22 = OpConstantComposite %10 %3 %3
|
||||
%23 = OpConstantComposite %13 %3 %3 %3
|
||||
%24 = OpConstantComposite %8 %5 %5 %5 %5
|
||||
%26 = OpTypePointer UniformConstant %7
|
||||
%25 = OpVariable %26 UniformConstant
|
||||
%28 = OpTypePointer UniformConstant %9
|
||||
%27 = OpVariable %28 UniformConstant
|
||||
%30 = OpTypePointer UniformConstant %11
|
||||
%29 = OpVariable %30 UniformConstant
|
||||
%32 = OpTypePointer UniformConstant %12
|
||||
%31 = OpVariable %32 UniformConstant
|
||||
%34 = OpTypePointer UniformConstant %14
|
||||
%33 = OpVariable %34 UniformConstant
|
||||
%36 = OpTypePointer UniformConstant %15
|
||||
%35 = OpVariable %36 UniformConstant
|
||||
%38 = OpTypePointer UniformConstant %16
|
||||
%37 = OpVariable %38 UniformConstant
|
||||
%40 = OpTypePointer UniformConstant %17
|
||||
%39 = OpVariable %40 UniformConstant
|
||||
%42 = OpTypePointer UniformConstant %18
|
||||
%41 = OpVariable %42 UniformConstant
|
||||
%44 = OpTypePointer UniformConstant %19
|
||||
%43 = OpVariable %44 UniformConstant
|
||||
%46 = OpTypePointer UniformConstant %20
|
||||
%45 = OpVariable %46 UniformConstant
|
||||
%48 = OpTypePointer UniformConstant %21
|
||||
%47 = OpVariable %48 UniformConstant
|
||||
%53 = OpTypeFunction %8 %4 %4
|
||||
%56 = OpTypeBool
|
||||
%57 = OpConstantNull %8
|
||||
%71 = OpTypeFunction %8 %10 %4
|
||||
%74 = OpConstantNull %8
|
||||
%80 = OpTypeVector %56 2
|
||||
%91 = OpTypeFunction %8 %10 %4 %4
|
||||
%95 = OpConstantNull %8
|
||||
%101 = OpTypeVector %56 3
|
||||
%111 = OpTypeFunction %8 %13 %4
|
||||
%114 = OpConstantNull %8
|
||||
%131 = OpConstantNull %8
|
||||
%146 = OpTypeFunction %6 %10 %4
|
||||
%149 = OpConstantNull %8
|
||||
%166 = OpTypeFunction %6 %10 %4 %4
|
||||
%170 = OpConstantNull %8
|
||||
%188 = OpConstantNull %8
|
||||
%204 = OpTypeFunction %2 %4 %8
|
||||
%215 = OpTypeFunction %2 %10 %8
|
||||
%228 = OpTypeFunction %2 %10 %4 %8
|
||||
%241 = OpTypeFunction %2 %13 %8
|
||||
%251 = OpTypePointer Output %8
|
||||
%250 = OpVariable %251 Output
|
||||
%253 = OpTypeFunction %2
|
||||
%52 = OpFunction %8 None %53
|
||||
%50 = OpFunctionParameter %4
|
||||
%51 = OpFunctionParameter %4
|
||||
%49 = OpLabel
|
||||
%54 = OpLoad %7 %25
|
||||
OpBranch %55
|
||||
%55 = OpLabel
|
||||
%61 = OpPhi %6 %52 %50 %52 %56 %60 %59
|
||||
OpReturnValue %61
|
||||
%58 = OpImageQueryLevels %4 %54
|
||||
%59 = OpULessThan %56 %51 %58
|
||||
OpSelectionMerge %60 None
|
||||
OpBranchConditional %59 %61 %60
|
||||
%61 = OpLabel
|
||||
%62 = OpImageQuerySizeLod %4 %54 %51
|
||||
%63 = OpULessThan %56 %50 %62
|
||||
OpBranchConditional %63 %64 %60
|
||||
%64 = OpLabel
|
||||
%65 = OpImageFetch %8 %54 %50 Lod %51
|
||||
OpBranch %60
|
||||
%60 = OpLabel
|
||||
%66 = OpPhi %8 %57 %55 %57 %61 %65 %64
|
||||
OpReturnValue %66
|
||||
OpFunctionEnd
|
||||
%65 = OpFunction %6 None %66
|
||||
%63 = OpFunctionParameter %8
|
||||
%64 = OpFunctionParameter %5
|
||||
%62 = OpLabel
|
||||
%67 = OpLoad %7 %22
|
||||
OpBranch %68
|
||||
%68 = OpLabel
|
||||
%70 = OpImageQueryLevels %5 %67
|
||||
%71 = OpULessThan %51 %64 %70
|
||||
OpSelectionMerge %72 None
|
||||
OpBranchConditional %71 %73 %72
|
||||
%70 = OpFunction %8 None %71
|
||||
%68 = OpFunctionParameter %10
|
||||
%69 = OpFunctionParameter %4
|
||||
%67 = OpLabel
|
||||
%72 = OpLoad %9 %27
|
||||
OpBranch %73
|
||||
%73 = OpLabel
|
||||
%74 = OpImageQuerySizeLod %8 %67 %64
|
||||
%76 = OpULessThan %75 %63 %74
|
||||
%77 = OpAll %51 %76
|
||||
OpBranchConditional %77 %78 %72
|
||||
%75 = OpImageQueryLevels %4 %72
|
||||
%76 = OpULessThan %56 %69 %75
|
||||
OpSelectionMerge %77 None
|
||||
OpBranchConditional %76 %78 %77
|
||||
%78 = OpLabel
|
||||
%79 = OpImageFetch %6 %67 %63 Lod %64
|
||||
OpBranch %72
|
||||
%72 = OpLabel
|
||||
%80 = OpPhi %6 %69 %68 %69 %73 %79 %78
|
||||
OpReturnValue %80
|
||||
%79 = OpImageQuerySizeLod %10 %72 %69
|
||||
%81 = OpULessThan %80 %68 %79
|
||||
%82 = OpAll %56 %81
|
||||
OpBranchConditional %82 %83 %77
|
||||
%83 = OpLabel
|
||||
%84 = OpImageFetch %8 %72 %68 Lod %69
|
||||
OpBranch %77
|
||||
%77 = OpLabel
|
||||
%85 = OpPhi %8 %74 %73 %74 %78 %84 %83
|
||||
OpReturnValue %85
|
||||
OpFunctionEnd
|
||||
%85 = OpFunction %6 None %86
|
||||
%82 = OpFunctionParameter %8
|
||||
%83 = OpFunctionParameter %5
|
||||
%84 = OpFunctionParameter %5
|
||||
%81 = OpLabel
|
||||
%87 = OpLoad %9 %24
|
||||
OpBranch %88
|
||||
%88 = OpLabel
|
||||
%89 = OpCompositeConstruct %11 %82 %83
|
||||
%91 = OpImageQueryLevels %5 %87
|
||||
%92 = OpULessThan %51 %84 %91
|
||||
OpSelectionMerge %93 None
|
||||
OpBranchConditional %92 %94 %93
|
||||
%94 = OpLabel
|
||||
%95 = OpImageQuerySizeLod %11 %87 %84
|
||||
%97 = OpULessThan %96 %89 %95
|
||||
%98 = OpAll %51 %97
|
||||
OpBranchConditional %98 %99 %93
|
||||
%99 = OpLabel
|
||||
%100 = OpImageFetch %6 %87 %89 Lod %84
|
||||
%90 = OpFunction %8 None %91
|
||||
%87 = OpFunctionParameter %10
|
||||
%88 = OpFunctionParameter %4
|
||||
%89 = OpFunctionParameter %4
|
||||
%86 = OpLabel
|
||||
%92 = OpLoad %11 %29
|
||||
OpBranch %93
|
||||
%93 = OpLabel
|
||||
%101 = OpPhi %6 %90 %88 %90 %94 %100 %99
|
||||
OpReturnValue %101
|
||||
%94 = OpCompositeConstruct %13 %87 %88
|
||||
%96 = OpImageQueryLevels %4 %92
|
||||
%97 = OpULessThan %56 %89 %96
|
||||
OpSelectionMerge %98 None
|
||||
OpBranchConditional %97 %99 %98
|
||||
%99 = OpLabel
|
||||
%100 = OpImageQuerySizeLod %13 %92 %89
|
||||
%102 = OpULessThan %101 %94 %100
|
||||
%103 = OpAll %56 %102
|
||||
OpBranchConditional %103 %104 %98
|
||||
%104 = OpLabel
|
||||
%105 = OpImageFetch %8 %92 %94 Lod %89
|
||||
OpBranch %98
|
||||
%98 = OpLabel
|
||||
%106 = OpPhi %8 %95 %93 %95 %99 %105 %104
|
||||
OpReturnValue %106
|
||||
OpFunctionEnd
|
||||
%105 = OpFunction %6 None %106
|
||||
%103 = OpFunctionParameter %11
|
||||
%104 = OpFunctionParameter %5
|
||||
%102 = OpLabel
|
||||
%107 = OpLoad %10 %26
|
||||
OpBranch %108
|
||||
%108 = OpLabel
|
||||
%110 = OpImageQueryLevels %5 %107
|
||||
%111 = OpULessThan %51 %104 %110
|
||||
OpSelectionMerge %112 None
|
||||
OpBranchConditional %111 %113 %112
|
||||
%110 = OpFunction %8 None %111
|
||||
%108 = OpFunctionParameter %13
|
||||
%109 = OpFunctionParameter %4
|
||||
%107 = OpLabel
|
||||
%112 = OpLoad %12 %31
|
||||
OpBranch %113
|
||||
%113 = OpLabel
|
||||
%114 = OpImageQuerySizeLod %11 %107 %104
|
||||
%115 = OpULessThan %96 %103 %114
|
||||
%116 = OpAll %51 %115
|
||||
OpBranchConditional %116 %117 %112
|
||||
%115 = OpImageQueryLevels %4 %112
|
||||
%116 = OpULessThan %56 %109 %115
|
||||
OpSelectionMerge %117 None
|
||||
OpBranchConditional %116 %118 %117
|
||||
%118 = OpLabel
|
||||
%119 = OpImageQuerySizeLod %13 %112 %109
|
||||
%120 = OpULessThan %101 %108 %119
|
||||
%121 = OpAll %56 %120
|
||||
OpBranchConditional %121 %122 %117
|
||||
%122 = OpLabel
|
||||
%123 = OpImageFetch %8 %112 %108 Lod %109
|
||||
OpBranch %117
|
||||
%117 = OpLabel
|
||||
%118 = OpImageFetch %6 %107 %103 Lod %104
|
||||
OpBranch %112
|
||||
%112 = OpLabel
|
||||
%119 = OpPhi %6 %109 %108 %109 %113 %118 %117
|
||||
OpReturnValue %119
|
||||
%124 = OpPhi %8 %114 %113 %114 %118 %123 %122
|
||||
OpReturnValue %124
|
||||
OpFunctionEnd
|
||||
%123 = OpFunction %6 None %66
|
||||
%121 = OpFunctionParameter %8
|
||||
%122 = OpFunctionParameter %5
|
||||
%120 = OpLabel
|
||||
%124 = OpLoad %12 %28
|
||||
OpBranch %125
|
||||
%128 = OpFunction %8 None %71
|
||||
%126 = OpFunctionParameter %10
|
||||
%127 = OpFunctionParameter %4
|
||||
%125 = OpLabel
|
||||
%127 = OpImageQuerySamples %5 %124
|
||||
%128 = OpULessThan %51 %122 %127
|
||||
OpSelectionMerge %129 None
|
||||
OpBranchConditional %128 %130 %129
|
||||
%129 = OpLoad %14 %33
|
||||
OpBranch %130
|
||||
%130 = OpLabel
|
||||
%131 = OpImageQuerySize %8 %124
|
||||
%132 = OpULessThan %75 %121 %131
|
||||
%133 = OpAll %51 %132
|
||||
OpBranchConditional %133 %134 %129
|
||||
%132 = OpImageQuerySamples %4 %129
|
||||
%133 = OpULessThan %56 %127 %132
|
||||
OpSelectionMerge %134 None
|
||||
OpBranchConditional %133 %135 %134
|
||||
%135 = OpLabel
|
||||
%136 = OpImageQuerySize %10 %129
|
||||
%137 = OpULessThan %80 %126 %136
|
||||
%138 = OpAll %56 %137
|
||||
OpBranchConditional %138 %139 %134
|
||||
%139 = OpLabel
|
||||
%140 = OpImageFetch %8 %129 %126 Sample %127
|
||||
OpBranch %134
|
||||
%134 = OpLabel
|
||||
%135 = OpImageFetch %6 %124 %121 Sample %122
|
||||
OpBranch %129
|
||||
%129 = OpLabel
|
||||
%136 = OpPhi %6 %126 %125 %126 %130 %135 %134
|
||||
OpReturnValue %136
|
||||
%141 = OpPhi %8 %131 %130 %131 %135 %140 %139
|
||||
OpReturnValue %141
|
||||
OpFunctionEnd
|
||||
%140 = OpFunction %4 None %141
|
||||
%138 = OpFunctionParameter %8
|
||||
%139 = OpFunctionParameter %5
|
||||
%137 = OpLabel
|
||||
%142 = OpLoad %13 %30
|
||||
OpBranch %143
|
||||
%143 = OpLabel
|
||||
%145 = OpImageQueryLevels %5 %142
|
||||
%146 = OpULessThan %51 %139 %145
|
||||
OpSelectionMerge %147 None
|
||||
OpBranchConditional %146 %148 %147
|
||||
%145 = OpFunction %6 None %146
|
||||
%143 = OpFunctionParameter %10
|
||||
%144 = OpFunctionParameter %4
|
||||
%142 = OpLabel
|
||||
%147 = OpLoad %15 %35
|
||||
OpBranch %148
|
||||
%148 = OpLabel
|
||||
%149 = OpImageQuerySizeLod %8 %142 %139
|
||||
%150 = OpULessThan %75 %138 %149
|
||||
%151 = OpAll %51 %150
|
||||
OpBranchConditional %151 %152 %147
|
||||
%150 = OpImageQueryLevels %4 %147
|
||||
%151 = OpULessThan %56 %144 %150
|
||||
OpSelectionMerge %152 None
|
||||
OpBranchConditional %151 %153 %152
|
||||
%153 = OpLabel
|
||||
%154 = OpImageQuerySizeLod %10 %147 %144
|
||||
%155 = OpULessThan %80 %143 %154
|
||||
%156 = OpAll %56 %155
|
||||
OpBranchConditional %156 %157 %152
|
||||
%157 = OpLabel
|
||||
%158 = OpImageFetch %8 %147 %143 Lod %144
|
||||
OpBranch %152
|
||||
%152 = OpLabel
|
||||
%153 = OpImageFetch %6 %142 %138 Lod %139
|
||||
OpBranch %147
|
||||
%147 = OpLabel
|
||||
%154 = OpPhi %6 %144 %143 %144 %148 %153 %152
|
||||
%155 = OpCompositeExtract %4 %154 0
|
||||
OpReturnValue %155
|
||||
%159 = OpPhi %8 %149 %148 %149 %153 %158 %157
|
||||
%160 = OpCompositeExtract %6 %159 0
|
||||
OpReturnValue %160
|
||||
OpFunctionEnd
|
||||
%160 = OpFunction %4 None %161
|
||||
%157 = OpFunctionParameter %8
|
||||
%158 = OpFunctionParameter %5
|
||||
%159 = OpFunctionParameter %5
|
||||
%156 = OpLabel
|
||||
%162 = OpLoad %14 %32
|
||||
OpBranch %163
|
||||
%163 = OpLabel
|
||||
%164 = OpCompositeConstruct %11 %157 %158
|
||||
%166 = OpImageQueryLevels %5 %162
|
||||
%167 = OpULessThan %51 %159 %166
|
||||
OpSelectionMerge %168 None
|
||||
OpBranchConditional %167 %169 %168
|
||||
%169 = OpLabel
|
||||
%170 = OpImageQuerySizeLod %11 %162 %159
|
||||
%171 = OpULessThan %96 %164 %170
|
||||
%172 = OpAll %51 %171
|
||||
OpBranchConditional %172 %173 %168
|
||||
%173 = OpLabel
|
||||
%174 = OpImageFetch %6 %162 %164 Lod %159
|
||||
%165 = OpFunction %6 None %166
|
||||
%162 = OpFunctionParameter %10
|
||||
%163 = OpFunctionParameter %4
|
||||
%164 = OpFunctionParameter %4
|
||||
%161 = OpLabel
|
||||
%167 = OpLoad %16 %37
|
||||
OpBranch %168
|
||||
%168 = OpLabel
|
||||
%175 = OpPhi %6 %165 %163 %165 %169 %174 %173
|
||||
%176 = OpCompositeExtract %4 %175 0
|
||||
OpReturnValue %176
|
||||
%169 = OpCompositeConstruct %13 %162 %163
|
||||
%171 = OpImageQueryLevels %4 %167
|
||||
%172 = OpULessThan %56 %164 %171
|
||||
OpSelectionMerge %173 None
|
||||
OpBranchConditional %172 %174 %173
|
||||
%174 = OpLabel
|
||||
%175 = OpImageQuerySizeLod %13 %167 %164
|
||||
%176 = OpULessThan %101 %169 %175
|
||||
%177 = OpAll %56 %176
|
||||
OpBranchConditional %177 %178 %173
|
||||
%178 = OpLabel
|
||||
%179 = OpImageFetch %8 %167 %169 Lod %164
|
||||
OpBranch %173
|
||||
%173 = OpLabel
|
||||
%180 = OpPhi %8 %170 %168 %170 %174 %179 %178
|
||||
%181 = OpCompositeExtract %6 %180 0
|
||||
OpReturnValue %181
|
||||
OpFunctionEnd
|
||||
%180 = OpFunction %4 None %141
|
||||
%178 = OpFunctionParameter %8
|
||||
%179 = OpFunctionParameter %5
|
||||
%177 = OpLabel
|
||||
%181 = OpLoad %15 %34
|
||||
OpBranch %182
|
||||
%185 = OpFunction %6 None %146
|
||||
%183 = OpFunctionParameter %10
|
||||
%184 = OpFunctionParameter %4
|
||||
%182 = OpLabel
|
||||
%184 = OpImageQuerySamples %5 %181
|
||||
%185 = OpULessThan %51 %179 %184
|
||||
OpSelectionMerge %186 None
|
||||
OpBranchConditional %185 %187 %186
|
||||
%186 = OpLoad %17 %39
|
||||
OpBranch %187
|
||||
%187 = OpLabel
|
||||
%188 = OpImageQuerySize %8 %181
|
||||
%189 = OpULessThan %75 %178 %188
|
||||
%190 = OpAll %51 %189
|
||||
OpBranchConditional %190 %191 %186
|
||||
%189 = OpImageQuerySamples %4 %186
|
||||
%190 = OpULessThan %56 %184 %189
|
||||
OpSelectionMerge %191 None
|
||||
OpBranchConditional %190 %192 %191
|
||||
%192 = OpLabel
|
||||
%193 = OpImageQuerySize %10 %186
|
||||
%194 = OpULessThan %80 %183 %193
|
||||
%195 = OpAll %56 %194
|
||||
OpBranchConditional %195 %196 %191
|
||||
%196 = OpLabel
|
||||
%197 = OpImageFetch %8 %186 %183 Sample %184
|
||||
OpBranch %191
|
||||
%191 = OpLabel
|
||||
%192 = OpImageFetch %6 %181 %178 Sample %179
|
||||
OpBranch %186
|
||||
%186 = OpLabel
|
||||
%193 = OpPhi %6 %183 %182 %183 %187 %192 %191
|
||||
%194 = OpCompositeExtract %4 %193 0
|
||||
OpReturnValue %194
|
||||
%198 = OpPhi %8 %188 %187 %188 %192 %197 %196
|
||||
%199 = OpCompositeExtract %6 %198 0
|
||||
OpReturnValue %199
|
||||
OpFunctionEnd
|
||||
%198 = OpFunction %2 None %199
|
||||
%196 = OpFunctionParameter %5
|
||||
%197 = OpFunctionParameter %6
|
||||
%195 = OpLabel
|
||||
%200 = OpLoad %16 %36
|
||||
OpBranch %201
|
||||
%201 = OpLabel
|
||||
%202 = OpImageQuerySize %5 %200
|
||||
%203 = OpULessThan %51 %196 %202
|
||||
OpSelectionMerge %204 None
|
||||
OpBranchConditional %203 %205 %204
|
||||
%205 = OpLabel
|
||||
OpImageWrite %200 %196 %197
|
||||
OpBranch %204
|
||||
%204 = OpLabel
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%209 = OpFunction %2 None %210
|
||||
%207 = OpFunctionParameter %8
|
||||
%208 = OpFunctionParameter %6
|
||||
%203 = OpFunction %2 None %204
|
||||
%201 = OpFunctionParameter %4
|
||||
%202 = OpFunctionParameter %8
|
||||
%200 = OpLabel
|
||||
%205 = OpLoad %18 %41
|
||||
OpBranch %206
|
||||
%206 = OpLabel
|
||||
%211 = OpLoad %17 %38
|
||||
OpBranch %212
|
||||
%212 = OpLabel
|
||||
%213 = OpImageQuerySize %8 %211
|
||||
%214 = OpULessThan %75 %207 %213
|
||||
%215 = OpAll %51 %214
|
||||
OpSelectionMerge %216 None
|
||||
OpBranchConditional %215 %217 %216
|
||||
%217 = OpLabel
|
||||
OpImageWrite %211 %207 %208
|
||||
OpBranch %216
|
||||
%216 = OpLabel
|
||||
%207 = OpImageQuerySize %4 %205
|
||||
%208 = OpULessThan %56 %201 %207
|
||||
OpSelectionMerge %209 None
|
||||
OpBranchConditional %208 %210 %209
|
||||
%210 = OpLabel
|
||||
OpImageWrite %205 %201 %202
|
||||
OpBranch %209
|
||||
%209 = OpLabel
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%222 = OpFunction %2 None %223
|
||||
%219 = OpFunctionParameter %8
|
||||
%220 = OpFunctionParameter %5
|
||||
%221 = OpFunctionParameter %6
|
||||
%218 = OpLabel
|
||||
%224 = OpLoad %18 %40
|
||||
OpBranch %225
|
||||
%225 = OpLabel
|
||||
%226 = OpCompositeConstruct %11 %219 %220
|
||||
%227 = OpImageQuerySize %11 %224
|
||||
%228 = OpULessThan %96 %226 %227
|
||||
%229 = OpAll %51 %228
|
||||
OpSelectionMerge %230 None
|
||||
OpBranchConditional %229 %231 %230
|
||||
%231 = OpLabel
|
||||
OpImageWrite %224 %226 %221
|
||||
%214 = OpFunction %2 None %215
|
||||
%212 = OpFunctionParameter %10
|
||||
%213 = OpFunctionParameter %8
|
||||
%211 = OpLabel
|
||||
%216 = OpLoad %19 %43
|
||||
OpBranch %217
|
||||
%217 = OpLabel
|
||||
%218 = OpImageQuerySize %10 %216
|
||||
%219 = OpULessThan %80 %212 %218
|
||||
%220 = OpAll %56 %219
|
||||
OpSelectionMerge %221 None
|
||||
OpBranchConditional %220 %222 %221
|
||||
%222 = OpLabel
|
||||
OpImageWrite %216 %212 %213
|
||||
OpBranch %221
|
||||
%221 = OpLabel
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%227 = OpFunction %2 None %228
|
||||
%224 = OpFunctionParameter %10
|
||||
%225 = OpFunctionParameter %4
|
||||
%226 = OpFunctionParameter %8
|
||||
%223 = OpLabel
|
||||
%229 = OpLoad %20 %45
|
||||
OpBranch %230
|
||||
%230 = OpLabel
|
||||
%231 = OpCompositeConstruct %13 %224 %225
|
||||
%232 = OpImageQuerySize %13 %229
|
||||
%233 = OpULessThan %101 %231 %232
|
||||
%234 = OpAll %56 %233
|
||||
OpSelectionMerge %235 None
|
||||
OpBranchConditional %234 %236 %235
|
||||
%236 = OpLabel
|
||||
OpImageWrite %229 %231 %226
|
||||
OpBranch %235
|
||||
%235 = OpLabel
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%235 = OpFunction %2 None %236
|
||||
%233 = OpFunctionParameter %11
|
||||
%234 = OpFunctionParameter %6
|
||||
%232 = OpLabel
|
||||
%237 = OpLoad %19 %42
|
||||
OpBranch %238
|
||||
%238 = OpLabel
|
||||
%239 = OpImageQuerySize %11 %237
|
||||
%240 = OpULessThan %96 %233 %239
|
||||
%241 = OpAll %51 %240
|
||||
OpSelectionMerge %242 None
|
||||
OpBranchConditional %241 %243 %242
|
||||
%240 = OpFunction %2 None %241
|
||||
%238 = OpFunctionParameter %13
|
||||
%239 = OpFunctionParameter %8
|
||||
%237 = OpLabel
|
||||
%242 = OpLoad %21 %47
|
||||
OpBranch %243
|
||||
%243 = OpLabel
|
||||
OpImageWrite %237 %233 %234
|
||||
OpBranch %242
|
||||
%242 = OpLabel
|
||||
%244 = OpImageQuerySize %13 %242
|
||||
%245 = OpULessThan %101 %238 %244
|
||||
%246 = OpAll %56 %245
|
||||
OpSelectionMerge %247 None
|
||||
OpBranchConditional %246 %248 %247
|
||||
%248 = OpLabel
|
||||
OpImageWrite %242 %238 %239
|
||||
OpBranch %247
|
||||
%247 = OpLabel
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%252 = OpFunction %2 None %253
|
||||
%249 = OpLabel
|
||||
%254 = OpLoad %7 %25
|
||||
%255 = OpLoad %9 %27
|
||||
%256 = OpLoad %11 %29
|
||||
%257 = OpLoad %12 %31
|
||||
%258 = OpLoad %14 %33
|
||||
%259 = OpLoad %18 %41
|
||||
%260 = OpLoad %19 %43
|
||||
%261 = OpLoad %20 %45
|
||||
%262 = OpLoad %21 %47
|
||||
OpBranch %263
|
||||
%263 = OpLabel
|
||||
%264 = OpFunctionCall %8 %52 %3 %3
|
||||
%265 = OpFunctionCall %8 %70 %22 %3
|
||||
%266 = OpFunctionCall %8 %90 %22 %3 %3
|
||||
%267 = OpFunctionCall %8 %110 %23 %3
|
||||
%268 = OpFunctionCall %8 %128 %22 %3
|
||||
%269 = OpFunctionCall %2 %203 %3 %24
|
||||
%270 = OpFunctionCall %2 %214 %22 %24
|
||||
%271 = OpFunctionCall %2 %227 %22 %3 %24
|
||||
%272 = OpFunctionCall %2 %240 %23 %24
|
||||
%273 = OpCompositeConstruct %8 %5 %5 %5 %5
|
||||
OpStore %250 %273
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@ -153,7 +153,16 @@ fn check_targets(module: &naga::Module, name: &str, targets: Targets) {
|
||||
if params.glsl_exclude_list.contains(&ep.name) {
|
||||
continue;
|
||||
}
|
||||
write_output_glsl(module, &info, &dest, name, ep.stage, &ep.name, ¶ms.glsl);
|
||||
write_output_glsl(
|
||||
module,
|
||||
&info,
|
||||
&dest,
|
||||
name,
|
||||
ep.stage,
|
||||
&ep.name,
|
||||
¶ms.glsl,
|
||||
params.bounds_check_policies,
|
||||
);
|
||||
}
|
||||
}
|
||||
}
|
||||
@ -273,6 +282,7 @@ fn write_output_glsl(
|
||||
stage: naga::ShaderStage,
|
||||
ep_name: &str,
|
||||
options: &naga::back::glsl::Options,
|
||||
bounds_check_policies: naga::proc::BoundsCheckPolicies,
|
||||
) {
|
||||
use naga::back::glsl;
|
||||
|
||||
@ -284,8 +294,15 @@ fn write_output_glsl(
|
||||
};
|
||||
|
||||
let mut buffer = String::new();
|
||||
let mut writer = glsl::Writer::new(&mut buffer, module, info, options, &pipeline_options)
|
||||
.expect("GLSL init failed");
|
||||
let mut writer = glsl::Writer::new(
|
||||
&mut buffer,
|
||||
module,
|
||||
info,
|
||||
options,
|
||||
&pipeline_options,
|
||||
bounds_check_policies,
|
||||
)
|
||||
.expect("GLSL init failed");
|
||||
writer.write().expect("GLSL write failed");
|
||||
|
||||
fs::write(
|
||||
@ -489,9 +506,12 @@ fn convert_wgsl() {
|
||||
("bounds-check-restrict", Targets::SPIRV | Targets::METAL),
|
||||
(
|
||||
"bounds-check-image-restrict",
|
||||
Targets::SPIRV | Targets::METAL,
|
||||
Targets::SPIRV | Targets::METAL | Targets::GLSL,
|
||||
),
|
||||
(
|
||||
"bounds-check-image-rzsw",
|
||||
Targets::SPIRV | Targets::METAL | Targets::GLSL,
|
||||
),
|
||||
("bounds-check-image-rzsw", Targets::SPIRV | Targets::METAL),
|
||||
("policy-mix", Targets::SPIRV | Targets::METAL),
|
||||
(
|
||||
"texture-arg",
|
||||
|
||||
Loading…
Reference in New Issue
Block a user