diff --git a/crates/voltex_renderer/src/rt_shadow.rs b/crates/voltex_renderer/src/rt_shadow.rs
new file mode 100644
index 0000000..d4832de
--- /dev/null
+++ b/crates/voltex_renderer/src/rt_shadow.rs
@@ -0,0 +1,89 @@
+use bytemuck::{Pod, Zeroable};
+
+/// Texture format used for the RT shadow output (single-channel float).
+pub const RT_SHADOW_FORMAT: wgpu::TextureFormat = wgpu::TextureFormat::R32Float;
+
+/// Uniform buffer for the RT shadow compute pass.
+///
+/// Layout (32 bytes, 16-byte aligned):
+///   light_direction [f32; 3] + _pad0: f32   → 16 bytes
+///   width: u32, height: u32, _pad1: [u32; 2] → 16 bytes
+#[repr(C)]
+#[derive(Copy, Clone, Debug, Pod, Zeroable)]
+pub struct RtShadowUniform {
+    pub light_direction: [f32; 3],
+    pub _pad0: f32,
+    pub width: u32,
+    pub height: u32,
+    pub _pad1: [u32; 2],
+}
+
+/// GPU resources for the RT shadow compute pass.
+pub struct RtShadowResources {
+    pub shadow_texture: wgpu::Texture,
+    pub shadow_view: wgpu::TextureView,
+    pub uniform_buffer: wgpu::Buffer,
+    pub width: u32,
+    pub height: u32,
+}
+
+impl RtShadowResources {
+    pub fn new(device: &wgpu::Device, width: u32, height: u32) -> Self {
+        let (shadow_texture, shadow_view) = create_shadow_texture(device, width, height);
+
+        let uniform_buffer = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("RT Shadow Uniform Buffer"),
+            size: std::mem::size_of::<RtShadowUniform>() as u64,
+            usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST,
+            mapped_at_creation: false,
+        });
+
+        Self { shadow_texture, shadow_view, uniform_buffer, width, height }
+    }
+
+    /// Recreate the shadow texture when the window is resized.
+    pub fn resize(&mut self, device: &wgpu::Device, width: u32, height: u32) {
+        let (shadow_texture, shadow_view) = create_shadow_texture(device, width, height);
+        self.shadow_texture = shadow_texture;
+        self.shadow_view = shadow_view;
+        self.width = width;
+        self.height = height;
+    }
+}
+
+// ── Helpers ──────────────────────────────────────────────────────────────────
+
+fn create_shadow_texture(
+    device: &wgpu::Device,
+    width: u32,
+    height: u32,
+) -> (wgpu::Texture, wgpu::TextureView) {
+    let texture = device.create_texture(&wgpu::TextureDescriptor {
+        label: Some("RT Shadow Texture"),
+        size: wgpu::Extent3d {
+            width,
+            height,
+            depth_or_array_layers: 1,
+        },
+        mip_level_count: 1,
+        sample_count: 1,
+        dimension: wgpu::TextureDimension::D2,
+        format: RT_SHADOW_FORMAT,
+        usage: wgpu::TextureUsages::STORAGE_BINDING | wgpu::TextureUsages::TEXTURE_BINDING,
+        view_formats: &[],
+    });
+    let view = texture.create_view(&wgpu::TextureViewDescriptor::default());
+    (texture, view)
+}
+
+// ── Tests ─────────────────────────────────────────────────────────────────────
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    #[test]
+    fn test_rt_shadow_uniform_size() {
+        assert_eq!(std::mem::size_of::<RtShadowUniform>(), 32);
+    }
+}
diff --git a/crates/voltex_renderer/src/rt_shadow_shader.wgsl b/crates/voltex_renderer/src/rt_shadow_shader.wgsl
new file mode 100644
index 0000000..15fe1a0
--- /dev/null
+++ b/crates/voltex_renderer/src/rt_shadow_shader.wgsl
@@ -0,0 +1,79 @@
+// RT Shadow compute shader.
+// Reads world-space position and normal from the G-Buffer, then fires a
+// shadow ray against the TLAS to determine per-pixel visibility.
+// Output is 1.0 (lit) or 0.0 (shadowed) stored in an R32Float texture.
+
+// ── Group 0: G-Buffer inputs ──────────────────────────────────────────────────
+
+@group(0) @binding(0) var t_position: texture_2d<f32>;
+@group(0) @binding(1) var t_normal:   texture_2d<f32>;
+
+// ── Group 1: RT data ─────────────────────────────────────────────────────────
+
+@group(1) @binding(0) var tlas:          acceleration_structure;
+@group(1) @binding(1) var t_shadow_out:  texture_storage_2d<r32float, write>;
+
+struct RtShadowUniform {
+    light_direction: vec3<f32>,
+    _pad0: f32,
+    width: u32,
+    height: u32,
+    _pad1: vec2<u32>,
+};
+
+@group(1) @binding(2) var<uniform> uniforms: RtShadowUniform;
+
+// ── Compute entry point ───────────────────────────────────────────────────────
+
+@compute @workgroup_size(8, 8)
+fn cs_main(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let coord = vec2<i32>(i32(gid.x), i32(gid.y));
+
+    // Bounds check
+    if gid.x >= uniforms.width || gid.y >= uniforms.height {
+        return;
+    }
+
+    // Read world position from G-Buffer
+    let world_pos = textureLoad(t_position, coord, 0).xyz;
+
+    // Background pixel: skip (position is (0,0,0) for skybox pixels)
+    if dot(world_pos, world_pos) < 0.001 {
+        textureStore(t_shadow_out, coord, vec4<f32>(1.0, 0.0, 0.0, 0.0));
+        return;
+    }
+
+    // Read and decode normal — G-Buffer stores N * 0.5 + 0.5
+    let normal_encoded = textureLoad(t_normal, coord, 0).rgb;
+    let N = normalize(normal_encoded * 2.0 - 1.0);
+
+    // Ray: from surface towards the light, biased along normal to avoid self-intersection
+    let ray_origin = world_pos + N * 0.01;
+    let ray_dir    = normalize(-uniforms.light_direction);
+
+    // Ray query: check for any occluder between surface and "infinity"
+    var rq: ray_query;
+    rayQueryInitialize(
+        &rq,
+        tlas,
+        RayDesc(
+            RAY_FLAG_TERMINATE_ON_FIRST_HIT | RAY_FLAG_SKIP_CLOSEST_HIT_SHADER,
+            0xFF,
+            0.0001,
+            1.0e6,
+            ray_origin,
+            ray_dir,
+        ),
+    );
+
+    // Advance until the query is done (with TERMINATE_ON_FIRST_HIT this is at most one step)
+    while rayQueryProceed(&rq) {}
+
+    // If anything was hit, the pixel is in shadow
+    var shadow: f32 = 1.0;
+    if rayQueryGetCommittedIntersectionType(&rq) != RAY_QUERY_INTERSECTION_NONE {
+        shadow = 0.0;
+    }
+
+    textureStore(t_shadow_out, coord, vec4<f32>(shadow, 0.0, 0.0, 0.0));
+}