nimlgen/tinygrad
1
0
Fork 0
forked from tinygrad/tinygrad
Code Pull Requests Activity
14,341 Commits 334 Branches 15 Tags
master
Commit Graph
386 Commits
Author SHA1 Message Date
chenyuandGitHub 138fb4a783 delete dead DType.scalar [PR] (#17561) 2026-08-16 21:12:17 -04:00
chenyuandGitHub 057a18a07c fix emulated long cast to double (#17559) 2026-08-16 20:52:42 -04:00
George HotzandGitHub e688e07758 add max_shape/max_numel to mixins + pad_to (#17553) 2026-08-15 20:05:10 -07:00
chenyuandGitHub 417563ca20 fix webgpu is_nan [pr] (#17551) 2026-08-15 16:31:19 -04:00
chenyuandGitHub fae893753b no pm_fold_cast_const in get_kernel_graph [pr] (#17546) 
* no pm_fold_cast_const in get_kernel_graph [pr]

* maybe
 2026-08-15 09:57:05 -04:00
chenyuandGitHub ac7067ac60 fix deconstruct_function for python 3.11 (#17534) 2026-08-14 13:38:34 -04:00
nimlgenandGitHub 6b35220622 cpu hcq2 (#17503) 
* cpu hcq2

* temp

* slop

* test with backpressure

* x

* x

* x

* x

* x

* x

* Dx

* save reverts

* um?

* x

* x

* call from py

* x?

* x

* submitters gone

* x

* x

* z

* Dx

* Dx

* x

* x

* fixes

* repl

* x

* f

* for now keep hcqbuffer
 2026-08-14 15:06:03 +03:00
George HotzandGitHub b1859805b1 remove Ops.BIND (#17511) 
* remove Ops.BIND

* param arg

* simplify that

* simplify

* cleaner

* props, not functions

* param and buffer can share
 2026-08-13 23:52:16 -07:00
qazalandGitHub 2e5a9a4121 no hardcoded device names in test_sliced_buffer_function (#17501) 2026-08-12 15:19:25 +09:00
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