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

Buffer zero init test for binding (and use in a compute shader) #1191

Merged
merged 3 commits into from
Feb 6, 2021
Merged

Buffer zero init test for binding (and use in a compute shader) #1191

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
5e488d4
Added player test with compute shader reading & writing an uninitiali…
Wumpf Feb 3, 2021
52aee1a
Better error output for player test
Wumpf Feb 3, 2021
fd810f8
merged buffer-zero-init-for-binding and buffer-zero-init-for-mapping-…
Wumpf Feb 6, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Jump to
Jump to file
Failed to load files.
Diff view
Diff view
15 changes: 15 additions & 0 deletions player/tests/data/buffer-zero-init-for-binding.wgsl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,15 @@
[[builtin(global_invocation_id)]]
var global_id: vec3<u32>;

[[block]]
struct InOutBuffer {
data: [[stride(4)]] array<u32>;
};

[[group(0), binding(0)]]
var<storage> buffer: [[access(read_write)]] InOutBuffer;

[[stage(compute), workgroup_size(1)]]
fn main() {
buffer.data[global_id.x] = buffer.data[global_id.x] + global_id.x;
}
237 changes: 167 additions & 70 deletions player/tests/data/buffer-zero-init.ron
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,70 +1,167 @@
(
features: (bits: 0x0),
expectations: [
(
name: "mapped_at_creation: false, with MAP_WRITE",
buffer: (index: 0, epoch: 1),
offset: 0,
data: Raw([0x00, 0x00, 0x00, 0x00]),
),
(
name: "mapped_at_creation: false, without MAP_WRITE",
buffer: (index: 1, epoch: 1),
offset: 0,
data: Raw([0x00, 0x00, 0x00, 0x00]),
),
(
name: "partially written buffer",
buffer: (index: 2, epoch: 1),
offset: 0,
data: Raw([0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x80, 0xBF, 0x00, 0x00, 0x80, 0xBF, 0x00, 0x00, 0x80, 0x3F, 0x00, 0x00, 0x80, 0x3F, 0x00, 0x00, 0x00, 0x00]),
),
],
actions: [
CreateBuffer(
Id(0, 1, Empty),
(
label: Some("mapped_at_creation: false, with MAP_WRITE"),
size: 16,
usage: (
bits: 131, // STORAGE + MAP_READ + MAP_WRITE
),
mapped_at_creation: false,
),
),
CreateBuffer(
Id(1, 1, Empty),
(
label: Some("mapped_at_creation: false, without MAP_WRITE"),
size: 16,
usage: (
bits: 129, // STORAGE + MAP_READ
),
mapped_at_creation: false,
),
),
CreateBuffer(
Id(2, 1, Empty),
(
label: Some("partially written"),
size: 24,
usage: (
bits: 9, // MAP_READ + COPY_DST
),
mapped_at_creation: false,
),
),
WriteBuffer(
id: Id(2, 1, Empty),
data: "data1.bin",
range: (
start: 4,
end: 20,
),
queued: true,
),

// TODO: Add a buffer with `mapped_at_creation: true`. As of writing there's no unmap action we could insert here.
Submit(1, []),
],
)
(
features: (bits: 0x0),
expectations: [
// Ensuring that mapping zero-inits buffers.
(
name: "mapped_at_creation: false, with MAP_WRITE",
buffer: (index: 0, epoch: 1),
offset: 0,
data: Raw([0x00, 0x00, 0x00, 0x00]),
),
(
name: "mapped_at_creation: false, without MAP_WRITE",
buffer: (index: 1, epoch: 1),
offset: 0,
data: Raw([0x00, 0x00, 0x00, 0x00]),
),
(
name: "partially written buffer",
buffer: (index: 2, epoch: 1),
offset: 0,
data: Raw([0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x80, 0xBF,
0x00, 0x00, 0x80, 0xBF,
0x00, 0x00, 0x80, 0x3F,
0x00, 0x00, 0x80, 0x3F,
0x00, 0x00, 0x00, 0x00]),
),
// Ensuring that binding zero-inits buffers
// (by observing correct side effects of compute shader reading & writing values)
(
name: "buffer has correct values",
buffer: (index: 3, epoch: 1),
offset: 0,
data: Raw([0x00, 0x00, 0x00, 0x00,
0x01, 0x00, 0x00, 0x00,
0x02, 0x00, 0x00, 0x00,
0x03, 0x00, 0x00, 0x00]),
)
],
actions: [
CreateBuffer(
Id(0, 1, Empty),
(
label: Some("mapped_at_creation: false, with MAP_WRITE"),
size: 16,
usage: (
bits: 131, // STORAGE + MAP_READ + MAP_WRITE
),
mapped_at_creation: false,
),
),
CreateBuffer(
Id(1, 1, Empty),
(
label: Some("mapped_at_creation: false, without MAP_WRITE"),
size: 16,
usage: (
bits: 129, // STORAGE + MAP_READ
),
mapped_at_creation: false,
),
),
CreateBuffer(
Id(2, 1, Empty),
(
label: Some("partially written"),
size: 24,
usage: (
bits: 9, // MAP_READ + COPY_DST
),
mapped_at_creation: false,
),
),
WriteBuffer(
id: Id(2, 1, Empty),
data: "data1.bin",
range: (
start: 4,
end: 20,
),
queued: true,
),
CreateShaderModule(
id: Id(0, 1, Empty),
desc: (
label: None,
experimental_translation: true,
),
data: "buffer-zero-init-for-binding.wgsl",
),
CreateBuffer(Id(3, 1, Empty), (
label: Some("used in binding"),
size: 16,
usage: (
bits: 129, // STORAGE + MAP_READ
),
mapped_at_creation: false,
)),
CreateBindGroupLayout(Id(0, 1, Empty), (
label: None,
entries: [
(
binding: 0,
visibility: (
bits: 4,
),
ty: Buffer(
ty: Storage(
read_only: false,
),
has_dynamic_offset: false,
min_binding_size: Some(16),
),
count: None,
),
],
)),
CreateBindGroup(Id(0, 1, Empty), (
label: None,
layout: Id(0, 1, Empty),
entries: [
(
binding: 0,
resource: Buffer((
buffer_id: Id(3, 1, Empty),
offset: 0,
size: Some(16),
)),
),
],
)),
CreatePipelineLayout(Id(0, 1, Empty), (
label: None,
bind_group_layouts: [
Id(0, 1, Empty),
],
push_constant_ranges: [],
)),
CreateComputePipeline(Id(0, 1, Empty), (
label: None,
layout: Some(Id(0, 1, Empty)),
stage: (
module: Id(0, 1, Empty),
entry_point: "main",
),
)),
Submit(1, [
RunComputePass(
base: (
label: None,
commands: [
SetPipeline(Id(0, 1, Empty)),
SetBindGroup(
index: 0,
num_dynamic_offsets: 0,
bind_group_id: Id(0, 1, Empty),
),
Dispatch((4, 1, 1)),
],
dynamic_offsets: [],
string_data: [],
push_constant_data: [],
),
)
]),
]
)
5 changes: 4 additions & 1 deletion player/tests/test.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -144,7 +144,10 @@ impl Test<'_> {
};

if &expected_data[..] != contents {
panic!("Test expectation is not met!");
panic!(
"Test expectation is not met!\nBuffer content was:\n{:?}\nbut expected:\n{:?}",
contents, expected_data
);
}
}

Expand Down