diff --git a/blade-graphics/Cargo.toml b/blade-graphics/Cargo.toml
index 1144c9ce..1241c63c 100644
--- a/blade-graphics/Cargo.toml
+++ b/blade-graphics/Cargo.toml
@@ -9,6 +9,11 @@ repository = "https://github.com/kvark/blade"
 
 [lib]
 
+[features]
+default = ["gles"]
+gles = ["dep:glow", "dep:egl", "dep:libloading", "naga/glsl-out"]
+derive = ["dep:blade-macros"]
+
 [dependencies]
 bitflags = { workspace = true }
 bytemuck = { workspace = true }
@@ -18,6 +23,8 @@ log = { workspace = true }
 mint = { workspace = true }
 naga = { workspace = true }
 raw-window-handle = "0.6"
+glow = { version = "0.13", optional = true }
+blade-macros = { "path" = "../blade-macros", optional = true }
 
 [target.'cfg(any(target_os = "ios", target_os = "macos"))'.dependencies]
 block = "0.1"
@@ -34,13 +41,11 @@ gpu-alloc-ash = "0.7"
 naga = { workspace = true, features = ["spv-out"] }
 slab = { workspace = true }
 
-[target.'cfg(any(gles, target_arch = "wasm32"))'.dependencies]
-glow = "0.13"
-naga = { workspace = true, features = ["glsl-out"] }
-
-[target.'cfg(all(gles, not(target_arch = "wasm32")))'.dependencies]
-egl = { package = "khronos-egl", version = "5.0", features = ["dynamic"] }
-libloading = { version = "0.8" }
+[target.'cfg(not(target_arch = "wasm32"))'.dependencies]
+egl = { package = "khronos-egl", version = "5.0", features = [
+    "dynamic",
+], optional = true }
+libloading = { version = "0.8", optional = true }
 
 [target.'cfg(all(target_arch = "wasm32"))'.dependencies]
 wasm-bindgen = "0.2.83"
diff --git a/blade-graphics/src/common/derive.rs b/blade-graphics/src/common/derive.rs
new file mode 100644
index 00000000..e86bafc3
--- /dev/null
+++ b/blade-graphics/src/common/derive.rs
@@ -0,0 +1,20 @@
+use crate::derive::HasShaderBinding;
+
+impl HasShaderBinding for super::TextureView {
+    const TYPE: crate::ShaderBinding = crate::ShaderBinding::Texture;
+}
+impl HasShaderBinding for super::Sampler {
+    const TYPE: crate::ShaderBinding = crate::ShaderBinding::Sampler;
+}
+impl HasShaderBinding for super::BufferPiece {
+    const TYPE: crate::ShaderBinding = crate::ShaderBinding::Buffer;
+}
+impl<'a, const N: crate::ResourceIndex> HasShaderBinding for &'a super::BufferArray<N> {
+    const TYPE: crate::ShaderBinding = crate::ShaderBinding::BufferArray { count: N };
+}
+impl<'a, const N: crate::ResourceIndex> HasShaderBinding for &'a super::TextureArray<N> {
+    const TYPE: crate::ShaderBinding = crate::ShaderBinding::TextureArray { count: N };
+}
+impl HasShaderBinding for super::AccelerationStructure {
+    const TYPE: crate::ShaderBinding = crate::ShaderBinding::AccelerationStructure;
+}
diff --git a/blade-graphics/src/common/lib.rs b/blade-graphics/src/common/lib.rs
new file mode 100644
index 00000000..0ee031c3
--- /dev/null
+++ b/blade-graphics/src/common/lib.rs
@@ -0,0 +1,107 @@
+pub type BufferPiece = crate::GenericBufferPiece<Buffer>;
+
+#[derive(Clone, Copy, Debug)]
+pub struct TexturePiece {
+    pub texture: Texture,
+    pub mip_level: u32,
+    pub array_layer: u32,
+    pub origin: [u32; 3],
+}
+
+impl From<Texture> for TexturePiece {
+    fn from(texture: Texture) -> Self {
+        Self {
+            texture,
+            mip_level: 0,
+            array_layer: 0,
+            origin: [0; 3],
+        }
+    }
+}
+
+pub type BufferArray<const N: crate::ResourceIndex> = crate::ResourceArray<BufferPiece, N>;
+pub type TextureArray<const N: crate::ResourceIndex> = crate::ResourceArray<TextureView, N>;
+
+#[derive(Clone, Debug)]
+pub struct AccelerationStructureMesh {
+    pub vertex_data: BufferPiece,
+    pub vertex_format: crate::VertexFormat,
+    pub vertex_stride: u32,
+    pub vertex_count: u32,
+    pub index_data: BufferPiece,
+    pub index_type: Option<crate::IndexType>,
+    pub triangle_count: u32,
+    pub transform_data: BufferPiece,
+    pub is_opaque: bool,
+}
+
+pub trait ShaderBindable: Clone + Copy + crate::derive::HasShaderBinding {
+    fn bind_to(&self, context: &mut PipelineContext, index: u32);
+}
+
+#[derive(Clone, Copy, Debug)]
+pub enum FinishOp {
+    Store,
+    Discard,
+    ResolveTo(TextureView),
+    Ignore,
+}
+
+#[derive(Debug)]
+pub struct RenderTarget {
+    pub view: TextureView,
+    pub init_op: crate::InitOp,
+    pub finish_op: FinishOp,
+}
+
+#[derive(Debug)]
+pub struct RenderTargetSet<'a> {
+    pub colors: &'a [RenderTarget],
+    pub depth_stencil: Option<RenderTarget>,
+}
+
+impl Context {
+    pub fn try_create_shader(
+        &self,
+        desc: super::ShaderDesc,
+    ) -> Result<super::Shader, &'static str> {
+        let module = naga::front::wgsl::parse_str(desc.source).map_err(|e| {
+            e.emit_to_stderr_with_path(desc.source, "");
+            "compilation failed"
+        })?;
+
+        let device_caps = self.capabilities();
+
+        // Bindings are set up at pipeline creation, ignore here
+        let flags = naga::valid::ValidationFlags::all() ^ naga::valid::ValidationFlags::BINDINGS;
+        let mut caps = naga::valid::Capabilities::empty();
+        caps.set(
+            naga::valid::Capabilities::RAY_QUERY | naga::valid::Capabilities::SAMPLED_TEXTURE_AND_STORAGE_BUFFER_ARRAY_NON_UNIFORM_INDEXING,
+            !device_caps.ray_query.is_empty(),
+        );
+        let info = naga::valid::Validator::new(flags, caps)
+            .validate(&module)
+            .map_err(|e| {
+                crate::util::emit_annotated_error(&e, "", desc.source);
+                crate::util::print_err(&e);
+                "validation failed"
+            })?;
+
+        Ok(super::Shader { module, info })
+    }
+
+    pub fn create_shader(&self, desc: super::ShaderDesc) -> super::Shader {
+        self.try_create_shader(desc).unwrap()
+    }
+}
+
+pub trait ShaderData {
+    fn layout() -> crate::ShaderDataLayout;
+    fn fill(&self, context: PipelineContext);
+}
+
+mod traits {
+    pub trait PipelineEncoder {
+        fn bind<D: super::ShaderData>(&mut self, group: u32, data: &D);
+    }
+}
diff --git a/blade-graphics/src/common/util.rs b/blade-graphics/src/common/util.rs
new file mode 100644
index 00000000..37e6bd0c
--- /dev/null
+++ b/blade-graphics/src/common/util.rs
@@ -0,0 +1,11 @@
+impl super::ComputePipeline {
+    /// Return the dispatch group counts sufficient to cover the given extent.
+    pub fn get_dispatch_for(&self, extent: crate::Extent) -> [u32; 3] {
+        let wg_size = self.get_workgroup_size();
+        [
+            (extent.width + wg_size[0] - 1) / wg_size[0],
+            (extent.height + wg_size[1] - 1) / wg_size[1],
+            (extent.depth + wg_size[2] - 1) / wg_size[2],
+        ]
+    }
+}
diff --git a/blade-graphics/src/derive.rs b/blade-graphics/src/derive.rs
index 6fabfb31..5352dd3a 100644
--- a/blade-graphics/src/derive.rs
+++ b/blade-graphics/src/derive.rs
@@ -10,24 +10,6 @@ impl<T: bytemuck::Pod> HasShaderBinding for T {
         size: mem::size_of::<T>() as u32,
     };
 }
-impl HasShaderBinding for super::TextureView {
-    const TYPE: ShaderBinding = ShaderBinding::Texture;
-}
-impl HasShaderBinding for super::Sampler {
-    const TYPE: ShaderBinding = ShaderBinding::Sampler;
-}
-impl HasShaderBinding for super::BufferPiece {
-    const TYPE: ShaderBinding = ShaderBinding::Buffer;
-}
-impl<'a, const N: ResourceIndex> HasShaderBinding for &'a super::BufferArray<N> {
-    const TYPE: ShaderBinding = ShaderBinding::BufferArray { count: N };
-}
-impl<'a, const N: ResourceIndex> HasShaderBinding for &'a super::TextureArray<N> {
-    const TYPE: ShaderBinding = ShaderBinding::TextureArray { count: N };
-}
-impl HasShaderBinding for super::AccelerationStructure {
-    const TYPE: ShaderBinding = ShaderBinding::AccelerationStructure;
-}
 
 pub trait HasVertexAttribute {
     const FORMAT: VertexFormat;
diff --git a/blade-graphics/src/gles/command.rs b/blade-graphics/src/gles/command.rs
index 0e2d6470..a1ebb32b 100644
--- a/blade-graphics/src/gles/command.rs
+++ b/blade-graphics/src/gles/command.rs
@@ -5,7 +5,7 @@ const COLOR_ATTACHMENTS: &[u32] = &[
     glow::COLOR_ATTACHMENT3,
 ];
 
-impl<T: bytemuck::Pod> crate::ShaderBindable for T {
+impl<T: bytemuck::Pod> super::ShaderBindable for T {
     fn bind_to(&self, ctx: &mut super::PipelineContext, index: u32) {
         let self_slice = bytemuck::bytes_of(self);
         let alignment = ctx.limits.uniform_buffer_alignment as usize;
@@ -26,7 +26,7 @@ impl<T: bytemuck::Pod> crate::ShaderBindable for T {
         }
     }
 }
-impl crate::ShaderBindable for super::TextureView {
+impl super::ShaderBindable for super::TextureView {
     fn bind_to(&self, ctx: &mut super::PipelineContext, index: u32) {
         let (texture, target) = self.inner.as_native();
         for &slot in ctx.targets[index as usize].iter() {
@@ -38,12 +38,12 @@ impl crate::ShaderBindable for super::TextureView {
         }
     }
 }
-impl<'a, const N: crate::ResourceIndex> crate::ShaderBindable for &'a crate::TextureArray<N> {
+impl<'a, const N: crate::ResourceIndex> super::ShaderBindable for &'a super::TextureArray<N> {
     fn bind_to(&self, _ctx: &mut super::PipelineContext, _index: u32) {
         unimplemented!()
     }
 }
-impl crate::ShaderBindable for super::Sampler {
+impl super::ShaderBindable for super::Sampler {
     fn bind_to(&self, ctx: &mut super::PipelineContext, index: u32) {
         for &slot in ctx.targets[index as usize].iter() {
             ctx.commands.push(super::Command::BindSampler {
@@ -53,7 +53,7 @@ impl crate::ShaderBindable for super::Sampler {
         }
     }
 }
-impl crate::ShaderBindable for crate::BufferPiece {
+impl super::ShaderBindable for super::BufferPiece {
     fn bind_to(&self, ctx: &mut super::PipelineContext, index: u32) {
         for &slot in ctx.targets[index as usize].iter() {
             ctx.commands.push(super::Command::BindBuffer {
@@ -65,12 +65,12 @@ impl crate::ShaderBindable for crate::BufferPiece {
         }
     }
 }
-impl<'a, const N: crate::ResourceIndex> crate::ShaderBindable for &'a crate::BufferArray<N> {
+impl<'a, const N: crate::ResourceIndex> super::ShaderBindable for &'a super::BufferArray<N> {
     fn bind_to(&self, _ctx: &mut super::PipelineContext, _index: u32) {
         unimplemented!()
     }
 }
-impl crate::ShaderBindable for super::AccelerationStructure {
+impl super::ShaderBindable for super::AccelerationStructure {
     fn bind_to(&self, ctx: &mut super::PipelineContext, index: u32) {
         for _ in ctx.targets[index as usize].iter() {
             unimplemented!()
@@ -119,7 +119,7 @@ impl super::CommandEncoder {
 
     pub fn render(
         &mut self,
-        targets: crate::RenderTargetSet,
+        targets: super::RenderTargetSet,
     ) -> super::PassEncoder<super::RenderPipeline> {
         let mut target_size = [0u16; 2];
         let mut invalidate_attachments = Vec::new();
@@ -130,7 +130,7 @@ impl super::CommandEncoder {
                 attachment,
                 view: rt.view,
             });
-            if let crate::FinishOp::Discard = rt.finish_op {
+            if let super::FinishOp::Discard = rt.finish_op {
                 invalidate_attachments.push(attachment);
             }
         }
@@ -145,7 +145,7 @@ impl super::CommandEncoder {
                 attachment,
                 view: rt.view,
             });
-            if let crate::FinishOp::Discard = rt.finish_op {
+            if let super::FinishOp::Discard = rt.finish_op {
                 invalidate_attachments.push(attachment);
             }
         }
@@ -271,10 +271,10 @@ impl<T> Drop for super::PassEncoder<'_, T> {
 
 #[hidden_trait::expose]
 impl crate::traits::TransferEncoder for super::PassEncoder<'_, ()> {
-    type BufferPiece = crate::BufferPiece;
-    type TexturePiece = crate::TexturePiece;
+    type BufferPiece = super::BufferPiece;
+    type TexturePiece = super::TexturePiece;
 
-    fn fill_buffer(&mut self, dst: crate::BufferPiece, size: u64, value: u8) {
+    fn fill_buffer(&mut self, dst: super::BufferPiece, size: u64, value: u8) {
         self.commands.push(super::Command::FillBuffer {
             dst: dst.into(),
             size,
@@ -284,8 +284,8 @@ impl crate::traits::TransferEncoder for super::PassEncoder<'_, ()> {
 
     fn copy_buffer_to_buffer(
         &mut self,
-        src: crate::BufferPiece,
-        dst: crate::BufferPiece,
+        src: super::BufferPiece,
+        dst: super::BufferPiece,
         size: u64,
     ) {
         self.commands.push(super::Command::CopyBufferToBuffer {
@@ -296,8 +296,8 @@ impl crate::traits::TransferEncoder for super::PassEncoder<'_, ()> {
     }
     fn copy_texture_to_texture(
         &mut self,
-        src: crate::TexturePiece,
-        dst: crate::TexturePiece,
+        src: super::TexturePiece,
+        dst: super::TexturePiece,
         size: crate::Extent,
     ) {
         self.commands.push(super::Command::CopyTextureToTexture {
@@ -309,9 +309,9 @@ impl crate::traits::TransferEncoder for super::PassEncoder<'_, ()> {
 
     fn copy_buffer_to_texture(
         &mut self,
-        src: crate::BufferPiece,
+        src: super::BufferPiece,
         bytes_per_row: u32,
-        dst: crate::TexturePiece,
+        dst: super::TexturePiece,
         size: crate::Extent,
     ) {
         self.commands.push(super::Command::CopyBufferToTexture {
@@ -324,8 +324,8 @@ impl crate::traits::TransferEncoder for super::PassEncoder<'_, ()> {
 
     fn copy_texture_to_buffer(
         &mut self,
-        src: crate::TexturePiece,
-        dst: crate::BufferPiece,
+        src: super::TexturePiece,
+        dst: super::BufferPiece,
         bytes_per_row: u32,
         size: crate::Extent,
     ) {
@@ -340,15 +340,15 @@ impl crate::traits::TransferEncoder for super::PassEncoder<'_, ()> {
 
 #[hidden_trait::expose]
 impl crate::traits::AccelerationStructureEncoder for super::PassEncoder<'_, ()> {
-    type AccelerationStructure = crate::AccelerationStructure;
-    type AccelerationStructureMesh = crate::AccelerationStructureMesh;
-    type BufferPiece = crate::BufferPiece;
+    type AccelerationStructure = super::AccelerationStructure;
+    type AccelerationStructureMesh = super::AccelerationStructureMesh;
+    type BufferPiece = super::BufferPiece;
 
     fn build_bottom_level(
         &mut self,
         _acceleration_structure: super::AccelerationStructure,
-        _meshes: &[crate::AccelerationStructureMesh],
-        _scratch_data: crate::BufferPiece,
+        _meshes: &[super::AccelerationStructureMesh],
+        _scratch_data: super::BufferPiece,
     ) {
         unimplemented!()
     }
@@ -358,16 +358,16 @@ impl crate::traits::AccelerationStructureEncoder for super::PassEncoder<'_, ()>
         _acceleration_structure: super::AccelerationStructure,
         _bottom_level: &[super::AccelerationStructure],
         _instance_count: u32,
-        _instance_data: crate::BufferPiece,
-        _scratch_data: crate::BufferPiece,
+        _instance_data: super::BufferPiece,
+        _scratch_data: super::BufferPiece,
     ) {
         unimplemented!()
     }
 }
 
 #[hidden_trait::expose]
-impl crate::traits::PipelineEncoder for super::PipelineEncoder<'_> {
-    fn bind<D: crate::ShaderData>(&mut self, group: u32, data: &D) {
+impl super::traits::PipelineEncoder for super::PipelineEncoder<'_> {
+    fn bind<D: super::ShaderData>(&mut self, group: u32, data: &D) {
         data.fill(super::PipelineContext {
             commands: self.commands,
             plain_data: self.plain_data,
@@ -386,13 +386,13 @@ impl crate::traits::ComputePipelineEncoder for super::PipelineEncoder<'_> {
 
 #[hidden_trait::expose]
 impl crate::traits::RenderPipelineEncoder for super::PipelineEncoder<'_> {
-    type BufferPiece = crate::BufferPiece;
+    type BufferPiece = super::BufferPiece;
 
     fn set_scissor_rect(&mut self, rect: &crate::ScissorRect) {
         self.commands.push(super::Command::SetScissor(rect.clone()));
     }
 
-    fn bind_vertex(&mut self, index: u32, vertex_buf: crate::BufferPiece) {
+    fn bind_vertex(&mut self, index: u32, vertex_buf: super::BufferPiece) {
         self.commands.push(super::Command::BindVertex {
             buffer: vertex_buf.buffer.raw,
         });
@@ -427,7 +427,7 @@ impl crate::traits::RenderPipelineEncoder for super::PipelineEncoder<'_> {
 
     fn draw_indexed(
         &mut self,
-        index_buf: crate::BufferPiece,
+        index_buf: super::BufferPiece,
         index_type: crate::IndexType,
         index_count: u32,
         base_vertex: i32,
@@ -445,15 +445,15 @@ impl crate::traits::RenderPipelineEncoder for super::PipelineEncoder<'_> {
         });
     }
 
-    fn draw_indirect(&mut self, _indirect_buf: crate::BufferPiece) {
+    fn draw_indirect(&mut self, _indirect_buf: super::BufferPiece) {
         unimplemented!()
     }
 
     fn draw_indexed_indirect(
         &mut self,
-        _index_buf: crate::BufferPiece,
+        _index_buf: super::BufferPiece,
         _index_type: crate::IndexType,
-        _indirect_buf: crate::BufferPiece,
+        _indirect_buf: super::BufferPiece,
     ) {
         unimplemented!()
     }
diff --git a/blade-graphics/src/gles/mod.rs b/blade-graphics/src/gles/mod.rs
index f2fcbe1e..112848ed 100644
--- a/blade-graphics/src/gles/mod.rs
+++ b/blade-graphics/src/gles/mod.rs
@@ -34,10 +34,18 @@ pub struct Buffer {
 unsafe impl Send for Buffer {}
 unsafe impl Sync for Buffer {}
 
-impl Buffer {
-    pub fn data(&self) -> *mut u8 {
+#[hidden_trait::expose]
+impl crate::traits::Buffer for Buffer {
+    fn data(&self) -> *mut u8 {
         self.data
     }
+
+    fn at(self, offset: u64) -> crate::GenericBufferPiece<Self> {
+        crate::GenericBufferPiece {
+            buffer: self,
+            offset,
+        }
+    }
 }
 
 #[derive(Clone, Copy, Debug, Hash, PartialEq)]
@@ -143,8 +151,8 @@ struct BufferPart {
     raw: glow::Buffer,
     offset: u64,
 }
-impl From<crate::BufferPiece> for BufferPart {
-    fn from(piece: crate::BufferPiece) -> Self {
+impl From<BufferPiece> for BufferPart {
+    fn from(piece: BufferPiece) -> Self {
         Self {
             raw: piece.buffer.raw,
             offset: piece.offset,
@@ -161,8 +169,8 @@ struct TexturePart {
     array_layer: u32,
     origin: [u32; 3],
 }
-impl From<crate::TexturePiece> for TexturePart {
-    fn from(piece: crate::TexturePiece) -> Self {
+impl From<TexturePiece> for TexturePart {
+    fn from(piece: TexturePiece) -> Self {
         let (raw, target) = piece.texture.inner.as_native();
         Self {
             raw,
@@ -560,3 +568,14 @@ fn map_compare_func(fun: crate::CompareFunction) -> u32 {
         Cf::Always => glow::ALWAYS,
     }
 }
+
+include!("../common/lib.rs");
+
+#[path = "../common/derive.rs"]
+mod derive;
+
+#[path = "../common/util.rs"]
+mod util;
+
+#[cfg(feature = "derive")]
+pub use blade_macros::ShaderDataGles as ShaderData;
diff --git a/blade-graphics/src/gles/resource.rs b/blade-graphics/src/gles/resource.rs
index 426a7759..3539e645 100644
--- a/blade-graphics/src/gles/resource.rs
+++ b/blade-graphics/src/gles/resource.rs
@@ -4,7 +4,7 @@ use std::{ptr, slice};
 impl super::Context {
     pub fn get_bottom_level_acceleration_structure_sizes(
         &self,
-        _meshes: &[crate::AccelerationStructureMesh],
+        _meshes: &[super::AccelerationStructureMesh],
     ) -> crate::AccelerationStructureSizes {
         unimplemented!()
     }
diff --git a/blade-graphics/src/lib.rs b/blade-graphics/src/lib.rs
index 32e1f16a..8ff7a1fc 100644
--- a/blade-graphics/src/lib.rs
+++ b/blade-graphics/src/lib.rs
@@ -48,6 +48,8 @@ pub const IDENTITY_TRANSFORM: Transform = mint::RowMatrix3x4 {
 };
 
 pub mod derive;
+#[cfg(feature = "gles")]
+pub mod gles;
 #[cfg_attr(
     all(not(vulkan), not(gles), any(target_os = "ios", target_os = "macos")),
     path = "metal/mod.rs"
@@ -65,10 +67,9 @@ pub mod derive;
     ),
     path = "vulkan/mod.rs"
 )]
-#[cfg_attr(any(gles, target_arch = "wasm32"), path = "gles/mod.rs")]
-mod hal;
+pub mod hal;
 mod shader;
-mod traits;
+pub mod traits;
 pub mod util;
 pub mod limits {
     pub const PLAIN_DATA_SIZE: u32 = 256;
@@ -76,7 +77,9 @@ pub mod limits {
     pub const STORAGE_BUFFER_ALIGNMENT: u64 = 256;
     pub const ACCELERATION_STRUCTURE_SCRATCH_ALIGNMENT: u64 = 256;
 }
-
+#[cfg(any(gles, target_arch = "wasm32"))]
+pub use gles::*;
+#[cfg(not(any(gles, target_arch = "wasm32")))]
 pub use hal::*;
 
 use std::{fmt, num::NonZeroU32};
@@ -117,9 +120,9 @@ pub enum NotSupportedError {
     ))]
     VulkanError(ash::vk::Result),
 
-    #[cfg(gles)]
+    #[cfg(feature = "gles")]
     GLESLoadingError(egl::LoadError<libloading::Error>),
-    #[cfg(gles)]
+    #[cfg(feature = "gles")]
     GLESError(egl::Error),
 
     NoSupportedDeviceFound,
@@ -171,18 +174,12 @@ pub struct BufferDesc<'a> {
 }
 
 #[derive(Clone, Copy, Debug)]
-pub struct BufferPiece {
-    pub buffer: Buffer,
+pub struct GenericBufferPiece<T: traits::Buffer> {
+    pub buffer: T,
     pub offset: u64,
 }
 
-impl From<Buffer> for BufferPiece {
-    fn from(buffer: Buffer) -> Self {
-        Self { buffer, offset: 0 }
-    }
-}
-
-impl BufferPiece {
+impl<T: traits::Buffer> GenericBufferPiece<T> {
     pub fn data(&self) -> *mut u8 {
         let base = self.buffer.data();
         assert!(!base.is_null());
@@ -190,12 +187,9 @@ impl BufferPiece {
     }
 }
 
-impl Buffer {
-    pub fn at(self, offset: u64) -> BufferPiece {
-        BufferPiece {
-            buffer: self,
-            offset,
-        }
+impl<T: traits::Buffer> From<T> for GenericBufferPiece<T> {
+    fn from(buffer: T) -> Self {
+        Self { buffer, offset: 0 }
     }
 }
 
@@ -243,27 +237,6 @@ impl<T, const N: ResourceIndex> std::ops::IndexMut<ResourceIndex> for ResourceAr
         &mut self.data[index as usize]
     }
 }
-pub type BufferArray<const N: ResourceIndex> = ResourceArray<BufferPiece, N>;
-pub type TextureArray<const N: ResourceIndex> = ResourceArray<TextureView, N>;
-
-#[derive(Clone, Copy, Debug)]
-pub struct TexturePiece {
-    pub texture: Texture,
-    pub mip_level: u32,
-    pub array_layer: u32,
-    pub origin: [u32; 3],
-}
-
-impl From<Texture> for TexturePiece {
-    fn from(texture: Texture) -> Self {
-        Self {
-            texture,
-            mip_level: 0,
-            array_layer: 0,
-            origin: [0; 3],
-        }
-    }
-}
 
 #[non_exhaustive]
 #[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
@@ -523,19 +496,6 @@ pub enum VertexFormat {
     I32Vec4,
 }
 
-#[derive(Clone, Debug)]
-pub struct AccelerationStructureMesh {
-    pub vertex_data: BufferPiece,
-    pub vertex_format: VertexFormat,
-    pub vertex_stride: u32,
-    pub vertex_count: u32,
-    pub index_data: BufferPiece,
-    pub index_type: Option<IndexType>,
-    pub triangle_count: u32,
-    pub transform_data: BufferPiece,
-    pub is_opaque: bool,
-}
-
 #[derive(Clone, Debug)]
 pub struct AccelerationStructureInstance {
     pub acceleration_structure_index: u32,
@@ -585,10 +545,6 @@ pub enum ShaderBinding {
     Plain { size: u32 },
 }
 
-pub trait ShaderBindable: Clone + Copy + derive::HasShaderBinding {
-    fn bind_to(&self, context: &mut PipelineContext, index: u32);
-}
-
 #[derive(Debug)]
 struct ShaderDataInfo {
     visibility: ShaderVisibility,
@@ -612,11 +568,6 @@ impl ShaderDataLayout {
     }
 }
 
-pub trait ShaderData {
-    fn layout() -> ShaderDataLayout;
-    fn fill(&self, context: PipelineContext);
-}
-
 #[derive(Copy, Clone, Debug, PartialEq)]
 pub struct VertexAttribute {
     pub offset: u32,
@@ -1016,27 +967,6 @@ pub enum InitOp {
     Clear(TextureColor),
 }
 
-#[derive(Clone, Copy, Debug)]
-pub enum FinishOp {
-    Store,
-    Discard,
-    ResolveTo(TextureView),
-    Ignore,
-}
-
-#[derive(Debug)]
-pub struct RenderTarget {
-    pub view: TextureView,
-    pub init_op: InitOp,
-    pub finish_op: FinishOp,
-}
-
-#[derive(Debug)]
-pub struct RenderTargetSet<'a> {
-    pub colors: &'a [RenderTarget],
-    pub depth_stencil: Option<RenderTarget>,
-}
-
 /// Mechanism used to acquire frames and display them on screen.
 #[derive(Copy, Clone, Debug, Default, Hash, Eq, PartialEq)]
 pub enum DisplaySync {
diff --git a/blade-graphics/src/shader.rs b/blade-graphics/src/shader.rs
index 6cc9f6b2..c90dda6f 100644
--- a/blade-graphics/src/shader.rs
+++ b/blade-graphics/src/shader.rs
@@ -8,41 +8,6 @@ impl From<naga::ShaderStage> for super::ShaderVisibility {
     }
 }
 
-impl super::Context {
-    pub fn try_create_shader(
-        &self,
-        desc: super::ShaderDesc,
-    ) -> Result<super::Shader, &'static str> {
-        let module = naga::front::wgsl::parse_str(desc.source).map_err(|e| {
-            e.emit_to_stderr_with_path(desc.source, "");
-            "compilation failed"
-        })?;
-
-        let device_caps = self.capabilities();
-
-        // Bindings are set up at pipeline creation, ignore here
-        let flags = naga::valid::ValidationFlags::all() ^ naga::valid::ValidationFlags::BINDINGS;
-        let mut caps = naga::valid::Capabilities::empty();
-        caps.set(
-            naga::valid::Capabilities::RAY_QUERY | naga::valid::Capabilities::SAMPLED_TEXTURE_AND_STORAGE_BUFFER_ARRAY_NON_UNIFORM_INDEXING,
-            !device_caps.ray_query.is_empty(),
-        );
-        let info = naga::valid::Validator::new(flags, caps)
-            .validate(&module)
-            .map_err(|e| {
-                crate::util::emit_annotated_error(&e, "", desc.source);
-                crate::util::print_err(&e);
-                "validation failed"
-            })?;
-
-        Ok(super::Shader { module, info })
-    }
-
-    pub fn create_shader(&self, desc: super::ShaderDesc) -> super::Shader {
-        self.try_create_shader(desc).unwrap()
-    }
-}
-
 impl super::Shader {
     pub fn at<'a>(&'a self, entry_point: &'a str) -> super::ShaderFunction<'a> {
         super::ShaderFunction {
diff --git a/blade-graphics/src/traits.rs b/blade-graphics/src/traits.rs
index d09013eb..537d4e47 100644
--- a/blade-graphics/src/traits.rs
+++ b/blade-graphics/src/traits.rs
@@ -1,6 +1,6 @@
 use std::{fmt::Debug, hash::Hash};
 pub trait ResourceDevice {
-    type Buffer: Send + Sync + Clone + Copy + Debug + Hash + PartialEq;
+    type Buffer: Buffer;
     type Texture: Send + Sync + Clone + Copy + Debug + Hash + PartialEq;
     type TextureView: Send + Sync + Clone + Copy + Debug + Hash + PartialEq;
     type Sampler: Send + Sync + Clone + Copy + Debug + Hash + PartialEq;
@@ -88,15 +88,11 @@ pub trait AccelerationStructureEncoder {
     );
 }
 
-pub trait PipelineEncoder {
-    fn bind<D: super::ShaderData>(&mut self, group: u32, data: &D);
-}
-
-pub trait ComputePipelineEncoder: PipelineEncoder {
+pub trait ComputePipelineEncoder {
     fn dispatch(&mut self, groups: [u32; 3]);
 }
 
-pub trait RenderPipelineEncoder: PipelineEncoder {
+pub trait RenderPipelineEncoder {
     type BufferPiece: Send + Sync + Clone + Copy + Debug;
 
     //TODO: reconsider exposing this here
@@ -126,3 +122,8 @@ pub trait RenderPipelineEncoder: PipelineEncoder {
         indirect_buf: Self::BufferPiece,
     );
 }
+
+pub trait Buffer: Send + Sync + Clone + Copy + Debug + Hash + PartialEq {
+    fn data(&self) -> *mut u8;
+    fn at(self, offset: u64) -> crate::GenericBufferPiece<Self>;
+}
diff --git a/blade-graphics/src/util.rs b/blade-graphics/src/util.rs
index 23e876e0..89027150 100644
--- a/blade-graphics/src/util.rs
+++ b/blade-graphics/src/util.rs
@@ -93,15 +93,3 @@ impl super::TextureFormat {
         }
     }
 }
-
-impl super::ComputePipeline {
-    /// Return the dispatch group counts sufficient to cover the given extent.
-    pub fn get_dispatch_for(&self, extent: super::Extent) -> [u32; 3] {
-        let wg_size = self.get_workgroup_size();
-        [
-            (extent.width + wg_size[0] - 1) / wg_size[0],
-            (extent.height + wg_size[1] - 1) / wg_size[1],
-            (extent.depth + wg_size[2] - 1) / wg_size[2],
-        ]
-    }
-}
diff --git a/blade-graphics/src/vulkan/command.rs b/blade-graphics/src/vulkan/command.rs
index 1206d9b2..02082862 100644
--- a/blade-graphics/src/vulkan/command.rs
+++ b/blade-graphics/src/vulkan/command.rs
@@ -50,12 +50,12 @@ impl super::PipelineContext<'_> {
     }
 }
 
-impl<T: bytemuck::Pod> crate::ShaderBindable for T {
+impl<T: bytemuck::Pod> super::ShaderBindable for T {
     fn bind_to(&self, ctx: &mut super::PipelineContext, index: u32) {
         ctx.write(index, *self);
     }
 }
-impl crate::ShaderBindable for super::TextureView {
+impl super::ShaderBindable for super::TextureView {
     fn bind_to(&self, ctx: &mut super::PipelineContext, index: u32) {
         ctx.write(
             index,
@@ -67,7 +67,7 @@ impl crate::ShaderBindable for super::TextureView {
         );
     }
 }
-impl<'a, const N: crate::ResourceIndex> crate::ShaderBindable for &'a crate::TextureArray<N> {
+impl<'a, const N: crate::ResourceIndex> super::ShaderBindable for &'a super::TextureArray<N> {
     fn bind_to(&self, ctx: &mut super::PipelineContext, index: u32) {
         ctx.write_array(
             index,
@@ -79,7 +79,7 @@ impl<'a, const N: crate::ResourceIndex> crate::ShaderBindable for &'a crate::Tex
         );
     }
 }
-impl crate::ShaderBindable for super::Sampler {
+impl super::ShaderBindable for super::Sampler {
     fn bind_to(&self, ctx: &mut super::PipelineContext, index: u32) {
         ctx.write(
             index,
@@ -91,7 +91,7 @@ impl crate::ShaderBindable for super::Sampler {
         );
     }
 }
-impl crate::ShaderBindable for crate::BufferPiece {
+impl super::ShaderBindable for super::BufferPiece {
     fn bind_to(&self, ctx: &mut super::PipelineContext, index: u32) {
         ctx.write(
             index,
@@ -103,7 +103,7 @@ impl crate::ShaderBindable for crate::BufferPiece {
         );
     }
 }
-impl<'a, const N: crate::ResourceIndex> crate::ShaderBindable for &'a crate::BufferArray<N> {
+impl<'a, const N: crate::ResourceIndex> super::ShaderBindable for &'a super::BufferArray<N> {
     fn bind_to(&self, ctx: &mut super::PipelineContext, index: u32) {
         ctx.write_array(
             index,
@@ -115,13 +115,13 @@ impl<'a, const N: crate::ResourceIndex> crate::ShaderBindable for &'a crate::Buf
         );
     }
 }
-impl crate::ShaderBindable for super::AccelerationStructure {
+impl super::ShaderBindable for super::AccelerationStructure {
     fn bind_to(&self, ctx: &mut super::PipelineContext, index: u32) {
         ctx.write(index, self.raw);
     }
 }
 
-impl crate::TexturePiece {
+impl super::TexturePiece {
     fn subresource_layers(&self, aspects: crate::TexelAspects) -> vk::ImageSubresourceLayers {
         vk::ImageSubresourceLayers {
             aspect_mask: super::map_aspects(aspects),
@@ -141,9 +141,9 @@ fn map_origin(origin: &[u32; 3]) -> vk::Offset3D {
 }
 
 fn make_buffer_image_copy(
-    buffer: &crate::BufferPiece,
+    buffer: &super::BufferPiece,
     bytes_per_row: u32,
-    texture: &crate::TexturePiece,
+    texture: &super::TexturePiece,
     size: &crate::Extent,
 ) -> vk::BufferImageCopy {
     let block_info = texture.texture.format.block_info();
@@ -158,7 +158,7 @@ fn make_buffer_image_copy(
     }
 }
 
-fn map_render_target(rt: &crate::RenderTarget) -> vk::RenderingAttachmentInfo<'static> {
+fn map_render_target(rt: &super::RenderTarget) -> vk::RenderingAttachmentInfo<'static> {
     let mut vk_info = vk::RenderingAttachmentInfo::default()
         .image_view(rt.view.raw)
         .image_layout(vk::ImageLayout::GENERAL)
@@ -358,7 +358,7 @@ impl super::CommandEncoder {
         }
     }
 
-    pub fn render(&mut self, targets: crate::RenderTargetSet) -> super::RenderCommandEncoder {
+    pub fn render(&mut self, targets: super::RenderTargetSet) -> super::RenderCommandEncoder {
         self.barrier();
         self.mark("pass/render");
 
@@ -450,8 +450,8 @@ impl super::CommandEncoder {
 
 #[hidden_trait::expose]
 impl crate::traits::TransferEncoder for super::TransferCommandEncoder<'_> {
-    type BufferPiece = crate::BufferPiece;
-    type TexturePiece = crate::TexturePiece;
+    type BufferPiece = super::BufferPiece;
+    type TexturePiece = super::TexturePiece;
 
     fn fill_buffer(&mut self, dst: crate::BufferPiece, size: u64, value: u8) {
         let value_u32 = (value as u32) * 0x1010101;
@@ -464,8 +464,8 @@ impl crate::traits::TransferEncoder for super::TransferCommandEncoder<'_> {
 
     fn copy_buffer_to_buffer(
         &mut self,
-        src: crate::BufferPiece,
-        dst: crate::BufferPiece,
+        src: super::BufferPiece,
+        dst: super::BufferPiece,
         size: u64,
     ) {
         let copy = vk::BufferCopy {
@@ -482,8 +482,8 @@ impl crate::traits::TransferEncoder for super::TransferCommandEncoder<'_> {
 
     fn copy_texture_to_texture(
         &mut self,
-        src: crate::TexturePiece,
-        dst: crate::TexturePiece,
+        src: super::TexturePiece,
+        dst: super::TexturePiece,
         size: crate::Extent,
     ) {
         let copy = vk::ImageCopy {
@@ -507,9 +507,9 @@ impl crate::traits::TransferEncoder for super::TransferCommandEncoder<'_> {
 
     fn copy_buffer_to_texture(
         &mut self,
-        src: crate::BufferPiece,
+        src: super::BufferPiece,
         bytes_per_row: u32,
-        dst: crate::TexturePiece,
+        dst: super::TexturePiece,
         size: crate::Extent,
     ) {
         let copy = make_buffer_image_copy(&src, bytes_per_row, &dst, &size);
@@ -526,8 +526,8 @@ impl crate::traits::TransferEncoder for super::TransferCommandEncoder<'_> {
 
     fn copy_texture_to_buffer(
         &mut self,
-        src: crate::TexturePiece,
-        dst: crate::BufferPiece,
+        src: super::TexturePiece,
+        dst: super::BufferPiece,
         bytes_per_row: u32,
         size: crate::Extent,
     ) {
@@ -548,9 +548,9 @@ impl crate::traits::TransferEncoder for super::TransferCommandEncoder<'_> {
 impl crate::traits::AccelerationStructureEncoder
     for super::AccelerationStructureCommandEncoder<'_>
 {
-    type AccelerationStructure = crate::AccelerationStructure;
-    type AccelerationStructureMesh = crate::AccelerationStructureMesh;
-    type BufferPiece = crate::BufferPiece;
+    type AccelerationStructure = super::AccelerationStructure;
+    type AccelerationStructureMesh = super::AccelerationStructureMesh;
+    type BufferPiece = super::BufferPiece;
 
     fn build_bottom_level(
         &mut self,
@@ -584,8 +584,8 @@ impl crate::traits::AccelerationStructureEncoder
         acceleration_structure: super::AccelerationStructure,
         _bottom_level: &[super::AccelerationStructure],
         instance_count: u32,
-        instance_data: crate::BufferPiece,
-        scratch_data: crate::BufferPiece,
+        instance_data: super::BufferPiece,
+        scratch_data: super::BufferPiece,
     ) {
         let build_range_info = vk::AccelerationStructureBuildRangeInfoKHR {
             primitive_count: instance_count,
@@ -700,8 +700,8 @@ impl super::PipelineEncoder<'_, '_> {
 }
 
 #[hidden_trait::expose]
-impl crate::traits::PipelineEncoder for super::PipelineEncoder<'_, '_> {
-    fn bind<D: crate::ShaderData>(&mut self, group: u32, data: &D) {
+impl super::traits::PipelineEncoder for super::PipelineEncoder<'_, '_> {
+    fn bind<D: super::ShaderData>(&mut self, group: u32, data: &D) {
         let dsl = &self.layout.descriptor_set_layouts[group as usize];
         self.update_data.clear();
         self.update_data.resize(dsl.template_size as usize, 0);
@@ -744,7 +744,7 @@ impl crate::traits::ComputePipelineEncoder for super::PipelineEncoder<'_, '_> {
 
 #[hidden_trait::expose]
 impl crate::traits::RenderPipelineEncoder for super::PipelineEncoder<'_, '_> {
-    type BufferPiece = crate::BufferPiece;
+    type BufferPiece = super::BufferPiece;
 
     fn set_scissor_rect(&mut self, rect: &crate::ScissorRect) {
         let vk_scissor = vk::Rect2D {
@@ -764,7 +764,7 @@ impl crate::traits::RenderPipelineEncoder for super::PipelineEncoder<'_, '_> {
         };
     }
 
-    fn bind_vertex(&mut self, index: u32, vertex_buf: crate::BufferPiece) {
+    fn bind_vertex(&mut self, index: u32, vertex_buf: super::BufferPiece) {
         unsafe {
             self.device.core.cmd_bind_vertex_buffers(
                 self.cmd_buf.raw,
@@ -795,7 +795,7 @@ impl crate::traits::RenderPipelineEncoder for super::PipelineEncoder<'_, '_> {
 
     fn draw_indexed(
         &mut self,
-        index_buf: crate::BufferPiece,
+        index_buf: super::BufferPiece,
         index_type: crate::IndexType,
         index_count: u32,
         base_vertex: i32,
@@ -821,7 +821,7 @@ impl crate::traits::RenderPipelineEncoder for super::PipelineEncoder<'_, '_> {
         }
     }
 
-    fn draw_indirect(&mut self, indirect_buf: crate::BufferPiece) {
+    fn draw_indirect(&mut self, indirect_buf: super::BufferPiece) {
         unsafe {
             self.device.core.cmd_draw_indirect(
                 self.cmd_buf.raw,
@@ -835,9 +835,9 @@ impl crate::traits::RenderPipelineEncoder for super::PipelineEncoder<'_, '_> {
 
     fn draw_indexed_indirect(
         &mut self,
-        index_buf: crate::BufferPiece,
+        index_buf: super::BufferPiece,
         index_type: crate::IndexType,
-        indirect_buf: crate::BufferPiece,
+        indirect_buf: super::BufferPiece,
     ) {
         let raw_index_type = super::map_index_type(index_type);
         unsafe {
diff --git a/blade-graphics/src/vulkan/mod.rs b/blade-graphics/src/vulkan/mod.rs
index 5329527d..7db09d04 100644
--- a/blade-graphics/src/vulkan/mod.rs
+++ b/blade-graphics/src/vulkan/mod.rs
@@ -127,10 +127,18 @@ impl Default for Buffer {
     }
 }
 
-impl Buffer {
-    pub fn data(&self) -> *mut u8 {
+#[hidden_trait::expose]
+impl crate::traits::Buffer for Buffer {
+    fn data(&self) -> *mut u8 {
         self.mapped_data
     }
+
+    fn at(self, offset: u64) -> crate::GenericBufferPiece<Self> {
+        crate::GenericBufferPiece {
+            buffer: self,
+            offset,
+        }
+    }
 }
 
 unsafe impl Send for Buffer {}
@@ -577,7 +585,7 @@ struct BottomLevelAccelerationStructureInput<'a> {
 }
 
 impl Device {
-    fn get_device_address(&self, piece: &crate::BufferPiece) -> u64 {
+    fn get_device_address(&self, piece: &BufferPiece) -> u64 {
         let vk_info = vk::BufferDeviceAddressInfo {
             buffer: piece.buffer.raw,
             ..Default::default()
@@ -588,7 +596,7 @@ impl Device {
 
     fn map_acceleration_structure_meshes(
         &self,
-        meshes: &[crate::AccelerationStructureMesh],
+        meshes: &[AccelerationStructureMesh],
     ) -> BottomLevelAccelerationStructureInput {
         let mut total_primitive_count = 0;
         let mut max_primitive_counts = Vec::with_capacity(meshes.len());
@@ -670,3 +678,14 @@ impl Device {
         }
     }
 }
+
+include!("../common/lib.rs");
+
+#[path = "../common/derive.rs"]
+mod derive;
+
+#[path = "../common/util.rs"]
+mod util;
+
+#[cfg(feature = "derive")]
+pub use blade_macros::ShaderDataHal as ShaderData;
diff --git a/blade-graphics/src/vulkan/resource.rs b/blade-graphics/src/vulkan/resource.rs
index d1dd7200..04b2db25 100644
--- a/blade-graphics/src/vulkan/resource.rs
+++ b/blade-graphics/src/vulkan/resource.rs
@@ -86,7 +86,7 @@ impl super::Context {
     //TODO: move these into `ResourceDevice` trait when ready
     pub fn get_bottom_level_acceleration_structure_sizes(
         &self,
-        meshes: &[crate::AccelerationStructureMesh],
+        meshes: &[super::AccelerationStructureMesh],
     ) -> crate::AccelerationStructureSizes {
         let blas_input = self.device.map_acceleration_structure_meshes(meshes);
         let rt = self.device.ray_tracing.as_ref().unwrap();
diff --git a/blade-macros/src/lib.rs b/blade-macros/src/lib.rs
index b59475ad..863ad892 100644
--- a/blade-macros/src/lib.rs
+++ b/blade-macros/src/lib.rs
@@ -4,6 +4,7 @@ mod shader_data;
 mod vertex;
 
 use proc_macro::TokenStream;
+use quote::quote;
 
 /// Derive the `ShaderData` trait for a struct
 ///
@@ -15,9 +16,28 @@ use proc_macro::TokenStream;
 ///   sm: blade_graphics::Sampler,
 /// }
 /// ```
-#[proc_macro_derive(ShaderData)]
-pub fn shader_data_derive(input: TokenStream) -> TokenStream {
-    let stream = match shader_data::generate(input) {
+#[proc_macro_derive(ShaderDataHal)]
+pub fn shader_data_hal_derive(input: TokenStream) -> TokenStream {
+    let stream = match shader_data::generate(input, quote! { blade_graphics::hal }) {
+        Ok(tokens) => tokens,
+        Err(err) => err.into_compile_error(),
+    };
+    stream.into()
+}
+
+/// Derive the `ShaderData` trait for a struct
+///
+/// ## Example
+///
+/// ```rust
+/// #[derive(blade_macros::ShaderData)]
+/// struct Test {
+///   sm: blade_graphics::Sampler,
+/// }
+/// ```
+#[proc_macro_derive(ShaderDataGles)]
+pub fn shader_data_gles_derive(input: TokenStream) -> TokenStream {
+    let stream = match shader_data::generate(input, quote! { blade_graphics::gles }) {
         Ok(tokens) => tokens,
         Err(err) => err.into_compile_error(),
     };
diff --git a/blade-macros/src/shader_data.rs b/blade-macros/src/shader_data.rs
index 13ffddcb..41350ff6 100644
--- a/blade-macros/src/shader_data.rs
+++ b/blade-macros/src/shader_data.rs
@@ -1,7 +1,7 @@
 use proc_macro::TokenStream;
 use quote::quote;
 
-pub fn generate(input_stream: TokenStream) -> syn::Result<proc_macro2::TokenStream> {
+pub fn generate(input_stream: TokenStream, namespace: proc_macro2::TokenStream) -> syn::Result<proc_macro2::TokenStream> {
     let item_struct = syn::parse::<syn::ItemStruct>(input_stream)?;
     let fields = match item_struct.fields {
         syn::Fields::Named(ref fields) => fields,
@@ -44,14 +44,14 @@ pub fn generate(input_stream: TokenStream) -> syn::Result<proc_macro2::TokenStre
 
     let struct_name = item_struct.ident;
     Ok(quote! {
-        impl<#(#generics),*> blade_graphics::ShaderData for #struct_name<#(#generics),*> {
+        impl<#(#generics,)*> #namespace::ShaderData for #struct_name<#(#generics),*> {
             fn layout() -> blade_graphics::ShaderDataLayout {
                 blade_graphics::ShaderDataLayout {
                     bindings: vec![#(#bindings),*],
                 }
             }
-            fn fill(&self, mut ctx: blade_graphics::PipelineContext) {
-                use blade_graphics::ShaderBindable as _;
+            fn fill(&self, mut ctx: #namespace::PipelineContext) {
+                use #namespace::ShaderBindable as _;
                 #(#assignments)*
             }
         }