b1tg and GitHub
e25bf77ce9
kimi delta attention ( #17281 )
...
* kimi delta attention
* config
* test
2026-07-30 08:23:56 -07:00
wozeparrot and GitHub
b5a2a5666a
gptoss moe routing ( #17284 )
2026-07-30 07:49:14 -07:00
nimlgen and GitHub
3f6f0a19ab
hcq2: less graph_rewrites ( #17300 )
2026-07-30 16:44:45 +03:00
qazal and GitHub
7b7b9c1ae0
spec tests for movement ops before custom_kernel ( #17299 )
...
* spec
* fixing
* xfail
* remove
* update tests
* add back
* change to max_kernels
* split and cleanup
2026-07-30 19:20:24 +09:00
qazal and GitHub
aba5ba4b40
update rangeify comment for mop after AFTER ( #17298 )
2026-07-30 17:09:32 +09:00
chenyu and GitHub
d52ef3077c
fix do_devectorize dtype ( #17295 )
...
each one derives its dtype from src, not shared
2026-07-29 23:04:25 -04:00
sirhcm and GitHub
bfc9fc6e0a
nicer TinyJit decorator ( #17293 )
2026-07-29 21:48:00 -04:00
sirhcm and GitHub
060f447db6
qcom: match cl for SP_CS_INSTR_SIZE ( #17289 )
2026-07-29 18:47:02 -04:00
chenyu and GitHub
fd912b348c
generic c0*x<c1 [pr] ( #17291 )
2026-07-29 18:45:23 -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
chenyu and GitHub
027907a544
fix c0*x<c1 symbolic [pr] ( #17290 )
...
use `//` and not math.ceil
2026-07-29 18:32:31 -04: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
sirhcm and GitHub
aab51fb7b6
64-bit UOp.variable support, try 2 ( #17256 )
2026-07-29 14:39:36 -04:00
chenyu and GitHub
dd86a30798
test case for weakfloat cast to weakint INDEX ( #17287 )
2026-07-29 14:15:52 -04:00
George Hotz and GitHub
52c9e5a99e
rename LOOP -> WEAK and STRONGLOOP -> LOOP ( #17283 )
2026-07-29 10:38:36 -07:00
chenyu and GitHub
6c2b9fac08
weak const cleanups [PR] ( #17282 )
2026-07-29 13:14:20 -04: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
Noah Schiro and GitHub
3df1b07c86
Patch half precision ops for older gpu archs ( #17274 )
2026-07-29 09:42:24 -07:00
chenyu and GitHub
3803f1583b
fix add_raw_barrier [pr] ( #17278 )
...
* fix add_raw_barrier [pr]
a kernel might need multiple barriers
* smaller?
2026-07-29 09:30:00 -07:00
qazal and GitHub
6ea7d366fa
test permuted input in custom Ops.PROGRAM test ( #17279 )
...
* test permuted input in custom Ops.PROGRAM test
* dtype
2026-07-29 18:29:05 +09:00
George Hotz and GitHub
527e57300c
fix smu reset for kernel >= 7 ( #17277 )
...
* fix smu reset for kernel >7
* match kernel
2026-07-28 23:11:21 -07:00
qazal and GitHub
e684fcc689
pm_reduce_collapse fix for re enabling stack for cat of same shape (codex) ( #17251 )
...
* change
* stack cat
* remove nop
* change
2026-07-29 14:47:28 +09: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
chenyu and GitHub
2f8f2d2d37
remove no-op explicit dtype= or cast [PR] ( #17276 )
2026-07-29 00:47:35 -04:00
George Hotz and GitHub
dd16d5aead
apply shrink bugfix for 3.11 ( #17271 )
2026-07-28 19:03:06 -07:00
chenyu and GitHub
291ee435bb
qwen3.6 for 27b and 35b-a3b ( #17268 )
2026-07-28 21:11:49 -04:00
George Hotz and GitHub
a17387d6a0
add UNSHARD to spec ( #17269 )
2026-07-28 17:00:21 -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
George Hotz and GitHub
1757067fa6
add device range as src[1] to multi (kimi) ( #17264 )
...
* add device range as src[1] to multi (kimi)
* cleanup
* fix tests
2026-07-28 16:39:34 -07:00
chenyu and GitHub
f11f884f49
update dtypes.md ( #17266 )
2026-07-28 19:30:13 -04:00
chenyu and GitHub
23e9e76e8c
DEFAULT_FLOAT/DEFAULT_INT ContextVar [pr] ( #17265 )
2026-07-28 19:08:18 -04: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
George Hotz and GitHub
0cdddf3495
remove unneeded default args in renderers ( #17262 )
2026-07-28 14:53:04 -07:00
chenyu and GitHub
fde3a8faff
CAPTURE_PROCESS_REPLAY=0 default chaging test [PR] ( #17261 )
...
* CAPTURE_PROCESS_REPLAY=0 default chaging test [PR]
quick patch before making default dtype context vars, which is quite annoying
* few more
2026-07-28 17:32:35 -04:00
George Hotz and GitHub
9ce65b71c0
delete dsp_pm_late ( #17260 )
2026-07-28 14:05:41 -07:00
George Hotz and GitHub
7d48926293
switch _device_num to AxisType.DEVICE range (kimi) ( #17252 )
...
* switch _device_num to AxisType.DEVICE range (kimi)
* improve
* don't exclude device for ranges
* current multi ends the range
2026-07-28 13:52:47 -07:00
drk and GitHub
d1c3ae0f91
preserve typed ranges through reshape ( #17259 )
2026-07-28 16:51:06 -04:00
chenyu and GitHub
749e002286
remove some dtype= when construct UOp [pr] ( #17257 )
...
prep for removing the dtype field
2026-07-28 16:05:35 -04:00
chenyu and GitHub
c9e11544df
delete explicit casts [pr] ( #17255 )
...
handled by broadcast
2026-07-28 13:57:13 -04:00
George Hotz and GitHub
a9ad08064e
make the github actions runners generic for gitea ( #17254 )
2026-07-28 10:48:49 -07:00
chenyu and GitHub
37cf15942b
UOp.const(dtype=None) infers from from_py [pr] ( #17253 )
...
eventually we drop the dtype arg once everything const dtype is inferred
2026-07-28 12:06:55 -04:00
qazal and GitHub
05e4727feb
arange stack regression test ( #17250 )
...
* simple failing test
* commend out stack
* a little smaller
* work
2026-07-28 20:00:22 +09:00
qazal and GitHub
e3b3eea1e2
test for extra copy in allreduce_cast ( #16740 )
...
* test for extra copy in allreduce_cast
* simpler
* cleanup
2026-07-28 19:02:54 +09:00
chenyu and GitHub
1380d6cc5a
pm_long_decomp doesn't depend on operand dtype [pr] ( #17249 )
...
* pm_long_decomp doesn't depend on operand dtype [pr]
the const the rules returned won't have strong dtype
* fix
2026-07-28 01:15:04 -04:00
chenyu and GitHub
4b7022e8f4
Revert "64-bit UOp.variable support ( #17246 )" ( #17248 )
...
This reverts commit ab8fb191b2 .
2026-07-28 00:00:26 -04:00
sirhcm and GitHub
ab8fb191b2
64-bit UOp.variable support ( #17246 )
2026-07-27 23:56:09 -04:00
chenyu and GitHub
f837ca3587
don't match const dtype in UPat [pr] ( #17247 )
2026-07-27 22:50:08 -04:00
chenyu and GitHub
654d475d37
update pm_long_decomp [pr] ( #17245 )
...
* update pm_long_decomp [pr]
instead of reading dtype from UOp (which might be const that's going to be weak), just pass output dtype in tag
* cleanup tag
2026-07-27 22:28:50 -04:00
chenyu and GitHub
550225f603
clean up pm_lower_weak [pr] ( #17243 )
2026-07-27 19:11:39 -04:00