diff --git a/naga/src/back/spv/image.rs b/naga/src/back/spv/image.rs index a76d015f3f..46b1dd22df 100644 --- a/naga/src/back/spv/image.rs +++ b/naga/src/back/spv/image.rs @@ -924,7 +924,43 @@ impl<'w> BlockContext<'w> { depth_id, ); - let lod_id = self.cached[lod_handle]; + let mut lod_id = self.cached[lod_handle]; + // SPIR-V expects the LOD to be a float for all image classes. + // lod_id, however, will be an integer for depth images, + // therefore we must do a conversion. + if matches!( + self.ir_module.types[image_type].inner, + crate::TypeInner::Image { + class: crate::ImageClass::Depth { .. }, + .. + } + ) { + let lod_f32_id = self.gen_id(); + let f32_type_id = self.get_type_id(LookupType::Local(LocalType::Numeric( + NumericType::Scalar(crate::Scalar::F32), + ))); + let convert_op = match *self.fun_info[lod_handle] + .ty + .inner_with(&self.ir_module.types) + { + crate::TypeInner::Scalar(crate::Scalar { + kind: crate::ScalarKind::Uint, + width: 4, + }) => spirv::Op::ConvertUToF, + crate::TypeInner::Scalar(crate::Scalar { + kind: crate::ScalarKind::Sint, + width: 4, + }) => spirv::Op::ConvertSToF, + _ => unreachable!(), + }; + block.body.push(Instruction::unary( + convert_op, + f32_type_id, + lod_f32_id, + lod_id, + )); + lod_id = lod_f32_id; + } mask |= spirv::ImageOperands::LOD; inst.add_operand(mask.bits()); inst.add_operand(lod_id); diff --git a/naga/src/front/spv/image.rs b/naga/src/front/spv/image.rs index 71ba79e1e4..1d2b8dc645 100644 --- a/naga/src/front/spv/image.rs +++ b/naga/src/front/spv/image.rs @@ -30,6 +30,7 @@ impl<'function> super::BlockContext<'function> { match self.expressions[handle] { crate::Expression::GlobalVariable(handle) => Ok(self.global_arena[handle].ty), crate::Expression::FunctionArgument(i) => Ok(self.arguments[i as usize].ty), + crate::Expression::Access { base, .. } => Ok(self.get_image_expr_ty(base)?), ref other => Err(Error::InvalidImageExpression(other.clone())), } } @@ -460,6 +461,7 @@ impl> super::Frontend { } else { None }; + let span = self.span_from_with_op(start); let mut image_ops = if words_left != 0 { words_left -= 1; @@ -486,9 +488,34 @@ impl> super::Frontend { let lod_lexp = self.lookup_expression.lookup(lod_expr)?; let lod_handle = self.get_expr_handle(lod_expr, lod_lexp, ctx, emitter, block, body_idx); + + let is_depth_image = { + let image_lexp = self.lookup_sampled_image.lookup(sampled_image_id)?; + let image_ty = ctx.get_image_expr_ty(image_lexp.image)?; + matches!( + ctx.type_arena[image_ty].inner, + crate::TypeInner::Image { + class: crate::ImageClass::Depth { .. }, + .. + } + ) + }; + level = if options.compare { log::debug!("Assuming {:?} is zero", lod_handle); crate::SampleLevel::Zero + } else if is_depth_image { + log::debug!( + "Assuming level {:?} converts losslessly to an integer", + lod_handle + ); + let expr = crate::Expression::As { + expr: lod_handle, + kind: crate::ScalarKind::Sint, + convert: Some(4), + }; + let s32_lod_handle = ctx.expressions.append(expr, span); + crate::SampleLevel::Exact(s32_lod_handle) } else { crate::SampleLevel::Exact(lod_handle) }; diff --git a/naga/src/valid/expression.rs b/naga/src/valid/expression.rs index 9a15034012..a32b134bb6 100644 --- a/naga/src/valid/expression.rs +++ b/naga/src/valid/expression.rs @@ -106,7 +106,7 @@ pub enum ExpressionError { InvalidGatherComponent(crate::SwizzleComponent), #[error("Gather can't be done for image dimension {0:?}")] InvalidGatherDimension(crate::ImageDimension), - #[error("Sample level (exact) type {0:?} is not a scalar float")] + #[error("Sample level (exact) type {0:?} has an invalid type")] InvalidSampleLevelExactType(Handle), #[error("Sample level (bias) type {0:?} is not a scalar float")] InvalidSampleLevelBiasType(Handle), @@ -530,11 +530,24 @@ impl super::Validator { crate::SampleLevel::Auto => ShaderStages::FRAGMENT, crate::SampleLevel::Zero => ShaderStages::all(), crate::SampleLevel::Exact(expr) => { - match resolver[expr] { - Ti::Scalar(Sc { - kind: Sk::Float, .. - }) => {} - _ => return Err(ExpressionError::InvalidSampleLevelExactType(expr)), + match class { + crate::ImageClass::Depth { .. } => match resolver[expr] { + Ti::Scalar(Sc { + kind: Sk::Sint | Sk::Uint, + .. + }) => {} + _ => { + return Err(ExpressionError::InvalidSampleLevelExactType(expr)) + } + }, + _ => match resolver[expr] { + Ti::Scalar(Sc { + kind: Sk::Float, .. + }) => {} + _ => { + return Err(ExpressionError::InvalidSampleLevelExactType(expr)) + } + }, } ShaderStages::all() } diff --git a/naga/tests/in/image.wgsl b/naga/tests/in/image.wgsl index e784801182..50eabcb6f9 100644 --- a/naga/tests/in/image.wgsl +++ b/naga/tests/in/image.wgsl @@ -185,7 +185,9 @@ fn gather() -> @location(0) vec4 { @fragment fn depth_no_comparison() -> @location(0) vec4 { let tc = vec2(0.5); + let level = 1; let s2d = textureSample(image_2d_depth, sampler_reg, tc); let s2d_gather = textureGather(image_2d_depth, sampler_reg, tc); - return s2d + s2d_gather; + let s2d_level = textureSampleLevel(image_2d_depth, sampler_reg, tc, level); + return s2d + s2d_gather + s2d_level; } diff --git a/naga/tests/out/hlsl/image.hlsl b/naga/tests/out/hlsl/image.hlsl index a81625b058..84e9f6d706 100644 --- a/naga/tests/out/hlsl/image.hlsl +++ b/naga/tests/out/hlsl/image.hlsl @@ -369,5 +369,6 @@ float4 depth_no_comparison() : SV_Target0 float2 tc_3 = (0.5).xx; float s2d_1 = image_2d_depth.Sample(sampler_reg, tc_3); float4 s2d_gather = image_2d_depth.Gather(sampler_reg, tc_3); - return ((s2d_1).xxxx + s2d_gather); + float s2d_level = image_2d_depth.SampleLevel(sampler_reg, tc_3, 1); + return (((s2d_1).xxxx + s2d_gather) + (s2d_level).xxxx); } diff --git a/naga/tests/out/msl/image.msl b/naga/tests/out/msl/image.msl index 9f0a24300a..114ed36553 100644 --- a/naga/tests/out/msl/image.msl +++ b/naga/tests/out/msl/image.msl @@ -265,5 +265,6 @@ fragment depth_no_comparisonOutput depth_no_comparison( metal::float2 tc_3 = metal::float2(0.5); float s2d_1 = image_2d_depth.sample(sampler_reg, tc_3); metal::float4 s2d_gather = image_2d_depth.gather(sampler_reg, tc_3); - return depth_no_comparisonOutput { metal::float4(s2d_1) + s2d_gather }; + float s2d_level = image_2d_depth.sample(sampler_reg, tc_3, metal::level(1)); + return depth_no_comparisonOutput { (metal::float4(s2d_1) + s2d_gather) + metal::float4(s2d_level) }; } diff --git a/naga/tests/out/spv/image.spvasm b/naga/tests/out/spv/image.spvasm index 974f29e166..5ec2925dc6 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: 520 +; Bound: 526 OpCapability Shader OpCapability Image1D OpCapability Sampled1D @@ -687,8 +687,14 @@ OpBranch %512 %515 = OpCompositeExtract %7 %514 0 %516 = OpSampledImage %428 %511 %510 %517 = OpImageGather %23 %516 %280 %198 -%518 = OpCompositeConstruct %23 %515 %515 %515 %515 -%519 = OpFAdd %23 %518 %517 -OpStore %508 %519 +%518 = OpSampledImage %428 %511 %510 +%520 = OpConvertSToF %7 %29 +%519 = OpImageSampleExplicitLod %23 %518 %280 Lod %520 +%521 = OpCompositeExtract %7 %519 0 +%522 = OpCompositeConstruct %23 %515 %515 %515 %515 +%523 = OpFAdd %23 %522 %517 +%524 = OpCompositeConstruct %23 %521 %521 %521 %521 +%525 = OpFAdd %23 %523 %524 +OpStore %508 %525 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/wgsl/image.wgsl b/naga/tests/out/wgsl/image.wgsl index dddc794045..0c10eda857 100644 --- a/naga/tests/out/wgsl/image.wgsl +++ b/naga/tests/out/wgsl/image.wgsl @@ -235,5 +235,6 @@ fn depth_no_comparison() -> @location(0) vec4 { const tc_3 = vec2(0.5f); let s2d_1 = textureSample(image_2d_depth, sampler_reg, tc_3); let s2d_gather = textureGather(image_2d_depth, sampler_reg, tc_3); - return (vec4(s2d_1) + s2d_gather); + let s2d_level = textureSampleLevel(image_2d_depth, sampler_reg, tc_3, 1i); + return ((vec4(s2d_1) + s2d_gather) + vec4(s2d_level)); }