Commit Graph

11089 Commits

Author SHA1 Message Date
George Hotz
1afa3c0877 vmap on full model (#13340)
* vmap on full model

* vmap gemm

* reduce sums on end

* outer reduce

* only if there's ranges

* put those rules in symbolic

* ranges

* do opt later

* add zero range
2025-11-18 16:06:06 -08:00
chenyu
46cb65e692 delete rules from sym [pr] (#13339) 2025-11-18 14:57:35 -05:00
George Hotz
9c59b3d19e vmap grad needs reduce_backward (#13336)
* vmap grad needs reduce_backward

* fuse and outer
2025-11-18 10:08:30 -08:00
qazal
a647c9eca6 sqtt ui minor fixes (#13335)
* roc.py cleanups

* direct append

* viz index cleanup

* simd row details
2025-11-19 01:27:56 +08:00
George Hotz
06e39a88a9 outer vmap works (#13334)
* outer vmap works

* fuse works

* vmap outer works

* outer ranges work

* grad work

* should be good to merge
2025-11-18 09:27:48 -08:00
chenyu
805de27e07 no load substitute in uop_given_valid [pr] (#13333) 2025-11-18 11:47:58 -05:00
chenyu
05294bc648 fix some mypy cast [pr] (#13331) 2025-11-18 09:23:42 -05:00
qazal
5623e765c8 VIZ=2 enables SQTT (#13330) 2025-11-18 22:20:31 +08:00
nimlgen
331f70aa75 roc: ctrlc (#13255)
* roc: ctrl-c works

* rm
2025-11-18 19:29:28 +08:00
George Hotz
583560ab72 this is the right way to write vmap (#13328) 2025-11-17 20:20:52 -08:00
Christopher Milan
8e8e53c886 int8_t is c_byte (#13326) 2025-11-17 21:25:40 -05:00
George Hotz
e4fead8a86 write scan in uops (#13321)
* write scan in uops

* ops range

* no need for variable

* meh, later

* shorter
2025-11-17 16:58:08 -08:00
wozeparrot
8894a5409d feat: hipcc compiler (#13319) 2025-11-17 15:13:32 -08:00
George Hotz
6d3385c284 print special ops in postrange (#13318)
* print special ops in postrange

* fix on OSX
2025-11-17 14:43:23 -08:00
chenyu
b637093be9 remove a few rules in pm_lower_index_dtype [pr] (#13317) 2025-11-17 17:04:56 -05:00
George Hotz
98e9e73286 hotfix: amd_uop_matmul getenvs 2025-11-17 13:26:01 -08:00
qazal
e7e1935225 cleanup sqtt/test_timing (#13315) 2025-11-18 04:28:05 +08:00
wozeparrot
33773fda87 tk initial mi350 (#13289) 2025-11-17 11:46:32 -08:00
nimlgen
e2cee64050 Revert "hcq: add tag to exec events (#13311)" (#13314)
This reverts commit f63ded5817.
2025-11-17 22:15:31 +03:00
chenyu
646372490c move tiktoken import in llama3 (#13316)
only Tokenizer requires that
2025-11-17 14:09:37 -05:00
qazal
a37f221e44 viz: visualize waves in the timeline (#13292)
* viz: visualize waves in the timeline

* timeline in format

* per step

* rm that
2025-11-17 22:04:21 +08:00
nimlgen
f63ded5817 hcq: add tag to exec events (#13311)
* hcq: add tag to exec events

* f

* fix

* fix
2025-11-17 16:59:30 +03:00
qazal
50a443f558 viz: add shader engine to wave exec payload (#13310)
* viz: show sqtt shader engine

* order it from smallest unit

* easier to config
2025-11-17 19:11:34 +08:00
nimlgen
9bb17c53ea amd: timer fix (#13267) 2025-11-17 13:59:03 +03:00
George Hotz
55be95da15 cleanup sqtt raw parser (#13309)
* cleanup sqtt raw parser

* better names (don't merge yet)

* clean up amd

* a few more names

* one more filter
2025-11-16 13:11:51 -08:00
George Hotz
cabd4add48 more work parsing SQTT, separate VIZ/PROFILE (#13308)
* more work parsing SQTT

* more minimal runner

* sep VIZ/PROFILE

* parse print new

* improve parser

* more filter

* that

* split them

* lil cleanup

* skip flaky test

* AQL in mmapeak
2025-11-16 10:40:39 -08:00
qazal
13efdf8c31 test s_nop stall (#13307) 2025-11-17 00:59:39 +08:00
George Hotz
295600dc5a saturday coffee shop work parsing the att format (#13295)
* saturday coffee shop work parsing the att format

* add examples

* parser

* classes of packets

* fully vibe coded parser

* vibing

* empty

* some vibe names

* vibes

* most of these are wrong

* more vibes

* better names

* parsing

* parse

* cleanup parser

* touchups
2025-11-16 08:25:51 -08:00
Christopher Milan
a9ed241172 properly suppress NIRRenderer.__del__ error (#13299) 2025-11-16 18:58:04 +03:00
qazal
c70b06ec19 sqtt test_timing work (#13304)
* sqtt test_timing cleanups

* only the instruction

* v_mfma_f32_16x16x32_f16 16 cycles, only after second one though
2025-11-16 23:49:24 +08:00
chenyu
8f0e747b3a Tensor._tri with arange (#13297) 2025-11-16 10:21:16 -05:00
chenyu
6372c95094 disable benchmark MobileNetV2 on DSP (#13305)
failed on tinyc2
2025-11-16 09:42:52 -05:00
Christopher Milan
61625a3898 fix objc finalizing bug (#13296) 2025-11-16 12:43:04 +03:00
nimlgen
acbe6361ab qcom: suppress_finalizing to free (#13294) 2025-11-16 11:49:23 +03:00
wozeparrot
ef42334239 tk: load store cleanup (#13290) 2025-11-15 17:08:23 -08:00
chenyu
e8844853ed Tensor.eye with arange (#13287)
with rangify we can write these with arange
2025-11-15 12:32:27 -05:00
Christopher Milan
5b823af696 Remove (pypi) clang dep for autogen (#13284)
* no more clang

* regen comgr_3

* ci doesn't need pypi clang

* fix objc

* REGEN for libclang

---------

Co-authored-by: George Hotz <72895+geohot@users.noreply.github.com>
2025-11-15 09:05:11 -08:00
George Hotz
df53c62a9f bump line count 2025-11-15 08:16:20 -08:00
nimlgen
d37e1fe065 nv: wait for wpr to reset (#13282)
* nv: wait for wpr to reet

* fix

* comment

* wai

* f

* fix
2025-11-15 20:00:49 +08:00
George Hotz
22c08b470c fold using outerworld range (#13286)
* scan using outerworld range

* almost

* sched

* simple range

* mypy

* woooo outer range

* spec passes

* print the numbers

* lol it runs

* real test
2025-11-14 20:43:41 -08:00
George Hotz
567066f51f tests for cast there and back (#13195)
* fix cast folding in llama

* dtypes that work everywhere

* Skip test_cast_there_and_back for backend casts

Skip test due to backend casting issues.
2025-11-14 16:56:09 -08:00
George Hotz
6c5fa349e1 add (unused) outer range (#13285) 2025-11-14 16:47:52 -08:00
Christopher Milan
d1bb08c5a1 In-tree autogen: objective c (#13223)
* checkout changes from autogen branch

* move assert

---------

Co-authored-by: George Hotz <72895+geohot@users.noreply.github.com>
2025-11-14 14:08:42 -08:00
George Hotz
e5351699bd openpilot warp (#13283)
* openpilot image warp test

* 0.4 ms on metal, 1 ms on CPU

* new inputs each time

* reshape
2025-11-14 13:55:32 -08:00
qazal
7c110e1a57 viz: minor cleanups for sqtt (#13275)
* small prg cleanup

* test_timing
2025-11-15 01:08:56 +08:00
chenyu
888aaab151 test_tiny cleanup (#13276) 2025-11-14 11:11:32 -05:00
nimlgen
3e63831b98 nv: support 580+ drivers (#13269)
* nv: 580+ support

* start

* f

* fake

* fix
2025-11-14 21:44:16 +08:00
qazal
2ee701a009 roc: fix CEnum access (#13270)
* roc: add decoder to ci

* also add installer

* use CEnum syntax

* try 2

* add to setup

* revert ci change

* the other enum too
2025-11-14 21:41:24 +08:00
nimlgen
c80d459d99 autogen: fix packed args structs (#13274)
* autogen: fix packed args structs

* and test this
2025-11-14 20:24:06 +08:00
nimlgen
14eb48b13a autogen: rename nv_gpu to nv_570 (#13273)
* autogen: rename nv_gpu to nv_570

* rename
2025-11-14 20:07:19 +08:00