feat: GPU procedural Phase 4 - texture composition

Multi-input composite shaders with sampler support.
- Dynamic bind group layouts (N input textures + 1 sampler)
- dispatch_composite() for multi-input compute dispatch
- create_gpu_composite_texture() API
- gen_blend.wgsl and gen_mask.wgsl shaders

Guarded with #if !defined(STRIP_GPU_COMPOSITE) for easy removal.

Tests:
- Blend two noise textures
- Mask noise with grid
- Multi-stage composite (composite of composites)

Size: ~830 bytes (2 shaders + dispatch logic)

handoff(Claude): GPU procedural Phase 4 complete

9 files changed, 460 insertions, 7 deletions

diff --git a/CMakeLists.txt b/CMakeLists.txt
index 26818c3..fb6beef 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -601,6 +601,14 @@ if(DEMO_BUILD_TESTS)
     target_link_libraries(test_gpu_procedural PRIVATE 3d gpu audio procedural util ${DEMO_LIBS})
     add_dependencies(test_gpu_procedural generate_demo_assets)
 
+    # GPU Composite Texture Test (Phase 4)
+    add_demo_test(test_gpu_composite GpuCompositeTest
+        src/tests/test_gpu_composite.cc
+        ${PLATFORM_SOURCES}
+        ${GEN_DEMO_CC})
+    target_link_libraries(test_gpu_composite PRIVATE 3d gpu audio procedural util ${DEMO_LIBS})
+    add_dependencies(test_gpu_composite generate_demo_assets)
+
     # Gantt chart output test (bash script)
     add_test(
         NAME GanttOutputTest
diff --git a/assets/final/demo_assets.txt b/assets/final/demo_assets.txt
index 0f7c267..819c9c9 100644
--- a/assets/final/demo_assets.txt
+++ b/assets/final/demo_assets.txt
@@ -55,6 +55,8 @@ SHADER_VIGNETTE, NONE, shaders/vignette.wgsl, "Vignette Shader"
 SHADER_COMPUTE_GEN_NOISE, NONE, shaders/compute/gen_noise.wgsl, "GPU Noise Compute Shader"
 SHADER_COMPUTE_GEN_PERLIN, NONE, shaders/compute/gen_perlin.wgsl, "GPU Perlin Noise Compute Shader"
 SHADER_COMPUTE_GEN_GRID, NONE, shaders/compute/gen_grid.wgsl, "GPU Grid Compute Shader"
+SHADER_COMPUTE_GEN_BLEND, NONE, shaders/compute/gen_blend.wgsl, "GPU Blend Composite Shader"
+SHADER_COMPUTE_GEN_MASK, NONE, shaders/compute/gen_mask.wgsl, "GPU Mask Composite Shader"
 CIRCLE_MASK_COMPUTE_SHADER, NONE, shaders/circle_mask_compute.wgsl, "Circle mask compute shader"
 CIRCLE_MASK_RENDER_SHADER, NONE, shaders/circle_mask_render.wgsl, "Circle mask render shader"
 MASKED_CUBE_SHADER, NONE, shaders/masked_cube.wgsl, "Masked cube shader"
diff --git a/assets/final/shaders/compute/gen_blend.wgsl b/assets/final/shaders/compute/gen_blend.wgsl
new file mode 100644
index 0000000..9fc9e1e
--- /dev/null
+++ b/assets/final/shaders/compute/gen_blend.wgsl
@@ -0,0 +1,29 @@
+// This file is part of the 64k demo project.
+// GPU composite shader: Blend two textures.
+
+struct BlendParams {
+  width: u32,
+  height: u32,
+  blend_factor: f32,
+  _pad0: f32,
+}
+
+@group(0) @binding(0) var output_tex: texture_storage_2d<rgba8unorm, write>;
+@group(0) @binding(1) var<uniform> params: BlendParams;
+@group(0) @binding(2) var input_a: texture_2d<f32>;
+@group(0) @binding(3) var input_b: texture_2d<f32>;
+@group(0) @binding(4) var tex_sampler: sampler;
+
+@compute @workgroup_size(8, 8, 1)
+fn main(@builtin(global_invocation_id) id: vec3<u32>) {
+  if (id.x >= params.width || id.y >= params.height) { return; }
+
+  let uv = vec2<f32>(f32(id.x) / f32(params.width),
+                      f32(id.y) / f32(params.height));
+
+  let color_a = textureSampleLevel(input_a, tex_sampler, uv, 0.0);
+  let color_b = textureSampleLevel(input_b, tex_sampler, uv, 0.0);
+  let blended = mix(color_a, color_b, params.blend_factor);
+
+  textureStore(output_tex, id.xy, blended);
+}
diff --git a/assets/final/shaders/compute/gen_mask.wgsl b/assets/final/shaders/compute/gen_mask.wgsl
new file mode 100644
index 0000000..1ce9f52
--- /dev/null
+++ b/assets/final/shaders/compute/gen_mask.wgsl
@@ -0,0 +1,27 @@
+// This file is part of the 64k demo project.
+// GPU composite shader: Multiply texture A by texture B (masking).
+
+struct MaskParams {
+  width: u32,
+  height: u32,
+}
+
+@group(0) @binding(0) var output_tex: texture_storage_2d<rgba8unorm, write>;
+@group(0) @binding(1) var<uniform> params: MaskParams;
+@group(0) @binding(2) var input_a: texture_2d<f32>;
+@group(0) @binding(3) var input_b: texture_2d<f32>;
+@group(0) @binding(4) var tex_sampler: sampler;
+
+@compute @workgroup_size(8, 8, 1)
+fn main(@builtin(global_invocation_id) id: vec3<u32>) {
+  if (id.x >= params.width || id.y >= params.height) { return; }
+
+  let uv = vec2<f32>(f32(id.x) / f32(params.width),
+                      f32(id.y) / f32(params.height));
+
+  let color_a = textureSampleLevel(input_a, tex_sampler, uv, 0.0);
+  let mask_b = textureSampleLevel(input_b, tex_sampler, uv, 0.0);
+  let masked = color_a * mask_b;
+
+  textureStore(output_tex, id.xy, masked);
+}
diff --git a/src/gpu/effects/shaders.cc b/src/gpu/effects/shaders.cc
index 6ed82d5..625c5b6 100644
--- a/src/gpu/effects/shaders.cc
+++ b/src/gpu/effects/shaders.cc
@@ -111,6 +111,16 @@ const char* gen_grid_compute_wgsl =
 
     SafeGetAsset(AssetId::ASSET_SHADER_COMPUTE_GEN_GRID);
 
+#if !defined(STRIP_GPU_COMPOSITE)
+const char* gen_blend_compute_wgsl =
+
+    SafeGetAsset(AssetId::ASSET_SHADER_COMPUTE_GEN_BLEND);
+
+const char* gen_mask_compute_wgsl =
+
+    SafeGetAsset(AssetId::ASSET_SHADER_COMPUTE_GEN_MASK);
+#endif
+
 const char* vignette_shader_wgsl =
 
     SafeGetAsset(AssetId::ASSET_SHADER_VIGNETTE);
diff --git a/src/gpu/effects/shaders.h b/src/gpu/effects/shaders.h
index a0f91da..68b8834 100644
--- a/src/gpu/effects/shaders.h
+++ b/src/gpu/effects/shaders.h
@@ -21,3 +21,7 @@ extern const char* vignette_shader_wgsl;
 extern const char* gen_noise_compute_wgsl;
 extern const char* gen_perlin_compute_wgsl;
 extern const char* gen_grid_compute_wgsl;
+#if !defined(STRIP_GPU_COMPOSITE)
+extern const char* gen_blend_compute_wgsl;
+extern const char* gen_mask_compute_wgsl;
+#endif
diff --git a/src/gpu/texture_manager.cc b/src/gpu/texture_manager.cc
index 2b83f63..7aeb67a 100644
--- a/src/gpu/texture_manager.cc
+++ b/src/gpu/texture_manager.cc
@@ -21,6 +21,20 @@
 void TextureManager::init(WGPUDevice device, WGPUQueue queue) {
   device_ = device;
   queue_ = queue;
+
+#if !defined(STRIP_GPU_COMPOSITE)
+  // Create linear sampler for composite shaders
+  WGPUSamplerDescriptor sampler_desc = {};
+  sampler_desc.addressModeU = WGPUAddressMode_ClampToEdge;
+  sampler_desc.addressModeV = WGPUAddressMode_ClampToEdge;
+  sampler_desc.magFilter = WGPUFilterMode_Linear;
+  sampler_desc.minFilter = WGPUFilterMode_Linear;
+  sampler_desc.mipmapFilter = WGPUMipmapFilterMode_Linear;
+  sampler_desc.lodMinClamp = 0.0f;
+  sampler_desc.lodMaxClamp = 1.0f;
+  sampler_desc.maxAnisotropy = 1;
+  linear_sampler_ = wgpuDeviceCreateSampler(device_, &sampler_desc);
+#endif
 }
 
 void TextureManager::shutdown() {
@@ -36,6 +50,13 @@ void TextureManager::shutdown() {
     }
   }
   compute_pipelines_.clear();
+
+#if !defined(STRIP_GPU_COMPOSITE)
+  if (linear_sampler_) {
+    wgpuSamplerRelease(linear_sampler_);
+    linear_sampler_ = nullptr;
+  }
+#endif
 }
 
 void TextureManager::create_procedural_texture(
@@ -125,7 +146,7 @@ WGPUTextureView TextureManager::get_texture_view(const std::string& name) {
 
 WGPUComputePipeline TextureManager::get_or_create_compute_pipeline(
     const std::string& func_name, const char* shader_code,
-    size_t uniform_size) {
+    size_t uniform_size, int num_input_textures) {
   auto it = compute_pipelines_.find(func_name);
   if (it != compute_pipelines_.end()) {
     return it->second.pipeline;
@@ -143,22 +164,45 @@ WGPUComputePipeline TextureManager::get_or_create_compute_pipeline(
   WGPUShaderModule shader_module =
       wgpuDeviceCreateShaderModule(device_, &shader_desc);
 
-  // Bind group layout (storage texture + uniform)
-  WGPUBindGroupLayoutEntry bgl_entries[2] = {};
+  // Dynamic bind group layout
+  // Binding 0: output storage texture
+  // Binding 1: uniform buffer
+  // Binding 2 to (2 + num_input_textures - 1): input textures
+  // Binding (2 + num_input_textures): sampler (if inputs > 0)
+  const int max_entries = 2 + num_input_textures + (num_input_textures > 0 ? 1 : 0);
+  std::vector<WGPUBindGroupLayoutEntry> bgl_entries(max_entries);
+
+  // Binding 0: Output storage texture
   bgl_entries[0].binding = 0;
   bgl_entries[0].visibility = WGPUShaderStage_Compute;
   bgl_entries[0].storageTexture.access = WGPUStorageTextureAccess_WriteOnly;
   bgl_entries[0].storageTexture.format = WGPUTextureFormat_RGBA8Unorm;
   bgl_entries[0].storageTexture.viewDimension = WGPUTextureViewDimension_2D;
 
+  // Binding 1: Uniform buffer
   bgl_entries[1].binding = 1;
   bgl_entries[1].visibility = WGPUShaderStage_Compute;
   bgl_entries[1].buffer.type = WGPUBufferBindingType_Uniform;
   bgl_entries[1].buffer.minBindingSize = uniform_size;
 
+  // Binding 2+: Input textures
+  for (int i = 0; i < num_input_textures; ++i) {
+    bgl_entries[2 + i].binding = 2 + i;
+    bgl_entries[2 + i].visibility = WGPUShaderStage_Compute;
+    bgl_entries[2 + i].texture.sampleType = WGPUTextureSampleType_Float;
+    bgl_entries[2 + i].texture.viewDimension = WGPUTextureViewDimension_2D;
+  }
+
+  // Binding N: Sampler (if inputs exist)
+  if (num_input_textures > 0) {
+    bgl_entries[2 + num_input_textures].binding = 2 + num_input_textures;
+    bgl_entries[2 + num_input_textures].visibility = WGPUShaderStage_Compute;
+    bgl_entries[2 + num_input_textures].sampler.type = WGPUSamplerBindingType_Filtering;
+  }
+
   WGPUBindGroupLayoutDescriptor bgl_desc = {};
-  bgl_desc.entryCount = 2;
-  bgl_desc.entries = bgl_entries;
+  bgl_desc.entryCount = max_entries;
+  bgl_desc.entries = bgl_entries.data();
   WGPUBindGroupLayout bind_group_layout =
       wgpuDeviceCreateBindGroupLayout(device_, &bgl_desc);
 
@@ -181,7 +225,7 @@ WGPUComputePipeline TextureManager::get_or_create_compute_pipeline(
   wgpuShaderModuleRelease(shader_module);
 
   // Cache pipeline
-  ComputePipelineInfo info = {pipeline, shader_code, uniform_size};
+  ComputePipelineInfo info = {pipeline, shader_code, uniform_size, num_input_textures};
   compute_pipelines_[func_name] = info;
 
   return pipeline;
@@ -424,6 +468,188 @@ void TextureManager::create_gpu_grid_texture(
 #endif
 }
 
+#if !defined(STRIP_GPU_COMPOSITE)
+void TextureManager::dispatch_composite(
+    const std::string& func_name, WGPUTexture target,
+    const GpuProceduralParams& params, const void* uniform_data,
+    size_t uniform_size, const std::vector<WGPUTextureView>& input_views) {
+  auto it = compute_pipelines_.find(func_name);
+  if (it == compute_pipelines_.end()) {
+    return; // Pipeline not created yet
+  }
+
+  WGPUComputePipeline pipeline = it->second.pipeline;
+  int num_inputs = (int)input_views.size();
+
+  // Create uniform buffer
+  WGPUBufferDescriptor buf_desc = {};
+  buf_desc.size = uniform_size;
+  buf_desc.usage = WGPUBufferUsage_Uniform | WGPUBufferUsage_CopyDst;
+  buf_desc.mappedAtCreation = WGPUOptionalBool_True;
+  WGPUBuffer uniform_buf = wgpuDeviceCreateBuffer(device_, &buf_desc);
+  void* mapped = wgpuBufferGetMappedRange(uniform_buf, 0, uniform_size);
+  memcpy(mapped, uniform_data, uniform_size);
+  wgpuBufferUnmap(uniform_buf);
+
+  // Create storage texture view
+  WGPUTextureViewDescriptor view_desc = {};
+  view_desc.format = WGPUTextureFormat_RGBA8Unorm;
+  view_desc.dimension = WGPUTextureViewDimension_2D;
+  view_desc.mipLevelCount = 1;
+  view_desc.arrayLayerCount = 1;
+  WGPUTextureView target_view = wgpuTextureCreateView(target, &view_desc);
+
+  // Dynamic bind group
+  const int max_entries = 2 + num_inputs + (num_inputs > 0 ? 1 : 0);
+  std::vector<WGPUBindGroupEntry> bg_entries(max_entries);
+
+  // Binding 0: Output texture
+  bg_entries[0].binding = 0;
+  bg_entries[0].textureView = target_view;
+
+  // Binding 1: Uniform buffer
+  bg_entries[1].binding = 1;
+  bg_entries[1].buffer = uniform_buf;
+  bg_entries[1].size = uniform_size;
+
+  // Binding 2+: Input textures
+  for (int i = 0; i < num_inputs; ++i) {
+    bg_entries[2 + i].binding = 2 + i;
+    bg_entries[2 + i].textureView = input_views[i];
+  }
+
+  // Binding N: Sampler
+  if (num_inputs > 0) {
+    bg_entries[2 + num_inputs].binding = 2 + num_inputs;
+    bg_entries[2 + num_inputs].sampler = linear_sampler_;
+  }
+
+  // Create bind group layout (must match pipeline)
+  const int layout_entries_count = 2 + num_inputs + (num_inputs > 0 ? 1 : 0);
+  std::vector<WGPUBindGroupLayoutEntry> bgl_entries(layout_entries_count);
+
+  bgl_entries[0].binding = 0;
+  bgl_entries[0].visibility = WGPUShaderStage_Compute;
+  bgl_entries[0].storageTexture.access = WGPUStorageTextureAccess_WriteOnly;
+  bgl_entries[0].storageTexture.format = WGPUTextureFormat_RGBA8Unorm;
+  bgl_entries[0].storageTexture.viewDimension = WGPUTextureViewDimension_2D;
+
+  bgl_entries[1].binding = 1;
+  bgl_entries[1].visibility = WGPUShaderStage_Compute;
+  bgl_entries[1].buffer.type = WGPUBufferBindingType_Uniform;
+  bgl_entries[1].buffer.minBindingSize = uniform_size;
+
+  for (int i = 0; i < num_inputs; ++i) {
+    bgl_entries[2 + i].binding = 2 + i;
+    bgl_entries[2 + i].visibility = WGPUShaderStage_Compute;
+    bgl_entries[2 + i].texture.sampleType = WGPUTextureSampleType_Float;
+    bgl_entries[2 + i].texture.viewDimension = WGPUTextureViewDimension_2D;
+  }
+
+  if (num_inputs > 0) {
+    bgl_entries[2 + num_inputs].binding = 2 + num_inputs;
+    bgl_entries[2 + num_inputs].visibility = WGPUShaderStage_Compute;
+    bgl_entries[2 + num_inputs].sampler.type = WGPUSamplerBindingType_Filtering;
+  }
+
+  WGPUBindGroupLayoutDescriptor bgl_desc = {};
+  bgl_desc.entryCount = layout_entries_count;
+  bgl_desc.entries = bgl_entries.data();
+  WGPUBindGroupLayout bind_group_layout =
+      wgpuDeviceCreateBindGroupLayout(device_, &bgl_desc);
+
+  WGPUBindGroupDescriptor bg_desc = {};
+  bg_desc.layout = bind_group_layout;
+  bg_desc.entryCount = max_entries;
+  bg_desc.entries = bg_entries.data();
+  WGPUBindGroup bind_group = wgpuDeviceCreateBindGroup(device_, &bg_desc);
+
+  // Dispatch compute
+  WGPUCommandEncoderDescriptor enc_desc = {};
+  WGPUCommandEncoder encoder =
+      wgpuDeviceCreateCommandEncoder(device_, &enc_desc);
+  WGPUComputePassEncoder pass =
+      wgpuCommandEncoderBeginComputePass(encoder, nullptr);
+  wgpuComputePassEncoderSetPipeline(pass, pipeline);
+  wgpuComputePassEncoderSetBindGroup(pass, 0, bind_group, 0, nullptr);
+  wgpuComputePassEncoderDispatchWorkgroups(pass, (params.width + 7) / 8,
+                                           (params.height + 7) / 8, 1);
+  wgpuComputePassEncoderEnd(pass);
+
+  WGPUCommandBufferDescriptor cmd_desc = {};
+  WGPUCommandBuffer cmd = wgpuCommandEncoderFinish(encoder, &cmd_desc);
+  wgpuQueueSubmit(queue_, 1, &cmd);
+
+  // Cleanup
+  wgpuCommandBufferRelease(cmd);
+  wgpuCommandEncoderRelease(encoder);
+  wgpuComputePassEncoderRelease(pass);
+  wgpuBindGroupRelease(bind_group);
+  wgpuBindGroupLayoutRelease(bind_group_layout);
+  wgpuBufferRelease(uniform_buf);
+  wgpuTextureViewRelease(target_view);
+}
+
+void TextureManager::create_gpu_composite_texture(
+    const std::string& name, const std::string& shader_func,
+    const char* shader_code, const void* uniform_data, size_t uniform_size,
+    int width, int height, const std::vector<std::string>& input_names) {
+  // Create pipeline if needed
+  get_or_create_compute_pipeline(shader_func, shader_code, uniform_size,
+                                 (int)input_names.size());
+
+  // Resolve input texture views
+  std::vector<WGPUTextureView> input_views;
+  input_views.reserve(input_names.size());
+  for (const auto& input_name : input_names) {
+    WGPUTextureView view = get_texture_view(input_name);
+    if (!view) {
+      fprintf(stderr, "Error: Input texture not found: %s\n",
+              input_name.c_str());
+      return;
+    }
+    input_views.push_back(view);
+  }
+
+  // Create output texture
+  WGPUTextureDescriptor tex_desc = {};
+  tex_desc.usage =
+      WGPUTextureUsage_StorageBinding | WGPUTextureUsage_TextureBinding;
+  tex_desc.dimension = WGPUTextureDimension_2D;
+  tex_desc.size = {(uint32_t)width, (uint32_t)height, 1};
+  tex_desc.format = WGPUTextureFormat_RGBA8Unorm;
+  tex_desc.mipLevelCount = 1;
+  tex_desc.sampleCount = 1;
+  WGPUTexture texture = wgpuDeviceCreateTexture(device_, &tex_desc);
+
+  // Dispatch composite shader
+  GpuProceduralParams params = {width, height, nullptr, 0};
+  dispatch_composite(shader_func, texture, params, uniform_data, uniform_size,
+                    input_views);
+
+  // Create view
+  WGPUTextureViewDescriptor view_desc = {};
+  view_desc.format = WGPUTextureFormat_RGBA8Unorm;
+  view_desc.dimension = WGPUTextureViewDimension_2D;
+  view_desc.mipLevelCount = 1;
+  view_desc.arrayLayerCount = 1;
+  WGPUTextureView view = wgpuTextureCreateView(texture, &view_desc);
+
+  // Store
+  GpuTexture gpu_tex;
+  gpu_tex.texture = texture;
+  gpu_tex.view = view;
+  gpu_tex.width = width;
+  gpu_tex.height = height;
+  textures_[name] = gpu_tex;
+
+#if !defined(STRIP_ALL)
+  printf("Generated GPU composite texture: %s (%dx%d, %zu inputs)\n",
+         name.c_str(), width, height, input_names.size());
+#endif
+}
+#endif  // !defined(STRIP_GPU_COMPOSITE)
+
 #if !defined(STRIP_ALL)
 WGPUTextureView TextureManager::get_or_generate_gpu_texture(
     const std::string& name, const GpuProceduralParams& params) {
diff --git a/src/gpu/texture_manager.h b/src/gpu/texture_manager.h
index 63c2947..86d1f63 100644
--- a/src/gpu/texture_manager.h
+++ b/src/gpu/texture_manager.h
@@ -51,6 +51,17 @@ class TextureManager {
   void create_gpu_grid_texture(const std::string& name,
                                const GpuProceduralParams& params);
 
+#if !defined(STRIP_GPU_COMPOSITE)
+  // GPU composite generation (multi-input textures)
+  void create_gpu_composite_texture(const std::string& name,
+                                    const std::string& shader_func,
+                                    const char* shader_code,
+                                    const void* uniform_data,
+                                    size_t uniform_size,
+                                    int width, int height,
+                                    const std::vector<std::string>& input_names);
+#endif
+
 #if !defined(STRIP_ALL)
   // On-demand lazy generation (stripped in final builds)
   WGPUTextureView get_or_generate_gpu_texture(const std::string& name,
@@ -65,17 +76,29 @@ class TextureManager {
     WGPUComputePipeline pipeline;
     const char* shader_code;
     size_t uniform_size;
+    int num_input_textures;
   };
 
   WGPUComputePipeline get_or_create_compute_pipeline(const std::string& func_name,
                                                       const char* shader_code,
-                                                      size_t uniform_size);
+                                                      size_t uniform_size,
+                                                      int num_input_textures = 0);
   void dispatch_compute(const std::string& func_name, WGPUTexture target,
                        const GpuProceduralParams& params, const void* uniform_data,
                        size_t uniform_size);
 
+#if !defined(STRIP_GPU_COMPOSITE)
+  void dispatch_composite(const std::string& func_name, WGPUTexture target,
+                         const GpuProceduralParams& params,
+                         const void* uniform_data, size_t uniform_size,
+                         const std::vector<WGPUTextureView>& input_views);
+#endif
+
   WGPUDevice device_;
   WGPUQueue queue_;
   std::map<std::string, GpuTexture> textures_;
   std::map<std::string, ComputePipelineInfo> compute_pipelines_;
+#if !defined(STRIP_GPU_COMPOSITE)
+  WGPUSampler linear_sampler_;
+#endif
 };
diff --git a/src/tests/test_gpu_composite.cc b/src/tests/test_gpu_composite.cc
new file mode 100644
index 0000000..e5ac788
--- /dev/null
+++ b/src/tests/test_gpu_composite.cc
@@ -0,0 +1,124 @@
+// This file is part of the 64k demo project.
+// Tests GPU composite texture generation (Phase 4).
+
+#include "gpu/gpu.h"
+#include "gpu/texture_manager.h"
+#include "platform/platform.h"
+#include <cstdint>
+#include <cstdio>
+#include <vector>
+
+#if !defined(STRIP_GPU_COMPOSITE)
+
+int main() {
+  printf("GPU Composite Test: Starting...\n");
+
+  // Initialize GPU
+  PlatformState platform = platform_init(false, 256, 256);
+  if (!platform.window) {
+    fprintf(stderr, "Error: Failed to create window\n");
+    return 1;
+  }
+
+  gpu_init(&platform);
+  const GpuContext* ctx = gpu_get_context();
+
+  extern void InitShaderComposer();
+  InitShaderComposer();
+
+  TextureManager tex_mgr;
+  tex_mgr.init(ctx->device, ctx->queue);
+
+  // Create base textures
+  float noise_params_a[2] = {1234.0f, 4.0f};
+  GpuProceduralParams noise_a = {256, 256, noise_params_a, 2};
+  tex_mgr.create_gpu_noise_texture("noise_a", noise_a);
+
+  float noise_params_b[2] = {5678.0f, 8.0f};
+  GpuProceduralParams noise_b = {256, 256, noise_params_b, 2};
+  tex_mgr.create_gpu_noise_texture("noise_b", noise_b);
+
+  float grid_params[2] = {32.0f, 2.0f};
+  GpuProceduralParams grid = {256, 256, grid_params, 2};
+  tex_mgr.create_gpu_grid_texture("grid", grid);
+
+  printf("SUCCESS: Base textures created (noise_a, noise_b, grid)\n");
+
+  // Test blend composite
+  extern const char* gen_blend_compute_wgsl;
+  struct {
+    uint32_t width, height;
+    float blend_factor, _pad0;
+  } blend_uni = {256, 256, 0.5f, 0.0f};
+
+  std::vector<std::string> blend_inputs = {"noise_a", "noise_b"};
+  tex_mgr.create_gpu_composite_texture("blended", "gen_blend",
+                                       gen_blend_compute_wgsl, &blend_uni,
+                                       sizeof(blend_uni), 256, 256, blend_inputs);
+
+  WGPUTextureView blended_view = tex_mgr.get_texture_view("blended");
+  if (!blended_view) {
+    fprintf(stderr, "Error: Blended texture not created\n");
+    tex_mgr.shutdown();
+    gpu_shutdown();
+    return 1;
+  }
+  printf("SUCCESS: Blend composite created (noise_a + noise_b)\n");
+
+  // Test mask composite
+  extern const char* gen_mask_compute_wgsl;
+  struct {
+    uint32_t width, height;
+  } mask_uni = {256, 256};
+
+  std::vector<std::string> mask_inputs = {"noise_a", "grid"};
+  tex_mgr.create_gpu_composite_texture("masked", "gen_mask", gen_mask_compute_wgsl,
+                                       &mask_uni, sizeof(mask_uni), 256, 256,
+                                       mask_inputs);
+
+  WGPUTextureView masked_view = tex_mgr.get_texture_view("masked");
+  if (!masked_view) {
+    fprintf(stderr, "Error: Masked texture not created\n");
+    tex_mgr.shutdown();
+    gpu_shutdown();
+    return 1;
+  }
+  printf("SUCCESS: Mask composite created (noise_a * grid)\n");
+
+  // Test multi-stage composite (composite of composite)
+  struct {
+    uint32_t width, height;
+    float blend_factor, _pad0;
+  } blend2_uni = {256, 256, 0.7f, 0.0f};
+
+  std::vector<std::string> blend2_inputs = {"blended", "masked"};
+  tex_mgr.create_gpu_composite_texture("final", "gen_blend",
+                                       gen_blend_compute_wgsl, &blend2_uni,
+                                       sizeof(blend2_uni), 256, 256, blend2_inputs);
+
+  WGPUTextureView final_view = tex_mgr.get_texture_view("final");
+  if (!final_view) {
+    fprintf(stderr, "Error: Multi-stage composite not created\n");
+    tex_mgr.shutdown();
+    gpu_shutdown();
+    return 1;
+  }
+  printf("SUCCESS: Multi-stage composite (composite of composites)\n");
+
+  // Cleanup
+  tex_mgr.shutdown();
+  gpu_shutdown();
+  platform_shutdown(&platform);
+
+  printf("All GPU composite tests passed!\n");
+  return 0;
+}
+
+#else
+
+int main() {
+  printf("GPU Composite Test: SKIPPED (STRIP_GPU_COMPOSITE defined)\n");
+  return 0;
+}
+
+#endif