GPU: portable unified memory, launch caching, shared memory, and a Metal benchmark (rebase of #266)#396
Draft
PhilippGrulich wants to merge 9 commits into
Draft
GPU: portable unified memory, launch caching, shared memory, and a Metal benchmark (rebase of #266)#396PhilippGrulich wants to merge 9 commits into
PhilippGrulich wants to merge 9 commits into
Conversation
… memory Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
… size) Replaces the two hardcoded GPU memory models (Metal's fixed NAUTILUS_BUFFER_SIZE copy + CUDA's raw managed-pointer pass-through) with one portable unified model: - gpu::allocUnified/copy/wrap return a gpu::Array<T>, a single-slot newtype over val<T*> that kernels and host entries accept for pointer arguments. - A small arg_factory hook in Engine.hpp lets a parameter type build itself from its value ref (default unchanged for every existing val<T>). - A runtime size table records each allocation's page-rounded length; generated Metal host code wraps the caller's memory with newBufferWithBytesNoCopy (zero copy, exact size) resolving the length via nautilus_gpu_buffer_bytes at dlopen (-undefined dynamic_lookup, deterministic codegen). CUDA passes the managed pointer through unchanged. - Deletes NAUTILUS_BUFFER_SIZE and the gpu.metal.bufferSize option. Array reconstructs its val<T*> from value+state on copy to avoid traceCopy tag collisions that would otherwise alias distinct buffer arguments. Migrates the demo and GPU test kernels to gpu::Array; adds a 4096-element large-batch test (previously truncated at the 1024-float ceiling). Regenerates trace/SSA/IR and Metal/CUDA codegen references (cleaner traces, no fixed size). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The generated host dispatch rebuilt MTLDevice/library/pipeline/queue on every launch — the pipeline-state compile and on-disk metallib reload dominate looped and multi-kernel batch workloads. Hoist them into function-local statics (thread-safe one-time init); only the command buffer, encoder, buffer binding and dispatch run per call. Regenerates Metal host-code references. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Adds block-shared memory: gpu::sharedArray<T,N>() returns a pointer to a statically-sized block-shared array, lowering to threadgroup T[N] (Metal) and __shared__ T[N] (CUDA); combine with syncThreads() for block-cooperative kernels (tiled reduction/scan). CPU fallback uses a thread-local arena. Metal requires per-pointer address spaces, so the device lowering now runs an address-space analysis per function: it seeds from sharedArray results and propagates threadgroup through pointer arithmetic, casts, and block-argument phis to a fixpoint, then retags declarations, phi-copy temporaries, and load/store casts accordingly (replacing the blanket 'device' post-process; output is byte-identical for kernels without shared memory). Adds a block-reduction kernel exercised by the Metal execution test (sum of 1..256 = 32896) and Metal/CUDA codegen + tracing references. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…flow Adds nautilus-gpu-benchmarks (ENABLE_BENCHMARKS): steady-state execution throughput for saxpy/vecadd/block-reduction over 1M/4M/16M-float unified buffers, plus the one-time compile cost. On this M-series machine: ~3 ms dispatch floor scaling to ~7-30 ms at 16M; compilation ~640 ms (external metal/metallib/cc toolchain), which launch caching amortizes. The benchmark surfaced a latent Metal codegen bug: a scalar kernel arg passed as 'constant T&' (const) cannot be reassigned, but when such a value is live inside a conditional it becomes an SSA phi whose trampoline emits 'var = temp;' -> MSL compile error. Existing kernels only used scalar args in branch conditions, never inside a guarded block, so they never hit it. Fix: bind the buffer to a '<var>_arg' parameter and copy it into a mutable local of the SSA name. Regenerates the affected metal device references. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Rewrites the saxpy/vecadd benchmark kernels as grid-stride loops so the same launch function runs both ways from one definition: threads stride in parallel on Metal, and on a non-Metal backend the fallback's gridDim=blockDim=1 makes it a full serial loop over all N on the CPU. Benchmarks each at 1M/4M/16M floats on both mlir (optimized native CPU) and metal (GPU), an apples-to-apples comparison of identical work over the same unified buffers. Finding on this M-series machine: for these bandwidth-bound kernels the CPU and GPU are roughly on par (e.g. saxpy 16M ~8.5 ms CPU vs ~7.9 ms GPU) — unified memory gives both the same bandwidth and the GPU's ~3 ms dispatch floor isn't amortized by so little compute per element. The GPU advantage shows on compute-heavier kernels. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Adds a compute-bound benchmark: per element, iters logistic-map steps (v = 3.9*v*(1-v)) over 1 read + 1 write, so it is arithmetic-bound and the chaotic recurrence cannot be optimized away. Grid-stride, so it compares CPU (serial) vs Metal (parallel) like the other kernels. This is the regime where the GPU wins: on this M-series machine over 256K elements the GPU is ~6.7x faster at 256 iters and ~13x at 1024 iters (the dispatch floor amortizes as arithmetic intensity grows), versus the roughly-on-par result for the bandwidth-bound saxpy/vecadd. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Main's tracer fixes (#390 et al.) changed how the v phi in the blockSum kernel is materialized: the assignment now happens on both branch paths and Block_8 takes the merged value ($26) as its block argument instead of reusing $12. Pure renumbering otherwise; SSA verification passes and both trace contexts agree. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01RA4hMgm5LhH5muupcH5vdZ
Both the CUDA and Metal lowerings of nautilus_gpu_shared_alloc silently emitted a zero-length shared array when the byte-size argument did not lower to a ConstIntOp, miscompiling the kernel. Throw NotImplementedException instead so the misuse surfaces at compile time. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01RA4hMgm5LhH5muupcH5vdZ
![github-actions[bot]](https://avatars.githubusercontent.com/in/15368?s=60&v=4){kind=small}
Contributor
There was a problem hiding this comment.
Tracing Benchmark
Details
| Benchmark suite | Current: 8d40d93 | Previous: 15d8286 | Ratio |
|---|---|---|---|
comp_mlir_add |
5.54811 ms (± 117.451) |
6.78023 ms (± 355.136) |
0.82 |
comp_mlir_ifThenElse |
6.2405 ms (± 345.339) |
7.04872 ms (± 397.01) |
0.89 |
comp_mlir_deeplyNestedIfElse |
5.10939 ms (± 159.806) |
5.70407 ms (± 388.637) |
0.90 |
comp_mlir_loop |
7.11055 ms (± 110.027) |
8.18717 ms (± 377.663) |
0.87 |
comp_mlir_ifInsideLoop |
28.5658 ms (± 216.145) |
31.346 ms (± 327.251) |
0.91 |
comp_mlir_loopDirectCall |
11.76 ms (± 207.277) |
13.6162 ms (± 657.925) |
0.86 |
comp_mlir_pointerLoop |
27.634 ms (± 204.394) |
30.2567 ms (± 709.761) |
0.91 |
comp_mlir_staticLoop |
4.9871 ms (± 139.453) |
5.35532 ms (± 252.284) |
0.93 |
comp_mlir_fibonacci |
10.3992 ms (± 71.0558) |
11.8157 ms (± 437.16) |
0.88 |
comp_mlir_gcd |
9.34438 ms (± 105.741) |
11.1637 ms (± 313.441) |
0.84 |
comp_mlir_nestedIf10 |
10.4485 ms (± 147.371) |
11.9196 ms (± 387.745) |
0.88 |
comp_mlir_nestedIf100 |
24.9132 ms (± 1.00534) |
27.5704 ms (± 570.607) |
0.90 |
comp_mlir_chainedIf10 |
9.51003 ms (± 150.213) |
10.9031 ms (± 495.219) |
0.87 |
comp_mlir_chainedIf100 |
20.1726 ms (± 605.968) |
23.6701 ms (± 404.568) |
0.85 |
comp_cpp_add |
24.5264 ms (± 324.811) |
28.407 ms (± 675.919) |
0.86 |
comp_cpp_ifThenElse |
25.0683 ms (± 367.332) |
28.9562 ms (± 610.909) |
0.87 |
comp_cpp_deeplyNestedIfElse |
25.8114 ms (± 340.522) |
29.9176 ms (± 618.581) |
0.86 |
comp_cpp_loop |
25.4636 ms (± 1.78502) |
28.7345 ms (± 806.526) |
0.89 |
comp_cpp_ifInsideLoop |
25.6126 ms (± 684.455) |
29.1447 ms (± 521.929) |
0.88 |
comp_cpp_loopDirectCall |
25.4671 ms (± 619.823) |
29.1267 ms (± 729.429) |
0.87 |
comp_cpp_pointerLoop |
25.3888 ms (± 449.115) |
29.2797 ms (± 703.427) |
0.87 |
comp_cpp_staticLoop |
25.1411 ms (± 407.993) |
28.9812 ms (± 674.451) |
0.87 |
comp_cpp_fibonacci |
25.0658 ms (± 327.639) |
28.9654 ms (± 799.357) |
0.87 |
comp_cpp_gcd |
25.1992 ms (± 363.664) |
27.9027 ms (± 719.009) |
0.90 |
comp_cpp_nestedIf10 |
27.6881 ms (± 369.19) |
29.6872 ms (± 643.661) |
0.93 |
comp_cpp_nestedIf100 |
53.0604 ms (± 1.17717) |
56.106 ms (± 1.15422) |
0.95 |
comp_cpp_chainedIf10 |
27.8823 ms (± 319.751) |
29.8068 ms (± 1.08393) |
0.94 |
comp_cpp_chainedIf100 |
57.1886 ms (± 565.868) |
60.5496 ms (± 972.437) |
0.94 |
comp_bc_add |
15.1442 us (± 1.778) |
17.3458 us (± 3.68086) |
0.87 |
comp_bc_ifThenElse |
20.2882 us (± 2.78747) |
21.4046 us (± 3.92389) |
0.95 |
comp_bc_deeplyNestedIfElse |
26.3348 us (± 4.30183) |
27.0946 us (± 4.52967) |
0.97 |
comp_bc_loop |
20.6406 us (± 2.92897) |
22.4391 us (± 3.96261) |
0.92 |
comp_bc_ifInsideLoop |
24.6648 us (± 4.81936) |
25.7951 us (± 4.12329) |
0.96 |
comp_bc_loopDirectCall |
20.9481 us (± 2.68099) |
22.9274 us (± 5.15774) |
0.91 |
comp_bc_pointerLoop |
22.8006 us (± 2.94052) |
24.132 us (± 4.82831) |
0.94 |
comp_bc_staticLoop |
18.7075 us (± 3.08169) |
20.0666 us (± 3.83513) |
0.93 |
comp_bc_fibonacci |
21.0426 us (± 2.86602) |
24.1432 us (± 6.31596) |
0.87 |
comp_bc_gcd |
20.0609 us (± 2.89935) |
21.0164 us (± 3.87911) |
0.95 |
comp_bc_nestedIf10 |
48.1208 us (± 6.10137) |
51.0067 us (± 12.2153) |
0.94 |
comp_bc_nestedIf100 |
273.689 us (± 10.276) |
283.828 us (± 15.4088) |
0.96 |
comp_bc_chainedIf10 |
62.9138 us (± 7.79651) |
71.4736 us (± 16.3587) |
0.88 |
comp_bc_chainedIf100 |
447.817 us (± 14.7371) |
468.496 us (± 24.16) |
0.96 |
comp_tbc_add |
1.99401 us (± 149.118) |
2.5099 us (± 555.747) |
0.79 |
comp_tbc_ifThenElse |
4.28865 us (± 401.948) |
4.7664 us (± 921.09) |
0.90 |
comp_tbc_deeplyNestedIfElse |
9.68472 us (± 1.1891) |
10.3845 us (± 2.41359) |
0.93 |
comp_tbc_loop |
4.33685 us (± 363.146) |
5.29407 us (± 1.26893) |
0.82 |
comp_tbc_ifInsideLoop |
7.45953 us (± 716.22) |
8.67493 us (± 2.31165) |
0.86 |
comp_tbc_loopDirectCall |
4.70774 us (± 517.045) |
5.66191 us (± 1.20281) |
0.83 |
comp_tbc_pointerLoop |
5.72482 us (± 622.182) |
5.81812 us (± 830.219) |
0.98 |
comp_tbc_staticLoop |
4.73453 us (± 487.269) |
6.0704 us (± 1.43761) |
0.78 |
comp_tbc_fibonacci |
4.72232 us (± 386.23) |
5.88885 us (± 1.54545) |
0.80 |
comp_tbc_gcd |
3.9765 us (± 375.078) |
4.57403 us (± 1.21873) |
0.87 |
comp_tbc_nestedIf10 |
26.127 us (± 2.86188) |
26.9305 us (± 4.74139) |
0.97 |
comp_tbc_nestedIf100 |
271.147 us (± 17.5587) |
279.783 us (± 15.0568) |
0.97 |
comp_tbc_chainedIf10 |
33.5432 us (± 3.75229) |
37.2041 us (± 8.31885) |
0.90 |
comp_tbc_chainedIf100 |
392.254 us (± 14.334) |
412.06 us (± 23.7445) |
0.95 |
comp_asmjit_add |
24.2067 us (± 3.56452) |
26.3694 us (± 5.86651) |
0.92 |
comp_asmjit_ifThenElse |
32.8267 us (± 3.7333) |
34.8958 us (± 6.56155) |
0.94 |
comp_asmjit_deeplyNestedIfElse |
49.3564 us (± 8.20671) |
55.0873 us (± 10.5758) |
0.90 |
comp_asmjit_loop |
32.7808 us (± 4.15829) |
34.4945 us (± 4.61838) |
0.95 |
comp_asmjit_ifInsideLoop |
48.3433 us (± 9.56705) |
53.3692 us (± 13.6794) |
0.91 |
comp_asmjit_loopDirectCall |
43.879 us (± 8.29265) |
45.5818 us (± 9.40818) |
0.96 |
comp_asmjit_pointerLoop |
45.1909 us (± 6.9588) |
49.6444 us (± 11.9579) |
0.91 |
comp_asmjit_staticLoop |
29.6866 us (± 3.48642) |
32.6602 us (± 6.8584) |
0.91 |
comp_asmjit_fibonacci |
34.6889 us (± 4.02937) |
38.1365 us (± 7.38136) |
0.91 |
comp_asmjit_gcd |
34.2554 us (± 4.06285) |
36.2773 us (± 6.94975) |
0.94 |
comp_asmjit_nestedIf10 |
76.4057 us (± 8.35796) |
82.9894 us (± 14.6) |
0.92 |
comp_asmjit_nestedIf100 |
490.219 us (± 12.1833) |
514.564 us (± 27.0355) |
0.95 |
comp_asmjit_chainedIf10 |
93.1498 us (± 8.07485) |
97.7693 us (± 13.7997) |
0.95 |
comp_asmjit_chainedIf100 |
630.548 us (± 12.7202) |
701.072 us (± 60.0719) |
0.90 |
ssa_add |
183.086 ns (± 15.1116) |
206.466 ns (± 41.2689) |
0.89 |
ssa_ifThenElse |
454.888 ns (± 27.7195) |
665.799 ns (± 125.333) |
0.68 |
ssa_deeplyNestedIfElse |
1.23063 us (± 181.978) |
1.68077 us (± 421.609) |
0.73 |
ssa_loop |
491.002 ns (± 34.8242) |
722.469 ns (± 149.423) |
0.68 |
ssa_ifInsideLoop |
940.156 ns (± 101.831) |
1120.42 ns (± 280259) |
0.84 |
ssa_loopDirectCall |
540.916 ns (± 93.3865) |
682.505 ns (± 250.704) |
0.79 |
ssa_pointerLoop |
634.157 ns (± 82.7281) |
745.36 ns (± 138.973) |
0.85 |
ssa_staticLoop |
432.537 ns (± 55.1174) |
706.926 ns (± 116.865) |
0.61 |
ssa_fibonacci |
503.983 ns (± 29.6389) |
764.676 ns (± 133.441) |
0.66 |
ssa_gcd |
490.256 ns (± 95.8628) |
604.492 ns (± 155.34) |
0.81 |
tiered_twotier_addOne |
3.34257 ms (± 52.2711) |
3.73435 ms (± 252.332) |
0.90 |
tiered_singletier_addOne |
3.24792 ms (± 38.495) |
3.51224 ms (± 247.741) |
0.92 |
single_compile_mlir_addOne |
3.26151 ms (± 72.8515) |
3.47959 ms (± 270.919) |
0.94 |
single_compile_cpp_addOne |
24.6996 ms (± 398.509) |
||
single_compile_bc_addOne |
75.6803 us (± 10.5024) |
||
tiered_twotier_sumLoop |
5.39094 ms (± 64.0782) |
||
tiered_singletier_sumLoop |
5.30062 ms (± 57.4466) |
||
single_compile_mlir_sumLoop |
5.31787 ms (± 69.6818) |
||
single_compile_cpp_sumLoop |
25.2264 ms (± 457.137) |
||
single_compile_bc_sumLoop |
134.155 us (± 12.1365) |
||
exec_bc_addOne |
39.6508 ns (± 5.42634) |
56.1928 ns (± 13.5346) |
0.71 |
exec_mlir_addOne |
280.864 ns (± 5.49431) |
331.5 ns (± 11.6556) |
0.85 |
exec_cpp_addOne |
4.02052 ns (± 0.73054) |
5.43939 ns (± 1.21141) |
0.74 |
exec_interpreted_addOne |
36.6245 ns (± 1.91993) |
51.1662 ns (± 9.63193) |
0.72 |
trace_add |
2.79176 us (± 544.218) |
3.82151 us (± 700.19) |
0.73 |
completing_trace_add |
2.55816 us (± 598.934) |
3.5965 us (± 741.07) |
0.71 |
trace_ifThenElse |
9.99008 us (± 1.5304) |
12.3087 us (± 2.74762) |
0.81 |
completing_trace_ifThenElse |
4.71233 us (± 357.279) |
7.29592 us (± 1.28302) |
0.65 |
trace_deeplyNestedIfElse |
27.2534 us (± 2.27025) |
35.0023 us (± 7.95209) |
0.78 |
completing_trace_deeplyNestedIfElse |
14.0415 us (± 2.09868) |
19.4258 us (± 3.98081) |
0.72 |
trace_loop |
9.15401 us (± 693.742) |
11.2303 us (± 2.30823) |
0.82 |
completing_trace_loop |
5.03428 us (± 586.332) |
6.97141 us (± 1.33993) |
0.72 |
trace_ifInsideLoop |
17.8675 us (± 1.56579) |
23.7106 us (± 5.17687) |
0.75 |
completing_trace_ifInsideLoop |
9.29707 us (± 1.22006) |
12.3956 us (± 2.49794) |
0.75 |
trace_loopDirectCall |
9.34638 us (± 1.02505) |
12.567 us (± 2.74059) |
0.74 |
completing_trace_loopDirectCall |
5.17191 us (± 599.597) |
6.55171 us (± 1.47167) |
0.79 |
trace_pointerLoop |
15.8999 us (± 2.02098) |
19.6787 us (± 6.08244) |
0.81 |
completing_trace_pointerLoop |
10.3409 us (± 869.008) |
14.8336 us (± 3.42047) |
0.70 |
trace_staticLoop |
8.36803 us (± 951.998) |
11.7925 us (± 2.23154) |
0.71 |
completing_trace_staticLoop |
7.90377 us (± 538.492) |
11.0747 us (± 2.07558) |
0.71 |
trace_fibonacci |
11.5602 us (± 1.75882) |
14.1998 us (± 3.47067) |
0.81 |
completing_trace_fibonacci |
6.28076 us (± 884.126) |
9.17402 us (± 1.8285) |
0.68 |
trace_gcd |
8.39536 us (± 769.677) |
11.3968 us (± 2.66044) |
0.74 |
completing_trace_gcd |
4.2329 us (± 359.131) |
6.1137 us (± 1.19334) |
0.69 |
trace_nestedIf10 |
44.2088 us (± 2.98784) |
69.5752 us (± 15.854) |
0.64 |
completing_trace_nestedIf10 |
45.0476 us (± 3.85326) |
62.9243 us (± 12.6442) |
0.72 |
trace_nestedIf100 |
1.96238 ms (± 21.7605) |
2.46732 ms (± 181.167) |
0.80 |
completing_trace_nestedIf100 |
2.01427 ms (± 18.408) |
2.10385 ms (± 206.112) |
0.96 |
trace_chainedIf10 |
106.248 us (± 4.55614) |
118.662 us (± 19.7498) |
0.90 |
completing_trace_chainedIf10 |
49.3447 us (± 4.35545) |
67.7192 us (± 18.1959) |
0.73 |
trace_chainedIf100 |
5.09744 ms (± 36.1779) |
5.63886 ms (± 300.561) |
0.90 |
completing_trace_chainedIf100 |
2.38397 ms (± 48.9621) |
2.50887 ms (± 186.352) |
0.95 |
ir_add |
777.142 ns (± 43.9244) |
953.482 ns (± 172.922) |
0.82 |
ir_ifThenElse |
1.7701 us (± 210.781) |
1.99287 us (± 588.775) |
0.89 |
ir_deeplyNestedIfElse |
3.59727 us (± 453.63) |
3.98684 us (± 1.0069) |
0.90 |
ir_loop |
1.77412 us (± 186.365) |
2.23964 us (± 537.526) |
0.79 |
ir_ifInsideLoop |
2.91357 us (± 208.974) |
3.61333 us (± 963.34) |
0.81 |
ir_loopDirectCall |
1.89677 us (± 136.839) |
2.48056 us (± 625.135) |
0.76 |
ir_pointerLoop |
2.08745 us (± 145.884) |
2.19 us (± 331.903) |
0.95 |
ir_staticLoop |
1.52196 us (± 117.873) |
2.16406 us (± 400.311) |
0.70 |
ir_fibonacci |
1.82825 us (± 185.857) |
2.54834 us (± 508.7) |
0.72 |
ir_gcd |
1.65785 us (± 238.626) |
1.73318 us (± 435.617) |
0.96 |
ir_nestedIf10 |
7.8575 us (± 589.29) |
8.73798 us (± 2.27089) |
0.90 |
ir_nestedIf100 |
88.0977 us (± 4.80222) |
98.7813 us (± 9.72234) |
0.89 |
ir_chainedIf10 |
11.8863 us (± 826.691) |
13.5602 us (± 2.97844) |
0.88 |
ir_chainedIf100 |
173.955 us (± 43.8007) |
167.704 us (± 17.7661) |
1.04 |
exec_mlir_add |
10.3845 ns (± 1.06394) |
12.2741 ns (± 1.86458) |
0.85 |
exec_mlir_fibonacci |
13.1853 us (± 1.33867) |
19.3109 us (± 3.09933) |
0.68 |
exec_mlir_sum |
529.747 us (± 17.99) |
606.883 us (± 24.2837) |
0.87 |
exec_cpp_add |
4.63959 ns (± 0.591795) |
6.50009 ns (± 1.51639) |
0.71 |
exec_cpp_fibonacci |
49.1434 us (± 6.58764) |
51.6654 us (± 7.11851) |
0.95 |
exec_cpp_sum |
11.3684 ms (± 96.3255) |
11.7215 ms (± 238.367) |
0.97 |
exec_bc_add |
48.175 ns (± 7.38137) |
66.0013 ns (± 14.7408) |
0.73 |
exec_bc_fibonacci |
434.433 us (± 15.594) |
445.876 us (± 29.5903) |
0.97 |
exec_bc_sum |
96.6889 ms (± 1.27013) |
96.7427 ms (± 1.04745) |
1.00 |
exec_tbc_add |
26.9866 ns (± 2.91428) |
38.8675 ns (± 8.35499) |
0.69 |
exec_tbc_fibonacci |
142.625 us (± 7.31575) |
145.39 us (± 11.0718) |
0.98 |
exec_tbc_sum |
38.0636 ms (± 1.01597) |
36.7909 ms (± 329.122) |
1.03 |
exec_asmjit_add |
3.53545 ns (± 0.653873) |
4.31765 ns (± 0.952263) |
0.82 |
exec_asmjit_fibonacci |
14.4508 us (± 1.59877) |
15.8853 us (± 2.74874) |
0.91 |
exec_asmjit_sum |
2.8029 ms (± 16.5388) |
3.11154 ms (± 300.665) |
0.90 |
exec_bc_add_passesOff |
47.03 ns (± 6.41301) |
63.4828 ns (± 13.7551) |
0.74 |
exec_bc_add_passesOn |
47.8621 ns (± 10.1905) |
65.6955 ns (± 12.5891) |
0.73 |
exec_bc_fibonacci_passesOff |
433.05 us (± 12.8084) |
444.33 us (± 46.0674) |
0.97 |
exec_bc_fibonacci_passesOn |
390.893 us (± 20.1852) |
405.403 us (± 45.528) |
0.96 |
exec_bc_sum_passesOff |
95.8664 ms (± 1.16616) |
96.3159 ms (± 887.015) |
1.00 |
exec_bc_sum_passesOn |
81.1464 ms (± 1.05197) |
79.0925 ms (± 1.23581) |
1.03 |
exec_tbc_add_passesOff |
28.1147 ns (± 4.59805) |
38.7324 ns (± 8.41232) |
0.73 |
exec_tbc_add_passesOn |
28.1919 ns (± 5.24626) |
40.8541 ns (± 7.73348) |
0.69 |
exec_tbc_fibonacci_passesOff |
142.771 us (± 8.09063) |
144.768 us (± 10.6403) |
0.99 |
exec_tbc_fibonacci_passesOn |
142.99 us (± 7.75493) |
144.729 us (± 11.1878) |
0.99 |
exec_tbc_sum_passesOff |
37.3679 ms (± 1.02939) |
36.7382 ms (± 256.87) |
1.02 |
exec_tbc_sum_passesOn |
34.0335 ms (± 143.406) |
34.1868 ms (± 243.385) |
1.00 |
exec_asmjit_add_passesOff |
3.54176 ns (± 0.517992) |
4.16268 ns (± 0.868263) |
0.85 |
exec_asmjit_add_passesOn |
3.49447 ns (± 0.357778) |
4.26302 ns (± 0.893631) |
0.82 |
exec_asmjit_fibonacci_passesOff |
14.2656 us (± 1.23409) |
16.2878 us (± 3.31987) |
0.88 |
exec_asmjit_fibonacci_passesOn |
14.1673 us (± 1.57416) |
16.1049 us (± 2.72819) |
0.88 |
exec_asmjit_sum_passesOff |
2.8166 ms (± 33.6849) |
3.09123 ms (± 350.247) |
0.91 |
exec_asmjit_sum_passesOn |
2.80871 ms (± 35.0146) |
3.09956 ms (± 374.155) |
0.91 |
exec_bc_add_noRegAlloc |
47.6284 ns (± 5.61401) |
64.9796 ns (± 12.3101) |
0.73 |
exec_bc_add_regAlloc |
46.9443 ns (± 7.92655) |
66.5212 ns (± 14.3839) |
0.71 |
exec_bc_fibonacci_noRegAlloc |
434.288 us (± 17.406) |
444.038 us (± 41.6562) |
0.98 |
exec_bc_fibonacci_regAlloc |
433.151 us (± 16.0153) |
454.985 us (± 38.7697) |
0.95 |
exec_bc_sum_noRegAlloc |
97.9572 ms (± 6.0151) |
96.839 ms (± 1.13443) |
1.01 |
exec_bc_sum_regAlloc |
96.1723 ms (± 1.18395) |
97.0714 ms (± 1.02696) |
0.99 |
exec_bc_add_call |
47.315 ns (± 6.05281) |
62.8661 ns (± 16.3417) |
0.75 |
exec_bc_add_switch |
48.3034 ns (± 10.1205) |
61.9497 ns (± 13.6949) |
0.78 |
exec_bc_add_threaded |
39.5183 ns (± 2.71853) |
57.5087 ns (± 13.8967) |
0.69 |
exec_bc_fibonacci_call |
432.893 us (± 12.1629) |
448.365 us (± 50.8711) |
0.97 |
exec_bc_fibonacci_switch |
395.898 us (± 21.716) |
407.58 us (± 47.6067) |
0.97 |
exec_bc_fibonacci_threaded |
438.215 us (± 13.896) |
455.035 us (± 15.7199) |
0.96 |
exec_bc_sum_call |
97.2283 ms (± 2.38347) |
97.3987 ms (± 9.243) |
1.00 |
exec_bc_sum_switch |
90.8867 ms (± 692.526) |
92.5574 ms (± 11.4172) |
0.98 |
exec_bc_sum_threaded |
87.5987 ms (± 595.354) |
87.4714 ms (± 341.966) |
1.00 |
exec_bc_add_threaded_noReuse |
39.62 ns (± 2.77004) |
58.4138 ns (± 14.353) |
0.68 |
exec_bc_add_threaded_reuse |
33.0916 ns (± 5.16896) |
48.4056 ns (± 9.69617) |
0.68 |
exec_bc_fibonacci_threaded_noReuse |
438.33 us (± 14.0961) |
449.045 us (± 17.5501) |
0.98 |
exec_bc_fibonacci_threaded_reuse |
439.869 us (± 14.0184) |
451.092 us (± 20.342) |
0.98 |
exec_bc_sum_threaded_noReuse |
87.6952 ms (± 969.257) |
87.5593 ms (± 814.441) |
1.00 |
exec_bc_sum_threaded_reuse |
87.3634 ms (± 181.647) |
87.4849 ms (± 275.835) |
1.00 |
exec_bc_add_threaded_noSuperinstr |
40.5836 ns (± 5.56593) |
59.9914 ns (± 11.52) |
0.68 |
exec_bc_add_threaded_superinstr |
39.8457 ns (± 4.68443) |
56.9183 ns (± 10.0123) |
0.70 |
e2e_tiered_bc_to_mlir |
3347.44 us (± 80833.4) |
3647.5 us (± 212251) |
0.92 |
e2e_single_mlir |
5.49586 ms (± 85.6125) |
6.37394 ms (± 299.81) |
0.86 |
This comment was automatically generated by workflow using github-action-benchmark.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Rebase of #266 onto current
main(57 commits ahead of the old merge base), with post-rebase verification and two follow-up commits.What's in here
The original 7 commits from #266, rebased cleanly onto
main:gpu::Array<T>: data-driven unified-memory model (no fixed buffer size), wired into tracing via a newarg_factoryhook inEngine.hppgpu::sharedArray<T, N>: portable static block-shared memory (threadgroupon Metal,__shared__on CUDA) with a threadgroup address-space propagation pass for MetalPlus two new commits from the rebase:
gpuLaunchBlockSumgoldens:main's tracer fixes (Fix tracer dropping Return when a repeated tag assigns a different target #390 et al.) changed how thevphi in the blockSum kernel is materialized — the assignment now happens on both branch paths and the merge block takes the merged value as its block argument. Pure renumbering otherwise; SSA verification passes and both trace contexts agree. These 5 golden files were the only semantic conflict withmain.sharedArraysize: both the CUDA and Metal lowerings ofnautilus_gpu_shared_allocsilently emitted a zero-length shared array when the byte-size argument did not lower to a constant, miscompiling the kernel. They now throwNotImplementedException.Verification (Linux, clang-18, Release,
-DENABLE_GPU_PLUGIN=ON, MLIR backend disabled locally)DemoGPU.cppandGPUMetalBenchmark.cppsyntax-check clean (Metal execution paths need Apple hardware, covered by the macOS CI leg)Note: the Metal
*.host.cppfiles underplugins/gpu/test/data/are byte-exact codegen goldens and are intentionally not clang-formatted (matching how they exist onmain).Supersedes #266.
🤖 Generated with Claude Code
https://claude.ai/code/session_01RA4hMgm5LhH5muupcH5vdZ
Generated by Claude Code