Commit Graph
364 Commits
Author SHA1 Message Date
qazalandGitHub c9cd44b751 more custom kernel contig input edge case tests (#17387) 2026-08-04 18:20:50 +09:00
geohot 87289a7410 hotfix: revert test_scalar_alu_index, violates spec 2026-08-03 15:25:51 -07:00
George HotzandGitHub c2625c78cb scalar ALU index fix + llm: preserve_thinking (#17381)
* cstyle: scalar ALU index fix, serve: preserve_thinking, test: fix Handler import

- cstyle.py: return scalar directly when ALU buffer has 1 element
- cli.py: add preserve_thinking param to FallbackTemplate.render
- serve.py: pass preserve_thinking=True when rendering chat completions
- test_llm_server.py: fix import to use Handler from llm.serve

* real fix
2026-08-03 15:10:21 -07:00
chenyuandGitHub b502fc1367 more const arg -> val (#17350) 2026-08-01 02:18:10 -04:00
qazalandGitHub 161783d8f7 add _device_num back to ast.variables (kimi) (#17327) 2026-08-01 14:56:50 +09:00
chenyuandGitHub ad24750487 const(value, dtype) -> const(value).cast(dtype) in tests (#17335) 2026-07-31 13:24:31 -04:00
chenyuandGitHub 7f4dbb8090 remove shape= from UOp.const [PR] (#17331)
inlined to const_like
2026-07-31 11:18:59 -04:00
chenyuandGitHub 8e2f175542 const(dtype, b) -> const(b, dtype) [PR] (#17328)
prep for dtype removal
2026-07-31 09:46:37 -04:00
qazalandGitHub 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
qazalandGitHub 0a3325f9c2 add mxfp4 quantize and layout kernels (#17320) 2026-07-31 14:50:18 +09:00
qazalandGitHub 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 HotzandGitHub 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 SchiroandGitHub f2c2f4456b Add softmin (#17292)
* Add softmin

* Remove extra tests
2026-07-30 20:05:11 -04:00
qazalandGitHub 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
George HotzandGitHub 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
George HotzandGitHub 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
sirhcmandGitHub aab51fb7b6 64-bit UOp.variable support, try 2 (#17256) 2026-07-29 14:39:36 -04:00
chenyuandGitHub 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
qazalandGitHub 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 HotzandGitHub 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 HotzandGitHub 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
chenyuandGitHub 23e9e76e8c DEFAULT_FLOAT/DEFAULT_INT ContextVar [pr] (#17265) 2026-07-28 19:08:18 -04:00
chenyuandGitHub 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
drkandGitHub d1c3ae0f91 preserve typed ranges through reshape (#17259) 2026-07-28 16:51:06 -04:00
qazalandGitHub 05e4727feb arange stack regression test (#17250)
* simple failing test

* commend out stack

* a little smaller

* work
2026-07-28 20:00:22 +09:00
chenyuandGitHub 4b7022e8f4 Revert "64-bit UOp.variable support (#17246)" (#17248)
This reverts commit ab8fb191b2.
2026-07-28 00:00:26 -04:00
sirhcmandGitHub ab8fb191b2 64-bit UOp.variable support (#17246) 2026-07-27 23:56:09 -04:00
chenyuandGitHub f45fc4c566 test updates from weak flip (#17231) 2026-07-27 01:27:42 -04:00
qazalandGitHub 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
chenyuandGitHub 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
chenyuandGitHub acc2374b6f minor normalize cleanup [pr] (#17215)
and add a few no-op clones in test_ops to prep for weak flip
2026-07-26 00:39:53 -04:00
chenyuandGitHub 492dc6d5fb delete _index_to_concrete_int [pr] (#17205)
staying weak is okay
2026-07-25 16:05:08 -04:00
chenyuandGitHub 0a6125e571 Invalid is bool, put weakint in promo lattice [pr] (#17188) 2026-07-25 00:06:53 -04:00
George HotzandGitHub f65001e297 cleanups because copy is not allowed in rangeify (kimi) (#17186)
* cleanups because copy is not allowed in rangeify (kimi)

* gpt cleanups
2026-07-24 14:17:10 -07:00
George HotzandGitHub dcad119415 convert COPY -> STORE early (#17172)
* convert COPY -> STORE early

* finish it (glm)

* no slop

* fixes

* not None

* copy one

* bugs fixed, just kernel count i think

* fix disk to disk

* move out of realize

* revert that

* simpler

* fixes

* add copy kernel simplification

* fixes

* dead copy

* remove reshape on store

* fix mstack lowering
2026-07-24 13:49:17 -07:00
chenyuandGitHub 38c04fa648 data_src concept in run_rangeify [pr] (#17183)
fixed a skip based on `dtypes.weakint`, it was pulling shape args
2026-07-24 14:20:41 -04:00
RaineandGitHub fcb1d3fc9f add x86 loops (#17179)
* add x86 loops

* lint
2026-07-24 07:26:09 -07:00
chenyuandGitHub 2dd2e3d786 don't support Tensor(list(np.array)) [pr] (#17175) 2026-07-24 03:13:28 -04:00
chenyuandGitHub a292dab8ab corner cases from weak const branch (#17173) 2026-07-24 00:09:56 -04:00
chenyuandGitHub f5d9c31d15 arange upcast to int64 with big N (#17167)
simplifies _one_hot_along_dim
2026-07-23 20:11:39 -04:00
George HotzandGitHub bf5989ea48 support loops in nir (kimi) (#17166)
* support loops in nir

* tpying
2026-07-23 17:05:15 -07:00
chenyuandGitHub 3e031d6c4d AND OR XOR SHL SHR cannot have float operands [PR] (#17164)
* AND OR XOR SHL SHR cannot have float operands [PR]

* dont
2026-07-23 18:51:56 -04:00
chenyuandGitHub 2983987321 clean up some tests (#17154)
duplicated or typo or slow
2026-07-23 11:22:18 -04:00
chenyuandGitHub 6b82c0cb95 minor test fixes (#17138)
TestCustomKernel square().max() masks NaN, and fixed test_tensor_cores_multi_reduce
2026-07-22 15:08:52 -04:00
sirhcmandGitHub c7027db715 lvp: skip test_masked_select (#17137) 2026-07-22 14:53:09 -04:00
nimlgenandGitHub a6fda6b102 add volatile to paramarg (#17127)
* volatile

* x

* x

* fo
2026-07-22 15:47:37 +03:00
nimlgenandGitHub 6e979b879b call inside c (#17131) 2026-07-22 13:49:06 +03:00
George HotzandGitHub 39924387b1 LOOP is srcless RANGE (kimi) (#17129)
* LOOP is srcless RANGE (kimi)

* upd

* cleanups

* cleanups

* we don't float anymore
2026-07-21 18:41:31 -07:00
chenyuandGitHub 92f9c850b4 fix pow(int, float) (#17126)
* fix pow(int, float)

* onnx
2026-07-21 18:48:24 -04:00
chenyuandGitHub f19a2ad771 single where mixin [pr] (#17118)
* single where mixin [pr]

no shape broadcasting in ufix and _broadcasted anymore

* QCOM vectorized bool is broken
2026-07-21 17:38:36 -04:00