Upgrade to wgpu version 26.0

crates/bevy_anti_aliasing/src/contrast_adaptive_sharpening/node.rs

@@ -98,6 +98,7 @@ impl Node for CasNode {
             label: Some("contrast_adaptive_sharpening"),
             color_attachments: &[Some(RenderPassColorAttachment {
                 view: destination,
+                depth_slice: None,
                 resolve_target: None,
                 ops: Operations::default(),
             })],

crates/bevy_anti_aliasing/src/fxaa/node.rs

@@ -64,6 +64,7 @@ impl ViewNode for FxaaNode {
             label: Some("fxaa_pass"),
             color_attachments: &[Some(RenderPassColorAttachment {
                 view: destination,
+                depth_slice: None,
                 resolve_target: None,
                 ops: Operations::default(),
             })],

crates/bevy_anti_aliasing/src/smaa/mod.rs

@@ -896,6 +896,7 @@ fn perform_edge_detection(
         label: Some("SMAA edge detection pass"),
         color_attachments: &[Some(RenderPassColorAttachment {
             view: &smaa_textures.edge_detection_color_texture.default_view,
+            depth_slice: None,
             resolve_target: None,
             ops: default(),
         })],
@@ -951,6 +952,7 @@ fn perform_blending_weight_calculation(
         label: Some("SMAA blending weight calculation pass"),
         color_attachments: &[Some(RenderPassColorAttachment {
             view: &smaa_textures.blend_texture.default_view,
+            depth_slice: None,
             resolve_target: None,
             ops: default(),
         })],
@@ -1007,6 +1009,7 @@ fn perform_neighborhood_blending(
         label: Some("SMAA neighborhood blending pass"),
         color_attachments: &[Some(RenderPassColorAttachment {
             view: destination,
+            depth_slice: None,
             resolve_target: None,
             ops: default(),
         })],

crates/bevy_anti_aliasing/src/taa/mod.rs

@@ -202,11 +202,13 @@ impl ViewNode for TemporalAntiAliasNode {
                 color_attachments: &[
                     Some(RenderPassColorAttachment {
                         view: view_target.destination,
+                        depth_slice: None,
                         resolve_target: None,
                         ops: Operations::default(),
                     }),
                     Some(RenderPassColorAttachment {
                         view: &taa_history_textures.write.default_view,
+                        depth_slice: None,
                         resolve_target: None,
                         ops: Operations::default(),
                     }),

crates/bevy_camera/Cargo.toml

@@ -26,7 +26,7 @@ bevy_color = { path = "../bevy_color", version = "0.17.0-dev", features = [
 bevy_window = { path = "../bevy_window", version = "0.17.0-dev" }
 
 # other
-wgpu-types = { version = "25", default-features = false }
+wgpu-types = { version = "26", default-features = false }
 serde = { version = "1", default-features = false, features = ["derive"] }
 thiserror = { version = "2", default-features = false }
 downcast-rs = { version = "2", default-features = false, features = ["std"] }

crates/bevy_color/Cargo.toml

@@ -20,7 +20,7 @@ serde = { version = "1.0", features = [
 ], default-features = false, optional = true }
 thiserror = { version = "2", default-features = false }
 derive_more = { version = "2", default-features = false, features = ["from"] }
-wgpu-types = { version = "25", default-features = false, optional = true }
+wgpu-types = { version = "26", default-features = false, optional = true }
 encase = { version = "0.10", default-features = false, optional = true }
 
 [features]

crates/bevy_core_pipeline/src/bloom/mod.rs

@@ -186,6 +186,7 @@ impl ViewNode for BloomNode {
                         label: Some("bloom_downsampling_first_pass"),
                         color_attachments: &[Some(RenderPassColorAttachment {
                             view,
+                            depth_slice: None,
                             resolve_target: None,
                             ops: Operations::default(),
                         })],
@@ -210,6 +211,7 @@ impl ViewNode for BloomNode {
                         label: Some("bloom_downsampling_pass"),
                         color_attachments: &[Some(RenderPassColorAttachment {
                             view,
+                            depth_slice: None,
                             resolve_target: None,
                             ops: Operations::default(),
                         })],
@@ -234,6 +236,7 @@ impl ViewNode for BloomNode {
                         label: Some("bloom_upsampling_pass"),
                         color_attachments: &[Some(RenderPassColorAttachment {
                             view,
+                            depth_slice: None,
                             resolve_target: None,
                             ops: Operations {
                                 load: LoadOp::Load,

crates/bevy_core_pipeline/src/deferred/node.rs

@@ -176,6 +176,7 @@ fn run_deferred_prepass<'w>(
                                 load: bevy_render::render_resource::LoadOp::Load,
                                 store: StoreOp::Store,
                             },
+                            depth_slice: None,
                         }
                     }
                     #[cfg(any(

crates/bevy_core_pipeline/src/dof/mod.rs

@@ -409,6 +409,7 @@ impl ViewNode for DepthOfFieldNode {
             let mut color_attachments: SmallVec<[_; 2]> = SmallVec::new();
             color_attachments.push(Some(RenderPassColorAttachment {
                 view: postprocess.destination,
+                depth_slice: None,
                 resolve_target: None,
                 ops: Operations {
                     load: LoadOp::Clear(default()),
@@ -429,6 +430,7 @@ impl ViewNode for DepthOfFieldNode {
                 };
                 color_attachments.push(Some(RenderPassColorAttachment {
                     view: &auxiliary_dof_texture.default_view,
+                    depth_slice: None,
                     resolve_target: None,
                     ops: Operations {
                         load: LoadOp::Clear(default()),

crates/bevy_core_pipeline/src/motion_blur/node.rs

@@ -84,6 +84,7 @@ impl ViewNode for MotionBlurNode {
             label: Some("motion_blur_pass"),
             color_attachments: &[Some(RenderPassColorAttachment {
                 view: post_process.destination,
+                depth_slice: None,
                 resolve_target: None,
                 ops: Operations::default(),
             })],

crates/bevy_core_pipeline/src/msaa_writeback.rs

@@ -87,6 +87,7 @@ impl ViewNode for MsaaWritebackNode {
             color_attachments: &[Some(RenderPassColorAttachment {
                 // If MSAA is enabled, then the sampled texture will always exist
                 view: target.sampled_main_texture_view().unwrap(),
+                depth_slice: None,
                 resolve_target: Some(post_process.destination),
                 ops: Operations {
                     load: LoadOp::Clear(LinearRgba::BLACK.into()),

crates/bevy_core_pipeline/src/post_process/mod.rs

@@ -379,6 +379,7 @@ impl ViewNode for PostProcessingNode {
             label: Some("postprocessing pass"),
             color_attachments: &[Some(RenderPassColorAttachment {
                 view: post_process.destination,
+                depth_slice: None,
                 resolve_target: None,
                 ops: Operations::default(),
             })],

crates/bevy_core_pipeline/src/tonemapping/node.rs

@@ -117,6 +117,7 @@ impl ViewNode for TonemappingNode {
             label: Some("tonemapping_pass"),
             color_attachments: &[Some(RenderPassColorAttachment {
                 view: destination,
+                depth_slice: None,
                 resolve_target: None,
                 ops: Operations {
                     load: LoadOp::Clear(Default::default()), // TODO shouldn't need to be cleared

crates/bevy_image/Cargo.toml

@@ -70,7 +70,7 @@ image = { version = "0.25.2", default-features = false }
 # misc
 bitflags = { version = "2.3", features = ["serde"] }
 bytemuck = { version = "1.5" }
-wgpu-types = { version = "25", default-features = false }
+wgpu-types = { version = "26", default-features = false }
 serde = { version = "1", features = ["derive"] }
 thiserror = { version = "2", default-features = false }
 futures-lite = "2.0.1"

crates/bevy_mesh/Cargo.toml

@@ -26,7 +26,7 @@ bevy_platform = { path = "../bevy_platform", version = "0.17.0-dev", default-fea
 # other
 bitflags = { version = "2.3", features = ["serde"] }
 bytemuck = { version = "1.5" }
-wgpu-types = { version = "25", default-features = false }
+wgpu-types = { version = "26", default-features = false }
 serde = { version = "1", default-features = false, features = [
   "derive",
 ], optional = true }

crates/bevy_pbr/src/light_probe/downsample.wgsl

@@ -50,22 +50,16 @@ var<workgroup> spd_intermediate_a: array<array<f32, 16>, 16>;
 @workgroup_size(256, 1, 1)
 fn downsample_first(
     @builtin(workgroup_id) workgroup_id: vec3u,
-    @builtin(local_invocation_index) local_invocation_index: u32,
-#ifdef SUBGROUP_SUPPORT
-    @builtin(subgroup_invocation_id) subgroup_invocation_id: u32,
-#endif
+    @builtin(local_invocation_index) local_invocation_index: u32
 ) {
-#ifndef SUBGROUP_SUPPORT
-    let subgroup_invocation_id = 0u;
-#endif
 
     let sub_xy = remap_for_wave_reduction(local_invocation_index % 64u);
     let x = sub_xy.x + 8u * ((local_invocation_index >> 6u) % 2u);
     let y = sub_xy.y + 8u * (local_invocation_index >> 7u);
 
-    spd_downsample_mips_0_1(x, y, workgroup_id.xy, local_invocation_index, constants.mips, workgroup_id.z, subgroup_invocation_id);
+    spd_downsample_mips_0_1(x, y, workgroup_id.xy, local_invocation_index, constants.mips, workgroup_id.z);
 
-    spd_downsample_next_four(x, y, workgroup_id.xy, local_invocation_index, 2u, constants.mips, workgroup_id.z, subgroup_invocation_id);
+    spd_downsample_next_four(x, y, workgroup_id.xy, local_invocation_index, 2u, constants.mips, workgroup_id.z);
 }
 
 // TODO: Once wgpu supports globallycoherent buffers, make it actually a single pass
@@ -74,24 +68,17 @@ fn downsample_first(
 fn downsample_second(
     @builtin(workgroup_id) workgroup_id: vec3u,
     @builtin(local_invocation_index) local_invocation_index: u32,
-#ifdef SUBGROUP_SUPPORT
-    @builtin(subgroup_invocation_id) subgroup_invocation_id: u32,
-#endif
 ) {
-#ifndef SUBGROUP_SUPPORT
-    let subgroup_invocation_id = 0u;
-#endif
-
     let sub_xy = remap_for_wave_reduction(local_invocation_index % 64u);
     let x = sub_xy.x + 8u * ((local_invocation_index >> 6u) % 2u);
     let y = sub_xy.y + 8u * (local_invocation_index >> 7u);
 
     spd_downsample_mips_6_7(x, y, constants.mips, workgroup_id.z);
 
-    spd_downsample_next_four(x, y, vec2(0u), local_invocation_index, 8u, constants.mips, workgroup_id.z, subgroup_invocation_id);
+    spd_downsample_next_four(x, y, vec2(0u), local_invocation_index, 8u, constants.mips, workgroup_id.z);
 }
 
-fn spd_downsample_mips_0_1(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, mips: u32, slice: u32, subgroup_invocation_id: u32) {
+fn spd_downsample_mips_0_1(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, mips: u32, slice: u32) {
     var v: array<vec4f, 4>;
 
     var tex = (workgroup_id * 64u) + vec2(x * 2u, y * 2u);
@@ -117,10 +104,10 @@ fn spd_downsample_mips_0_1(x: u32, y: u32, workgroup_id: vec2u, local_invocation
     if mips <= 1u { return; }
 
 #ifdef SUBGROUP_SUPPORT
-    v[0] = spd_reduce_quad(v[0], subgroup_invocation_id);
-    v[1] = spd_reduce_quad(v[1], subgroup_invocation_id);
-    v[2] = spd_reduce_quad(v[2], subgroup_invocation_id);
-    v[3] = spd_reduce_quad(v[3], subgroup_invocation_id);
+    v[0] = spd_reduce_quad(v[0]);
+    v[1] = spd_reduce_quad(v[1]);
+    v[2] = spd_reduce_quad(v[2]);
+    v[3] = spd_reduce_quad(v[3]);
 
     if local_invocation_index % 4u == 0u {
         spd_store((workgroup_id * 16u) + vec2(x / 2u, y / 2u), v[0], 1u, slice);
@@ -160,28 +147,28 @@ fn spd_downsample_mips_0_1(x: u32, y: u32, workgroup_id: vec2u, local_invocation
 #endif
 }
 
-fn spd_downsample_next_four(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, mips: u32, slice: u32, subgroup_invocation_id: u32) {
+fn spd_downsample_next_four(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, mips: u32, slice: u32) {
     if mips <= base_mip { return; }
     workgroupBarrier();
-    spd_downsample_mip_2(x, y, workgroup_id, local_invocation_index, base_mip, slice, subgroup_invocation_id);
+    spd_downsample_mip_2(x, y, workgroup_id, local_invocation_index, base_mip, slice);
 
     if mips <= base_mip + 1u { return; }
     workgroupBarrier();
-    spd_downsample_mip_3(x, y, workgroup_id, local_invocation_index, base_mip + 1u, slice, subgroup_invocation_id);
+    spd_downsample_mip_3(x, y, workgroup_id, local_invocation_index, base_mip + 1u, slice);
 
     if mips <= base_mip + 2u { return; }
     workgroupBarrier();
-    spd_downsample_mip_4(x, y, workgroup_id, local_invocation_index, base_mip + 2u, slice, subgroup_invocation_id);
+    spd_downsample_mip_4(x, y, workgroup_id, local_invocation_index, base_mip + 2u, slice);
 
     if mips <= base_mip + 3u { return; }
     workgroupBarrier();
-    spd_downsample_mip_5(x, y, workgroup_id, local_invocation_index, base_mip + 3u, slice, subgroup_invocation_id);
+    spd_downsample_mip_5(x, y, workgroup_id, local_invocation_index, base_mip + 3u, slice);
 }
 
-fn spd_downsample_mip_2(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, slice: u32, subgroup_invocation_id: u32) {
+fn spd_downsample_mip_2(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, slice: u32) {
 #ifdef SUBGROUP_SUPPORT
     var v = spd_load_intermediate(x, y);
-    v = spd_reduce_quad(v, subgroup_invocation_id);
+    v = spd_reduce_quad(v);
     if local_invocation_index % 4u == 0u {
         spd_store((workgroup_id * 8u) + vec2(x / 2u, y / 2u), v, base_mip, slice);
         spd_store_intermediate(x + (y / 2u) % 2u, y, v);
@@ -200,11 +187,11 @@ fn spd_downsample_mip_2(x: u32, y: u32, workgroup_id: vec2u, local_invocation_in
 #endif
 }
 
-fn spd_downsample_mip_3(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, slice: u32, subgroup_invocation_id: u32) {
+fn spd_downsample_mip_3(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, slice: u32) {
 #ifdef SUBGROUP_SUPPORT
     if local_invocation_index < 64u {
         var v = spd_load_intermediate(x * 2u + y % 2u, y * 2u);
-        v = spd_reduce_quad(v, subgroup_invocation_id);
+        v = spd_reduce_quad(v);
         if local_invocation_index % 4u == 0u {
             spd_store((workgroup_id * 4u) + vec2(x / 2u, y / 2u), v, base_mip, slice);
             spd_store_intermediate(x * 2u + y / 2u, y * 2u, v);
@@ -224,11 +211,11 @@ fn spd_downsample_mip_3(x: u32, y: u32, workgroup_id: vec2u, local_invocation_in
 #endif
 }
 
-fn spd_downsample_mip_4(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, slice: u32, subgroup_invocation_id: u32) {
+fn spd_downsample_mip_4(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, slice: u32) {
 #ifdef SUBGROUP_SUPPORT
     if local_invocation_index < 16u {
         var v = spd_load_intermediate(x * 4u + y, y * 4u);
-        v = spd_reduce_quad(v, subgroup_invocation_id);
+        v = spd_reduce_quad(v);
         if local_invocation_index % 4u == 0u {
             spd_store((workgroup_id * 2u) + vec2(x / 2u, y / 2u), v, base_mip, slice);
             spd_store_intermediate(x / 2u + y, 0u, v);
@@ -248,11 +235,11 @@ fn spd_downsample_mip_4(x: u32, y: u32, workgroup_id: vec2u, local_invocation_in
 #endif
 }
 
-fn spd_downsample_mip_5(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, slice: u32, subgroup_invocation_id: u32) {
+fn spd_downsample_mip_5(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, slice: u32) {
 #ifdef SUBGROUP_SUPPORT
     if local_invocation_index < 4u {
         var v = spd_load_intermediate(local_invocation_index, 0u);
-        v = spd_reduce_quad(v, subgroup_invocation_id);
+        v = spd_reduce_quad(v);
         if local_invocation_index % 4u == 0u {
             spd_store(workgroup_id, v, base_mip, slice);
         }
@@ -436,20 +423,12 @@ fn spd_reduce_4(v0: vec4f, v1: vec4f, v2: vec4f, v3: vec4f) -> vec4f {
 }
 
 #ifdef SUBGROUP_SUPPORT
-fn spd_reduce_quad(v: vec4f, subgroup_invocation_id: u32) -> vec4f {
-    let quad = subgroup_invocation_id & (~0x3u);
+fn spd_reduce_quad(v: vec4f) -> vec4f {
     let v0 = v;
-    let v1 = subgroupBroadcast(v, quad | 1u);
-    let v2 = subgroupBroadcast(v, quad | 2u);
-    let v3 = subgroupBroadcast(v, quad | 3u);
+    let v1 = quadSwapX(v);
+    let v2 = quadSwapY(v);
+    let v3 = quadSwapDiagonal(v);
     return spd_reduce_4(v0, v1, v2, v3);
-
-    // TODO: Use subgroup quad operations once wgpu supports them
-    // let v0 = v;
-    // let v1 = quadSwapX(v);
-    // let v2 = quadSwapY(v);
-    // let v3 = quadSwapDiagonal(v);
-    // return spd_reduce_4(v0, v1, v2, v3);
 }
 #endif
 

crates/bevy_pbr/src/meshlet/visibility_buffer_raster_node.rs

@@ -595,6 +595,7 @@ fn raster_pass(
         }),
         color_attachments: &[Some(RenderPassColorAttachment {
             view: dummy_render_target,
+            depth_slice: None,
             resolve_target: None,
             ops: Operations {
                 load: LoadOp::Clear(LinearRgba::BLACK.into()),

crates/bevy_pbr/src/ssr/mod.rs

@@ -302,6 +302,7 @@ impl ViewNode for ScreenSpaceReflectionsNode {
             label: Some("SSR pass"),
             color_attachments: &[Some(RenderPassColorAttachment {
                 view: postprocess.destination,
+                depth_slice: None,
                 resolve_target: None,
                 ops: Operations::default(),
             })],

crates/bevy_pbr/src/volumetric_fog/render.rs

@@ -431,6 +431,7 @@ impl ViewNode for VolumetricFogNode {
                 label: Some("volumetric lighting pass"),
                 color_attachments: &[Some(RenderPassColorAttachment {
                     view: view_target.main_texture_view(),
+                    depth_slice: None,
                     resolve_target: None,
                     ops: Operations {
                         load: LoadOp::Load,

crates/bevy_reflect/Cargo.toml

@@ -109,7 +109,7 @@ uuid = { version = "1.13.1", default-features = false, optional = true, features
   "serde",
 ] }
 variadics_please = "1.1"
-wgpu-types = { version = "25", features = [
+wgpu-types = { version = "26", features = [
   "serde",
 ], optional = true, default-features = false }
 

crates/bevy_render/Cargo.toml

@@ -91,7 +91,7 @@ codespan-reporting = "0.12.0"
 # It is enabled for now to avoid having to do a significant overhaul of the renderer just for wasm.
 # When the 'atomics' feature is enabled `fragile-send-sync-non-atomic` does nothing
 # and Bevy instead wraps `wgpu` types to verify they are not used off their origin thread.
-wgpu = { version = "25", default-features = false, features = [
+wgpu = { version = "26", default-features = false, features = [
   "wgsl",
   "dx12",
   "metal",
@@ -100,7 +100,7 @@ wgpu = { version = "25", default-features = false, features = [
   "naga-ir",
   "fragile-send-sync-non-atomic-wasm",
 ] }
-naga = { version = "25", features = ["wgsl-in"] }
+naga = { version = "26", features = ["wgsl-in"] }
 bytemuck = { version = "1.5", features = ["derive", "must_cast"] }
 downcast-rs = { version = "2", default-features = false, features = ["std"] }
 thiserror = { version = "2", default-features = false }