It's been nearly half a year since I last wrote publicly about crowley, my experiments in high-performance JSON stream querying. A few quick updates:
- Most of the speculations in May's post came to fruition, including the 'dual-stream' system, which since became the 'pipelined' execution mode
- I found that you can reliably query JSON streams at over 2GB/s on a single thread, and each additional thread you add can buy you around 1GB/s, up to a limit around 5GB/s with three extra threads. This optimizes per-source latency, but is worse for overall system throughput.
- I gave a brief lightning talk on
crowleyat RustConf 2026 in Montreal, and was encouraged to clean it up and submit as a main-track talk for next year's conference in Vancouver. - While it is not yet open source, I published Python bindings on PyPI for people to experiment with.
- Those bindings are very featureful, but messy, and I'm currently in the process of another ground-up rewrite oriented towards optimizing
crowleyfor high-throughput server-side work, rather than as a consumer data processing tool.
As part of shifting crowley's focus from 'install bindings and use it in a notebook/script' to 'deploy it as a long-running service', I've been chasing down two broad ideas.
The first borrows from TigerBeetle's style document, themselves borrowing from NASA's Rule of Ten: put hard limits on everything, and do not use dynamic memory allocation after initialization.
When I first read this principle, I dismissed it as a ridiculously control-freak thing to do in commercial software, outside of embedded systems or literal rockets.
But the more I experimented with it, the more I came to appreciate it (or perhaps I'm learning to accept my inner control freak).
This has spurred an interest in alternate allocators, arenas, and the like, which is well-timed considering that generic trait allocators will reach rustc v1.100 in a week or so, and I expect to present on the subject at RustNYC's UnConf in December.
The second big idea, which we'll spend the rest of this post discussing, is querying JSON on the GPU.
From SIMD parsing to GPU parsing
If you've read my previous post, or are on top of the SIMDJSON literature, it will probably be fairly straightforward to see how this would work. For everyone else, a refresher:
While JSON parsing has historically been a heavily-serial affair, Langdale & Lemire' 2019 Paper 'Parsing Gigabytes of JSON per Second' demonstrated that much of the per-byte work could be eliminated by performing mostly branchless up-front byte classification with SIMD operations, determining which bytes are necessary for parsing and which are superfluous in batches.
Earlier this year with crowley, I adapted the SIMDJSON model to stream parsing and querying with constant memory use, rather than building an in-memory DOM.
The process as I've adapted it works roughly like so:
- Take a 64-byte chunks of JSON, such as:
{ "\\\"Nam[{": [ 116,"\\\\" , 234, "true", false ], "t":"\\\"" }
- Reduce this down to a set of 64-bit bitmasks representing things like quotation marks, backslashes, structural characters ('[', '(', ':' and so on), whitespace, and string interiors:
1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_ : Even bits (for bit-twiddling hacks)
_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1_1 : Odd bits (for bit-twiddling hacks)
__1___1_____1________1____1________1____1___________1_1_1___11__ : Q (naive quote locations)
___111________________1111_______________________________111____ : B (Is this byte a backslash?)
___1__________________1__________________________________1______ : S = B & ~(B << 1) (Is this the start of a chain of backslashes?)
______________________1_________________________________________ : ES = S & E (get all the slash starts that begin on an even offset)
___111____________________1______________________________111____ : EC = B + ES (yields carries on backslash sequences with even starts)
__________________________1_____________________________________ : ECE = EC & ~B (filter out backslashes to get carries only)
________________________________________________________________ : OD1 = ECE & ~E (select only end of sequences ending on an odd offset)
___1_____________________________________________________1______ : OS = S & O (get all the slash starts that begin on an odd offset)
______1_______________1111__________________________________1___ : OC = B + OS (otherwise the same as above)
______1_____________________________________________________1___ : OCE = OC & ~B
______1_____________________________________________________1___ : OD2 = OCE & E
______1_____________________________________________________1___ : OD = OD1 | OD2 (merge to get the ends of all odd-length sequences of backslashes)
__1_________1________1____1________1____1___________1_1_1____1__ : Q &= ~OD (remove escaped quotes from naive quote search)
__1111111111_________11111_________11111____________11__11111___ : R = CLMUL (Q ,~0) (carry-less multiplication by all 1s to find the range of strings, including the starting quote but excluding the ending quote)
- Do some bit-twiddling on these masks to produce a final relevance mask of structural and pseudo-structural characters
__1_________1________1____1________1____1___________1_1_1____1__ : Q (quote locations, no longer naive)
__1111111111_________11111_________11111____________11__11111___ : R (string range `[)`)
1_________11_1_1____1_______1____1_______1_______11____1_______1 : S (structural)
_1____________1_1__________1_1____1_______1_____1__1__________1_ : W (whitespace of any kind)
1 ____________1_1____1_______1____1_______1_______11____1_______1 : S = S & ~R (eliminate quoted regions from structural characters)
1 _1_________11_1____11____1_1____1_1____11_______11_1_111____1_1 : S = S|Q (restore ending quotes to pseudo-structural characters)
111 _________11111___11____1111___111____111_____11111_111____111 : P = S|W (to find pseudo-structural, first find all structural or whitespace characters)
_111_________11111___11____1111___111____111_____11111_111____11 : P = P << 1 (then advance the structural-whitespace mask forward by one)
_____________1_1_1__________1_1__________1_1_____11____1_______1 : P &= ~W & ~R (now eliminate remaining whitespace and quoted characters - voila, we have pseudo-structural characters!)
1 _1_________11_1_1__11____1_1_1__1_1____11_1_____11_1_111____1_1 : S = S|P (now merge pseudo-structural and structural character masks together)
1 _1__________1_1_1__11______1_1__1_1_____1_1_____11_1__11______1 : S & ~(Q & ~R) (eliminate ending quotes from the structural character list)
-
Also construct another couple of masks for opens and closes, which we can use for fast skipping in non-matching regions.
-
In a serial pass, use the results from the previous chunks to adjust the mask in case we ended inside a string, on a backslash run, etc.
-
Use the masks to jump through the original byte chunk, skipping irrelevant bytes and parsing directly on the SPS bytes.
-
Query the rest of the owl
This works really well! Beyond adding the open and close masks for additional downstream information, I wasn't really able to optimize this step much further: all the optimization work on crowley from that point on involved figuring out the best way to split up the work on one or more threads.
But as I optimized the rest of the system further and further, it turned out that the SIMD processing became our bottleneck again.
In particular, processing the 64-byte chunk into the 64-bit bitmask is actually quite expensive!
Porting to AVX-512, which has larger registers and work in larger batches, helped, but not enough to change the calculus.
So the question became: can we do this work not just in wide CPU registers, but on the GPU?
In principle, there's nothing stopping us from doing so. But developing primarily on an M-series Mac, I figured earlier this year that doing this would be a lot of work with Metal GPU internals, and I put it out of my mind.
That was, until the last day of RustConf 2026, when Manuel Drehwald announced at the Scientific Computing Birds of a Feather, as if it were an afterthought, that he had merged std::offload into the nightly compiler the previous evening.
The std::offload module is one of many 'Rust on the GPU' initiatives which were the hot thing of this year's conference, with varying scopes and ambitions. This one intends to let Rust programmers write portable GPU code directly in Rust, which can then compile down to native GPU code with performance characteristics comparable to hand-written PTX. Here's the example in the docs:
#[offload_kernel] fn kernel(x: *mut [f64; 256]) { // SAFETY: // calling our `arch` functions and dereferencing a raw pointer is unsafe unsafe { let n = (*x).len(); let i = (thread_idx_x() + block_idx_x() * block_dim_x()) as usize; if i < n { (*x)[i] = i as f64; } } } let mut x = [0.0f64; 256]; core::offload::offload! { kernel = kernel, workgroup_dim = [256, 1, 1], args = (&mut x as *mut [f64; 256],), }
This remains EXTREMELY experimental. Docs are sparse, and last I checked it only works on x86 Linux systems with NVIDIA chips.
But it does work.
Offloading SIMD Logic
For our initial extraction, we have six masks, which we'll store together:
const STRUCTURAL: u8 = 1 << 0; const WS: u8 = 1 << 1; const QUOTE: u8 = 1 << 2; const BS: u8 = 1 << 3; const OPEN: u8 = 1 << 4; const CLOSE: u8 = 1 << 5; #[repr(C)] struct Masks { s: u64, w: u64, q: u64, bs: u64, opens: u64, closes: u64, }
And we can straightforwardly classify our masks as so: note that some of the structural bytes are also open/close bytes.
#[inline] fn class_of_match(b: u8) -> u8 { match b { b'{' | b'[' => STRUCTURAL | OPEN, b'}' | b']' => STRUCTURAL | CLOSE, b',' | b':' => STRUCTURAL, b' ' | b'\t' | b'\n' | b'\r' => WS, b'"' => QUOTE, b'\\' => BS, _ => 0, } }
But this isn't quite enough: we want to perform this classification branchlessly. While the code above may produce branchless assembly after the optimizing compiler is done with it, I'd rather have certainty.
/// Produces 0xFF when a == b, otherwise 0x00 #[inline] fn eq_mask(a: u8, b: u8) -> u8 { ((a == b) as u8).wrapping_neg() }
Which in aarch compiles to
and w8, w0, #0xff
cmp w8, w1, uxtb
csetm w0, eq ; 0xFF if equal, else 0x00
ret
And the PTX is simply
setp.eq.b16 %p1, %rs1, %rs2;
selp.b32 %r1, -1, 0, %p1;
Which we can put all together as follows:
/// Branchlessly classify one byte #[inline] fn class_of_branchless(b: u8) -> u8 { let open = eq_mask(b, b'{') | eq_mask(b, b'['); let close = eq_mask(b, b'}') | eq_mask(b, b']'); let sep = eq_mask(b, b',') | eq_mask(b, b':'); let ws = eq_mask(b, b' ') | eq_mask(b, b'\t') | eq_mask(b, b'\n') | eq_mask(b, b'\r'); (open & OPEN) | (close & CLOSE) | ((open | close | sep) & STRUCTURAL) | (ws & WS) | (eq_mask(b, b'"') & QUOTE) | (eq_mask(b, b'\\') & BS) }
This does more work overall, but is guaranteed branchless, so we don't have to worry about any mispredictions or blocking vectorization. Since we have several byte classes, which in JSON are going to be heavily interspersed and thus quite difficult to predict, this is very useful.
I also tried to get clever, doing the equality masking with a bit trick
fn eq_mask(a: u8, b: u8) -> u8 { let x = a ^ b; let differs = ((x | x.wrapping_neg()) >> 7); (differs ^ 1).wrapping_neg() }
It turns out that, yes, LLVM is smart and mostly identifies that these two are the same thing and compiles them both down to the same asm/PTX... EXCEPT in the let ws = ... mask, where it turns out the clever-clogs bit-hack messes with the compiler, resulting in a nontrivial divergence in instruction counts:
.visible .func _RNvCsl7DvhbEQTqD_7___rustc17rust_begin_unwind(
.param .b64 _RNvCsl7DvhbEQTqD_7___rustc17rust_begin_unwind_param_0
)
.noreturn
{
$L__BB0_1:
bra.uni $L__BB0_1;
}
; The simple version
.visible .func (.param .b32 func_retval0) eq_simple(
.param .b32 eq_simple_param_0,
.param .b32 eq_simple_param_1
)
{
.reg .pred %p<2>;
.reg .b16 %rs<3>;
.reg .b32 %r<2>;
ld.param.b8 %rs1, [eq_simple_param_0];
ld.param.b8 %rs2, [eq_simple_param_1];
setp.eq.b16 %p1, %rs1, %rs2;
selp.b32 %r1, -1, 0, %p1;
st.param.b32 [func_retval0], %r1;
ret;
}
.visible .func (.param .b32 func_retval0) ws_simple(
.param .b32 ws_simple_param_0
)
{
.reg .pred %p<6>;
.reg .b16 %rs<4>;
.reg .b32 %r<2>;
ld.param.b8 %rs1, [ws_simple_param_0];
setp.eq.b16 %p1, %rs1, 32;
add.s16 %rs2, %rs1, -9;
and.b16 %rs3, %rs2, 255;
setp.lt.u16 %p2, %rs3, 2;
setp.eq.b16 %p3, %rs1, 13;
or.pred %p4, %p1, %p3;
or.pred %p5, %p4, %p2;
selp.b32 %r1, -1, 0, %p5;
st.param.b32 [func_retval0], %r1;
ret;
}
; My clever-clogs version
.visible .func (.param .b32 func_retval0) eq_clever(
.param .b32 eq_clever_param_0,
.param .b32 eq_clever_param_1
)
{
.reg .pred %p<2>;
.reg .b16 %rs<3>;
.reg .b32 %r<2>;
ld.param.b8 %rs1, [eq_clever_param_0];
ld.param.b8 %rs2, [eq_clever_param_1];
setp.eq.b16 %p1, %rs1, %rs2;
selp.b32 %r1, -1, 0, %p1;
st.param.b32 [func_retval0], %r1;
ret;
}
.visible .func (.param .b32 func_retval0) ws_clever(
.param .b32 ws_clever_param_0
)
{
.reg .pred %p<2>;
.reg .b16 %rs<15>;
.reg .b32 %r<2>;
ld.param.b8 %rs1, [ws_clever_param_0];
xor.b16 %rs2, %rs1, 479;
add.s16 %rs3, %rs2, 1;
xor.b16 %rs4, %rs1, 502;
add.s16 %rs5, %rs4, 1;
and.b16 %rs6, %rs3, %rs5;
xor.b16 %rs7, %rs1, 245;
add.s16 %rs8, %rs7, 1;
and.b16 %rs9, %rs6, %rs8;
xor.b16 %rs10, %rs1, 242;
add.s16 %rs11, %rs10, 1;
and.b16 %rs12, %rs9, %rs11;
or.b16 %rs13, %rs12, %rs1;
cvt.s16.s8 %rs14, %rs13;
setp.gt.s16 %p1, %rs14, -1;
selp.b32 %r1, -1, 0, %p1;
st.param.b32 [func_retval0], %r1;
ret;
}
Which is REALLY WEIRD!
Turns out, this is a quirk of how inlining and optimization are ordered.
Because they're tiny functions used all over both the eq_mask implementations are inlined. I've seen LLVM's inlining heuristics glossed as 'yes'.
And because they both are ultimately doing the same operation, the compiler eventually compiles them down to the same operation.
But the optimization process takes multiple passes in LLVM IR, and the inlining process happens partway through, such that the version that gets inlined isn't fully optimized... and so the end results diverge.
In particular, LLVM is able to notice "oh, there's several comparisons here, let's combine them" and checks for \t and \n in the same step, BUT ONLY IF THE COMPARISONS ARE CLEAR, and the bitwise tricks obscure them too much.
Moral of the story: try to outsmart the compiler at your own peril.
Anyway.
Having done the above, executing it with an offload kernel is quite straightforward:
fn classify_branchless(chunk: *const [u8; 64], out: *mut [u8; 64]) { // SAFETY GPU position access: i is always checked against array length // SAFETY raw pointer deref: please don't fuck up unsafe { let n = (*chunk).len(); let i = (thread_idx_x() + block_idx_x() * block_dim_x()) as usize; if i < n { (*out)[i] = class_of_branchless((*chunk)[i]); } } }
And... that's it!
At least, this gets us a correct result. To handle large quantities of data efficiently, we'll have to tweak it, but this is far enough for our current purposes.
But Is It Worth It?
Broadly, profiling the single-threaded implementation found that the parallel classification stage took up ~55% of runtime, the serial reconciliation pass took 5-10%, and the bitwalking+parsing+querying took the remaining 35-40%. This is obviously contingent on the exact dataset, query, and SIMD level, but the pattern broadly holds. In particular, going to AVX-512 speeds up the parallel classification stage, but not enough to make it faster than the serial querying stage.
So, in theory, using a GPU for the classification stage could get us an approximately 2x throughput increase under ideal conditions, so that we could query JSON streams at roughtly 5GB/s with a single thread per source doing the serial stages and the GPU plus a coordinating master thread handling the parallel stage.
In practice... we don't see that.
Not because the GPU is too slow or the model doesn't work, but because the PCIe bridge isn't wide enough.
Testing this on a T4 GPU, the data transfer peaked at 6GB/s in each direction, such that the GPU could only keep up with the throughput demands of 2-3 threads, not nearly enough to make this scheme worthwhile!
Of course, if we were to run this on an integrated memory GPU, we could bypass this issue entirely. I happen to be developing on just such a system!
Unfortunately, the std::offload module remains very experimental. It only works on x86 Linux machines with NVIDIA GPUs, and while a handful of systems like that with integrated GPUs do exist, like the NVIDIA Jetson Orin, they're not readily available on cloud offerings and I'm not spending that much money just for this experiment.
So for the moment, GPU-boosted JSON querying is on the backburner.
In the (hopefully near) future, when I can use std::offload on more diverse systems and/or powerful nodes with integrated memory GPUs can be rented on standard clouds, I expect to come back to this question and potentially get it in production-ready form.