Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Binding arrays play nice with bounds checks #1855

Merged
merged 3 commits into from
Apr 25, 2022
Merged

Binding arrays play nice with bounds checks #1855

Show file tree
Hide file tree
Changes from 2 commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
1ac8dae
Binding arrays play nice with bounds checks
cwfitzgerald Apr 19, 2022
7ec8ff3
Fix msl binding array issue
cwfitzgerald Apr 20, 2022
fd6aaba
Clarify docs
cwfitzgerald Apr 25, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
23 changes: 16 additions & 7 deletions src/back/msl/writer.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1270,13 +1270,14 @@ impl<W: Write> Writer<W> {
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,
Expand Down Expand Up @@ -3339,11 +3340,19 @@ impl<W: Write> Writer<W> {
}
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(),
}
}
Comment on lines -3342 to +3355
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This seems like an unrelated fix. It's fine to include it in this PR, but if it has something to do with bounds checks, let me know, because I'm not understanding the patch.

Copy link
Member Author

@cwfitzgerald cwfitzgerald Apr 25, 2022 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Ah yeah, sorry this should have been titled "fixes needed to get the binding indexing wgpu PR working" :)

tl;dr: Because of the change above this, it is now possible to get the type of binding array, and we need to "punch through" the binding array to find the actual type of binding we're looking at.

It's only related to bounds checks as this is code that is only hit when bounds checks are on.

None => false,
};
if !good {
Expand Down
12 changes: 11 additions & 1 deletion src/proc/index.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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`.
Expand All @@ -140,7 +144,13 @@ impl BoundsCheckPolicies {
types: &UniqueArena<crate::Type>,
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() {
jimblandy marked this conversation as resolved.
Show resolved Hide resolved
Some(crate::AddressSpace::Storage { access: _ } | crate::AddressSpace::Uniform) => {
self.buffer
}
Expand Down
5 changes: 5 additions & 0 deletions tests/in/binding-arrays.param.ron
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -35,5 +35,10 @@
binding_map: {
(group: 0, binding: 0): (binding_array_size: Some(10)),
},
),
bounds_check_policies: (
index: ReadZeroSkipWrite,
buffer: ReadZeroSkipWrite,
image: ReadZeroSkipWrite,
)
)
24 changes: 18 additions & 6 deletions tests/out/msl/binding-arrays.msl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,6 +4,12 @@

using metal::uint;

struct DefaultConstructible {
template<typename T>
operator T() && {
return T {};
}
};
struct UniformIndex {
uint index;
};
Expand Down Expand Up @@ -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());
Expand Down Expand Up @@ -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));
}
jimblandy marked this conversation as resolved.
Show resolved Hide resolved
metal::int2 _e253 = i2_;
int _e254 = i1_;
metal::float2 v2_ = static_cast<metal::float2>(_e253 + metal::int2(_e254));
Expand Down
Loading