Basic working circle render pipeline

build.zig

@@ -56,4 +56,39 @@ pub fn build(b: *std.Build) !void {
         const run_cmd = b.addRunArtifact(exe);
         run_step.dependOn(&run_cmd.step);
     }
+
+    const exe = b.addExecutable(.{
+        .name = "circle",
+        .root_module = b.createModule(.{
+            .root_source_file = b.path("src/circle.zig"),
+            .target = target,
+            .optimize = optimize,
+            .imports = &.{},
+        }),
+    });
+
+    exe.root_module.addIncludePath(b.path("libs/wgpu-native/include"));
+    exe.root_module.addLibraryPath(b.path("libs/wgpu-native/lib"));
+    exe.root_module.addObjectFile(b.path("libs/wgpu-native/lib/libwgpu_native.a"));
+
+    // Platform-specific system frameworks needed by wgpu-native
+    if (t.os.tag == .macos) {
+        exe.root_module.linkFramework("Metal", .{});
+        exe.root_module.linkFramework("QuartzCore", .{});
+        exe.root_module.linkFramework("Foundation", .{});
+        exe.root_module.linkFramework("CoreGraphics", .{});
+    } else if (t.os.tag == .windows) {
+        exe.root_module.linkSystemLibrary("d3d12", .{});
+        exe.root_module.linkSystemLibrary("dxgi", .{});
+        exe.root_module.linkSystemLibrary("user32", .{});
+    } else {
+        exe.root_module.linkSystemLibrary("vulkan", .{});
+        exe.root_module.linkSystemLibrary("gcc_s", .{});
+    }
+
+    b.installArtifact(exe);
+
+    const run_step = b.step("circle", "Run circle");
+    const run_cmd = b.addRunArtifact(exe);
+    run_step.dependOn(&run_cmd.step);
 }

src/GpuRender.zig

@@ -0,0 +1,168 @@
+const std = @import("std");
+const c = @import("utils.zig").c;
+const sv = @import("utils.zig").sv;
+const GpuAllocator = @import("GpuAllocator.zig");
+const GpuBuffer = @import("GpuBuffer.zig");
+const GpuDevice = @import("GpuDevice.zig");
+
+pub const Binding = struct {
+    element_size: u32 = 0,
+};
+
+pub const RenderDef = struct {
+    bindings: []const Binding = &.{},
+    /// The surface texture format we are rendering to (e.g., BGRA8Unorm)
+    texture_format: c.WGPUTextureFormat,
+    /// The names of the entry points inside your WGSL code
+    vertex_entry: []const u8 = "vs_main",
+    fragment_entry: []const u8 = "fs_main",
+    /// Primitive topology, default to triangle list
+    topology: c.WGPUPrimitiveTopology = c.WGPUPrimitiveTopology_TriangleList,
+};
+
+pip: c.WGPURenderPipeline,
+def: RenderDef,
+
+pub fn init(device: GpuDevice, wgsl: []const u8, def: RenderDef) !@This() {
+    var wgsl_src = c.WGPUShaderSourceWGSL{
+        .chain = .{ .sType = c.WGPUSType_ShaderSourceWGSL },
+        .code = sv(wgsl),
+    };
+    const shader = c.wgpuDeviceCreateShaderModule(device.device, &.{
+        .nextInChain = @ptrCast(&wgsl_src),
+    }) orelse return error.Shader;
+    defer c.wgpuShaderModuleRelease(shader);
+
+    // 1. Setup the Color Target State (where the fragment shader outputs)
+    const blend = c.WGPUBlendState{
+        .color = .{ .operation = c.WGPUBlendOperation_Add, .srcFactor = c.WGPUBlendFactor_SrcAlpha, .dstFactor = c.WGPUBlendFactor_OneMinusSrcAlpha },
+        .alpha = .{ .operation = c.WGPUBlendOperation_Add, .srcFactor = c.WGPUBlendFactor_One, .dstFactor = c.WGPUBlendFactor_Zero },
+    };
+
+    const color_target = c.WGPUColorTargetState{
+        .format = def.texture_format,
+        .blend = &blend,
+        .writeMask = c.WGPUColorWriteMask_All,
+    };
+
+    // 2. Setup the Fragment State
+    const fragment_state = c.WGPUFragmentState{
+        .module = shader,
+        .entryPoint = sv(def.fragment_entry),
+        .targetCount = 1,
+        .targets = &color_target,
+    };
+
+    // 3. Compile the Complete Render Pipeline
+    const pip = c.wgpuDeviceCreateRenderPipeline(device.device, &.{
+        .vertex = .{
+            .module = shader,
+            .entryPoint = sv(def.vertex_entry),
+            .bufferCount = 0, // Assuming procedural drawing (like our circle!)
+        },
+        .primitive = .{
+            .topology = def.topology,
+            .stripIndexFormat = c.WGPUIndexFormat_Undefined,
+            .frontFace = c.WGPUFrontFace_CCW,
+            .cullMode = c.WGPUCullMode_None,
+        },
+        .multisample = .{
+            .count = 1,
+            .mask = 0xFFFFFFFF,
+            .alphaToCoverageEnabled = 0,
+        },
+        .fragment = &fragment_state,
+    }) orelse return error.Pipeline;
+
+    return .{
+        .pip = pip,
+        .def = def,
+    };
+}
+
+pub fn deinit(self: @This()) void {
+    c.wgpuRenderPipelineRelease(self.pip);
+}
+
+/// Execute the render pass targeting a specific frame texture view.
+/// Passes bind groups via a tuple exactly like your original compute setup.
+pub fn draw(
+    self: @This(),
+    gloc: GpuAllocator,
+    target_view: c.WGPUTextureView,
+    vertex_count: u32,
+    args: anytype,
+) !void {
+    const type_info = @typeInfo(@TypeOf(args));
+    if (type_info != .@"struct" or !type_info.@"struct".is_tuple)
+        @compileError("Expected a tuple of GpuBuffers for args. E.g. .{ uniform_buf }");
+
+    const fields = type_info.@"struct".fields;
+    if (fields.len != self.def.bindings.len)
+        return error.InvalidArgumentCount;
+
+    var entries_buf: [32]c.WGPUBindGroupEntry = undefined;
+
+    inline for (fields, 0..) |field, i| {
+        const buf = @field(args, field.name);
+        if (@TypeOf(buf) != GpuBuffer) {
+            @compileError("All arguments in the tuple must be of type GpuBuffer");
+        }
+        entries_buf[i] = .{
+            .binding = @intCast(i),
+            .buffer = buf.raw,
+            .offset = 0,
+            .size = buf.size,
+        };
+    }
+
+    const entries = entries_buf[0..fields.len];
+
+    // Create Render Bind Group from layout
+    const bgl = c.wgpuRenderPipelineGetBindGroupLayout(self.pip, 0);
+    defer c.wgpuBindGroupLayoutRelease(bgl);
+
+    const bg = c.wgpuDeviceCreateBindGroup(gloc.device.device, &.{
+        .layout = bgl,
+        .entries = entries.ptr,
+        .entryCount = @intCast(entries.len),
+    }) orelse return error.BindGroup;
+    defer c.wgpuBindGroupRelease(bg);
+
+    // Encode Render Command
+    const enc = c.wgpuDeviceCreateCommandEncoder(gloc.device.device, null) orelse return error.Encoder;
+    defer c.wgpuCommandEncoderRelease(enc);
+
+    const color_attachment = c.WGPURenderPassColorAttachment{
+        .view = target_view,
+        .resolveTarget = null,
+        .loadOp = c.WGPULoadOp_Clear,
+        .storeOp = c.WGPUStoreOp_Store,
+        .clearValue = .{ .r = 0.1, .g = 0.1, .b = 0.1, .a = 1.0 },
+        .depthSlice = c.WGPU_DEPTH_SLICE_UNDEFINED,
+    };
+
+    const pass_desc = c.WGPURenderPassDescriptor{
+        .colorAttachmentCount = 1,
+        .colorAttachments = &color_attachment,
+        .depthStencilAttachment = null,
+    };
+
+    const pass = c.wgpuCommandEncoderBeginRenderPass(enc, &pass_desc);
+    c.wgpuRenderPassEncoderSetPipeline(pass, self.pip);
+
+    if (fields.len > 0) {
+        c.wgpuRenderPassEncoderSetBindGroup(pass, 0, bg, 0, null);
+    }
+
+    // Draw! (Instead of Compute Dispatch)
+    c.wgpuRenderPassEncoderDraw(pass, vertex_count, 1, 0, 0);
+
+    c.wgpuRenderPassEncoderEnd(pass);
+    c.wgpuRenderPassEncoderRelease(pass);
+
+    const cmd = c.wgpuCommandEncoderFinish(enc, null);
+    defer c.wgpuCommandBufferRelease(cmd);
+
+    c.wgpuQueueSubmit(gloc.device.queue, 1, &cmd);
+}

src/circle.zig

@@ -0,0 +1,120 @@
+const std = @import("std");
+const gpu = @import("lib.zig");
+const c = @import("utils.zig").c;
+const sv = @import("utils.zig").sv;
+const GpuDevice = gpu.GpuDevice;
+const GpuArena = gpu.GpuArena;
+const GpuBuffer = gpu.GpuBuffer;
+const GpuRender = gpu.GpuRender;
+
+pub fn main(init: std.process.Init) !void {
+    const allocator = init.gpa;
+
+    // 1. Open the raw headless GPU Device you shared
+    const device = try GpuDevice.init(.{});
+    defer device.deinit();
+
+    var grena = GpuArena.init(allocator, device);
+    defer grena.deinit();
+    const gloc = grena.gpuAllocator();
+
+    const width: u32 = 512;
+    const height: u32 = 512;
+    // We use standard RGBA8Unorm format for an offscreen image target
+    const render_format = c.WGPUTextureFormat_RGBA8Unorm;
+
+    // 2. Load our Render Pipeline (Procedural Triangle Strip)
+    const circle_rp = try GpuRender.init(
+        device,
+        @embedFile("shaders/circle.wgsl"),
+        .{
+            .bindings = &.{},
+            .texture_format = render_format,
+            .topology = c.WGPUPrimitiveTopology_TriangleStrip,
+        },
+    );
+    defer circle_rp.deinit();
+
+    // 3. Create the offscreen VRAM texture to render into
+    const texture_desc = c.WGPUTextureDescriptor{
+        .nextInChain = null,
+        .label = sv("Offscreen Render Target"),
+        .usage = c.WGPUTextureUsage_RenderAttachment | c.WGPUTextureUsage_CopySrc,
+        .dimension = c.WGPUTextureDimension_2D,
+        .size = .{ .width = width, .height = height, .depthOrArrayLayers = 1 },
+        .format = render_format,
+        .mipLevelCount = 1,
+        .sampleCount = 1,
+        .viewFormatCount = 0,
+        .viewFormats = null,
+    };
+    const target_texture = c.wgpuDeviceCreateTexture(device.device, &texture_desc) orelse return error.Texture;
+    defer c.wgpuTextureRelease(target_texture);
+
+    const target_view = c.wgpuTextureCreateView(target_texture, null) orelse return error.View;
+    defer c.wgpuTextureViewRelease(target_view);
+
+    // 4. Create a staging buffer to pull pixels from VRAM to CPU
+    // 4 bytes per pixel (RGBA8)
+    const row_bytes = width * 4;
+    const buffer_bytes = row_bytes * height;
+
+    // Create a regular GpuBuffer set up to receive texture copy transfers
+    const cpu_staging_buf = try GpuBuffer.init(gloc, buffer_bytes, .initMany(&.{ .CopyDst, .CopySrc }));
+
+    // 5. Draw the Circle Frame into the texture view!
+    try circle_rp.draw(gloc, target_view, 4, .{});
+
+    // 6. Copy the texture data into our CPU staging buffer
+    const enc = c.wgpuDeviceCreateCommandEncoder(device.device, null) orelse return error.Encoder;
+    defer c.wgpuCommandEncoderRelease(enc);
+
+    const src_copy = c.WGPUTexelCopyTextureInfo{
+        .texture = target_texture,
+        .mipLevel = 0,
+        .origin = .{ .x = 0, .y = 0, .z = 0 },
+        .aspect = c.WGPUTextureAspect_All,
+    };
+    const dst_copy = c.WGPUTexelCopyBufferInfo{
+        .buffer = cpu_staging_buf.raw,
+        .layout = .{
+            .offset = 0,
+            .bytesPerRow = row_bytes,
+            .rowsPerImage = height,
+        },
+    };
+    const copy_size = c.WGPUExtent3D{ .width = width, .height = height, .depthOrArrayLayers = 1 };
+
+    c.wgpuCommandEncoderCopyTextureToBuffer(enc, &src_copy, &dst_copy, &copy_size);
+
+    const cmd = c.wgpuCommandEncoderFinish(enc, null);
+    defer c.wgpuCommandBufferRelease(cmd);
+    c.wgpuQueueSubmit(device.queue, 1, &cmd);
+
+    // 7. Map and read the raw image bytes back to CPU
+    // (This uses whatever slice-reading helpers your `GpuBuffer` wrapper provides)
+    const pixels = try cpu_staging_buf.read(allocator, u8);
+    defer allocator.free(pixels);
+
+    // Now you have the raw binary image data! Let's output a simple Netpbm PPM image file
+    // so you can actually open and look at your rendered circle.
+    try savePpm(init.io, "circle.ppm", width, height, pixels);
+    std.debug.print("Successfully rendered circle to circle.ppm!\n", .{});
+}
+
+fn savePpm(io: std.Io, filename: []const u8, w: u32, h: u32, rgba_pixels: []const u8) !void {
+    const file = try std.Io.Dir.cwd().createFile(io, filename, .{});
+    defer file.close(io);
+
+    var buf: [255]u8 = undefined;
+    var writer = file.writer(io, &buf);
+
+    // PPM Header: P6 format means raw RGB bytes
+    try writer.interface.print("P6\n{d} {d}\n255\n", .{ w, h });
+
+    // Strip Alpha channel when writing out to standard RGB PPM format
+    var i: usize = 0;
+    while (i < rgba_pixels.len) : (i += 4) {
+        try writer.interface.writeAll(rgba_pixels[i .. i + 3]);
+    }
+}

src/lib.zig

@@ -3,3 +3,4 @@ pub const GpuArena = @import("GpuArena.zig");
 pub const GpuBuffer = @import("GpuBuffer.zig");
 pub const GpuDevice = @import("GpuDevice.zig");
 pub const GpuCompute = @import("GpuCompute.zig");
+pub const GpuRender = @import("GpuRender.zig");

src/shaders/add.wgsl

@@ -0,0 +1,24 @@
+enable f16;
+
+@group(0) @binding(0) var<storage, read> A: array<f16>;
+@group(0) @binding(1) var<storage, read> B: array<f16>;
+@group(0) @binding(2) var<storage, read_write> C: array<f16>;
+@group(0) @binding(3) var<uniform> size: u32; 
+
+@compute @workgroup_size(256)
+fn main(
+    @builtin(global_invocation_id) global_id : vec3<u32>,
+    @builtin(num_workgroups) num_workgroups: vec3<u32>
+) {
+    // 1. Calculate the total number of threads across the entire grid
+    let total_threads = num_workgroups.x * 256u; 
+    
+    // 2. Start at this thread's unique global ID
+    var index = global_id.x;
+    
+    // 3. Stride through the tensor elements
+    while (index < size) {
+        C[index] = A[index] + B[index];
+        index += total_threads; // Jump forward by the total thread count
+    }
+}

src/shaders/circle.wgsl

@@ -0,0 +1,39 @@
+struct VertexOutput {
+    @builtin(position) position: vec4f,
+    @location(0) uv: vec2f,
+};
+
+@vertex
+fn vs_main(@builtin(vertex_index) vertex_index: u32) -> VertexOutput {
+    var output: VertexOutput;
+    // Hardcoded fullscreen quad layout using 4 vertices (Triangle Strip)
+    // Indexes: 0: Top-Left, 1: Bottom-Left, 2: Top-Right, 3: Bottom-Right
+    var pos = array<vec2f, 4>(
+        vec2f(-1.0,  1.0),
+        vec2f(-1.0, -1.0),
+        vec2f( 1.0,  1.0),
+        vec2f( 1.0, -1.0)
+    );
+    
+    output.position = vec4f(pos[vertex_index], 0.0, 1.0);
+    output.uv = pos[vertex_index]; // Ranges cleanly from -1.0 to 1.0
+    return output;
+}
+
+@fragment
+fn fs_main(input: VertexOutput) -> @location(0) vec4f {
+    // Distance from the center (0,0)
+    let distance = length(input.uv);
+    let radius = 0.5;
+    
+    // Smooth out pixel edges (anti-aliasing)
+    let edge_softness = 0.005;
+    let alpha = 1.0 - smoothstep(radius - edge_softness, radius + edge_softness, distance);
+    
+    if (alpha <= 0.0) {
+        discard; 
+    }
+    
+    // Draw a sharp/smooth red circle
+    return vec4f(1.0, 0.3, 0.3, alpha);
+}