Add a test

tests/tests/bgra8unorm_storage.rs

@@ -0,0 +1,158 @@
+//! Tests for texture copy bounds checks.
+
+use std::borrow::Cow;
+
+use wasm_bindgen_test::*;
+use wgpu_test::{initialize_test, TestParameters};
+
+const SHADER_SRC: &str = "
+@group(0) @binding(0) var tex: texture_storage_2d<bgra8unorm, write>;
+@compute @workgroup_size(256)
+fn main(@builtin(workgroup_id) wgid: vec3<u32>) {
+    var texel = vec4f(0.0, 0.0, 1.0, 1.0);
+    textureStore(tex, wgid.xy, texel);
+}
+";
+
+#[test]
+#[wasm_bindgen_test]
+fn bgra8unorm_storage() {
+    let parameters = TestParameters::default()
+        .limits(wgpu::Limits {
+            max_storage_textures_per_shader_stage: 1,
+            ..Default::default()
+        })
+        .features(wgpu::Features::BGRA8UNORM_STORAGE);
+
+    initialize_test(parameters, |ctx| {
+        let device = &ctx.device;
+        let texture = ctx.device.create_texture(&wgpu::TextureDescriptor {
+            label: None,
+            size: wgpu::Extent3d {
+                width: 256,
+                height: 256,
+                depth_or_array_layers: 1,
+            },
+            mip_level_count: 1,
+            sample_count: 1,
+            dimension: wgpu::TextureDimension::D2,
+            format: wgpu::TextureFormat::Bgra8Unorm,
+            usage: wgpu::TextureUsages::STORAGE_BINDING | wgpu::TextureUsages::COPY_SRC,
+            view_formats: &[],
+        });
+
+        let view = texture.create_view(&wgpu::TextureViewDescriptor {
+            label: None,
+            format: None,
+            dimension: None,
+            aspect: wgpu::TextureAspect::All,
+            base_mip_level: 0,
+            base_array_layer: 0,
+            mip_level_count: Some(1),
+            array_layer_count: Some(1),
+        });
+
+        let readback_buffer = device.create_buffer(&wgpu::BufferDescriptor {
+            label: None,
+            size: 256 * 256 * 4,
+            usage: wgpu::BufferUsages::COPY_DST | wgpu::BufferUsages::MAP_READ,
+            mapped_at_creation: false,
+        });
+
+        let bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+            label: None,
+            entries: &[wgpu::BindGroupLayoutEntry {
+                binding: 0,
+                visibility: wgpu::ShaderStages::COMPUTE,
+                ty: wgpu::BindingType::StorageTexture {
+                    access: wgpu::StorageTextureAccess::WriteOnly,
+                    format: wgpu::TextureFormat::Bgra8Unorm,
+                    view_dimension: wgpu::TextureViewDimension::D2,
+                },
+                count: None,
+            }],
+        });
+
+        let bg = device.create_bind_group(&wgpu::BindGroupDescriptor {
+            label: None,
+            layout: &bgl,
+            entries: &[wgpu::BindGroupEntry {
+                binding: 0,
+                resource: wgpu::BindingResource::TextureView(&view),
+            }],
+        });
+
+        let pl = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+            label: None,
+            bind_group_layouts: &[&bgl],
+            push_constant_ranges: &[],
+        });
+
+        let module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
+            label: None,
+            source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(SHADER_SRC)),
+        });
+
+        let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+            label: None,
+            layout: Some(&pl),
+            entry_point: "main",
+            module: &module,
+        });
+
+        let mut encoder =
+            device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
+
+        {
+            let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
+                label: None,
+                timestamp_writes: None,
+            });
+
+            pass.set_bind_group(0, &bg, &[]);
+            pass.set_pipeline(&pipeline);
+            pass.dispatch_workgroups(256, 256, 1);
+        }
+
+        encoder.copy_texture_to_buffer(
+            wgpu::ImageCopyTexture {
+                texture: &texture,
+                mip_level: 0,
+                origin: wgpu::Origin3d { x: 0, y: 0, z: 0 },
+                aspect: wgpu::TextureAspect::All,
+            },
+            wgpu::ImageCopyBuffer {
+                buffer: &readback_buffer,
+                layout: wgpu::ImageDataLayout {
+                    offset: 0,
+                    bytes_per_row: Some(256 * 4),
+                    rows_per_image: Some(256),
+                },
+            },
+            wgpu::Extent3d {
+                width: 256,
+                height: 256,
+                depth_or_array_layers: 1,
+            },
+        );
+
+        ctx.queue.submit(Some(encoder.finish()));
+
+        let buffer_slice = readback_buffer.slice(..);
+        buffer_slice.map_async(wgpu::MapMode::Read, Result::unwrap);
+        device.poll(wgpu::Maintain::Wait);
+
+        {
+            let texels = buffer_slice.get_mapped_range();
+            assert_eq!(texels.len(), 256 * 256 * 4);
+            for texel in texels.chunks(4) {
+                assert_eq!(texel[0], 255); // b
+                assert_eq!(texel[1], 0); // g
+                assert_eq!(texel[2], 0); // r
+                assert_eq!(texel[3], 255); // a
+            }
+        }
+
+        readback_buffer.unmap();
+    });
+}

tests/tests/root.rs

@@ -6,6 +6,7 @@ mod regression {
     mod issue_4122;
 }
 
+mod bgra8unorm_storage;
 mod bind_group_layout_dedup;
 mod buffer;
 mod buffer_copy;

wgpu-hal/src/vulkan/adapter.rs

@@ -1780,13 +1780,14 @@ fn supports_bgra8unorm_storage(
             buffer_features: vk::FormatFeatureFlags2::empty(),
         };
 
+        let p_next: *mut ash::vk::FormatProperties3 = &mut properties3;
         let mut properties2 = vk::FormatProperties2 {
             s_type: vk::StructureType::FORMAT_PROPERTIES_2,
-            p_next: &mut properties3 as *mut ash::vk::FormatProperties3 as *mut _,
+            p_next: p_next as *mut _,
             format_properties: vk::FormatProperties {
-                linear_tiling_features: FormatFeatureFlags::empty(),
-                optimal_tiling_features: FormatFeatureFlags::empty(),
-                buffer_features: FormatFeatureFlags::empty(),
+                linear_tiling_features: vk::FormatFeatureFlags::empty(),
+                optimal_tiling_features: vk::FormatFeatureFlags::empty(),
+                buffer_features: vk::FormatFeatureFlags::empty(),
             },
         };