[WIP] compute shader support

[floooh]

Related sokol-shdc PR: floooh/sokol-tools#173

• merge:
  • sokol-gfx compute support sokol-zig#100
  • sokol-gfx compute support sokol-odin#28
  • sokol-gfx compute support sokol-nim#35
  • sokol-gfx compute support sokol-rust#33
  • sokol-gfx compute shader support. sokol-c3#1
  • sokol-gfx compute shader support. kassane/sokol-d#56
  • sokol-gfx compute support colinbellino/sokol-jai#5
  • [WIP] compute shader support sokol-tools#173
  • New samples for sokol-gfx compute support. sokol-samples#162
• bindings:
  • zig
    • sokol-shdc
    • example
  • odin
    • sokol-shdc
    • example
  • c3
    • sokol-shdc
    • example
  • nim
    • sokol-shdc
    • example
  • rust
    • sokol-shdc
    • example
  • d
    • sokol-shdc
    • example
  • jai
    • sokol-shdc
    • example
  • yaml
• general stuff:
  • sokol-shdc documentation
  • sokol-gfx documentation
  • blog post
  • sokok-shdc changelog
  • sokol changelog
  • sample webpage
    • webgl2
    • webgpu
  • GL backend: ifdef-compute-not-available
  • GL backend: use glMemoryBarrier (see: https://arm-software.github.io/opengl-es-sdk-for-android/compute_intro.html)
    • track a 'gpu_dirty' bool in sg_buffer.gl, in sg_apply_bindings issue a memory barrier if any of the storage buffers in the bindings is 'gpu_dirty', and reset the dirty flag, do this across passes and frames
  • Metal backend: remove the managed-buffer-synchronization and use the same granular approach as the memory barrier code in GL => this wouldn't work because the buffer sync needs to happen outside a compute pass
  • change the internal names of the various *_info structs to not be abbreviated so much
  • rename sg_features.storage_buffer to sg_features.compute
  • write a compute demo which samples from regular textures (just to check if texture sampling works in compute passes)
  • port wgpu-compute-boids
  • drive-by: remove return value from _sg_*_apply_bindings(), all implementations always return true anyway => it's actually used in the wgpu backend
  • investigate D3D11 Refcounting Leaks on shutdown (doh, missing d3d11_shutdown()
• update buffer init behaviour:
  • no longer initialize desc.data size from desc.size (and stricter validation for data.size) - code cleanup, not directly related to compute shader
  • GL specific hazard tracking:
    • track a per-buffer flag when a storage buffer was bound for read/write (cleared
      cleared when bound for read-only
    • if a storage buffer is bound in any pass-type and no matter if read/only or read/write and the read/write flag in the buffer was set, issue a glMemoryBarrier()) - don't need other glMemoryBarrier types since storage buffers cannot currently be re-used as vertex- or index-buffer.
    • remove the other unneeded current hazard-tracking code
  • allow creating immutable buffers without content (will be initialized to zero)
    • metal
    • d3d11
    • gl
    • wgpu
• compute shader objects
  • metal
  • d3d11
    • on sg_shader_storage_buffer rename hlsl_register_t_n to hlsl_register_t_or_u_n
  • gl
  • wgpu
• compute pipelines
  • metal
    • investigate: setThreadgroupMemoryLength (this is only need for dynamically sized threadgroup memory, not for statically allocated threadgroup variables (e.g. SPIRVCross turns GLSL shared variables into statically allocated threadgroup variables)
    • semi-related: set bound buffer immutability:
      • in render pipelines: all bound vertex- and storage-buffers are MTLMutabilityImmutable
      • in compute pipelines: read-only storage buffers are MTLMutabilityImmutable
  • d3d11
  • gl
  • wgpu
• compute passes
  • postponed NOTE: concurrent compute pass encoders require bumping the minimal supported macOS base version to 10.14
  • track gpu-written storage-buffers in compute passes
  • metal
    • synchronize written storage buffers (managed buffers only - this blits the gpu-written data back into the cpu-side managed buffer, do this at end of a compute pass)
  • d3d11
  • gl
  • wgpu
• apply_bindings
  • metal
    • figure out if samplers are supposed to work in compute passes, e.g. MTLComputeCommandEncoder allows to bind samplers but how does that work without the partial-derivatives?
    • apply images and samplers in compute passes
  • d3d11
    • ??? - works without d3d11 validation layer warnings, maybe doesn't apply to "raw buffers"? hmm.. "Buffer resources used for output from the compute shader must be created with the D3D11_BIND_RENDER_TARGET flag. Such resources may be read from, however."
    • ??? works without d3d11 validation layer warnings, maybe doesn't apply to "raw buffers"? "Buffer resources created with the D3D11_BIND_SHADER_RESOURCE flag may only be used as inputs to the compute shader."
  • gl
  • wgpu
• dispatch
  • metal
    • assume iPhone8 (e.g. non-uniform thread size feature) (probably don't need this with the below changes)
    • rewrite: take workgroup size from shader reflection
    • rewrite: dispatch takes num_groups, not num_threads
  • d3d11
  • gl
  • wgpu
• validation layer / assert:
  • render/compute shader vs SG_SHADERSTAGE_COMPUTE
  • validate sg_shader_desc.mtl_threads_per_threadgroup (must be > 0, question is just: what's the upper bound?)
  • don't call sg_dispatch() in render passes
  • check that the num_groups_* in the dispatch call doesn't exceed 0xFFFF (see D3D11 functional spec)
  • don't call sg_draw(), sg_apply_viewport, sg_apply_scissor_rect in compute passes
  • storage buffer bindings on vertex- or fragment-stage must be readonly
  • in sg_begin_pass() don't set attachments or swapchain in compute passes
  • in sg_apply_bindings() don't bind vertex/index buffers, textures and samplers (textures and samplers should be fine)
  • in sg_make_shader(): don't mix texture, sampler and writable storage buffer bindings
  • don't use read/write storage resource bindings in render pipeline objects
  • ...?
• trace hooks and stats tracking
  • stats: num_dispatch
  • update sokol_gfx_imgui.h
  • ...?

Questions and open problems:

• [YES] support thread-shared memory? (requires special CPU side support in Metal: https://developer.apple.com/documentation/metal/mtlcomputecommandencoder/setthreadgroupmemorylength(_:index:)?language=objc) => setThreadGroupMemoryLength is only needed when the length isn't known to the shader (e.g. the shared memory isn't statically allocated in the shader itself, we don't need that since SPIRVCross always statically allocates)

• [YES] should we just allow multiple writes to the same resource and insert the necessary barriers in the APIs which require it?

• [DONE] ||in HLSL, GLSL and WGSL the shader defines the workgroup size, but in Metal that's defined on the CPU side in the dispatch call.

  • GLSL: layout(local_size_x = X​, local_size_y = Y​, local_size_z = Z​) in;
  • HLSL: numthreads[x, y, z];
  • WGSL: @workgroup_size(x,y,z)
  • then the dispatch call differs by either providing the number of threads, or the number of workgroup invocation in each dimension (the latter implies that number of threads is an integer multiple of number of workgroups)
  • Solution:
    • follow WebGPU convention
    • need to provide workgroup size as shader reflection information in sg_shader_desc, and this will be passed into the Metal dispatch call... (but needs to be evaluated against the max allowed sizes)
    • need to figure out the max-safe-workgroup-size values (again, what does WebGPU do?)
    • ok, WebGPU does the dispatchWorkgroups thing, e.g. if workgroup_size is 64, then: passEncoder.dispatchWorkgroups(Math.ceil(numParticles / 64));
    • ...still need to figure the max safe values for workgroup size though.

• [LATER] storage textures now or later?

• [LATER] split SG_BUFFERTYPE into a bool-flags struct? (to allow storage+vertex buffers), alternatively: just do vertex pulling from storage buffers...?