sirhcm/tinygrad
1
0
Fork 0
forked from tinygrad/tinygrad
Code Pull Requests 2 Actions 8 Activity
14,300 Commits 336 Branches 15 Tags
simple_rangeify
Commit Graph
377 Commits
Author SHA1 Message Date
RaineandGitHub d41ca5e60f Fix WMMA CI (#17479) 
* init

* split into sub tests

* trigger ci
 2026-08-10 08:39:14 -07:00
qazalandGitHub 44f1f45cd5 llama: custom silu kernels (#17462) 
* start by copying the C

* uop kernel

* cleanup tests

* estimates is part of SPEC
 2026-08-10 16:43:01 +09:00
Robert JosephandGitHub 0c96cdc300 fix prod gradients at zero (#17404) 2026-08-07 09:56:52 -04:00
qazalGitHubgeohot
 		f253c4469d remove contiguous from custom_kernel (#17149) 
* no user contig on custom_kernel

* clean up

* non removable

* test MXFP4 llama without hipcc

* use compiler-free HIPCC renderer in llama CI

* move llama coverage to AMD tests

* run llama coverage in AMD test matrix

* respect configured ROCm path in llama profile

* work

* clean up

* fix

* add views back

* remove that

* update test

* test_double_permute one less kernel

* test_shrink less kernels

---------

Co-authored-by: George Hotz <[email protected]>
 2026-08-07 11:20:59 +09:00
chenyuandGitHub 28195d51fb fix f2f from fp8e5m2fnuz to half (#17442) 
* fix f2f from fp8e5m2fnuz to half

* it works if it's supported
 2026-08-06 21:19:46 -04:00
chenyuandGitHub 9020a88f03 truncate float in DType.const [pr] (#17439) 2026-08-06 20:19:08 -04:00
qazalandGitHub f258708d7d llama: custom quantize_mxfp4+transpose kernel (codex) (#17434) 
* llama: custom quantize_mxfp4+transpose kernel (codex)

* rename to cpp

* inline

* cleanup

* lds load_bf16x4

* more tests, add Estimates
 2026-08-06 16:13:28 +09:00
George HotzandGitHub b4372df9c6 revert wrong custom kernel fix (#17427) 2026-08-05 18:31:35 -07:00
George HotzandGitHub 470c032a5e fix slice + non contig kernels (#17423) 
* movement: resolve negative int slice bounds against symbolic sizes

negative int bounds in a slice against a symbolic dim were passed through
unresolved, giving wrong views. resolve them against the (possibly
symbolic) size, like slice.indices does for int dims

* schedule: realize custom kernel inputs that don't resolve to a buffer state

rangeify assigns ranges backward from consumers and CALL contributes none,
so the subgraph above a custom kernel input gets no ranges unless something
in it is realized, and reduce conversion crashes with a KeyError. realize
call inputs that don't resolve to a buffer state.

only view-only movement ops preserve the underlying buffer: anything
computed (ALU, REDUCE, ...) must be realized even if one of its sources
resolves to a buffer, since the whole subgraph above the call has no
ranges. unwrapping src[0] unconditionally missed const branches hanging
off non-src[0] children and silently resolved REDUCEs to their source
buffer. includes regression tests for pure const, mixed buffer+const, and
view-over-buffer inputs
 2026-08-05 16:18:17 -07:00
George HotzandGitHub 46f0003776 more KernelCountException (#17407) 2026-08-04 22:55:38 -07:00
George HotzandGitHub e1f42681fa add new schedule tests + format better (#17402) 
* add new schedule tests + format better

* assert_kernel_count
 2026-08-04 18:46:38 -07:00
George HotzandGitHub 6122b3c98f use check_schedule in tests where possible (#17400) 2026-08-04 18:17:29 -07:00
chenyuandGitHub d79772f057 fix pow on extreme inputs (#17397) 
* fix pow on extreme inputs

* WEBGPU
 2026-08-04 19:30:43 -04:00
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