George Hotz and GitHub
9082ecef5d
use .val to access the value of Ops.CONST ( #17347 )
2026-07-31 19:51:42 -07:00
George Hotz and GitHub
099d69ff7d
ci: split macos unit test into metal and mock runners ( #17346 )
2026-07-31 19:41:18 -07:00
George Hotz and GitHub
a88f832f0c
remove UOp.val ( #17345 )
2026-07-31 18:58:38 -07:00
chenyu and GitHub
850989115d
__int__ and __float__ work for weak ( #17342 )
2026-07-31 19:56:37 -04:00
sirhcm and GitHub
15d515299e
heuristics: try multiple TC axes ( #17341 )
2026-07-31 19:25:18 -04:00
sirhcm and GitHub
85ced44db6
tc: don't allow reduce over output dims ( #17340 )
2026-07-31 16:48:51 -04:00
chenyu and GitHub
277433259e
fix sym_infer for CAST ( #17338 )
2026-07-31 14:38:38 -04:00
chenyu and GitHub
8dc225e28e
dtype_from_uop(INS) is None [PR] ( #17337 )
2026-07-31 14:33:15 -04:00
chenyu and GitHub
ad24750487
const(value, dtype) -> const(value).cast(dtype) in tests ( #17335 )
2026-07-31 13:24:31 -04:00
qazal and GitHub
a11ee26bb8
viz: prep for faster cli DEBUG=3 ( #17334 )
...
* move data
* split
2026-08-01 02:16:30 +09:00
sirhcm and GitHub
b95bd5b2a5
don't reset chestnut in benchmark ( #17333 )
2026-07-31 12:39:04 -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
7f4dbb8090
remove shape= from UOp.const [PR] ( #17331 )
...
inlined to const_like
2026-07-31 11:18:59 -04: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
nimlgen and GitHub
155b84ee80
hcq2 faster schedule ( #17324 )
...
* avoid quadratic STACK dtype promotion
* build HCQ patch stacks directly
* pack HCQ command buffers linearly
* remove HCQ command buffer simplification
2026-07-31 16:08:18 +03:00
wozeparrot and GitHub
4f5cadd15d
gptoss ci ( #17325 )
2026-07-31 05:51:58 -07:00
nimlgen and GitHub
0c4bfaeb48
coalesce ints ( #17323 )
...
* merge ints
* fix z3 validation of coalesced loads
* fix uint vector names in CUDA and Metal
2026-07-31 15:49:56 +03:00
qazal and GitHub
6d2700f0b7
failing test for unbound _device_num err in BEAM ( #17326 )
...
* min failing test
* switch to cpu
* err
2026-07-31 20:24:44 +09:00
wozeparrot and GitHub
93b74c75fc
gptoss: grouped moe ( #17322 )
2026-07-31 03:25:32 -07: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
sirhcm and GitHub
13452b3775
benchmark comma big model ( #17312 )
2026-07-31 01:37:34 -04: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
chenyu and GitHub
6608b9d8e2
adjust lower weak order [pr] ( #17314 )
...
* adjust lower weak order [pr]
* no symbolic_simple
2026-07-30 23:36:12 -04:00
Noah Schiro and GitHub
f2c2f4456b
Add softmin ( #17292 )
...
* Add softmin
* Remove extra tests
2026-07-30 20:05:11 -04:00
chenyu and GitHub
fe8ece7efa
fix weak for image gate fusion [pr] ( #17310 )
2026-07-30 19:47:47 -04:00
chenyu and GitHub
ce500c1946
broadcast doesn't cast CONST [pr] ( #17308 )
...
* broadcast doesn't cast CONST [pr]
it might change arg from python int to python float, but not cast the CONST, so we don't need or_casted in symbolic. we lower the weak dtype right before decomp
* fix
2026-07-30 18:51:59 -04:00
sirhcm and GitHub
b290372121
ftdi reset chestnut before running comma benchmark ( #17306 )
2026-07-30 14:28:45 -04:00
sirhcm and GitHub
da15c43e51
update openpilot benchmarks ( #17304 )
2026-07-30 14:22:20 -04:00
chenyu and GitHub
b488cc7df2
update dtype_from_uop for SHR/SHL ( #17307 )
...
certain backend cast the distance, and we should not re-broadcast it
2026-07-30 13:57:58 -04:00
chenyu and GitHub
341c4ed4f5
rewrite writes new dtype with dtype_from_uop try 2 [pr] ( #17305 )
...
with NIR fix which is INDEX related
some Ops currently depends on explicitly set dtype and not dtype_from_uop
2026-07-30 13:47:21 -04:00
chenyu and GitHub
d05a3e6c0b
Revert "rewrite writes new dtype with dtype_from_uop [pr] ( #17302 )" ( #17303 )
...
This reverts commit 417245ab26 .
2026-07-30 13:08:26 -04:00
chenyu and GitHub
417245ab26
rewrite writes new dtype with dtype_from_uop [pr] ( #17302 )
...
some Ops currently depends on explicitly set dtype and not dtype_from_uop
2026-07-30 12:46:24 -04:00
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