Snap piet-gpu-hal, add readme.

This snaps to the latest piet-gpu-hal, which fixes some bugs. Also make
the buffer read-write (even though it's only readonly) to work around
limitations in DX12 resource binding. Add a README.

Author: raphlinus

README.md

@@ -18,6 +18,8 @@ The second subdirectory is a simple GUI application that runs a compute shader a
 
 Recommended activity: find an existing shadertoy that implements some interesting algorithm or visual effect, and port it to run in a compute shader. Is there some limitation of the original that could be improved by compute capabilities?
 
+A great place to find shaders to adapt is [The Book of Shaders].
+
 ## A note on the choice of runtime
 
 Your compute shader code cannot run on its own, but rather needs a *runtime* to connect to the GPU, set up resources such as buffers and compiled shader code, and manage the submission of that work to the GPU. There is, as of this writing, no standard runtime for such things, but I hope that will change in time.
@@ -66,3 +68,4 @@ Licensed under either of
 [Cg]: https://www.khronos.org/opengl/wiki/Cg
 [rust-gpu]: https://github.com/EmbarkStudios/rust-gpu
 [IREE]: https://google.github.io/iree/
+[The Book of Shaders]: https://thebookofshaders.com/

piet-compute-toy/Cargo.lock

@@ -9,4 +9,4 @@ edition = "2018"
 
 [dependencies]
 winit = "0.25"
-piet-gpu-hal = { git = "https://github.com/linebender/piet-gpu", rev = "bae185" }
+piet-gpu-hal = { git = "https://github.com/linebender/piet-gpu", rev = "074fafa" }

piet-compute-toy/Cargo.toml

@@ -0,0 +1,15 @@
+# Compute shader toy based on piet-gpu
+
+To run it without changing the shaders, just do `cargo run`.
+
+To change the shaders, you will need to recompile those. You need [ninja] and spirv tools (glslangValidator and spirv-cross) in your path. The easiest way to get those is to install the [Vulkan SDK]. Then, to recompile the shaders, do:
+
+```shell
+(cd shader && ninja) && cargo run
+```
+
+This version is based on piet-gpu-hal, which is the runtime for [piet-gpu]. It's still very new so there are limitations and things that don't work yet. If you run into something, please file an issue!
+
+[Vulkan SDK]: https://www.lunarg.com/vulkan-sdk/
+[ninja]: https://ninja-build.org/
+[piet-gpu]: https://github.com/linebender/piet-gpu

piet-compute-toy/README.md

@@ -9,12 +9,11 @@ rule glsl
   command = $glslang_validator -V -o $out $in
 
 rule hlsl
-  command = $spirv_cross --hlsl $in --output $out
+  command = $spirv_cross --hlsl --shader-model 50 $in --output $out
 
 rule msl
   command = $spirv_cross --msl $in --output $out
 
 build gen/shader.spv: glsl shader.comp
-# TODO: fix "Separate image and samplers not supported in legacy HLSL" error
-#build gen/shader.hlsl: hlsl gen/shader.spv
+build gen/shader.hlsl: hlsl gen/shader.spv
 build gen/shader.msl: msl gen/shader.spv

piet-compute-toy/shader/build.ninja

@@ -0,0 +1,25 @@
+static const uint3 gl_WorkGroupSize = uint3(16u, 16u, 1u);
+
+RWByteAddressBuffer _24 : register(u0);
+RWTexture2D<unorm float4> image : register(u1);
+
+static uint3 gl_GlobalInvocationID;
+struct SPIRV_Cross_Input
+{
+    uint3 gl_GlobalInvocationID : SV_DispatchThreadID;
+};
+
+void comp_main()
+{
+    uint2 xy = gl_GlobalInvocationID.xy;
+    float2 fragCoord = (float2(gl_GlobalInvocationID.xy) / float2(float(_24.Load(0)), float(_24.Load(4)))) - 0.5f.xx;
+    float4 fragColor = float4(fragCoord.x + 0.5f, fragCoord.y + 0.5f, 0.5f + (0.5f * sin(asfloat(_24.Load(8)))), 1.0f);
+    image[int2(xy)] = fragColor;
+}
+
+[numthreads(16, 16, 1)]
+void main(SPIRV_Cross_Input stage_input)
+{
+    gl_GlobalInvocationID = stage_input.gl_GlobalInvocationID;
+    comp_main();
+}

piet-compute-toy/shader/gen/shader.hlsl

@@ -12,7 +12,7 @@ struct Params
 
 constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(16u, 16u, 1u);
 
-kernel void main0(const device Params& _24 [[buffer(0)]], texture2d<float, access::write> image [[texture(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+kernel void main0(device Params& _24 [[buffer(0)]], texture2d<float, access::write> image [[texture(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
     uint2 xy = gl_GlobalInvocationID.xy;
     float2 fragCoord = (float2(gl_GlobalInvocationID.xy) / float2(float(_24.width), float(_24.height))) - float2(0.5);

piet-compute-toy/shader/gen/shader.msl

@@ -20,7 +20,7 @@
 #version 450
 layout(local_size_x = 16, local_size_y = 16) in;
 
-layout(set = 0, binding = 0) readonly restrict buffer Params {
+layout(set = 0, binding = 0) restrict buffer Params {
     uint width;
     uint height;
     float iTime;