From 1ac8dae8aaedc513734e49ba1e4e8b32e00e3209 Mon Sep 17 00:00:00 2001 From: Connor Fitzgerald Date: Tue, 19 Apr 2022 19:06:18 -0400 Subject: [PATCH 1/3] Binding arrays play nice with bounds checks --- src/back/msl/writer.rs | 5 +- src/proc/index.rs | 12 +- tests/in/binding-arrays.param.ron | 5 + tests/out/msl/binding-arrays.msl | 24 +- tests/out/spv/binding-arrays.spvasm | 542 ++++++++++++++++------------ tests/snapshots.rs | 4 +- 6 files changed, 344 insertions(+), 248 deletions(-) diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index e365f70d4e..55dc198fdb 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -1270,13 +1270,14 @@ impl Writer { let expression = &context.function.expressions[expr_handle]; log::trace!("expression {:?} = {:?}", expr_handle, expression); match *expression { - crate::Expression::Access { .. } | crate::Expression::AccessIndex { .. } => { + crate::Expression::Access { base, .. } + | crate::Expression::AccessIndex { base, .. } => { // This is an acceptable place to generate a `ReadZeroSkipWrite` check. // Since `put_bounds_checks` and `put_access_chain` handle an entire // access chain at a time, recursing back through `put_expression` only // for index expressions and the base object, we will never see intermediate // `Access` or `AccessIndex` expressions here. - let policy = context.choose_bounds_check_policy(expr_handle); + let policy = context.choose_bounds_check_policy(base); if policy == index::BoundsCheckPolicy::ReadZeroSkipWrite && self.put_bounds_checks( expr_handle, diff --git a/src/proc/index.rs b/src/proc/index.rs index ead07052b5..60204fd3c2 100644 --- a/src/proc/index.rs +++ b/src/proc/index.rs @@ -117,6 +117,10 @@ pub struct BoundsCheckPolicies { /// [`ImageStore`]: crate::Statement::ImageStore #[cfg_attr(feature = "deserialize", serde(default))] pub image: BoundsCheckPolicy, + + /// How should the generated code handle binding array indexes that are out of bounds. + #[cfg_attr(feature = "deserialize", serde(default))] + pub binding_array: BoundsCheckPolicy, } /// The default `BoundsCheckPolicy` is `Unchecked`. @@ -140,7 +144,13 @@ impl BoundsCheckPolicies { types: &UniqueArena, info: &valid::FunctionInfo, ) -> BoundsCheckPolicy { - match info[access].ty.inner_with(types).pointer_space() { + let ty = info[access].ty.inner_with(types); + + if let crate::TypeInner::BindingArray { .. } = *ty { + return self.binding_array; + } + + match ty.pointer_space() { Some(crate::AddressSpace::Storage { access: _ } | crate::AddressSpace::Uniform) => { self.buffer } diff --git a/tests/in/binding-arrays.param.ron b/tests/in/binding-arrays.param.ron index b2f1dabaa2..2d3e15263b 100644 --- a/tests/in/binding-arrays.param.ron +++ b/tests/in/binding-arrays.param.ron @@ -35,5 +35,10 @@ binding_map: { (group: 0, binding: 0): (binding_array_size: Some(10)), }, + ), + bounds_check_policies: ( + index: ReadZeroSkipWrite, + buffer: ReadZeroSkipWrite, + image: ReadZeroSkipWrite, ) ) diff --git a/tests/out/msl/binding-arrays.msl b/tests/out/msl/binding-arrays.msl index 7291e77f24..62e0bb6b16 100644 --- a/tests/out/msl/binding-arrays.msl +++ b/tests/out/msl/binding-arrays.msl @@ -4,6 +4,12 @@ using metal::uint; +struct DefaultConstructible { + template + operator T() && { + return T {}; + } +}; struct UniformIndex { uint index; }; @@ -65,13 +71,13 @@ fragment main_Output main_( metal::float4 _e75 = texture_array_depth[non_uniform_index].gather_compare(samp_comp[non_uniform_index], uv, 0.0); v4_ = _e71 + _e75; metal::float4 _e77 = v4_; - metal::float4 _e81 = texture_array_unbounded[0].read(metal::uint2(pix), 0); + metal::float4 _e81 = (uint(0) < texture_array_unbounded[0].get_num_mip_levels() && metal::all(metal::uint2(pix) < metal::uint2(texture_array_unbounded[0].get_width(0), texture_array_unbounded[0].get_height(0))) ? texture_array_unbounded[0].read(metal::uint2(pix), 0): DefaultConstructible()); v4_ = _e77 + _e81; metal::float4 _e83 = v4_; - metal::float4 _e86 = texture_array_unbounded[uniform_index].read(metal::uint2(pix), 0); + metal::float4 _e86 = (uint(0) < texture_array_unbounded[uniform_index].get_num_mip_levels() && metal::all(metal::uint2(pix) < metal::uint2(texture_array_unbounded[uniform_index].get_width(0), texture_array_unbounded[uniform_index].get_height(0))) ? texture_array_unbounded[uniform_index].read(metal::uint2(pix), 0): DefaultConstructible()); v4_ = _e83 + _e86; metal::float4 _e88 = v4_; - metal::float4 _e91 = texture_array_unbounded[non_uniform_index].read(metal::uint2(pix), 0); + metal::float4 _e91 = (uint(0) < texture_array_unbounded[non_uniform_index].get_num_mip_levels() && metal::all(metal::uint2(pix) < metal::uint2(texture_array_unbounded[non_uniform_index].get_width(0), texture_array_unbounded[non_uniform_index].get_height(0))) ? texture_array_unbounded[non_uniform_index].read(metal::uint2(pix), 0): DefaultConstructible()); v4_ = _e88 + _e91; int _e93 = i1_; i1_ = _e93 + int(texture_array_2darray[0].get_array_size()); @@ -146,11 +152,17 @@ fragment main_Output main_( metal::float4 _e244 = texture_array_bounded[non_uniform_index].sample(samp[non_uniform_index], uv, metal::level(0.0)); v4_ = _e240 + _e244; metal::float4 _e248 = v4_; - texture_array_storage[0].write(_e248, metal::uint2(pix)); + if (metal::all(metal::uint2(pix) < metal::uint2(texture_array_storage[0].get_width(), texture_array_storage[0].get_height()))) { + texture_array_storage[0].write(_e248, metal::uint2(pix)); + } metal::float4 _e250 = v4_; - texture_array_storage[uniform_index].write(_e250, metal::uint2(pix)); + if (metal::all(metal::uint2(pix) < metal::uint2(texture_array_storage[uniform_index].get_width(), texture_array_storage[uniform_index].get_height()))) { + texture_array_storage[uniform_index].write(_e250, metal::uint2(pix)); + } metal::float4 _e252 = v4_; - texture_array_storage[non_uniform_index].write(_e252, metal::uint2(pix)); + if (metal::all(metal::uint2(pix) < metal::uint2(texture_array_storage[non_uniform_index].get_width(), texture_array_storage[non_uniform_index].get_height()))) { + texture_array_storage[non_uniform_index].write(_e252, metal::uint2(pix)); + } metal::int2 _e253 = i2_; int _e254 = i1_; metal::float2 v2_ = static_cast(_e253 + metal::int2(_e254)); diff --git a/tests/out/spv/binding-arrays.spvasm b/tests/out/spv/binding-arrays.spvasm index 71f13c3d99..f68a755cc6 100644 --- a/tests/out/spv/binding-arrays.spvasm +++ b/tests/out/spv/binding-arrays.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 382 +; Bound: 429 OpCapability Shader OpCapability ImageQuery OpCapability ShaderNonUniform @@ -40,23 +40,23 @@ OpDecorate %114 NonUniform OpDecorate %116 NonUniform OpDecorate %141 NonUniform OpDecorate %143 NonUniform -OpDecorate %159 NonUniform -OpDecorate %178 NonUniform -OpDecorate %194 NonUniform +OpDecorate %181 NonUniform OpDecorate %210 NonUniform -OpDecorate %231 NonUniform -OpDecorate %233 NonUniform -OpDecorate %255 NonUniform -OpDecorate %257 NonUniform -OpDecorate %279 NonUniform -OpDecorate %281 NonUniform -OpDecorate %303 NonUniform -OpDecorate %305 NonUniform -OpDecorate %327 NonUniform -OpDecorate %329 NonUniform -OpDecorate %351 NonUniform -OpDecorate %353 NonUniform -OpDecorate %365 NonUniform +OpDecorate %226 NonUniform +OpDecorate %242 NonUniform +OpDecorate %263 NonUniform +OpDecorate %265 NonUniform +OpDecorate %287 NonUniform +OpDecorate %289 NonUniform +OpDecorate %311 NonUniform +OpDecorate %313 NonUniform +OpDecorate %335 NonUniform +OpDecorate %337 NonUniform +OpDecorate %359 NonUniform +OpDecorate %361 NonUniform +OpDecorate %383 NonUniform +OpDecorate %385 NonUniform +OpDecorate %407 NonUniform %2 = OpTypeVoid %4 = OpTypeInt 32 1 %3 = OpConstant %4 5 @@ -122,10 +122,15 @@ OpDecorate %365 NonUniform %121 = OpTypePointer UniformConstant %17 %124 = OpTypePointer UniformConstant %21 %127 = OpTypeSampledImage %17 -%163 = OpTypePointer UniformConstant %13 -%166 = OpTypeVector %4 3 -%198 = OpTypePointer UniformConstant %15 -%357 = OpTypePointer UniformConstant %19 +%150 = OpTypeBool +%151 = OpConstantNull %25 +%157 = OpTypeVector %150 2 +%167 = OpConstantNull %25 +%182 = OpConstantNull %25 +%195 = OpTypePointer UniformConstant %13 +%198 = OpTypeVector %4 3 +%230 = OpTypePointer UniformConstant %15 +%389 = OpTypePointer UniformConstant %19 %63 = OpFunction %2 None %64 %56 = OpLabel %50 = OpVariable %51 Function @@ -221,267 +226,330 @@ OpStore %54 %146 %147 = OpLoad %25 %54 %148 = OpAccessChain %79 %27 %66 %149 = OpLoad %10 %148 -%150 = OpImageFetch %25 %149 %77 Lod %5 -%151 = OpFAdd %25 %147 %150 -OpStore %54 %151 -%152 = OpLoad %25 %54 -%153 = OpAccessChain %79 %27 %71 -%154 = OpLoad %10 %153 -%155 = OpImageFetch %25 %154 %77 Lod %5 -%156 = OpFAdd %25 %152 %155 -OpStore %54 %156 -%157 = OpLoad %25 %54 -%158 = OpAccessChain %79 %27 %72 -%159 = OpLoad %10 %158 -%160 = OpImageFetch %25 %159 %77 Lod %5 -%161 = OpFAdd %25 %157 %160 -OpStore %54 %161 -%162 = OpLoad %4 %48 -%164 = OpAccessChain %163 %33 %66 -%165 = OpLoad %13 %164 -%167 = OpImageQuerySizeLod %166 %165 %66 -%168 = OpCompositeExtract %4 %167 2 -%169 = OpIAdd %4 %162 %168 -OpStore %48 %169 -%170 = OpLoad %4 %48 -%171 = OpAccessChain %163 %33 %71 -%172 = OpLoad %13 %171 -%173 = OpImageQuerySizeLod %166 %172 %66 -%174 = OpCompositeExtract %4 %173 2 -%175 = OpIAdd %4 %170 %174 -OpStore %48 %175 -%176 = OpLoad %4 %48 -%177 = OpAccessChain %163 %33 %72 -%178 = OpLoad %13 %177 -%179 = OpImageQuerySizeLod %166 %178 %66 -%180 = OpCompositeExtract %4 %179 2 -%181 = OpIAdd %4 %176 %180 -OpStore %48 %181 -%182 = OpLoad %4 %48 -%183 = OpAccessChain %79 %31 %66 -%184 = OpLoad %10 %183 -%185 = OpImageQueryLevels %4 %184 -%186 = OpIAdd %4 %182 %185 -OpStore %48 %186 -%187 = OpLoad %4 %48 -%188 = OpAccessChain %79 %31 %71 -%189 = OpLoad %10 %188 -%190 = OpImageQueryLevels %4 %189 -%191 = OpIAdd %4 %187 %190 -OpStore %48 %191 -%192 = OpLoad %4 %48 -%193 = OpAccessChain %79 %31 %72 -%194 = OpLoad %10 %193 -%195 = OpImageQueryLevels %4 %194 -%196 = OpIAdd %4 %192 %195 -OpStore %48 %196 -%197 = OpLoad %4 %48 -%199 = OpAccessChain %198 %35 %66 -%200 = OpLoad %15 %199 -%201 = OpImageQuerySamples %4 %200 -%202 = OpIAdd %4 %197 %201 -OpStore %48 %202 -%203 = OpLoad %4 %48 -%204 = OpAccessChain %198 %35 %71 -%205 = OpLoad %15 %204 -%206 = OpImageQuerySamples %4 %205 -%207 = OpIAdd %4 %203 %206 +%152 = OpImageQueryLevels %4 %149 +%153 = OpULessThan %150 %5 %152 +OpSelectionMerge %154 None +OpBranchConditional %153 %155 %154 +%155 = OpLabel +%156 = OpImageQuerySizeLod %26 %149 %5 +%158 = OpULessThan %157 %77 %156 +%159 = OpAll %150 %158 +OpBranchConditional %159 %160 %154 +%160 = OpLabel +%161 = OpImageFetch %25 %149 %77 Lod %5 +OpBranch %154 +%154 = OpLabel +%162 = OpPhi %25 %151 %68 %151 %155 %161 %160 +%163 = OpFAdd %25 %147 %162 +OpStore %54 %163 +%164 = OpLoad %25 %54 +%165 = OpAccessChain %79 %27 %71 +%166 = OpLoad %10 %165 +%168 = OpImageQueryLevels %4 %166 +%169 = OpULessThan %150 %5 %168 +OpSelectionMerge %170 None +OpBranchConditional %169 %171 %170 +%171 = OpLabel +%172 = OpImageQuerySizeLod %26 %166 %5 +%173 = OpULessThan %157 %77 %172 +%174 = OpAll %150 %173 +OpBranchConditional %174 %175 %170 +%175 = OpLabel +%176 = OpImageFetch %25 %166 %77 Lod %5 +OpBranch %170 +%170 = OpLabel +%177 = OpPhi %25 %167 %154 %167 %171 %176 %175 +%178 = OpFAdd %25 %164 %177 +OpStore %54 %178 +%179 = OpLoad %25 %54 +%180 = OpAccessChain %79 %27 %72 +%181 = OpLoad %10 %180 +%183 = OpImageQueryLevels %4 %181 +%184 = OpULessThan %150 %5 %183 +OpSelectionMerge %185 None +OpBranchConditional %184 %186 %185 +%186 = OpLabel +%187 = OpImageQuerySizeLod %26 %181 %5 +%188 = OpULessThan %157 %77 %187 +%189 = OpAll %150 %188 +OpBranchConditional %189 %190 %185 +%190 = OpLabel +%191 = OpImageFetch %25 %181 %77 Lod %5 +OpBranch %185 +%185 = OpLabel +%192 = OpPhi %25 %182 %170 %182 %186 %191 %190 +%193 = OpFAdd %25 %179 %192 +OpStore %54 %193 +%194 = OpLoad %4 %48 +%196 = OpAccessChain %195 %33 %66 +%197 = OpLoad %13 %196 +%199 = OpImageQuerySizeLod %198 %197 %66 +%200 = OpCompositeExtract %4 %199 2 +%201 = OpIAdd %4 %194 %200 +OpStore %48 %201 +%202 = OpLoad %4 %48 +%203 = OpAccessChain %195 %33 %71 +%204 = OpLoad %13 %203 +%205 = OpImageQuerySizeLod %198 %204 %66 +%206 = OpCompositeExtract %4 %205 2 +%207 = OpIAdd %4 %202 %206 OpStore %48 %207 %208 = OpLoad %4 %48 -%209 = OpAccessChain %198 %35 %72 -%210 = OpLoad %15 %209 -%211 = OpImageQuerySamples %4 %210 -%212 = OpIAdd %4 %208 %211 -OpStore %48 %212 -%213 = OpLoad %25 %54 -%214 = OpAccessChain %79 %31 %66 -%215 = OpLoad %10 %214 -%216 = OpAccessChain %97 %41 %66 -%217 = OpLoad %21 %216 -%218 = OpSampledImage %100 %215 %217 -%219 = OpImageSampleImplicitLod %25 %218 %76 -%220 = OpFAdd %25 %213 %219 -OpStore %54 %220 -%221 = OpLoad %25 %54 -%222 = OpAccessChain %79 %31 %71 -%223 = OpLoad %10 %222 -%224 = OpAccessChain %97 %41 %71 -%225 = OpLoad %21 %224 -%226 = OpSampledImage %100 %223 %225 -%227 = OpImageSampleImplicitLod %25 %226 %76 -%228 = OpFAdd %25 %221 %227 -OpStore %54 %228 -%229 = OpLoad %25 %54 -%230 = OpAccessChain %79 %31 %72 -%231 = OpLoad %10 %230 -%232 = OpAccessChain %97 %41 %72 -%233 = OpLoad %21 %232 -%234 = OpSampledImage %100 %231 %233 -%235 = OpImageSampleImplicitLod %25 %234 %76 -%236 = OpFAdd %25 %229 %235 -OpStore %54 %236 -%237 = OpLoad %25 %54 -%238 = OpAccessChain %79 %31 %66 -%239 = OpLoad %10 %238 -%240 = OpAccessChain %97 %41 %66 -%241 = OpLoad %21 %240 -%242 = OpSampledImage %100 %239 %241 -%243 = OpImageSampleImplicitLod %25 %242 %76 Bias %6 -%244 = OpFAdd %25 %237 %243 -OpStore %54 %244 +%209 = OpAccessChain %195 %33 %72 +%210 = OpLoad %13 %209 +%211 = OpImageQuerySizeLod %198 %210 %66 +%212 = OpCompositeExtract %4 %211 2 +%213 = OpIAdd %4 %208 %212 +OpStore %48 %213 +%214 = OpLoad %4 %48 +%215 = OpAccessChain %79 %31 %66 +%216 = OpLoad %10 %215 +%217 = OpImageQueryLevels %4 %216 +%218 = OpIAdd %4 %214 %217 +OpStore %48 %218 +%219 = OpLoad %4 %48 +%220 = OpAccessChain %79 %31 %71 +%221 = OpLoad %10 %220 +%222 = OpImageQueryLevels %4 %221 +%223 = OpIAdd %4 %219 %222 +OpStore %48 %223 +%224 = OpLoad %4 %48 +%225 = OpAccessChain %79 %31 %72 +%226 = OpLoad %10 %225 +%227 = OpImageQueryLevels %4 %226 +%228 = OpIAdd %4 %224 %227 +OpStore %48 %228 +%229 = OpLoad %4 %48 +%231 = OpAccessChain %230 %35 %66 +%232 = OpLoad %15 %231 +%233 = OpImageQuerySamples %4 %232 +%234 = OpIAdd %4 %229 %233 +OpStore %48 %234 +%235 = OpLoad %4 %48 +%236 = OpAccessChain %230 %35 %71 +%237 = OpLoad %15 %236 +%238 = OpImageQuerySamples %4 %237 +%239 = OpIAdd %4 %235 %238 +OpStore %48 %239 +%240 = OpLoad %4 %48 +%241 = OpAccessChain %230 %35 %72 +%242 = OpLoad %15 %241 +%243 = OpImageQuerySamples %4 %242 +%244 = OpIAdd %4 %240 %243 +OpStore %48 %244 %245 = OpLoad %25 %54 -%246 = OpAccessChain %79 %31 %71 +%246 = OpAccessChain %79 %31 %66 %247 = OpLoad %10 %246 -%248 = OpAccessChain %97 %41 %71 +%248 = OpAccessChain %97 %41 %66 %249 = OpLoad %21 %248 %250 = OpSampledImage %100 %247 %249 -%251 = OpImageSampleImplicitLod %25 %250 %76 Bias %6 +%251 = OpImageSampleImplicitLod %25 %250 %76 %252 = OpFAdd %25 %245 %251 OpStore %54 %252 %253 = OpLoad %25 %54 -%254 = OpAccessChain %79 %31 %72 +%254 = OpAccessChain %79 %31 %71 %255 = OpLoad %10 %254 -%256 = OpAccessChain %97 %41 %72 +%256 = OpAccessChain %97 %41 %71 %257 = OpLoad %21 %256 %258 = OpSampledImage %100 %255 %257 -%259 = OpImageSampleImplicitLod %25 %258 %76 Bias %6 +%259 = OpImageSampleImplicitLod %25 %258 %76 %260 = OpFAdd %25 %253 %259 OpStore %54 %260 -%261 = OpLoad %7 %52 -%262 = OpAccessChain %121 %37 %66 -%263 = OpLoad %17 %262 -%264 = OpAccessChain %124 %43 %66 +%261 = OpLoad %25 %54 +%262 = OpAccessChain %79 %31 %72 +%263 = OpLoad %10 %262 +%264 = OpAccessChain %97 %41 %72 %265 = OpLoad %21 %264 -%266 = OpSampledImage %127 %263 %265 -%267 = OpImageSampleDrefImplicitLod %7 %266 %76 %6 -%268 = OpFAdd %7 %261 %267 -OpStore %52 %268 -%269 = OpLoad %7 %52 -%270 = OpAccessChain %121 %37 %71 -%271 = OpLoad %17 %270 -%272 = OpAccessChain %124 %43 %71 +%266 = OpSampledImage %100 %263 %265 +%267 = OpImageSampleImplicitLod %25 %266 %76 +%268 = OpFAdd %25 %261 %267 +OpStore %54 %268 +%269 = OpLoad %25 %54 +%270 = OpAccessChain %79 %31 %66 +%271 = OpLoad %10 %270 +%272 = OpAccessChain %97 %41 %66 %273 = OpLoad %21 %272 -%274 = OpSampledImage %127 %271 %273 -%275 = OpImageSampleDrefImplicitLod %7 %274 %76 %6 -%276 = OpFAdd %7 %269 %275 -OpStore %52 %276 -%277 = OpLoad %7 %52 -%278 = OpAccessChain %121 %37 %72 -%279 = OpLoad %17 %278 -%280 = OpAccessChain %124 %43 %72 +%274 = OpSampledImage %100 %271 %273 +%275 = OpImageSampleImplicitLod %25 %274 %76 Bias %6 +%276 = OpFAdd %25 %269 %275 +OpStore %54 %276 +%277 = OpLoad %25 %54 +%278 = OpAccessChain %79 %31 %71 +%279 = OpLoad %10 %278 +%280 = OpAccessChain %97 %41 %71 %281 = OpLoad %21 %280 -%282 = OpSampledImage %127 %279 %281 -%283 = OpImageSampleDrefImplicitLod %7 %282 %76 %6 -%284 = OpFAdd %7 %277 %283 -OpStore %52 %284 -%285 = OpLoad %7 %52 -%286 = OpAccessChain %121 %37 %66 -%287 = OpLoad %17 %286 -%288 = OpAccessChain %124 %43 %66 +%282 = OpSampledImage %100 %279 %281 +%283 = OpImageSampleImplicitLod %25 %282 %76 Bias %6 +%284 = OpFAdd %25 %277 %283 +OpStore %54 %284 +%285 = OpLoad %25 %54 +%286 = OpAccessChain %79 %31 %72 +%287 = OpLoad %10 %286 +%288 = OpAccessChain %97 %41 %72 %289 = OpLoad %21 %288 -%290 = OpSampledImage %127 %287 %289 -%291 = OpImageSampleDrefExplicitLod %7 %290 %76 %6 Lod %6 -%292 = OpFAdd %7 %285 %291 -OpStore %52 %292 +%290 = OpSampledImage %100 %287 %289 +%291 = OpImageSampleImplicitLod %25 %290 %76 Bias %6 +%292 = OpFAdd %25 %285 %291 +OpStore %54 %292 %293 = OpLoad %7 %52 -%294 = OpAccessChain %121 %37 %71 +%294 = OpAccessChain %121 %37 %66 %295 = OpLoad %17 %294 -%296 = OpAccessChain %124 %43 %71 +%296 = OpAccessChain %124 %43 %66 %297 = OpLoad %21 %296 %298 = OpSampledImage %127 %295 %297 -%299 = OpImageSampleDrefExplicitLod %7 %298 %76 %6 Lod %6 +%299 = OpImageSampleDrefImplicitLod %7 %298 %76 %6 %300 = OpFAdd %7 %293 %299 OpStore %52 %300 %301 = OpLoad %7 %52 -%302 = OpAccessChain %121 %37 %72 +%302 = OpAccessChain %121 %37 %71 %303 = OpLoad %17 %302 -%304 = OpAccessChain %124 %43 %72 +%304 = OpAccessChain %124 %43 %71 %305 = OpLoad %21 %304 %306 = OpSampledImage %127 %303 %305 -%307 = OpImageSampleDrefExplicitLod %7 %306 %76 %6 Lod %6 +%307 = OpImageSampleDrefImplicitLod %7 %306 %76 %6 %308 = OpFAdd %7 %301 %307 OpStore %52 %308 -%309 = OpLoad %25 %54 -%310 = OpAccessChain %79 %31 %66 -%311 = OpLoad %10 %310 -%312 = OpAccessChain %97 %41 %66 +%309 = OpLoad %7 %52 +%310 = OpAccessChain %121 %37 %72 +%311 = OpLoad %17 %310 +%312 = OpAccessChain %124 %43 %72 %313 = OpLoad %21 %312 -%314 = OpSampledImage %100 %311 %313 -%315 = OpImageSampleExplicitLod %25 %314 %76 Grad %76 %76 -%316 = OpFAdd %25 %309 %315 -OpStore %54 %316 -%317 = OpLoad %25 %54 -%318 = OpAccessChain %79 %31 %71 -%319 = OpLoad %10 %318 -%320 = OpAccessChain %97 %41 %71 +%314 = OpSampledImage %127 %311 %313 +%315 = OpImageSampleDrefImplicitLod %7 %314 %76 %6 +%316 = OpFAdd %7 %309 %315 +OpStore %52 %316 +%317 = OpLoad %7 %52 +%318 = OpAccessChain %121 %37 %66 +%319 = OpLoad %17 %318 +%320 = OpAccessChain %124 %43 %66 %321 = OpLoad %21 %320 -%322 = OpSampledImage %100 %319 %321 -%323 = OpImageSampleExplicitLod %25 %322 %76 Grad %76 %76 -%324 = OpFAdd %25 %317 %323 -OpStore %54 %324 -%325 = OpLoad %25 %54 -%326 = OpAccessChain %79 %31 %72 -%327 = OpLoad %10 %326 -%328 = OpAccessChain %97 %41 %72 +%322 = OpSampledImage %127 %319 %321 +%323 = OpImageSampleDrefExplicitLod %7 %322 %76 %6 Lod %6 +%324 = OpFAdd %7 %317 %323 +OpStore %52 %324 +%325 = OpLoad %7 %52 +%326 = OpAccessChain %121 %37 %71 +%327 = OpLoad %17 %326 +%328 = OpAccessChain %124 %43 %71 %329 = OpLoad %21 %328 -%330 = OpSampledImage %100 %327 %329 -%331 = OpImageSampleExplicitLod %25 %330 %76 Grad %76 %76 -%332 = OpFAdd %25 %325 %331 -OpStore %54 %332 -%333 = OpLoad %25 %54 -%334 = OpAccessChain %79 %31 %66 -%335 = OpLoad %10 %334 -%336 = OpAccessChain %97 %41 %66 +%330 = OpSampledImage %127 %327 %329 +%331 = OpImageSampleDrefExplicitLod %7 %330 %76 %6 Lod %6 +%332 = OpFAdd %7 %325 %331 +OpStore %52 %332 +%333 = OpLoad %7 %52 +%334 = OpAccessChain %121 %37 %72 +%335 = OpLoad %17 %334 +%336 = OpAccessChain %124 %43 %72 %337 = OpLoad %21 %336 -%338 = OpSampledImage %100 %335 %337 -%339 = OpImageSampleExplicitLod %25 %338 %76 Lod %6 -%340 = OpFAdd %25 %333 %339 -OpStore %54 %340 +%338 = OpSampledImage %127 %335 %337 +%339 = OpImageSampleDrefExplicitLod %7 %338 %76 %6 Lod %6 +%340 = OpFAdd %7 %333 %339 +OpStore %52 %340 %341 = OpLoad %25 %54 -%342 = OpAccessChain %79 %31 %71 +%342 = OpAccessChain %79 %31 %66 %343 = OpLoad %10 %342 -%344 = OpAccessChain %97 %41 %71 +%344 = OpAccessChain %97 %41 %66 %345 = OpLoad %21 %344 %346 = OpSampledImage %100 %343 %345 -%347 = OpImageSampleExplicitLod %25 %346 %76 Lod %6 +%347 = OpImageSampleExplicitLod %25 %346 %76 Grad %76 %76 %348 = OpFAdd %25 %341 %347 OpStore %54 %348 %349 = OpLoad %25 %54 -%350 = OpAccessChain %79 %31 %72 +%350 = OpAccessChain %79 %31 %71 %351 = OpLoad %10 %350 -%352 = OpAccessChain %97 %41 %72 +%352 = OpAccessChain %97 %41 %71 %353 = OpLoad %21 %352 %354 = OpSampledImage %100 %351 %353 -%355 = OpImageSampleExplicitLod %25 %354 %76 Lod %6 +%355 = OpImageSampleExplicitLod %25 %354 %76 Grad %76 %76 %356 = OpFAdd %25 %349 %355 OpStore %54 %356 -%358 = OpAccessChain %357 %39 %66 -%359 = OpLoad %19 %358 -%360 = OpLoad %25 %54 -OpImageWrite %359 %77 %360 -%361 = OpAccessChain %357 %39 %71 -%362 = OpLoad %19 %361 -%363 = OpLoad %25 %54 -OpImageWrite %362 %77 %363 -%364 = OpAccessChain %357 %39 %72 -%365 = OpLoad %19 %364 -%366 = OpLoad %25 %54 -OpImageWrite %365 %77 %366 -%367 = OpLoad %26 %50 -%368 = OpLoad %4 %48 -%369 = OpCompositeConstruct %26 %368 %368 -%370 = OpIAdd %26 %367 %369 -%371 = OpConvertSToF %75 %370 -%372 = OpLoad %25 %54 -%373 = OpCompositeExtract %7 %371 0 -%374 = OpCompositeExtract %7 %371 1 -%375 = OpCompositeExtract %7 %371 0 -%376 = OpCompositeExtract %7 %371 1 -%377 = OpCompositeConstruct %25 %373 %374 %375 %376 -%378 = OpFAdd %25 %372 %377 -%379 = OpLoad %7 %52 -%380 = OpCompositeConstruct %25 %379 %379 %379 %379 -%381 = OpFAdd %25 %378 %380 -OpStore %61 %381 +%357 = OpLoad %25 %54 +%358 = OpAccessChain %79 %31 %72 +%359 = OpLoad %10 %358 +%360 = OpAccessChain %97 %41 %72 +%361 = OpLoad %21 %360 +%362 = OpSampledImage %100 %359 %361 +%363 = OpImageSampleExplicitLod %25 %362 %76 Grad %76 %76 +%364 = OpFAdd %25 %357 %363 +OpStore %54 %364 +%365 = OpLoad %25 %54 +%366 = OpAccessChain %79 %31 %66 +%367 = OpLoad %10 %366 +%368 = OpAccessChain %97 %41 %66 +%369 = OpLoad %21 %368 +%370 = OpSampledImage %100 %367 %369 +%371 = OpImageSampleExplicitLod %25 %370 %76 Lod %6 +%372 = OpFAdd %25 %365 %371 +OpStore %54 %372 +%373 = OpLoad %25 %54 +%374 = OpAccessChain %79 %31 %71 +%375 = OpLoad %10 %374 +%376 = OpAccessChain %97 %41 %71 +%377 = OpLoad %21 %376 +%378 = OpSampledImage %100 %375 %377 +%379 = OpImageSampleExplicitLod %25 %378 %76 Lod %6 +%380 = OpFAdd %25 %373 %379 +OpStore %54 %380 +%381 = OpLoad %25 %54 +%382 = OpAccessChain %79 %31 %72 +%383 = OpLoad %10 %382 +%384 = OpAccessChain %97 %41 %72 +%385 = OpLoad %21 %384 +%386 = OpSampledImage %100 %383 %385 +%387 = OpImageSampleExplicitLod %25 %386 %76 Lod %6 +%388 = OpFAdd %25 %381 %387 +OpStore %54 %388 +%390 = OpAccessChain %389 %39 %66 +%391 = OpLoad %19 %390 +%392 = OpLoad %25 %54 +%393 = OpImageQuerySize %26 %391 +%394 = OpULessThan %157 %77 %393 +%395 = OpAll %150 %394 +OpSelectionMerge %396 None +OpBranchConditional %395 %397 %396 +%397 = OpLabel +OpImageWrite %391 %77 %392 +OpBranch %396 +%396 = OpLabel +%398 = OpAccessChain %389 %39 %71 +%399 = OpLoad %19 %398 +%400 = OpLoad %25 %54 +%401 = OpImageQuerySize %26 %399 +%402 = OpULessThan %157 %77 %401 +%403 = OpAll %150 %402 +OpSelectionMerge %404 None +OpBranchConditional %403 %405 %404 +%405 = OpLabel +OpImageWrite %399 %77 %400 +OpBranch %404 +%404 = OpLabel +%406 = OpAccessChain %389 %39 %72 +%407 = OpLoad %19 %406 +%408 = OpLoad %25 %54 +%409 = OpImageQuerySize %26 %407 +%410 = OpULessThan %157 %77 %409 +%411 = OpAll %150 %410 +OpSelectionMerge %412 None +OpBranchConditional %411 %413 %412 +%413 = OpLabel +OpImageWrite %407 %77 %408 +OpBranch %412 +%412 = OpLabel +%414 = OpLoad %26 %50 +%415 = OpLoad %4 %48 +%416 = OpCompositeConstruct %26 %415 %415 +%417 = OpIAdd %26 %414 %416 +%418 = OpConvertSToF %75 %417 +%419 = OpLoad %25 %54 +%420 = OpCompositeExtract %7 %418 0 +%421 = OpCompositeExtract %7 %418 1 +%422 = OpCompositeExtract %7 %418 0 +%423 = OpCompositeExtract %7 %418 1 +%424 = OpCompositeConstruct %25 %420 %421 %422 %423 +%425 = OpFAdd %25 %419 %424 +%426 = OpLoad %7 %52 +%427 = OpCompositeConstruct %25 %426 %426 %426 %426 +%428 = OpFAdd %25 %425 %427 +OpStore %61 %428 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/snapshots.rs b/tests/snapshots.rs index 727149d986..f6a3749c3c 100644 --- a/tests/snapshots.rs +++ b/tests/snapshots.rs @@ -252,8 +252,8 @@ fn write_output_msl( let mut options = options.clone(); options.bounds_check_policies = bounds_check_policies; - let (string, tr_info) = - msl::write_string(module, info, &options, pipeline_options).expect("Metal write failed"); + let (string, tr_info) = msl::write_string(module, info, &options, pipeline_options) + .unwrap_or_else(|err| panic!("Metal write failed: {}", err)); for (ep, result) in module.entry_points.iter().zip(tr_info.entry_point_names) { if let Err(error) = result { From 7ec8ff3c0692f208f2c23e1475fec37415c8d748 Mon Sep 17 00:00:00 2001 From: Connor Fitzgerald Date: Wed, 20 Apr 2022 01:36:24 -0400 Subject: [PATCH 2/3] Fix msl binding array issue --- src/back/msl/writer.rs | 18 +++++++++++++----- 1 file changed, 13 insertions(+), 5 deletions(-) diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index 55dc198fdb..836985846c 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -3340,11 +3340,19 @@ impl Writer { } if let Some(ref br) = var.binding { let good = match options.per_stage_map[ep.stage].resources.get(br) { - Some(target) => match module.types[var.ty].inner { - crate::TypeInner::Image { .. } => target.texture.is_some(), - crate::TypeInner::Sampler { .. } => target.sampler.is_some(), - _ => target.buffer.is_some(), - }, + Some(target) => { + let binding_ty = match module.types[var.ty].inner { + crate::TypeInner::BindingArray { base, .. } => { + &module.types[base].inner + } + ref ty => ty, + }; + match *binding_ty { + crate::TypeInner::Image { .. } => target.texture.is_some(), + crate::TypeInner::Sampler { .. } => target.sampler.is_some(), + _ => target.buffer.is_some(), + } + } None => false, }; if !good { From fd6aaba1364459d8a0aa307aac9d0a20b97c9d9c Mon Sep 17 00:00:00 2001 From: Connor Fitzgerald Date: Mon, 25 Apr 2022 00:13:13 -0400 Subject: [PATCH 3/3] Clarify docs --- src/proc/index.rs | 11 ++++++----- 1 file changed, 6 insertions(+), 5 deletions(-) diff --git a/src/proc/index.rs b/src/proc/index.rs index 60204fd3c2..3fea79ec01 100644 --- a/src/proc/index.rs +++ b/src/proc/index.rs @@ -131,20 +131,21 @@ impl Default for BoundsCheckPolicy { } impl BoundsCheckPolicies { - /// Determine which policy applies to `access`. + /// Determine which policy applies to `base`. /// - /// `access` is a subtree of `Access` and `AccessIndex` expressions, - /// operating either on a pointer to a value, or on a value directly. + /// `base` is the "base" expression (the expression being indexed) of a `Access` + /// and `AccessIndex` expression. This is either a pointer, a value, being directly + /// indexed, or a binding array. /// /// See the documentation for [`BoundsCheckPolicy`] for details about /// when each policy applies. pub fn choose_policy( &self, - access: Handle, + base: Handle, types: &UniqueArena, info: &valid::FunctionInfo, ) -> BoundsCheckPolicy { - let ty = info[access].ty.inner_with(types); + let ty = info[base].ty.inner_with(types); if let crate::TypeInner::BindingArray { .. } = *ty { return self.binding_array;