From 7de47a8f2d008d7254e1774ec096331c26c7ecb1 Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Lucas=20Roman=C3=B3?=
 <9062026+lucasromanosantos@users.noreply.github.com>
Date: Sun, 17 Jul 2022 12:54:51 -0300
Subject: [PATCH] examples: add map-async example
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

Signed-off-by: Lucas Romanó <9062026+lucasromanosantos@users.noreply.github.com>
---
 build.zig                    |  1 +
 examples/LICENSE             |  1 +
 examples/map-async/main.wgsl | 16 +++++++
 examples/map-async/main.zig  | 86 ++++++++++++++++++++++++++++++++++++
 4 files changed, 104 insertions(+)
 create mode 100644 examples/map-async/main.wgsl
 create mode 100644 examples/map-async/main.zig

diff --git a/build.zig b/build.zig
index 5da56433..09c47880 100644
--- a/build.zig
+++ b/build.zig
@@ -46,6 +46,7 @@ pub fn build(b: *std.build.Builder) void {
         .{ .name = "textured-cube", .packages = &[_]Pkg{ Packages.zmath, Packages.zigimg } },
         .{ .name = "ecs-app", .packages = &[_]Pkg{} },
         .{ .name = "image-blur", .packages = &[_]Pkg{Packages.zigimg} },
+        .{ .name = "map-async", .packages = &[_]Pkg{} },
         // NOTE: examples with std_platform_only should be placed at last
         .{ .name = "gkurve", .packages = &[_]Pkg{ Packages.zmath, Packages.zigimg, freetype.pkg }, .std_platform_only = true },
     }) |example| {
diff --git a/examples/LICENSE b/examples/LICENSE
index c7e6fc5d..e30d676f 100644
--- a/examples/LICENSE
+++ b/examples/LICENSE
@@ -7,3 +7,4 @@ The following examples have been ported from https://github.com/austinEng/webgpu
 * ./textured-cube
 * ./fractal-cube
 * ./image-blur
+* ./map-async
diff --git a/examples/map-async/main.wgsl b/examples/map-async/main.wgsl
new file mode 100644
index 00000000..dd867eaf
--- /dev/null
+++ b/examples/map-async/main.wgsl
@@ -0,0 +1,16 @@
+@group(0) @binding(0) var<storage, write> output: array<f32>;
+
+@stage(compute) @workgroup_size(64, 1, 1)
+fn main(
+  @builtin(global_invocation_id)
+  global_id : vec3<u32>,
+
+  @builtin(local_invocation_id)
+  local_id : vec3<u32>,
+) {
+  if (global_id.x >= arrayLength(&output)) {
+    return;
+  }
+  output[global_id.x] =
+    f32(global_id.x) * 1000. + f32(local_id.x);
+}
diff --git a/examples/map-async/main.zig b/examples/map-async/main.zig
new file mode 100644
index 00000000..00b918eb
--- /dev/null
+++ b/examples/map-async/main.zig
@@ -0,0 +1,86 @@
+const std = @import("std");
+const mach = @import("mach");
+const gpu = @import("gpu");
+
+pub const App = @This();
+
+const workgroup_size = 64;
+const buffer_size = 1000;
+
+pub fn init(_: *App, core: *mach.Core) !void {
+    const output = core.device.createBuffer(&.{
+        .usage = .{ .storage = true, .copy_src = true },
+        .size = buffer_size * @sizeOf(f32),
+        .mapped_at_creation = false,
+    });
+
+    const staging = core.device.createBuffer(&.{
+        .usage = .{ .map_read = true, .copy_dst = true },
+        .size = buffer_size * @sizeOf(f32),
+        .mapped_at_creation = false,
+    });
+
+    const compute_module = core.device.createShaderModule(&.{
+        .label = "shader module",
+        .code = .{ .wgsl = @embedFile("main.wgsl") },
+    });
+
+    const compute_pipeline = core.device.createComputePipeline(&gpu.ComputePipeline.Descriptor{ .compute = gpu.ProgrammableStageDescriptor{
+        .module = compute_module,
+        .entry_point = "main",
+    } });
+
+    const compute_bind_group = core.device.createBindGroup(&gpu.BindGroup.Descriptor{
+        .layout = compute_pipeline.getBindGroupLayout(0),
+        .entries = &[_]gpu.BindGroup.Entry{
+            gpu.BindGroup.Entry.buffer(0, output, 0, buffer_size),
+        },
+    });
+
+    compute_module.release();
+
+    const encoder = core.device.createCommandEncoder(null);
+
+    const compute_pass = encoder.beginComputePass(null);
+    compute_pass.setPipeline(compute_pipeline);
+    compute_pass.setBindGroup(0, compute_bind_group, &.{});
+    compute_pass.dispatch(try std.math.divCeil(u32, buffer_size, workgroup_size), 1, 1);
+    compute_pass.end();
+
+    encoder.copyBufferToBuffer(output, 0, staging, 0, buffer_size);
+
+    var command = encoder.finish(null);
+    encoder.release();
+
+    var response: gpu.Buffer.MapAsyncStatus = undefined;
+    var callback = gpu.Buffer.MapCallback.init(*gpu.Buffer.MapAsyncStatus, &response, (struct {
+        pub fn callback(ctx: *gpu.Buffer.MapAsyncStatus, callback_response: gpu.Buffer.MapAsyncStatus) void {
+            ctx.* = callback_response;
+        }
+    }).callback);
+
+    var queue = core.device.getQueue();
+    queue.submit(&.{command});
+
+    staging.mapAsync(gpu.Buffer.MapMode.read, 0, buffer_size, &callback);
+    while (true) {
+        if (response == gpu.Buffer.MapAsyncStatus.success) {
+            break;
+        } else {
+            core.device.tick();
+        }
+    }
+
+    const staging_mapped = staging.getConstMappedRange(f32, 0, buffer_size / @sizeOf(f32));
+    for (staging_mapped) |v| {
+        std.debug.print("{d} ", .{v});
+    }
+    std.debug.print("\n", .{});
+    staging.unmap();
+}
+
+pub fn deinit(_: *App, _: *mach.Core) void {}
+
+pub fn update(_: *App, core: *mach.Core) !void {
+    core.setShouldClose(true);
+}