nimlgen/tinygrad
1
0
Fork 0
forked from tinygrad/tinygrad
Code Pull Requests Activity
14,227 Commits 334 Branches 15 Tags
tensor_graph_single_sink
Commit Graph
6046 Commits
Author SHA1 Message Date
chenyuandGitHub 6cb419b9b7 regression test for bert nan with weak (#17412) 2026-08-05 12:14:26 -04:00
chenyuandGitHub 3bf9e70b19 Revert "don't cast weak in _broadcasted [pr] (#17408)" (#17409) 
This reverts commit b45058b5ec.
 2026-08-05 02:40:33 -04:00
RaineGitHubGeorge Hotz
 		77e124e455 fix AMD WMMA emulation and test in CI (#17184) 
* fix SPEC=1 test_tensor_cores

* implement i32 WMMA for RDNA3, add regression test to CI

* gfx950 scaled mfma llvmir fix

* detect VOP3PX2 in emu

* start cdna4 scaled mfam emu

* fix gfx12 llvmir signatures

* oops

* fix

* fix src2 const field extraction (gpt)

* scaled mfma fixes cdna4 (kimi)

* fp8 out breaks mfma (glm)

* fix const signature

* another

* fix f string for linter

* lint

* clean

* and a final lint

* .

* fix mypy

* skip slow tests on ci

* reduce unroll tensor shape -> 64x64

---------

Co-authored-by: George Hotz <[email protected]>
 2026-08-04 23:25:32 -07:00
chenyuandGitHub b45058b5ec don't cast weak in _broadcasted [pr] (#17408) 
* don't cast weak in _broadcasted [pr]

* fine now?
 2026-08-05 02:19:31 -04:00
George HotzandGitHub 46f0003776 more KernelCountException (#17407) 2026-08-04 22:55:38 -07:00
chenyuandGitHub 9b508dfafc remove invalid special case in cast [PR] (#17405) 2026-08-04 23:08:07 -04: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
chenyuandGitHub 3eab809e06 update minimum to not create strong type const [PR] (#17401) 2026-08-04 21:29:24 -04: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
chenyuandGitHub c1a10e0726 fix _min_max for CAST from float to int [pr] (#17396) 
* fix _min_max for CAST from float to int [pr]

* fix
 2026-08-04 18:06:50 -04:00
George HotzandGitHub f489f4bc36 add test_eye + color INDEX (#17394) 2026-08-04 11:08:08 -07:00
b1tgandGitHub 0796853845 support symbolic shapes in allreduce (#17364) 2026-08-04 07:53:12 -07:00
qazalandGitHub c9cd44b751 more custom kernel contig input edge case tests (#17387) 2026-08-04 18:20:50 +09:00
George HotzandGitHub c21a552f3d llm: bugfixes + warmup (#17384) 2026-08-03 18:23:14 -07:00
George HotzandGitHub 3cb786f447 llm: update test_llm_server tests (#17382) 2026-08-03 16:06:17 -07: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 314df72b5f deflake test_hcq with MOCKGPU (#17370) 
for MOCKGPU we compare with e2e wall time which would be device agnostic
 2026-08-03 13:22:59 +03:00
George HotzandGitHub 05bc7c6994 fix llm reasoning and Linear import (#17372) 2026-08-02 12:31:52 -07:00
chenyuandGitHub 59df317b12 use UOp.const to create new consts [PR] (#17368) 
replace arg won't work with ConstArg
 2026-08-02 13:31:55 -04:00
chenyuandGitHub 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
chenyuandGitHub 6c0ec39279 UOp.is_invalid [PR] (#17351) 
helper to prep ConstArg
 2026-08-01 02:36:32 -04: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
George HotzandGitHub a88f832f0c remove UOp.val (#17345) 2026-07-31 18:58:38 -07:00
chenyuandGitHub 850989115d __int__ and __float__ work for weak (#17342) 2026-07-31 19:56:37 -04:00
sirhcmandGitHub 85ced44db6 tc: don't allow reduce over output dims (#17340) 2026-07-31 16:48:51 -04:00
chenyuandGitHub 277433259e fix sym_infer for CAST (#17338) 2026-07-31 14:38:38 -04: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
nimlgenandGitHub 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
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
chenyuandGitHub fe8ece7efa fix weak for image gate fusion [pr] (#17310) 2026-07-30 19:47:47 -04:00
chenyuandGitHub 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
chenyuandGitHub 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
b1tgandGitHub e25bf77ce9 kimi delta attention (#17281) 
* kimi delta attention

* config

* test
 2026-07-30 08:23:56 -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
chenyuandGitHub d52ef3077c fix do_devectorize dtype (#17295) 
each one derives its dtype from src, not shared
 2026-07-29 23:04:25 -04:00
chenyuandGitHub fd912b348c generic c0*x<c1 [pr] (#17291) 2026-07-29 18:45:23 -04: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
chenyuandGitHub 027907a544 fix c0*x<c1 symbolic [pr] (#17290) 
use `//` and not math.ceil
 2026-07-29 18:32:31 -04: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 dd86a30798 test case for weakfloat cast to weakint INDEX (#17287) 2026-07-29 14:15:52 -04:00