geohot/tinygrad
1
0
Fork 0
forked from tinygrad/tinygrad
Code Pull Requests Actions Activity
14,474 Commits 4 Branches 15 Tags
master
Commit Graph
45 Commits
Author SHA1 Message Date
George HotzandGitHub c31038ff37 use KernelCountException when kernel count is being compared (#17584) 2026-08-18 16:06:03 -07:00
qazalandGitHub 2e5a9a4121 no hardcoded device names in test_sliced_buffer_function (#17501) 2026-08-12 15:19:25 +09: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
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 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
qazalandGitHub c9cd44b751 more custom kernel contig input edge case tests (#17387) 2026-08-04 18:20:50 +09: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
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
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
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
drkandGitHub d1c3ae0f91 preserve typed ranges through reshape (#17259) 2026-07-28 16:51:06 -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 0a6125e571 Invalid is bool, put weakint in promo lattice [pr] (#17188) 2026-07-25 00:06:53 -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
George HotzandGitHub fdffc6c0c8 remove dtypes base (#16931) 
* remove dtypes base

* find/replace bug

* vcount is junk too
 2026-07-08 11:48:11 -07:00
George HotzandGitHub a145fdce4b remove the .gep method (#16867) 2026-07-04 14:07:09 -07:00
chenyuandGitHub 3a9b87e48f remove Ops.DEVICE from Ops.PROGRAM [PR] (#16797) 2026-06-29 17:41:31 -04:00
chenyuandGitHub 1e08c0a07c remove NOOP from AFTER with multiple srcs (#16646) 2026-06-17 14:35:02 -04:00
chenyuandGitHub 1acc40600d indexing an after with all fully invalid stores is invalid (#16643) 
* indexing an after with all fully invalid stores is invalid

* typing cast
 2026-06-17 11:06:36 -04:00
qazalandGitHub 4d34590b7d llama: less E kernels (#16517) 2026-06-12 19:49:25 +09:00
qazalandGitHub 34481830f1 rangeify: fix cost function for AFTER(out, CALL) (#16559) 
* simple failing test

* fix rangeify cost function

* new ops count
 2026-06-10 17:30:50 +09:00
chenyuandGitHub fa14cde05c test update for arange and eye (#16394) 
these will need explicit clone to make a buffer
 2026-05-26 22:48:34 -04:00
chenyuandGitHub 8631b6f17d remove use of requires_grad in test/ (#16237) 2026-05-16 17:21:07 -04:00
wozeparrotandGitHub 2d48d7ab09 remove more invalid (#16227) 2026-05-16 02:52:27 -07:00
George HotzandGitHub 64c81dfd24 add all codegen stages to spec_tensor (#16163) 2026-05-12 10:35:38 -07:00
qazalandGitHub fe39cf148a add Ops.SOURCE test (#16155) 
* simple failing test

* raises

* change
 2026-05-12 22:49:32 +09:00
George HotzandGitHub 8294d105a7 Update the spec in spec.py to match the current state (#16132) 
* start work on specv2

* more spec

* more spec

* fix amd emulator

* more spec

* more

* fix test_uop_graph

* move those

* spec=2

* skip those questionable tests

* ptx fix

* more spec=2

* store

* allow custom function in tensor

* spec 2

* fix beam search for tensor cores

* delete the old specs

* fix import
 2026-05-11 20:07:47 -07:00
qazalandGitHub 8b147a9ed5 minimal repro for llama copies 2 (#16011) 2026-05-01 22:23:47 +09:00
qazalandGitHub 58b34e71bd failing test for llama useless copies (#15989) 2026-05-01 00:55:29 +09:00
nimlgenandGitHub d2ab6ea7a6 remove schedule batch 3 (#15924) 
* remove shcedule batch 3

* batch 6

* batch 7
 2026-04-25 11:53:16 +03:00
chenyuandGitHub 7a1adfd2aa update Tensor.allclose to return Tensor (#15904) 
matches jax
 2026-04-24 08:27:17 -04:00
chenyuandGitHub 9192c93b7e Tensor.invalid -> Tesnor.invalids (#15849) 
matches ones and zeros, and to not share name with UOp.invalid
 2026-04-21 11:19:51 -04:00
chenyuandGitHub 8bd4fead26 UOp.size -> prod(max_shape) (#15755) 
and more test updates
 2026-04-15 22:41:30 -04:00
chenyuandGitHub 10c262ced8 update tests that use UOp.size (#15753) 2026-04-15 21:58:27 -04:00
chenyuandGitHub 8fdef2d3e4 mean/std/var to mixin (#15593) 2026-04-03 10:42:41 -04:00
chenyuandGitHub fc5b94b902 fix UOp.where(const, const) (#15560) 
* fix UOp.where(const, const)

* fix
 2026-04-01 05:28:49 -04:00
chenyuandGitHub f0eaac4235 reduce mixin (#15523) 2026-03-30 05:23:58 -04:00
George HotzandGitHub 6109117af1 anonymous buffers are Invalid (#15336) 
* anonymous buffers are Invalid

* unique_const

* work

* remove invalid writes

* test_anonymous_buffers_in_function
 2026-03-18 14:52:56 +08:00
George HotzandGitHub 677145b393 all consts have shapes (#14959) 
* all consts have shapes

* vconst has shape too

* use normal schedule

* cast ptrdtype

* image

* bitcast issue + hack
 2026-02-23 10:26:50 +08:00
George HotzandGitHub 2611907afb start ripping out old scheduler -- no maps (#14909) 
* start ripping out old scheduler -- no maps

* no more metadata
 2026-02-20 21:05:04 +08:00
George HotzandGitHub c331798201 move tests to test/backend (#14691) 
* move tests to test/backend

* fix imports

* fix CI

* revert that one

* Fix formatting in README for test command
 2026-02-12 11:09:44 +08:00