Skip to content

Commit 828a4cc

Browse files
author
Mike Solomon
committed
Begins work on x binding (somewhat buggy)
1 parent 4bf38c5 commit 828a4cc

File tree

1 file changed

+80
-8
lines changed

1 file changed

+80
-8
lines changed

sandbox/Sandbox.purs

Lines changed: 80 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -262,6 +262,8 @@ main = do
262262
timeData :: Float32Array <- fromArray $ hackyFloatConv [ 0.0 ]
263263
rotateZData :: Float32Array <- freshIdentityMatrix
264264
rotateZResultData :: Float32Array <- freshIdentityMatrix
265+
rotateXData :: Float32Array <- freshIdentityMatrix
266+
rotateXResultData :: Float32Array <- freshIdentityMatrix
265267
-- 📇 Index Buffer Data
266268
indices :: Uint16Array <- fromArray $ hackyIntConv
267269
[
@@ -364,10 +366,14 @@ main = do
364366
timeBuffer <- liftEffect $ createBufferF timeData
365367
(GPUBufferUsage.storage .|. GPUBufferUsage.copyDst)
366368
scaleBuffer <- liftEffect $ createBufferF scaleData
367-
(GPUBufferUsage.storage .|. GPUBufferUsage.copySrc)
369+
(GPUBufferUsage.storage )
368370
rotateZBuffer <- liftEffect $ createBufferF rotateZData
369-
(GPUBufferUsage.storage .|. GPUBufferUsage.copySrc)
371+
(GPUBufferUsage.storage )
370372
rotateZResultBuffer <- liftEffect $ createBufferF rotateZResultData
373+
(GPUBufferUsage.storage)
374+
rotateXBuffer <- liftEffect $ createBufferF rotateXData
375+
(GPUBufferUsage.storage )
376+
rotateXResultBuffer <- liftEffect $ createBufferF rotateXResultData
371377
(GPUBufferUsage.storage .|. GPUBufferUsage.copySrc)
372378
-- 🖍️ Shaders
373379
let
@@ -448,6 +454,32 @@ fn main(@builtin(global_invocation_id) global_id : vec3<u32>) {
448454
}"""
449455
}
450456
rotateZModule <- liftEffect $ createShaderModule device rotateZDesc
457+
let
458+
rotateXDesc = x
459+
{ code:
460+
"""
461+
@group(0) @binding(0) var<storage, read_write> resultMatrix : mat4x4<f32>;
462+
@group(1) @binding(0) var<storage, read> time : f32;
463+
fn xyt2trig(x: u32, y: u32, time: f32) -> f32 {
464+
const pi = 3.14159;
465+
var o = (x << 1) + y;
466+
return sin(((time / 8.0) * pi) + (f32((o + 1) % 3) * (pi / 2.0)));
467+
}
468+
469+
@compute @workgroup_size(2,2)
470+
fn main(@builtin(global_invocation_id) global_id : vec3<u32>) {
471+
// Guard against out-of-bounds work group sizes
472+
if (global_id.x >= 2u || global_id.y >= 2u) {
473+
return;
474+
}
475+
// 1,1 is cos(a) 0 1
476+
// 1,2 is -sin(a) 1 2
477+
// 2,1 is sin(a) 2 0
478+
// 2,2 is cos(a) 3 1
479+
resultMatrix[global_id.x + 1][global_id.y + 1] = xyt2trig(global_id.x, global_id.y, time);
480+
}"""
481+
}
482+
rotateXModule <- liftEffect $ createShaderModule device rotateXDesc
451483
let
452484
matrixMultiplicationDesc = x
453485
{ code:
@@ -632,6 +664,24 @@ fn main(@location(0) inColor: vec3<f32>) -> @location(0) vec4<f32> {
632664
(x { buffer: rotateZResultBuffer } :: GPUBufferBinding)
633665
]
634666
}
667+
rotateXComputeBindGroup <- liftEffect $ createBindGroup device $ x
668+
{ layout: simpleIOComputeBindGroupLayout
669+
, entries:
670+
[ gpuBindGroupEntry 0
671+
(x { buffer: rotateXBuffer } :: GPUBufferBinding)
672+
]
673+
}
674+
rotateXResultIOComputeBindGroup <- liftEffect $ createBindGroup device $ x
675+
{ layout: matrixMultiplicationComputeBindGroupLayout
676+
, entries:
677+
[ gpuBindGroupEntry 0
678+
(x { buffer: rotateXBuffer } :: GPUBufferBinding)
679+
, gpuBindGroupEntry 1
680+
(x { buffer: rotateZResultBuffer } :: GPUBufferBinding)
681+
, gpuBindGroupEntry 2
682+
(x { buffer: rotateXResultBuffer } :: GPUBufferBinding)
683+
]
684+
}
635685
timeComputeBindGroup <- liftEffect $ createBindGroup device $ x
636686
{ layout: timeBindGroupLayout
637687
, entries:
@@ -670,6 +720,10 @@ fn main(@location(0) inColor: vec3<f32>) -> @location(0) vec4<f32> {
670720
{ "module": rotateZModule
671721
, entryPoint: "main"
672722
}
723+
(rotateXState :: GPUProgrammableStage) = x
724+
{ "module": rotateXModule
725+
, entryPoint: "main"
726+
}
673727
(matrixMultiplicationState :: GPUProgrammableStage) = x
674728
{ "module": matrixMultiplicationModule
675729
, entryPoint: "main"
@@ -725,6 +779,10 @@ fn main(@location(0) inColor: vec3<f32>) -> @location(0) vec4<f32> {
725779
{ layout: ioPlusTComputeLayout
726780
, compute: rotateZState
727781
}
782+
rotateXPipeline <- liftEffect $ createComputePipeline device $ x
783+
{ layout: ioPlusTComputeLayout
784+
, compute: rotateXState
785+
}
728786
matrixMultiplicationPipeline <- liftEffect $ createComputePipeline device $
729787
x
730788
{ layout: matrixMultiplicationLayout
@@ -813,9 +871,27 @@ fn main(@location(0) inColor: vec3<f32>) -> @location(0) vec4<f32> {
813871
matrixMultiplicationPipeline
814872
GPUComputePassEncoder.setBindGroup rotateZMulPassEncoder 0
815873
rotateZResultIOComputeBindGroup
816-
GPUComputePassEncoder.dispatchWorkgroupsXY rotateZMulPassEncoder 2 2
874+
GPUComputePassEncoder.dispatchWorkgroupsXY rotateZMulPassEncoder 4 4
817875
GPUComputePassEncoder.end rotateZMulPassEncoder
818876
---------------------
877+
rotateXPassEncoder <- beginComputePass commandEncoder (x {})
878+
GPUComputePassEncoder.setPipeline rotateXPassEncoder
879+
rotateXPipeline
880+
GPUComputePassEncoder.setBindGroup rotateXPassEncoder 0
881+
rotateXComputeBindGroup
882+
GPUComputePassEncoder.setBindGroup rotateXPassEncoder 1
883+
timeComputeBindGroup
884+
GPUComputePassEncoder.dispatchWorkgroupsXY rotateXPassEncoder 2 2
885+
GPUComputePassEncoder.end rotateXPassEncoder
886+
---------------------
887+
rotateXMulPassEncoder <- beginComputePass commandEncoder (x {})
888+
GPUComputePassEncoder.setPipeline rotateXMulPassEncoder
889+
matrixMultiplicationPipeline
890+
GPUComputePassEncoder.setBindGroup rotateXMulPassEncoder 0
891+
rotateXResultIOComputeBindGroup
892+
GPUComputePassEncoder.dispatchWorkgroupsXY rotateXMulPassEncoder 4 4
893+
GPUComputePassEncoder.end rotateXMulPassEncoder
894+
---------------------
819895
-- xTransTestPassEncoder <- beginComputePass commandEncoder (x {})
820896
-- GPUComputePassEncoder.setPipeline xTransTestPassEncoder
821897
-- xTransTestPipeline
@@ -824,7 +900,7 @@ fn main(@location(0) inColor: vec3<f32>) -> @location(0) vec4<f32> {
824900
-- GPUComputePassEncoder.dispatchWorkgroupsX xTransTestPassEncoder 1
825901
-- GPUComputePassEncoder.end xTransTestPassEncoder
826902
--
827-
copyBufferToBuffer commandEncoder rotateZResultBuffer 0 uniformBuffer 0
903+
copyBufferToBuffer commandEncoder rotateXResultBuffer 0 uniformBuffer 0
828904
(4 * 16)
829905
-- 🖌️ Encode drawing commands
830906
let
@@ -861,10 +937,6 @@ fn main(@location(0) inColor: vec3<f32>) -> @location(0) vec4<f32> {
861937
GPUComputePassEncoder.setBindGroup resetPassEncoder 0
862938
scaleBindGroup
863939
GPUComputePassEncoder.dispatchWorkgroupsXY resetPassEncoder 4 4
864-
-- reset the multiplier
865-
GPUComputePassEncoder.setBindGroup resetPassEncoder 0
866-
rotateZComputeBindGroup
867-
GPUComputePassEncoder.dispatchWorkgroupsXY resetPassEncoder 4 4
868940
GPUComputePassEncoder.end resetPassEncoder
869941
-- 🙌 finish commandEncoder
870942
toSubmit <- finish commandEncoder

0 commit comments

Comments
 (0)