Skip to content

Commit 470bd39

Browse files
author
Mike Solomon
committed
Squashed commit of the following:
commit 52cf537bffc3e2b866c4f7447efc1258bee49986 Author: Mike Solomon <[email protected]> Date: Mon Jan 9 23:00:54 2023 +0200 Fixes shaders commit 5544fcb Author: Mike Solomon <[email protected]> Date: Mon Jan 9 21:59:10 2023 +0200 Uses arrays only in compute shaders
1 parent eb5188f commit 470bd39

File tree

1 file changed

+57
-59
lines changed

1 file changed

+57
-59
lines changed

sandbox/Sandbox.purs

Lines changed: 57 additions & 59 deletions
Original file line numberDiff line numberDiff line change
@@ -21,6 +21,7 @@ import Data.UInt (UInt)
2121
import Effect (Effect)
2222
import Effect.Aff (error, launchAff_, throwError)
2323
import Effect.Class (liftEffect)
24+
2425
import Unsafe.Coerce (unsafeCoerce)
2526
import Web.DOM.Element (setAttribute)
2627
import Web.DOM.NonElementParentNode (getElementById)
@@ -51,6 +52,7 @@ import Web.GPU.GPUFragmentState (GPUFragmentState)
5152
import Web.GPU.GPUFrontFace (cw)
5253
import Web.GPU.GPUIndexFormat (uint16)
5354
import Web.GPU.GPULoadOp as GPULoadOp
55+
5456
import Web.GPU.GPUPrimitiveState (GPUPrimitiveState)
5557
import Web.GPU.GPUPrimitiveTopology (triangleList)
5658
import Web.GPU.GPUProgrammableStage (GPUProgrammableStage)
@@ -288,6 +290,8 @@ main = do
288290
translateZResultData :: Float32Array <- freshIdentityMatrix
289291
perspectiveData :: Float32Array <- getPerspectiveMatrix
290292
perspectiveResultData :: Float32Array <- freshIdentityMatrix
293+
-- msdelta
294+
hackyData :: Float32Array <- freshIdentityMatrix
291295
-- 📇 Index Buffer Data
292296
indices :: Uint16Array <- fromArray $ hackyIntConv
293297
[
@@ -389,45 +393,36 @@ main = do
389393
timeBuffer <- liftEffect $ createBufferF timeData
390394
(GPUBufferUsage.storage .|. GPUBufferUsage.copyDst)
391395
scaleBuffer <- liftEffect $ createBufferF scaleData
392-
(GPUBufferUsage.storage)
396+
(GPUBufferUsage.storage .|. GPUBufferUsage.copySrc)
393397
rotateZBuffer <- liftEffect $ createBufferF rotateZData
394-
(GPUBufferUsage.storage)
398+
(GPUBufferUsage.storage .|. GPUBufferUsage.copySrc)
395399
rotateZResultBuffer <- liftEffect $ createBufferF rotateZResultData
396-
(GPUBufferUsage.storage)
400+
(GPUBufferUsage.storage .|. GPUBufferUsage.copySrc)
397401
rotateXBuffer <- liftEffect $ createBufferF rotateXData
398-
(GPUBufferUsage.storage)
402+
(GPUBufferUsage.storage .|. GPUBufferUsage.copySrc)
399403
rotateXResultBuffer <- liftEffect $ createBufferF rotateXResultData
400-
(GPUBufferUsage.storage)
404+
(GPUBufferUsage.storage .|. GPUBufferUsage.copySrc)
401405
rotateYBuffer <- liftEffect $ createBufferF rotateYData
402-
(GPUBufferUsage.storage)
406+
(GPUBufferUsage.storage .|. GPUBufferUsage.copySrc)
403407
rotateYResultBuffer <- liftEffect $ createBufferF rotateYResultData
404-
(GPUBufferUsage.storage)
408+
(GPUBufferUsage.storage .|. GPUBufferUsage.copySrc)
405409
translateZBuffer <- liftEffect $ createBufferF translateZData
406-
(GPUBufferUsage.storage)
410+
(GPUBufferUsage.storage .|. GPUBufferUsage.copySrc)
407411
translateZResultBuffer <- liftEffect $ createBufferF translateZResultData
408-
(GPUBufferUsage.storage)
412+
(GPUBufferUsage.storage .|. GPUBufferUsage.copySrc)
409413
perspectiveBuffer <- liftEffect $ createBufferF perspectiveData
410-
(GPUBufferUsage.storage)
414+
(GPUBufferUsage.storage .|. GPUBufferUsage.copySrc)
411415
perspectiveResultBuffer <- liftEffect $ createBufferF perspectiveResultData
412416
(GPUBufferUsage.storage .|. GPUBufferUsage.copySrc)
417+
-- msdelta
418+
hackyBuffer <- liftEffect $ createBufferF hackyData
419+
(GPUBufferUsage.copyDst .|. GPUBufferUsage.mapRead)
413420
-- 🖍️ Shaders
414-
let
415-
resetDesc = x
416-
{ code:
417-
"""
418-
@group(0) @binding(0) var<storage, read_write> resultMatrix : mat4x4<f32>;
419-
420-
@compute @workgroup_size(4,4)
421-
fn main(@builtin(global_invocation_id) global_id : vec3<u32>) {
422-
resultMatrix[global_id.x][global_id.y] = select(0.0, 1.0, global_id.x == global_id.y);
423-
}"""
424-
}
425-
resetModule <- liftEffect $ createShaderModule device resetDesc
426421
let
427422
initialScaleDesc = x
428423
{ code:
429424
"""
430-
@group(0) @binding(0) var<storage, read_write> resultMatrix : mat4x4<f32>;
425+
@group(0) @binding(0) var<storage, read_write> resultMatrix : array<f32>;
431426
432427
@compute @workgroup_size(4)
433428
fn main(@builtin(global_invocation_id) global_id : vec3<u32>) {
@@ -437,7 +432,7 @@ fn main(@builtin(global_invocation_id) global_id : vec3<u32>) {
437432
return;
438433
}
439434
440-
resultMatrix[global_id.x][global_id.x] = resultMatrix[global_id.x][global_id.x] * 0.25;
435+
resultMatrix[global_id.x*4 + global_id.x] = 0.25;
441436
}"""
442437
}
443438
initialScaleModule <- liftEffect $ createShaderModule device
@@ -446,75 +441,82 @@ fn main(@builtin(global_invocation_id) global_id : vec3<u32>) {
446441
rotateZDesc = x
447442
{ code:
448443
"""
449-
@group(0) @binding(0) var<storage, read_write> resultMatrix : mat4x4<f32>;
444+
@group(0) @binding(0) var<storage, read_write> resultMatrix : array<f32>;
450445
@group(1) @binding(0) var<storage, read> time : f32;
451446
fn xyt2trig(x: u32, y: u32, time: f32) -> f32 {
452447
const pi = 3.14159;
453448
var o = (x << 1) + y;
454449
return sin((time * pi) + (f32(2 - ((o + 1) % 3)) * (pi / 2.0)));
455450
}
456451
457-
@compute @workgroup_size(2,2)
452+
@compute @workgroup_size(4)
458453
fn main(@builtin(global_invocation_id) global_id : vec3<u32>) {
459-
resultMatrix[global_id.x][global_id.y] = xyt2trig(global_id.x, global_id.y, time);
454+
var ixx = global_id.x / 2;
455+
var ixy = global_id.x % 2;
456+
resultMatrix[ixx*4 + ixy] = xyt2trig(ixx, ixy, time);
460457
}"""
461458
}
462459
rotateZModule <- liftEffect $ createShaderModule device rotateZDesc
463460
let
464461
rotateYDesc = x
465462
{ code:
466463
"""
467-
@group(0) @binding(0) var<storage, read_write> resultMatrix : mat4x4<f32>;
464+
@group(0) @binding(0) var<storage, read_write> resultMatrix : array<f32>;
468465
@group(1) @binding(0) var<storage, read> time : f32;
469466
fn xyt2trig(x: u32, y: u32, time: f32) -> f32 {
470467
const pi = 3.14159;
471468
var o = (x << 1) + y;
472469
return sin((time * pi) + (f32(2 - ((o + 1) % 3)) * (pi / 2.0)));
473470
}
474471
475-
@compute @workgroup_size(2,2)
472+
@compute @workgroup_size(4)
476473
fn main(@builtin(global_invocation_id) global_id : vec3<u32>) {
477-
resultMatrix[global_id.x * 2][global_id.y * 2] = xyt2trig(global_id.x, global_id.y, time);
474+
var ixx = global_id.x / 2;
475+
var ixy = global_id.x % 2;
476+
477+
resultMatrix[(ixx*8) + (ixy*2)] = xyt2trig(ixx, ixy, time);
478478
}"""
479479
}
480480
rotateYModule <- liftEffect $ createShaderModule device rotateYDesc
481481
let
482482
rotateXDesc = x
483483
{ code:
484484
"""
485-
@group(0) @binding(0) var<storage, read_write> resultMatrix : mat4x4<f32>;
485+
@group(0) @binding(0) var<storage, read_write> resultMatrix : array<f32>;
486486
@group(1) @binding(0) var<storage, read> time : f32;
487487
fn xyt2trig(x: u32, y: u32, time: f32) -> f32 {
488-
const pi = 3.14159;
488+
const pi = 3.14159;
489489
var o = (x << 1) + y;
490490
return sin((time * pi) + (f32((o + 1) % 3) * (pi / 2.0)));
491491
}
492492
493-
@compute @workgroup_size(2,2)
493+
@compute @workgroup_size(4)
494494
fn main(@builtin(global_invocation_id) global_id : vec3<u32>) {
495495
// Guard against out-of-bounds work group sizes
496-
if (global_id.x >= 2u || global_id.y >= 2u) {
497-
return;
498-
}
499-
resultMatrix[global_id.x + 1][global_id.y + 1] = xyt2trig(global_id.x, global_id.y, time);
496+
var ixx = global_id.x / 2;
497+
var ixy = global_id.x % 2;
498+
499+
resultMatrix[1 + ((ixx + 1)*4) + ixy] = xyt2trig(ixx, ixy, time);
500500
}"""
501501
}
502502
rotateXModule <- liftEffect $ createShaderModule device rotateXDesc
503503
let
504504
matrixMultiplicationDesc = x
505505
{ code:
506506
"""
507-
@group(0) @binding(0) var<storage, read> matrixL : mat4x4<f32>;
508-
@group(0) @binding(1) var<storage, read> matrixR : mat4x4<f32>;
509-
@group(0) @binding(2) var<storage, read_write> resultMatrix : mat4x4<f32>;
510-
@compute @workgroup_size(4,4)
507+
@group(0) @binding(0) var<storage, read> matrixL : array<f32>;
508+
@group(0) @binding(1) var<storage, read> matrixR : array<f32>;
509+
@group(0) @binding(2) var<storage, read_write> resultMatrix : array<f32>;
510+
@compute @workgroup_size(16)
511511
fn main(@builtin(global_invocation_id) global_id : vec3<u32>) {
512512
var result = 0.0;
513+
var ixx = global_id.x / 4;
514+
var ixy = global_id.x % 4;
513515
for (var i = 0u; i < 4u; i = i + 1u) {
514-
result = result + (matrixL[i][global_id.y] * matrixR[global_id.x][i]);
516+
result = result + (matrixL[(i * 4 )+ ixy] * matrixR[(ixx * 4) + i]);
515517
}
516518
517-
resultMatrix[global_id.x][global_id.y] = result;
519+
resultMatrix[ixx*4 + ixy] = result;
518520
}"""
519521
}
520522
matrixMultiplicationModule <- liftEffect $ createShaderModule device
@@ -762,10 +764,6 @@ fn main(@location(0) inColor: vec3<f32>) -> @location(0) vec4<f32> {
762764
}
763765
-- 🎭 Shader Stages
764766
let
765-
(resetState :: GPUProgrammableStage) = x
766-
{ "module": resetModule
767-
, entryPoint: "main"
768-
}
769767
(initialScaleState :: GPUProgrammableStage) = x
770768
{ "module": initialScaleModule
771769
, entryPoint: "main"
@@ -821,10 +819,6 @@ fn main(@location(0) inColor: vec3<f32>) -> @location(0) vec4<f32> {
821819
, primitive
822820
, depthStencil
823821
}
824-
resetPipeline <- liftEffect $ createComputePipeline device $ x
825-
{ layout: simpleIOComputeLayout
826-
, compute: resetState
827-
}
828822
initialScalePipeline <- liftEffect $ createComputePipeline device $ x
829823
{ layout: simpleIOComputeLayout
830824
, compute: initialScaleState
@@ -1016,17 +1010,20 @@ fn main(@location(0) inColor: vec3<f32>) -> @location(0) vec4<f32> {
10161010
setBindGroup passEncoder 0 uniformBindGroup
10171011
drawIndexedWithInstanceCount passEncoder 36 1
10181012
end passEncoder
1019-
-- 💻 Encode compute commands for resetting buffer
1020-
resetPassEncoder <- beginComputePass commandEncoder (x {})
1021-
GPUComputePassEncoder.setPipeline resetPassEncoder resetPipeline
1022-
-- reset the main storage
1023-
GPUComputePassEncoder.setBindGroup resetPassEncoder 0
1024-
scaleBindGroup
1025-
GPUComputePassEncoder.dispatchWorkgroups resetPassEncoder 1
1026-
GPUComputePassEncoder.end resetPassEncoder
1013+
--------
1014+
copyBufferToBuffer commandEncoder rotateXBuffer 0
1015+
hackyBuffer
1016+
0
1017+
(4 * 16)
10271018
-- 🙌 finish commandEncoder
10281019
toSubmit <- finish commandEncoder
10291020
submit queue [ toSubmit ]
1021+
-- launchAff_ do
1022+
-- toAffE $ convertPromise <$> mapAsync hackyBuffer GPUMapMode.read
1023+
-- liftEffect do
1024+
-- mr <- getMappedRange hackyBuffer
1025+
-- arr <- (whole mr :: Effect Float32Array) >>= toArray
1026+
-- logShow arr
10301027
let
10311028
render = unit # fix \f _ -> do
10321029
-- ⏭ Acquire next image from context
@@ -1037,6 +1034,7 @@ fn main(@location(0) inColor: vec3<f32>) -> @location(0) vec4<f32> {
10371034
encodeCommands colorTextureView
10381035

10391036
-- ➿ Refresh canvas
1037+
-- msdelta
10401038
window >>= void <<< requestAnimationFrame (f unit)
10411039

10421040
liftEffect render

0 commit comments

Comments
 (0)