Skip to content

Commit

Permalink
get index of vertices in the view
Browse files Browse the repository at this point in the history
  • Loading branch information
@sid597
sid597 committed Sep 12, 2024
1 parent e50d180 commit 995476b
Show file tree
Hide file tree
Showing 3 changed files with 110 additions and 50 deletions.
73 changes: 61 additions & 12 deletions src/app/client/webgpu/core.cljs
View file
Original file line number Diff line number Diff line change
@@ -1,12 +1,17 @@
(ns app.client.webgpu.core
(:require [app.client.webgpu.shader :refer [shader-descriptor add-new-rects-shader-descriptor]]))
(:require
[global-flow :refer [!visible-rects !old-visible-rects]]
[app.client.webgpu.shader :refer [shader-descriptor add-new-rects-shader-descriptor]]
[hyperfiddle.incseq :as i]))



(defn upload-vertices [from data device fformat context config]
(println "upload vertices ::" from "::" config)
(defn upload-vertices [from data device fformat context config ids]
;(println "upload vertices ::" from "::" config "")
(let [varray (js/Float32Array. (clj->js data))
settings-array (js/Float32Array. (clj->js config))
ids-array (js/Float32Array. (clj->js ids))
ids-array-length (.-byteLength ids-array)
settings-array (js/Float32Array. (clj->js config))
num-rectangles (count data)
output-size (* num-rectangles 12)
shader-module (.createShaderModule device add-new-rects-shader-descriptor)
Expand All @@ -30,6 +35,18 @@
:size (.-byteLength settings-array)
:usage (bit-or js/GPUBufferUsage.UNIFORM
js/GPUBufferUsage.COPY_DST)}))
id-buffer (.createBuffer
device
(clj->js {:label "id buffer"
:size ids-array-length
:usage (bit-or js/GPUBufferUsage.STORAGE
js/GPUBufferUsage.COPY_DST)}))
rendered-ids-buffer (.createBuffer
device
(clj->js {:label "rendered ids buffer"
:size ids-array-length
:usage (bit-or js/GPUBufferUsage.STORAGE
js/GPUBufferUsage.COPY_SRC)}))
binding-group-layout (.createBindGroupLayout
device
(clj->js {:label "compute bind group layout"
Expand All @@ -41,8 +58,13 @@
:buffer {:type "storage"}}
{:binding 2
:visibility js/GPUShaderStage.COMPUTE
:buffer {:type "uniform"}}])}))

:buffer {:type "uniform"}}
{:binding 3
:visibility js/GPUShaderStage.COMPUTE
:buffer {:type "read-only-storage"}}
{:binding 4
:visibility js/GPUShaderStage.COMPUTE
:buffer {:type "storage"}}])}))
bind-group (.createBindGroup
device
(clj->js {:layout binding-group-layout
Expand All @@ -51,8 +73,11 @@
{:binding 1
:resource {:buffer output-buffer}}
{:binding 2
:resource {:buffer settings-uniform-buffer}}])}))

:resource {:buffer settings-uniform-buffer}}
{:binding 3
:resource {:buffer id-buffer}}
{:binding 4
:resource {:buffer rendered-ids-buffer}}])}))
pipeline-layout (.createPipelineLayout
device
(clj->js {:label "compute pipeline layout"
Expand All @@ -73,14 +98,14 @@
:visibility js/GPUShaderStage.VERTEX
:buffer {:type "uniform"}}])}))


render-bind-group (.createBindGroup
device
(clj->js {:layout render-binding-group-layout
:entries (clj->js [{:binding 1
:resource {:buffer output-buffer}}
:resource {:buffer output-buffer}}
{:binding 2
:resource {:buffer settings-uniform-buffer}}])}))

render-pipeline-layout (.createPipelineLayout
device
(clj->js {:label "compute pipeline layout"
Expand All @@ -104,6 +129,8 @@
(.then (fn [info] (js/console.log "compute shader info:" info))))
(.writeBuffer (.-queue device) input-buffer 0 varray)
(.writeBuffer (.-queue device) settings-uniform-buffer 0 settings-array)
(.writeBuffer (.-queue device) id-buffer 0 ids-array)


(let [encoder (.createCommandEncoder device)
compute-pass (.beginComputePass encoder)]
Expand All @@ -112,7 +139,29 @@
(.setBindGroup compute-pass 0 bind-group)
(.dispatchWorkgroups compute-pass (/ num-rectangles 64))
(.end compute-pass)
(.submit (.-queue device) [(.finish encoder)]))
;(.submit (.-queue device) [(.finish encoder)])
(let [staging-buffer (.createBuffer
device
(clj->js {:label "staging buffer"
:size ids-array-length
:usage (bit-or js/GPUBufferUsage.MAP_READ
js/GPUBufferUsage.COPY_DST)}))]
(.copyBufferToBuffer encoder rendered-ids-buffer 0 staging-buffer 0 ids-array-length)

(.submit (.-queue device) [(.finish encoder)])

; Read the staging buffer
(-> (.mapAsync staging-buffer js/GPUMapMode.READ)
(.then (fn []
(let [mapped-range (.getMappedRange staging-buffer)
num-rendered (js/Float32Array. mapped-range)
rendered-ids (sort (into-array num-rendered))
new-rects (into-array (filter (complement zero?) rendered-ids))]
(when (= "initial" from)
(reset! !old-visible-rects new-rects))
(reset! !visible-rects new-rects)
(.unmap staging-buffer)))))))


;; Render pipeline
(let [encoder (.createCommandEncoder device)
Expand All @@ -129,4 +178,4 @@
(.setBindGroup render-pass 0 render-bind-group)
(.draw render-pass (* num-rectangles 2))
(.end render-pass)
(.submit (.-queue device) [(.finish encoder)]))))
(.submit (.-queue device) [(.finish encoder)]))))
83 changes: 45 additions & 38 deletions src/app/client/webgpu/shader.cljs
View file
Original file line number Diff line number Diff line change
Expand Up @@ -33,9 +33,11 @@
zoomFactor: f32,
}
// Constants for screen dimensions
@group(0) @binding(0) var<storage, read> rectangles: array<f32>; // Flattened input array
@group(0) @binding(0) var<storage, read> rectangles: array<f32>; // Flattened input array
@group(0) @binding(1) var<storage, read_write> vertices: array<f32>; // Output vertices
@group(0) @binding(2) var<uniform> canvas_settings: CanvasSettings; // Canvas settings uniform\n
@group(0) @binding(2) var<uniform> canvas_settings: CanvasSettings; // Canvas settings uniform\n
@group(0) @binding(3) var<storage, read> id_buffer: array<u32>;
@group(0) @binding(4) var<storage, read_write> rendered_ids: array<atomic<u32>>;
@compute @workgroup_size(64)
Expand All @@ -50,52 +52,57 @@
let width = rectangles[base_index + 3];
// Calculate the four corners of the rectangle in clip space
let left = (x / canvas_settings.width ) * 2 - 1 ;
let right = ((x + width) / canvas_settings.width ) * 2 - 1;
let top = 1 - (y / canvas_settings.height ) * 2 ;
let bottom = 1 - ((y + height) / canvas_settings.height ) * 2 ;
// Create 6 vertices for two triangles (12 float values)
let vertex_index = index * 12; // 6 vertices * 2 components each
// Triangle 1
vertices[vertex_index + 0] = left;
vertices[vertex_index + 1] = top;
vertices[vertex_index + 2] = right;
vertices[vertex_index + 3] = top;
vertices[vertex_index + 4] = left;
vertices[vertex_index + 5] = bottom;
// Triangle 2
vertices[vertex_index + 6] = right;
vertices[vertex_index + 7] = top;
vertices[vertex_index + 8] = right;
vertices[vertex_index + 9] = bottom;
vertices[vertex_index + 10] = left;
vertices[vertex_index + 11] = bottom;
}
let left = ((x / canvas_settings.width ) * 2 - 1) * canvas_settings.zoomFactor + canvas_settings.panX;
let right = (((x + width) / canvas_settings.width ) * 2 - 1) * canvas_settings.zoomFactor + canvas_settings.panX;
let top = (1 - (y / canvas_settings.height ) * 2) * canvas_settings.zoomFactor + canvas_settings.panY ;
let bottom = (1 - ((y + height) / canvas_settings.height ) * 2) * canvas_settings.zoomFactor + canvas_settings.panY ;
if (max(left, right) >= -1.0 && min(left, right) <= 1.0 && max(top, bottom) >= -1.0 && min(top, bottom) <= 1.0) {\n
let rect_id = id_buffer[index];
let pos = atomicAdd(&rendered_ids[0], 1u);
atomicStore(&rendered_ids[pos], rect_id);
// Create 6 vertices for two triangles (12 float values)
let vertex_index = index * 12; // 6 vertices * 2 components each
// Triangle 1
vertices[vertex_index + 0] = left;
vertices[vertex_index + 1] = top;
vertices[vertex_index + 2] = right;
vertices[vertex_index + 3] = top;
vertices[vertex_index + 4] = left;
vertices[vertex_index + 5] = bottom;
// Triangle 2
vertices[vertex_index + 6] = right;
vertices[vertex_index + 7] = top;
vertices[vertex_index + 8] = right;
vertices[vertex_index + 9] = bottom;
vertices[vertex_index + 10] = left;
vertices[vertex_index + 11] = bottom;
}}
struct VertexOutput {
@builtin(position) position: vec4f,
@location(0) fragPos: vec2f
}
@group(0) @binding(1) var<storage, read> vertex_buffer: array<f32>; // Changed to read-only\n
@group(0) @binding(2) var<uniform> can_settings: CanvasSettings; // Canvas settings uniform\\n\n
@group(0) @binding(1) var<storage, read> vertex_buffer: array<f32>; // Changed to read-only
@vertex
fn renderVertices(@builtin(vertex_index) vertexIndex: u32) -> VertexOutput {
let index = vertexIndex * 2;
let px = vertex_buffer[index] * can_settings.zoomFactor + can_settings.panX;
let py = vertex_buffer[index + 1] * can_settings.zoomFactor + can_settings.panY;
// Pass position to the fragment shader
var output: VertexOutput;
output.position = vec4f(px, py, 0.0, 1.0);
output.fragPos = vec2f(px, py);
return output;
let px = vertex_buffer[index] ;
let py = vertex_buffer[index + 1];
var output: VertexOutput;
output.position = vec4f(px, py, 0.0, 1.0);
output.fragPos = vec2f(px, py);
return output;
}
@fragment
Expand Down
4 changes: 4 additions & 0 deletions src/global_flow.cljs
View file
Original file line number Diff line number Diff line change
Expand Up @@ -53,6 +53,9 @@
(def !canvas-x (atom nil))
(defonce !offset (atom nil))
(defonce !zoom-factor (atom nil))
(def !visible-rects (atom nil))
(def !old-visible-rects (atom nil))
(def !global-event (atom nil))


(defn mouse-down?> [node]
Expand All @@ -63,6 +66,7 @@
(m/relieve {})))



(defn await-promise
"Returns a task completing with the result of given promise"
[p]
Expand Down

0 comments on commit 995476b

Please sign in to comment.