Raph Levien raphlinus

Considerations for subgroups

One feature that is clearly out of scope for WebGPU 1.0 but is desired for the near future is subgroups. It is a way to move data between threads within a workgroup with less overhead and latency than workgroup shared memory, but poses more challenges for portability. While almost all modern GPU hardware supports subgroup operations, the feature poses significant compatibility challenges. In particular, while workgroup size is determined by the programmer within generous ranges (WebGPU requires a minimum maximum of 256), subgroup sizes vary by hardware and also compiler heuristics. Shaders need be written in a way that adapts to a wide range of subgroup sizes, which is quite challenging.

This issue will be written largely from the perspective of accelerating prefix sum operations (an important primitive within Vello), but there are many potential applications. One relatively recent development is cooperative matrix operations, which are supported in most newer GPU hardware a

V: A post-truth language project

This is ordinarily something I'd write on my blog, and I might still, but I'm feeling cautious for reasons that will become clear.

It's often said that we live in a "post-truth" society. This is usually said in the context of politics, fake reviews, etc., and is adjacent to discussions of disinformation, also often connected to the threat of deepfakes and the potential of large language models to produce persuasive text. Even in science there is a replication crisis, and a disturbing amount of straight-up fraud, as documented by Dr. Elisabeth Bik and friends. But it's far more unusual for a programming language to be so characterized. Meet V.

The programming language space is known for passionate discussions which often come down to values, because the space is one of tradeoffs. Want safety (which I believe is related to religious-like concepts of purity and contamination)? Sure, but you have to decide whether you're willing to give up performance

The not-too-clever programmer

This is a translation of grugbrain.dev into clear English. All props to the original author.

Introduction

This is a collection of thoughts on software development, originally written by an pseudonymous author styling themselves the "grug brain developer," but then translated into clear English by Raph Levien.

I am not an extremely smart developer, but I have many years of experience and have learned some things, although still don't know everything.

	-------------------------------------
	Translated Report (Full Report Below)
	-------------------------------------

	Process: Xcode [50519]
	Path: /Applications/Xcode.app/Contents/MacOS/Xcode
	Identifier: com.apple.dt.Xcode
	Version: 14.2 (21534)
	Build Info: IDEFrameworks-21534000000000000~49 (14C18)
	App Item ID: 497799835

	import Foundation
	import Metal

	// This is translated by naga from tile_alloc.wgsl
	// naga --buffer-bounds-check-policy ReadZeroSkipWrite tile_alloc.wgsl tile_alloc.metal
	let metalProgram = """
	// language: metal2.0
	#include <metal_stdlib>
	#include <simd/simd.h>


	@group(0) @binding(0)
	var<storage, read_write> bump: atomic<u32>;

	@group(0) @binding(1)
	var<storage, read_write> paths: array<u32>;

	let WG_SIZE = 256u;

	var<workgroup> sh_tile_offset: u32;

	<html>
	<body>
	<p>WebGPU test:</p>

	</body><div id="out"></div>
	<div>Expected: [1, 1, 1]</div>
	<script>
	function display(text) {
	document.getElementById("out").textContent = text;
	}

	@compute @workgroup_size(256)
	fn main(@builtin(local_invocation_id) local_id: vec3<u32>) {
	if local_id.x < 5u {
	workgroupBarrier();
	}
	}


	struct Alloc {
	offset : u32,
	}

	struct BinInstanceRef {
	offset : u32,
	}

	struct BinInstance {

	// Copyright 2022 Google LLC.
	// SPDX-License-Identifier: Apache-2.0

	// Scrolling test case by Raph Levien

	import SwiftUI
	import CryptoKit

	struct ContentView: View {
	var body: some View {