chenyu and GitHub
8e524ca467
CONST related cleanups [pr] ( #17352 )
...
ConstFloat(nan) != nan should be False, and some Invalid bool cleanups
2026-08-01 09:54:28 -04:00
chenyu and GitHub
6c0ec39279
UOp.is_invalid [PR] ( #17351 )
...
helper to prep ConstArg
2026-08-01 02:36:32 -04:00
chenyu and GitHub
b502fc1367
more const arg -> val ( #17350 )
2026-08-01 02:18:10 -04:00
qazal and GitHub
161783d8f7
add _device_num back to ast.variables (kimi) ( #17327 )
2026-08-01 14:56:50 +09: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
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
ad24750487
const(value, dtype) -> const(value).cast(dtype) in tests ( #17335 )
2026-07-31 13:24:31 -04: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
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
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
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
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
b1tg and GitHub
e25bf77ce9
kimi delta attention ( #17281 )
...
* kimi delta attention
* config
* test
2026-07-30 08:23:56 -07: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
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
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
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
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
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
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
23e9e76e8c
DEFAULT_FLOAT/DEFAULT_INT ContextVar [pr] ( #17265 )
2026-07-28 19:08:18 -04: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
drk and GitHub
d1c3ae0f91
preserve typed ranges through reshape ( #17259 )
2026-07-28 16:51:06 -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
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
0bb36c9989
make Tensor(None) weakfloat ( #17241 )
...
match other float consts
2026-07-27 16:45:28 -04:00
nimlgen and GitHub
8b9ef157d1
run_linear in external_test_gpu_crash ( #17239 )
2026-07-27 22:00:11 +03:00
b1tg and GitHub
bdbb1d702f
fix shard axis through symbolic reshape ( #17238 )
...
* fix shard axis through symbolic reshape
* bind
2026-07-27 11:32:15 -04:00
chenyu and GitHub
818a892ebc
flip from_py to use weak dtypes [pr] ( #17229 )
2026-07-27 10:07:00 -04:00
chenyu and GitHub
f45fc4c566
test updates from weak flip ( #17231 )
2026-07-27 01:27:42 -04:00
qazal and GitHub
95e3b0066f
webgpu failing test for duplicate PARAM in CALL [pr] ( #17230 )
...
* webgpu failing test for duplicate PARAM in CALL [pr]
* typo
2026-07-27 14:06:06 +09:00
chenyu and GitHub
0f98212e80
skip test_float_to_fp8e4m3_extreme_values ( #17228 )
...
fp8 overflow behavior changed in torch 2.13.0, skipped the test for now
2026-07-26 21:01:51 -04:00
chenyu and GitHub
456b5b5060
fix python_alu inf ( #17227 )
2026-07-26 20:04:12 -04:00