From 0aa7681165f9ae74f2ddd109f4bdfa73a4e094cd Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Jo=C3=A3o=20Capucho?= Date: Mon, 30 May 2022 20:13:58 +0100 Subject: [PATCH] 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 --- cli/src/main.rs | 1 + src/back/glsl/features.rs | 83 +- src/back/glsl/mod.rs | 469 +++++++++-- .../in/bounds-check-image-restrict.param.ron | 5 + tests/in/bounds-check-image-restrict.wgsl | 43 +- tests/in/bounds-check-image-rzsw.param.ron | 5 + tests/in/bounds-check-image-rzsw.wgsl | 43 +- ...age-restrict.fragment_shader.Fragment.glsl | 86 ++ ...k-image-rzsw.fragment_shader.Fragment.glsl | 81 ++ tests/out/glsl/image.main.Compute.glsl | 4 +- tests/out/msl/bounds-check-image-restrict.msl | 29 + tests/out/msl/bounds-check-image-rzsw.msl | 29 + .../spv/bounds-check-image-restrict.spvasm | 665 ++++++++------- tests/out/spv/bounds-check-image-rzsw.spvasm | 773 +++++++++--------- tests/snapshots.rs | 30 +- 15 files changed, 1556 insertions(+), 790 deletions(-) create mode 100644 tests/out/glsl/bounds-check-image-restrict.fragment_shader.Fragment.glsl create mode 100644 tests/out/glsl/bounds-check-image-rzsw.fragment_shader.Fragment.glsl diff --git a/cli/src/main.rs b/cli/src/main.rs index 2e4b1b3c35..707f95f065 100644 --- a/cli/src/main.rs +++ b/cli/src/main.rs @@ -455,6 +455,7 @@ fn run() -> Result<(), Box> { ))?, ¶ms.glsl, &pipeline_options, + params.bounds_check_policies, ) .unwrap_pretty(); writer.write()?; diff --git a/src/back/glsl/features.rs b/src/back/glsl/features.rs index d3f117045f..82718b040e 100644 --- a/src/back/glsl/features.rs +++ b/src/back/glsl/features.rs @@ -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) diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index abce4d4abd..9e534df539 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -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; @@ -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 { // 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, - coordinate: Handle, - array_index: Option>, - dim: crate::ImageDimension, ctx: &back::FunctionCtx, + expr: Handle, + image: Handle, + level_expr: Handle, ) -> Result<(), Error> { - use crate::ImageDimension as IDim; + // 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);")?; - let tex_1d_hack = dim == IDim::D1 && self.options.version.is_es(); + 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, + array_index: Option>, + // Emulate 1D images as 2D for profiles that don't support it (glsl es) + tex_1d_hack: bool, + ) -> Result<(), Error> { 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, + coordinate: Handle, + array_index: Option>, + value: Handle, + ) -> 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, + ctx: &back::FunctionCtx, + image: Handle, + coordinate: Handle, + array_index: Option>, + sample: Option>, + level: Option>, + ) -> 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(()) } diff --git a/tests/in/bounds-check-image-restrict.param.ron b/tests/in/bounds-check-image-restrict.param.ron index d6af131dfb..e91f7fc24d 100644 --- a/tests/in/bounds-check-image-restrict.param.ron +++ b/tests/in/bounds-check-image-restrict.param.ron @@ -6,4 +6,9 @@ version: (1, 1), debug: true, ), + glsl: ( + version: Desktop(430), + writer_flags: (bits: 0), + binding_map: { }, + ), ) diff --git a/tests/in/bounds-check-image-restrict.wgsl b/tests/in/bounds-check-image-restrict.wgsl index 3be562ff25..0cce20eafe 100644 --- a/tests/in/bounds-check-image-restrict.wgsl +++ b/tests/in/bounds-check-image-restrict.wgsl @@ -5,79 +5,100 @@ fn test_textureLoad_1d(coords: i32, level: i32) -> vec4 { return textureLoad(image_1d, coords, level); } -@group(0) @binding(0) +@group(0) @binding(1) var image_2d: texture_2d; fn test_textureLoad_2d(coords: vec2, level: i32) -> vec4 { return textureLoad(image_2d, coords, level); } -@group(0) @binding(0) +@group(0) @binding(2) var image_2d_array: texture_2d_array; fn test_textureLoad_2d_array(coords: vec2, index: i32, level: i32) -> vec4 { return textureLoad(image_2d_array, coords, index, level); } -@group(0) @binding(0) +@group(0) @binding(3) var image_3d: texture_3d; fn test_textureLoad_3d(coords: vec3, level: i32) -> vec4 { return textureLoad(image_3d, coords, level); } -@group(0) @binding(0) +@group(0) @binding(4) var image_multisampled_2d: texture_multisampled_2d; fn test_textureLoad_multisampled_2d(coords: vec2, _sample: i32) -> vec4 { 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, 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, 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, _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; fn test_textureStore_1d(coords: i32, value: vec4) { textureStore(image_storage_1d, coords, value); } -@group(0) @binding(0) +@group(0) @binding(9) var image_storage_2d: texture_storage_2d; fn test_textureStore_2d(coords: vec2, value: vec4) { textureStore(image_storage_2d, coords, value); } -@group(0) @binding(0) +@group(0) @binding(10) var image_storage_2d_array: texture_storage_2d_array; fn test_textureStore_2d_array(coords: vec2, array_index: i32, value: vec4) { textureStore(image_storage_2d_array, coords, array_index, value); } -@group(0) @binding(0) +@group(0) @binding(11) var image_storage_3d: texture_storage_3d; fn test_textureStore_3d(coords: vec3, value: vec4) { 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 { + test_textureLoad_1d(0, 0); + test_textureLoad_2d(vec2(), 0); + test_textureLoad_2d_array(vec2(), 0, 0); + test_textureLoad_3d(vec3(), 0); + test_textureLoad_multisampled_2d(vec2(), 0); + // Not yet implemented for GLSL: + // test_textureLoad_depth_2d(vec2(), 0); + // test_textureLoad_depth_2d_array(vec2(), 0, 0); + // test_textureLoad_depth_multisampled_2d(vec2(), 0); + test_textureStore_1d(0, vec4()); + test_textureStore_2d(vec2(), vec4()); + test_textureStore_2d_array(vec2(), 0, vec4()); + test_textureStore_3d(vec3(), vec4()); + + return vec4(0.,0.,0.,0.); +} diff --git a/tests/in/bounds-check-image-rzsw.param.ron b/tests/in/bounds-check-image-rzsw.param.ron index 2d00b645c3..72ccb27e9d 100644 --- a/tests/in/bounds-check-image-rzsw.param.ron +++ b/tests/in/bounds-check-image-rzsw.param.ron @@ -6,4 +6,9 @@ version: (1, 1), debug: true, ), + glsl: ( + version: Desktop(430), + writer_flags: (bits: 0), + binding_map: { }, + ), ) diff --git a/tests/in/bounds-check-image-rzsw.wgsl b/tests/in/bounds-check-image-rzsw.wgsl index 3be562ff25..0cce20eafe 100644 --- a/tests/in/bounds-check-image-rzsw.wgsl +++ b/tests/in/bounds-check-image-rzsw.wgsl @@ -5,79 +5,100 @@ fn test_textureLoad_1d(coords: i32, level: i32) -> vec4 { return textureLoad(image_1d, coords, level); } -@group(0) @binding(0) +@group(0) @binding(1) var image_2d: texture_2d; fn test_textureLoad_2d(coords: vec2, level: i32) -> vec4 { return textureLoad(image_2d, coords, level); } -@group(0) @binding(0) +@group(0) @binding(2) var image_2d_array: texture_2d_array; fn test_textureLoad_2d_array(coords: vec2, index: i32, level: i32) -> vec4 { return textureLoad(image_2d_array, coords, index, level); } -@group(0) @binding(0) +@group(0) @binding(3) var image_3d: texture_3d; fn test_textureLoad_3d(coords: vec3, level: i32) -> vec4 { return textureLoad(image_3d, coords, level); } -@group(0) @binding(0) +@group(0) @binding(4) var image_multisampled_2d: texture_multisampled_2d; fn test_textureLoad_multisampled_2d(coords: vec2, _sample: i32) -> vec4 { 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, 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, 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, _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; fn test_textureStore_1d(coords: i32, value: vec4) { textureStore(image_storage_1d, coords, value); } -@group(0) @binding(0) +@group(0) @binding(9) var image_storage_2d: texture_storage_2d; fn test_textureStore_2d(coords: vec2, value: vec4) { textureStore(image_storage_2d, coords, value); } -@group(0) @binding(0) +@group(0) @binding(10) var image_storage_2d_array: texture_storage_2d_array; fn test_textureStore_2d_array(coords: vec2, array_index: i32, value: vec4) { textureStore(image_storage_2d_array, coords, array_index, value); } -@group(0) @binding(0) +@group(0) @binding(11) var image_storage_3d: texture_storage_3d; fn test_textureStore_3d(coords: vec3, value: vec4) { 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 { + test_textureLoad_1d(0, 0); + test_textureLoad_2d(vec2(), 0); + test_textureLoad_2d_array(vec2(), 0, 0); + test_textureLoad_3d(vec3(), 0); + test_textureLoad_multisampled_2d(vec2(), 0); + // Not yet implemented for GLSL: + // test_textureLoad_depth_2d(vec2(), 0); + // test_textureLoad_depth_2d_array(vec2(), 0, 0); + // test_textureLoad_depth_multisampled_2d(vec2(), 0); + test_textureStore_1d(0, vec4()); + test_textureStore_2d(vec2(), vec4()); + test_textureStore_2d_array(vec2(), 0, vec4()); + test_textureStore_3d(vec3(), vec4()); + + return vec4(0.,0.,0.,0.); +} diff --git a/tests/out/glsl/bounds-check-image-restrict.fragment_shader.Fragment.glsl b/tests/out/glsl/bounds-check-image-restrict.fragment_shader.Fragment.glsl new file mode 100644 index 0000000000..a1d8b5b055 --- /dev/null +++ b/tests/out/glsl/bounds-check-image-restrict.fragment_shader.Fragment.glsl @@ -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; +} + diff --git a/tests/out/glsl/bounds-check-image-rzsw.fragment_shader.Fragment.glsl b/tests/out/glsl/bounds-check-image-rzsw.fragment_shader.Fragment.glsl new file mode 100644 index 0000000000..9c2bef19ab --- /dev/null +++ b/tests/out/glsl/bounds-check-image-rzsw.fragment_shader.Fragment.glsl @@ -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; +} + diff --git a/tests/out/glsl/image.main.Compute.glsl b/tests/out/glsl/image.main.Compute.glsl index 08c6a32e96..96d4bb7753 100644 --- a/tests/out/glsl/image.main.Compute.glsl +++ b/tests/out/glsl/image.main.Compute.glsl @@ -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; } diff --git a/tests/out/msl/bounds-check-image-restrict.msl b/tests/out/msl/bounds-check-image-restrict.msl index 6a45484754..b2f6eca4d8 100644 --- a/tests/out/msl/bounds-check-image-restrict.msl +++ b/tests/out/msl/bounds-check-image-restrict.msl @@ -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 image_1d [[user(fake0)]] +, metal::texture2d image_2d [[user(fake0)]] +, metal::texture2d_array image_2d_array [[user(fake0)]] +, metal::texture3d image_3d [[user(fake0)]] +, metal::texture2d_ms image_multisampled_2d [[user(fake0)]] +, metal::texture1d image_storage_1d [[user(fake0)]] +, metal::texture2d image_storage_2d [[user(fake0)]] +, metal::texture2d_array image_storage_2d_array [[user(fake0)]] +, metal::texture3d 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) }; +} diff --git a/tests/out/msl/bounds-check-image-rzsw.msl b/tests/out/msl/bounds-check-image-rzsw.msl index 78bef45eed..16cbf5c0a6 100644 --- a/tests/out/msl/bounds-check-image-rzsw.msl +++ b/tests/out/msl/bounds-check-image-rzsw.msl @@ -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 image_1d [[user(fake0)]] +, metal::texture2d image_2d [[user(fake0)]] +, metal::texture2d_array image_2d_array [[user(fake0)]] +, metal::texture3d image_3d [[user(fake0)]] +, metal::texture2d_ms image_multisampled_2d [[user(fake0)]] +, metal::texture1d image_storage_1d [[user(fake0)]] +, metal::texture2d image_storage_2d [[user(fake0)]] +, metal::texture2d_array image_storage_2d_array [[user(fake0)]] +, metal::texture3d 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) }; +} diff --git a/tests/out/spv/bounds-check-image-restrict.spvasm b/tests/out/spv/bounds-check-image-restrict.spvasm index fe3bf1759f..51d53f3d03 100644 --- a/tests/out/spv/bounds-check-image-restrict.spvasm +++ b/tests/out/spv/bounds-check-image-restrict.spvasm @@ -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 +%187 = OpFunction %2 None %188 +%185 = OpFunctionParameter %10 +%186 = OpFunctionParameter %8 +%184 = OpLabel +%189 = OpLoad %19 %43 +OpBranch %190 +%190 = OpLabel +%191 = OpImageQuerySize %10 %189 +%193 = OpISub %10 %191 %192 +%194 = OpExtInst %10 %1 UMin %185 %193 +OpImageWrite %189 %194 %186 OpReturn OpFunctionEnd -%194 = OpFunction %2 None %195 -%191 = OpFunctionParameter %8 -%192 = OpFunctionParameter %5 -%193 = OpFunctionParameter %6 -%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 +%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 -%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 +%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 \ No newline at end of file diff --git a/tests/out/spv/bounds-check-image-rzsw.spvasm b/tests/out/spv/bounds-check-image-rzsw.spvasm index 981355a419..3c8e79009a 100644 --- a/tests/out/spv/bounds-check-image-rzsw.spvasm +++ b/tests/out/spv/bounds-check-image-rzsw.spvasm @@ -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 +%203 = OpFunction %2 None %204 +%201 = OpFunctionParameter %4 +%202 = OpFunctionParameter %8 +%200 = OpLabel +%205 = OpLoad %18 %41 +OpBranch %206 +%206 = 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 -%209 = OpFunction %2 None %210 -%207 = OpFunctionParameter %8 -%208 = OpFunctionParameter %6 -%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 +%214 = OpFunction %2 None %215 +%212 = OpFunctionParameter %10 +%213 = OpFunctionParameter %8 +%211 = OpLabel +%216 = OpLoad %19 %43 +OpBranch %217 %217 = OpLabel -OpImageWrite %211 %207 %208 -OpBranch %216 -%216 = 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 -%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 +%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 \ No newline at end of file diff --git a/tests/snapshots.rs b/tests/snapshots.rs index f6a3749c3c..459cdf2221 100644 --- a/tests/snapshots.rs +++ b/tests/snapshots.rs @@ -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",