add some more pipelined-rendering shader examples (#3041)

based on #3031 

Adds some examples showing of how to use the new pipelined rendering for custom shaders.

- a minimal shader example which doesn't use render assets
- the same but using glsl
- an example showing how to render instanced data
- a shader which uses the seconds since startup to animate some textures


Instancing shader:
![grafik](https://user-images.githubusercontent.com/22177966/139299294-e176b62a-53d1-4287-9a66-02fb55affc02.png)
Animated shader:
![animate_shader](https://user-images.githubusercontent.com/22177966/139299718-2940c0f3-8480-4ee0-98d7-b6ba40dc1472.gif)
(the gif makes it look a bit ugly)

Co-authored-by: Carter Anderson <mcanders1@gmail.com>

Cargo.toml

@@ -102,6 +102,7 @@ anyhow = "1.0.4"
 rand = "0.8.0"
 ron = "0.7.0"
 serde = { version = "1", features = ["derive"] }
+bytemuck = "1.7"
 # Needed to poll Task examples
 futures-lite = "1.11.3"
 
@@ -418,6 +419,22 @@ path = "examples/shader/shader_defs.rs"
 name = "shader_material"
 path = "examples/shader/shader_material.rs"
 
+[[example]]
+name = "shader_material_glsl"
+path = "examples/shader/shader_material_glsl.rs"
+
+[[example]]
+name = "shader_instancing"
+path = "examples/shader/shader_instancing.rs"
+
+[[example]]
+name = "animate_shader"
+path = "examples/shader/animate_shader.rs"
+
+[[example]]
+name = "compute_shader_game_of_life"
+path = "examples/shader/compute_shader_game_of_life.rs"
+
 # Tools
 [[example]]
 name = "bevymark"

assets/shaders/animate_shader.wgsl

@@ -0,0 +1,72 @@
+#import bevy_pbr::mesh_view_bind_group
+#import bevy_pbr::mesh_struct
+
+[[group(1), binding(0)]]
+var<uniform> mesh: Mesh;
+
+struct Vertex {
+    [[location(0)]] position: vec3<f32>;
+    [[location(1)]] normal: vec3<f32>;
+    [[location(2)]] uv: vec2<f32>;
+};
+
+struct VertexOutput {
+    [[builtin(position)]] clip_position: vec4<f32>;
+    [[location(0)]] uv: vec2<f32>;
+};
+
+[[stage(vertex)]]
+fn vertex(vertex: Vertex) -> VertexOutput {
+    let world_position = mesh.model * vec4<f32>(vertex.position, 1.0);
+
+    var out: VertexOutput;
+    out.clip_position = view.view_proj * world_position;
+    out.uv = vertex.uv;
+    return out;
+}
+
+
+struct Time {
+    time_since_startup: f32;
+};
+[[group(2), binding(0)]]
+var<uniform> time: Time;
+
+
+fn oklab_to_linear_srgb(c: vec3<f32>) -> vec3<f32> {
+    let L = c.x;
+    let a = c.y;
+    let b = c.z;
+
+    let l_ = L + 0.3963377774 * a + 0.2158037573 * b;
+    let m_ = L - 0.1055613458 * a - 0.0638541728 * b;
+    let s_ = L - 0.0894841775 * a - 1.2914855480 * b;
+
+    let l = l_*l_*l_;
+    let m = m_*m_*m_;
+    let s = s_*s_*s_;
+
+    return vec3<f32>(
+		 4.0767416621 * l - 3.3077115913 * m + 0.2309699292 * s,
+		-1.2684380046 * l + 2.6097574011 * m - 0.3413193965 * s,
+		-0.0041960863 * l - 0.7034186147 * m + 1.7076147010 * s,
+    );
+}
+
+[[stage(fragment)]]
+fn fragment(in: VertexOutput) -> [[location(0)]] vec4<f32> {
+    let speed = 2.0;
+    let t_1 = sin(time.time_since_startup * speed) * 0.5 + 0.5;
+    let t_2 = cos(time.time_since_startup * speed);
+
+    let distance_to_center = distance(in.uv, vec2<f32>(0.5)) * 1.4;
+
+    // blending is done in a perceptual color space: https://bottosson.github.io/posts/oklab/
+    let red = vec3<f32>(0.627955, 0.224863, 0.125846);
+    let green = vec3<f32>(0.86644, -0.233887, 0.179498);
+    let blue = vec3<f32>(0.701674, 0.274566, -0.169156);
+    let white = vec3<f32>(1.0, 0.0, 0.0);
+    let mixed = mix(mix(red, blue, t_1), mix(green, white, t_2), distance_to_center);
+
+    return vec4<f32>(oklab_to_linear_srgb(mixed), 1.0);
+}

assets/shaders/custom_material.frag

@@ -0,0 +1,11 @@
+#version 450
+
+layout(location = 0) out vec4 o_Target;
+
+layout(set = 1, binding = 0) uniform CustomMaterial {
+    vec4 Color;
+};
+
+void main() {
+    o_Target = Color;
+}

assets/shaders/custom_material.vert

@@ -0,0 +1,26 @@
+#version 450
+
+layout(location = 0) in vec3 Vertex_Position;
+layout(location = 1) in vec3 Vertex_Normal;
+layout(location = 2) in vec2 Vertex_Uv;
+
+layout(set = 0, binding = 0) uniform CameraViewProj {
+    mat4 ViewProj;
+    mat4 InverseView;
+    mat4 Projection;
+    vec3 WorldPosition;
+    float near;
+    float far;
+    float width;
+    float height;
+};
+
+layout(set = 2, binding = 0) uniform Mesh {
+    mat4 Model;
+    mat4 InverseTransposeModel;
+    uint flags;
+};
+
+void main() {
+    gl_Position = ViewProj * Model * vec4(Vertex_Position, 1.0);
+}

assets/shaders/game_of_life.wgsl

@@ -0,0 +1,67 @@
+[[group(0), binding(0)]]
+var texture: texture_storage_2d<rgba8unorm, read_write>;
+
+fn hash(value: u32) -> u32 {
+    var state = value;
+    state = state ^ 2747636419u;
+    state = state * 2654435769u;
+    state = state ^ state >> 16u;
+    state = state * 2654435769u;
+    state = state ^ state >> 16u;
+    state = state * 2654435769u;
+    return state;
+}
+fn randomFloat(value: u32) -> f32 {
+    return f32(hash(value)) / 4294967295.0;
+}
+
+[[stage(compute), workgroup_size(8, 8, 1)]]
+fn init([[builtin(global_invocation_id)]] invocation_id: vec3<u32>, [[builtin(num_workgroups)]] num_workgroups: vec3<u32>) {
+    let location = vec2<i32>(i32(invocation_id.x), i32(invocation_id.y));
+    let location_f32 = vec2<f32>(f32(invocation_id.x), f32(invocation_id.y));
+
+    let randomNumber = randomFloat(invocation_id.y * num_workgroups.x + invocation_id.x);
+    let alive = randomNumber > 0.9;
+    let color = vec4<f32>(f32(alive));
+
+    textureStore(texture, location, color);
+}
+
+
+fn get(location: vec2<i32>, offset_x: i32, offset_y: i32) -> i32 {
+    let value: vec4<f32> = textureLoad(texture, location + vec2<i32>(offset_x, offset_y));
+    return i32(value.x);
+}
+
+fn count_alive(location: vec2<i32>) -> i32 {
+    return get(location, -1, -1) +
+           get(location, -1,  0) +
+           get(location, -1,  1) +
+           get(location,  0, -1) +
+           get(location,  0,  1) +
+           get(location,  1, -1) +
+           get(location,  1,  0) +
+           get(location,  1,  1);
+}
+
+[[stage(compute), workgroup_size(8, 8, 1)]]
+fn update([[builtin(global_invocation_id)]] invocation_id: vec3<u32>) {
+    let location = vec2<i32>(i32(invocation_id.x), i32(invocation_id.y));
+
+    let n_alive = count_alive(location);
+    let color = vec4<f32>(f32(n_alive) / 8.0);
+
+    var alive: bool;
+    if (n_alive == 3) {
+        alive = true;
+    } else if (n_alive == 2) {
+        let currently_alive = get(location, 0, 0);
+        alive = bool(currently_alive);
+    } else {
+        alive = false;
+    }
+
+    storageBarrier();
+
+    textureStore(texture, location, vec4<f32>(f32(alive)));
+}

assets/shaders/hot.frag

@@ -1,11 +0,0 @@
-#version 450
-
-layout(location = 0) out vec4 o_Target;
-
-layout(set = 2, binding = 0) uniform MyMaterial_color {
-    vec4 color;
-};
-
-void main() {
-    o_Target = color * 0.5;
-}

assets/shaders/hot.vert

@@ -1,15 +0,0 @@
-#version 450
-
-layout(location = 0) in vec3 Vertex_Position;
-
-layout(set = 0, binding = 0) uniform CameraViewProj {
-    mat4 ViewProj;
-};
-
-layout(set = 1, binding = 0) uniform Transform {
-    mat4 Model;
-};
-
-void main() {
-    gl_Position = ViewProj * Model * vec4(Vertex_Position, 1.0);
-}

assets/shaders/instancing.wgsl

@@ -0,0 +1,35 @@
+#import bevy_pbr::mesh_view_bind_group
+#import bevy_pbr::mesh_struct
+
+[[group(1), binding(0)]]
+var<uniform> mesh: Mesh;
+
+struct Vertex {
+    [[location(0)]] position: vec3<f32>;
+    [[location(1)]] normal: vec3<f32>;
+    [[location(2)]] uv: vec2<f32>;
+
+    [[location(3)]] i_pos_scale: vec4<f32>;
+    [[location(4)]] i_color: vec4<f32>;
+};
+
+struct VertexOutput {
+    [[builtin(position)]] clip_position: vec4<f32>;
+    [[location(0)]] color: vec4<f32>;
+};
+
+[[stage(vertex)]]
+fn vertex(vertex: Vertex) -> VertexOutput {
+    let position = vertex.position * vertex.i_pos_scale.w + vertex.i_pos_scale.xyz;
+    let world_position = mesh.model * vec4<f32>(position, 1.0);
+
+    var out: VertexOutput;
+    out.clip_position = view.view_proj * world_position;
+    out.color = vertex.i_color;
+    return out;
+}
+
+[[stage(fragment)]]
+fn fragment(in: VertexOutput) -> [[location(0)]] vec4<f32> {
+    return in.color;
+}

examples/README.md

@@ -215,6 +215,10 @@ Example | File | Description
 Example | File | Description
 --- | --- | ---
 `shader_material` | [`shader/shader_material.rs`](./shader/shader_material.rs) | Illustrates creating a custom material and a shader that uses it
+`shader_material_glsl` | [`shader/shader_material_glsl.rs`](./shader/shader_material_glsl.rs) | A custom shader using the GLSL shading language.
+`shader_instancing` | [`shader/shader_instancing.rs`](./shader/shader_instancing.rs) | A custom shader showing off rendering a mesh multiple times in one draw call.
+`animate_shader` | [`shader/animate_shader.rs`](./shader/animate_shader.rs) | Shows how to pass changing data like the time since startup into a shader.
+`compute_shader_game_of_life` | [`shader/compute_shader_game_of_life.rs`](./shader/compute_shader_game_of_life.rs) | A compute shader simulating Conway's Game of Life
 `shader_defs` | [`shader/shader_defs.rs`](./shader/shader_defs.rs) | Demonstrates creating a custom material that uses "shaders defs" (a tool to selectively toggle parts of a shader)
 
 ## Tests

examples/shader/animate_shader.rs

@@ -0,0 +1,249 @@
+use bevy::{
+    core_pipeline::Transparent3d,
+    ecs::system::{lifetimeless::SRes, SystemParamItem},
+    pbr::{
+        DrawMesh, MeshPipeline, MeshPipelineKey, MeshUniform, SetMeshBindGroup,
+        SetMeshViewBindGroup,
+    },
+    prelude::*,
+    render::{
+        render_phase::{
+            AddRenderCommand, DrawFunctions, EntityRenderCommand, RenderCommandResult, RenderPhase,
+            SetItemPipeline, TrackedRenderPass,
+        },
+        render_resource::*,
+        renderer::{RenderDevice, RenderQueue},
+        view::{ComputedVisibility, ExtractedView, Msaa, Visibility},
+        RenderApp, RenderStage,
+    },
+};
+
+fn main() {
+    App::new()
+        .add_plugins(DefaultPlugins)
+        .add_plugin(CustomMaterialPlugin)
+        .add_startup_system(setup)
+        .run();
+}
+
+fn setup(mut commands: Commands, mut meshes: ResMut<Assets<Mesh>>) {
+    // cube
+    commands.spawn().insert_bundle((
+        meshes.add(Mesh::from(shape::Cube { size: 1.0 })),
+        Transform::from_xyz(0.0, 0.5, 0.0),
+        GlobalTransform::default(),
+        CustomMaterial,
+        Visibility::default(),
+        ComputedVisibility::default(),
+    ));
+
+    // camera
+    commands.spawn_bundle(PerspectiveCameraBundle {
+        transform: Transform::from_xyz(-2.0, 2.5, 5.0).looking_at(Vec3::ZERO, Vec3::Y),
+        ..Default::default()
+    });
+}
+
+#[derive(Component)]
+struct CustomMaterial;
+
+pub struct CustomMaterialPlugin;
+
+impl Plugin for CustomMaterialPlugin {
+    fn build(&self, app: &mut App) {
+        let render_device = app.world.get_resource::<RenderDevice>().unwrap();
+        let buffer = render_device.create_buffer(&BufferDescriptor {
+            label: Some("time uniform buffer"),
+            size: std::mem::size_of::<f32>() as u64,
+            usage: BufferUsages::UNIFORM | BufferUsages::COPY_DST,
+            mapped_at_creation: false,
+        });
+
+        app.sub_app_mut(RenderApp)
+            .add_render_command::<Transparent3d, DrawCustom>()
+            .insert_resource(TimeMeta {
+                buffer,
+                bind_group: None,
+            })
+            .init_resource::<CustomPipeline>()
+            .init_resource::<SpecializedPipelines<CustomPipeline>>()
+            .add_system_to_stage(RenderStage::Extract, extract_time)
+            .add_system_to_stage(RenderStage::Extract, extract_custom_material)
+            .add_system_to_stage(RenderStage::Prepare, prepare_time)
+            .add_system_to_stage(RenderStage::Queue, queue_custom)
+            .add_system_to_stage(RenderStage::Queue, queue_time_bind_group);
+    }
+}
+
+// extract the `CustomMaterial` component into the render world
+fn extract_custom_material(
+    mut commands: Commands,
+    mut previous_len: Local<usize>,
+    mut query: Query<Entity, With<CustomMaterial>>,
+) {
+    let mut values = Vec::with_capacity(*previous_len);
+    for entity in query.iter_mut() {
+        values.push((entity, (CustomMaterial,)));
+    }
+    *previous_len = values.len();
+    commands.insert_or_spawn_batch(values);
+}
+
+// add each entity with a mesh and a `CustomMaterial` to every view's `Transparent3d` render phase using the `CustomPipeline`
+fn queue_custom(
+    transparent_3d_draw_functions: Res<DrawFunctions<Transparent3d>>,
+    custom_pipeline: Res<CustomPipeline>,
+    msaa: Res<Msaa>,
+    mut pipelines: ResMut<SpecializedPipelines<CustomPipeline>>,
+    mut pipeline_cache: ResMut<RenderPipelineCache>,
+    material_meshes: Query<(Entity, &MeshUniform), (With<Handle<Mesh>>, With<CustomMaterial>)>,
+    mut views: Query<(&ExtractedView, &mut RenderPhase<Transparent3d>)>,
+) {
+    let draw_custom = transparent_3d_draw_functions
+        .read()
+        .get_id::<DrawCustom>()
+        .unwrap();
+
+    let key = MeshPipelineKey::from_msaa_samples(msaa.samples)
+        | MeshPipelineKey::from_primitive_topology(PrimitiveTopology::TriangleList);
+    let pipeline = pipelines.specialize(&mut pipeline_cache, &custom_pipeline, key);
+
+    for (view, mut transparent_phase) in views.iter_mut() {
+        let view_matrix = view.transform.compute_matrix();
+        let view_row_2 = view_matrix.row(2);
+        for (entity, mesh_uniform) in material_meshes.iter() {
+            transparent_phase.add(Transparent3d {
+                entity,
+                pipeline,
+                draw_function: draw_custom,
+                distance: view_row_2.dot(mesh_uniform.transform.col(3)),
+            });
+        }
+    }
+}
+
+#[derive(Default)]
+struct ExtractedTime {
+    seconds_since_startup: f32,
+}
+
+// extract the passed time into a resource in the render world
+fn extract_time(mut commands: Commands, time: Res<Time>) {
+    commands.insert_resource(ExtractedTime {
+        seconds_since_startup: time.seconds_since_startup() as f32,
+    });
+}
+
+struct TimeMeta {
+    buffer: Buffer,
+    bind_group: Option<BindGroup>,
+}
+
+// write the extracted time into the corresponding uniform buffer
+fn prepare_time(
+    time: Res<ExtractedTime>,
+    time_meta: ResMut<TimeMeta>,
+    render_queue: Res<RenderQueue>,
+) {
+    render_queue.write_buffer(
+        &time_meta.buffer,
+        0,
+        bevy::core::cast_slice(&[time.seconds_since_startup]),
+    );
+}
+
+// create a bind group for the time uniform buffer
+fn queue_time_bind_group(
+    render_device: Res<RenderDevice>,
+    mut time_meta: ResMut<TimeMeta>,
+    pipeline: Res<CustomPipeline>,
+) {
+    let bind_group = render_device.create_bind_group(&BindGroupDescriptor {
+        label: None,
+        layout: &pipeline.time_bind_group_layout,
+        entries: &[BindGroupEntry {
+            binding: 0,
+            resource: time_meta.buffer.as_entire_binding(),
+        }],
+    });
+    time_meta.bind_group = Some(bind_group);
+}
+
+pub struct CustomPipeline {
+    shader: Handle<Shader>,
+    mesh_pipeline: MeshPipeline,
+    time_bind_group_layout: BindGroupLayout,
+}
+
+impl FromWorld for CustomPipeline {
+    fn from_world(world: &mut World) -> Self {
+        let world = world.cell();
+        let asset_server = world.get_resource::<AssetServer>().unwrap();
+        let shader = asset_server.load("shaders/animate_shader.wgsl");
+
+        let render_device = world.get_resource_mut::<RenderDevice>().unwrap();
+        let time_bind_group_layout =
+            render_device.create_bind_group_layout(&BindGroupLayoutDescriptor {
+                label: Some("time bind group"),
+                entries: &[BindGroupLayoutEntry {
+                    binding: 0,
+                    visibility: ShaderStages::FRAGMENT,
+                    ty: BindingType::Buffer {
+                        ty: BufferBindingType::Uniform,
+                        has_dynamic_offset: false,
+                        min_binding_size: BufferSize::new(std::mem::size_of::<f32>() as u64),
+                    },
+                    count: None,
+                }],
+            });
+
+        let mesh_pipeline = world.get_resource::<MeshPipeline>().unwrap();
+
+        CustomPipeline {
+            shader,
+            mesh_pipeline: mesh_pipeline.clone(),
+            time_bind_group_layout,
+        }
+    }
+}
+
+impl SpecializedPipeline for CustomPipeline {
+    type Key = MeshPipelineKey;
+
+    fn specialize(&self, key: Self::Key) -> RenderPipelineDescriptor {
+        let mut descriptor = self.mesh_pipeline.specialize(key);
+        descriptor.vertex.shader = self.shader.clone();
+        descriptor.fragment.as_mut().unwrap().shader = self.shader.clone();
+        descriptor.layout = Some(vec![
+            self.mesh_pipeline.view_layout.clone(),
+            self.mesh_pipeline.mesh_layout.clone(),
+            self.time_bind_group_layout.clone(),
+        ]);
+        descriptor
+    }
+}
+
+type DrawCustom = (
+    SetItemPipeline,
+    SetMeshViewBindGroup<0>,
+    SetMeshBindGroup<1>,
+    SetTimeBindGroup<2>,
+    DrawMesh,
+);
+
+struct SetTimeBindGroup<const I: usize>;
+impl<const I: usize> EntityRenderCommand for SetTimeBindGroup<I> {
+    type Param = SRes<TimeMeta>;
+
+    fn render<'w>(
+        _view: Entity,
+        _item: Entity,
+        time_meta: SystemParamItem<'w, '_, Self::Param>,
+        pass: &mut TrackedRenderPass<'w>,
+    ) -> RenderCommandResult {
+        let time_bind_group = time_meta.into_inner().bind_group.as_ref().unwrap();
+        pass.set_bind_group(I, time_bind_group, &[]);
+
+        RenderCommandResult::Success
+    }
+}

examples/shader/compute_shader_game_of_life.rs

@@ -0,0 +1,209 @@
+use bevy::{
+    core_pipeline::node::MAIN_PASS_DEPENDENCIES,
+    prelude::*,
+    render::{
+        render_asset::RenderAssets,
+        render_graph::{self, RenderGraph},
+        render_resource::*,
+        renderer::{RenderContext, RenderDevice},
+        RenderApp, RenderStage,
+    },
+    window::WindowDescriptor,
+};
+
+const SIZE: (u32, u32) = (1280, 720);
+const WORKGROUP_SIZE: u32 = 8;
+
+fn main() {
+    App::new()
+        .insert_resource(ClearColor(Color::BLACK))
+        .insert_resource(WindowDescriptor {
+            // uncomment for unthrottled FPS
+            // vsync: false,
+            ..Default::default()
+        })
+        .add_plugins(DefaultPlugins)
+        .add_plugin(GameOfLifeComputePlugin)
+        .add_startup_system(setup)
+        .run();
+}
+
+fn setup(mut commands: Commands, mut images: ResMut<Assets<Image>>) {
+    let mut image = Image::new_fill(
+        Extent3d {
+            width: SIZE.0,
+            height: SIZE.1,
+            depth_or_array_layers: 1,
+        },
+        TextureDimension::D2,
+        &[0, 0, 0, 255],
+        TextureFormat::Rgba8Unorm,
+    );
+    image.texture_descriptor.usage =
+        TextureUsages::COPY_DST | TextureUsages::STORAGE_BINDING | TextureUsages::TEXTURE_BINDING;
+    let image = images.add(image);
+
+    commands.spawn_bundle(SpriteBundle {
+        sprite: Sprite {
+            custom_size: Some(Vec2::new(SIZE.0 as f32, SIZE.1 as f32)),
+            ..Default::default()
+        },
+        texture: image.clone(),
+        ..Default::default()
+    });
+    commands.spawn_bundle(OrthographicCameraBundle::new_2d());
+
+    commands.insert_resource(GameOfLifeImage(image))
+}
+
+pub struct GameOfLifeComputePlugin;
+
+impl Plugin for GameOfLifeComputePlugin {
+    fn build(&self, app: &mut App) {
+        let render_app = app.sub_app_mut(RenderApp);
+        render_app
+            .init_resource::<GameOfLifePipeline>()
+            .add_system_to_stage(RenderStage::Extract, extract_game_of_life_image)
+            .add_system_to_stage(RenderStage::Queue, queue_bind_group);
+
+        let mut render_graph = render_app.world.get_resource_mut::<RenderGraph>().unwrap();
+        render_graph.add_node("game_of_life", DispatchGameOfLife::default());
+        render_graph
+            .add_node_edge("game_of_life", MAIN_PASS_DEPENDENCIES)
+            .unwrap();
+    }
+}
+
+struct GameOfLifeImage(Handle<Image>);
+struct GameOfLifeImageBindGroup(BindGroup);
+
+fn extract_game_of_life_image(mut commands: Commands, image: Res<GameOfLifeImage>) {
+    commands.insert_resource(GameOfLifeImage(image.0.clone()));
+}
+fn queue_bind_group(
+    mut commands: Commands,
+    pipeline: Res<GameOfLifePipeline>,
+    gpu_images: Res<RenderAssets<Image>>,
+    game_of_life_image: Res<GameOfLifeImage>,
+    render_device: Res<RenderDevice>,
+) {
+    let view = &gpu_images[&game_of_life_image.0];
+    let bind_group = render_device.create_bind_group(&BindGroupDescriptor {
+        label: None,
+        layout: &pipeline.texture_bind_group_layout,
+        entries: &[BindGroupEntry {
+            binding: 0,
+            resource: BindingResource::TextureView(&view.texture_view),
+        }],
+    });
+    commands.insert_resource(GameOfLifeImageBindGroup(bind_group));
+}
+
+pub struct GameOfLifePipeline {
+    sim_pipeline: ComputePipeline,
+    init_pipeline: ComputePipeline,
+    texture_bind_group_layout: BindGroupLayout,
+}
+
+impl FromWorld for GameOfLifePipeline {
+    fn from_world(world: &mut World) -> Self {
+        let render_device = world.get_resource::<RenderDevice>().unwrap();
+
+        let shader_source = include_str!("../../assets/shaders/game_of_life.wgsl");
+        let shader = render_device.create_shader_module(&ShaderModuleDescriptor {
+            label: None,
+            source: ShaderSource::Wgsl(shader_source.into()),
+        });
+
+        let texture_bind_group_layout =
+            render_device.create_bind_group_layout(&BindGroupLayoutDescriptor {
+                label: None,
+                entries: &[BindGroupLayoutEntry {
+                    binding: 0,
+                    visibility: ShaderStages::COMPUTE,
+                    ty: BindingType::StorageTexture {
+                        access: StorageTextureAccess::ReadWrite,
+                        format: TextureFormat::Rgba8Unorm,
+                        view_dimension: TextureViewDimension::D2,
+                    },
+                    count: None,
+                }],
+            });
+
+        let pipeline_layout = render_device.create_pipeline_layout(&PipelineLayoutDescriptor {
+            label: None,
+            bind_group_layouts: &[&texture_bind_group_layout],
+            push_constant_ranges: &[],
+        });
+        let init_pipeline = render_device.create_compute_pipeline(&ComputePipelineDescriptor {
+            label: None,
+            layout: Some(&pipeline_layout),
+            module: &shader,
+            entry_point: "init",
+        });
+        let sim_pipeline = render_device.create_compute_pipeline(&ComputePipelineDescriptor {
+            label: None,
+            layout: Some(&pipeline_layout),
+            module: &shader,
+            entry_point: "update",
+        });
+
+        GameOfLifePipeline {
+            sim_pipeline,
+            init_pipeline,
+            texture_bind_group_layout,
+        }
+    }
+}
+
+enum Initialized {
+    Default,
+    No,
+    Yes,
+}
+
+struct DispatchGameOfLife {
+    initialized: Initialized,
+}
+impl Default for DispatchGameOfLife {
+    fn default() -> Self {
+        Self {
+            initialized: Initialized::Default,
+        }
+    }
+}
+impl render_graph::Node for DispatchGameOfLife {
+    fn update(&mut self, _world: &mut World) {
+        match self.initialized {
+            Initialized::Default => self.initialized = Initialized::No,
+            Initialized::No => self.initialized = Initialized::Yes,
+            Initialized::Yes => {}
+        }
+    }
+
+    fn run(
+        &self,
+        _graph: &mut render_graph::RenderGraphContext,
+        render_context: &mut RenderContext,
+        world: &World,
+    ) -> Result<(), render_graph::NodeRunError> {
+        let pipeline = world.get_resource::<GameOfLifePipeline>().unwrap();
+        let texture_bind_group = &world.get_resource::<GameOfLifeImageBindGroup>().unwrap().0;
+
+        let mut pass = render_context
+            .command_encoder
+            .begin_compute_pass(&ComputePassDescriptor::default());
+
+        if let Initialized::No = self.initialized {
+            pass.set_pipeline(&pipeline.init_pipeline);
+            pass.set_bind_group(0, texture_bind_group, &[]);
+            pass.dispatch(SIZE.0 / WORKGROUP_SIZE, SIZE.1 / WORKGROUP_SIZE, 1);
+        }
+
+        pass.set_pipeline(&pipeline.sim_pipeline);
+        pass.set_bind_group(0, texture_bind_group, &[]);
+        pass.dispatch(SIZE.0 / WORKGROUP_SIZE, SIZE.1 / WORKGROUP_SIZE, 1);
+
+        Ok(())
+    }
+}

examples/shader/shader_instancing.rs

@@ -0,0 +1,250 @@
+use bevy::{
+    core_pipeline::Transparent3d,
+    ecs::system::{lifetimeless::*, SystemParamItem},
+    math::prelude::*,
+    pbr::{MeshPipeline, MeshPipelineKey, MeshUniform, SetMeshBindGroup, SetMeshViewBindGroup},
+    prelude::*,
+    render::{
+        mesh::GpuBufferInfo,
+        render_asset::RenderAssets,
+        render_component::{ExtractComponent, ExtractComponentPlugin},
+        render_phase::{
+            AddRenderCommand, DrawFunctions, EntityRenderCommand, RenderCommandResult, RenderPhase,
+            SetItemPipeline, TrackedRenderPass,
+        },
+        render_resource::*,
+        renderer::RenderDevice,
+        view::{ComputedVisibility, ExtractedView, Msaa, Visibility},
+        RenderApp, RenderStage,
+    },
+};
+use bytemuck::{Pod, Zeroable};
+
+fn main() {
+    App::new()
+        .add_plugins(DefaultPlugins)
+        .add_plugin(CustomMaterialPlugin)
+        .add_startup_system(setup)
+        .run();
+}
+
+fn setup(mut commands: Commands, mut meshes: ResMut<Assets<Mesh>>) {
+    commands.spawn().insert_bundle((
+        meshes.add(Mesh::from(shape::Cube { size: 0.5 })),
+        Transform::from_xyz(0.0, 0.0, 0.0),
+        GlobalTransform::default(),
+        InstanceMaterialData(
+            (1..=10)
+                .flat_map(|x| (1..=10).map(move |y| (x as f32 / 10.0, y as f32 / 10.0)))
+                .map(|(x, y)| InstanceData {
+                    position: Vec3::new(x * 10.0 - 5.0, y * 10.0 - 5.0, 0.0),
+                    scale: 1.0,
+                    color: Color::hsla(x * 360., y, 0.5, 1.0).as_rgba_f32(),
+                })
+                .collect(),
+        ),
+        Visibility::default(),
+        ComputedVisibility::default(),
+    ));
+
+    // camera
+    commands.spawn_bundle(PerspectiveCameraBundle {
+        transform: Transform::from_xyz(0.0, 0.0, 15.0).looking_at(Vec3::ZERO, Vec3::Y),
+        ..Default::default()
+    });
+}
+
+#[derive(Component)]
+struct InstanceMaterialData(Vec<InstanceData>);
+impl ExtractComponent for InstanceMaterialData {
+    type Query = &'static InstanceMaterialData;
+    type Filter = ();
+
+    fn extract_component(item: bevy::ecs::query::QueryItem<Self::Query>) -> Self {
+        InstanceMaterialData(item.0.clone())
+    }
+}
+
+pub struct CustomMaterialPlugin;
+
+impl Plugin for CustomMaterialPlugin {
+    fn build(&self, app: &mut App) {
+        app.add_plugin(ExtractComponentPlugin::<InstanceMaterialData>::default());
+        app.sub_app_mut(RenderApp)
+            .add_render_command::<Transparent3d, DrawCustom>()
+            .init_resource::<CustomPipeline>()
+            .init_resource::<SpecializedPipelines<CustomPipeline>>()
+            .add_system_to_stage(RenderStage::Queue, queue_custom)
+            .add_system_to_stage(RenderStage::Prepare, prepare_instance_buffers);
+    }
+}
+
+#[derive(Clone, Copy, Pod, Zeroable)]
+#[repr(C)]
+struct InstanceData {
+    position: Vec3,
+    scale: f32,
+    color: [f32; 4],
+}
+
+fn queue_custom(
+    transparent_3d_draw_functions: Res<DrawFunctions<Transparent3d>>,
+    custom_pipeline: Res<CustomPipeline>,
+    msaa: Res<Msaa>,
+    mut pipelines: ResMut<SpecializedPipelines<CustomPipeline>>,
+    mut pipeline_cache: ResMut<RenderPipelineCache>,
+    material_meshes: Query<
+        (Entity, &MeshUniform),
+        (With<Handle<Mesh>>, With<InstanceMaterialData>),
+    >,
+    mut views: Query<(&ExtractedView, &mut RenderPhase<Transparent3d>)>,
+) {
+    let draw_custom = transparent_3d_draw_functions
+        .read()
+        .get_id::<DrawCustom>()
+        .unwrap();
+
+    let key = MeshPipelineKey::from_msaa_samples(msaa.samples)
+        | MeshPipelineKey::from_primitive_topology(PrimitiveTopology::TriangleList);
+    let pipeline = pipelines.specialize(&mut pipeline_cache, &custom_pipeline, key);
+
+    for (view, mut transparent_phase) in views.iter_mut() {
+        let view_matrix = view.transform.compute_matrix();
+        let view_row_2 = view_matrix.row(2);
+        for (entity, mesh_uniform) in material_meshes.iter() {
+            transparent_phase.add(Transparent3d {
+                entity,
+                pipeline,
+                draw_function: draw_custom,
+                distance: view_row_2.dot(mesh_uniform.transform.col(3)),
+            });
+        }
+    }
+}
+
+#[derive(Component)]
+pub struct InstanceBuffer {
+    buffer: Buffer,
+    length: usize,
+}
+
+fn prepare_instance_buffers(
+    mut commands: Commands,
+    query: Query<(Entity, &InstanceMaterialData)>,
+    render_device: Res<RenderDevice>,
+) {
+    for (entity, instance_data) in query.iter() {
+        let buffer = render_device.create_buffer_with_data(&BufferInitDescriptor {
+            label: Some("instance data buffer"),
+            contents: bytemuck::cast_slice(instance_data.0.as_slice()),
+            usage: BufferUsages::VERTEX | BufferUsages::COPY_DST,
+        });
+        commands.entity(entity).insert(InstanceBuffer {
+            buffer,
+            length: instance_data.0.len(),
+        });
+    }
+}
+
+pub struct CustomPipeline {
+    shader: Handle<Shader>,
+    mesh_pipeline: MeshPipeline,
+}
+
+impl FromWorld for CustomPipeline {
+    fn from_world(world: &mut World) -> Self {
+        let world = world.cell();
+        let asset_server = world.get_resource::<AssetServer>().unwrap();
+        asset_server.watch_for_changes().unwrap();
+        let shader = asset_server.load("shaders/instancing.wgsl");
+
+        let mesh_pipeline = world.get_resource::<MeshPipeline>().unwrap();
+
+        CustomPipeline {
+            shader,
+            mesh_pipeline: mesh_pipeline.clone(),
+        }
+    }
+}
+
+impl SpecializedPipeline for CustomPipeline {
+    type Key = MeshPipelineKey;
+
+    fn specialize(&self, key: Self::Key) -> RenderPipelineDescriptor {
+        let mut descriptor = self.mesh_pipeline.specialize(key);
+        descriptor.vertex.shader = self.shader.clone();
+        descriptor.vertex.buffers.push(VertexBufferLayout {
+            array_stride: std::mem::size_of::<InstanceData>() as u64,
+            step_mode: VertexStepMode::Instance,
+            attributes: vec![
+                VertexAttribute {
+                    format: VertexFormat::Float32x4,
+                    offset: 0,
+                    shader_location: 3, // shader locations 0-2 are taken up by Position, Normal and UV attributes
+                },
+                VertexAttribute {
+                    format: VertexFormat::Float32x4,
+                    offset: VertexFormat::Float32x4.size(),
+                    shader_location: 4,
+                },
+            ],
+        });
+        descriptor.fragment.as_mut().unwrap().shader = self.shader.clone();
+        descriptor.layout = Some(vec![
+            self.mesh_pipeline.view_layout.clone(),
+            self.mesh_pipeline.mesh_layout.clone(),
+        ]);
+
+        descriptor
+    }
+}
+
+type DrawCustom = (
+    SetItemPipeline,
+    SetMeshViewBindGroup<0>,
+    SetMeshBindGroup<1>,
+    DrawMeshInstanced,
+);
+
+pub struct DrawMeshInstanced;
+impl EntityRenderCommand for DrawMeshInstanced {
+    type Param = (
+        SRes<RenderAssets<Mesh>>,
+        SQuery<Read<Handle<Mesh>>>,
+        SQuery<Read<InstanceBuffer>>,
+    );
+    #[inline]
+    fn render<'w>(
+        _view: Entity,
+        item: Entity,
+        (meshes, mesh_query, instance_buffer_query): SystemParamItem<'w, '_, Self::Param>,
+        pass: &mut TrackedRenderPass<'w>,
+    ) -> RenderCommandResult {
+        let mesh_handle = mesh_query.get(item).unwrap();
+        let instance_buffer = instance_buffer_query.get(item).unwrap();
+
+        let gpu_mesh = match meshes.into_inner().get(mesh_handle) {
+            Some(gpu_mesh) => gpu_mesh,
+            None => return RenderCommandResult::Failure,
+        };
+
+        pass.set_vertex_buffer(0, gpu_mesh.vertex_buffer.slice(..));
+        pass.set_vertex_buffer(1, instance_buffer.buffer.slice(..));
+
+        pass.set_vertex_buffer(0, gpu_mesh.vertex_buffer.slice(..));
+        match &gpu_mesh.buffer_info {
+            GpuBufferInfo::Indexed {
+                buffer,
+                index_format,
+                count,
+            } => {
+                pass.set_index_buffer(buffer.slice(..), 0, *index_format);
+                pass.draw_indexed(0..*count, 0, 0..instance_buffer.length as u32);
+            }
+            GpuBufferInfo::NonIndexed { vertex_count } => {
+                pass.draw_indexed(0..*vertex_count, 0, 0..instance_buffer.length as u32);
+            }
+        }
+        RenderCommandResult::Success
+    }
+}

examples/shader/shader_material_glsl.rs

@@ -0,0 +1,130 @@
+use bevy::{
+    ecs::system::{lifetimeless::SRes, SystemParamItem},
+    pbr::{MaterialPipeline, SpecializedMaterial},
+    prelude::*,
+    reflect::TypeUuid,
+    render::{
+        render_asset::{PrepareAssetError, RenderAsset},
+        render_resource::{
+            std140::{AsStd140, Std140},
+            *,
+        },
+        renderer::RenderDevice,
+    },
+};
+
+fn main() {
+    App::new()
+        .add_plugins(DefaultPlugins)
+        .add_plugin(MaterialPlugin::<CustomMaterial>::default())
+        .add_startup_system(setup)
+        .run();
+}
+
+/// set up a simple 3D scene
+fn setup(
+    mut commands: Commands,
+    mut meshes: ResMut<Assets<Mesh>>,
+    mut materials: ResMut<Assets<CustomMaterial>>,
+) {
+    // cube
+    commands.spawn().insert_bundle(MaterialMeshBundle {
+        mesh: meshes.add(Mesh::from(shape::Cube { size: 1.0 })),
+        transform: Transform::from_xyz(0.0, 0.5, 0.0),
+        material: materials.add(CustomMaterial {
+            color: Color::GREEN,
+        }),
+        ..Default::default()
+    });
+
+    // camera
+    commands.spawn_bundle(PerspectiveCameraBundle {
+        transform: Transform::from_xyz(-2.0, 2.5, 5.0).looking_at(Vec3::ZERO, Vec3::Y),
+        ..Default::default()
+    });
+}
+
+#[derive(Debug, Clone, TypeUuid)]
+#[uuid = "4ee9c363-1124-4113-890e-199d81b00281"]
+pub struct CustomMaterial {
+    color: Color,
+}
+
+#[derive(Clone)]
+pub struct GpuCustomMaterial {
+    _buffer: Buffer,
+    bind_group: BindGroup,
+}
+
+impl RenderAsset for CustomMaterial {
+    type ExtractedAsset = CustomMaterial;
+    type PreparedAsset = GpuCustomMaterial;
+    type Param = (SRes<RenderDevice>, SRes<MaterialPipeline<Self>>);
+    fn extract_asset(&self) -> Self::ExtractedAsset {
+        self.clone()
+    }
+
+    fn prepare_asset(
+        extracted_asset: Self::ExtractedAsset,
+        (render_device, material_pipeline): &mut SystemParamItem<Self::Param>,
+    ) -> Result<Self::PreparedAsset, PrepareAssetError<Self::ExtractedAsset>> {
+        let color = Vec4::from_slice(&extracted_asset.color.as_linear_rgba_f32());
+        let buffer = render_device.create_buffer_with_data(&BufferInitDescriptor {
+            contents: color.as_std140().as_bytes(),
+            label: None,
+            usage: BufferUsages::UNIFORM | BufferUsages::COPY_DST,
+        });
+        let bind_group = render_device.create_bind_group(&BindGroupDescriptor {
+            entries: &[BindGroupEntry {
+                binding: 0,
+                resource: buffer.as_entire_binding(),
+            }],
+            label: None,
+            layout: &material_pipeline.material_layout,
+        });
+
+        Ok(GpuCustomMaterial {
+            _buffer: buffer,
+            bind_group,
+        })
+    }
+}
+
+impl SpecializedMaterial for CustomMaterial {
+    type Key = ();
+
+    fn key(_: &<CustomMaterial as RenderAsset>::PreparedAsset) -> Self::Key {}
+
+    fn specialize(_: Self::Key, descriptor: &mut RenderPipelineDescriptor) {
+        descriptor.vertex.entry_point = "main".into();
+        descriptor.fragment.as_mut().unwrap().entry_point = "main".into();
+    }
+
+    fn vertex_shader(asset_server: &AssetServer) -> Option<Handle<Shader>> {
+        Some(asset_server.load("shaders/custom_material.vert"))
+    }
+
+    fn fragment_shader(asset_server: &AssetServer) -> Option<Handle<Shader>> {
+        Some(asset_server.load("shaders/custom_material.frag"))
+    }
+
+    fn bind_group(render_asset: &<Self as RenderAsset>::PreparedAsset) -> &BindGroup {
+        &render_asset.bind_group
+    }
+
+    fn bind_group_layout(render_device: &RenderDevice) -> BindGroupLayout {
+        render_device.create_bind_group_layout(&BindGroupLayoutDescriptor {
+            entries: &[BindGroupLayoutEntry {
+                binding: 0,
+                visibility: ShaderStages::FRAGMENT,
+                ty: BindingType::Buffer {
+                    ty: BufferBindingType::Uniform,
+                    has_dynamic_offset: false,
+                    min_binding_size: BufferSize::new(Vec4::std140_size_static() as u64),
+                },
+                count: None,
+            }],
+            label: None,
+        })
+    }
+}