diff --git a/naga/src/back/hlsl/help.rs b/naga/src/back/hlsl/help.rs index 00d64faa4dd..058d3969237 100644 --- a/naga/src/back/hlsl/help.rs +++ b/naga/src/back/hlsl/help.rs @@ -32,8 +32,8 @@ use core::fmt::Write; use super::{ super::FunctionCtx, writer::{ - ABS_FUNCTION, DIV_FUNCTION, EXTRACT_BITS_FUNCTION, INSERT_BITS_FUNCTION, MOD_FUNCTION, - NEG_FUNCTION, + ABS_FUNCTION, DIV_FUNCTION, EXTRACT_BITS_FUNCTION, F2I32_FUNCTION, F2I64_FUNCTION, + F2U32_FUNCTION, F2U64_FUNCTION, INSERT_BITS_FUNCTION, MOD_FUNCTION, NEG_FUNCTION, }, BackendResult, WrappedType, }; @@ -97,6 +97,15 @@ pub(super) struct WrappedBinaryOp { pub(super) right_ty: (Option, crate::Scalar), } +#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)] +pub(super) struct WrappedCast { + // This can only represent scalar or vector types. If we ever need to wrap + // casts with other types, we'll need a better representation. + pub(super) vector_size: Option, + pub(super) src_scalar: crate::Scalar, + pub(super) dst_scalar: crate::Scalar, +} + /// HLSL backend requires its own `ImageQuery` enum. /// /// It is used inside `WrappedImageQuery` and should be unique per ImageQuery function. @@ -1355,6 +1364,97 @@ impl super::Writer<'_, W> { Ok(()) } + fn write_wrapped_cast_functions( + &mut self, + module: &crate::Module, + func_ctx: &FunctionCtx, + ) -> BackendResult { + for (_, expression) in func_ctx.expressions.iter() { + if let crate::Expression::As { + expr, + kind, + convert: Some(width), + } = *expression + { + // Avoid undefined behaviour when casting from a float to integer + // when the value is out of range for the target type. Additionally + // ensure we clamp to the correct value as per the WGSL spec. + // + // https://www.w3.org/TR/WGSL/#floating-point-conversion: + // * If X is exactly representable in the target type T, then the + // result is that value. + // * Otherwise, the result is the value in T closest to + // truncate(X) and also exactly representable in the original + // floating point type. + let src_ty = func_ctx.resolve_type(expr, &module.types); + let Some((vector_size, src_scalar)) = src_ty.vector_size_and_scalar() else { + continue; + }; + let dst_scalar = crate::Scalar { kind, width }; + if src_scalar.kind != ScalarKind::Float + || (dst_scalar.kind != ScalarKind::Sint && dst_scalar.kind != ScalarKind::Uint) + { + continue; + } + + let wrapped = WrappedCast { + src_scalar, + vector_size, + dst_scalar, + }; + if !self.wrapped.insert(WrappedType::Cast(wrapped)) { + continue; + } + + let (src_ty, dst_ty) = match vector_size { + None => ( + crate::TypeInner::Scalar(src_scalar), + crate::TypeInner::Scalar(dst_scalar), + ), + Some(vector_size) => ( + crate::TypeInner::Vector { + scalar: src_scalar, + size: vector_size, + }, + crate::TypeInner::Vector { + scalar: dst_scalar, + size: vector_size, + }, + ), + }; + let (min, max) = + crate::proc::min_max_float_representable_by(src_scalar, dst_scalar); + let cast_str = format!( + "{}{}", + dst_scalar.to_hlsl_str()?, + vector_size + .map(crate::common::vector_size_str) + .unwrap_or(""), + ); + let fun_name = match dst_scalar { + crate::Scalar::I32 => F2I32_FUNCTION, + crate::Scalar::U32 => F2U32_FUNCTION, + crate::Scalar::I64 => F2I64_FUNCTION, + crate::Scalar::U64 => F2U64_FUNCTION, + _ => unreachable!(), + }; + self.write_value_type(module, &dst_ty)?; + write!(self.out, " {fun_name}(")?; + self.write_value_type(module, &src_ty)?; + writeln!(self.out, " value) {{")?; + let level = crate::back::Level(1); + write!(self.out, "{level}return {cast_str}(clamp(value, ")?; + self.write_literal(min)?; + write!(self.out, ", ")?; + self.write_literal(max)?; + writeln!(self.out, "));",)?; + writeln!(self.out, "}}")?; + writeln!(self.out)?; + } + } + Ok(()) + } + /// Helper function that writes various wrapped functions pub(super) fn write_wrapped_functions( &mut self, @@ -1366,6 +1466,7 @@ impl super::Writer<'_, W> { self.write_wrapped_binary_ops(module, func_ctx)?; self.write_wrapped_expression_functions(module, func_ctx.expressions, Some(func_ctx))?; self.write_wrapped_zero_value_functions(module, func_ctx.expressions)?; + self.write_wrapped_cast_functions(module, func_ctx)?; for (handle, _) in func_ctx.expressions.iter() { match func_ctx.expressions[handle] { diff --git a/naga/src/back/hlsl/keywords.rs b/naga/src/back/hlsl/keywords.rs index ac28797d8b9..d3b9ad8cf0d 100644 --- a/naga/src/back/hlsl/keywords.rs +++ b/naga/src/back/hlsl/keywords.rs @@ -830,6 +830,10 @@ pub const RESERVED: &[&str] = &[ super::writer::DIV_FUNCTION, super::writer::MOD_FUNCTION, super::writer::NEG_FUNCTION, + super::writer::F2I32_FUNCTION, + super::writer::F2U32_FUNCTION, + super::writer::F2I64_FUNCTION, + super::writer::F2U64_FUNCTION, ]; // DXC scalar types, from https://github.com/microsoft/DirectXShaderCompiler/blob/18c9e114f9c314f93e68fbc72ce207d4ed2e65ae/tools/clang/lib/AST/ASTContextHLSL.cpp#L48-L254 diff --git a/naga/src/back/hlsl/mod.rs b/naga/src/back/hlsl/mod.rs index 60a4bb4857b..9e041ff73f8 100644 --- a/naga/src/back/hlsl/mod.rs +++ b/naga/src/back/hlsl/mod.rs @@ -462,6 +462,7 @@ enum WrappedType { Math(help::WrappedMath), UnaryOp(help::WrappedUnaryOp), BinaryOp(help::WrappedBinaryOp), + Cast(help::WrappedCast), } #[derive(Default)] diff --git a/naga/src/back/hlsl/writer.rs b/naga/src/back/hlsl/writer.rs index f0ceceed549..dd79174ac61 100644 --- a/naga/src/back/hlsl/writer.rs +++ b/naga/src/back/hlsl/writer.rs @@ -38,6 +38,10 @@ pub(crate) const ABS_FUNCTION: &str = "naga_abs"; pub(crate) const DIV_FUNCTION: &str = "naga_div"; pub(crate) const MOD_FUNCTION: &str = "naga_mod"; pub(crate) const NEG_FUNCTION: &str = "naga_neg"; +pub(crate) const F2I32_FUNCTION: &str = "naga_f2i32"; +pub(crate) const F2U32_FUNCTION: &str = "naga_f2u32"; +pub(crate) const F2I64_FUNCTION: &str = "naga_f2i64"; +pub(crate) const F2U64_FUNCTION: &str = "naga_f2u64"; struct EpStructMember { name: String, @@ -2612,6 +2616,28 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { }) } + pub(super) fn write_literal(&mut self, literal: crate::Literal) -> BackendResult { + match literal { + crate::Literal::F64(value) => write!(self.out, "{value:?}L")?, + crate::Literal::F32(value) => write!(self.out, "{value:?}")?, + crate::Literal::F16(value) => write!(self.out, "{value:?}h")?, + crate::Literal::U32(value) => write!(self.out, "{value}u")?, + // HLSL has no suffix for explicit i32 literals, but not using any suffix + // makes the type ambiguous which prevents overload resolution from + // working. So we explicitly use the int() constructor syntax. + crate::Literal::I32(value) => write!(self.out, "int({value})")?, + crate::Literal::U64(value) => write!(self.out, "{value}uL")?, + crate::Literal::I64(value) => write!(self.out, "{value}L")?, + crate::Literal::Bool(value) => write!(self.out, "{value}")?, + crate::Literal::AbstractInt(_) | crate::Literal::AbstractFloat(_) => { + return Err(Error::Custom( + "Abstract types should not appear in IR presented to backends".into(), + )); + } + } + Ok(()) + } + fn write_possibly_const_expression( &mut self, module: &Module, @@ -2625,26 +2651,9 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { use crate::Expression; match expressions[expr] { - Expression::Literal(literal) => match literal { - // Floats are written using `Debug` instead of `Display` because it always appends the - // decimal part even it's zero - crate::Literal::F64(value) => write!(self.out, "{value:?}L")?, - crate::Literal::F32(value) => write!(self.out, "{value:?}")?, - crate::Literal::F16(value) => write!(self.out, "{value:?}h")?, - crate::Literal::U32(value) => write!(self.out, "{value}u")?, - // HLSL has no suffix for explicit i32 literals, but not using any suffix - // makes the type ambiguous which prevents overload resolution from - // working. So we explicitly use the int() constructor syntax. - crate::Literal::I32(value) => write!(self.out, "int({value})")?, - crate::Literal::U64(value) => write!(self.out, "{value}uL")?, - crate::Literal::I64(value) => write!(self.out, "{value}L")?, - crate::Literal::Bool(value) => write!(self.out, "{value}")?, - crate::Literal::AbstractInt(_) | crate::Literal::AbstractFloat(_) => { - return Err(Error::Custom( - "Abstract types should not appear in IR presented to backends".into(), - )); - } - }, + Expression::Literal(literal) => { + self.write_literal(literal)?; + } Expression::Constant(handle) => { let constant = &module.constants[handle]; if constant.name.is_some() { @@ -3320,53 +3329,72 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { convert, } => { let inner = func_ctx.resolve_type(expr, &module.types); - let close_paren = match convert { - Some(dst_width) => { - let scalar = Scalar { - kind, - width: dst_width, - }; - match *inner { - TypeInner::Vector { size, .. } => { - write!( - self.out, - "{}{}(", - scalar.to_hlsl_str()?, - common::vector_size_str(size) - )?; - } - TypeInner::Scalar(_) => { - write!(self.out, "{}(", scalar.to_hlsl_str()?,)?; - } - TypeInner::Matrix { columns, rows, .. } => { - write!( - self.out, - "{}{}x{}(", - scalar.to_hlsl_str()?, - common::vector_size_str(columns), - common::vector_size_str(rows) - )?; - } - _ => { - return Err(Error::Unimplemented(format!( - "write_expr expression::as {inner:?}" - ))); - } - }; - true - } - None => { - if inner.scalar_width() == Some(8) { - false - } else { - write!(self.out, "{}(", kind.to_hlsl_cast(),)?; + if inner.scalar_kind() == Some(ScalarKind::Float) + && (kind == ScalarKind::Sint || kind == ScalarKind::Uint) + && convert.is_some() + { + // Use helper functions for float to int casts in order to + // avoid undefined behaviour when value is out of range for + // the target type. + let fun_name = match (kind, convert) { + (ScalarKind::Sint, Some(4)) => F2I32_FUNCTION, + (ScalarKind::Uint, Some(4)) => F2U32_FUNCTION, + (ScalarKind::Sint, Some(8)) => F2I64_FUNCTION, + (ScalarKind::Uint, Some(8)) => F2U64_FUNCTION, + _ => unreachable!(), + }; + write!(self.out, "{fun_name}(")?; + self.write_expr(module, expr, func_ctx)?; + write!(self.out, ")")?; + } else { + let close_paren = match convert { + Some(dst_width) => { + let scalar = Scalar { + kind, + width: dst_width, + }; + match *inner { + TypeInner::Vector { size, .. } => { + write!( + self.out, + "{}{}(", + scalar.to_hlsl_str()?, + common::vector_size_str(size) + )?; + } + TypeInner::Scalar(_) => { + write!(self.out, "{}(", scalar.to_hlsl_str()?,)?; + } + TypeInner::Matrix { columns, rows, .. } => { + write!( + self.out, + "{}{}x{}(", + scalar.to_hlsl_str()?, + common::vector_size_str(columns), + common::vector_size_str(rows) + )?; + } + _ => { + return Err(Error::Unimplemented(format!( + "write_expr expression::as {inner:?}" + ))); + } + }; true } + None => { + if inner.scalar_width() == Some(8) { + false + } else { + write!(self.out, "{}(", kind.to_hlsl_cast(),)?; + true + } + } + }; + self.write_expr(module, expr, func_ctx)?; + if close_paren { + write!(self.out, ")")?; } - }; - self.write_expr(module, expr, func_ctx)?; - if close_paren { - write!(self.out, ")")?; } } Expression::Math { diff --git a/naga/src/back/msl/keywords.rs b/naga/src/back/msl/keywords.rs index 2228ef17648..4feb6004fc0 100644 --- a/naga/src/back/msl/keywords.rs +++ b/naga/src/back/msl/keywords.rs @@ -349,6 +349,10 @@ pub const RESERVED: &[&str] = &[ super::writer::DIV_FUNCTION, super::writer::MOD_FUNCTION, super::writer::NEG_FUNCTION, + super::writer::F2I32_FUNCTION, + super::writer::F2U32_FUNCTION, + super::writer::F2I64_FUNCTION, + super::writer::F2U64_FUNCTION, super::writer::ARGUMENT_BUFFER_WRAPPER_STRUCT, ]; diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs index c2042b12523..83b8bd43d7d 100644 --- a/naga/src/back/msl/writer.rs +++ b/naga/src/back/msl/writer.rs @@ -52,6 +52,10 @@ pub(crate) const ABS_FUNCTION: &str = "naga_abs"; pub(crate) const DIV_FUNCTION: &str = "naga_div"; pub(crate) const MOD_FUNCTION: &str = "naga_mod"; pub(crate) const NEG_FUNCTION: &str = "naga_neg"; +pub(crate) const F2I32_FUNCTION: &str = "naga_f2i32"; +pub(crate) const F2U32_FUNCTION: &str = "naga_f2u32"; +pub(crate) const F2I64_FUNCTION: &str = "naga_f2i64"; +pub(crate) const F2U64_FUNCTION: &str = "naga_f2u64"; /// For some reason, Metal does not let you have `metal::texture<..>*` as a buffer argument. /// However, if you put that texture inside a struct, everything is totally fine. This /// baffles me to no end. @@ -405,6 +409,11 @@ enum WrappedFunction { fun: crate::MathFunction, arg_ty: (Option, crate::Scalar), }, + Cast { + src_scalar: crate::Scalar, + vector_size: Option, + dst_scalar: crate::Scalar, + }, } pub struct Writer { @@ -1449,6 +1458,61 @@ impl Writer { ) } + fn put_literal(&mut self, literal: crate::Literal) -> BackendResult { + match literal { + crate::Literal::F64(_) => { + return Err(Error::CapabilityNotSupported(valid::Capabilities::FLOAT64)) + } + crate::Literal::F16(value) => { + if value.is_infinite() { + let sign = if value.is_sign_negative() { "-" } else { "" }; + write!(self.out, "{sign}INFINITY")?; + } else if value.is_nan() { + write!(self.out, "NAN")?; + } else { + let suffix = if value.fract() == f16::from_f32(0.0) { + ".0h" + } else { + "h" + }; + write!(self.out, "{value}{suffix}")?; + } + } + crate::Literal::F32(value) => { + if value.is_infinite() { + let sign = if value.is_sign_negative() { "-" } else { "" }; + write!(self.out, "{sign}INFINITY")?; + } else if value.is_nan() { + write!(self.out, "NAN")?; + } else { + let suffix = if value.fract() == 0.0 { ".0" } else { "" }; + write!(self.out, "{value}{suffix}")?; + } + } + crate::Literal::U32(value) => { + write!(self.out, "{value}u")?; + } + crate::Literal::I32(value) => { + write!(self.out, "{value}")?; + } + crate::Literal::U64(value) => { + write!(self.out, "{value}uL")?; + } + crate::Literal::I64(value) => { + write!(self.out, "{value}L")?; + } + crate::Literal::Bool(value) => { + write!(self.out, "{value}")?; + } + crate::Literal::AbstractInt(_) | crate::Literal::AbstractFloat(_) => { + return Err(Error::GenericValidation( + "Unsupported abstract literal".into(), + )); + } + } + Ok(()) + } + #[allow(clippy::too_many_arguments)] fn put_possibly_const_expression( &mut self, @@ -1465,57 +1529,9 @@ impl Writer { E: Fn(&mut Self, &C, Handle) -> BackendResult, { match expressions[expr_handle] { - crate::Expression::Literal(literal) => match literal { - crate::Literal::F64(_) => { - return Err(Error::CapabilityNotSupported(valid::Capabilities::FLOAT64)) - } - crate::Literal::F16(value) => { - if value.is_infinite() { - let sign = if value.is_sign_negative() { "-" } else { "" }; - write!(self.out, "{sign}INFINITY")?; - } else if value.is_nan() { - write!(self.out, "NAN")?; - } else { - let suffix = if value.fract() == f16::from_f32(0.0) { - ".0h" - } else { - "h" - }; - write!(self.out, "{value}{suffix}")?; - } - } - crate::Literal::F32(value) => { - if value.is_infinite() { - let sign = if value.is_sign_negative() { "-" } else { "" }; - write!(self.out, "{sign}INFINITY")?; - } else if value.is_nan() { - write!(self.out, "NAN")?; - } else { - let suffix = if value.fract() == 0.0 { ".0" } else { "" }; - write!(self.out, "{value}{suffix}")?; - } - } - crate::Literal::U32(value) => { - write!(self.out, "{value}u")?; - } - crate::Literal::I32(value) => { - write!(self.out, "{value}")?; - } - crate::Literal::U64(value) => { - write!(self.out, "{value}uL")?; - } - crate::Literal::I64(value) => { - write!(self.out, "{value}L")?; - } - crate::Literal::Bool(value) => { - write!(self.out, "{value}")?; - } - crate::Literal::AbstractInt(_) | crate::Literal::AbstractFloat(_) => { - return Err(Error::GenericValidation( - "Unsupported abstract literal".into(), - )); - } - }, + crate::Expression::Literal(literal) => { + self.put_literal(literal)?; + } crate::Expression::Constant(handle) => { let constant = &module.constants[handle]; if constant.name.is_some() { @@ -2315,24 +2331,43 @@ impl Writer { convert, } => match *context.resolve_type(expr) { crate::TypeInner::Scalar(src) | crate::TypeInner::Vector { scalar: src, .. } => { - let target_scalar = crate::Scalar { - kind, - width: convert.unwrap_or(src.width), - }; - let op = match convert { - Some(_) => "static_cast", - None => "as_type", - }; - write!(self.out, "{op}<")?; - match *context.resolve_type(expr) { - crate::TypeInner::Vector { size, .. } => { - put_numeric_type(&mut self.out, target_scalar, &[size])? - } - _ => put_numeric_type(&mut self.out, target_scalar, &[])?, - }; - write!(self.out, ">(")?; - self.put_expression(expr, context, true)?; - write!(self.out, ")")?; + if src.kind == crate::ScalarKind::Float + && (kind == crate::ScalarKind::Sint || kind == crate::ScalarKind::Uint) + && convert.is_some() + { + // Use helper functions for float to int casts in order to avoid + // undefined behaviour when value is out of range for the target + // type. + let fun_name = match (kind, convert) { + (crate::ScalarKind::Sint, Some(4)) => F2I32_FUNCTION, + (crate::ScalarKind::Uint, Some(4)) => F2U32_FUNCTION, + (crate::ScalarKind::Sint, Some(8)) => F2I64_FUNCTION, + (crate::ScalarKind::Uint, Some(8)) => F2U64_FUNCTION, + _ => unreachable!(), + }; + write!(self.out, "{fun_name}(")?; + self.put_expression(expr, context, true)?; + write!(self.out, ")")?; + } else { + let target_scalar = crate::Scalar { + kind, + width: convert.unwrap_or(src.width), + }; + let op = match convert { + Some(_) => "static_cast", + None => "as_type", + }; + write!(self.out, "{op}<")?; + match *context.resolve_type(expr) { + crate::TypeInner::Vector { size, .. } => { + put_numeric_type(&mut self.out, target_scalar, &[size])? + } + _ => put_numeric_type(&mut self.out, target_scalar, &[])?, + }; + write!(self.out, ">(")?; + self.put_expression(expr, context, true)?; + write!(self.out, ")")?; + } } crate::TypeInner::Matrix { columns, @@ -5344,6 +5379,76 @@ template _ => {} } } + crate::Expression::As { + expr, + kind, + convert: Some(width), + } => { + // Avoid undefined behaviour when casting from a float to integer + // when the value is out of range for the target type. Additionally + // ensure we clamp to the correct value as per the WGSL spec. + // + // https://www.w3.org/TR/WGSL/#floating-point-conversion: + // * If X is exactly representable in the target type T, then the + // result is that value. + // * Otherwise, the result is the value in T closest to + // truncate(X) and also exactly representable in the original + // floating point type. + let src_ty = func_ctx.resolve_type(expr, &module.types); + let Some((vector_size, src_scalar)) = src_ty.vector_size_and_scalar() else { + continue; + }; + let dst_scalar = crate::Scalar { kind, width }; + if src_scalar.kind != crate::ScalarKind::Float + || (dst_scalar.kind != crate::ScalarKind::Sint + && dst_scalar.kind != crate::ScalarKind::Uint) + { + continue; + } + let wrapped = WrappedFunction::Cast { + src_scalar, + vector_size, + dst_scalar, + }; + if !self.wrapped_functions.insert(wrapped) { + continue; + } + let (min, max) = proc::min_max_float_representable_by(src_scalar, dst_scalar); + + let mut src_type_name = String::new(); + match vector_size { + None => put_numeric_type(&mut src_type_name, src_scalar, &[])?, + Some(size) => put_numeric_type(&mut src_type_name, src_scalar, &[size])?, + }; + let mut dst_type_name = String::new(); + match vector_size { + None => put_numeric_type(&mut dst_type_name, dst_scalar, &[])?, + Some(size) => put_numeric_type(&mut dst_type_name, dst_scalar, &[size])?, + }; + let fun_name = match dst_scalar { + crate::Scalar::I32 => F2I32_FUNCTION, + crate::Scalar::U32 => F2U32_FUNCTION, + crate::Scalar::I64 => F2I64_FUNCTION, + crate::Scalar::U64 => F2U64_FUNCTION, + _ => unreachable!(), + }; + + writeln!( + self.out, + "{dst_type_name} {fun_name}({src_type_name} value) {{" + )?; + let level = back::Level(1); + write!( + self.out, + "{level}return static_cast<{dst_type_name}>({NAMESPACE}::clamp(value, " + )?; + self.put_literal(min)?; + write!(self.out, ", ")?; + self.put_literal(max)?; + writeln!(self.out, "));")?; + writeln!(self.out, "}}")?; + writeln!(self.out)?; + } _ => {} } } diff --git a/naga/src/back/spv/block.rs b/naga/src/back/spv/block.rs index 71c592bd222..f021dc6d813 100644 --- a/naga/src/back/spv/block.rs +++ b/naga/src/back/spv/block.rs @@ -1876,10 +1876,10 @@ impl BlockContext<'_> { }; enum Cast { - Identity, - Unary(spirv::Op), - Binary(spirv::Op, Word), - Ternary(spirv::Op, Word, Word), + Identity(Word), + Unary(spirv::Op, Word), + Binary(spirv::Op, Word, Word), + Ternary(spirv::Op, Word, Word, Word), } let cast = match (src_scalar.kind, kind, convert) { // Filter out identity casts. Some Adreno drivers are @@ -1888,10 +1888,10 @@ impl BlockContext<'_> { if src_kind == kind && convert.filter(|&width| width != src_scalar.width).is_none() => { - Cast::Identity + Cast::Identity(expr_id) } - (Sk::Bool, Sk::Bool, _) => Cast::Unary(spirv::Op::CopyObject), - (_, _, None) => Cast::Unary(spirv::Op::Bitcast), + (Sk::Bool, Sk::Bool, _) => Cast::Unary(spirv::Op::CopyObject, expr_id), + (_, _, None) => Cast::Unary(spirv::Op::Bitcast, expr_id), // casting to a bool - generate `OpXxxNotEqual` (_, Sk::Bool, Some(_)) => { let op = match src_scalar.kind { @@ -1916,7 +1916,7 @@ impl BlockContext<'_> { None => zero_scalar_id, }; - Cast::Binary(op, zero_id) + Cast::Binary(op, expr_id, zero_id) } // casting from a bool - generate `OpSelect` (Sk::Bool, _, Some(dst_width)) => { @@ -1948,60 +1948,99 @@ impl BlockContext<'_> { None => (one_scalar_id, zero_scalar_id), }; - Cast::Ternary(spirv::Op::Select, accept_id, reject_id) + Cast::Ternary(spirv::Op::Select, expr_id, accept_id, reject_id) + } + // Avoid undefined behaviour when casting from a float to integer + // when the value is out of range for the target type. Additionally + // ensure we clamp to the correct value as per the WGSL spec. + // + // https://www.w3.org/TR/WGSL/#floating-point-conversion: + // * If X is exactly representable in the target type T, then the + // result is that value. + // * Otherwise, the result is the value in T closest to + // truncate(X) and also exactly representable in the original + // floating point type. + (Sk::Float, Sk::Sint | Sk::Uint, Some(width)) => { + let dst_scalar = crate::Scalar { kind, width }; + let (min, max) = + crate::proc::min_max_float_representable_by(src_scalar, dst_scalar); + let expr_type_id = self.get_expression_type_id(&self.fun_info[expr].ty); + + let maybe_splat_const = |writer: &mut Writer, const_id| match src_size { + None => const_id, + Some(size) => { + let constituent_ids = [const_id; crate::VectorSize::MAX]; + writer.get_constant_composite( + LookupType::Local(LocalType::Numeric(NumericType::Vector { + size, + scalar: src_scalar, + })), + &constituent_ids[..size as usize], + ) + } + }; + let min_const_id = self.writer.get_constant_scalar(min); + let min_const_id = maybe_splat_const(self.writer, min_const_id); + let max_const_id = self.writer.get_constant_scalar(max); + let max_const_id = maybe_splat_const(self.writer, max_const_id); + + let clamp_id = self.gen_id(); + block.body.push(Instruction::ext_inst( + self.writer.gl450_ext_inst_id, + spirv::GLOp::FClamp, + expr_type_id, + clamp_id, + &[expr_id, min_const_id, max_const_id], + )); + + let op = match dst_scalar.kind { + crate::ScalarKind::Sint => spirv::Op::ConvertFToS, + crate::ScalarKind::Uint => spirv::Op::ConvertFToU, + _ => unreachable!(), + }; + Cast::Unary(op, clamp_id) } - (Sk::Float, Sk::Uint, Some(_)) => Cast::Unary(spirv::Op::ConvertFToU), - (Sk::Float, Sk::Sint, Some(_)) => Cast::Unary(spirv::Op::ConvertFToS), (Sk::Float, Sk::Float, Some(dst_width)) if src_scalar.width != dst_width => { - Cast::Unary(spirv::Op::FConvert) + Cast::Unary(spirv::Op::FConvert, expr_id) } - (Sk::Sint, Sk::Float, Some(_)) => Cast::Unary(spirv::Op::ConvertSToF), + (Sk::Sint, Sk::Float, Some(_)) => Cast::Unary(spirv::Op::ConvertSToF, expr_id), (Sk::Sint, Sk::Sint, Some(dst_width)) if src_scalar.width != dst_width => { - Cast::Unary(spirv::Op::SConvert) + Cast::Unary(spirv::Op::SConvert, expr_id) } - (Sk::Uint, Sk::Float, Some(_)) => Cast::Unary(spirv::Op::ConvertUToF), + (Sk::Uint, Sk::Float, Some(_)) => Cast::Unary(spirv::Op::ConvertUToF, expr_id), (Sk::Uint, Sk::Uint, Some(dst_width)) if src_scalar.width != dst_width => { - Cast::Unary(spirv::Op::UConvert) + Cast::Unary(spirv::Op::UConvert, expr_id) } (Sk::Uint, Sk::Sint, Some(dst_width)) if src_scalar.width != dst_width => { - Cast::Unary(spirv::Op::SConvert) + Cast::Unary(spirv::Op::SConvert, expr_id) } (Sk::Sint, Sk::Uint, Some(dst_width)) if src_scalar.width != dst_width => { - Cast::Unary(spirv::Op::UConvert) + Cast::Unary(spirv::Op::UConvert, expr_id) } // We assume it's either an identity cast, or int-uint. - _ => Cast::Unary(spirv::Op::Bitcast), + _ => Cast::Unary(spirv::Op::Bitcast, expr_id), }; Ok(match cast { - Cast::Identity => expr_id, - Cast::Unary(op) => { + Cast::Identity(expr) => expr, + Cast::Unary(op, op1) => { let id = self.gen_id(); block .body - .push(Instruction::unary(op, result_type_id, id, expr_id)); + .push(Instruction::unary(op, result_type_id, id, op1)); id } - Cast::Binary(op, operand) => { + Cast::Binary(op, op1, op2) => { let id = self.gen_id(); - block.body.push(Instruction::binary( - op, - result_type_id, - id, - expr_id, - operand, - )); + block + .body + .push(Instruction::binary(op, result_type_id, id, op1, op2)); id } - Cast::Ternary(op, op1, op2) => { + Cast::Ternary(op, op1, op2, op3) => { let id = self.gen_id(); - block.body.push(Instruction::ternary( - op, - result_type_id, - id, - expr_id, - op1, - op2, - )); + block + .body + .push(Instruction::ternary(op, result_type_id, id, op1, op2, op3)); id } }) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index c0a8253b99f..406dbbef52d 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -1697,13 +1697,15 @@ impl<'a> ConstantEvaluator<'a> { Err(ConstantEvaluatorError::InvalidCastArg { from, to }) }; + use crate::proc::type_methods::IntFloatLimits; + let expr = match self.expressions[expr] { Expression::Literal(literal) => { let literal = match target { Sc::I32 => Literal::I32(match literal { Literal::I32(v) => v, Literal::U32(v) => v as i32, - Literal::F32(v) => v as i32, + Literal::F32(v) => v.clamp(i32::min_float(), i32::max_float()) as i32, Literal::F16(v) => f16::to_i32(&v).unwrap(), //Only None on NaN or Inf Literal::Bool(v) => v as i32, Literal::F64(_) | Literal::I64(_) | Literal::U64(_) => { @@ -1715,7 +1717,7 @@ impl<'a> ConstantEvaluator<'a> { Sc::U32 => Literal::U32(match literal { Literal::I32(v) => v as u32, Literal::U32(v) => v, - Literal::F32(v) => v as u32, + Literal::F32(v) => v.clamp(u32::min_float(), u32::max_float()) as u32, // max(0) avoids None due to negative, therefore only None on NaN or Inf Literal::F16(v) => f16::to_u32(&v.max(f16::ZERO)).unwrap(), Literal::Bool(v) => v as u32, @@ -1728,9 +1730,9 @@ impl<'a> ConstantEvaluator<'a> { Sc::I64 => Literal::I64(match literal { Literal::I32(v) => v as i64, Literal::U32(v) => v as i64, - Literal::F32(v) => v as i64, + Literal::F32(v) => v.clamp(i64::min_float(), i64::max_float()) as i64, Literal::Bool(v) => v as i64, - Literal::F64(v) => v as i64, + Literal::F64(v) => v.clamp(i64::min_float(), i64::max_float()) as i64, Literal::I64(v) => v, Literal::U64(v) => v as i64, Literal::F16(v) => f16::to_i64(&v).unwrap(), //Only None on NaN or Inf @@ -1740,9 +1742,9 @@ impl<'a> ConstantEvaluator<'a> { Sc::U64 => Literal::U64(match literal { Literal::I32(v) => v as u64, Literal::U32(v) => v as u64, - Literal::F32(v) => v as u64, + Literal::F32(v) => v.clamp(u64::min_float(), u64::max_float()) as u64, Literal::Bool(v) => v as u64, - Literal::F64(v) => v as u64, + Literal::F64(v) => v.clamp(u64::min_float(), u64::max_float()) as u64, Literal::I64(v) => v as u64, Literal::U64(v) => v, // max(0) avoids None due to negative, therefore only None on NaN or Inf @@ -2753,19 +2755,19 @@ impl TryFromAbstract for u32 { impl TryFromAbstract for i64 { fn try_from_abstract(value: f64) -> Result { - // As above, except i64::MIN and i64::MAX are not exactly representable - // by f64. i64 is not part of the WGSL spec, however, so we're free to - // ignore that requirement. - Ok(value as i64) + // As above, except we clamp to the minimum and maximum values + // representable by both f64 and i64. + use crate::proc::type_methods::IntFloatLimits; + Ok(value.clamp(i64::min_float(), i64::max_float()) as i64) } } impl TryFromAbstract for u64 { fn try_from_abstract(value: f64) -> Result { - // As above, except u64::MAX is not exactly representable by f64. u64 - // is not part of the WGSL spec, however, so we're free to ignore that - // requirement. - Ok(value as u64) + // As above, this time clamping to the minimum and maximum values + // representable by both f64 and u64. + use crate::proc::type_methods::IntFloatLimits; + Ok(value.clamp(u64::min_float(), u64::max_float()) as u64) } } diff --git a/naga/src/proc/mod.rs b/naga/src/proc/mod.rs index 118ed19f5ee..787a2f1a63c 100644 --- a/naga/src/proc/mod.rs +++ b/naga/src/proc/mod.rs @@ -20,6 +20,7 @@ pub use layouter::{Alignment, LayoutError, LayoutErrorInner, Layouter, TypeLayou pub use namer::{EntryPointIndex, NameKey, Namer}; pub use terminator::ensure_block_returns; use thiserror::Error; +pub use type_methods::min_max_float_representable_by; pub use typifier::{ResolveContext, ResolveError, TypeResolution}; impl From for super::Scalar { diff --git a/naga/src/proc/type_methods.rs b/naga/src/proc/type_methods.rs index d4089bc1fab..9e99fb29a09 100644 --- a/naga/src/proc/type_methods.rs +++ b/naga/src/proc/type_methods.rs @@ -308,3 +308,114 @@ impl crate::TypeInner { } } } + +/// Helper trait for providing the min and max values exactly representable by +/// the integer type `Self` and floating point type `F`. +pub trait IntFloatLimits +where + F: num_traits::Float, +{ + /// Returns the minimum value exactly representable by the integer type + /// `Self` and floating point type `F`. + fn min_float() -> F; + /// Returns the maximum value exactly representable by the integer type + /// `Self` and floating point type `F`. + fn max_float() -> F; +} + +macro_rules! define_int_float_limits { + ($int:ty, $float:ty, $min:expr, $max:expr) => { + impl IntFloatLimits<$float> for $int { + fn min_float() -> $float { + $min + } + fn max_float() -> $float { + $max + } + } + }; +} + +define_int_float_limits!(i32, half::f16, half::f16::MIN, half::f16::MAX); +define_int_float_limits!(u32, half::f16, half::f16::ZERO, half::f16::MAX); +define_int_float_limits!(i64, half::f16, half::f16::MIN, half::f16::MAX); +define_int_float_limits!(u64, half::f16, half::f16::ZERO, half::f16::MAX); +define_int_float_limits!(i32, f32, -2147483648.0f32, 2147483520.0f32); +define_int_float_limits!(u32, f32, 0.0f32, 4294967040.0f32); +define_int_float_limits!( + i64, + f32, + -9223372036854775808.0f32, + 9223371487098961920.0f32 +); +define_int_float_limits!(u64, f32, 0.0f32, 18446742974197923840.0f32); +define_int_float_limits!(i32, f64, -2147483648.0f64, 2147483647.0f64); +define_int_float_limits!(u32, f64, 0.0f64, 4294967295.0f64); +define_int_float_limits!( + i64, + f64, + -9223372036854775808.0f64, + 9223372036854774784.0f64 +); +define_int_float_limits!(u64, f64, 0.0f64, 18446744073709549568.0f64); + +/// Returns a tuple of [`crate::Literal`]s representing the minimum and maximum +/// float values exactly representable by the provided float and integer types. +/// Panics if `float` is not one of `F16`, `F32`, or `F64`, or `int` is +/// not one of `I32`, `U32`, `I64`, or `U64`. +pub fn min_max_float_representable_by( + float: crate::Scalar, + int: crate::Scalar, +) -> (crate::Literal, crate::Literal) { + match (float, int) { + (crate::Scalar::F16, crate::Scalar::I32) => ( + crate::Literal::F16(i32::min_float()), + crate::Literal::F16(i32::max_float()), + ), + (crate::Scalar::F16, crate::Scalar::U32) => ( + crate::Literal::F16(u32::min_float()), + crate::Literal::F16(u32::max_float()), + ), + (crate::Scalar::F16, crate::Scalar::I64) => ( + crate::Literal::F16(i64::min_float()), + crate::Literal::F16(i64::max_float()), + ), + (crate::Scalar::F16, crate::Scalar::U64) => ( + crate::Literal::F16(u64::min_float()), + crate::Literal::F16(u64::max_float()), + ), + (crate::Scalar::F32, crate::Scalar::I32) => ( + crate::Literal::F32(i32::min_float()), + crate::Literal::F32(i32::max_float()), + ), + (crate::Scalar::F32, crate::Scalar::U32) => ( + crate::Literal::F32(u32::min_float()), + crate::Literal::F32(u32::max_float()), + ), + (crate::Scalar::F32, crate::Scalar::I64) => ( + crate::Literal::F32(i64::min_float()), + crate::Literal::F32(i64::max_float()), + ), + (crate::Scalar::F32, crate::Scalar::U64) => ( + crate::Literal::F32(u64::min_float()), + crate::Literal::F32(u64::max_float()), + ), + (crate::Scalar::F64, crate::Scalar::I32) => ( + crate::Literal::F64(i32::min_float()), + crate::Literal::F64(i32::max_float()), + ), + (crate::Scalar::F64, crate::Scalar::U32) => ( + crate::Literal::F64(u32::min_float()), + crate::Literal::F64(u32::max_float()), + ), + (crate::Scalar::F64, crate::Scalar::I64) => ( + crate::Literal::F64(i64::min_float()), + crate::Literal::F64(i64::max_float()), + ), + (crate::Scalar::F64, crate::Scalar::U64) => ( + crate::Literal::F64(u64::min_float()), + crate::Literal::F64(u64::max_float()), + ), + _ => unreachable!(), + } +} diff --git a/naga/tests/in/wgsl/conversion-float-to-int-no-f64.toml b/naga/tests/in/wgsl/conversion-float-to-int-no-f64.toml new file mode 100644 index 00000000000..bd00460e530 --- /dev/null +++ b/naga/tests/in/wgsl/conversion-float-to-int-no-f64.toml @@ -0,0 +1,4 @@ +# Other backends are tested by conversion-float-to-int.wgsl which is a +# superset of this test. +targets = "METAL" +god_mode = true diff --git a/naga/tests/in/wgsl/conversion-float-to-int-no-f64.wgsl b/naga/tests/in/wgsl/conversion-float-to-int-no-f64.wgsl new file mode 100644 index 00000000000..bf9eed402c8 --- /dev/null +++ b/naga/tests/in/wgsl/conversion-float-to-int-no-f64.wgsl @@ -0,0 +1,109 @@ +// conversion-float-to-int.wgsl test with `f64` types removed. +// This means it can be used with the Metal backend. + +enable f16; + +const MIN_F16 = -65504h; +const MAX_F16 = 65504h; +const MIN_F32 = -3.40282347E+38f; +const MAX_F32 = 3.40282347E+38f; +const MIN_ABSTRACT_FLOAT = -1.7976931348623157E+308; +const MAX_ABSTRACT_FLOAT = 1.7976931348623157E+308; + +// conversion from float to int during const evaluation. If value is out of +// range of the destination type we must clamp to a value that is representable +// by both the source and destination types. +fn test_const_eval() { + var min_f16_to_i32 = i32(MIN_F16); + var max_f16_to_i32 = i32(MAX_F16); + var min_f16_to_u32 = u32(MIN_F16); + var max_f16_to_u32 = u32(MAX_F16); + var min_f16_to_i64 = i64(MIN_F16); + var max_f16_to_i64 = i64(MAX_F16); + var min_f16_to_u64 = u64(MIN_F16); + var max_f16_to_u64 = u64(MAX_F16); + var min_f32_to_i32 = i32(MIN_F32); + var max_f32_to_i32 = i32(MAX_F32); + var min_f32_to_u32 = u32(MIN_F32); + var max_f32_to_u32 = u32(MAX_F32); + var min_f32_to_i64 = i64(MIN_F32); + var max_f32_to_i64 = i64(MAX_F32); + var min_f32_to_u64 = u64(MIN_F32); + var max_f32_to_u64 = u64(MAX_F32); + var min_abstract_float_to_i32 = i32(MIN_ABSTRACT_FLOAT); + var max_abstract_float_to_i32 = i32(MAX_ABSTRACT_FLOAT); + var min_abstract_float_to_u32 = u32(MIN_ABSTRACT_FLOAT); + var max_abstract_float_to_u32 = u32(MAX_ABSTRACT_FLOAT); + var min_abstract_float_to_i64 = i64(MIN_ABSTRACT_FLOAT); + var max_abstract_float_to_i64 = i64(MAX_ABSTRACT_FLOAT); + var min_abstract_float_to_u64 = u64(MIN_ABSTRACT_FLOAT); + var max_abstract_float_to_u64 = u64(MAX_ABSTRACT_FLOAT); +} + +// conversion from float to int at runtime. Generated code must ensure we avoid +// undefined behaviour due to casting a value that is out of range of the +// destination type, and that in such cases the result is a value that is +// representable by both the source and destination types. +fn test_f16_to_i32(f: f16) -> i32 { + return i32(f); +} + +fn test_f16_to_u32(f: f16) -> u32 { + return u32(f); +} + +fn test_f16_to_i64(f: f16) -> i64 { + return i64(f); +} + +fn test_f16_to_u64(f: f16) -> u64 { + return u64(f); +} + +fn test_f32_to_i32(f: f32) -> i32 { + return i32(f); +} + +fn test_f32_to_u32(f: f32) -> u32 { + return u32(f); +} + +fn test_f32_to_i64(f: f32) -> i64 { + return i64(f); +} + +fn test_f32_to_u64(f: f32) -> u64 { + return u64(f); +} + +fn test_f16_to_i32_vec(f: vec2) -> vec2 { + return vec2(f); +} + +fn test_f16_to_u32_vec(f: vec2) -> vec2 { + return vec2(f); +} + +fn test_f16_to_i64_vec(f: vec2) -> vec2 { + return vec2(f); +} + +fn test_f16_to_u64_vec(f: vec2) -> vec2 { + return vec2(f); +} + +fn test_f32_to_i32_vec(f: vec2) -> vec2 { + return vec2(f); +} + +fn test_f32_to_u32_vec(f: vec2) -> vec2 { + return vec2(f); +} + +fn test_f32_to_i64_vec(f: vec2) -> vec2 { + return vec2(f); +} + +fn test_f32_to_u64_vec(f: vec2) -> vec2 { + return vec2(f); +} diff --git a/naga/tests/in/wgsl/conversion-float-to-int.toml b/naga/tests/in/wgsl/conversion-float-to-int.toml new file mode 100644 index 00000000000..87d10e02284 --- /dev/null +++ b/naga/tests/in/wgsl/conversion-float-to-int.toml @@ -0,0 +1,8 @@ +# GLSL does not support i64 or f16, but we don't handle float to integer +# conversion specially there anyway. +# MSL does not support f64. See conversion-float-to-int-no-f64.wgsl. +targets = "SPIRV | HLSL | WGSL" +god_mode = true + +[hlsl] +shader_model = "V6_0" diff --git a/naga/tests/in/wgsl/conversion-float-to-int.wgsl b/naga/tests/in/wgsl/conversion-float-to-int.wgsl new file mode 100644 index 00000000000..ccec1528073 --- /dev/null +++ b/naga/tests/in/wgsl/conversion-float-to-int.wgsl @@ -0,0 +1,147 @@ +enable f16; + +const MIN_F16 = -65504h; +const MAX_F16 = 65504h; +const MIN_F32 = -3.40282347E+38f; +const MAX_F32 = 3.40282347E+38f; +const MIN_F64 = -1.7976931348623157E+308lf; +const MAX_F64 = 1.7976931348623157E+308lf; +const MIN_ABSTRACT_FLOAT = -1.7976931348623157E+308; +const MAX_ABSTRACT_FLOAT = 1.7976931348623157E+308; + +// conversion from float to int during const evaluation. If value is out of +// range of the destination type we must clamp to a value that is representable +// by both the source and destination types. +fn test_const_eval() { + var min_f16_to_i32 = i32(MIN_F16); + var max_f16_to_i32 = i32(MAX_F16); + var min_f16_to_u32 = u32(MIN_F16); + var max_f16_to_u32 = u32(MAX_F16); + var min_f16_to_i64 = i64(MIN_F16); + var max_f16_to_i64 = i64(MAX_F16); + var min_f16_to_u64 = u64(MIN_F16); + var max_f16_to_u64 = u64(MAX_F16); + var min_f32_to_i32 = i32(MIN_F32); + var max_f32_to_i32 = i32(MAX_F32); + var min_f32_to_u32 = u32(MIN_F32); + var max_f32_to_u32 = u32(MAX_F32); + var min_f32_to_i64 = i64(MIN_F32); + var max_f32_to_i64 = i64(MAX_F32); + var min_f32_to_u64 = u64(MIN_F32); + var max_f32_to_u64 = u64(MAX_F32); + var min_f64_to_i64 = i64(MIN_F64); + var max_f64_to_i64 = i64(MAX_F64); + var min_f64_to_u64 = u64(MIN_F64); + var max_f64_to_u64 = u64(MAX_F64); + var min_abstract_float_to_i32 = i32(MIN_ABSTRACT_FLOAT); + var max_abstract_float_to_i32 = i32(MAX_ABSTRACT_FLOAT); + var min_abstract_float_to_u32 = u32(MIN_ABSTRACT_FLOAT); + var max_abstract_float_to_u32 = u32(MAX_ABSTRACT_FLOAT); + var min_abstract_float_to_i64 = i64(MIN_ABSTRACT_FLOAT); + var max_abstract_float_to_i64 = i64(MAX_ABSTRACT_FLOAT); + var min_abstract_float_to_u64 = u64(MIN_ABSTRACT_FLOAT); + var max_abstract_float_to_u64 = u64(MAX_ABSTRACT_FLOAT); +} + +// conversion from float to int at runtime. Generated code must ensure we avoid +// undefined behaviour due to casting a value that is out of range of the +// destination type, and that in such cases the result is a value that is +// representable by both the source and destination types. +fn test_f16_to_i32(f: f16) -> i32 { + return i32(f); +} + +fn test_f16_to_u32(f: f16) -> u32 { + return u32(f); +} + +fn test_f16_to_i64(f: f16) -> i64 { + return i64(f); +} + +fn test_f16_to_u64(f: f16) -> u64 { + return u64(f); +} + +fn test_f32_to_i32(f: f32) -> i32 { + return i32(f); +} + +fn test_f32_to_u32(f: f32) -> u32 { + return u32(f); +} + +fn test_f32_to_i64(f: f32) -> i64 { + return i64(f); +} + +fn test_f32_to_u64(f: f32) -> u64 { + return u64(f); +} + +fn test_f64_to_i32(f: f64) -> i32 { + return i32(f); +} + +fn test_f64_to_u32(f: f64) -> u32 { + return u32(f); +} + +fn test_f64_to_i64(f: f64) -> i64 { + return i64(f); +} + +fn test_f64_to_u64(f: f64) -> u64 { + return u64(f); +} + +fn test_f16_to_i32_vec(f: vec2) -> vec2 { + return vec2(f); +} + +fn test_f16_to_u32_vec(f: vec2) -> vec2 { + return vec2(f); +} + +fn test_f16_to_i64_vec(f: vec2) -> vec2 { + return vec2(f); +} + +fn test_f16_to_u64_vec(f: vec2) -> vec2 { + return vec2(f); +} + +fn test_f32_to_i32_vec(f: vec2) -> vec2 { + return vec2(f); +} + +fn test_f32_to_u32_vec(f: vec2) -> vec2 { + return vec2(f); +} + +fn test_f32_to_i64_vec(f: vec2) -> vec2 { + return vec2(f); +} + +fn test_f32_to_u64_vec(f: vec2) -> vec2 { + return vec2(f); +} + +fn test_f64_to_i32_vec(f: vec2) -> vec2 { + return vec2(f); +} + +fn test_f64_to_u32_vec(f: vec2) -> vec2 { + return vec2(f); +} + +fn test_f64_to_i64_vec(f: vec2) -> vec2 { + return vec2(f); +} + +fn test_f64_to_u64_vec(f: vec2) -> vec2 { + return vec2(f); +} + +@compute @workgroup_size(1) +fn main() {} diff --git a/naga/tests/out/hlsl/access.hlsl b/naga/tests/out/hlsl/access.hlsl index 167046fb30f..b9f03865687 100644 --- a/naga/tests/out/hlsl/access.hlsl +++ b/naga/tests/out/hlsl/access.hlsl @@ -342,6 +342,10 @@ ret_ZeroValuearray5_array10_float__ ZeroValuearray5_array10_float__() { return (float[5][10])0; } +int naga_f2i32(float value) { + return int(clamp(value, -2147483600.0, 2147483500.0)); +} + typedef uint2 ret_Constructarray2_uint2_[2]; ret_Constructarray2_uint2_ Constructarray2_uint2_(uint2 arg0, uint2 arg1) { uint2 ret[2] = { arg0, arg1 }; @@ -370,7 +374,7 @@ float4 foo_vert(uint vi : SV_VertexID) : SV_Position int a_2 = asint(bar.Load(0+(((NagaBufferLengthRW(bar) - 160) / 8) - 2u)*8+160)); int2 c = asint(qux.Load2(0)); const float _e33 = read_from_private(foo); - c2_ = Constructarray5_int_(a_2, int(b), int(3), int(4), int(5)); + c2_ = Constructarray5_int_(a_2, naga_f2i32(b), int(3), int(4), int(5)); c2_[min(uint((vi + 1u)), 4u)] = int(42); int value_1 = c2_[min(uint(vi), 4u)]; const float _e47 = test_arr_as_arg(ZeroValuearray5_array10_float__()); diff --git a/naga/tests/out/hlsl/conversion-float-to-int.hlsl b/naga/tests/out/hlsl/conversion-float-to-int.hlsl new file mode 100644 index 00000000000..ca0a79ce691 --- /dev/null +++ b/naga/tests/out/hlsl/conversion-float-to-int.hlsl @@ -0,0 +1,262 @@ +static const half MIN_F16_ = -65504.0h; +static const half MAX_F16_ = 65504.0h; +static const float MIN_F32_ = -3.4028235e38; +static const float MAX_F32_ = 3.4028235e38; +static const double MIN_F64_ = -1.7976931348623157e308L; +static const double MAX_F64_ = 1.7976931348623157e308L; + +void test_const_eval() +{ + int min_f16_to_i32_ = int(-65504); + int max_f16_to_i32_ = int(65504); + uint min_f16_to_u32_ = 0u; + uint max_f16_to_u32_ = 65504u; + int64_t min_f16_to_i64_ = -65504L; + int64_t max_f16_to_i64_ = 65504L; + uint64_t min_f16_to_u64_ = 0uL; + uint64_t max_f16_to_u64_ = 65504uL; + int min_f32_to_i32_ = int(-2147483648); + int max_f32_to_i32_ = int(2147483520); + uint min_f32_to_u32_ = 0u; + uint max_f32_to_u32_ = 4294967040u; + int64_t min_f32_to_i64_ = -9223372036854775808L; + int64_t max_f32_to_i64_ = 9223371487098961920L; + uint64_t min_f32_to_u64_ = 0uL; + uint64_t max_f32_to_u64_ = 18446742974197923840uL; + int64_t min_f64_to_i64_ = -9223372036854775808L; + int64_t max_f64_to_i64_ = 9223372036854774784L; + uint64_t min_f64_to_u64_ = 0uL; + uint64_t max_f64_to_u64_ = 18446744073709549568uL; + int min_abstract_float_to_i32_ = int(-2147483648); + int max_abstract_float_to_i32_ = int(2147483647); + uint min_abstract_float_to_u32_ = 0u; + uint max_abstract_float_to_u32_ = 4294967295u; + int64_t min_abstract_float_to_i64_ = -9223372036854775808L; + int64_t max_abstract_float_to_i64_ = 9223372036854774784L; + uint64_t min_abstract_float_to_u64_ = 0uL; + uint64_t max_abstract_float_to_u64_ = 18446744073709549568uL; + + return; +} + +int naga_f2i32(half value) { + return int(clamp(value, -65504.0h, 65504.0h)); +} + +int test_f16_to_i32_(half f) +{ + return naga_f2i32(f); +} + +uint naga_f2u32(half value) { + return uint(clamp(value, 0.0h, 65504.0h)); +} + +uint test_f16_to_u32_(half f_1) +{ + return naga_f2u32(f_1); +} + +int64_t naga_f2i64(half value) { + return int64_t(clamp(value, -65504.0h, 65504.0h)); +} + +int64_t test_f16_to_i64_(half f_2) +{ + return naga_f2i64(f_2); +} + +uint64_t naga_f2u64(half value) { + return uint64_t(clamp(value, 0.0h, 65504.0h)); +} + +uint64_t test_f16_to_u64_(half f_3) +{ + return naga_f2u64(f_3); +} + +int naga_f2i32(float value) { + return int(clamp(value, -2147483600.0, 2147483500.0)); +} + +int test_f32_to_i32_(float f_4) +{ + return naga_f2i32(f_4); +} + +uint naga_f2u32(float value) { + return uint(clamp(value, 0.0, 4294967000.0)); +} + +uint test_f32_to_u32_(float f_5) +{ + return naga_f2u32(f_5); +} + +int64_t naga_f2i64(float value) { + return int64_t(clamp(value, -9.223372e18, 9.2233715e18)); +} + +int64_t test_f32_to_i64_(float f_6) +{ + return naga_f2i64(f_6); +} + +uint64_t naga_f2u64(float value) { + return uint64_t(clamp(value, 0.0, 1.8446743e19)); +} + +uint64_t test_f32_to_u64_(float f_7) +{ + return naga_f2u64(f_7); +} + +int naga_f2i32(double value) { + return int(clamp(value, -2147483648.0L, 2147483647.0L)); +} + +int test_f64_to_i32_(double f_8) +{ + return naga_f2i32(f_8); +} + +uint naga_f2u32(double value) { + return uint(clamp(value, 0.0L, 4294967295.0L)); +} + +uint test_f64_to_u32_(double f_9) +{ + return naga_f2u32(f_9); +} + +int64_t naga_f2i64(double value) { + return int64_t(clamp(value, -9.223372036854776e18L, 9.223372036854775e18L)); +} + +int64_t test_f64_to_i64_(double f_10) +{ + return naga_f2i64(f_10); +} + +uint64_t naga_f2u64(double value) { + return uint64_t(clamp(value, 0.0L, 1.844674407370955e19L)); +} + +uint64_t test_f64_to_u64_(double f_11) +{ + return naga_f2u64(f_11); +} + +int2 naga_f2i32(half2 value) { + return int2(clamp(value, -65504.0h, 65504.0h)); +} + +int2 test_f16_to_i32_vec(half2 f_12) +{ + return naga_f2i32(f_12); +} + +uint2 naga_f2u32(half2 value) { + return uint2(clamp(value, 0.0h, 65504.0h)); +} + +uint2 test_f16_to_u32_vec(half2 f_13) +{ + return naga_f2u32(f_13); +} + +int64_t2 naga_f2i64(half2 value) { + return int64_t2(clamp(value, -65504.0h, 65504.0h)); +} + +int64_t2 test_f16_to_i64_vec(half2 f_14) +{ + return naga_f2i64(f_14); +} + +uint64_t2 naga_f2u64(half2 value) { + return uint64_t2(clamp(value, 0.0h, 65504.0h)); +} + +uint64_t2 test_f16_to_u64_vec(half2 f_15) +{ + return naga_f2u64(f_15); +} + +int2 naga_f2i32(float2 value) { + return int2(clamp(value, -2147483600.0, 2147483500.0)); +} + +int2 test_f32_to_i32_vec(float2 f_16) +{ + return naga_f2i32(f_16); +} + +uint2 naga_f2u32(float2 value) { + return uint2(clamp(value, 0.0, 4294967000.0)); +} + +uint2 test_f32_to_u32_vec(float2 f_17) +{ + return naga_f2u32(f_17); +} + +int64_t2 naga_f2i64(float2 value) { + return int64_t2(clamp(value, -9.223372e18, 9.2233715e18)); +} + +int64_t2 test_f32_to_i64_vec(float2 f_18) +{ + return naga_f2i64(f_18); +} + +uint64_t2 naga_f2u64(float2 value) { + return uint64_t2(clamp(value, 0.0, 1.8446743e19)); +} + +uint64_t2 test_f32_to_u64_vec(float2 f_19) +{ + return naga_f2u64(f_19); +} + +int2 naga_f2i32(double2 value) { + return int2(clamp(value, -2147483648.0L, 2147483647.0L)); +} + +int2 test_f64_to_i32_vec(double2 f_20) +{ + return naga_f2i32(f_20); +} + +uint2 naga_f2u32(double2 value) { + return uint2(clamp(value, 0.0L, 4294967295.0L)); +} + +uint2 test_f64_to_u32_vec(double2 f_21) +{ + return naga_f2u32(f_21); +} + +int64_t2 naga_f2i64(double2 value) { + return int64_t2(clamp(value, -9.223372036854776e18L, 9.223372036854775e18L)); +} + +int64_t2 test_f64_to_i64_vec(double2 f_22) +{ + return naga_f2i64(f_22); +} + +uint64_t2 naga_f2u64(double2 value) { + return uint64_t2(clamp(value, 0.0L, 1.844674407370955e19L)); +} + +uint64_t2 test_f64_to_u64_vec(double2 f_23) +{ + return naga_f2u64(f_23); +} + +[numthreads(1, 1, 1)] +void main() +{ + return; +} diff --git a/naga/tests/out/hlsl/conversion-float-to-int.ron b/naga/tests/out/hlsl/conversion-float-to-int.ron new file mode 100644 index 00000000000..b973fe3da13 --- /dev/null +++ b/naga/tests/out/hlsl/conversion-float-to-int.ron @@ -0,0 +1,12 @@ +( + vertex:[ + ], + fragment:[ + ], + compute:[ + ( + entry_point:"main", + target_profile:"cs_6_0", + ), + ], +) diff --git a/naga/tests/out/hlsl/image.hlsl b/naga/tests/out/hlsl/image.hlsl index 5319a1ca367..28c4b326d25 100644 --- a/naga/tests/out/hlsl/image.hlsl +++ b/naga/tests/out/hlsl/image.hlsl @@ -59,13 +59,17 @@ void main(uint3 local_id : SV_GroupThreadID) return; } +uint naga_f2u32(float value) { + return uint(clamp(value, 0.0, 4294967000.0)); +} + [numthreads(16, 1, 1)] void depth_load(uint3 local_id_1 : SV_GroupThreadID) { uint2 dim_1 = NagaRWDimensions2D(image_storage_src); int2 itc_1 = naga_mod(int2((dim_1 * local_id_1.xy)), int2(int(10), int(20))); float val = image_depth_multisampled_src.Load(itc_1, int(local_id_1.z)).x; - image_dst[itc_1.x] = (uint(val)).xxxx; + image_dst[itc_1.x] = (naga_f2u32(val)).xxxx; return; } diff --git a/naga/tests/out/hlsl/int64.hlsl b/naga/tests/out/hlsl/int64.hlsl index a20c8e53bb3..ef13f9c623e 100644 --- a/naga/tests/out/hlsl/int64.hlsl +++ b/naga/tests/out/hlsl/int64.hlsl @@ -53,6 +53,10 @@ ByteAddressBuffer input_arrays : register(t2); RWByteAddressBuffer output : register(u3); RWByteAddressBuffer output_arrays : register(u4); +int64_t naga_f2i64(float value) { + return int64_t(clamp(value, -9.223372e18, 9.2233715e18)); +} + typedef int64_t ret_Constructarray2_int64_t_[2]; ret_Constructarray2_int64_t_ Constructarray2_int64_t_(int64_t arg0, int64_t arg1) { int64_t ret[2] = { arg0, arg1 }; @@ -79,7 +83,7 @@ int64_t int64_function(int64_t x) float _e35 = input_uniform.val_f32_; int64_t _e36 = val; int64_t _e40 = val; - val = (_e40 + int64_t((_e35 + float(_e36)))); + val = (_e40 + naga_f2i64((_e35 + float(_e36)))); int64_t _e44 = input_uniform.val_i64_; int64_t _e47 = val; val = (_e47 + (_e44).xxx.z); @@ -142,6 +146,10 @@ int64_t int64_function(int64_t x) return _e153; } +uint64_t naga_f2u64(float value) { + return uint64_t(clamp(value, 0.0, 1.8446743e19)); +} + typedef uint64_t ret_Constructarray2_uint64_t_[2]; ret_Constructarray2_uint64_t_ Constructarray2_uint64_t_(uint64_t arg0, uint64_t arg1) { uint64_t ret[2] = { arg0, arg1 }; @@ -168,7 +176,7 @@ uint64_t uint64_function(uint64_t x_1) float _e35 = input_uniform.val_f32_; uint64_t _e36 = val_1; uint64_t _e40 = val_1; - val_1 = (_e40 + uint64_t((_e35 + float(_e36)))); + val_1 = (_e40 + naga_f2u64((_e35 + float(_e36)))); uint64_t _e44 = input_uniform.val_u64_; uint64_t _e47 = val_1; val_1 = (_e47 + (_e44).xxx.z); diff --git a/naga/tests/out/msl/access.msl b/naga/tests/out/msl/access.msl index eb0808f80e3..8a56f0eeae5 100644 --- a/naga/tests/out/msl/access.msl +++ b/naga/tests/out/msl/access.msl @@ -247,6 +247,10 @@ int var_members_of_members( int _e15 = thing.om_nom_nom.delicious; return _e15; } +int naga_f2i32(float value) { + return static_cast(metal::clamp(value, -2147483600.0, 2147483500.0)); +} + struct foo_vertInput { }; @@ -273,7 +277,7 @@ vertex foo_vertOutput foo_vert( int a_2 = bar.data[(1 + (_buffer_sizes.size1 - 160 - 8) / 8) - 2u].value; metal::int2 c = qux; float _e33 = read_from_private(foo); - c2_ = type_20 {a_2, static_cast(b), 3, 4, 5}; + c2_ = type_20 {a_2, naga_f2i32(b), 3, 4, 5}; c2_.inner[vi + 1u] = 42; int value_1 = c2_.inner[vi]; float _e47 = test_arr_as_arg(type_18 {}); diff --git a/naga/tests/out/msl/bounds-check-restrict.msl b/naga/tests/out/msl/bounds-check-restrict.msl index a8bf3715555..dc5315ec02d 100644 --- a/naga/tests/out/msl/bounds-check-restrict.msl +++ b/naga/tests/out/msl/bounds-check-restrict.msl @@ -73,12 +73,16 @@ float index_twice( return _e6; } +int naga_f2i32(float value) { + return static_cast(metal::clamp(value, -2147483600.0, 2147483500.0)); +} + float index_expensive( int i_6, device Globals const& globals, constant _mslBufferSizes& _buffer_sizes ) { - float _e11 = globals.a.inner[metal::min(unsigned(static_cast(metal::sin(static_cast(i_6) / 100.0) * 100.0)), 9u)]; + float _e11 = globals.a.inner[metal::min(unsigned(naga_f2i32(metal::sin(static_cast(i_6) / 100.0) * 100.0)), 9u)]; return _e11; } @@ -149,7 +153,7 @@ void set_expensive( device Globals& globals, constant _mslBufferSizes& _buffer_sizes ) { - globals.a.inner[metal::min(unsigned(static_cast(metal::sin(static_cast(i_12) / 100.0) * 100.0)), 9u)] = v_6; + globals.a.inner[metal::min(unsigned(naga_f2i32(metal::sin(static_cast(i_12) / 100.0) * 100.0)), 9u)] = v_6; return; } diff --git a/naga/tests/out/msl/bounds-check-zero.msl b/naga/tests/out/msl/bounds-check-zero.msl index 8269d4bf70f..7488002f76d 100644 --- a/naga/tests/out/msl/bounds-check-zero.msl +++ b/naga/tests/out/msl/bounds-check-zero.msl @@ -79,12 +79,16 @@ float index_twice( return _e6; } +int naga_f2i32(float value) { + return static_cast(metal::clamp(value, -2147483600.0, 2147483500.0)); +} + float index_expensive( int i_6, device Globals const& globals, constant _mslBufferSizes& _buffer_sizes ) { - int _e9 = static_cast(metal::sin(static_cast(i_6) / 100.0) * 100.0); + int _e9 = naga_f2i32(metal::sin(static_cast(i_6) / 100.0) * 100.0); float _e11 = uint(_e9) < 10 ? globals.a.inner[_e9] : DefaultConstructible(); return _e11; } @@ -166,7 +170,7 @@ void set_expensive( device Globals& globals, constant _mslBufferSizes& _buffer_sizes ) { - int _e10 = static_cast(metal::sin(static_cast(i_12) / 100.0) * 100.0); + int _e10 = naga_f2i32(metal::sin(static_cast(i_12) / 100.0) * 100.0); if (uint(_e10) < 10) { globals.a.inner[_e10] = v_6; } diff --git a/naga/tests/out/msl/conversion-float-to-int-no-f64.msl b/naga/tests/out/msl/conversion-float-to-int-no-f64.msl new file mode 100644 index 00000000000..b6a5f7521f9 --- /dev/null +++ b/naga/tests/out/msl/conversion-float-to-int-no-f64.msl @@ -0,0 +1,199 @@ +// language: metal1.0 +#include +#include + +using metal::uint; + +constant half MIN_F16_ = -65504.0h; +constant half MAX_F16_ = 65504.0h; +constant float MIN_F32_ = -340282350000000000000000000000000000000.0; +constant float MAX_F32_ = 340282350000000000000000000000000000000.0; + +void test_const_eval( +) { + int min_f16_to_i32_ = -65504; + int max_f16_to_i32_ = 65504; + uint min_f16_to_u32_ = 0u; + uint max_f16_to_u32_ = 65504u; + long min_f16_to_i64_ = -65504L; + long max_f16_to_i64_ = 65504L; + ulong min_f16_to_u64_ = 0uL; + ulong max_f16_to_u64_ = 65504uL; + int min_f32_to_i32_ = -2147483648; + int max_f32_to_i32_ = 2147483520; + uint min_f32_to_u32_ = 0u; + uint max_f32_to_u32_ = 4294967040u; + long min_f32_to_i64_ = -9223372036854775808L; + long max_f32_to_i64_ = 9223371487098961920L; + ulong min_f32_to_u64_ = 0uL; + ulong max_f32_to_u64_ = 18446742974197923840uL; + int min_abstract_float_to_i32_ = -2147483648; + int max_abstract_float_to_i32_ = 2147483647; + uint min_abstract_float_to_u32_ = 0u; + uint max_abstract_float_to_u32_ = 4294967295u; + long min_abstract_float_to_i64_ = -9223372036854775808L; + long max_abstract_float_to_i64_ = 9223372036854774784L; + ulong min_abstract_float_to_u64_ = 0uL; + ulong max_abstract_float_to_u64_ = 18446744073709549568uL; + return; +} + +int naga_f2i32(half value) { + return static_cast(metal::clamp(value, -65504.0h, 65504.0h)); +} + +int test_f16_to_i32_( + half f +) { + return naga_f2i32(f); +} + +uint naga_f2u32(half value) { + return static_cast(metal::clamp(value, 0.0h, 65504.0h)); +} + +uint test_f16_to_u32_( + half f_1 +) { + return naga_f2u32(f_1); +} + +long naga_f2i64(half value) { + return static_cast(metal::clamp(value, -65504.0h, 65504.0h)); +} + +long test_f16_to_i64_( + half f_2 +) { + return naga_f2i64(f_2); +} + +ulong naga_f2u64(half value) { + return static_cast(metal::clamp(value, 0.0h, 65504.0h)); +} + +ulong test_f16_to_u64_( + half f_3 +) { + return naga_f2u64(f_3); +} + +int naga_f2i32(float value) { + return static_cast(metal::clamp(value, -2147483600.0, 2147483500.0)); +} + +int test_f32_to_i32_( + float f_4 +) { + return naga_f2i32(f_4); +} + +uint naga_f2u32(float value) { + return static_cast(metal::clamp(value, 0.0, 4294967000.0)); +} + +uint test_f32_to_u32_( + float f_5 +) { + return naga_f2u32(f_5); +} + +long naga_f2i64(float value) { + return static_cast(metal::clamp(value, -9223372000000000000.0, 9223371500000000000.0)); +} + +long test_f32_to_i64_( + float f_6 +) { + return naga_f2i64(f_6); +} + +ulong naga_f2u64(float value) { + return static_cast(metal::clamp(value, 0.0, 18446743000000000000.0)); +} + +ulong test_f32_to_u64_( + float f_7 +) { + return naga_f2u64(f_7); +} + +metal::int2 naga_f2i32(metal::half2 value) { + return static_cast(metal::clamp(value, -65504.0h, 65504.0h)); +} + +metal::int2 test_f16_to_i32_vec( + metal::half2 f_8 +) { + return naga_f2i32(f_8); +} + +metal::uint2 naga_f2u32(metal::half2 value) { + return static_cast(metal::clamp(value, 0.0h, 65504.0h)); +} + +metal::uint2 test_f16_to_u32_vec( + metal::half2 f_9 +) { + return naga_f2u32(f_9); +} + +metal::long2 naga_f2i64(metal::half2 value) { + return static_cast(metal::clamp(value, -65504.0h, 65504.0h)); +} + +metal::long2 test_f16_to_i64_vec( + metal::half2 f_10 +) { + return naga_f2i64(f_10); +} + +metal::ulong2 naga_f2u64(metal::half2 value) { + return static_cast(metal::clamp(value, 0.0h, 65504.0h)); +} + +metal::ulong2 test_f16_to_u64_vec( + metal::half2 f_11 +) { + return naga_f2u64(f_11); +} + +metal::int2 naga_f2i32(metal::float2 value) { + return static_cast(metal::clamp(value, -2147483600.0, 2147483500.0)); +} + +metal::int2 test_f32_to_i32_vec( + metal::float2 f_12 +) { + return naga_f2i32(f_12); +} + +metal::uint2 naga_f2u32(metal::float2 value) { + return static_cast(metal::clamp(value, 0.0, 4294967000.0)); +} + +metal::uint2 test_f32_to_u32_vec( + metal::float2 f_13 +) { + return naga_f2u32(f_13); +} + +metal::long2 naga_f2i64(metal::float2 value) { + return static_cast(metal::clamp(value, -9223372000000000000.0, 9223371500000000000.0)); +} + +metal::long2 test_f32_to_i64_vec( + metal::float2 f_14 +) { + return naga_f2i64(f_14); +} + +metal::ulong2 naga_f2u64(metal::float2 value) { + return static_cast(metal::clamp(value, 0.0, 18446743000000000000.0)); +} + +metal::ulong2 test_f32_to_u64_vec( + metal::float2 f_15 +) { + return naga_f2u64(f_15); +} diff --git a/naga/tests/out/msl/image.msl b/naga/tests/out/msl/image.msl index 920d20f6c30..723ebd9871e 100644 --- a/naga/tests/out/msl/image.msl +++ b/naga/tests/out/msl/image.msl @@ -41,6 +41,10 @@ kernel void main_( return; } +uint naga_f2u32(float value) { + return static_cast(metal::clamp(value, 0.0, 4294967000.0)); +} + struct depth_loadInput { }; @@ -53,7 +57,7 @@ kernel void depth_load( metal::uint2 dim_1 = metal::uint2(image_storage_src.get_width(), image_storage_src.get_height()); metal::int2 itc_1 = naga_mod(static_cast(dim_1 * local_id_1.xy), metal::int2(10, 20)); float val = image_depth_multisampled_src.read(metal::uint2(itc_1), static_cast(local_id_1.z)); - image_dst.write(metal::uint4(static_cast(val)), uint(itc_1.x)); + image_dst.write(metal::uint4(naga_f2u32(val)), uint(itc_1.x)); return; } diff --git a/naga/tests/out/msl/int64.msl b/naga/tests/out/msl/int64.msl index 3e21cd0ae52..e026cac919b 100644 --- a/naga/tests/out/msl/int64.msl +++ b/naga/tests/out/msl/int64.msl @@ -34,6 +34,10 @@ struct StorageCompatible { }; constant ulong constant_variable = 20uL; +long naga_f2i64(float value) { + return static_cast(metal::clamp(value, -9223372000000000000.0, 9223371500000000000.0)); +} + long naga_abs(long val) { return metal::select(as_type(-as_type(val)), val, val >= 0); } @@ -63,7 +67,7 @@ long int64_function( float _e35 = input_uniform.val_f32_; long _e36 = val; long _e40 = val; - val = as_type(as_type(_e40) + as_type(static_cast(_e35 + static_cast(_e36)))); + val = as_type(as_type(_e40) + as_type(naga_f2i64(_e35 + static_cast(_e36)))); long _e44 = input_uniform.val_i64_; long _e47 = val; val = as_type(as_type(_e47) + as_type(metal::long3(_e44).z)); @@ -124,6 +128,10 @@ long int64_function( return _e153; } +ulong naga_f2u64(float value) { + return static_cast(metal::clamp(value, 0.0, 18446743000000000000.0)); +} + ulong uint64_function( ulong x_1, constant UniformCompatible& input_uniform, @@ -149,7 +157,7 @@ ulong uint64_function( float _e35 = input_uniform.val_f32_; ulong _e36 = val_1; ulong _e40 = val_1; - val_1 = _e40 + static_cast(_e35 + static_cast(_e36)); + val_1 = _e40 + naga_f2u64(_e35 + static_cast(_e36)); ulong _e44 = input_uniform.val_u64_; ulong _e47 = val_1; val_1 = _e47 + metal::ulong3(_e44).z; diff --git a/naga/tests/out/spv/access.spvasm b/naga/tests/out/spv/access.spvasm index 2aa23e832b9..ac979146f5b 100644 --- a/naga/tests/out/spv/access.spvasm +++ b/naga/tests/out/spv/access.spvasm @@ -1,18 +1,18 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 402 +; Bound: 405 OpCapability Shader OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 OpEntryPoint Vertex %311 "foo_vert" %306 %309 -OpEntryPoint Fragment %363 "foo_frag" %362 -OpEntryPoint GLCompute %381 "assign_through_ptr" -OpEntryPoint GLCompute %392 "assign_to_ptr_components" -OpExecutionMode %363 OriginUpperLeft -OpExecutionMode %381 LocalSize 1 1 1 -OpExecutionMode %392 LocalSize 1 1 1 +OpEntryPoint Fragment %366 "foo_frag" %365 +OpEntryPoint GLCompute %384 "assign_through_ptr" +OpEntryPoint GLCompute %395 "assign_to_ptr_components" +OpExecutionMode %366 OriginUpperLeft +OpExecutionMode %384 LocalSize 1 1 1 +OpExecutionMode %395 LocalSize 1 1 1 %3 = OpString "access.wgsl" OpSource Unknown 0 %3 "// This snapshot tests accessing various containers, dereferencing pointers. @@ -326,13 +326,13 @@ OpName %306 "vi" OpName %311 "foo_vert" OpName %322 "foo" OpName %323 "c2" -OpName %363 "foo_frag" -OpName %381 "assign_through_ptr" -OpName %386 "val" -OpName %387 "arr" -OpName %392 "assign_to_ptr_components" -OpName %393 "s1" -OpName %395 "a1" +OpName %366 "foo_frag" +OpName %384 "assign_through_ptr" +OpName %389 "val" +OpName %390 "arr" +OpName %395 "assign_to_ptr_components" +OpName %396 "s1" +OpName %398 "a1" OpMemberDecorate %7 0 Offset 0 OpMemberDecorate %7 1 Offset 16 OpMemberDecorate %7 2 Offset 28 @@ -386,7 +386,7 @@ OpDecorate %63 Block OpMemberDecorate %63 0 Offset 0 OpDecorate %306 BuiltIn VertexIndex OpDecorate %309 BuiltIn Position -OpDecorate %362 Location 0 +OpDecorate %365 Location 0 %2 = OpTypeVoid %4 = OpTypeInt 32 0 %5 = OpTypeVector %4 3 @@ -540,24 +540,26 @@ OpDecorate %362 Location 0 %340 = OpTypePointer StorageBuffer %20 %343 = OpTypePointer StorageBuffer %8 %344 = OpTypePointer StorageBuffer %6 -%356 = OpTypeVector %6 4 -%362 = OpVariable %310 Output -%365 = OpConstantComposite %11 %316 %316 %316 -%366 = OpConstantComposite %11 %71 %71 %71 -%367 = OpConstantComposite %11 %73 %73 %73 -%368 = OpConstantComposite %11 %75 %75 %75 -%369 = OpConstantComposite %10 %365 %366 %367 %368 -%370 = OpConstantComposite %18 %48 %48 -%371 = OpConstantComposite %18 %44 %44 -%372 = OpConstantComposite %19 %370 %371 -%373 = OpConstantNull %24 -%374 = OpConstantComposite %32 %316 %316 %316 %316 -%382 = OpConstant %4 33 -%383 = OpConstantComposite %32 %79 %79 %79 %79 -%384 = OpConstantComposite %32 %147 %147 %147 %147 -%385 = OpConstantComposite %35 %383 %384 -%394 = OpConstantNull %37 -%396 = OpConstantNull %39 +%349 = OpConstant %9 -2147483600.0 +%350 = OpConstant %9 2147483500.0 +%359 = OpTypeVector %6 4 +%365 = OpVariable %310 Output +%368 = OpConstantComposite %11 %316 %316 %316 +%369 = OpConstantComposite %11 %71 %71 %71 +%370 = OpConstantComposite %11 %73 %73 %73 +%371 = OpConstantComposite %11 %75 %75 %75 +%372 = OpConstantComposite %10 %368 %369 %370 %371 +%373 = OpConstantComposite %18 %48 %48 +%374 = OpConstantComposite %18 %44 %44 +%375 = OpConstantComposite %19 %373 %374 +%376 = OpConstantNull %24 +%377 = OpConstantComposite %32 %316 %316 %316 %316 +%385 = OpConstant %4 33 +%386 = OpConstantComposite %32 %79 %79 %79 %79 +%387 = OpConstantComposite %32 %147 %147 %147 %147 +%388 = OpConstantComposite %35 %386 %387 +%397 = OpConstantNull %37 +%399 = OpConstantNull %39 %66 = OpFunction %2 None %67 %65 = OpLabel %94 = OpVariable %95 Function %70 @@ -974,94 +976,95 @@ OpLine %3 129 53 OpLine %3 130 18 %348 = OpFunctionCall %9 %198 %322 OpLine %3 133 28 -%349 = OpConvertFToS %6 %339 +%351 = OpExtInst %9 %1 FClamp %339 %349 %350 +%352 = OpConvertFToS %6 %351 OpLine %3 133 11 -%350 = OpCompositeConstruct %33 %346 %349 %318 %319 %320 +%353 = OpCompositeConstruct %33 %346 %352 %318 %319 %320 OpLine %3 133 2 -OpStore %323 %350 +OpStore %323 %353 OpLine %3 134 2 -%351 = OpIAdd %4 %308 %44 +%354 = OpIAdd %4 %308 %44 OpLine %3 134 2 -%352 = OpAccessChain %95 %323 %351 -OpStore %352 %264 +%355 = OpAccessChain %95 %323 %354 +OpStore %355 %264 OpLine %3 135 14 -%353 = OpAccessChain %95 %323 %308 -%354 = OpLoad %6 %353 +%356 = OpAccessChain %95 %323 %308 +%357 = OpLoad %6 %356 OpLine %3 137 2 -%355 = OpFunctionCall %9 %204 %321 +%358 = OpFunctionCall %9 %204 %321 OpLine %3 139 19 -%357 = OpCompositeConstruct %356 %354 %354 %354 %354 -%358 = OpConvertSToF %32 %357 -%359 = OpMatrixTimesVector %11 %332 %358 +%360 = OpCompositeConstruct %359 %357 %357 %357 %357 +%361 = OpConvertSToF %32 %360 +%362 = OpMatrixTimesVector %11 %332 %361 OpLine %3 139 9 -%360 = OpCompositeConstruct %32 %359 %73 -OpStore %309 %360 +%363 = OpCompositeConstruct %32 %362 %73 +OpStore %309 %363 OpReturn OpFunctionEnd -%363 = OpFunction %2 None %67 -%361 = OpLabel -%364 = OpAccessChain %313 %59 %48 -OpBranch %375 -%375 = OpLabel +%366 = OpFunction %2 None %67 +%364 = OpLabel +%367 = OpAccessChain %313 %59 %48 +OpBranch %378 +%378 = OpLabel OpLine %3 145 2 OpLine %3 145 2 OpLine %3 145 2 -%376 = OpAccessChain %337 %54 %48 %44 %15 -OpStore %376 %71 +%379 = OpAccessChain %337 %54 %48 %44 %15 +OpStore %379 %71 OpLine %3 146 2 OpLine %3 146 28 OpLine %3 146 44 OpLine %3 146 60 OpLine %3 146 16 OpLine %3 146 2 -%377 = OpAccessChain %330 %54 %48 -OpStore %377 %369 +%380 = OpAccessChain %330 %54 %48 +OpStore %380 %372 OpLine %3 147 2 OpLine %3 147 32 OpLine %3 147 12 OpLine %3 147 2 -%378 = OpAccessChain %333 %54 %40 -OpStore %378 %372 +%381 = OpAccessChain %333 %54 %40 +OpStore %381 %375 OpLine %3 148 2 OpLine %3 148 2 OpLine %3 148 2 -%379 = OpAccessChain %344 %54 %31 %44 %48 -OpStore %379 %70 +%382 = OpAccessChain %344 %54 %31 %44 %48 +OpStore %382 %70 OpLine %3 149 2 -OpStore %364 %373 +OpStore %367 %376 OpLine %3 151 9 -OpStore %362 %374 +OpStore %365 %377 OpReturn OpFunctionEnd -%381 = OpFunction %2 None %67 -%380 = OpLabel -%386 = OpVariable %34 Function %382 -%387 = OpVariable %36 Function %385 -OpBranch %388 -%388 = OpLabel +%384 = OpFunction %2 None %67 +%383 = OpLabel +%389 = OpVariable %34 Function %385 +%390 = OpVariable %36 Function %388 +OpBranch %391 +%391 = OpLabel OpLine %3 165 5 -%389 = OpFunctionCall %2 %211 %386 +%392 = OpFunctionCall %2 %211 %389 OpLine %3 167 32 OpLine %3 167 43 OpLine %3 167 32 OpLine %3 167 12 OpLine %3 168 5 -%390 = OpFunctionCall %2 %217 %387 +%393 = OpFunctionCall %2 %217 %390 OpReturn OpFunctionEnd -%392 = OpFunction %2 None %67 -%391 = OpLabel -%393 = OpVariable %38 Function %394 -%395 = OpVariable %41 Function %396 -OpBranch %397 -%397 = OpLabel +%395 = OpFunction %2 None %67 +%394 = OpLabel +%396 = OpVariable %38 Function %397 +%398 = OpVariable %41 Function %399 +OpBranch %400 +%400 = OpLabel OpLine %3 194 4 -%398 = OpFunctionCall %2 %232 %393 +%401 = OpFunctionCall %2 %232 %396 OpLine %3 195 4 -%399 = OpFunctionCall %4 %225 %393 +%402 = OpFunctionCall %4 %225 %396 OpLine %3 198 4 -%400 = OpFunctionCall %2 %245 %395 +%403 = OpFunctionCall %2 %245 %398 OpLine %3 199 4 -%401 = OpFunctionCall %4 %238 %395 +%404 = OpFunctionCall %4 %238 %398 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/spv/bounds-check-restrict.spvasm b/naga/tests/out/spv/bounds-check-restrict.spvasm index d831ddf942d..571d7b61dbf 100644 --- a/naga/tests/out/spv/bounds-check-restrict.spvasm +++ b/naga/tests/out/spv/bounds-check-restrict.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 178 +; Bound: 182 OpCapability Shader OpCapability Linkage OpExtension "SPV_KHR_storage_buffer_storage_class" @@ -45,12 +45,14 @@ OpDecorate %12 Binding 0 %60 = OpConstant %6 2 %68 = OpTypeFunction %3 %11 %11 %77 = OpConstant %3 100.0 -%89 = OpTypeFunction %3 -%103 = OpTypeFunction %2 %11 %3 -%127 = OpTypeFunction %2 %11 %7 -%136 = OpTypeFunction %2 %11 %11 %3 -%156 = OpTypeFunction %2 %3 -%166 = OpConstant %6 1000 +%83 = OpConstant %3 -2147483600.0 +%84 = OpConstant %3 2147483500.0 +%92 = OpTypeFunction %3 +%106 = OpTypeFunction %2 %11 %3 +%130 = OpTypeFunction %2 %11 %7 +%139 = OpTypeFunction %2 %11 %11 %3 +%160 = OpTypeFunction %2 %3 +%170 = OpConstant %6 1000 %16 = OpFunction %3 None %17 %15 = OpFunctionParameter %11 %14 = OpLabel @@ -124,134 +126,136 @@ OpBranch %78 %80 = OpFDiv %3 %79 %77 %81 = OpExtInst %3 %1 Sin %80 %82 = OpFMul %3 %81 %77 -%83 = OpConvertFToS %11 %82 -%84 = OpExtInst %6 %1 UMin %83 %21 -%85 = OpAccessChain %20 %12 %23 %84 -%86 = OpLoad %3 %85 -OpReturnValue %86 +%85 = OpExtInst %3 %1 FClamp %82 %83 %84 +%86 = OpConvertFToS %11 %85 +%87 = OpExtInst %6 %1 UMin %86 %21 +%88 = OpAccessChain %20 %12 %23 %87 +%89 = OpLoad %3 %88 +OpReturnValue %89 OpFunctionEnd -%88 = OpFunction %3 None %89 -%87 = OpLabel -OpBranch %90 +%91 = OpFunction %3 None %92 %90 = OpLabel -%91 = OpAccessChain %20 %12 %23 %21 -%92 = OpLoad %3 %91 -%93 = OpAccessChain %20 %12 %32 %35 -%94 = OpLoad %3 %93 -%95 = OpFAdd %3 %92 %94 -%96 = OpAccessChain %20 %12 %60 %60 %35 +OpBranch %93 +%93 = OpLabel +%94 = OpAccessChain %20 %12 %23 %21 +%95 = OpLoad %3 %94 +%96 = OpAccessChain %20 %12 %32 %35 %97 = OpLoad %3 %96 %98 = OpFAdd %3 %95 %97 -OpReturnValue %98 +%99 = OpAccessChain %20 %12 %60 %60 %35 +%100 = OpLoad %3 %99 +%101 = OpFAdd %3 %98 %100 +OpReturnValue %101 OpFunctionEnd -%102 = OpFunction %2 None %103 -%100 = OpFunctionParameter %11 -%101 = OpFunctionParameter %3 -%99 = OpLabel -OpBranch %104 -%104 = OpLabel -%105 = OpExtInst %6 %1 UMin %100 %21 -%106 = OpAccessChain %20 %12 %23 %105 -OpStore %106 %101 -OpReturn -OpFunctionEnd -%110 = OpFunction %2 None %103 -%108 = OpFunctionParameter %11 -%109 = OpFunctionParameter %3 +%105 = OpFunction %2 None %106 +%103 = OpFunctionParameter %11 +%104 = OpFunctionParameter %3 +%102 = OpLabel +OpBranch %107 %107 = OpLabel -OpBranch %111 -%111 = OpLabel -%112 = OpArrayLength %6 %12 3 -%113 = OpISub %6 %112 %32 -%114 = OpExtInst %6 %1 UMin %108 %113 -%115 = OpAccessChain %20 %12 %35 %114 -OpStore %115 %109 +%108 = OpExtInst %6 %1 UMin %103 %21 +%109 = OpAccessChain %20 %12 %23 %108 +OpStore %109 %104 OpReturn OpFunctionEnd -%119 = OpFunction %2 None %103 -%117 = OpFunctionParameter %11 -%118 = OpFunctionParameter %3 -%116 = OpLabel -OpBranch %120 -%120 = OpLabel -%121 = OpExtInst %6 %1 UMin %117 %35 -%122 = OpAccessChain %20 %12 %32 %121 -OpStore %122 %118 +%113 = OpFunction %2 None %106 +%111 = OpFunctionParameter %11 +%112 = OpFunctionParameter %3 +%110 = OpLabel +OpBranch %114 +%114 = OpLabel +%115 = OpArrayLength %6 %12 3 +%116 = OpISub %6 %115 %32 +%117 = OpExtInst %6 %1 UMin %111 %116 +%118 = OpAccessChain %20 %12 %35 %117 +OpStore %118 %112 OpReturn OpFunctionEnd -%126 = OpFunction %2 None %127 -%124 = OpFunctionParameter %11 -%125 = OpFunctionParameter %7 +%122 = OpFunction %2 None %106 +%120 = OpFunctionParameter %11 +%121 = OpFunctionParameter %3 +%119 = OpLabel +OpBranch %123 %123 = OpLabel -OpBranch %128 -%128 = OpLabel -%129 = OpExtInst %6 %1 UMin %124 %60 -%130 = OpAccessChain %42 %12 %60 %129 -OpStore %130 %125 +%124 = OpExtInst %6 %1 UMin %120 %35 +%125 = OpAccessChain %20 %12 %32 %124 +OpStore %125 %121 OpReturn OpFunctionEnd -%135 = OpFunction %2 None %136 -%132 = OpFunctionParameter %11 -%133 = OpFunctionParameter %11 -%134 = OpFunctionParameter %3 +%129 = OpFunction %2 None %130 +%127 = OpFunctionParameter %11 +%128 = OpFunctionParameter %7 +%126 = OpLabel +OpBranch %131 %131 = OpLabel -OpBranch %137 -%137 = OpLabel -%138 = OpExtInst %6 %1 UMin %133 %35 -%139 = OpExtInst %6 %1 UMin %132 %60 -%140 = OpAccessChain %20 %12 %60 %139 %138 -OpStore %140 %134 +%132 = OpExtInst %6 %1 UMin %127 %60 +%133 = OpAccessChain %42 %12 %60 %132 +OpStore %133 %128 OpReturn OpFunctionEnd -%144 = OpFunction %2 None %103 -%142 = OpFunctionParameter %11 -%143 = OpFunctionParameter %3 -%141 = OpLabel -OpBranch %145 -%145 = OpLabel -%146 = OpConvertSToF %3 %142 -%147 = OpFDiv %3 %146 %77 -%148 = OpExtInst %3 %1 Sin %147 -%149 = OpFMul %3 %148 %77 -%150 = OpConvertFToS %11 %149 -%151 = OpExtInst %6 %1 UMin %150 %21 -%152 = OpAccessChain %20 %12 %23 %151 -OpStore %152 %143 +%138 = OpFunction %2 None %139 +%135 = OpFunctionParameter %11 +%136 = OpFunctionParameter %11 +%137 = OpFunctionParameter %3 +%134 = OpLabel +OpBranch %140 +%140 = OpLabel +%141 = OpExtInst %6 %1 UMin %136 %35 +%142 = OpExtInst %6 %1 UMin %135 %60 +%143 = OpAccessChain %20 %12 %60 %142 %141 +OpStore %143 %137 OpReturn OpFunctionEnd -%155 = OpFunction %2 None %156 -%154 = OpFunctionParameter %3 -%153 = OpLabel -OpBranch %157 -%157 = OpLabel -%158 = OpAccessChain %20 %12 %23 %21 -OpStore %158 %154 -%159 = OpAccessChain %20 %12 %32 %35 -OpStore %159 %154 -%160 = OpAccessChain %20 %12 %60 %60 %35 -OpStore %160 %154 +%147 = OpFunction %2 None %106 +%145 = OpFunctionParameter %11 +%146 = OpFunctionParameter %3 +%144 = OpLabel +OpBranch %148 +%148 = OpLabel +%149 = OpConvertSToF %3 %145 +%150 = OpFDiv %3 %149 %77 +%151 = OpExtInst %3 %1 Sin %150 +%152 = OpFMul %3 %151 %77 +%153 = OpExtInst %3 %1 FClamp %152 %83 %84 +%154 = OpConvertFToS %11 %153 +%155 = OpExtInst %6 %1 UMin %154 %21 +%156 = OpAccessChain %20 %12 %23 %155 +OpStore %156 %146 OpReturn OpFunctionEnd -%162 = OpFunction %3 None %89 +%159 = OpFunction %2 None %160 +%158 = OpFunctionParameter %3 +%157 = OpLabel +OpBranch %161 %161 = OpLabel -OpBranch %163 -%163 = OpLabel -%164 = OpArrayLength %6 %12 3 -%165 = OpISub %6 %164 %32 -%167 = OpExtInst %6 %1 UMin %166 %165 -%168 = OpAccessChain %20 %12 %35 %167 -%169 = OpLoad %3 %168 -OpReturnValue %169 +%162 = OpAccessChain %20 %12 %23 %21 +OpStore %162 %158 +%163 = OpAccessChain %20 %12 %32 %35 +OpStore %163 %158 +%164 = OpAccessChain %20 %12 %60 %60 %35 +OpStore %164 %158 +OpReturn +OpFunctionEnd +%166 = OpFunction %3 None %92 +%165 = OpLabel +OpBranch %167 +%167 = OpLabel +%168 = OpArrayLength %6 %12 3 +%169 = OpISub %6 %168 %32 +%171 = OpExtInst %6 %1 UMin %170 %169 +%172 = OpAccessChain %20 %12 %35 %171 +%173 = OpLoad %3 %172 +OpReturnValue %173 OpFunctionEnd -%172 = OpFunction %2 None %156 -%171 = OpFunctionParameter %3 -%170 = OpLabel -OpBranch %173 -%173 = OpLabel -%174 = OpArrayLength %6 %12 3 -%175 = OpISub %6 %174 %32 -%176 = OpExtInst %6 %1 UMin %166 %175 -%177 = OpAccessChain %20 %12 %35 %176 -OpStore %177 %171 +%176 = OpFunction %2 None %160 +%175 = OpFunctionParameter %3 +%174 = OpLabel +OpBranch %177 +%177 = OpLabel +%178 = OpArrayLength %6 %12 3 +%179 = OpISub %6 %178 %32 +%180 = OpExtInst %6 %1 UMin %170 %179 +%181 = OpAccessChain %20 %12 %35 %180 +OpStore %181 %175 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/spv/bounds-check-zero.spvasm b/naga/tests/out/spv/bounds-check-zero.spvasm index 835362b8d3d..3854ef69ca4 100644 --- a/naga/tests/out/spv/bounds-check-zero.spvasm +++ b/naga/tests/out/spv/bounds-check-zero.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 218 +; Bound: 222 OpCapability Shader OpCapability Linkage OpExtension "SPV_KHR_storage_buffer_storage_class" @@ -48,13 +48,15 @@ OpDecorate %12 Binding 0 %76 = OpConstantNull %7 %85 = OpTypeFunction %3 %11 %11 %98 = OpConstant %3 100.0 -%113 = OpTypeFunction %3 -%115 = OpConstant %6 9 -%128 = OpTypeFunction %2 %11 %3 -%157 = OpTypeFunction %2 %11 %7 -%168 = OpTypeFunction %2 %11 %11 %3 -%193 = OpTypeFunction %2 %3 -%202 = OpConstant %6 1000 +%104 = OpConstant %3 -2147483600.0 +%105 = OpConstant %3 2147483500.0 +%116 = OpTypeFunction %3 +%118 = OpConstant %6 9 +%131 = OpTypeFunction %2 %11 %3 +%160 = OpTypeFunction %2 %11 %7 +%171 = OpTypeFunction %2 %11 %11 %3 +%197 = OpTypeFunction %2 %3 +%206 = OpConstant %6 1000 %16 = OpFunction %3 None %17 %15 = OpFunctionParameter %11 %14 = OpLabel @@ -164,179 +166,181 @@ OpBranch %99 %101 = OpFDiv %3 %100 %98 %102 = OpExtInst %3 %1 Sin %101 %103 = OpFMul %3 %102 %98 -%104 = OpConvertFToS %11 %103 -%105 = OpULessThan %22 %104 %5 -OpSelectionMerge %107 None -OpBranchConditional %105 %108 %107 -%108 = OpLabel -%106 = OpAccessChain %20 %12 %23 %104 -%109 = OpLoad %3 %106 -OpBranch %107 -%107 = OpLabel -%110 = OpPhi %3 %25 %99 %109 %108 -OpReturnValue %110 -OpFunctionEnd -%112 = OpFunction %3 None %113 +%106 = OpExtInst %3 %1 FClamp %103 %104 %105 +%107 = OpConvertFToS %11 %106 +%108 = OpULessThan %22 %107 %5 +OpSelectionMerge %110 None +OpBranchConditional %108 %111 %110 %111 = OpLabel -OpBranch %114 +%109 = OpAccessChain %20 %12 %23 %107 +%112 = OpLoad %3 %109 +OpBranch %110 +%110 = OpLabel +%113 = OpPhi %3 %25 %99 %112 %111 +OpReturnValue %113 +OpFunctionEnd +%115 = OpFunction %3 None %116 %114 = OpLabel -%116 = OpAccessChain %20 %12 %23 %115 -%117 = OpLoad %3 %116 -%118 = OpAccessChain %20 %12 %50 %37 -%119 = OpLoad %3 %118 -%120 = OpFAdd %3 %117 %119 -%121 = OpAccessChain %20 %12 %74 %74 %37 +OpBranch %117 +%117 = OpLabel +%119 = OpAccessChain %20 %12 %23 %118 +%120 = OpLoad %3 %119 +%121 = OpAccessChain %20 %12 %50 %37 %122 = OpLoad %3 %121 %123 = OpFAdd %3 %120 %122 -OpReturnValue %123 +%124 = OpAccessChain %20 %12 %74 %74 %37 +%125 = OpLoad %3 %124 +%126 = OpFAdd %3 %123 %125 +OpReturnValue %126 OpFunctionEnd -%127 = OpFunction %2 None %128 -%125 = OpFunctionParameter %11 -%126 = OpFunctionParameter %3 -%124 = OpLabel -OpBranch %129 -%129 = OpLabel -%130 = OpULessThan %22 %125 %5 -OpSelectionMerge %132 None -OpBranchConditional %130 %133 %132 -%133 = OpLabel -%131 = OpAccessChain %20 %12 %23 %125 -OpStore %131 %126 +%130 = OpFunction %2 None %131 +%128 = OpFunctionParameter %11 +%129 = OpFunctionParameter %3 +%127 = OpLabel OpBranch %132 %132 = OpLabel +%133 = OpULessThan %22 %128 %5 +OpSelectionMerge %135 None +OpBranchConditional %133 %136 %135 +%136 = OpLabel +%134 = OpAccessChain %20 %12 %23 %128 +OpStore %134 %129 +OpBranch %135 +%135 = OpLabel OpReturn OpFunctionEnd -%137 = OpFunction %2 None %128 -%135 = OpFunctionParameter %11 -%136 = OpFunctionParameter %3 -%134 = OpLabel -OpBranch %138 -%138 = OpLabel -%139 = OpArrayLength %6 %12 3 -%140 = OpULessThan %22 %135 %139 -OpSelectionMerge %142 None -OpBranchConditional %140 %143 %142 -%143 = OpLabel -%141 = OpAccessChain %20 %12 %37 %135 -OpStore %141 %136 -OpBranch %142 -%142 = OpLabel +%140 = OpFunction %2 None %131 +%138 = OpFunctionParameter %11 +%139 = OpFunctionParameter %3 +%137 = OpLabel +OpBranch %141 +%141 = OpLabel +%142 = OpArrayLength %6 %12 3 +%143 = OpULessThan %22 %138 %142 +OpSelectionMerge %145 None +OpBranchConditional %143 %146 %145 +%146 = OpLabel +%144 = OpAccessChain %20 %12 %37 %138 +OpStore %144 %139 +OpBranch %145 +%145 = OpLabel OpReturn OpFunctionEnd -%147 = OpFunction %2 None %128 -%145 = OpFunctionParameter %11 -%146 = OpFunctionParameter %3 -%144 = OpLabel -OpBranch %148 -%148 = OpLabel -%149 = OpULessThan %22 %145 %48 -OpSelectionMerge %151 None -OpBranchConditional %149 %152 %151 -%152 = OpLabel -%150 = OpAccessChain %20 %12 %50 %145 -OpStore %150 %146 +%150 = OpFunction %2 None %131 +%148 = OpFunctionParameter %11 +%149 = OpFunctionParameter %3 +%147 = OpLabel OpBranch %151 %151 = OpLabel +%152 = OpULessThan %22 %148 %48 +OpSelectionMerge %154 None +OpBranchConditional %152 %155 %154 +%155 = OpLabel +%153 = OpAccessChain %20 %12 %50 %148 +OpStore %153 %149 +OpBranch %154 +%154 = OpLabel OpReturn OpFunctionEnd -%156 = OpFunction %2 None %157 -%154 = OpFunctionParameter %11 -%155 = OpFunctionParameter %7 -%153 = OpLabel -OpBranch %158 -%158 = OpLabel -%159 = OpULessThan %22 %154 %37 -OpSelectionMerge %161 None -OpBranchConditional %159 %162 %161 -%162 = OpLabel -%160 = OpAccessChain %47 %12 %74 %154 -OpStore %160 %155 +%159 = OpFunction %2 None %160 +%157 = OpFunctionParameter %11 +%158 = OpFunctionParameter %7 +%156 = OpLabel OpBranch %161 %161 = OpLabel +%162 = OpULessThan %22 %157 %37 +OpSelectionMerge %164 None +OpBranchConditional %162 %165 %164 +%165 = OpLabel +%163 = OpAccessChain %47 %12 %74 %157 +OpStore %163 %158 +OpBranch %164 +%164 = OpLabel OpReturn OpFunctionEnd -%167 = OpFunction %2 None %168 -%164 = OpFunctionParameter %11 -%165 = OpFunctionParameter %11 -%166 = OpFunctionParameter %3 -%163 = OpLabel -OpBranch %169 -%169 = OpLabel -%170 = OpULessThan %22 %165 %48 -%171 = OpULessThan %22 %164 %37 -%172 = OpLogicalAnd %22 %170 %171 -OpSelectionMerge %174 None -OpBranchConditional %172 %175 %174 -%175 = OpLabel -%173 = OpAccessChain %20 %12 %74 %164 %165 -OpStore %173 %166 -OpBranch %174 -%174 = OpLabel +%170 = OpFunction %2 None %171 +%167 = OpFunctionParameter %11 +%168 = OpFunctionParameter %11 +%169 = OpFunctionParameter %3 +%166 = OpLabel +OpBranch %172 +%172 = OpLabel +%173 = OpULessThan %22 %168 %48 +%174 = OpULessThan %22 %167 %37 +%175 = OpLogicalAnd %22 %173 %174 +OpSelectionMerge %177 None +OpBranchConditional %175 %178 %177 +%178 = OpLabel +%176 = OpAccessChain %20 %12 %74 %167 %168 +OpStore %176 %169 +OpBranch %177 +%177 = OpLabel OpReturn OpFunctionEnd -%179 = OpFunction %2 None %128 -%177 = OpFunctionParameter %11 -%178 = OpFunctionParameter %3 -%176 = OpLabel -OpBranch %180 -%180 = OpLabel -%181 = OpConvertSToF %3 %177 -%182 = OpFDiv %3 %181 %98 -%183 = OpExtInst %3 %1 Sin %182 -%184 = OpFMul %3 %183 %98 -%185 = OpConvertFToS %11 %184 -%186 = OpULessThan %22 %185 %5 -OpSelectionMerge %188 None -OpBranchConditional %186 %189 %188 -%189 = OpLabel -%187 = OpAccessChain %20 %12 %23 %185 -OpStore %187 %178 -OpBranch %188 -%188 = OpLabel +%182 = OpFunction %2 None %131 +%180 = OpFunctionParameter %11 +%181 = OpFunctionParameter %3 +%179 = OpLabel +OpBranch %183 +%183 = OpLabel +%184 = OpConvertSToF %3 %180 +%185 = OpFDiv %3 %184 %98 +%186 = OpExtInst %3 %1 Sin %185 +%187 = OpFMul %3 %186 %98 +%188 = OpExtInst %3 %1 FClamp %187 %104 %105 +%189 = OpConvertFToS %11 %188 +%190 = OpULessThan %22 %189 %5 +OpSelectionMerge %192 None +OpBranchConditional %190 %193 %192 +%193 = OpLabel +%191 = OpAccessChain %20 %12 %23 %189 +OpStore %191 %181 +OpBranch %192 +%192 = OpLabel OpReturn OpFunctionEnd -%192 = OpFunction %2 None %193 -%191 = OpFunctionParameter %3 -%190 = OpLabel -OpBranch %194 +%196 = OpFunction %2 None %197 +%195 = OpFunctionParameter %3 %194 = OpLabel -%195 = OpAccessChain %20 %12 %23 %115 -OpStore %195 %191 -%196 = OpAccessChain %20 %12 %50 %37 -OpStore %196 %191 -%197 = OpAccessChain %20 %12 %74 %74 %37 -OpStore %197 %191 -OpReturn -OpFunctionEnd -%199 = OpFunction %3 None %113 +OpBranch %198 %198 = OpLabel -OpBranch %200 -%200 = OpLabel -%201 = OpArrayLength %6 %12 3 -%203 = OpULessThan %22 %202 %201 -OpSelectionMerge %205 None -OpBranchConditional %203 %206 %205 -%206 = OpLabel -%204 = OpAccessChain %20 %12 %37 %202 -%207 = OpLoad %3 %204 -OpBranch %205 -%205 = OpLabel -%208 = OpPhi %3 %25 %200 %207 %206 -OpReturnValue %208 +%199 = OpAccessChain %20 %12 %23 %118 +OpStore %199 %195 +%200 = OpAccessChain %20 %12 %50 %37 +OpStore %200 %195 +%201 = OpAccessChain %20 %12 %74 %74 %37 +OpStore %201 %195 +OpReturn OpFunctionEnd -%211 = OpFunction %2 None %193 -%210 = OpFunctionParameter %3 +%203 = OpFunction %3 None %116 +%202 = OpLabel +OpBranch %204 +%204 = OpLabel +%205 = OpArrayLength %6 %12 3 +%207 = OpULessThan %22 %206 %205 +OpSelectionMerge %209 None +OpBranchConditional %207 %210 %209 +%210 = OpLabel +%208 = OpAccessChain %20 %12 %37 %206 +%211 = OpLoad %3 %208 +OpBranch %209 %209 = OpLabel -OpBranch %212 -%212 = OpLabel -%213 = OpArrayLength %6 %12 3 -%214 = OpULessThan %22 %202 %213 -OpSelectionMerge %216 None -OpBranchConditional %214 %217 %216 -%217 = OpLabel -%215 = OpAccessChain %20 %12 %37 %202 -OpStore %215 %210 +%212 = OpPhi %3 %25 %204 %211 %210 +OpReturnValue %212 +OpFunctionEnd +%215 = OpFunction %2 None %197 +%214 = OpFunctionParameter %3 +%213 = OpLabel OpBranch %216 %216 = OpLabel +%217 = OpArrayLength %6 %12 3 +%218 = OpULessThan %22 %206 %217 +OpSelectionMerge %220 None +OpBranchConditional %218 %221 %220 +%221 = OpLabel +%219 = OpAccessChain %20 %12 %37 %206 +OpStore %219 %214 +OpBranch %220 +%220 = OpLabel OpReturn OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/spv/conversion-float-to-int.spvasm b/naga/tests/out/spv/conversion-float-to-int.spvasm new file mode 100644 index 00000000000..5b523c5989b --- /dev/null +++ b/naga/tests/out/spv/conversion-float-to-int.spvasm @@ -0,0 +1,372 @@ +; SPIR-V +; Version: 1.1 +; Generator: rspirv +; Bound: 280 +OpCapability Shader +OpCapability Float16 +OpCapability StorageBuffer16BitAccess +OpCapability UniformAndStorageBuffer16BitAccess +OpCapability StorageInputOutput16 +OpCapability Float64 +OpCapability Int64 +OpExtension "SPV_KHR_16bit_storage" +%1 = OpExtInstImport "GLSL.std.450" +OpMemoryModel Logical GLSL450 +OpEntryPoint GLCompute %278 "main" +OpExecutionMode %278 LocalSize 1 1 1 +%2 = OpTypeVoid +%3 = OpTypeFloat 16 +%4 = OpTypeFloat 32 +%5 = OpTypeFloat 64 +%6 = OpTypeInt 32 1 +%7 = OpTypeInt 32 0 +%8 = OpTypeInt 64 1 +%9 = OpTypeInt 64 0 +%10 = OpTypeVector %3 2 +%11 = OpTypeVector %6 2 +%12 = OpTypeVector %7 2 +%13 = OpTypeVector %8 2 +%14 = OpTypeVector %9 2 +%15 = OpTypeVector %4 2 +%16 = OpTypeVector %5 2 +%17 = OpConstant %3 9.0399e-41 +%18 = OpConstant %3 4.4481e-41 +%19 = OpConstant %4 -3.4028235e38 +%20 = OpConstant %4 3.4028235e38 +%21 = OpConstant %5 -1.7976931348623157e308 +%22 = OpConstant %5 1.7976931348623157e308 +%25 = OpTypeFunction %2 +%26 = OpConstant %6 -65504 +%27 = OpConstant %6 65504 +%28 = OpConstant %7 0 +%29 = OpConstant %7 65504 +%30 = OpConstant %8 -65504 +%31 = OpConstant %8 65504 +%32 = OpConstant %9 0 +%33 = OpConstant %9 65504 +%34 = OpConstant %6 -2147483648 +%35 = OpConstant %6 2147483520 +%36 = OpConstant %7 4294967040 +%37 = OpConstant %8 -9223372036854775808 +%38 = OpConstant %8 9223371487098961920 +%39 = OpConstant %9 18446742974197923840 +%40 = OpConstant %8 9223372036854774784 +%41 = OpConstant %9 18446744073709549568 +%42 = OpConstant %6 2147483647 +%43 = OpConstant %7 4294967295 +%45 = OpTypePointer Function %6 +%48 = OpTypePointer Function %7 +%51 = OpTypePointer Function %8 +%54 = OpTypePointer Function %9 +%80 = OpTypeFunction %6 %3 +%87 = OpTypeFunction %7 %3 +%89 = OpConstant %3 0.0 +%95 = OpTypeFunction %8 %3 +%102 = OpTypeFunction %9 %3 +%109 = OpTypeFunction %6 %4 +%111 = OpConstant %4 -2147483600.0 +%112 = OpConstant %4 2147483500.0 +%118 = OpTypeFunction %7 %4 +%120 = OpConstant %4 0.0 +%121 = OpConstant %4 4294967000.0 +%127 = OpTypeFunction %8 %4 +%129 = OpConstant %4 -9.223372e18 +%130 = OpConstant %4 9.2233715e18 +%136 = OpTypeFunction %9 %4 +%138 = OpConstant %4 1.8446743e19 +%144 = OpTypeFunction %6 %5 +%146 = OpConstant %5 -2147483648.0 +%147 = OpConstant %5 2147483647.0 +%153 = OpTypeFunction %7 %5 +%155 = OpConstant %5 0.0 +%156 = OpConstant %5 4294967295.0 +%162 = OpTypeFunction %8 %5 +%164 = OpConstant %5 -9.223372036854776e18 +%165 = OpConstant %5 9.223372036854775e18 +%171 = OpTypeFunction %9 %5 +%173 = OpConstant %5 1.844674407370955e19 +%179 = OpTypeFunction %11 %10 +%181 = OpConstantComposite %10 %17 %17 +%182 = OpConstantComposite %10 %18 %18 +%188 = OpTypeFunction %12 %10 +%190 = OpConstantComposite %10 %89 %89 +%196 = OpTypeFunction %13 %10 +%203 = OpTypeFunction %14 %10 +%210 = OpTypeFunction %11 %15 +%212 = OpConstantComposite %15 %111 %111 +%213 = OpConstantComposite %15 %112 %112 +%219 = OpTypeFunction %12 %15 +%221 = OpConstantComposite %15 %120 %120 +%222 = OpConstantComposite %15 %121 %121 +%228 = OpTypeFunction %13 %15 +%230 = OpConstantComposite %15 %129 %129 +%231 = OpConstantComposite %15 %130 %130 +%237 = OpTypeFunction %14 %15 +%239 = OpConstantComposite %15 %138 %138 +%245 = OpTypeFunction %11 %16 +%247 = OpConstantComposite %16 %146 %146 +%248 = OpConstantComposite %16 %147 %147 +%254 = OpTypeFunction %12 %16 +%256 = OpConstantComposite %16 %155 %155 +%257 = OpConstantComposite %16 %156 %156 +%263 = OpTypeFunction %13 %16 +%265 = OpConstantComposite %16 %164 %164 +%266 = OpConstantComposite %16 %165 %165 +%272 = OpTypeFunction %14 %16 +%274 = OpConstantComposite %16 %173 %173 +%24 = OpFunction %2 None %25 +%23 = OpLabel +%73 = OpVariable %51 Function %40 +%70 = OpVariable %48 Function %28 +%67 = OpVariable %54 Function %41 +%64 = OpVariable %51 Function %37 +%61 = OpVariable %51 Function %38 +%58 = OpVariable %48 Function %28 +%55 = OpVariable %54 Function %33 +%50 = OpVariable %51 Function %30 +%46 = OpVariable %45 Function %27 +%75 = OpVariable %54 Function %41 +%72 = OpVariable %51 Function %37 +%69 = OpVariable %45 Function %42 +%66 = OpVariable %54 Function %32 +%63 = OpVariable %54 Function %39 +%60 = OpVariable %51 Function %37 +%57 = OpVariable %45 Function %35 +%53 = OpVariable %54 Function %32 +%49 = OpVariable %48 Function %29 +%44 = OpVariable %45 Function %26 +%74 = OpVariable %54 Function %32 +%71 = OpVariable %48 Function %43 +%68 = OpVariable %45 Function %34 +%65 = OpVariable %51 Function %40 +%62 = OpVariable %54 Function %32 +%59 = OpVariable %48 Function %36 +%56 = OpVariable %45 Function %34 +%52 = OpVariable %51 Function %31 +%47 = OpVariable %48 Function %28 +OpBranch %76 +%76 = OpLabel +OpReturn +OpFunctionEnd +%79 = OpFunction %6 None %80 +%78 = OpFunctionParameter %3 +%77 = OpLabel +OpBranch %81 +%81 = OpLabel +%82 = OpExtInst %3 %1 FClamp %78 %17 %18 +%83 = OpConvertFToS %6 %82 +OpReturnValue %83 +OpFunctionEnd +%86 = OpFunction %7 None %87 +%85 = OpFunctionParameter %3 +%84 = OpLabel +OpBranch %88 +%88 = OpLabel +%90 = OpExtInst %3 %1 FClamp %85 %89 %18 +%91 = OpConvertFToU %7 %90 +OpReturnValue %91 +OpFunctionEnd +%94 = OpFunction %8 None %95 +%93 = OpFunctionParameter %3 +%92 = OpLabel +OpBranch %96 +%96 = OpLabel +%97 = OpExtInst %3 %1 FClamp %93 %17 %18 +%98 = OpConvertFToS %8 %97 +OpReturnValue %98 +OpFunctionEnd +%101 = OpFunction %9 None %102 +%100 = OpFunctionParameter %3 +%99 = OpLabel +OpBranch %103 +%103 = OpLabel +%104 = OpExtInst %3 %1 FClamp %100 %89 %18 +%105 = OpConvertFToU %9 %104 +OpReturnValue %105 +OpFunctionEnd +%108 = OpFunction %6 None %109 +%107 = OpFunctionParameter %4 +%106 = OpLabel +OpBranch %110 +%110 = OpLabel +%113 = OpExtInst %4 %1 FClamp %107 %111 %112 +%114 = OpConvertFToS %6 %113 +OpReturnValue %114 +OpFunctionEnd +%117 = OpFunction %7 None %118 +%116 = OpFunctionParameter %4 +%115 = OpLabel +OpBranch %119 +%119 = OpLabel +%122 = OpExtInst %4 %1 FClamp %116 %120 %121 +%123 = OpConvertFToU %7 %122 +OpReturnValue %123 +OpFunctionEnd +%126 = OpFunction %8 None %127 +%125 = OpFunctionParameter %4 +%124 = OpLabel +OpBranch %128 +%128 = OpLabel +%131 = OpExtInst %4 %1 FClamp %125 %129 %130 +%132 = OpConvertFToS %8 %131 +OpReturnValue %132 +OpFunctionEnd +%135 = OpFunction %9 None %136 +%134 = OpFunctionParameter %4 +%133 = OpLabel +OpBranch %137 +%137 = OpLabel +%139 = OpExtInst %4 %1 FClamp %134 %120 %138 +%140 = OpConvertFToU %9 %139 +OpReturnValue %140 +OpFunctionEnd +%143 = OpFunction %6 None %144 +%142 = OpFunctionParameter %5 +%141 = OpLabel +OpBranch %145 +%145 = OpLabel +%148 = OpExtInst %5 %1 FClamp %142 %146 %147 +%149 = OpConvertFToS %6 %148 +OpReturnValue %149 +OpFunctionEnd +%152 = OpFunction %7 None %153 +%151 = OpFunctionParameter %5 +%150 = OpLabel +OpBranch %154 +%154 = OpLabel +%157 = OpExtInst %5 %1 FClamp %151 %155 %156 +%158 = OpConvertFToU %7 %157 +OpReturnValue %158 +OpFunctionEnd +%161 = OpFunction %8 None %162 +%160 = OpFunctionParameter %5 +%159 = OpLabel +OpBranch %163 +%163 = OpLabel +%166 = OpExtInst %5 %1 FClamp %160 %164 %165 +%167 = OpConvertFToS %8 %166 +OpReturnValue %167 +OpFunctionEnd +%170 = OpFunction %9 None %171 +%169 = OpFunctionParameter %5 +%168 = OpLabel +OpBranch %172 +%172 = OpLabel +%174 = OpExtInst %5 %1 FClamp %169 %155 %173 +%175 = OpConvertFToU %9 %174 +OpReturnValue %175 +OpFunctionEnd +%178 = OpFunction %11 None %179 +%177 = OpFunctionParameter %10 +%176 = OpLabel +OpBranch %180 +%180 = OpLabel +%183 = OpExtInst %10 %1 FClamp %177 %181 %182 +%184 = OpConvertFToS %11 %183 +OpReturnValue %184 +OpFunctionEnd +%187 = OpFunction %12 None %188 +%186 = OpFunctionParameter %10 +%185 = OpLabel +OpBranch %189 +%189 = OpLabel +%191 = OpExtInst %10 %1 FClamp %186 %190 %182 +%192 = OpConvertFToU %12 %191 +OpReturnValue %192 +OpFunctionEnd +%195 = OpFunction %13 None %196 +%194 = OpFunctionParameter %10 +%193 = OpLabel +OpBranch %197 +%197 = OpLabel +%198 = OpExtInst %10 %1 FClamp %194 %181 %182 +%199 = OpConvertFToS %13 %198 +OpReturnValue %199 +OpFunctionEnd +%202 = OpFunction %14 None %203 +%201 = OpFunctionParameter %10 +%200 = OpLabel +OpBranch %204 +%204 = OpLabel +%205 = OpExtInst %10 %1 FClamp %201 %190 %182 +%206 = OpConvertFToU %14 %205 +OpReturnValue %206 +OpFunctionEnd +%209 = OpFunction %11 None %210 +%208 = OpFunctionParameter %15 +%207 = OpLabel +OpBranch %211 +%211 = OpLabel +%214 = OpExtInst %15 %1 FClamp %208 %212 %213 +%215 = OpConvertFToS %11 %214 +OpReturnValue %215 +OpFunctionEnd +%218 = OpFunction %12 None %219 +%217 = OpFunctionParameter %15 +%216 = OpLabel +OpBranch %220 +%220 = OpLabel +%223 = OpExtInst %15 %1 FClamp %217 %221 %222 +%224 = OpConvertFToU %12 %223 +OpReturnValue %224 +OpFunctionEnd +%227 = OpFunction %13 None %228 +%226 = OpFunctionParameter %15 +%225 = OpLabel +OpBranch %229 +%229 = OpLabel +%232 = OpExtInst %15 %1 FClamp %226 %230 %231 +%233 = OpConvertFToS %13 %232 +OpReturnValue %233 +OpFunctionEnd +%236 = OpFunction %14 None %237 +%235 = OpFunctionParameter %15 +%234 = OpLabel +OpBranch %238 +%238 = OpLabel +%240 = OpExtInst %15 %1 FClamp %235 %221 %239 +%241 = OpConvertFToU %14 %240 +OpReturnValue %241 +OpFunctionEnd +%244 = OpFunction %11 None %245 +%243 = OpFunctionParameter %16 +%242 = OpLabel +OpBranch %246 +%246 = OpLabel +%249 = OpExtInst %16 %1 FClamp %243 %247 %248 +%250 = OpConvertFToS %11 %249 +OpReturnValue %250 +OpFunctionEnd +%253 = OpFunction %12 None %254 +%252 = OpFunctionParameter %16 +%251 = OpLabel +OpBranch %255 +%255 = OpLabel +%258 = OpExtInst %16 %1 FClamp %252 %256 %257 +%259 = OpConvertFToU %12 %258 +OpReturnValue %259 +OpFunctionEnd +%262 = OpFunction %13 None %263 +%261 = OpFunctionParameter %16 +%260 = OpLabel +OpBranch %264 +%264 = OpLabel +%267 = OpExtInst %16 %1 FClamp %261 %265 %266 +%268 = OpConvertFToS %13 %267 +OpReturnValue %268 +OpFunctionEnd +%271 = OpFunction %14 None %272 +%270 = OpFunctionParameter %16 +%269 = OpLabel +OpBranch %273 +%273 = OpLabel +%275 = OpExtInst %16 %1 FClamp %270 %256 %274 +%276 = OpConvertFToU %14 %275 +OpReturnValue %276 +OpFunctionEnd +%278 = OpFunction %2 None %25 +%277 = OpLabel +OpBranch %279 +%279 = OpLabel +OpReturn +OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/spv/debug-symbol-large-source.spvasm b/naga/tests/out/spv/debug-symbol-large-source.spvasm index 5b61cb96085..ee6ed420922 100644 --- a/naga/tests/out/spv/debug-symbol-large-source.spvasm +++ b/naga/tests/out/spv/debug-symbol-large-source.spvasm @@ -1,19 +1,19 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 672 +; Bound: 676 OpCapability Shader OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %367 "gen_terrain_compute" %364 OpEntryPoint Vertex %444 "gen_terrain_vertex" %435 %438 %440 %442 -OpEntryPoint Fragment %493 "gen_terrain_fragment" %483 %485 %488 %491 %492 -OpEntryPoint Vertex %586 "vs_main" %577 %580 %582 %583 %585 -OpEntryPoint Fragment %611 "fs_main" %604 %606 %608 %610 +OpEntryPoint Fragment %495 "gen_terrain_fragment" %485 %487 %490 %493 %494 +OpEntryPoint Vertex %590 "vs_main" %581 %584 %586 %587 %589 +OpEntryPoint Fragment %615 "fs_main" %608 %610 %612 %614 OpExecutionMode %367 LocalSize 64 1 1 -OpExecutionMode %493 OriginUpperLeft -OpExecutionMode %611 OriginUpperLeft +OpExecutionMode %495 OriginUpperLeft +OpExecutionMode %615 OriginUpperLeft %3 = OpString "debug-symbol-large-source.wgsl" OpSource Unknown 0 %3 "//This is a test 这是测试 これはテストだ 테스트입니다 This is a test 这是测试 これはテストだ 테스트입니다 This is a test 这是测试 これはテストだ 테스트입니다 //This is a test 这是测试 これはテストだ 테스트입니다 This is a test 这是测试 これはテストだ 테스트입니다 This is a test 这是测试 これはテストだ 테스트입니다 @@ -7561,25 +7561,25 @@ OpName %438 "index" OpName %440 "position" OpName %442 "uv" OpName %444 "gen_terrain_vertex" -OpName %483 "index" -OpName %485 "position" -OpName %488 "uv" -OpName %491 "vert_component" -OpName %492 "index" -OpName %493 "gen_terrain_fragment" -OpName %496 "vert_component" -OpName %497 "index" -OpName %577 "position" -OpName %580 "normal" -OpName %582 "clip_position" -OpName %583 "normal" -OpName %585 "world_pos" -OpName %586 "vs_main" -OpName %604 "clip_position" -OpName %606 "normal" -OpName %608 "world_pos" -OpName %611 "fs_main" -OpName %620 "color" +OpName %485 "index" +OpName %487 "position" +OpName %490 "uv" +OpName %493 "vert_component" +OpName %494 "index" +OpName %495 "gen_terrain_fragment" +OpName %498 "vert_component" +OpName %499 "index" +OpName %581 "position" +OpName %584 "normal" +OpName %586 "clip_position" +OpName %587 "normal" +OpName %589 "world_pos" +OpName %590 "vs_main" +OpName %608 "clip_position" +OpName %610 "normal" +OpName %612 "world_pos" +OpName %615 "fs_main" +OpName %624 "color" OpMemberDecorate %13 0 Offset 0 OpMemberDecorate %13 1 Offset 8 OpMemberDecorate %13 2 Offset 16 @@ -7644,21 +7644,21 @@ OpDecorate %438 Location 0 OpDecorate %438 Flat OpDecorate %440 BuiltIn Position OpDecorate %442 Location 1 -OpDecorate %483 Location 0 -OpDecorate %483 Flat -OpDecorate %485 BuiltIn FragCoord -OpDecorate %488 Location 1 -OpDecorate %491 Location 0 -OpDecorate %492 Location 1 -OpDecorate %577 Location 0 -OpDecorate %580 Location 1 -OpDecorate %582 BuiltIn Position -OpDecorate %583 Location 0 -OpDecorate %585 Location 1 -OpDecorate %604 BuiltIn FragCoord -OpDecorate %606 Location 0 -OpDecorate %608 Location 1 +OpDecorate %485 Location 0 +OpDecorate %485 Flat +OpDecorate %487 BuiltIn FragCoord +OpDecorate %490 Location 1 +OpDecorate %493 Location 0 +OpDecorate %494 Location 1 +OpDecorate %581 Location 0 +OpDecorate %584 Location 1 +OpDecorate %586 BuiltIn Position +OpDecorate %587 Location 0 +OpDecorate %589 Location 1 +OpDecorate %608 BuiltIn FragCoord OpDecorate %610 Location 0 +OpDecorate %612 Location 1 +OpDecorate %614 Location 0 %2 = OpTypeVoid %4 = OpTypeFloat 32 %5 = OpTypeVector %4 3 @@ -7800,36 +7800,37 @@ OpDecorate %610 Location 0 %445 = OpTypePointer Uniform %20 %447 = OpConstant %4 -1.0 %448 = OpConstantComposite %6 %447 %447 -%483 = OpVariable %436 Input -%486 = OpTypePointer Input %7 -%485 = OpVariable %486 Input -%489 = OpTypePointer Input %6 -%488 = OpVariable %489 Input -%491 = OpVariable %439 Output -%492 = OpVariable %439 Output -%495 = OpConstant %4 6.0 -%578 = OpTypePointer Input %5 -%577 = OpVariable %578 Input -%580 = OpVariable %578 Input -%582 = OpVariable %441 Output -%584 = OpTypePointer Output %5 -%583 = OpVariable %584 Output -%585 = OpVariable %584 Output -%587 = OpTypePointer Uniform %24 -%590 = OpTypePointer Uniform %23 -%604 = OpVariable %486 Input -%606 = OpVariable %578 Input -%608 = OpVariable %578 Input -%610 = OpVariable %441 Output -%613 = OpTypePointer Uniform %25 -%615 = OpConstantComposite %5 %286 %286 %286 -%616 = OpConstant %4 0.7 -%617 = OpConstantComposite %5 %78 %286 %616 -%618 = OpConstant %4 0.2 -%619 = OpConstantComposite %5 %618 %618 %618 -%621 = OpConstantNull %5 -%636 = OpTypePointer Uniform %5 -%645 = OpTypePointer Uniform %7 +%473 = OpConstant %4 4294967000.0 +%485 = OpVariable %436 Input +%488 = OpTypePointer Input %7 +%487 = OpVariable %488 Input +%491 = OpTypePointer Input %6 +%490 = OpVariable %491 Input +%493 = OpVariable %439 Output +%494 = OpVariable %439 Output +%497 = OpConstant %4 6.0 +%582 = OpTypePointer Input %5 +%581 = OpVariable %582 Input +%584 = OpVariable %582 Input +%586 = OpVariable %441 Output +%588 = OpTypePointer Output %5 +%587 = OpVariable %588 Output +%589 = OpVariable %588 Output +%591 = OpTypePointer Uniform %24 +%594 = OpTypePointer Uniform %23 +%608 = OpVariable %488 Input +%610 = OpVariable %582 Input +%612 = OpVariable %582 Input +%614 = OpVariable %441 Output +%617 = OpTypePointer Uniform %25 +%619 = OpConstantComposite %5 %286 %286 %286 +%620 = OpConstant %4 0.7 +%621 = OpConstantComposite %5 %78 %286 %620 +%622 = OpConstant %4 0.2 +%623 = OpConstantComposite %5 %622 %622 %622 +%625 = OpConstantNull %5 +%640 = OpTypePointer Uniform %5 +%649 = OpTypePointer Uniform %7 %53 = OpFunction %5 None %54 %52 = OpFunctionParameter %5 %51 = OpLabel @@ -8383,286 +8384,289 @@ OpLine %3 7311 17 %470 = OpConvertUToF %4 %469 %471 = OpFMul %4 %467 %470 %472 = OpFAdd %4 %466 %471 -%473 = OpConvertFToU %8 %472 +%474 = OpExtInst %4 %1 FClamp %472 %74 %473 +%475 = OpConvertFToU %8 %474 OpLine %3 7311 17 -%474 = OpAccessChain %393 %446 %374 -%475 = OpLoad %8 %474 -%476 = OpIAdd %8 %473 %475 +%476 = OpAccessChain %393 %446 %374 +%477 = OpLoad %8 %476 +%478 = OpIAdd %8 %475 %477 OpLine %3 7313 12 -%477 = OpCompositeConstruct %21 %476 %461 %458 -%478 = OpCompositeExtract %8 %477 0 -OpStore %438 %478 -%479 = OpCompositeExtract %7 %477 1 -OpStore %440 %479 -%480 = OpCompositeExtract %6 %477 2 -OpStore %442 %480 +%479 = OpCompositeConstruct %21 %478 %461 %458 +%480 = OpCompositeExtract %8 %479 0 +OpStore %438 %480 +%481 = OpCompositeExtract %7 %479 1 +OpStore %440 %481 +%482 = OpCompositeExtract %6 %479 2 +OpStore %442 %482 OpReturn OpFunctionEnd -%493 = OpFunction %2 None %368 -%481 = OpLabel -%496 = OpVariable %125 Function %74 -%497 = OpVariable %217 Function %135 -%484 = OpLoad %8 %483 -%487 = OpLoad %7 %485 -%490 = OpLoad %6 %488 -%482 = OpCompositeConstruct %21 %484 %487 %490 -%494 = OpAccessChain %445 %36 %135 -OpBranch %498 -%498 = OpLabel +%495 = OpFunction %2 None %368 +%483 = OpLabel +%498 = OpVariable %125 Function %74 +%499 = OpVariable %217 Function %135 +%486 = OpLoad %8 %485 +%489 = OpLoad %7 %487 +%492 = OpLoad %6 %490 +%484 = OpCompositeConstruct %21 %486 %489 %492 +%496 = OpAccessChain %445 %36 %135 +OpBranch %500 +%500 = OpLabel OpLine %3 7324 17 -%499 = OpCompositeExtract %6 %482 2 -%500 = OpCompositeExtract %4 %499 0 +%501 = OpCompositeExtract %6 %484 2 +%502 = OpCompositeExtract %4 %501 0 OpLine %3 7324 17 -%501 = OpAccessChain %393 %494 %373 -%502 = OpLoad %8 %501 -%503 = OpConvertUToF %4 %502 -%504 = OpFMul %4 %500 %503 -%505 = OpCompositeExtract %6 %482 2 -%506 = OpCompositeExtract %4 %505 1 +%503 = OpAccessChain %393 %496 %373 +%504 = OpLoad %8 %503 +%505 = OpConvertUToF %4 %504 +%506 = OpFMul %4 %502 %505 +%507 = OpCompositeExtract %6 %484 2 +%508 = OpCompositeExtract %4 %507 1 OpLine %3 7324 70 -%507 = OpAccessChain %393 %494 %373 -%508 = OpLoad %8 %507 -OpLine %3 7324 13 -%509 = OpAccessChain %393 %494 %373 +%509 = OpAccessChain %393 %496 %373 %510 = OpLoad %8 %509 -%511 = OpIMul %8 %508 %510 -%512 = OpConvertUToF %4 %511 -%513 = OpFMul %4 %506 %512 -%514 = OpFAdd %4 %504 %513 -%515 = OpConvertFToU %8 %514 OpLine %3 7324 13 -%516 = OpAccessChain %393 %494 %374 -%517 = OpLoad %8 %516 -%518 = OpIAdd %8 %515 %517 +%511 = OpAccessChain %393 %496 %373 +%512 = OpLoad %8 %511 +%513 = OpIMul %8 %510 %512 +%514 = OpConvertUToF %4 %513 +%515 = OpFMul %4 %508 %514 +%516 = OpFAdd %4 %506 %515 +%517 = OpExtInst %4 %1 FClamp %516 %74 %473 +%518 = OpConvertFToU %8 %517 +OpLine %3 7324 13 +%519 = OpAccessChain %393 %496 %374 +%520 = OpLoad %8 %519 +%521 = OpIAdd %8 %518 %520 OpLine %3 7325 32 -%519 = OpConvertUToF %4 %518 +%522 = OpConvertUToF %4 %521 OpLine %3 7325 22 -%520 = OpFDiv %4 %519 %495 -%521 = OpExtInst %4 %1 Floor %520 -%522 = OpConvertFToU %8 %521 +%523 = OpFDiv %4 %522 %497 +%524 = OpExtInst %4 %1 Floor %523 +%525 = OpExtInst %4 %1 FClamp %524 %74 %473 +%526 = OpConvertFToU %8 %525 OpLine %3 7326 22 -%523 = OpFunctionCall %8 %427 %518 %371 +%527 = OpFunctionCall %8 %427 %521 %371 OpLine %3 7328 36 -%524 = OpAccessChain %377 %494 %135 -%525 = OpLoad %10 %524 +%528 = OpAccessChain %377 %496 %135 +%529 = OpLoad %10 %528 OpLine %3 7328 57 -%526 = OpAccessChain %380 %494 %126 -%527 = OpLoad %11 %526 +%530 = OpAccessChain %380 %496 %126 +%531 = OpLoad %11 %530 OpLine %3 7328 13 -%528 = OpFunctionCall %6 %325 %522 %525 %527 +%532 = OpFunctionCall %6 %325 %526 %529 %531 OpLine %3 7329 31 -%529 = OpAccessChain %386 %494 %372 -%530 = OpLoad %6 %529 +%533 = OpAccessChain %386 %496 %372 +%534 = OpLoad %6 %533 OpLine %3 7329 13 -%531 = OpFunctionCall %14 %284 %528 %530 +%535 = OpFunctionCall %14 %284 %532 %534 OpLine %3 7333 5 -OpSelectionMerge %532 None -OpSwitch %523 %539 0 %533 1 %534 2 %535 3 %536 4 %537 5 %538 -%533 = OpLabel +OpSelectionMerge %536 None +OpSwitch %527 %543 0 %537 1 %538 2 %539 3 %540 4 %541 5 %542 +%537 = OpLabel OpLine %3 7334 37 -%540 = OpCompositeExtract %5 %531 0 -%541 = OpCompositeExtract %4 %540 0 +%544 = OpCompositeExtract %5 %535 0 +%545 = OpCompositeExtract %4 %544 0 OpLine %3 7334 20 -OpStore %496 %541 -OpBranch %532 -%534 = OpLabel +OpStore %498 %545 +OpBranch %536 +%538 = OpLabel OpLine %3 7335 37 -%542 = OpCompositeExtract %5 %531 0 -%543 = OpCompositeExtract %4 %542 1 +%546 = OpCompositeExtract %5 %535 0 +%547 = OpCompositeExtract %4 %546 1 OpLine %3 7335 20 -OpStore %496 %543 -OpBranch %532 -%535 = OpLabel +OpStore %498 %547 +OpBranch %536 +%539 = OpLabel OpLine %3 7336 37 -%544 = OpCompositeExtract %5 %531 0 -%545 = OpCompositeExtract %4 %544 2 +%548 = OpCompositeExtract %5 %535 0 +%549 = OpCompositeExtract %4 %548 2 OpLine %3 7336 20 -OpStore %496 %545 -OpBranch %532 -%536 = OpLabel +OpStore %498 %549 +OpBranch %536 +%540 = OpLabel OpLine %3 7337 37 -%546 = OpCompositeExtract %5 %531 1 -%547 = OpCompositeExtract %4 %546 0 +%550 = OpCompositeExtract %5 %535 1 +%551 = OpCompositeExtract %4 %550 0 OpLine %3 7337 20 -OpStore %496 %547 -OpBranch %532 -%537 = OpLabel +OpStore %498 %551 +OpBranch %536 +%541 = OpLabel OpLine %3 7338 37 -%548 = OpCompositeExtract %5 %531 1 -%549 = OpCompositeExtract %4 %548 1 +%552 = OpCompositeExtract %5 %535 1 +%553 = OpCompositeExtract %4 %552 1 OpLine %3 7338 20 -OpStore %496 %549 -OpBranch %532 -%538 = OpLabel +OpStore %498 %553 +OpBranch %536 +%542 = OpLabel OpLine %3 7339 37 -%550 = OpCompositeExtract %5 %531 1 -%551 = OpCompositeExtract %4 %550 2 +%554 = OpCompositeExtract %5 %535 1 +%555 = OpCompositeExtract %4 %554 2 OpLine %3 7339 20 -OpStore %496 %551 -OpBranch %532 -%539 = OpLabel -OpBranch %532 -%532 = OpLabel +OpStore %498 %555 +OpBranch %536 +%543 = OpLabel +OpBranch %536 +%536 = OpLabel OpLine %3 7343 15 -%552 = OpAccessChain %393 %494 %135 %135 -%553 = OpLoad %8 %552 -%554 = OpFunctionCall %8 %313 %522 %553 -%555 = OpIAdd %8 %522 %554 +%556 = OpAccessChain %393 %496 %135 %135 +%557 = OpLoad %8 %556 +%558 = OpFunctionCall %8 %313 %526 %557 +%559 = OpIAdd %8 %526 %558 OpLine %3 7344 15 -%556 = OpIAdd %8 %555 %126 +%560 = OpIAdd %8 %559 %126 OpLine %3 7345 15 -%557 = OpAccessChain %393 %494 %135 %135 -%558 = OpLoad %8 %557 -%559 = OpIAdd %8 %555 %558 +%561 = OpAccessChain %393 %496 %135 %135 +%562 = OpLoad %8 %561 +%563 = OpIAdd %8 %559 %562 OpLine %3 7345 15 -%560 = OpIAdd %8 %559 %126 +%564 = OpIAdd %8 %563 %126 OpLine %3 7346 15 -%561 = OpIAdd %8 %560 %126 +%565 = OpIAdd %8 %564 %126 OpLine %3 7349 5 -OpSelectionMerge %562 None -OpSwitch %523 %567 0 %563 3 %563 2 %564 4 %564 1 %565 5 %566 -%563 = OpLabel +OpSelectionMerge %566 None +OpSwitch %527 %571 0 %567 3 %567 2 %568 4 %568 1 %569 5 %570 +%567 = OpLabel OpLine %3 7350 24 -OpStore %497 %555 -OpBranch %562 -%564 = OpLabel +OpStore %499 %559 +OpBranch %566 +%568 = OpLabel OpLine %3 7351 24 -OpStore %497 %561 -OpBranch %562 -%565 = OpLabel +OpStore %499 %565 +OpBranch %566 +%569 = OpLabel OpLine %3 7352 20 -OpStore %497 %560 -OpBranch %562 -%566 = OpLabel +OpStore %499 %564 +OpBranch %566 +%570 = OpLabel OpLine %3 7353 20 -OpStore %497 %556 -OpBranch %562 -%567 = OpLabel -OpBranch %562 -%562 = OpLabel +OpStore %499 %560 +OpBranch %566 +%571 = OpLabel +OpBranch %566 +%566 = OpLabel OpLine %3 7356 13 -%568 = OpCompositeExtract %8 %482 0 +%572 = OpCompositeExtract %8 %484 0 OpLine %3 7356 5 -OpStore %497 %568 +OpStore %499 %572 OpLine %3 7365 27 -%569 = OpLoad %4 %496 -%570 = OpBitcast %8 %569 +%573 = OpLoad %4 %498 +%574 = OpBitcast %8 %573 OpLine %3 7366 12 -%571 = OpLoad %8 %497 -%572 = OpCompositeConstruct %22 %570 %571 -%573 = OpCompositeExtract %8 %572 0 -OpStore %491 %573 -%574 = OpCompositeExtract %8 %572 1 -OpStore %492 %574 +%575 = OpLoad %8 %499 +%576 = OpCompositeConstruct %22 %574 %575 +%577 = OpCompositeExtract %8 %576 0 +OpStore %493 %577 +%578 = OpCompositeExtract %8 %576 1 +OpStore %494 %578 OpReturn OpFunctionEnd -%586 = OpFunction %2 None %368 -%575 = OpLabel -%579 = OpLoad %5 %577 -%581 = OpLoad %5 %580 -%576 = OpCompositeConstruct %14 %579 %581 -%588 = OpAccessChain %587 %39 %135 -OpBranch %589 -%589 = OpLabel +%590 = OpFunction %2 None %368 +%579 = OpLabel +%583 = OpLoad %5 %581 +%585 = OpLoad %5 %584 +%580 = OpCompositeConstruct %14 %583 %585 +%592 = OpAccessChain %591 %39 %135 +OpBranch %593 +%593 = OpLabel OpLine %3 7397 25 -%591 = OpAccessChain %590 %588 %126 -%592 = OpLoad %23 %591 -%593 = OpCompositeExtract %5 %576 0 +%595 = OpAccessChain %594 %592 %126 +%596 = OpLoad %23 %595 +%597 = OpCompositeExtract %5 %580 0 OpLine %3 7397 25 -%594 = OpCompositeConstruct %7 %593 %56 -%595 = OpMatrixTimesVector %7 %592 %594 +%598 = OpCompositeConstruct %7 %597 %56 +%599 = OpMatrixTimesVector %7 %596 %598 OpLine %3 7398 18 -%596 = OpCompositeExtract %5 %576 1 +%600 = OpCompositeExtract %5 %580 1 OpLine %3 7399 12 -%597 = OpCompositeExtract %5 %576 0 -%598 = OpCompositeConstruct %26 %595 %596 %597 -%599 = OpCompositeExtract %7 %598 0 -OpStore %582 %599 -%600 = OpCompositeExtract %5 %598 1 -OpStore %583 %600 -%601 = OpCompositeExtract %5 %598 2 -OpStore %585 %601 +%601 = OpCompositeExtract %5 %580 0 +%602 = OpCompositeConstruct %26 %599 %600 %601 +%603 = OpCompositeExtract %7 %602 0 +OpStore %586 %603 +%604 = OpCompositeExtract %5 %602 1 +OpStore %587 %604 +%605 = OpCompositeExtract %5 %602 2 +OpStore %589 %605 OpReturn OpFunctionEnd -%611 = OpFunction %2 None %368 -%602 = OpLabel -%620 = OpVariable %95 Function %621 -%605 = OpLoad %7 %604 -%607 = OpLoad %5 %606 -%609 = OpLoad %5 %608 -%603 = OpCompositeConstruct %26 %605 %607 %609 -%612 = OpAccessChain %587 %39 %135 -%614 = OpAccessChain %613 %42 %135 -OpBranch %622 -%622 = OpLabel +%615 = OpFunction %2 None %368 +%606 = OpLabel +%624 = OpVariable %95 Function %625 +%609 = OpLoad %7 %608 +%611 = OpLoad %5 %610 +%613 = OpLoad %5 %612 +%607 = OpCompositeConstruct %26 %609 %611 %613 +%616 = OpAccessChain %591 %39 %135 +%618 = OpAccessChain %617 %42 %135 +OpBranch %626 +%626 = OpLabel OpLine %3 7421 28 OpLine %3 7421 17 -%623 = OpCompositeExtract %5 %603 2 -%624 = OpExtInst %5 %1 Fract %623 -%625 = OpExtInst %5 %1 SmoothStep %80 %615 %624 +%627 = OpCompositeExtract %5 %607 2 +%628 = OpExtInst %5 %1 Fract %627 +%629 = OpExtInst %5 %1 SmoothStep %80 %619 %628 OpLine %3 7421 5 -OpStore %620 %625 +OpStore %624 %629 OpLine %3 7422 17 OpLine %3 7422 13 -%626 = OpAccessChain %125 %620 %135 -%627 = OpLoad %4 %626 -%628 = OpAccessChain %125 %620 %126 -%629 = OpLoad %4 %628 -%630 = OpFMul %4 %627 %629 -%631 = OpAccessChain %125 %620 %372 -%632 = OpLoad %4 %631 -%633 = OpFMul %4 %630 %632 -%634 = OpCompositeConstruct %5 %633 %633 %633 -%635 = OpExtInst %5 %1 FMix %617 %619 %634 +%630 = OpAccessChain %125 %624 %135 +%631 = OpLoad %4 %630 +%632 = OpAccessChain %125 %624 %126 +%633 = OpLoad %4 %632 +%634 = OpFMul %4 %631 %633 +%635 = OpAccessChain %125 %624 %372 +%636 = OpLoad %4 %635 +%637 = OpFMul %4 %634 %636 +%638 = OpCompositeConstruct %5 %637 %637 %637 +%639 = OpExtInst %5 %1 FMix %621 %623 %638 OpLine %3 7422 5 -OpStore %620 %635 +OpStore %624 %639 OpLine %3 7425 25 -%637 = OpAccessChain %636 %614 %126 -%638 = OpLoad %5 %637 -%639 = OpVectorTimesScalar %5 %638 %286 +%641 = OpAccessChain %640 %618 %126 +%642 = OpLoad %5 %641 +%643 = OpVectorTimesScalar %5 %642 %286 OpLine %3 7427 21 -%640 = OpAccessChain %636 %614 %135 -%641 = OpLoad %5 %640 -%642 = OpCompositeExtract %5 %603 2 -%643 = OpFSub %5 %641 %642 -%644 = OpExtInst %5 %1 Normalize %643 +%644 = OpAccessChain %640 %618 %135 +%645 = OpLoad %5 %644 +%646 = OpCompositeExtract %5 %607 2 +%647 = OpFSub %5 %645 %646 +%648 = OpExtInst %5 %1 Normalize %647 OpLine %3 7428 20 -%646 = OpAccessChain %645 %612 %135 -%647 = OpLoad %7 %646 -%648 = OpVectorShuffle %5 %647 %647 0 1 2 -%649 = OpCompositeExtract %5 %603 2 -%650 = OpFSub %5 %648 %649 -%651 = OpExtInst %5 %1 Normalize %650 +%650 = OpAccessChain %649 %616 %135 +%651 = OpLoad %7 %650 +%652 = OpVectorShuffle %5 %651 %651 0 1 2 +%653 = OpCompositeExtract %5 %607 2 +%654 = OpFSub %5 %652 %653 +%655 = OpExtInst %5 %1 Normalize %654 OpLine %3 7429 20 -%652 = OpFAdd %5 %651 %644 -%653 = OpExtInst %5 %1 Normalize %652 +%656 = OpFAdd %5 %655 %648 +%657 = OpExtInst %5 %1 Normalize %656 OpLine %3 7431 32 -%654 = OpCompositeExtract %5 %603 1 -%655 = OpDot %4 %654 %644 +%658 = OpCompositeExtract %5 %607 1 +%659 = OpDot %4 %658 %648 OpLine %3 7431 28 -%656 = OpExtInst %4 %1 FMax %655 %74 +%660 = OpExtInst %4 %1 FMax %659 %74 OpLine %3 7432 25 -%657 = OpAccessChain %636 %614 %126 -%658 = OpLoad %5 %657 -%659 = OpVectorTimesScalar %5 %658 %656 +%661 = OpAccessChain %640 %618 %126 +%662 = OpLoad %5 %661 +%663 = OpVectorTimesScalar %5 %662 %660 OpLine %3 7434 37 -%660 = OpCompositeExtract %5 %603 1 -%661 = OpDot %4 %660 %653 +%664 = OpCompositeExtract %5 %607 1 +%665 = OpDot %4 %664 %657 OpLine %3 7434 33 -%662 = OpExtInst %4 %1 FMax %661 %74 +%666 = OpExtInst %4 %1 FMax %665 %74 OpLine %3 7434 29 -%663 = OpExtInst %4 %1 Pow %662 %345 +%667 = OpExtInst %4 %1 Pow %666 %345 OpLine %3 7435 26 -%664 = OpAccessChain %636 %614 %126 -%665 = OpLoad %5 %664 -%666 = OpVectorTimesScalar %5 %665 %663 +%668 = OpAccessChain %640 %618 %126 +%669 = OpLoad %5 %668 +%670 = OpVectorTimesScalar %5 %669 %667 OpLine %3 7437 18 -%667 = OpFAdd %5 %639 %659 -%668 = OpFAdd %5 %667 %666 -%669 = OpLoad %5 %620 -%670 = OpFMul %5 %668 %669 +%671 = OpFAdd %5 %643 %663 +%672 = OpFAdd %5 %671 %670 +%673 = OpLoad %5 %624 +%674 = OpFMul %5 %672 %673 OpLine %3 7439 12 -%671 = OpCompositeConstruct %7 %670 %56 -OpStore %610 %671 +%675 = OpCompositeConstruct %7 %674 %56 +OpStore %614 %675 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/spv/debug-symbol-terrain.spvasm b/naga/tests/out/spv/debug-symbol-terrain.spvasm index 8f637e929f9..42236277ffa 100644 --- a/naga/tests/out/spv/debug-symbol-terrain.spvasm +++ b/naga/tests/out/spv/debug-symbol-terrain.spvasm @@ -1,19 +1,19 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 672 +; Bound: 676 OpCapability Shader OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %367 "gen_terrain_compute" %364 OpEntryPoint Vertex %444 "gen_terrain_vertex" %435 %438 %440 %442 -OpEntryPoint Fragment %493 "gen_terrain_fragment" %483 %485 %488 %491 %492 -OpEntryPoint Vertex %586 "vs_main" %577 %580 %582 %583 %585 -OpEntryPoint Fragment %611 "fs_main" %604 %606 %608 %610 +OpEntryPoint Fragment %495 "gen_terrain_fragment" %485 %487 %490 %493 %494 +OpEntryPoint Vertex %590 "vs_main" %581 %584 %586 %587 %589 +OpEntryPoint Fragment %615 "fs_main" %608 %610 %612 %614 OpExecutionMode %367 LocalSize 64 1 1 -OpExecutionMode %493 OriginUpperLeft -OpExecutionMode %611 OriginUpperLeft +OpExecutionMode %495 OriginUpperLeft +OpExecutionMode %615 OriginUpperLeft %3 = OpString "debug-symbol-terrain.wgsl" OpSource Unknown 0 %3 "// Taken from https://github.com/sotrh/learn-wgpu/blob/11820796f5e1dbce42fb1119f04ddeb4b167d2a0/code/intermediate/tutorial13-terrain/src/terrain.wgsl // ============================ @@ -396,25 +396,25 @@ OpName %438 "index" OpName %440 "position" OpName %442 "uv" OpName %444 "gen_terrain_vertex" -OpName %483 "index" -OpName %485 "position" -OpName %488 "uv" -OpName %491 "vert_component" -OpName %492 "index" -OpName %493 "gen_terrain_fragment" -OpName %496 "vert_component" -OpName %497 "index" -OpName %577 "position" -OpName %580 "normal" -OpName %582 "clip_position" -OpName %583 "normal" -OpName %585 "world_pos" -OpName %586 "vs_main" -OpName %604 "clip_position" -OpName %606 "normal" -OpName %608 "world_pos" -OpName %611 "fs_main" -OpName %620 "color" +OpName %485 "index" +OpName %487 "position" +OpName %490 "uv" +OpName %493 "vert_component" +OpName %494 "index" +OpName %495 "gen_terrain_fragment" +OpName %498 "vert_component" +OpName %499 "index" +OpName %581 "position" +OpName %584 "normal" +OpName %586 "clip_position" +OpName %587 "normal" +OpName %589 "world_pos" +OpName %590 "vs_main" +OpName %608 "clip_position" +OpName %610 "normal" +OpName %612 "world_pos" +OpName %615 "fs_main" +OpName %624 "color" OpMemberDecorate %13 0 Offset 0 OpMemberDecorate %13 1 Offset 8 OpMemberDecorate %13 2 Offset 16 @@ -479,21 +479,21 @@ OpDecorate %438 Location 0 OpDecorate %438 Flat OpDecorate %440 BuiltIn Position OpDecorate %442 Location 1 -OpDecorate %483 Location 0 -OpDecorate %483 Flat -OpDecorate %485 BuiltIn FragCoord -OpDecorate %488 Location 1 -OpDecorate %491 Location 0 -OpDecorate %492 Location 1 -OpDecorate %577 Location 0 -OpDecorate %580 Location 1 -OpDecorate %582 BuiltIn Position -OpDecorate %583 Location 0 -OpDecorate %585 Location 1 -OpDecorate %604 BuiltIn FragCoord -OpDecorate %606 Location 0 -OpDecorate %608 Location 1 +OpDecorate %485 Location 0 +OpDecorate %485 Flat +OpDecorate %487 BuiltIn FragCoord +OpDecorate %490 Location 1 +OpDecorate %493 Location 0 +OpDecorate %494 Location 1 +OpDecorate %581 Location 0 +OpDecorate %584 Location 1 +OpDecorate %586 BuiltIn Position +OpDecorate %587 Location 0 +OpDecorate %589 Location 1 +OpDecorate %608 BuiltIn FragCoord OpDecorate %610 Location 0 +OpDecorate %612 Location 1 +OpDecorate %614 Location 0 %2 = OpTypeVoid %4 = OpTypeFloat 32 %5 = OpTypeVector %4 3 @@ -635,36 +635,37 @@ OpDecorate %610 Location 0 %445 = OpTypePointer Uniform %20 %447 = OpConstant %4 -1.0 %448 = OpConstantComposite %6 %447 %447 -%483 = OpVariable %436 Input -%486 = OpTypePointer Input %7 -%485 = OpVariable %486 Input -%489 = OpTypePointer Input %6 -%488 = OpVariable %489 Input -%491 = OpVariable %439 Output -%492 = OpVariable %439 Output -%495 = OpConstant %4 6.0 -%578 = OpTypePointer Input %5 -%577 = OpVariable %578 Input -%580 = OpVariable %578 Input -%582 = OpVariable %441 Output -%584 = OpTypePointer Output %5 -%583 = OpVariable %584 Output -%585 = OpVariable %584 Output -%587 = OpTypePointer Uniform %24 -%590 = OpTypePointer Uniform %23 -%604 = OpVariable %486 Input -%606 = OpVariable %578 Input -%608 = OpVariable %578 Input -%610 = OpVariable %441 Output -%613 = OpTypePointer Uniform %25 -%615 = OpConstantComposite %5 %286 %286 %286 -%616 = OpConstant %4 0.7 -%617 = OpConstantComposite %5 %78 %286 %616 -%618 = OpConstant %4 0.2 -%619 = OpConstantComposite %5 %618 %618 %618 -%621 = OpConstantNull %5 -%636 = OpTypePointer Uniform %5 -%645 = OpTypePointer Uniform %7 +%473 = OpConstant %4 4294967000.0 +%485 = OpVariable %436 Input +%488 = OpTypePointer Input %7 +%487 = OpVariable %488 Input +%491 = OpTypePointer Input %6 +%490 = OpVariable %491 Input +%493 = OpVariable %439 Output +%494 = OpVariable %439 Output +%497 = OpConstant %4 6.0 +%582 = OpTypePointer Input %5 +%581 = OpVariable %582 Input +%584 = OpVariable %582 Input +%586 = OpVariable %441 Output +%588 = OpTypePointer Output %5 +%587 = OpVariable %588 Output +%589 = OpVariable %588 Output +%591 = OpTypePointer Uniform %24 +%594 = OpTypePointer Uniform %23 +%608 = OpVariable %488 Input +%610 = OpVariable %582 Input +%612 = OpVariable %582 Input +%614 = OpVariable %441 Output +%617 = OpTypePointer Uniform %25 +%619 = OpConstantComposite %5 %286 %286 %286 +%620 = OpConstant %4 0.7 +%621 = OpConstantComposite %5 %78 %286 %620 +%622 = OpConstant %4 0.2 +%623 = OpConstantComposite %5 %622 %622 %622 +%625 = OpConstantNull %5 +%640 = OpTypePointer Uniform %5 +%649 = OpTypePointer Uniform %7 %53 = OpFunction %5 None %54 %52 = OpFunctionParameter %5 %51 = OpLabel @@ -1218,286 +1219,289 @@ OpLine %3 168 17 %470 = OpConvertUToF %4 %469 %471 = OpFMul %4 %467 %470 %472 = OpFAdd %4 %466 %471 -%473 = OpConvertFToU %8 %472 +%474 = OpExtInst %4 %1 FClamp %472 %74 %473 +%475 = OpConvertFToU %8 %474 OpLine %3 168 17 -%474 = OpAccessChain %393 %446 %374 -%475 = OpLoad %8 %474 -%476 = OpIAdd %8 %473 %475 +%476 = OpAccessChain %393 %446 %374 +%477 = OpLoad %8 %476 +%478 = OpIAdd %8 %475 %477 OpLine %3 170 12 -%477 = OpCompositeConstruct %21 %476 %461 %458 -%478 = OpCompositeExtract %8 %477 0 -OpStore %438 %478 -%479 = OpCompositeExtract %7 %477 1 -OpStore %440 %479 -%480 = OpCompositeExtract %6 %477 2 -OpStore %442 %480 +%479 = OpCompositeConstruct %21 %478 %461 %458 +%480 = OpCompositeExtract %8 %479 0 +OpStore %438 %480 +%481 = OpCompositeExtract %7 %479 1 +OpStore %440 %481 +%482 = OpCompositeExtract %6 %479 2 +OpStore %442 %482 OpReturn OpFunctionEnd -%493 = OpFunction %2 None %368 -%481 = OpLabel -%496 = OpVariable %125 Function %74 -%497 = OpVariable %217 Function %135 -%484 = OpLoad %8 %483 -%487 = OpLoad %7 %485 -%490 = OpLoad %6 %488 -%482 = OpCompositeConstruct %21 %484 %487 %490 -%494 = OpAccessChain %445 %36 %135 -OpBranch %498 -%498 = OpLabel +%495 = OpFunction %2 None %368 +%483 = OpLabel +%498 = OpVariable %125 Function %74 +%499 = OpVariable %217 Function %135 +%486 = OpLoad %8 %485 +%489 = OpLoad %7 %487 +%492 = OpLoad %6 %490 +%484 = OpCompositeConstruct %21 %486 %489 %492 +%496 = OpAccessChain %445 %36 %135 +OpBranch %500 +%500 = OpLabel OpLine %3 181 17 -%499 = OpCompositeExtract %6 %482 2 -%500 = OpCompositeExtract %4 %499 0 +%501 = OpCompositeExtract %6 %484 2 +%502 = OpCompositeExtract %4 %501 0 OpLine %3 181 17 -%501 = OpAccessChain %393 %494 %373 -%502 = OpLoad %8 %501 -%503 = OpConvertUToF %4 %502 -%504 = OpFMul %4 %500 %503 -%505 = OpCompositeExtract %6 %482 2 -%506 = OpCompositeExtract %4 %505 1 +%503 = OpAccessChain %393 %496 %373 +%504 = OpLoad %8 %503 +%505 = OpConvertUToF %4 %504 +%506 = OpFMul %4 %502 %505 +%507 = OpCompositeExtract %6 %484 2 +%508 = OpCompositeExtract %4 %507 1 OpLine %3 181 70 -%507 = OpAccessChain %393 %494 %373 -%508 = OpLoad %8 %507 -OpLine %3 181 13 -%509 = OpAccessChain %393 %494 %373 +%509 = OpAccessChain %393 %496 %373 %510 = OpLoad %8 %509 -%511 = OpIMul %8 %508 %510 -%512 = OpConvertUToF %4 %511 -%513 = OpFMul %4 %506 %512 -%514 = OpFAdd %4 %504 %513 -%515 = OpConvertFToU %8 %514 OpLine %3 181 13 -%516 = OpAccessChain %393 %494 %374 -%517 = OpLoad %8 %516 -%518 = OpIAdd %8 %515 %517 +%511 = OpAccessChain %393 %496 %373 +%512 = OpLoad %8 %511 +%513 = OpIMul %8 %510 %512 +%514 = OpConvertUToF %4 %513 +%515 = OpFMul %4 %508 %514 +%516 = OpFAdd %4 %506 %515 +%517 = OpExtInst %4 %1 FClamp %516 %74 %473 +%518 = OpConvertFToU %8 %517 +OpLine %3 181 13 +%519 = OpAccessChain %393 %496 %374 +%520 = OpLoad %8 %519 +%521 = OpIAdd %8 %518 %520 OpLine %3 182 32 -%519 = OpConvertUToF %4 %518 +%522 = OpConvertUToF %4 %521 OpLine %3 182 22 -%520 = OpFDiv %4 %519 %495 -%521 = OpExtInst %4 %1 Floor %520 -%522 = OpConvertFToU %8 %521 +%523 = OpFDiv %4 %522 %497 +%524 = OpExtInst %4 %1 Floor %523 +%525 = OpExtInst %4 %1 FClamp %524 %74 %473 +%526 = OpConvertFToU %8 %525 OpLine %3 183 22 -%523 = OpFunctionCall %8 %427 %518 %371 +%527 = OpFunctionCall %8 %427 %521 %371 OpLine %3 185 36 -%524 = OpAccessChain %377 %494 %135 -%525 = OpLoad %10 %524 +%528 = OpAccessChain %377 %496 %135 +%529 = OpLoad %10 %528 OpLine %3 185 57 -%526 = OpAccessChain %380 %494 %126 -%527 = OpLoad %11 %526 +%530 = OpAccessChain %380 %496 %126 +%531 = OpLoad %11 %530 OpLine %3 185 13 -%528 = OpFunctionCall %6 %325 %522 %525 %527 +%532 = OpFunctionCall %6 %325 %526 %529 %531 OpLine %3 186 31 -%529 = OpAccessChain %386 %494 %372 -%530 = OpLoad %6 %529 +%533 = OpAccessChain %386 %496 %372 +%534 = OpLoad %6 %533 OpLine %3 186 13 -%531 = OpFunctionCall %14 %284 %528 %530 +%535 = OpFunctionCall %14 %284 %532 %534 OpLine %3 190 5 -OpSelectionMerge %532 None -OpSwitch %523 %539 0 %533 1 %534 2 %535 3 %536 4 %537 5 %538 -%533 = OpLabel +OpSelectionMerge %536 None +OpSwitch %527 %543 0 %537 1 %538 2 %539 3 %540 4 %541 5 %542 +%537 = OpLabel OpLine %3 191 37 -%540 = OpCompositeExtract %5 %531 0 -%541 = OpCompositeExtract %4 %540 0 +%544 = OpCompositeExtract %5 %535 0 +%545 = OpCompositeExtract %4 %544 0 OpLine %3 191 20 -OpStore %496 %541 -OpBranch %532 -%534 = OpLabel +OpStore %498 %545 +OpBranch %536 +%538 = OpLabel OpLine %3 192 37 -%542 = OpCompositeExtract %5 %531 0 -%543 = OpCompositeExtract %4 %542 1 +%546 = OpCompositeExtract %5 %535 0 +%547 = OpCompositeExtract %4 %546 1 OpLine %3 192 20 -OpStore %496 %543 -OpBranch %532 -%535 = OpLabel +OpStore %498 %547 +OpBranch %536 +%539 = OpLabel OpLine %3 193 37 -%544 = OpCompositeExtract %5 %531 0 -%545 = OpCompositeExtract %4 %544 2 +%548 = OpCompositeExtract %5 %535 0 +%549 = OpCompositeExtract %4 %548 2 OpLine %3 193 20 -OpStore %496 %545 -OpBranch %532 -%536 = OpLabel +OpStore %498 %549 +OpBranch %536 +%540 = OpLabel OpLine %3 194 37 -%546 = OpCompositeExtract %5 %531 1 -%547 = OpCompositeExtract %4 %546 0 +%550 = OpCompositeExtract %5 %535 1 +%551 = OpCompositeExtract %4 %550 0 OpLine %3 194 20 -OpStore %496 %547 -OpBranch %532 -%537 = OpLabel +OpStore %498 %551 +OpBranch %536 +%541 = OpLabel OpLine %3 195 37 -%548 = OpCompositeExtract %5 %531 1 -%549 = OpCompositeExtract %4 %548 1 +%552 = OpCompositeExtract %5 %535 1 +%553 = OpCompositeExtract %4 %552 1 OpLine %3 195 20 -OpStore %496 %549 -OpBranch %532 -%538 = OpLabel +OpStore %498 %553 +OpBranch %536 +%542 = OpLabel OpLine %3 196 37 -%550 = OpCompositeExtract %5 %531 1 -%551 = OpCompositeExtract %4 %550 2 +%554 = OpCompositeExtract %5 %535 1 +%555 = OpCompositeExtract %4 %554 2 OpLine %3 196 20 -OpStore %496 %551 -OpBranch %532 -%539 = OpLabel -OpBranch %532 -%532 = OpLabel +OpStore %498 %555 +OpBranch %536 +%543 = OpLabel +OpBranch %536 +%536 = OpLabel OpLine %3 200 15 -%552 = OpAccessChain %393 %494 %135 %135 -%553 = OpLoad %8 %552 -%554 = OpFunctionCall %8 %313 %522 %553 -%555 = OpIAdd %8 %522 %554 +%556 = OpAccessChain %393 %496 %135 %135 +%557 = OpLoad %8 %556 +%558 = OpFunctionCall %8 %313 %526 %557 +%559 = OpIAdd %8 %526 %558 OpLine %3 201 15 -%556 = OpIAdd %8 %555 %126 +%560 = OpIAdd %8 %559 %126 OpLine %3 202 15 -%557 = OpAccessChain %393 %494 %135 %135 -%558 = OpLoad %8 %557 -%559 = OpIAdd %8 %555 %558 +%561 = OpAccessChain %393 %496 %135 %135 +%562 = OpLoad %8 %561 +%563 = OpIAdd %8 %559 %562 OpLine %3 202 15 -%560 = OpIAdd %8 %559 %126 +%564 = OpIAdd %8 %563 %126 OpLine %3 203 15 -%561 = OpIAdd %8 %560 %126 +%565 = OpIAdd %8 %564 %126 OpLine %3 206 5 -OpSelectionMerge %562 None -OpSwitch %523 %567 0 %563 3 %563 2 %564 4 %564 1 %565 5 %566 -%563 = OpLabel +OpSelectionMerge %566 None +OpSwitch %527 %571 0 %567 3 %567 2 %568 4 %568 1 %569 5 %570 +%567 = OpLabel OpLine %3 207 24 -OpStore %497 %555 -OpBranch %562 -%564 = OpLabel +OpStore %499 %559 +OpBranch %566 +%568 = OpLabel OpLine %3 208 24 -OpStore %497 %561 -OpBranch %562 -%565 = OpLabel +OpStore %499 %565 +OpBranch %566 +%569 = OpLabel OpLine %3 209 20 -OpStore %497 %560 -OpBranch %562 -%566 = OpLabel +OpStore %499 %564 +OpBranch %566 +%570 = OpLabel OpLine %3 210 20 -OpStore %497 %556 -OpBranch %562 -%567 = OpLabel -OpBranch %562 -%562 = OpLabel +OpStore %499 %560 +OpBranch %566 +%571 = OpLabel +OpBranch %566 +%566 = OpLabel OpLine %3 213 13 -%568 = OpCompositeExtract %8 %482 0 +%572 = OpCompositeExtract %8 %484 0 OpLine %3 213 5 -OpStore %497 %568 +OpStore %499 %572 OpLine %3 222 27 -%569 = OpLoad %4 %496 -%570 = OpBitcast %8 %569 +%573 = OpLoad %4 %498 +%574 = OpBitcast %8 %573 OpLine %3 223 12 -%571 = OpLoad %8 %497 -%572 = OpCompositeConstruct %22 %570 %571 -%573 = OpCompositeExtract %8 %572 0 -OpStore %491 %573 -%574 = OpCompositeExtract %8 %572 1 -OpStore %492 %574 +%575 = OpLoad %8 %499 +%576 = OpCompositeConstruct %22 %574 %575 +%577 = OpCompositeExtract %8 %576 0 +OpStore %493 %577 +%578 = OpCompositeExtract %8 %576 1 +OpStore %494 %578 OpReturn OpFunctionEnd -%586 = OpFunction %2 None %368 -%575 = OpLabel -%579 = OpLoad %5 %577 -%581 = OpLoad %5 %580 -%576 = OpCompositeConstruct %14 %579 %581 -%588 = OpAccessChain %587 %39 %135 -OpBranch %589 -%589 = OpLabel +%590 = OpFunction %2 None %368 +%579 = OpLabel +%583 = OpLoad %5 %581 +%585 = OpLoad %5 %584 +%580 = OpCompositeConstruct %14 %583 %585 +%592 = OpAccessChain %591 %39 %135 +OpBranch %593 +%593 = OpLabel OpLine %3 254 25 -%591 = OpAccessChain %590 %588 %126 -%592 = OpLoad %23 %591 -%593 = OpCompositeExtract %5 %576 0 +%595 = OpAccessChain %594 %592 %126 +%596 = OpLoad %23 %595 +%597 = OpCompositeExtract %5 %580 0 OpLine %3 254 25 -%594 = OpCompositeConstruct %7 %593 %56 -%595 = OpMatrixTimesVector %7 %592 %594 +%598 = OpCompositeConstruct %7 %597 %56 +%599 = OpMatrixTimesVector %7 %596 %598 OpLine %3 255 18 -%596 = OpCompositeExtract %5 %576 1 +%600 = OpCompositeExtract %5 %580 1 OpLine %3 256 12 -%597 = OpCompositeExtract %5 %576 0 -%598 = OpCompositeConstruct %26 %595 %596 %597 -%599 = OpCompositeExtract %7 %598 0 -OpStore %582 %599 -%600 = OpCompositeExtract %5 %598 1 -OpStore %583 %600 -%601 = OpCompositeExtract %5 %598 2 -OpStore %585 %601 +%601 = OpCompositeExtract %5 %580 0 +%602 = OpCompositeConstruct %26 %599 %600 %601 +%603 = OpCompositeExtract %7 %602 0 +OpStore %586 %603 +%604 = OpCompositeExtract %5 %602 1 +OpStore %587 %604 +%605 = OpCompositeExtract %5 %602 2 +OpStore %589 %605 OpReturn OpFunctionEnd -%611 = OpFunction %2 None %368 -%602 = OpLabel -%620 = OpVariable %95 Function %621 -%605 = OpLoad %7 %604 -%607 = OpLoad %5 %606 -%609 = OpLoad %5 %608 -%603 = OpCompositeConstruct %26 %605 %607 %609 -%612 = OpAccessChain %587 %39 %135 -%614 = OpAccessChain %613 %42 %135 -OpBranch %622 -%622 = OpLabel +%615 = OpFunction %2 None %368 +%606 = OpLabel +%624 = OpVariable %95 Function %625 +%609 = OpLoad %7 %608 +%611 = OpLoad %5 %610 +%613 = OpLoad %5 %612 +%607 = OpCompositeConstruct %26 %609 %611 %613 +%616 = OpAccessChain %591 %39 %135 +%618 = OpAccessChain %617 %42 %135 +OpBranch %626 +%626 = OpLabel OpLine %3 278 28 OpLine %3 278 17 -%623 = OpCompositeExtract %5 %603 2 -%624 = OpExtInst %5 %1 Fract %623 -%625 = OpExtInst %5 %1 SmoothStep %80 %615 %624 +%627 = OpCompositeExtract %5 %607 2 +%628 = OpExtInst %5 %1 Fract %627 +%629 = OpExtInst %5 %1 SmoothStep %80 %619 %628 OpLine %3 278 5 -OpStore %620 %625 +OpStore %624 %629 OpLine %3 279 17 OpLine %3 279 13 -%626 = OpAccessChain %125 %620 %135 -%627 = OpLoad %4 %626 -%628 = OpAccessChain %125 %620 %126 -%629 = OpLoad %4 %628 -%630 = OpFMul %4 %627 %629 -%631 = OpAccessChain %125 %620 %372 -%632 = OpLoad %4 %631 -%633 = OpFMul %4 %630 %632 -%634 = OpCompositeConstruct %5 %633 %633 %633 -%635 = OpExtInst %5 %1 FMix %617 %619 %634 +%630 = OpAccessChain %125 %624 %135 +%631 = OpLoad %4 %630 +%632 = OpAccessChain %125 %624 %126 +%633 = OpLoad %4 %632 +%634 = OpFMul %4 %631 %633 +%635 = OpAccessChain %125 %624 %372 +%636 = OpLoad %4 %635 +%637 = OpFMul %4 %634 %636 +%638 = OpCompositeConstruct %5 %637 %637 %637 +%639 = OpExtInst %5 %1 FMix %621 %623 %638 OpLine %3 279 5 -OpStore %620 %635 +OpStore %624 %639 OpLine %3 282 25 -%637 = OpAccessChain %636 %614 %126 -%638 = OpLoad %5 %637 -%639 = OpVectorTimesScalar %5 %638 %286 +%641 = OpAccessChain %640 %618 %126 +%642 = OpLoad %5 %641 +%643 = OpVectorTimesScalar %5 %642 %286 OpLine %3 284 21 -%640 = OpAccessChain %636 %614 %135 -%641 = OpLoad %5 %640 -%642 = OpCompositeExtract %5 %603 2 -%643 = OpFSub %5 %641 %642 -%644 = OpExtInst %5 %1 Normalize %643 +%644 = OpAccessChain %640 %618 %135 +%645 = OpLoad %5 %644 +%646 = OpCompositeExtract %5 %607 2 +%647 = OpFSub %5 %645 %646 +%648 = OpExtInst %5 %1 Normalize %647 OpLine %3 285 20 -%646 = OpAccessChain %645 %612 %135 -%647 = OpLoad %7 %646 -%648 = OpVectorShuffle %5 %647 %647 0 1 2 -%649 = OpCompositeExtract %5 %603 2 -%650 = OpFSub %5 %648 %649 -%651 = OpExtInst %5 %1 Normalize %650 +%650 = OpAccessChain %649 %616 %135 +%651 = OpLoad %7 %650 +%652 = OpVectorShuffle %5 %651 %651 0 1 2 +%653 = OpCompositeExtract %5 %607 2 +%654 = OpFSub %5 %652 %653 +%655 = OpExtInst %5 %1 Normalize %654 OpLine %3 286 20 -%652 = OpFAdd %5 %651 %644 -%653 = OpExtInst %5 %1 Normalize %652 +%656 = OpFAdd %5 %655 %648 +%657 = OpExtInst %5 %1 Normalize %656 OpLine %3 288 32 -%654 = OpCompositeExtract %5 %603 1 -%655 = OpDot %4 %654 %644 +%658 = OpCompositeExtract %5 %607 1 +%659 = OpDot %4 %658 %648 OpLine %3 288 28 -%656 = OpExtInst %4 %1 FMax %655 %74 +%660 = OpExtInst %4 %1 FMax %659 %74 OpLine %3 289 25 -%657 = OpAccessChain %636 %614 %126 -%658 = OpLoad %5 %657 -%659 = OpVectorTimesScalar %5 %658 %656 +%661 = OpAccessChain %640 %618 %126 +%662 = OpLoad %5 %661 +%663 = OpVectorTimesScalar %5 %662 %660 OpLine %3 291 37 -%660 = OpCompositeExtract %5 %603 1 -%661 = OpDot %4 %660 %653 +%664 = OpCompositeExtract %5 %607 1 +%665 = OpDot %4 %664 %657 OpLine %3 291 33 -%662 = OpExtInst %4 %1 FMax %661 %74 +%666 = OpExtInst %4 %1 FMax %665 %74 OpLine %3 291 29 -%663 = OpExtInst %4 %1 Pow %662 %345 +%667 = OpExtInst %4 %1 Pow %666 %345 OpLine %3 292 26 -%664 = OpAccessChain %636 %614 %126 -%665 = OpLoad %5 %664 -%666 = OpVectorTimesScalar %5 %665 %663 +%668 = OpAccessChain %640 %618 %126 +%669 = OpLoad %5 %668 +%670 = OpVectorTimesScalar %5 %669 %667 OpLine %3 294 18 -%667 = OpFAdd %5 %639 %659 -%668 = OpFAdd %5 %667 %666 -%669 = OpLoad %5 %620 -%670 = OpFMul %5 %668 %669 +%671 = OpFAdd %5 %643 %663 +%672 = OpFAdd %5 %671 %670 +%673 = OpLoad %5 %624 +%674 = OpFMul %5 %672 %673 OpLine %3 296 12 -%671 = OpCompositeConstruct %7 %670 %56 -OpStore %610 %671 +%675 = OpCompositeConstruct %7 %674 %56 +OpStore %614 %675 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/spv/image.spvasm b/naga/tests/out/spv/image.spvasm index f1303d23ab2..b0f33ef15d1 100644 --- a/naga/tests/out/spv/image.spvasm +++ b/naga/tests/out/spv/image.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 545 +; Bound: 547 OpCapability Shader OpCapability Image1D OpCapability Sampled1D @@ -11,18 +11,18 @@ OpCapability ImageQuery OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %96 "main" %93 OpEntryPoint GLCompute %188 "depth_load" %186 -OpEntryPoint Vertex %208 "queries" %206 -OpEntryPoint Vertex %260 "levels_queries" %259 -OpEntryPoint Fragment %291 "texture_sample" %290 -OpEntryPoint Fragment %438 "texture_sample_comparison" %436 -OpEntryPoint Fragment %494 "gather" %493 -OpEntryPoint Fragment %528 "depth_no_comparison" %527 +OpEntryPoint Vertex %211 "queries" %209 +OpEntryPoint Vertex %263 "levels_queries" %262 +OpEntryPoint Fragment %294 "texture_sample" %293 +OpEntryPoint Fragment %441 "texture_sample_comparison" %439 +OpEntryPoint Fragment %496 "gather" %495 +OpEntryPoint Fragment %530 "depth_no_comparison" %529 OpExecutionMode %96 LocalSize 16 1 1 OpExecutionMode %188 LocalSize 16 1 1 -OpExecutionMode %291 OriginUpperLeft -OpExecutionMode %438 OriginUpperLeft -OpExecutionMode %494 OriginUpperLeft -OpExecutionMode %528 OriginUpperLeft +OpExecutionMode %294 OriginUpperLeft +OpExecutionMode %441 OriginUpperLeft +OpExecutionMode %496 OriginUpperLeft +OpExecutionMode %530 OriginUpperLeft %3 = OpString "image.wgsl" OpSource Unknown 0 %3 "@group(0) @binding(0) var image_mipmapped_src: texture_2d; @@ -250,14 +250,14 @@ OpName %93 "local_id" OpName %96 "main" OpName %186 "local_id" OpName %188 "depth_load" -OpName %208 "queries" -OpName %260 "levels_queries" -OpName %291 "texture_sample" -OpName %306 "a" -OpName %438 "texture_sample_comparison" -OpName %443 "a" -OpName %494 "gather" -OpName %528 "depth_no_comparison" +OpName %211 "queries" +OpName %263 "levels_queries" +OpName %294 "texture_sample" +OpName %309 "a" +OpName %441 "texture_sample_comparison" +OpName %446 "a" +OpName %496 "gather" +OpName %530 "depth_no_comparison" OpDecorate %29 DescriptorSet 0 OpDecorate %29 Binding 0 OpDecorate %31 DescriptorSet 0 @@ -307,12 +307,12 @@ OpDecorate %68 DescriptorSet 1 OpDecorate %68 Binding 4 OpDecorate %93 BuiltIn LocalInvocationId OpDecorate %186 BuiltIn LocalInvocationId -OpDecorate %206 BuiltIn Position -OpDecorate %259 BuiltIn Position -OpDecorate %290 Location 0 -OpDecorate %436 Location 0 -OpDecorate %493 Location 0 -OpDecorate %527 Location 0 +OpDecorate %209 BuiltIn Position +OpDecorate %262 BuiltIn Position +OpDecorate %293 Location 0 +OpDecorate %439 Location 0 +OpDecorate %495 Location 0 +OpDecorate %529 Location 0 %2 = OpTypeVoid %5 = OpTypeInt 32 0 %4 = OpTypeImage %5 2D 0 0 0 1 Unknown @@ -401,41 +401,42 @@ OpDecorate %527 Location 0 %116 = OpTypeVector %5 4 %129 = OpTypeVector %15 3 %186 = OpVariable %94 Input -%207 = OpTypePointer Output %24 -%206 = OpVariable %207 Output -%217 = OpConstant %5 0 -%259 = OpVariable %207 Output -%290 = OpVariable %207 Output -%297 = OpConstant %8 0.5 -%298 = OpTypeVector %8 2 -%299 = OpConstantComposite %298 %297 %297 -%300 = OpTypeVector %8 3 -%301 = OpConstantComposite %300 %297 %297 %297 -%302 = OpConstant %15 3 -%303 = OpConstantComposite %14 %302 %88 -%304 = OpConstant %8 2.3 -%305 = OpConstant %8 2.0 -%307 = OpTypePointer Function %24 -%308 = OpConstantNull %24 -%310 = OpTypeSampledImage %16 -%315 = OpTypeSampledImage %17 -%336 = OpTypeSampledImage %19 -%397 = OpTypeSampledImage %21 -%437 = OpTypePointer Output %8 -%436 = OpVariable %437 Output -%444 = OpTypePointer Function %8 -%445 = OpConstantNull %8 -%447 = OpTypeSampledImage %26 -%452 = OpTypeSampledImage %27 -%465 = OpTypeSampledImage %28 -%472 = OpConstant %8 0.0 -%493 = OpVariable %207 Output -%504 = OpConstant %5 1 -%507 = OpConstant %5 3 -%512 = OpTypeSampledImage %4 -%515 = OpTypeVector %15 4 -%516 = OpTypeSampledImage %18 -%527 = OpVariable %207 Output +%203 = OpConstant %8 0.0 +%204 = OpConstant %8 4294967000.0 +%210 = OpTypePointer Output %24 +%209 = OpVariable %210 Output +%220 = OpConstant %5 0 +%262 = OpVariable %210 Output +%293 = OpVariable %210 Output +%300 = OpConstant %8 0.5 +%301 = OpTypeVector %8 2 +%302 = OpConstantComposite %301 %300 %300 +%303 = OpTypeVector %8 3 +%304 = OpConstantComposite %303 %300 %300 %300 +%305 = OpConstant %15 3 +%306 = OpConstantComposite %14 %305 %88 +%307 = OpConstant %8 2.3 +%308 = OpConstant %8 2.0 +%310 = OpTypePointer Function %24 +%311 = OpConstantNull %24 +%313 = OpTypeSampledImage %16 +%318 = OpTypeSampledImage %17 +%339 = OpTypeSampledImage %19 +%400 = OpTypeSampledImage %21 +%440 = OpTypePointer Output %8 +%439 = OpVariable %440 Output +%447 = OpTypePointer Function %8 +%448 = OpConstantNull %8 +%450 = OpTypeSampledImage %26 +%455 = OpTypeSampledImage %27 +%468 = OpTypeSampledImage %28 +%495 = OpVariable %210 Output +%506 = OpConstant %5 1 +%509 = OpConstant %5 3 +%514 = OpTypeSampledImage %4 +%517 = OpTypeVector %15 4 +%518 = OpTypeSampledImage %18 +%529 = OpVariable %210 Output %70 = OpFunction %14 None %71 %72 = OpFunctionParameter %14 %73 = OpFunctionParameter %14 @@ -585,475 +586,476 @@ OpLine %3 48 20 %201 = OpCompositeExtract %8 %200 0 OpLine %3 49 29 %202 = OpCompositeExtract %15 %197 0 -%203 = OpConvertFToU %5 %201 -%204 = OpCompositeConstruct %116 %203 %203 %203 %203 +%205 = OpExtInst %8 %1 FClamp %201 %203 %204 +%206 = OpConvertFToU %5 %205 +%207 = OpCompositeConstruct %116 %206 %206 %206 %206 OpLine %3 49 5 -OpImageWrite %191 %202 %204 +OpImageWrite %191 %202 %207 OpReturn OpFunctionEnd -%208 = OpFunction %2 None %97 -%205 = OpLabel -%209 = OpLoad %16 %44 -%210 = OpLoad %17 %46 -%211 = OpLoad %19 %51 -%212 = OpLoad %20 %53 -%213 = OpLoad %21 %55 -%214 = OpLoad %22 %57 -%215 = OpLoad %23 %59 -OpBranch %216 -%216 = OpLabel +%211 = OpFunction %2 None %97 +%208 = OpLabel +%212 = OpLoad %16 %44 +%213 = OpLoad %17 %46 +%214 = OpLoad %19 %51 +%215 = OpLoad %20 %53 +%216 = OpLoad %21 %55 +%217 = OpLoad %22 %57 +%218 = OpLoad %23 %59 +OpBranch %219 +%219 = OpLabel OpLine %3 74 18 -%218 = OpImageQuerySizeLod %5 %209 %217 +%221 = OpImageQuerySizeLod %5 %212 %220 OpLine %3 75 22 -%219 = OpBitcast %15 %218 -%220 = OpImageQuerySizeLod %5 %209 %219 +%222 = OpBitcast %15 %221 +%223 = OpImageQuerySizeLod %5 %212 %222 OpLine %3 76 18 -%221 = OpImageQuerySizeLod %108 %210 %217 +%224 = OpImageQuerySizeLod %108 %213 %220 OpLine %3 77 22 -%222 = OpImageQuerySizeLod %108 %210 %88 +%225 = OpImageQuerySizeLod %108 %213 %88 OpLine %3 78 24 -%223 = OpImageQuerySizeLod %13 %211 %217 -%224 = OpVectorShuffle %108 %223 %223 0 1 +%226 = OpImageQuerySizeLod %13 %214 %220 +%227 = OpVectorShuffle %108 %226 %226 0 1 OpLine %3 79 28 -%225 = OpImageQuerySizeLod %13 %211 %88 -%226 = OpVectorShuffle %108 %225 %225 0 1 +%228 = OpImageQuerySizeLod %13 %214 %88 +%229 = OpVectorShuffle %108 %228 %228 0 1 OpLine %3 80 20 -%227 = OpImageQuerySizeLod %108 %212 %217 +%230 = OpImageQuerySizeLod %108 %215 %220 OpLine %3 81 24 -%228 = OpImageQuerySizeLod %108 %212 %88 +%231 = OpImageQuerySizeLod %108 %215 %88 OpLine %3 82 26 -%229 = OpImageQuerySizeLod %13 %213 %217 -%230 = OpVectorShuffle %108 %229 %229 0 0 +%232 = OpImageQuerySizeLod %13 %216 %220 +%233 = OpVectorShuffle %108 %232 %232 0 0 OpLine %3 83 30 -%231 = OpImageQuerySizeLod %13 %213 %88 -%232 = OpVectorShuffle %108 %231 %231 0 0 +%234 = OpImageQuerySizeLod %13 %216 %88 +%235 = OpVectorShuffle %108 %234 %234 0 0 OpLine %3 84 18 -%233 = OpImageQuerySizeLod %13 %214 %217 +%236 = OpImageQuerySizeLod %13 %217 %220 OpLine %3 85 22 -%234 = OpImageQuerySizeLod %13 %214 %88 +%237 = OpImageQuerySizeLod %13 %217 %88 OpLine %3 86 21 -%235 = OpImageQuerySize %108 %215 +%238 = OpImageQuerySize %108 %218 OpLine %3 88 15 -%236 = OpCompositeExtract %5 %221 1 -%237 = OpIAdd %5 %218 %236 -%238 = OpCompositeExtract %5 %222 1 -%239 = OpIAdd %5 %237 %238 -%240 = OpCompositeExtract %5 %224 1 -%241 = OpIAdd %5 %239 %240 -%242 = OpCompositeExtract %5 %226 1 -%243 = OpIAdd %5 %241 %242 -%244 = OpCompositeExtract %5 %227 1 -%245 = OpIAdd %5 %243 %244 -%246 = OpCompositeExtract %5 %228 1 -%247 = OpIAdd %5 %245 %246 -%248 = OpCompositeExtract %5 %230 1 -%249 = OpIAdd %5 %247 %248 -%250 = OpCompositeExtract %5 %232 1 -%251 = OpIAdd %5 %249 %250 -%252 = OpCompositeExtract %5 %233 2 -%253 = OpIAdd %5 %251 %252 -%254 = OpCompositeExtract %5 %234 2 -%255 = OpIAdd %5 %253 %254 +%239 = OpCompositeExtract %5 %224 1 +%240 = OpIAdd %5 %221 %239 +%241 = OpCompositeExtract %5 %225 1 +%242 = OpIAdd %5 %240 %241 +%243 = OpCompositeExtract %5 %227 1 +%244 = OpIAdd %5 %242 %243 +%245 = OpCompositeExtract %5 %229 1 +%246 = OpIAdd %5 %244 %245 +%247 = OpCompositeExtract %5 %230 1 +%248 = OpIAdd %5 %246 %247 +%249 = OpCompositeExtract %5 %231 1 +%250 = OpIAdd %5 %248 %249 +%251 = OpCompositeExtract %5 %233 1 +%252 = OpIAdd %5 %250 %251 +%253 = OpCompositeExtract %5 %235 1 +%254 = OpIAdd %5 %252 %253 +%255 = OpCompositeExtract %5 %236 2 +%256 = OpIAdd %5 %254 %255 +%257 = OpCompositeExtract %5 %237 2 +%258 = OpIAdd %5 %256 %257 OpLine %3 91 12 -%256 = OpConvertUToF %8 %255 -%257 = OpCompositeConstruct %24 %256 %256 %256 %256 -OpStore %206 %257 +%259 = OpConvertUToF %8 %258 +%260 = OpCompositeConstruct %24 %259 %259 %259 %259 +OpStore %209 %260 OpReturn OpFunctionEnd -%260 = OpFunction %2 None %97 -%258 = OpLabel -%261 = OpLoad %17 %46 -%262 = OpLoad %19 %51 -%263 = OpLoad %20 %53 -%264 = OpLoad %21 %55 -%265 = OpLoad %22 %57 -%266 = OpLoad %23 %59 -OpBranch %267 -%267 = OpLabel +%263 = OpFunction %2 None %97 +%261 = OpLabel +%264 = OpLoad %17 %46 +%265 = OpLoad %19 %51 +%266 = OpLoad %20 %53 +%267 = OpLoad %21 %55 +%268 = OpLoad %22 %57 +%269 = OpLoad %23 %59 +OpBranch %270 +%270 = OpLabel OpLine %3 96 25 -%268 = OpImageQueryLevels %5 %261 +%271 = OpImageQueryLevels %5 %264 OpLine %3 97 25 -%269 = OpImageQuerySizeLod %13 %262 %217 -%270 = OpCompositeExtract %5 %269 2 +%272 = OpImageQuerySizeLod %13 %265 %220 +%273 = OpCompositeExtract %5 %272 2 OpLine %3 98 31 -%271 = OpImageQueryLevels %5 %262 +%274 = OpImageQueryLevels %5 %265 OpLine %3 99 31 -%272 = OpImageQuerySizeLod %13 %262 %217 -%273 = OpCompositeExtract %5 %272 2 +%275 = OpImageQuerySizeLod %13 %265 %220 +%276 = OpCompositeExtract %5 %275 2 OpLine %3 100 27 -%274 = OpImageQueryLevels %5 %263 +%277 = OpImageQueryLevels %5 %266 OpLine %3 101 33 -%275 = OpImageQueryLevels %5 %264 +%278 = OpImageQueryLevels %5 %267 OpLine %3 102 27 -%276 = OpImageQuerySizeLod %13 %264 %217 -%277 = OpCompositeExtract %5 %276 2 +%279 = OpImageQuerySizeLod %13 %267 %220 +%280 = OpCompositeExtract %5 %279 2 OpLine %3 103 25 -%278 = OpImageQueryLevels %5 %265 +%281 = OpImageQueryLevels %5 %268 OpLine %3 104 26 -%279 = OpImageQuerySamples %5 %266 +%282 = OpImageQuerySamples %5 %269 OpLine %3 106 15 -%280 = OpIAdd %5 %270 %277 -%281 = OpIAdd %5 %280 %279 -%282 = OpIAdd %5 %281 %268 -%283 = OpIAdd %5 %282 %271 -%284 = OpIAdd %5 %283 %278 -%285 = OpIAdd %5 %284 %274 -%286 = OpIAdd %5 %285 %275 +%283 = OpIAdd %5 %273 %280 +%284 = OpIAdd %5 %283 %282 +%285 = OpIAdd %5 %284 %271 +%286 = OpIAdd %5 %285 %274 +%287 = OpIAdd %5 %286 %281 +%288 = OpIAdd %5 %287 %277 +%289 = OpIAdd %5 %288 %278 OpLine %3 108 12 -%287 = OpConvertUToF %8 %286 -%288 = OpCompositeConstruct %24 %287 %287 %287 %287 -OpStore %259 %288 +%290 = OpConvertUToF %8 %289 +%291 = OpCompositeConstruct %24 %290 %290 %290 %290 +OpStore %262 %291 OpReturn OpFunctionEnd -%291 = OpFunction %2 None %97 -%289 = OpLabel -%306 = OpVariable %307 Function %308 -%292 = OpLoad %16 %44 -%293 = OpLoad %17 %46 -%294 = OpLoad %19 %51 -%295 = OpLoad %21 %55 -%296 = OpLoad %25 %61 -OpBranch %309 -%309 = OpLabel +%294 = OpFunction %2 None %97 +%292 = OpLabel +%309 = OpVariable %310 Function %311 +%295 = OpLoad %16 %44 +%296 = OpLoad %17 %46 +%297 = OpLoad %19 %51 +%298 = OpLoad %21 %55 +%299 = OpLoad %25 %61 +OpBranch %312 +%312 = OpLabel OpLine %3 116 16 OpLine %3 117 17 OpLine %3 118 20 OpLine %3 121 5 -%311 = OpSampledImage %310 %292 %296 -%312 = OpImageSampleImplicitLod %24 %311 %297 -%313 = OpLoad %24 %306 -%314 = OpFAdd %24 %313 %312 +%314 = OpSampledImage %313 %295 %299 +%315 = OpImageSampleImplicitLod %24 %314 %300 +%316 = OpLoad %24 %309 +%317 = OpFAdd %24 %316 %315 OpLine %3 121 5 -OpStore %306 %314 +OpStore %309 %317 OpLine %3 122 5 -%316 = OpSampledImage %315 %293 %296 -%317 = OpImageSampleImplicitLod %24 %316 %299 -%318 = OpLoad %24 %306 -%319 = OpFAdd %24 %318 %317 +%319 = OpSampledImage %318 %296 %299 +%320 = OpImageSampleImplicitLod %24 %319 %302 +%321 = OpLoad %24 %309 +%322 = OpFAdd %24 %321 %320 OpLine %3 122 5 -OpStore %306 %319 +OpStore %309 %322 OpLine %3 123 5 -%320 = OpSampledImage %315 %293 %296 -%321 = OpImageSampleImplicitLod %24 %320 %299 ConstOffset %303 -%322 = OpLoad %24 %306 -%323 = OpFAdd %24 %322 %321 +%323 = OpSampledImage %318 %296 %299 +%324 = OpImageSampleImplicitLod %24 %323 %302 ConstOffset %306 +%325 = OpLoad %24 %309 +%326 = OpFAdd %24 %325 %324 OpLine %3 123 5 -OpStore %306 %323 +OpStore %309 %326 OpLine %3 124 5 -%324 = OpSampledImage %315 %293 %296 -%325 = OpImageSampleExplicitLod %24 %324 %299 Lod %304 -%326 = OpLoad %24 %306 -%327 = OpFAdd %24 %326 %325 +%327 = OpSampledImage %318 %296 %299 +%328 = OpImageSampleExplicitLod %24 %327 %302 Lod %307 +%329 = OpLoad %24 %309 +%330 = OpFAdd %24 %329 %328 OpLine %3 124 5 -OpStore %306 %327 +OpStore %309 %330 OpLine %3 125 5 -%328 = OpSampledImage %315 %293 %296 -%329 = OpImageSampleExplicitLod %24 %328 %299 Lod|ConstOffset %304 %303 -%330 = OpLoad %24 %306 -%331 = OpFAdd %24 %330 %329 +%331 = OpSampledImage %318 %296 %299 +%332 = OpImageSampleExplicitLod %24 %331 %302 Lod|ConstOffset %307 %306 +%333 = OpLoad %24 %309 +%334 = OpFAdd %24 %333 %332 OpLine %3 125 5 -OpStore %306 %331 +OpStore %309 %334 OpLine %3 126 5 -%332 = OpSampledImage %315 %293 %296 -%333 = OpImageSampleImplicitLod %24 %332 %299 Bias|ConstOffset %305 %303 -%334 = OpLoad %24 %306 -%335 = OpFAdd %24 %334 %333 +%335 = OpSampledImage %318 %296 %299 +%336 = OpImageSampleImplicitLod %24 %335 %302 Bias|ConstOffset %308 %306 +%337 = OpLoad %24 %309 +%338 = OpFAdd %24 %337 %336 OpLine %3 126 5 -OpStore %306 %335 +OpStore %309 %338 OpLine %3 127 5 -%337 = OpConvertUToF %8 %217 -%338 = OpCompositeConstruct %300 %299 %337 -%339 = OpSampledImage %336 %294 %296 -%340 = OpImageSampleImplicitLod %24 %339 %338 -%341 = OpLoad %24 %306 -%342 = OpFAdd %24 %341 %340 +%340 = OpConvertUToF %8 %220 +%341 = OpCompositeConstruct %303 %302 %340 +%342 = OpSampledImage %339 %297 %299 +%343 = OpImageSampleImplicitLod %24 %342 %341 +%344 = OpLoad %24 %309 +%345 = OpFAdd %24 %344 %343 OpLine %3 127 5 -OpStore %306 %342 +OpStore %309 %345 OpLine %3 128 5 -%343 = OpConvertUToF %8 %217 -%344 = OpCompositeConstruct %300 %299 %343 -%345 = OpSampledImage %336 %294 %296 -%346 = OpImageSampleImplicitLod %24 %345 %344 ConstOffset %303 -%347 = OpLoad %24 %306 -%348 = OpFAdd %24 %347 %346 +%346 = OpConvertUToF %8 %220 +%347 = OpCompositeConstruct %303 %302 %346 +%348 = OpSampledImage %339 %297 %299 +%349 = OpImageSampleImplicitLod %24 %348 %347 ConstOffset %306 +%350 = OpLoad %24 %309 +%351 = OpFAdd %24 %350 %349 OpLine %3 128 5 -OpStore %306 %348 +OpStore %309 %351 OpLine %3 129 5 -%349 = OpConvertUToF %8 %217 -%350 = OpCompositeConstruct %300 %299 %349 -%351 = OpSampledImage %336 %294 %296 -%352 = OpImageSampleExplicitLod %24 %351 %350 Lod %304 -%353 = OpLoad %24 %306 -%354 = OpFAdd %24 %353 %352 +%352 = OpConvertUToF %8 %220 +%353 = OpCompositeConstruct %303 %302 %352 +%354 = OpSampledImage %339 %297 %299 +%355 = OpImageSampleExplicitLod %24 %354 %353 Lod %307 +%356 = OpLoad %24 %309 +%357 = OpFAdd %24 %356 %355 OpLine %3 129 5 -OpStore %306 %354 +OpStore %309 %357 OpLine %3 130 5 -%355 = OpConvertUToF %8 %217 -%356 = OpCompositeConstruct %300 %299 %355 -%357 = OpSampledImage %336 %294 %296 -%358 = OpImageSampleExplicitLod %24 %357 %356 Lod|ConstOffset %304 %303 -%359 = OpLoad %24 %306 -%360 = OpFAdd %24 %359 %358 +%358 = OpConvertUToF %8 %220 +%359 = OpCompositeConstruct %303 %302 %358 +%360 = OpSampledImage %339 %297 %299 +%361 = OpImageSampleExplicitLod %24 %360 %359 Lod|ConstOffset %307 %306 +%362 = OpLoad %24 %309 +%363 = OpFAdd %24 %362 %361 OpLine %3 130 5 -OpStore %306 %360 +OpStore %309 %363 OpLine %3 131 5 -%361 = OpConvertUToF %8 %217 -%362 = OpCompositeConstruct %300 %299 %361 -%363 = OpSampledImage %336 %294 %296 -%364 = OpImageSampleImplicitLod %24 %363 %362 Bias|ConstOffset %305 %303 -%365 = OpLoad %24 %306 -%366 = OpFAdd %24 %365 %364 +%364 = OpConvertUToF %8 %220 +%365 = OpCompositeConstruct %303 %302 %364 +%366 = OpSampledImage %339 %297 %299 +%367 = OpImageSampleImplicitLod %24 %366 %365 Bias|ConstOffset %308 %306 +%368 = OpLoad %24 %309 +%369 = OpFAdd %24 %368 %367 OpLine %3 131 5 -OpStore %306 %366 +OpStore %309 %369 OpLine %3 132 5 -%367 = OpConvertSToF %8 %77 -%368 = OpCompositeConstruct %300 %299 %367 -%369 = OpSampledImage %336 %294 %296 -%370 = OpImageSampleImplicitLod %24 %369 %368 -%371 = OpLoad %24 %306 -%372 = OpFAdd %24 %371 %370 +%370 = OpConvertSToF %8 %77 +%371 = OpCompositeConstruct %303 %302 %370 +%372 = OpSampledImage %339 %297 %299 +%373 = OpImageSampleImplicitLod %24 %372 %371 +%374 = OpLoad %24 %309 +%375 = OpFAdd %24 %374 %373 OpLine %3 132 5 -OpStore %306 %372 +OpStore %309 %375 OpLine %3 133 5 -%373 = OpConvertSToF %8 %77 -%374 = OpCompositeConstruct %300 %299 %373 -%375 = OpSampledImage %336 %294 %296 -%376 = OpImageSampleImplicitLod %24 %375 %374 ConstOffset %303 -%377 = OpLoad %24 %306 -%378 = OpFAdd %24 %377 %376 +%376 = OpConvertSToF %8 %77 +%377 = OpCompositeConstruct %303 %302 %376 +%378 = OpSampledImage %339 %297 %299 +%379 = OpImageSampleImplicitLod %24 %378 %377 ConstOffset %306 +%380 = OpLoad %24 %309 +%381 = OpFAdd %24 %380 %379 OpLine %3 133 5 -OpStore %306 %378 +OpStore %309 %381 OpLine %3 134 5 -%379 = OpConvertSToF %8 %77 -%380 = OpCompositeConstruct %300 %299 %379 -%381 = OpSampledImage %336 %294 %296 -%382 = OpImageSampleExplicitLod %24 %381 %380 Lod %304 -%383 = OpLoad %24 %306 -%384 = OpFAdd %24 %383 %382 +%382 = OpConvertSToF %8 %77 +%383 = OpCompositeConstruct %303 %302 %382 +%384 = OpSampledImage %339 %297 %299 +%385 = OpImageSampleExplicitLod %24 %384 %383 Lod %307 +%386 = OpLoad %24 %309 +%387 = OpFAdd %24 %386 %385 OpLine %3 134 5 -OpStore %306 %384 +OpStore %309 %387 OpLine %3 135 5 -%385 = OpConvertSToF %8 %77 -%386 = OpCompositeConstruct %300 %299 %385 -%387 = OpSampledImage %336 %294 %296 -%388 = OpImageSampleExplicitLod %24 %387 %386 Lod|ConstOffset %304 %303 -%389 = OpLoad %24 %306 -%390 = OpFAdd %24 %389 %388 +%388 = OpConvertSToF %8 %77 +%389 = OpCompositeConstruct %303 %302 %388 +%390 = OpSampledImage %339 %297 %299 +%391 = OpImageSampleExplicitLod %24 %390 %389 Lod|ConstOffset %307 %306 +%392 = OpLoad %24 %309 +%393 = OpFAdd %24 %392 %391 OpLine %3 135 5 -OpStore %306 %390 +OpStore %309 %393 OpLine %3 136 5 -%391 = OpConvertSToF %8 %77 -%392 = OpCompositeConstruct %300 %299 %391 -%393 = OpSampledImage %336 %294 %296 -%394 = OpImageSampleImplicitLod %24 %393 %392 Bias|ConstOffset %305 %303 -%395 = OpLoad %24 %306 -%396 = OpFAdd %24 %395 %394 +%394 = OpConvertSToF %8 %77 +%395 = OpCompositeConstruct %303 %302 %394 +%396 = OpSampledImage %339 %297 %299 +%397 = OpImageSampleImplicitLod %24 %396 %395 Bias|ConstOffset %308 %306 +%398 = OpLoad %24 %309 +%399 = OpFAdd %24 %398 %397 OpLine %3 136 5 -OpStore %306 %396 +OpStore %309 %399 OpLine %3 137 5 -%398 = OpConvertUToF %8 %217 -%399 = OpCompositeConstruct %24 %301 %398 -%400 = OpSampledImage %397 %295 %296 -%401 = OpImageSampleImplicitLod %24 %400 %399 -%402 = OpLoad %24 %306 -%403 = OpFAdd %24 %402 %401 +%401 = OpConvertUToF %8 %220 +%402 = OpCompositeConstruct %24 %304 %401 +%403 = OpSampledImage %400 %298 %299 +%404 = OpImageSampleImplicitLod %24 %403 %402 +%405 = OpLoad %24 %309 +%406 = OpFAdd %24 %405 %404 OpLine %3 137 5 -OpStore %306 %403 +OpStore %309 %406 OpLine %3 138 5 -%404 = OpConvertUToF %8 %217 -%405 = OpCompositeConstruct %24 %301 %404 -%406 = OpSampledImage %397 %295 %296 -%407 = OpImageSampleExplicitLod %24 %406 %405 Lod %304 -%408 = OpLoad %24 %306 -%409 = OpFAdd %24 %408 %407 +%407 = OpConvertUToF %8 %220 +%408 = OpCompositeConstruct %24 %304 %407 +%409 = OpSampledImage %400 %298 %299 +%410 = OpImageSampleExplicitLod %24 %409 %408 Lod %307 +%411 = OpLoad %24 %309 +%412 = OpFAdd %24 %411 %410 OpLine %3 138 5 -OpStore %306 %409 +OpStore %309 %412 OpLine %3 139 5 -%410 = OpConvertUToF %8 %217 -%411 = OpCompositeConstruct %24 %301 %410 -%412 = OpSampledImage %397 %295 %296 -%413 = OpImageSampleImplicitLod %24 %412 %411 Bias %305 -%414 = OpLoad %24 %306 -%415 = OpFAdd %24 %414 %413 +%413 = OpConvertUToF %8 %220 +%414 = OpCompositeConstruct %24 %304 %413 +%415 = OpSampledImage %400 %298 %299 +%416 = OpImageSampleImplicitLod %24 %415 %414 Bias %308 +%417 = OpLoad %24 %309 +%418 = OpFAdd %24 %417 %416 OpLine %3 139 5 -OpStore %306 %415 +OpStore %309 %418 OpLine %3 140 5 -%416 = OpConvertSToF %8 %77 -%417 = OpCompositeConstruct %24 %301 %416 -%418 = OpSampledImage %397 %295 %296 -%419 = OpImageSampleImplicitLod %24 %418 %417 -%420 = OpLoad %24 %306 -%421 = OpFAdd %24 %420 %419 +%419 = OpConvertSToF %8 %77 +%420 = OpCompositeConstruct %24 %304 %419 +%421 = OpSampledImage %400 %298 %299 +%422 = OpImageSampleImplicitLod %24 %421 %420 +%423 = OpLoad %24 %309 +%424 = OpFAdd %24 %423 %422 OpLine %3 140 5 -OpStore %306 %421 +OpStore %309 %424 OpLine %3 141 5 -%422 = OpConvertSToF %8 %77 -%423 = OpCompositeConstruct %24 %301 %422 -%424 = OpSampledImage %397 %295 %296 -%425 = OpImageSampleExplicitLod %24 %424 %423 Lod %304 -%426 = OpLoad %24 %306 -%427 = OpFAdd %24 %426 %425 +%425 = OpConvertSToF %8 %77 +%426 = OpCompositeConstruct %24 %304 %425 +%427 = OpSampledImage %400 %298 %299 +%428 = OpImageSampleExplicitLod %24 %427 %426 Lod %307 +%429 = OpLoad %24 %309 +%430 = OpFAdd %24 %429 %428 OpLine %3 141 5 -OpStore %306 %427 +OpStore %309 %430 OpLine %3 142 5 -%428 = OpConvertSToF %8 %77 -%429 = OpCompositeConstruct %24 %301 %428 -%430 = OpSampledImage %397 %295 %296 -%431 = OpImageSampleImplicitLod %24 %430 %429 Bias %305 -%432 = OpLoad %24 %306 -%433 = OpFAdd %24 %432 %431 +%431 = OpConvertSToF %8 %77 +%432 = OpCompositeConstruct %24 %304 %431 +%433 = OpSampledImage %400 %298 %299 +%434 = OpImageSampleImplicitLod %24 %433 %432 Bias %308 +%435 = OpLoad %24 %309 +%436 = OpFAdd %24 %435 %434 OpLine %3 142 5 -OpStore %306 %433 +OpStore %309 %436 OpLine %3 1 1 -%434 = OpLoad %24 %306 -OpStore %290 %434 +%437 = OpLoad %24 %309 +OpStore %293 %437 OpReturn OpFunctionEnd -%438 = OpFunction %2 None %97 -%435 = OpLabel -%443 = OpVariable %444 Function %445 -%439 = OpLoad %25 %63 -%440 = OpLoad %26 %64 -%441 = OpLoad %27 %66 -%442 = OpLoad %28 %68 -OpBranch %446 -%446 = OpLabel +%441 = OpFunction %2 None %97 +%438 = OpLabel +%446 = OpVariable %447 Function %448 +%442 = OpLoad %25 %63 +%443 = OpLoad %26 %64 +%444 = OpLoad %27 %66 +%445 = OpLoad %28 %68 +OpBranch %449 +%449 = OpLabel OpLine %3 157 14 OpLine %3 158 15 OpLine %3 161 5 -%448 = OpSampledImage %447 %440 %439 -%449 = OpImageSampleDrefImplicitLod %8 %448 %299 %297 -%450 = OpLoad %8 %443 -%451 = OpFAdd %8 %450 %449 +%451 = OpSampledImage %450 %443 %442 +%452 = OpImageSampleDrefImplicitLod %8 %451 %302 %300 +%453 = OpLoad %8 %446 +%454 = OpFAdd %8 %453 %452 OpLine %3 161 5 -OpStore %443 %451 +OpStore %446 %454 OpLine %3 162 5 -%453 = OpConvertUToF %8 %217 -%454 = OpCompositeConstruct %300 %299 %453 -%455 = OpSampledImage %452 %441 %439 -%456 = OpImageSampleDrefImplicitLod %8 %455 %454 %297 -%457 = OpLoad %8 %443 -%458 = OpFAdd %8 %457 %456 +%456 = OpConvertUToF %8 %220 +%457 = OpCompositeConstruct %303 %302 %456 +%458 = OpSampledImage %455 %444 %442 +%459 = OpImageSampleDrefImplicitLod %8 %458 %457 %300 +%460 = OpLoad %8 %446 +%461 = OpFAdd %8 %460 %459 OpLine %3 162 5 -OpStore %443 %458 +OpStore %446 %461 OpLine %3 163 5 -%459 = OpConvertSToF %8 %77 -%460 = OpCompositeConstruct %300 %299 %459 -%461 = OpSampledImage %452 %441 %439 -%462 = OpImageSampleDrefImplicitLod %8 %461 %460 %297 -%463 = OpLoad %8 %443 -%464 = OpFAdd %8 %463 %462 +%462 = OpConvertSToF %8 %77 +%463 = OpCompositeConstruct %303 %302 %462 +%464 = OpSampledImage %455 %444 %442 +%465 = OpImageSampleDrefImplicitLod %8 %464 %463 %300 +%466 = OpLoad %8 %446 +%467 = OpFAdd %8 %466 %465 OpLine %3 163 5 -OpStore %443 %464 +OpStore %446 %467 OpLine %3 164 5 -%466 = OpSampledImage %465 %442 %439 -%467 = OpImageSampleDrefImplicitLod %8 %466 %301 %297 -%468 = OpLoad %8 %443 -%469 = OpFAdd %8 %468 %467 +%469 = OpSampledImage %468 %445 %442 +%470 = OpImageSampleDrefImplicitLod %8 %469 %304 %300 +%471 = OpLoad %8 %446 +%472 = OpFAdd %8 %471 %470 OpLine %3 164 5 -OpStore %443 %469 +OpStore %446 %472 OpLine %3 165 5 -%470 = OpSampledImage %447 %440 %439 -%471 = OpImageSampleDrefExplicitLod %8 %470 %299 %297 Lod %472 -%473 = OpLoad %8 %443 -%474 = OpFAdd %8 %473 %471 +%473 = OpSampledImage %450 %443 %442 +%474 = OpImageSampleDrefExplicitLod %8 %473 %302 %300 Lod %203 +%475 = OpLoad %8 %446 +%476 = OpFAdd %8 %475 %474 OpLine %3 165 5 -OpStore %443 %474 +OpStore %446 %476 OpLine %3 166 5 -%475 = OpConvertUToF %8 %217 -%476 = OpCompositeConstruct %300 %299 %475 -%477 = OpSampledImage %452 %441 %439 -%478 = OpImageSampleDrefExplicitLod %8 %477 %476 %297 Lod %472 -%479 = OpLoad %8 %443 -%480 = OpFAdd %8 %479 %478 +%477 = OpConvertUToF %8 %220 +%478 = OpCompositeConstruct %303 %302 %477 +%479 = OpSampledImage %455 %444 %442 +%480 = OpImageSampleDrefExplicitLod %8 %479 %478 %300 Lod %203 +%481 = OpLoad %8 %446 +%482 = OpFAdd %8 %481 %480 OpLine %3 166 5 -OpStore %443 %480 +OpStore %446 %482 OpLine %3 167 5 -%481 = OpConvertSToF %8 %77 -%482 = OpCompositeConstruct %300 %299 %481 -%483 = OpSampledImage %452 %441 %439 -%484 = OpImageSampleDrefExplicitLod %8 %483 %482 %297 Lod %472 -%485 = OpLoad %8 %443 -%486 = OpFAdd %8 %485 %484 +%483 = OpConvertSToF %8 %77 +%484 = OpCompositeConstruct %303 %302 %483 +%485 = OpSampledImage %455 %444 %442 +%486 = OpImageSampleDrefExplicitLod %8 %485 %484 %300 Lod %203 +%487 = OpLoad %8 %446 +%488 = OpFAdd %8 %487 %486 OpLine %3 167 5 -OpStore %443 %486 +OpStore %446 %488 OpLine %3 168 5 -%487 = OpSampledImage %465 %442 %439 -%488 = OpImageSampleDrefExplicitLod %8 %487 %301 %297 Lod %472 -%489 = OpLoad %8 %443 -%490 = OpFAdd %8 %489 %488 +%489 = OpSampledImage %468 %445 %442 +%490 = OpImageSampleDrefExplicitLod %8 %489 %304 %300 Lod %203 +%491 = OpLoad %8 %446 +%492 = OpFAdd %8 %491 %490 OpLine %3 168 5 -OpStore %443 %490 +OpStore %446 %492 OpLine %3 1 1 -%491 = OpLoad %8 %443 -OpStore %436 %491 +%493 = OpLoad %8 %446 +OpStore %439 %493 OpReturn OpFunctionEnd -%494 = OpFunction %2 None %97 -%492 = OpLabel -%495 = OpLoad %17 %46 -%496 = OpLoad %4 %48 -%497 = OpLoad %18 %49 -%498 = OpLoad %25 %61 -%499 = OpLoad %25 %63 -%500 = OpLoad %26 %64 -OpBranch %501 -%501 = OpLabel +%496 = OpFunction %2 None %97 +%494 = OpLabel +%497 = OpLoad %17 %46 +%498 = OpLoad %4 %48 +%499 = OpLoad %18 %49 +%500 = OpLoad %25 %61 +%501 = OpLoad %25 %63 +%502 = OpLoad %26 %64 +OpBranch %503 +%503 = OpLabel OpLine %3 174 14 OpLine %3 176 15 -%502 = OpSampledImage %315 %495 %498 -%503 = OpImageGather %24 %502 %299 %504 +%504 = OpSampledImage %318 %497 %500 +%505 = OpImageGather %24 %504 %302 %506 OpLine %3 177 22 -%505 = OpSampledImage %315 %495 %498 -%506 = OpImageGather %24 %505 %299 %507 ConstOffset %303 +%507 = OpSampledImage %318 %497 %500 +%508 = OpImageGather %24 %507 %302 %509 ConstOffset %306 OpLine %3 178 21 -%508 = OpSampledImage %447 %500 %499 -%509 = OpImageDrefGather %24 %508 %299 %297 +%510 = OpSampledImage %450 %502 %501 +%511 = OpImageDrefGather %24 %510 %302 %300 OpLine %3 179 28 -%510 = OpSampledImage %447 %500 %499 -%511 = OpImageDrefGather %24 %510 %299 %297 ConstOffset %303 +%512 = OpSampledImage %450 %502 %501 +%513 = OpImageDrefGather %24 %512 %302 %300 ConstOffset %306 OpLine %3 181 13 -%513 = OpSampledImage %512 %496 %498 -%514 = OpImageGather %116 %513 %299 %217 +%515 = OpSampledImage %514 %498 %500 +%516 = OpImageGather %116 %515 %302 %220 OpLine %3 182 13 -%517 = OpSampledImage %516 %497 %498 -%518 = OpImageGather %515 %517 %299 %217 +%519 = OpSampledImage %518 %499 %500 +%520 = OpImageGather %517 %519 %302 %220 OpLine %3 183 13 -%519 = OpConvertUToF %24 %514 -%520 = OpConvertSToF %24 %518 -%521 = OpFAdd %24 %519 %520 +%521 = OpConvertUToF %24 %516 +%522 = OpConvertSToF %24 %520 +%523 = OpFAdd %24 %521 %522 OpLine %3 185 12 -%522 = OpFAdd %24 %503 %506 -%523 = OpFAdd %24 %522 %509 -%524 = OpFAdd %24 %523 %511 -%525 = OpFAdd %24 %524 %521 -OpStore %493 %525 +%524 = OpFAdd %24 %505 %508 +%525 = OpFAdd %24 %524 %511 +%526 = OpFAdd %24 %525 %513 +%527 = OpFAdd %24 %526 %523 +OpStore %495 %527 OpReturn OpFunctionEnd -%528 = OpFunction %2 None %97 -%526 = OpLabel -%529 = OpLoad %25 %61 -%530 = OpLoad %26 %64 -OpBranch %531 -%531 = OpLabel +%530 = OpFunction %2 None %97 +%528 = OpLabel +%531 = OpLoad %25 %61 +%532 = OpLoad %26 %64 +OpBranch %533 +%533 = OpLabel OpLine %3 190 14 OpLine %3 192 15 -%532 = OpSampledImage %447 %530 %529 -%533 = OpImageSampleImplicitLod %24 %532 %299 -%534 = OpCompositeExtract %8 %533 0 +%534 = OpSampledImage %450 %532 %531 +%535 = OpImageSampleImplicitLod %24 %534 %302 +%536 = OpCompositeExtract %8 %535 0 OpLine %3 193 22 -%535 = OpSampledImage %447 %530 %529 -%536 = OpImageGather %24 %535 %299 %217 +%537 = OpSampledImage %450 %532 %531 +%538 = OpImageGather %24 %537 %302 %220 OpLine %3 194 21 -%537 = OpSampledImage %447 %530 %529 -%539 = OpConvertSToF %8 %88 -%538 = OpImageSampleExplicitLod %24 %537 %299 Lod %539 -%540 = OpCompositeExtract %8 %538 0 +%539 = OpSampledImage %450 %532 %531 +%541 = OpConvertSToF %8 %88 +%540 = OpImageSampleExplicitLod %24 %539 %302 Lod %541 +%542 = OpCompositeExtract %8 %540 0 OpLine %3 192 15 -%541 = OpCompositeConstruct %24 %534 %534 %534 %534 -%542 = OpFAdd %24 %541 %536 -%543 = OpCompositeConstruct %24 %540 %540 %540 %540 -%544 = OpFAdd %24 %542 %543 -OpStore %527 %544 +%543 = OpCompositeConstruct %24 %536 %536 %536 %536 +%544 = OpFAdd %24 %543 %538 +%545 = OpCompositeConstruct %24 %542 %542 %542 %542 +%546 = OpFAdd %24 %544 %545 +OpStore %529 %546 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/spv/int64.spvasm b/naga/tests/out/spv/int64.spvasm index eae78d3e288..377247a2159 100644 --- a/naga/tests/out/spv/int64.spvasm +++ b/naga/tests/out/spv/int64.spvasm @@ -1,14 +1,14 @@ ; SPIR-V ; Version: 1.0 ; Generator: rspirv -; Bound: 378 +; Bound: 384 OpCapability Shader OpCapability Int64 OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %362 "main" -OpExecutionMode %362 LocalSize 1 1 1 +OpEntryPoint GLCompute %368 "main" +OpExecutionMode %368 LocalSize 1 1 1 OpMemberDecorate %14 0 Offset 0 OpMemberDecorate %14 1 Offset 4 OpMemberDecorate %14 2 Offset 8 @@ -98,43 +98,47 @@ OpMemberDecorate %36 0 Offset 0 %76 = OpTypePointer Uniform %6 %77 = OpConstant %5 1 %86 = OpTypePointer Uniform %7 -%95 = OpTypePointer Uniform %3 -%96 = OpConstant %5 7 -%103 = OpTypePointer Uniform %4 -%104 = OpConstant %5 3 -%110 = OpTypePointer Uniform %8 -%111 = OpConstant %5 4 -%118 = OpTypePointer Uniform %9 -%119 = OpConstant %5 5 -%126 = OpTypePointer Uniform %10 -%127 = OpConstant %5 6 -%137 = OpTypePointer StorageBuffer %3 -%144 = OpTypePointer StorageBuffer %11 -%145 = OpTypePointer Uniform %11 -%146 = OpConstant %5 8 -%153 = OpTypePointer StorageBuffer %12 -%154 = OpTypePointer Uniform %12 -%155 = OpConstant %5 9 -%162 = OpTypePointer StorageBuffer %13 -%163 = OpTypePointer Uniform %13 -%164 = OpConstant %5 10 -%171 = OpTypePointer StorageBuffer %17 -%191 = OpConstantNull %3 -%219 = OpTypeFunction %4 %4 -%225 = OpConstant %4 31 -%226 = OpConstant %4 18446744073709551615 -%227 = OpConstant %4 5 -%229 = OpTypePointer Function %4 -%292 = OpTypePointer StorageBuffer %4 -%299 = OpTypePointer StorageBuffer %8 -%306 = OpTypePointer StorageBuffer %9 -%313 = OpTypePointer StorageBuffer %10 -%320 = OpTypePointer StorageBuffer %15 -%340 = OpConstantNull %4 -%363 = OpTypeFunction %2 -%369 = OpConstant %4 67 -%370 = OpConstant %3 60 -%376 = OpConstant %5 11 +%92 = OpConstant %7 -9.223372e18 +%93 = OpConstant %7 9.2233715e18 +%98 = OpTypePointer Uniform %3 +%99 = OpConstant %5 7 +%106 = OpTypePointer Uniform %4 +%107 = OpConstant %5 3 +%113 = OpTypePointer Uniform %8 +%114 = OpConstant %5 4 +%121 = OpTypePointer Uniform %9 +%122 = OpConstant %5 5 +%129 = OpTypePointer Uniform %10 +%130 = OpConstant %5 6 +%140 = OpTypePointer StorageBuffer %3 +%147 = OpTypePointer StorageBuffer %11 +%148 = OpTypePointer Uniform %11 +%149 = OpConstant %5 8 +%156 = OpTypePointer StorageBuffer %12 +%157 = OpTypePointer Uniform %12 +%158 = OpConstant %5 9 +%165 = OpTypePointer StorageBuffer %13 +%166 = OpTypePointer Uniform %13 +%167 = OpConstant %5 10 +%174 = OpTypePointer StorageBuffer %17 +%194 = OpConstantNull %3 +%222 = OpTypeFunction %4 %4 +%228 = OpConstant %4 31 +%229 = OpConstant %4 18446744073709551615 +%230 = OpConstant %4 5 +%232 = OpTypePointer Function %4 +%263 = OpConstant %7 0.0 +%264 = OpConstant %7 1.8446743e19 +%298 = OpTypePointer StorageBuffer %4 +%305 = OpTypePointer StorageBuffer %8 +%312 = OpTypePointer StorageBuffer %9 +%319 = OpTypePointer StorageBuffer %10 +%326 = OpTypePointer StorageBuffer %15 +%346 = OpConstantNull %4 +%369 = OpTypeFunction %2 +%375 = OpConstant %4 67 +%376 = OpConstant %3 60 +%382 = OpConstant %5 11 %40 = OpFunction %3 None %41 %39 = OpFunctionParameter %3 %38 = OpLabel @@ -179,299 +183,301 @@ OpStore %56 %85 %89 = OpLoad %3 %56 %90 = OpConvertSToF %7 %89 %91 = OpFAdd %7 %88 %90 -%92 = OpConvertFToS %3 %91 -%93 = OpLoad %3 %56 -%94 = OpIAdd %3 %93 %92 -OpStore %56 %94 -%97 = OpAccessChain %95 %44 %96 -%98 = OpLoad %3 %97 -%99 = OpCompositeConstruct %12 %98 %98 %98 -%100 = OpCompositeExtract %3 %99 2 -%101 = OpLoad %3 %56 -%102 = OpIAdd %3 %101 %100 -OpStore %56 %102 -%105 = OpAccessChain %103 %44 %104 -%106 = OpLoad %4 %105 -%107 = OpBitcast %3 %106 -%108 = OpLoad %3 %56 -%109 = OpIAdd %3 %108 %107 -OpStore %56 %109 -%112 = OpAccessChain %110 %44 %111 -%113 = OpLoad %8 %112 -%114 = OpBitcast %11 %113 -%115 = OpCompositeExtract %3 %114 1 -%116 = OpLoad %3 %56 -%117 = OpIAdd %3 %116 %115 -OpStore %56 %117 -%120 = OpAccessChain %118 %44 %119 -%121 = OpLoad %9 %120 -%122 = OpBitcast %12 %121 -%123 = OpCompositeExtract %3 %122 2 -%124 = OpLoad %3 %56 -%125 = OpIAdd %3 %124 %123 -OpStore %56 %125 -%128 = OpAccessChain %126 %44 %127 -%129 = OpLoad %10 %128 -%130 = OpBitcast %13 %129 -%131 = OpCompositeExtract %3 %130 3 -%132 = OpLoad %3 %56 -%133 = OpIAdd %3 %132 %131 -OpStore %56 %133 -%134 = OpISub %3 %54 %19 +%94 = OpExtInst %7 %1 FClamp %91 %92 %93 +%95 = OpConvertFToS %3 %94 +%96 = OpLoad %3 %56 +%97 = OpIAdd %3 %96 %95 +OpStore %56 %97 +%100 = OpAccessChain %98 %44 %99 +%101 = OpLoad %3 %100 +%102 = OpCompositeConstruct %12 %101 %101 %101 +%103 = OpCompositeExtract %3 %102 2 +%104 = OpLoad %3 %56 +%105 = OpIAdd %3 %104 %103 +OpStore %56 %105 +%108 = OpAccessChain %106 %44 %107 +%109 = OpLoad %4 %108 +%110 = OpBitcast %3 %109 +%111 = OpLoad %3 %56 +%112 = OpIAdd %3 %111 %110 +OpStore %56 %112 +%115 = OpAccessChain %113 %44 %114 +%116 = OpLoad %8 %115 +%117 = OpBitcast %11 %116 +%118 = OpCompositeExtract %3 %117 1 +%119 = OpLoad %3 %56 +%120 = OpIAdd %3 %119 %118 +OpStore %56 %120 +%123 = OpAccessChain %121 %44 %122 +%124 = OpLoad %9 %123 +%125 = OpBitcast %12 %124 +%126 = OpCompositeExtract %3 %125 2 +%127 = OpLoad %3 %56 +%128 = OpIAdd %3 %127 %126 +OpStore %56 %128 +%131 = OpAccessChain %129 %44 %130 +%132 = OpLoad %10 %131 +%133 = OpBitcast %13 %132 +%134 = OpCompositeExtract %3 %133 3 %135 = OpLoad %3 %56 %136 = OpIAdd %3 %135 %134 OpStore %56 %136 -%138 = OpAccessChain %95 %44 %96 -%139 = OpLoad %3 %138 -%140 = OpAccessChain %137 %46 %96 -%141 = OpLoad %3 %140 -%142 = OpIAdd %3 %139 %141 -%143 = OpAccessChain %137 %49 %96 -OpStore %143 %142 -%147 = OpAccessChain %145 %44 %146 -%148 = OpLoad %11 %147 -%149 = OpAccessChain %144 %46 %146 -%150 = OpLoad %11 %149 -%151 = OpIAdd %11 %148 %150 -%152 = OpAccessChain %144 %49 %146 -OpStore %152 %151 -%156 = OpAccessChain %154 %44 %155 -%157 = OpLoad %12 %156 -%158 = OpAccessChain %153 %46 %155 -%159 = OpLoad %12 %158 -%160 = OpIAdd %12 %157 %159 -%161 = OpAccessChain %153 %49 %155 -OpStore %161 %160 -%165 = OpAccessChain %163 %44 %164 -%166 = OpLoad %13 %165 -%167 = OpAccessChain %162 %46 %164 -%168 = OpLoad %13 %167 -%169 = OpIAdd %13 %166 %168 -%170 = OpAccessChain %162 %49 %164 -OpStore %170 %169 -%172 = OpAccessChain %171 %48 %77 -%173 = OpLoad %17 %172 -%174 = OpAccessChain %171 %50 %77 -OpStore %174 %173 -%175 = OpLoad %3 %56 -%176 = OpExtInst %3 %1 SAbs %175 -%177 = OpLoad %3 %56 -%178 = OpIAdd %3 %177 %176 -OpStore %56 %178 -%179 = OpLoad %3 %56 +%137 = OpISub %3 %54 %19 +%138 = OpLoad %3 %56 +%139 = OpIAdd %3 %138 %137 +OpStore %56 %139 +%141 = OpAccessChain %98 %44 %99 +%142 = OpLoad %3 %141 +%143 = OpAccessChain %140 %46 %99 +%144 = OpLoad %3 %143 +%145 = OpIAdd %3 %142 %144 +%146 = OpAccessChain %140 %49 %99 +OpStore %146 %145 +%150 = OpAccessChain %148 %44 %149 +%151 = OpLoad %11 %150 +%152 = OpAccessChain %147 %46 %149 +%153 = OpLoad %11 %152 +%154 = OpIAdd %11 %151 %153 +%155 = OpAccessChain %147 %49 %149 +OpStore %155 %154 +%159 = OpAccessChain %157 %44 %158 +%160 = OpLoad %12 %159 +%161 = OpAccessChain %156 %46 %158 +%162 = OpLoad %12 %161 +%163 = OpIAdd %12 %160 %162 +%164 = OpAccessChain %156 %49 %158 +OpStore %164 %163 +%168 = OpAccessChain %166 %44 %167 +%169 = OpLoad %13 %168 +%170 = OpAccessChain %165 %46 %167 +%171 = OpLoad %13 %170 +%172 = OpIAdd %13 %169 %171 +%173 = OpAccessChain %165 %49 %167 +OpStore %173 %172 +%175 = OpAccessChain %174 %48 %77 +%176 = OpLoad %17 %175 +%177 = OpAccessChain %174 %50 %77 +OpStore %177 %176 +%178 = OpLoad %3 %56 +%179 = OpExtInst %3 %1 SAbs %178 %180 = OpLoad %3 %56 -%181 = OpLoad %3 %56 -%183 = OpExtInst %3 %1 SMax %179 %180 -%182 = OpExtInst %3 %1 SMin %183 %181 +%181 = OpIAdd %3 %180 %179 +OpStore %56 %181 +%182 = OpLoad %3 %56 +%183 = OpLoad %3 %56 %184 = OpLoad %3 %56 -%185 = OpIAdd %3 %184 %182 -OpStore %56 %185 -%186 = OpLoad %3 %56 -%187 = OpCompositeConstruct %11 %186 %186 -%188 = OpLoad %3 %56 -%189 = OpCompositeConstruct %11 %188 %188 -%192 = OpCompositeExtract %3 %187 0 -%193 = OpCompositeExtract %3 %189 0 -%194 = OpIMul %3 %192 %193 -%195 = OpIAdd %3 %191 %194 -%196 = OpCompositeExtract %3 %187 1 -%197 = OpCompositeExtract %3 %189 1 -%198 = OpIMul %3 %196 %197 -%190 = OpIAdd %3 %195 %198 -%199 = OpLoad %3 %56 -%200 = OpIAdd %3 %199 %190 -OpStore %56 %200 -%201 = OpLoad %3 %56 +%186 = OpExtInst %3 %1 SMax %182 %183 +%185 = OpExtInst %3 %1 SMin %186 %184 +%187 = OpLoad %3 %56 +%188 = OpIAdd %3 %187 %185 +OpStore %56 %188 +%189 = OpLoad %3 %56 +%190 = OpCompositeConstruct %11 %189 %189 +%191 = OpLoad %3 %56 +%192 = OpCompositeConstruct %11 %191 %191 +%195 = OpCompositeExtract %3 %190 0 +%196 = OpCompositeExtract %3 %192 0 +%197 = OpIMul %3 %195 %196 +%198 = OpIAdd %3 %194 %197 +%199 = OpCompositeExtract %3 %190 1 +%200 = OpCompositeExtract %3 %192 1 +%201 = OpIMul %3 %199 %200 +%193 = OpIAdd %3 %198 %201 %202 = OpLoad %3 %56 -%203 = OpExtInst %3 %1 SMax %201 %202 +%203 = OpIAdd %3 %202 %193 +OpStore %56 %203 %204 = OpLoad %3 %56 -%205 = OpIAdd %3 %204 %203 -OpStore %56 %205 -%206 = OpLoad %3 %56 +%205 = OpLoad %3 %56 +%206 = OpExtInst %3 %1 SMax %204 %205 %207 = OpLoad %3 %56 -%208 = OpExtInst %3 %1 SMin %206 %207 +%208 = OpIAdd %3 %207 %206 +OpStore %56 %208 %209 = OpLoad %3 %56 -%210 = OpIAdd %3 %209 %208 -OpStore %56 %210 -%211 = OpLoad %3 %56 -%212 = OpExtInst %3 %1 SSign %211 -%213 = OpLoad %3 %56 -%214 = OpIAdd %3 %213 %212 -OpStore %56 %214 -%215 = OpLoad %3 %56 -OpReturnValue %215 +%210 = OpLoad %3 %56 +%211 = OpExtInst %3 %1 SMin %209 %210 +%212 = OpLoad %3 %56 +%213 = OpIAdd %3 %212 %211 +OpStore %56 %213 +%214 = OpLoad %3 %56 +%215 = OpExtInst %3 %1 SSign %214 +%216 = OpLoad %3 %56 +%217 = OpIAdd %3 %216 %215 +OpStore %56 %217 +%218 = OpLoad %3 %56 +OpReturnValue %218 OpFunctionEnd -%218 = OpFunction %4 None %219 -%217 = OpFunctionParameter %4 -%216 = OpLabel -%228 = OpVariable %229 Function %20 -%220 = OpAccessChain %42 %23 %43 -%221 = OpAccessChain %45 %26 %43 -%222 = OpAccessChain %47 %29 %43 -%223 = OpAccessChain %45 %32 %43 -%224 = OpAccessChain %47 %35 %43 -OpBranch %230 -%230 = OpLabel -%231 = OpIAdd %4 %225 %226 -%232 = OpISub %4 %231 %226 -%233 = OpLoad %4 %228 -%234 = OpIAdd %4 %233 %232 -OpStore %228 %234 -%235 = OpLoad %4 %228 -%236 = OpIAdd %4 %235 %227 -%237 = OpLoad %4 %228 -%238 = OpIAdd %4 %237 %236 -OpStore %228 %238 -%239 = OpAccessChain %67 %220 %43 -%240 = OpLoad %5 %239 -%241 = OpLoad %4 %228 -%242 = OpUConvert %5 %241 -%243 = OpIAdd %5 %240 %242 -%244 = OpUConvert %4 %243 -%245 = OpLoad %4 %228 -%246 = OpIAdd %4 %245 %244 -OpStore %228 %246 -%247 = OpAccessChain %76 %220 %77 -%248 = OpLoad %6 %247 -%249 = OpLoad %4 %228 -%250 = OpSConvert %6 %249 -%251 = OpIAdd %6 %248 %250 -%252 = OpUConvert %4 %251 -%253 = OpLoad %4 %228 -%254 = OpIAdd %4 %253 %252 -OpStore %228 %254 -%255 = OpAccessChain %86 %220 %16 -%256 = OpLoad %7 %255 -%257 = OpLoad %4 %228 -%258 = OpConvertUToF %7 %257 -%259 = OpFAdd %7 %256 %258 -%260 = OpConvertFToU %4 %259 -%261 = OpLoad %4 %228 -%262 = OpIAdd %4 %261 %260 -OpStore %228 %262 -%263 = OpAccessChain %103 %220 %104 -%264 = OpLoad %4 %263 -%265 = OpCompositeConstruct %9 %264 %264 %264 -%266 = OpCompositeExtract %4 %265 2 -%267 = OpLoad %4 %228 +%221 = OpFunction %4 None %222 +%220 = OpFunctionParameter %4 +%219 = OpLabel +%231 = OpVariable %232 Function %20 +%223 = OpAccessChain %42 %23 %43 +%224 = OpAccessChain %45 %26 %43 +%225 = OpAccessChain %47 %29 %43 +%226 = OpAccessChain %45 %32 %43 +%227 = OpAccessChain %47 %35 %43 +OpBranch %233 +%233 = OpLabel +%234 = OpIAdd %4 %228 %229 +%235 = OpISub %4 %234 %229 +%236 = OpLoad %4 %231 +%237 = OpIAdd %4 %236 %235 +OpStore %231 %237 +%238 = OpLoad %4 %231 +%239 = OpIAdd %4 %238 %230 +%240 = OpLoad %4 %231 +%241 = OpIAdd %4 %240 %239 +OpStore %231 %241 +%242 = OpAccessChain %67 %223 %43 +%243 = OpLoad %5 %242 +%244 = OpLoad %4 %231 +%245 = OpUConvert %5 %244 +%246 = OpIAdd %5 %243 %245 +%247 = OpUConvert %4 %246 +%248 = OpLoad %4 %231 +%249 = OpIAdd %4 %248 %247 +OpStore %231 %249 +%250 = OpAccessChain %76 %223 %77 +%251 = OpLoad %6 %250 +%252 = OpLoad %4 %231 +%253 = OpSConvert %6 %252 +%254 = OpIAdd %6 %251 %253 +%255 = OpUConvert %4 %254 +%256 = OpLoad %4 %231 +%257 = OpIAdd %4 %256 %255 +OpStore %231 %257 +%258 = OpAccessChain %86 %223 %16 +%259 = OpLoad %7 %258 +%260 = OpLoad %4 %231 +%261 = OpConvertUToF %7 %260 +%262 = OpFAdd %7 %259 %261 +%265 = OpExtInst %7 %1 FClamp %262 %263 %264 +%266 = OpConvertFToU %4 %265 +%267 = OpLoad %4 %231 %268 = OpIAdd %4 %267 %266 -OpStore %228 %268 -%269 = OpAccessChain %95 %220 %96 -%270 = OpLoad %3 %269 -%271 = OpBitcast %4 %270 -%272 = OpLoad %4 %228 -%273 = OpIAdd %4 %272 %271 -OpStore %228 %273 -%274 = OpAccessChain %145 %220 %146 -%275 = OpLoad %11 %274 -%276 = OpBitcast %8 %275 -%277 = OpCompositeExtract %4 %276 1 -%278 = OpLoad %4 %228 +OpStore %231 %268 +%269 = OpAccessChain %106 %223 %107 +%270 = OpLoad %4 %269 +%271 = OpCompositeConstruct %9 %270 %270 %270 +%272 = OpCompositeExtract %4 %271 2 +%273 = OpLoad %4 %231 +%274 = OpIAdd %4 %273 %272 +OpStore %231 %274 +%275 = OpAccessChain %98 %223 %99 +%276 = OpLoad %3 %275 +%277 = OpBitcast %4 %276 +%278 = OpLoad %4 %231 %279 = OpIAdd %4 %278 %277 -OpStore %228 %279 -%280 = OpAccessChain %154 %220 %155 -%281 = OpLoad %12 %280 -%282 = OpBitcast %9 %281 -%283 = OpCompositeExtract %4 %282 2 -%284 = OpLoad %4 %228 +OpStore %231 %279 +%280 = OpAccessChain %148 %223 %149 +%281 = OpLoad %11 %280 +%282 = OpBitcast %8 %281 +%283 = OpCompositeExtract %4 %282 1 +%284 = OpLoad %4 %231 %285 = OpIAdd %4 %284 %283 -OpStore %228 %285 -%286 = OpAccessChain %163 %220 %164 -%287 = OpLoad %13 %286 -%288 = OpBitcast %10 %287 -%289 = OpCompositeExtract %4 %288 3 -%290 = OpLoad %4 %228 +OpStore %231 %285 +%286 = OpAccessChain %157 %223 %158 +%287 = OpLoad %12 %286 +%288 = OpBitcast %9 %287 +%289 = OpCompositeExtract %4 %288 2 +%290 = OpLoad %4 %231 %291 = OpIAdd %4 %290 %289 -OpStore %228 %291 -%293 = OpAccessChain %103 %220 %104 -%294 = OpLoad %4 %293 -%295 = OpAccessChain %292 %221 %104 -%296 = OpLoad %4 %295 -%297 = OpIAdd %4 %294 %296 -%298 = OpAccessChain %292 %223 %104 -OpStore %298 %297 -%300 = OpAccessChain %110 %220 %111 -%301 = OpLoad %8 %300 -%302 = OpAccessChain %299 %221 %111 -%303 = OpLoad %8 %302 -%304 = OpIAdd %8 %301 %303 -%305 = OpAccessChain %299 %223 %111 -OpStore %305 %304 -%307 = OpAccessChain %118 %220 %119 -%308 = OpLoad %9 %307 -%309 = OpAccessChain %306 %221 %119 -%310 = OpLoad %9 %309 -%311 = OpIAdd %9 %308 %310 -%312 = OpAccessChain %306 %223 %119 -OpStore %312 %311 -%314 = OpAccessChain %126 %220 %127 -%315 = OpLoad %10 %314 -%316 = OpAccessChain %313 %221 %127 -%317 = OpLoad %10 %316 -%318 = OpIAdd %10 %315 %317 -%319 = OpAccessChain %313 %223 %127 -OpStore %319 %318 -%321 = OpAccessChain %320 %222 %43 -%322 = OpLoad %15 %321 -%323 = OpAccessChain %320 %224 %43 -OpStore %323 %322 -%324 = OpLoad %4 %228 -%325 = OpCopyObject %4 %324 -%326 = OpLoad %4 %228 -%327 = OpIAdd %4 %326 %325 -OpStore %228 %327 -%328 = OpLoad %4 %228 -%329 = OpLoad %4 %228 -%330 = OpLoad %4 %228 -%332 = OpExtInst %4 %1 UMax %328 %329 -%331 = OpExtInst %4 %1 UMin %332 %330 -%333 = OpLoad %4 %228 -%334 = OpIAdd %4 %333 %331 -OpStore %228 %334 -%335 = OpLoad %4 %228 -%336 = OpCompositeConstruct %8 %335 %335 -%337 = OpLoad %4 %228 -%338 = OpCompositeConstruct %8 %337 %337 -%341 = OpCompositeExtract %4 %336 0 -%342 = OpCompositeExtract %4 %338 0 -%343 = OpIMul %4 %341 %342 -%344 = OpIAdd %4 %340 %343 -%345 = OpCompositeExtract %4 %336 1 -%346 = OpCompositeExtract %4 %338 1 -%347 = OpIMul %4 %345 %346 -%339 = OpIAdd %4 %344 %347 -%348 = OpLoad %4 %228 -%349 = OpIAdd %4 %348 %339 -OpStore %228 %349 -%350 = OpLoad %4 %228 -%351 = OpLoad %4 %228 -%352 = OpExtInst %4 %1 UMax %350 %351 -%353 = OpLoad %4 %228 -%354 = OpIAdd %4 %353 %352 -OpStore %228 %354 -%355 = OpLoad %4 %228 -%356 = OpLoad %4 %228 -%357 = OpExtInst %4 %1 UMin %355 %356 -%358 = OpLoad %4 %228 -%359 = OpIAdd %4 %358 %357 -OpStore %228 %359 -%360 = OpLoad %4 %228 -OpReturnValue %360 +OpStore %231 %291 +%292 = OpAccessChain %166 %223 %167 +%293 = OpLoad %13 %292 +%294 = OpBitcast %10 %293 +%295 = OpCompositeExtract %4 %294 3 +%296 = OpLoad %4 %231 +%297 = OpIAdd %4 %296 %295 +OpStore %231 %297 +%299 = OpAccessChain %106 %223 %107 +%300 = OpLoad %4 %299 +%301 = OpAccessChain %298 %224 %107 +%302 = OpLoad %4 %301 +%303 = OpIAdd %4 %300 %302 +%304 = OpAccessChain %298 %226 %107 +OpStore %304 %303 +%306 = OpAccessChain %113 %223 %114 +%307 = OpLoad %8 %306 +%308 = OpAccessChain %305 %224 %114 +%309 = OpLoad %8 %308 +%310 = OpIAdd %8 %307 %309 +%311 = OpAccessChain %305 %226 %114 +OpStore %311 %310 +%313 = OpAccessChain %121 %223 %122 +%314 = OpLoad %9 %313 +%315 = OpAccessChain %312 %224 %122 +%316 = OpLoad %9 %315 +%317 = OpIAdd %9 %314 %316 +%318 = OpAccessChain %312 %226 %122 +OpStore %318 %317 +%320 = OpAccessChain %129 %223 %130 +%321 = OpLoad %10 %320 +%322 = OpAccessChain %319 %224 %130 +%323 = OpLoad %10 %322 +%324 = OpIAdd %10 %321 %323 +%325 = OpAccessChain %319 %226 %130 +OpStore %325 %324 +%327 = OpAccessChain %326 %225 %43 +%328 = OpLoad %15 %327 +%329 = OpAccessChain %326 %227 %43 +OpStore %329 %328 +%330 = OpLoad %4 %231 +%331 = OpCopyObject %4 %330 +%332 = OpLoad %4 %231 +%333 = OpIAdd %4 %332 %331 +OpStore %231 %333 +%334 = OpLoad %4 %231 +%335 = OpLoad %4 %231 +%336 = OpLoad %4 %231 +%338 = OpExtInst %4 %1 UMax %334 %335 +%337 = OpExtInst %4 %1 UMin %338 %336 +%339 = OpLoad %4 %231 +%340 = OpIAdd %4 %339 %337 +OpStore %231 %340 +%341 = OpLoad %4 %231 +%342 = OpCompositeConstruct %8 %341 %341 +%343 = OpLoad %4 %231 +%344 = OpCompositeConstruct %8 %343 %343 +%347 = OpCompositeExtract %4 %342 0 +%348 = OpCompositeExtract %4 %344 0 +%349 = OpIMul %4 %347 %348 +%350 = OpIAdd %4 %346 %349 +%351 = OpCompositeExtract %4 %342 1 +%352 = OpCompositeExtract %4 %344 1 +%353 = OpIMul %4 %351 %352 +%345 = OpIAdd %4 %350 %353 +%354 = OpLoad %4 %231 +%355 = OpIAdd %4 %354 %345 +OpStore %231 %355 +%356 = OpLoad %4 %231 +%357 = OpLoad %4 %231 +%358 = OpExtInst %4 %1 UMax %356 %357 +%359 = OpLoad %4 %231 +%360 = OpIAdd %4 %359 %358 +OpStore %231 %360 +%361 = OpLoad %4 %231 +%362 = OpLoad %4 %231 +%363 = OpExtInst %4 %1 UMin %361 %362 +%364 = OpLoad %4 %231 +%365 = OpIAdd %4 %364 %363 +OpStore %231 %365 +%366 = OpLoad %4 %231 +OpReturnValue %366 OpFunctionEnd -%362 = OpFunction %2 None %363 -%361 = OpLabel -%364 = OpAccessChain %42 %23 %43 -%365 = OpAccessChain %45 %26 %43 -%366 = OpAccessChain %47 %29 %43 -%367 = OpAccessChain %45 %32 %43 -%368 = OpAccessChain %47 %35 %43 -OpBranch %371 -%371 = OpLabel -%372 = OpFunctionCall %4 %218 %369 -%373 = OpFunctionCall %3 %40 %370 -%374 = OpBitcast %4 %373 -%375 = OpIAdd %4 %372 %374 -%377 = OpAccessChain %292 %367 %376 -OpStore %377 %375 +%368 = OpFunction %2 None %369 +%367 = OpLabel +%370 = OpAccessChain %42 %23 %43 +%371 = OpAccessChain %45 %26 %43 +%372 = OpAccessChain %47 %29 %43 +%373 = OpAccessChain %45 %32 %43 +%374 = OpAccessChain %47 %35 %43 +OpBranch %377 +%377 = OpLabel +%378 = OpFunctionCall %4 %221 %375 +%379 = OpFunctionCall %3 %40 %376 +%380 = OpBitcast %4 %379 +%381 = OpIAdd %4 %378 %380 +%383 = OpAccessChain %298 %373 %382 +OpStore %383 %381 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/wgsl/conversion-float-to-int.wgsl b/naga/tests/out/wgsl/conversion-float-to-int.wgsl new file mode 100644 index 00000000000..a185b81dea1 --- /dev/null +++ b/naga/tests/out/wgsl/conversion-float-to-int.wgsl @@ -0,0 +1,142 @@ +enable f16; + +const MIN_F16_: f16 = -65504h; +const MAX_F16_: f16 = 65504h; +const MIN_F32_: f32 = -340282350000000000000000000000000000000f; +const MAX_F32_: f32 = 340282350000000000000000000000000000000f; +const MIN_F64_: f64 = -1.7976931348623157e308lf; +const MAX_F64_: f64 = 1.7976931348623157e308lf; + +fn test_const_eval() { + var min_f16_to_i32_: i32 = -65504i; + var max_f16_to_i32_: i32 = 65504i; + var min_f16_to_u32_: u32 = 0u; + var max_f16_to_u32_: u32 = 65504u; + var min_f16_to_i64_: i64 = -65504li; + var max_f16_to_i64_: i64 = 65504li; + var min_f16_to_u64_: u64 = 0lu; + var max_f16_to_u64_: u64 = 65504lu; + var min_f32_to_i32_: i32 = i32(-2147483648); + var max_f32_to_i32_: i32 = 2147483520i; + var min_f32_to_u32_: u32 = 0u; + var max_f32_to_u32_: u32 = 4294967040u; + var min_f32_to_i64_: i64 = -9223372036854775807li - 1li; + var max_f32_to_i64_: i64 = 9223371487098961920li; + var min_f32_to_u64_: u64 = 0lu; + var max_f32_to_u64_: u64 = 18446742974197923840lu; + var min_f64_to_i64_: i64 = -9223372036854775807li - 1li; + var max_f64_to_i64_: i64 = 9223372036854774784li; + var min_f64_to_u64_: u64 = 0lu; + var max_f64_to_u64_: u64 = 18446744073709549568lu; + var min_abstract_float_to_i32_: i32 = i32(-2147483648); + var max_abstract_float_to_i32_: i32 = 2147483647i; + var min_abstract_float_to_u32_: u32 = 0u; + var max_abstract_float_to_u32_: u32 = 4294967295u; + var min_abstract_float_to_i64_: i64 = -9223372036854775807li - 1li; + var max_abstract_float_to_i64_: i64 = 9223372036854774784li; + var min_abstract_float_to_u64_: u64 = 0lu; + var max_abstract_float_to_u64_: u64 = 18446744073709549568lu; + + return; +} + +fn test_f16_to_i32_(f: f16) -> i32 { + return i32(f); +} + +fn test_f16_to_u32_(f_1: f16) -> u32 { + return u32(f_1); +} + +fn test_f16_to_i64_(f_2: f16) -> i64 { + return i64(f_2); +} + +fn test_f16_to_u64_(f_3: f16) -> u64 { + return u64(f_3); +} + +fn test_f32_to_i32_(f_4: f32) -> i32 { + return i32(f_4); +} + +fn test_f32_to_u32_(f_5: f32) -> u32 { + return u32(f_5); +} + +fn test_f32_to_i64_(f_6: f32) -> i64 { + return i64(f_6); +} + +fn test_f32_to_u64_(f_7: f32) -> u64 { + return u64(f_7); +} + +fn test_f64_to_i32_(f_8: f64) -> i32 { + return i32(f_8); +} + +fn test_f64_to_u32_(f_9: f64) -> u32 { + return u32(f_9); +} + +fn test_f64_to_i64_(f_10: f64) -> i64 { + return i64(f_10); +} + +fn test_f64_to_u64_(f_11: f64) -> u64 { + return u64(f_11); +} + +fn test_f16_to_i32_vec(f_12: vec2) -> vec2 { + return vec2(f_12); +} + +fn test_f16_to_u32_vec(f_13: vec2) -> vec2 { + return vec2(f_13); +} + +fn test_f16_to_i64_vec(f_14: vec2) -> vec2 { + return vec2(f_14); +} + +fn test_f16_to_u64_vec(f_15: vec2) -> vec2 { + return vec2(f_15); +} + +fn test_f32_to_i32_vec(f_16: vec2) -> vec2 { + return vec2(f_16); +} + +fn test_f32_to_u32_vec(f_17: vec2) -> vec2 { + return vec2(f_17); +} + +fn test_f32_to_i64_vec(f_18: vec2) -> vec2 { + return vec2(f_18); +} + +fn test_f32_to_u64_vec(f_19: vec2) -> vec2 { + return vec2(f_19); +} + +fn test_f64_to_i32_vec(f_20: vec2) -> vec2 { + return vec2(f_20); +} + +fn test_f64_to_u32_vec(f_21: vec2) -> vec2 { + return vec2(f_21); +} + +fn test_f64_to_i64_vec(f_22: vec2) -> vec2 { + return vec2(f_22); +} + +fn test_f64_to_u64_vec(f_23: vec2) -> vec2 { + return vec2(f_23); +} + +@compute @workgroup_size(1, 1, 1) +fn main() { + return; +}