A Metal compute shader is a kernel that runs one thread per output element straight into a texture, skipping the vertex and raster pipeline entirely. In MSL, the Metal Shading Language, that means a kernel void function writing to a texture2d<float, access::write>, one thread per pixel. For procedural, audio-reactive visuals this is the fastest possible shape: there is no geometry to transform and no rasterizer to feed. RenderWave’s entire shader catalog is built this way, and this post is the long version of why.
If you only remember one thing: on Apple Silicon, a fullscreen generative visual is a compute dispatch, and everything interesting is in how you size it, feed it, and scale it. The rest of this post covers all three.
What is a compute shader, exactly?
Graphics programmers usually think in fragments: a vertex shader positions triangles, a fragment shader colors the pixels they cover. A compute kernel drops that model. You get a grid of threads, each with a position, and each thread does whatever work you assign it. Apple documents the two dispatch entry points on MTLComputeCommandEncoder: dispatchThreadgroups(_:threadsPerThreadgroup:) for grids aligned to threadgroup boundaries, and dispatchThreads(_:threadsPerThreadgroup:) for arbitrarily sized grids, which uses nonuniform threadgroups (supported since macOS 10.13) to handle edge threads without manual rounding.
A minimal MSL kernel that writes a fullscreen texture looks like this:
kernel void my_visual(
texture2d<float, access::write> output [[texture(0)]],
constant Uniforms& u [[buffer(0)]],
uint2 gid [[thread_position_in_grid]])
{
if (gid.x >= output.get_width() || gid.y >= output.get_height()) { return; }
float3 color = do_the_math(gid, u);
output.write(float4(color, 1.0), gid);
}One thread per pixel, no draw call plumbing, and the texture lands ready for the next stage of the pipeline. That is the whole mental model.
Is MSL just C++?
Close. MSL is a C++17-based language. The current Metal Shading Language Specification (version 4.1, dated June 2026) states that in Metal 4 and later the language is a C++17-based specification. In practice that means real headers, real templates, and real shared libraries between shaders.
This matters more than it sounds. RenderWave ships 77 Metal compute kernels in its catalog, and nearly all of them compile against one shared C++17 header library that provides the heavy math: signed distance fields, curl noise, ACES and club-style tone maps, and an HDR output sanitizer. None of that would survive copy-pasted into every shader file. If you come from GLSL, where shared code means string concatenation tricks, this is the single biggest quality-of-life upgrade MSL gives you.
One caveat worth knowing: per Apple’s spec, only Apple silicon supports the new language features in standard 3.2 and above, and Metal 4 hardware requirements start at M1. RenderWave targets Metal 3-era APIs on macOS 15 and later, which covers every Apple Silicon Mac while keeping the language surface stable.
How do you size the dispatch?
The wrong way is hardcoding a threadgroup size. The right way is querying the pipeline at runtime. Apple’s own sizing guide gives the recipe:
let w = pipeline.threadExecutionWidth
let h = pipeline.maxTotalThreadsPerThreadgroup / w
let threadsPerThreadgroup = MTLSize(width: w, height: h, depth: 1)
dispatchThreads(gridSize, threadsPerThreadgroup: threadsPerThreadgroup)Apple notes that maxTotalThreadsPerThreadgroup may differ between pipeline objects and is fixed once a pipeline is created, so you query per pipeline, not per device. RenderWave applies this exact pattern at every compute dispatch site in its pipeline: deck mixing, layer compositing, transitions, post effects, output mapping, and the Syphon handoff. One recipe, consistently, everywhere.
Does the audio ever reach the GPU?
In RenderWave, no, and that surprises people. The engine runs a 40-band mel-scale FFT on the CPU using Accelerate’s vDSP, groups the bands into bass, mid, mid-high, and high, and derives per-band transient hits, presence metrics, and onset strength from spectral flux. Then it compresses everything the shader can know into a single uniform block: time, render size, the shader’s parameter values (up to 26 per shader, already modulated by the audio bands), tempo fields, and a jitter offset. The kernel is a pure function of that block.
The design has three payoffs:
- No per-frame audio buffers crossing the PCIe-adjacent boundary. One small buffer write per frame, already bound.
- Determinism. Same uniforms, same frame, every time. That makes presets reproducible and debugging sane.
- The audio-routing complexity lives in Swift, where it belongs, instead of being re-invented per shader.
The counterintuitive conclusion: “audio-reactive shader” is a misnomer at the kernel level. The shader reacts to parameters. The engine decides what the audio did to those parameters.
What about MetalFX?
MetalFX is Apple’s upscaling framework: it renders at a lower internal resolution and upscales to the output resolution in less time than rendering directly would take. There are two paths. MTLFXSpatialScaler needs only the color texture. MTLFXTemporalScaler needs color, depth, and motion vectors, and it is where things get interesting for non-game content.
Games have G-buffers with motion data from the camera and geometry. A VJ engine has none. So RenderWave writes its own motion estimation: a brute-force per-pixel block-matching kernel that compares the current frame against the previous frame and writes motion vectors, plus a second kernel that fills the uniform motion and depth textures TemporalScaler expects. With those two feeds in place, temporal upscaling works on pure procedural content. The engine uses spatial scaling by default and gates temporal scaling to Apple9-family GPUs (M3 and later), with a user-facing toggle.
That is the quiet enabler behind 8K output from an internal render at 50 to 85 percent resolution, depending on the quality preset.
Hand-written kernels or Metal Performance Shaders?
Both. Metal Performance Shaders is Apple’s collection of compute and graphics kernels tuned per GPU family. RenderWave hand-writes every visual kernel (that is the product), but where a stock kernel wins, it uses the stock kernel: Gaussian blur and denoise come from MPS. The rule of thumb: write MSL when the math is the visual, use MPS when the math is a known primitive.
Compute vs fragment shader for this kind of work
If your content is procedural (math producing pixels, no geometry, no materials), compute is the simpler and often faster path: no raster pipeline, no fullscreen-quad boilerplate, and direct control over thread scheduling. If your content is textured meshes with lighting and depth testing, the fragment pipeline earns its keep. RenderWave’s catalog is mostly the former, which is why it is compute-first, but the 8 ISF-format 3D mesh shaders in the library use the traditional vertex and fragment path, because geometry needs it.
FAQ
Does MSL run on Intel Macs?
Older MSL versions do, on Intel Macs with Metal support. New language features from standard 3.2 onward are Apple-silicon-only per the spec, and Metal 4 requires M1 or later. RenderWave itself requires macOS 15 or later on Apple Silicon, so this is academic for RenderWave users but real for cross-platform engine authors.
What is the difference between dispatchThreads and dispatchThreadgroups?
dispatchThreadgroups expects a grid already rounded up to whole threadgroups, so you handle edge threads with bounds checks. dispatchThreads takes the true grid size and Metal synthesizes partial threadgroups at the edges (nonuniform threadgroups, macOS 10.13 and later). Modern code prefers dispatchThreads, ideally still with a bounds check as cheap insurance.
How much threadgroup memory do Apple GPUs have?
You query it at runtime with MTLDevice.maxThreadgroupMemoryLength, which returns the per-kernel maximum in bytes. The per-family numeric tables live in Apple’s Metal feature set documentation. Note that plenty of visual kernels, including all of RenderWave’s catalog, use zero threadgroup memory: a thread-per-pixel kernel with uniform inputs needs none.
Is Metal 4 required for any of this?
No. Compute dispatch, nonuniform threadgroups, argument buffers, and MetalFX all work on Metal 3-era APIs. Metal 4’s headline features target Apple’s newest stack. Staying on the Metal 3 surface maximizes the range of Macs a tool can support, which for a shipping product beats chasing the newest API suffix.
Can a compute pipeline really do an entire VJ frame?
Yes. In RenderWave the whole composite stage is compute: the A/B deck crossfader kernel, layer compositing, transitions, the post-processing effect rack, and the final output mapping are all kernels writing textures, with the last stop being either the presentation path or a Syphon server. Syphon shares the finished texture with MadMapper, OBS, or Resolume with no re-render.
Where do I learn MSL?
Start with Apple’s Metal documentation and the MSL Specification PDF. The calculating threadgroup and grid sizes article is the one every kernel author reads twice.
By Wesley Walz, builder of RenderWave. The engine described here is the one that ships: 77 compute kernels, one shared C++17 header library, and a 4-band FFT that never once touches the GPU. If you want to see what that architecture does on stage, RenderWave runs on any Apple Silicon Mac and has a free 720p tier.