chenyu and GitHub
b502fc1367
more const arg -> val ( #17350 )
2026-08-01 02:18:10 -04:00
nimlgen and GitHub
1095bbe409
hcq2: fix ib reuse ( #17330 )
...
* hcq2: initialize IB reuse counters at link
* x
2026-07-31 18:43:23 +03:00
chenyu and GitHub
8e2f175542
const(dtype, b) -> const(b, dtype) [PR] ( #17328 )
...
prep for dtype removal
2026-07-31 09:46:37 -04:00
qazal and GitHub
f7964acb64
llama with MXFP4 ( #17321 )
...
* mxfp4 in llama
* less
* name
2026-07-31 18:25:25 +09:00
qazal and GitHub
0a3325f9c2
add mxfp4 quantize and layout kernels ( #17320 )
2026-07-31 14:50:18 +09:00
qazal and GitHub
a8c1e89500
fp4 asm gemm 6+ pflops ( #17315 )
...
* fp4 gemm
* better kernargs structure
* move to .s files
* work
* work
* p2
* style
* use .py
* move to dsl
* cleanup
* add MFMA_SCALE_X2_ENCODING
* cleanup mfma
* fma docs
* gemm_mxfp4
* more cleanup
* move
* move to cdna_asm_gemm
* change
* rm
* change
* mx
2026-07-31 14:25:23 +09:00
George Hotz and GitHub
d65ea465ed
cleanup gemm fragment + add store unshard ( #17313 )
...
* cleanup gemm fragment + add store unshard
* multi
* fix
2026-07-30 20:53:28 -07:00
wozeparrot and GitHub
b5a2a5666a
gptoss moe routing ( #17284 )
2026-07-30 07:49:14 -07:00
sirhcm and GitHub
060f447db6
qcom: match cl for SP_CS_INSTR_SIZE ( #17289 )
2026-07-29 18:47:02 -04:00
George Hotz and GitHub
138676ab81
improve fragment example + index unshard (kimi) ( #17288 )
...
* fix dtypes in fragment example
* match tilelang
* flip locals
* fix index on unshard
* test fixes
* kimi needs more taste
2026-07-29 15:38:38 -07:00
nimlgen and GitHub
d4ba8b6e0f
hcq2: use stack ( #17286 )
2026-07-29 22:09:34 +03:00
George Hotz and GitHub
b30c7e00d4
support 2d on UNSHARD (kimi) ( #17285 )
...
* support 2d on UNSHARD
* fixes
* Fix test and spec
* single barrier
* 2d sharding works for devices too
* cleanups
* no _rewrap
2026-07-29 12:01:59 -07:00
George Hotz and GitHub
52c9e5a99e
rename LOOP -> WEAK and STRONGLOOP -> LOOP ( #17283 )
2026-07-29 10:38:36 -07:00
George Hotz and GitHub
bd296a7359
enable alloc_fragment support with UNSHARD (kimi) ( #17272 )
...
* enable alloc_fragment support with UNSHARD (kimi)
* cleaner with implicit barrier
* cleanups
* cleaner
* strongloop
* dcount cleanups
2026-07-29 09:46:45 -07:00
George Hotz and GitHub
451120c6e1
make .barrier implicit (kimi) ( #17275 )
...
* make .barrier implicit (kimi)
* simplier
* lil
* remove tinygrad stock barriers
* readable
* lil
2026-07-28 22:34:57 -07:00
George Hotz and GitHub
57ae1bc7a7
rename MULTI to UNSHARD ( #17267 )
...
* rename MULTI to UNSHARD
* comment updates (glm)
* rename method to unshard
2026-07-28 16:51:41 -07:00
755dfb243b
rename CPU_COUNT to NUM_CPU_THREADS with cgroup awareness ( #17263 )
...
Rename CPU_COUNT to NUM_CPU_THREADS so it can be overridden via env var.
Default uses _get_cpu_count() which respects cgroup limits:
- os.process_cpu_count() on Python 3.13+
- /sys/fs/cgroup/cpu.max on cgroup v2
- /sys/fs/cgroup/cpu/cpu.cfs_quota_us on cgroup v1
- os.sched_getaffinity(0) fallback
Use NUM_CPU_THREADS.value in the dataloader instead of cpu_count(),
and update export_model.py and all renderer references.
Co-authored-by: teeny-runner <runner@teeny>
2026-07-28 15:33:57 -07:00
nimlgen and GitHub
97a2265362
hcq2: amd indirect ( #17220 )
...
* ind
* mock
2026-07-26 21:36:01 +03:00
chenyu and GitHub
a60b5f77ac
fix torch backend out= into a view ( #17210 )
2026-07-25 21:04:42 -04:00
nimlgen and GitHub
3946df787d
hcq2, cpu is hcq2-ish ( #17197 )
...
* m
* i
* x
* x
* Df
* x
* x
2026-07-26 02:38:15 +03:00
chenyu and GitHub
076b37e1ae
failing batch norm test ( #17206 )
...
* failing batch norm test
running stats does not schedule in training now since there's no reader
* not that
2026-07-25 18:46:49 -04:00
chenyu and GitHub
f902513355
derive torch backend dispatch from the aten schema ( #17203 )
2026-07-25 14:21:38 -04:00
chenyu and GitHub
983ad3bd95
fix torch backend batchnorm backward ( #17201 )
2026-07-25 13:15:09 -04:00
chenyu and GitHub
9f78504304
checked cast in torch backend unwrap ( #17199 )
2026-07-25 12:34:08 -04:00
chenyu and GitHub
8a10892f5a
fix torch backend as_strided ( #17195 )
...
0 means 0 offset
2026-07-25 02:24:17 -04:00
chenyu and GitHub
4c58b260fb
less wrong calculate_storage_offset ( #17194 )
...
initially for speed, then realized it's just wrong
2026-07-25 01:53:54 -04:00
wozeparrot and GitHub
c9b60caf8c
gptoss: moe gemm kernels ( #17178 )
2026-07-24 07:25:15 -07:00
nimlgen and GitHub
0ecef210bb
hcq2 cleanup 2 ( #17177 )
...
* hcq uops
* x
2026-07-24 13:04:21 +03:00
nimlgen and GitHub
709d7b2343
hcq2 cleanups ( #17162 )
2026-07-24 10:14:09 +03:00
nimlgen and GitHub
62273d50fc
simpler hcq buffer ( #17148 )
...
* simpleer hcq buffer
* x
2026-07-23 13:28:51 +03:00
nimlgen and GitHub
787b2f2db2
hcq2: use ins for hcq ir ( #17120 )
2026-07-21 23:42:34 +03:00
chenyu and GitHub
5244d3cd2a
fix test_u32_to_f16 ( #17119 )
2026-07-21 16:06:05 -04:00
chenyu and GitHub
13ca9bd8a6
remove dtypes.index again ( #17104 )
...
also reverted some dtype change, the split made things needlessly complicated
2026-07-20 20:30:04 -04:00
nimlgen and GitHub
2b96d64496
hcq2: tiny opts and fixes ( #17092 )
2026-07-20 18:46:52 +03:00
qazal and GitHub
abba2aebda
llama: correct fused qkv shape assert ( #17086 )
2026-07-20 15:51:21 +09:00
qazal and GitHub
1cf8f2f68c
llama: inplace amax update ( #17064 )
...
* llama: inplace amax update
* remove amax_out return
* work
* fit
* work
* work
* keep
* diff cleanup
2026-07-20 15:05:41 +09:00
chenyu and GitHub
9970a0aad0
fix Tensor << Tensor for x86 ( #17082 )
...
* fix Tensor << Tensor for x86
* torch
2026-07-19 12:31:05 -04:00
nimlgen and GitHub
232529ce88
hcq2: simpler sync ( #17069 )
...
* x
* y
* n
2026-07-18 16:27:44 +03:00
qazal and GitHub
7dd3422c63
llama: replace two stage amax with atomics ( #17063 )
...
* atomic amax in c kernels
* quantize fp8 UOp kernel
* diff
2026-07-17 19:27:10 +09:00
wozeparrot and GitHub
a836c3822a
gptoss: 3d mx block scale ( #17062 )
2026-07-16 23:30:24 -07:00
George Hotz and GitHub
61e104bdfb
use UOp.wmma everywhere ( #17045 )
...
* use UOp.wmma everywhere
* fix
2026-07-16 10:40:48 -07:00
nimlgen and GitHub
e33e96415f
hcq2: tiny cleanupg ( #17042 )
2026-07-16 16:14:54 +03:00
George Hotz and GitHub
8b0dd870ce
use wmma helper ( #17038 )
2026-07-15 23:25:17 -07:00
nimlgen and GitHub
4257939e50
remove copyin/copyout from Buffer ( #17020 )
...
* remove copyin/copyout from Buffer
* x
* x
* x
* x
2026-07-14 19:47:22 +03:00
qazal and GitHub
939f28d571
fused qkv rope custom kernel ( #17021 )
...
* work
* fused qkv_norm
* work
* speed
* not that yet
* test cleanup
* just clone
* remove .realize()
* cleanup tests
2026-07-15 01:08:42 +09:00
nimlgen and GitHub
e4bdc529c4
hcq2 ci ( #17008 )
...
* hcq2 ci
* x
2026-07-13 19:29:08 +03:00
qazal and GitHub
62ad646d1c
llama: gemm/fa backward speedups (gpt 5.6) ( #17007 )
...
* fp8 atb gemm speedup
* work
* revert
* fa bw faster
2026-07-14 00:27:00 +09:00
nimlgen and GitHub
4d2becddf8
hcq2: spec=2 ( #17006 )
...
* hcq2: spec=2
* hcq: isolate HCQ spec rules
* chq
* move
2026-07-13 18:13:33 +03:00
George Hotz and GitHub
03ecad9486
full removal of dtype.vec ( #16996 )
...
* full removal of dtype.vec
* fix typo
2026-07-12 09:25:50 -07:00
qazal and GitHub
cae6696d75
llama: split current and next amax state ( #16993 )
2026-07-12 18:52:33 +09:00