Merge branch 'refs/heads/model-manager-improvements' into model-selector-multi-runner-debug-model-download

Enable model label button only when conditions are met
Previously, the button's state update was misplaced, leading to potential issues with its interactive availability. The logic has been adjusted to ensure it is properly enabled or disabled based on onroad status and download progress. This change improves UX consistency and prevents unintended actions.
2026-07-20 22:57:03 +08:00 · 2025-01-05 22:55:45 +01:00 · 2025-01-05 22:54:29 +01:00 · 2025-01-05 22:50:05 +01:00 · 2025-01-05 22:28:53 +01:00 · 2025-01-05 22:28:26 +01:00
215 changed files with 12280 additions and 1052 deletions
@@ -3,6 +3,8 @@
 # to move existing files into LFS:
 # git add --renormalize .
 *.onnx filter=lfs diff=lfs merge=lfs -text
+*.thneed filter=lfs diff=lfs merge=lfs -text
+*.pkl filter=lfs diff=lfs merge=lfs -text
 *.svg filter=lfs diff=lfs merge=lfs -text
 *.png filter=lfs diff=lfs merge=lfs -text
 *.gif filter=lfs diff=lfs merge=lfs -text
@@ -0,0 +1,15 @@
+---
+name: Bug fix
+about: For openpilot bug fixes
+title: ''
+labels: 'bugfix'
+assignees: ''
+---
+
+**Description**
+
+<!-- A description of the bug and the fix. Also link the issue if it exists. -->
+
+**Verification**
+
+<!-- Explain how you tested this bug fix. -->
@@ -0,0 +1,19 @@
+---
+name: Car Bug fix
+about: For vehicle/brand specific bug fixes
+title: ''
+labels: 'car bug fix'
+assignees: ''
+---
+
+**Description**
+
+<!-- A description of the bug and the fix. Also link the issue if it exists. -->
+
+**Verification**
+
+<!-- Explain how you tested this bug fix. -->
+
+**Route**
+
+Route: [a route with the bug fix]
@@ -0,0 +1,15 @@
+---
+name: Car port
+about: For new car ports
+title: ''
+labels: 'car port'
+assignees: ''
+---
+
+**Checklist**
+
+- [ ] added entry to CAR in selfdrive/car/*/values.py and ran `selfdrive/car/docs.py` to generate new docs
+- [ ] test route added to [routes.py](https://github.com/commaai/openpilot/blob/master/selfdrive/car/tests/routes.py)
+- [ ] route with openpilot:
+- [ ] route with stock system:
+- [ ] car harness used (if comma doesn't sell it, put N/A):
@@ -0,0 +1,13 @@
+---
+name: Fingerprint
+about: For adding fingerprints to existing cars
+title: ''
+labels: 'fingerprint'
+assignees: ''
+---
+
+**Car**
+Which car (make, model, year) this fingerprint is for
+
+**Route**
+A route with the fingerprint
@@ -0,0 +1,15 @@
+---
+name: Refactor
+about: For code refactors
+title: ''
+labels: 'refactor'
+assignees: ''
+---
+
+**Description**
+
+<!-- A description of the refactor, including the goals it accomplishes. -->
+
+**Verification**
+
+<!-- Explain how you tested the refactor for regressions. -->
@@ -0,0 +1,31 @@
+---
+name: Tuning
+about: For openpilot tuning changes
+title: ''
+labels: 'tuning'
+assignees: ''
+---
+
+**Description**
+
+<!-- A description of what is wrong with the current tuning and how the PR addresses this. -->
+
+**Verification**
+
+<!-- To verify tuning, capture the following scenarios (broadly, not exactly), with current tune and this tune.
+Use the PlotJuggler tuning layout to compare planned versus actual behavior.
+
+Run ./juggle.py <route> --layout layouts/tuning.xml , screenshot the full tab of interest, and paste into this PR.
+
+Longitudinal:
+* Maintaining speed at 25, 40, 65mph
+* Driving up and down hills
+* Accelerating from a stop
+* Decelerating to a stop
+* Following large changes in set speed
+* Coming to a stop behind a lead car
+
+Lateral:
+* Straight driving at ~25, ~45 and ~65mph
+* Turns driving at ~25, ~45 and ~65mph
+-->
@@ -0,0 +1,30 @@
+import pathlib
+
+GITHUB_FOLDER = pathlib.Path(__file__).parent
+
+PULL_REQUEST_TEMPLATES = (GITHUB_FOLDER / "PULL_REQUEST_TEMPLATE")
+
+order = ["fingerprint", "car_bugfix", "bugfix", "car_port", "refactor"]
+
+def create_pull_request_template():
+  with open(GITHUB_FOLDER / "pull_request_template.md", "w") as f:
+    f.write("<!-- Please copy and paste the relevant template -->\n\n")
+
+    for t in order:
+      template = PULL_REQUEST_TEMPLATES / f"{t}.md"
+      text = template.read_text()
+
+      # Remove metadata for GitHub
+      start = text.find("---")
+      end = text.find("---", start+1)
+      text = text[end + 4:]
+
+      # Remove comments
+      text = text.replace("<!-- ", "").replace("-->", "")
+
+      f.write(f"<!--- ***** Template: {template.stem.replace('_', ' ').title()} *****\n")
+      f.write(text)
+      f.write("\n\n")
+      f.write("-->\n\n")
+
+create_pull_request_template()
@@ -74,6 +74,7 @@ comma*.sh
 selfdrive/modeld/thneed/compile
 selfdrive/modeld/models/*.thneed
 selfdrive/modeld/models/*.pkl
+sunnypilot/modeld/thneed/compile

 *.bz2
 *.zst
@@ -15,4 +15,4 @@
  url = https://github.com/commaai/teleoprtc
 [submodule "tinygrad"]
  path = tinygrad_repo
-  url = https://github.com/tinygrad/tinygrad.git
+  url = https://github.com/commaai/tinygrad.git
@@ -49,6 +49,10 @@ AddOption('--ccflags',
          default='',
          help='pass arbitrary flags over the command line')

+AddOption('--snpe',
+          action='store_true',
+          help='use SNPE on PC')
+
 AddOption('--external-sconscript',
          action='store',
          metavar='FILE',
@@ -392,6 +396,8 @@ SConscript(['third_party/SConscript'])

 SConscript(['selfdrive/SConscript'])

+SConscript(['sunnypilot/SConscript'])
+
 if Dir('#tools/cabana/').exists() and GetOption('extras'):
  SConscript(['tools/replay/SConscript'])
  if arch != "larch64":
@@ -64,6 +64,11 @@ struct ModelManagerSP @0xaedffd8f31e7b55d {
    progress @1 :Float32;
    eta @2 :UInt32;
  }
+  
+  enum Runner {
+    snpe @0;
+    tinygrad @1;
+  }

  struct ModelBundle {
    index @0 :UInt32;
@@ -73,7 +78,6 @@ struct ModelManagerSP @0xaedffd8f31e7b55d {
    status @4 :DownloadStatus;
    generation @5 :UInt32;
    environment @6 :Text;
-    is20hz @7 :Bool;
  }
 }

@@ -0,0 +1,49 @@
+using Cxx = import "./include/c++.capnp";
+$Cxx.namespace("cereal");
+
+@0xa086df597ef5d7a0;
+
+# Geometry
+struct Point {
+  x @0: Float64;
+  y @1: Float64;
+  z @2: Float64;
+}
+
+struct PolyLine {
+  points @0: List(Point);
+}
+
+# Map features
+struct Lane {
+  id @0 :Text;
+
+  leftBoundary @1 :LaneBoundary;
+  rightBoundary @2 :LaneBoundary;
+
+  leftAdjacentId @3 :Text;
+  rightAdjacentId @4 :Text;
+
+  inboundIds @5 :List(Text);
+  outboundIds @6 :List(Text);
+
+  struct LaneBoundary {
+    polyLine @0 :PolyLine;
+    startHeading @1 :Float32; # WRT north
+  }
+}
+
+# Map tiles
+struct TileSummary {
+  version @0 :Text;
+  updatedAt @1 :UInt64; # Millis since epoch
+
+  level @2 :UInt8;
+  x @3 :UInt16;
+  y @4 :UInt16;
+}
+
+struct MapTile {
+  summary @0 :TileSummary;
+  lanes @1 :List(Lane);
+}
@@ -133,6 +133,7 @@ std::unordered_map<std::string, uint32_t> keys = {
    {"GsmRoaming", PERSISTENT | BACKUP},
    {"HardwareSerial", PERSISTENT},
    {"HasAcceptedTerms", PERSISTENT},
+    {"IMEI", PERSISTENT},
    {"InstallDate", PERSISTENT},
    {"IsDriverViewEnabled", CLEAR_ON_MANAGER_START},
    {"IsEngaged", PERSISTENT},
@@ -255,29 +255,6 @@ bool ends_with(const std::string& s, const std::string& suffix) {
         strcmp(s.c_str() + (s.size() - suffix.size()), suffix.c_str()) == 0;
 }

-std::string strip(const std::string &str) {
-  auto should_trim = [](unsigned char ch) {
-    // trim whitespace or a null character
-    return std::isspace(ch) || ch == '\0';
-  };
-
-  size_t start = 0;
-  while (start < str.size() && should_trim(static_cast<unsigned char>(str[start]))) {
-    start++;
-  }
-
-  if (start == str.size()) {
-    return "";
-  }
-
-  size_t end = str.size() - 1;
-  while (end > 0 && should_trim(static_cast<unsigned char>(str[end]))) {
-    end--;
-  }
-
-  return str.substr(start, end - start + 1);
-}
-
 std::string check_output(const std::string& command) {
  char buffer[128];
  std::string result;
@@ -74,7 +74,6 @@ float getenv(const char* key, float default_val);
 std::string hexdump(const uint8_t* in, const size_t size);
 bool starts_with(const std::string &s1, const std::string &s2);
 bool ends_with(const std::string &s, const std::string &suffix);
-std::string strip(const std::string &str);

 // ***** random helpers *****
 int random_int(int min, int max);
@@ -4,7 +4,7 @@

 A supported vehicle is one that just works when you install a comma device. All supported cars provide a better experience than any stock system. Supported vehicles reference the US market unless otherwise specified.

-# 290 Supported Cars
+# 289 Supported Cars

 |Make|Model|Supported Package|ACC|No ACC accel below|No ALC below|Steering Torque|Resume from stop|<a href="##"><img width=2000></a>Hardware Needed<br>&nbsp;|Video|
 |---|---|---|:---:|:---:|:---:|:---:|:---:|:---:|:---:|
@@ -56,7 +56,6 @@ A supported vehicle is one that just works when you install a comma device. All
 |Genesis|GV60 (Performance Trim) 2022-23[<sup>5</sup>](#footnotes)|All|openpilot available[<sup>1</sup>](#footnotes)|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 Hyundai K connector<br>- 1 RJ45 cable (7 ft)<br>- 1 comma 3X<br>- 1 comma power v2<br>- 1 harness box<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=Genesis&model=GV60 (Performance Trim) 2022-23">Buy Here</a></sub></details>||
 |Genesis|GV70 (2.5T Trim, without HDA II) 2022-23[<sup>5</sup>](#footnotes)|All|Stock|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 Hyundai L connector<br>- 1 RJ45 cable (7 ft)<br>- 1 comma 3X<br>- 1 comma power v2<br>- 1 harness box<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=Genesis&model=GV70 (2.5T Trim, without HDA II) 2022-23">Buy Here</a></sub></details>||
 |Genesis|GV70 (3.5T Trim, without HDA II) 2022-23[<sup>5</sup>](#footnotes)|All|Stock|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 Hyundai M connector<br>- 1 RJ45 cable (7 ft)<br>- 1 comma 3X<br>- 1 comma power v2<br>- 1 harness box<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=Genesis&model=GV70 (3.5T Trim, without HDA II) 2022-23">Buy Here</a></sub></details>||
-|Genesis|GV70 Electrified (Australia Only) 2022[<sup>5</sup>](#footnotes)|All|openpilot available[<sup>1</sup>](#footnotes)|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 Hyundai Q connector<br>- 1 RJ45 cable (7 ft)<br>- 1 comma 3X<br>- 1 comma power v2<br>- 1 harness box<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=Genesis&model=GV70 Electrified (Australia Only) 2022">Buy Here</a></sub></details>||
 |Genesis|GV70 Electrified (with HDA II) 2023[<sup>5</sup>](#footnotes)|Highway Driving Assist II|openpilot available[<sup>1</sup>](#footnotes)|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 Hyundai Q connector<br>- 1 RJ45 cable (7 ft)<br>- 1 comma 3X<br>- 1 comma power v2<br>- 1 harness box<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=Genesis&model=GV70 Electrified (with HDA II) 2023">Buy Here</a></sub></details>||
 |Genesis|GV80 2023[<sup>5</sup>](#footnotes)|All|Stock|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 Hyundai M connector<br>- 1 RJ45 cable (7 ft)<br>- 1 comma 3X<br>- 1 comma power v2<br>- 1 harness box<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=Genesis&model=GV80 2023">Buy Here</a></sub></details>||
 |GMC|Sierra 1500 2020-21|Driver Alert Package II|openpilot available[<sup>1</sup>](#footnotes)|0 mph|6 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-empty.svg)](##)|<details><summary>Parts</summary><sub>- 1 GM connector<br>- 1 comma 3X<br>- 1 harness box<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=GMC&model=Sierra 1500 2020-21">Buy Here</a></sub></details>|<a href="https://youtu.be/5HbNoBLzRwE" target="_blank"><img height="18px" src="assets/icon-youtube.svg"></img></a>|
@@ -80,7 +79,7 @@ A supported vehicle is one that just works when you install a comma device. All
 |Honda|Odyssey 2018-20|Honda Sensing|openpilot|26 mph|0 mph|[![star](assets/icon-star-empty.svg)](##)|[![star](assets/icon-star-empty.svg)](##)|<details><summary>Parts</summary><sub>- 1 Honda Nidec connector<br>- 1 RJ45 cable (7 ft)<br>- 1 comma 3X<br>- 1 comma power v2<br>- 1 harness box<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=Honda&model=Odyssey 2018-20">Buy Here</a></sub></details>||
 |Honda|Passport 2019-23|All|openpilot|26 mph|12 mph|[![star](assets/icon-star-empty.svg)](##)|[![star](assets/icon-star-empty.svg)](##)|<details><summary>Parts</summary><sub>- 1 Honda Nidec connector<br>- 1 RJ45 cable (7 ft)<br>- 1 comma 3X<br>- 1 comma power v2<br>- 1 harness box<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=Honda&model=Passport 2019-23">Buy Here</a></sub></details>||
 |Honda|Pilot 2016-22|Honda Sensing|openpilot|26 mph|12 mph|[![star](assets/icon-star-empty.svg)](##)|[![star](assets/icon-star-empty.svg)](##)|<details><summary>Parts</summary><sub>- 1 Honda Nidec connector<br>- 1 RJ45 cable (7 ft)<br>- 1 comma 3X<br>- 1 comma power v2<br>- 1 harness box<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=Honda&model=Pilot 2016-22">Buy Here</a></sub></details>||
-|Honda|Ridgeline 2017-25|Honda Sensing|openpilot|26 mph|12 mph|[![star](assets/icon-star-empty.svg)](##)|[![star](assets/icon-star-empty.svg)](##)|<details><summary>Parts</summary><sub>- 1 Honda Nidec connector<br>- 1 RJ45 cable (7 ft)<br>- 1 comma 3X<br>- 1 comma power v2<br>- 1 harness box<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=Honda&model=Ridgeline 2017-25">Buy Here</a></sub></details>||
+|Honda|Ridgeline 2017-24|Honda Sensing|openpilot|26 mph|12 mph|[![star](assets/icon-star-empty.svg)](##)|[![star](assets/icon-star-empty.svg)](##)|<details><summary>Parts</summary><sub>- 1 Honda Nidec connector<br>- 1 RJ45 cable (7 ft)<br>- 1 comma 3X<br>- 1 comma power v2<br>- 1 harness box<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=Honda&model=Ridgeline 2017-24">Buy Here</a></sub></details>||
 |Hyundai|Azera 2022|All|openpilot available[<sup>1</sup>](#footnotes)|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 Hyundai K connector<br>- 1 RJ45 cable (7 ft)<br>- 1 comma 3X<br>- 1 comma power v2<br>- 1 harness box<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=Hyundai&model=Azera 2022">Buy Here</a></sub></details>||
 |Hyundai|Azera Hybrid 2019|All|openpilot available[<sup>1</sup>](#footnotes)|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 Hyundai C connector<br>- 1 RJ45 cable (7 ft)<br>- 1 comma 3X<br>- 1 comma power v2<br>- 1 harness box<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=Hyundai&model=Azera Hybrid 2019">Buy Here</a></sub></details>||
 |Hyundai|Azera Hybrid 2020|All|openpilot available[<sup>1</sup>](#footnotes)|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 Hyundai K connector<br>- 1 RJ45 cable (7 ft)<br>- 1 comma 3X<br>- 1 comma power v2<br>- 1 harness box<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=Hyundai&model=Azera Hybrid 2020">Buy Here</a></sub></details>||
@@ -92,7 +91,7 @@ A supported vehicle is one that just works when you install a comma device. All
 |Hyundai|Elantra Hybrid 2021-23|Smart Cruise Control (SCC)|openpilot available[<sup>1</sup>](#footnotes)|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 Hyundai K connector<br>- 1 RJ45 cable (7 ft)<br>- 1 comma 3X<br>- 1 comma power v2<br>- 1 harness box<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=Hyundai&model=Elantra Hybrid 2021-23">Buy Here</a></sub></details>|<a href="https://youtu.be/_EdYQtV52-c" target="_blank"><img height="18px" src="assets/icon-youtube.svg"></img></a>|
 |Hyundai|Genesis 2015-16|Smart Cruise Control (SCC)|Stock|19 mph|37 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-empty.svg)](##)|<details><summary>Parts</summary><sub>- 1 Hyundai J connector<br>- 1 RJ45 cable (7 ft)<br>- 1 comma 3X<br>- 1 comma power v2<br>- 1 harness box<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=Hyundai&model=Genesis 2015-16">Buy Here</a></sub></details>||
 |Hyundai|i30 2017-19|Smart Cruise Control (SCC)|Stock|0 mph|32 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 Hyundai E connector<br>- 1 RJ45 cable (7 ft)<br>- 1 comma 3X<br>- 1 comma power v2<br>- 1 harness box<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=Hyundai&model=i30 2017-19">Buy Here</a></sub></details>||
-|Hyundai|Ioniq 5 (Southeast Asia and Europe only) 2022-24[<sup>5</sup>](#footnotes)|All|openpilot available[<sup>1</sup>](#footnotes)|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 Hyundai Q connector<br>- 1 RJ45 cable (7 ft)<br>- 1 comma 3X<br>- 1 comma power v2<br>- 1 harness box<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=Hyundai&model=Ioniq 5 (Southeast Asia and Europe only) 2022-24">Buy Here</a></sub></details>||
+|Hyundai|Ioniq 5 (Non-US only) 2022-24[<sup>5</sup>](#footnotes)|All|openpilot available[<sup>1</sup>](#footnotes)|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 Hyundai Q connector<br>- 1 RJ45 cable (7 ft)<br>- 1 comma 3X<br>- 1 comma power v2<br>- 1 harness box<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=Hyundai&model=Ioniq 5 (Non-US only) 2022-24">Buy Here</a></sub></details>||
 |Hyundai|Ioniq 5 (with HDA II) 2022-24[<sup>5</sup>](#footnotes)|Highway Driving Assist II|openpilot available[<sup>1</sup>](#footnotes)|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 Hyundai Q connector<br>- 1 RJ45 cable (7 ft)<br>- 1 comma 3X<br>- 1 comma power v2<br>- 1 harness box<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=Hyundai&model=Ioniq 5 (with HDA II) 2022-24">Buy Here</a></sub></details>||
 |Hyundai|Ioniq 5 (without HDA II) 2022-24[<sup>5</sup>](#footnotes)|Highway Driving Assist|openpilot available[<sup>1</sup>](#footnotes)|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 Hyundai K connector<br>- 1 RJ45 cable (7 ft)<br>- 1 comma 3X<br>- 1 comma power v2<br>- 1 harness box<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=Hyundai&model=Ioniq 5 (without HDA II) 2022-24">Buy Here</a></sub></details>||
 |Hyundai|Ioniq 6 (with HDA II) 2023-24[<sup>5</sup>](#footnotes)|Highway Driving Assist II|Stock|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 Hyundai P connector<br>- 1 RJ45 cable (7 ft)<br>- 1 comma 3X<br>- 1 comma power v2<br>- 1 harness box<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=Hyundai&model=Ioniq 6 (with HDA II) 2023-24">Buy Here</a></sub></details>||
@@ -295,7 +294,7 @@ A supported vehicle is one that just works when you install a comma device. All
 |Volkswagen|Teramont 2018-22|Adaptive Cruise Control (ACC) & Lane Assist|openpilot available[<sup>1,12</sup>](#footnotes)|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 USB-C coupler<br>- 1 VW J533 connector<br>- 1 comma 3X<br>- 1 harness box<br>- 1 long OBD-C cable<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=Volkswagen&model=Teramont 2018-22">Buy Here</a></sub></details>||
 |Volkswagen|Teramont Cross Sport 2021-22|Adaptive Cruise Control (ACC) & Lane Assist|openpilot available[<sup>1,12</sup>](#footnotes)|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 USB-C coupler<br>- 1 VW J533 connector<br>- 1 comma 3X<br>- 1 harness box<br>- 1 long OBD-C cable<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=Volkswagen&model=Teramont Cross Sport 2021-22">Buy Here</a></sub></details>||
 |Volkswagen|Teramont X 2021-22|Adaptive Cruise Control (ACC) & Lane Assist|openpilot available[<sup>1,12</sup>](#footnotes)|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 USB-C coupler<br>- 1 VW J533 connector<br>- 1 comma 3X<br>- 1 harness box<br>- 1 long OBD-C cable<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=Volkswagen&model=Teramont X 2021-22">Buy Here</a></sub></details>||
-|Volkswagen|Tiguan 2018-23|Adaptive Cruise Control (ACC) & Lane Assist|openpilot available[<sup>1,12</sup>](#footnotes)|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 USB-C coupler<br>- 1 VW J533 connector<br>- 1 comma 3X<br>- 1 harness box<br>- 1 long OBD-C cable<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=Volkswagen&model=Tiguan 2018-23">Buy Here</a></sub></details>||
+|Volkswagen|Tiguan 2018-24|Adaptive Cruise Control (ACC) & Lane Assist|openpilot available[<sup>1,12</sup>](#footnotes)|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 USB-C coupler<br>- 1 VW J533 connector<br>- 1 comma 3X<br>- 1 harness box<br>- 1 long OBD-C cable<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=Volkswagen&model=Tiguan 2018-24">Buy Here</a></sub></details>||
 |Volkswagen|Tiguan eHybrid 2021-23|Adaptive Cruise Control (ACC) & Lane Assist|openpilot available[<sup>1,12</sup>](#footnotes)|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 USB-C coupler<br>- 1 VW J533 connector<br>- 1 comma 3X<br>- 1 harness box<br>- 1 long OBD-C cable<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=Volkswagen&model=Tiguan eHybrid 2021-23">Buy Here</a></sub></details>||
 |Volkswagen|Touran 2016-23|Adaptive Cruise Control (ACC) & Lane Assist|openpilot available[<sup>1,12</sup>](#footnotes)|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 USB-C coupler<br>- 1 VW J533 connector<br>- 1 comma 3X<br>- 1 harness box<br>- 1 long OBD-C cable<br>- 1 mount<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=Volkswagen&model=Touran 2016-23">Buy Here</a></sub></details>||

@@ -103,6 +103,7 @@ dev = [
  "lru-dict",
  "matplotlib",
  "parameterized >=0.8, <0.9",
+  #"pprofile",
  "pyautogui",
  "pyopencl; platform_machine != 'aarch64'", # broken on arm64
  "pytools < 2024.1.11; platform_machine != 'aarch64'", # pyopencl use a broken version
@@ -111,10 +112,13 @@ dev = [
  "tabulate",
  "types-requests",
  "types-tabulate",
+
+  # this is only pinned since 5.15.11 is broken
+  "pyqt5 ==5.15.2; platform_machine == 'x86_64'", # no aarch64 wheels for macOS/linux
 ]

 tools = [
-  "metadrive-simulator @ https://github.com/commaai/metadrive/releases/download/MetaDrive-minimal-0.4.2.4/metadrive_simulator-0.4.2.4-py3-none-any.whl ; (platform_machine != 'aarch64')",
+  "metadrive-simulator @ https://github.com/commaai/metadrive/releases/download/MetaDrive-minimal/metadrive_simulator-0.4.2.3-py3-none-any.whl ; (platform_machine != 'aarch64')",
  "rerun-sdk >= 0.18",
 ]

@@ -133,7 +137,7 @@ allow-direct-references = true

 [tool.pytest.ini_options]
 minversion = "6.0"
-addopts = "--ignore=openpilot/ --ignore=opendbc/ --ignore=panda/ --ignore=rednose_repo/ --ignore=tinygrad_repo/ --ignore=teleoprtc_repo/ --ignore=msgq/  -Werror --strict-config --strict-markers --durations=10 -n auto --dist=loadgroup"
+addopts = "--ignore=openpilot/ --ignore=opendbc/ --ignore=panda/ --ignore=rednose_repo/ --ignore=tinygrad_repo/ --ignore=teleoprtc_repo/ --ignore=msgq/ --ignore=sunnypilot/tinygrad_repo/  -Werror --strict-config --strict-markers --durations=10 -n auto --dist=loadgroup"
 cpp_files = "test_*"
 cpp_harness = "selfdrive/test/cpp_harness.py"
 python_files = "test_*.py"
@@ -0,0 +1,39 @@
+#!/usr/bin/env bash
+set -ex
+
+DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" >/dev/null && pwd)"
+cd $DIR
+
+# git clone --mirror
+SRC=/tmp/openpilot.git/
+OUT=/tmp/smallpilot/
+
+echo "starting size $(du -hs .git/)"
+
+rm -rf $OUT
+
+cd $SRC
+git remote update
+
+# copy contents
+#rsync -a --exclude='.git/' $DIR $OUT
+
+cp -r $SRC $OUT
+
+cd $OUT
+
+# remove all tags
+git tag -l | xargs git tag -d
+
+# remove non-master branches
+BRANCHES="release2 release3 devel master-ci nightly"
+for branch in $BRANCHES; do
+  git branch -D $branch
+  git branch -D ${branch}-staging || true
+done
+
+#git gc
+git reflog expire --expire=now --all
+git gc --prune=now
+git gc --aggressive --prune=now
+echo "new one is $(du -hs .)"
@@ -0,0 +1,54 @@
+#!/usr/bin/env python3
+import os
+import ast
+import stat
+import subprocess
+
+fouts = {x.decode('utf-8') for x in subprocess.check_output(['git', 'ls-files']).strip().split()}
+
+pyf = []
+for d in ["cereal", "common", "scripts", "selfdrive", "tools"]:
+  for root, _, files in os.walk(d):
+    for f in files:
+      if f.endswith(".py"):
+        pyf.append(os.path.join(root, f))
+
+imps: set[str] = set()
+
+class Analyzer(ast.NodeVisitor):
+  def visit_Import(self, node):
+    for alias in node.names:
+      imps.add(alias.name)
+    self.generic_visit(node)
+
+  def visit_ImportFrom(self, node):
+    imps.add(node.module)
+    self.generic_visit(node)
+
+tlns = 0
+carlns = 0
+scriptlns = 0
+testlns = 0
+for f in sorted(pyf):
+  if f not in fouts:
+    continue
+  xbit = bool(os.stat(f)[stat.ST_MODE] & stat.S_IXUSR)
+  src = open(f).read()
+  lns = len(src.split("\n"))
+  tree = ast.parse(src)
+  Analyzer().visit(tree)
+  print(f"{lns:5d} {f} {xbit}")
+  if 'test' in f:
+    testlns += lns
+  elif f.startswith(('tools/', 'scripts/', 'selfdrive/debug')):
+    scriptlns += lns
+  elif f.startswith('selfdrive/car'):
+    carlns += lns
+  else:
+    tlns += lns
+
+print(f"{tlns} lines of openpilot python")
+print(f"{carlns} lines of car ports")
+print(f"{scriptlns} lines of tools/scripts/debug")
+print(f"{testlns} lines of tests")
+#print(sorted(list(imps)))
@@ -0,0 +1,11 @@
+#!/usr/bin/env python3
+from collections import Counter
+from pprint import pprint
+
+from opendbc.car.docs import get_all_car_docs
+
+if __name__ == "__main__":
+  cars = get_all_car_docs()
+  make_count = Counter(l.make for l in cars)
+  print("\n", "*" * 20, len(cars), "total", "*" * 20, "\n")
+  pprint(make_count)
@@ -0,0 +1,391 @@
+#!/usr/bin/env bash
+set -e
+
+SRC=/tmp/openpilot/
+SRC_CLONE=/tmp/openpilot-clone/
+OUT=/tmp/openpilot-tiny/
+
+REWRITE_IGNORE_BRANCHES=(
+  dashcam3
+  devel
+  master-ci
+  nightly
+  release2
+  release3
+  release3-staging
+)
+
+VALIDATE_IGNORE_FILES=(
+  ".github/ISSUE_TEMPLATE/bug_report.md"
+  ".github/pull_request_template.md"
+)
+
+DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" >/dev/null && pwd)"
+cd $DIR
+
+LOGS_DIR=$DIR/git-rewrite-$(date +"%Y-%m-%dT%H:%M:%S%z")
+mkdir -p $LOGS_DIR
+
+GIT_REWRITE_LOG=$LOGS_DIR/git-rewrite-log.txt
+BRANCH_DIFF_LOG=$LOGS_DIR/branch-diff-log.txt
+COMMIT_DIFF_LOG=$LOGS_DIR/commit-diff-log.txt
+
+START_TIME=$(date +%s)
+exec > >(while IFS= read -r line; do
+  CURRENT_TIME=$(date +%s)
+  ELAPSED_TIME=$((CURRENT_TIME - START_TIME))
+  echo "[${ELAPSED_TIME}s] $line"
+done | tee -a "$GIT_REWRITE_LOG") 2>&1
+
+# INSTALL git-filter-repo
+if [ ! -f /tmp/git-filter-repo ]; then
+  echo "Installing git-filter-repo..."
+  curl -sSo /tmp/git-filter-repo https://raw.githubusercontent.com/newren/git-filter-repo/main/git-filter-repo
+  chmod +x /tmp/git-filter-repo
+fi
+
+# MIRROR openpilot
+if [ ! -d $SRC ]; then
+  echo "Mirroring openpilot..."
+  git clone --mirror https://github.com/commaai/openpilot.git $SRC # 4.18 GiB (488034 objects)
+
+  cd $SRC
+
+  echo "Starting size $(du -sh .)"
+
+  git remote update
+
+  # the git-filter-repo analysis is bliss - can be found in the repo root/filter-repo/analysis
+  echo "Analyzing with git-filter-repo..."
+  /tmp/git-filter-repo --force --analyze
+
+  echo "Pushing to openpilot-archive..."
+  # push to archive repo - in smaller parts because the 2 GB push limit - https://docs.github.com/en/get-started/using-git/troubleshooting-the-2-gb-push-limit
+  ARCHIVE_REPO=git@github.com:commaai/openpilot-archive.git
+  git push --prune $ARCHIVE_REPO +refs/heads/master:refs/heads/master # push master first so it's the default branch (when openpilot-archive is an empty repo)
+  git push --prune $ARCHIVE_REPO +refs/heads/*:refs/heads/* # 956.39 MiB (110725 objects)
+  git push --prune $ARCHIVE_REPO +refs/tags/*:refs/tags/* # 1.75 GiB (21694 objects)
+  # git push --mirror $ARCHIVE_REPO || true # fails to push refs/pull/* (deny updating a hidden ref) for pull requests
+  # we fail and continue - more reading: https://stackoverflow.com/a/34266401/639708 and https://blog.plataformatec.com.br/2013/05/how-to-properly-mirror-a-git-repository/
+fi
+
+# REWRITE master and tags
+if [ ! -d $SRC_CLONE ]; then
+  echo "Cloning $SRC..."
+  GIT_LFS_SKIP_SMUDGE=1 git clone $SRC $SRC_CLONE
+
+  cd $SRC_CLONE
+
+  echo "Checking out old history..."
+
+  git checkout tags/v0.7.1 > /dev/null 2>&1
+  # checkout as main, since we need master ref later
+  git checkout -b main
+
+  echo "Creating setup commits..."
+
+  # rm these so we don't get conflicts later
+  git rm -r cereal opendbc panda selfdrive/ui/ui > /dev/null
+  git commit -m "removed conflicting files" > /dev/null
+
+  # skip-smudge to get rid of some lfs errors that it can't find the reference of some lfs files
+  # we don't care about fetching/pushing lfs right now
+  git lfs install --skip-smudge --local
+
+  # squash initial setup commits
+  git cherry-pick -n -X theirs 6c33a5c..59b3d06 > /dev/null
+  git commit -m "switching to master" > /dev/null
+
+  # squash the two commits
+  git reset --soft HEAD~2
+  git commit -m "switching to master" -m "$(git log --reverse --format=%B 6c33a5c..59b3d06)" -m "removed conflicting files" > /dev/null
+
+  # get commits we want to cherry-pick
+  # will start with the next commit after #59b3d06 tools is local now
+  COMMITS=$(git rev-list --reverse 59b3d06..master)
+
+  # we need this for logging
+  TOTAL_COMMITS=$(echo $COMMITS | wc -w | xargs)
+  CURRENT_COMMIT_NUMBER=0
+
+  # empty this file
+  > commit-map.txt
+
+  echo "Rewriting master commits..."
+
+  for COMMIT in $COMMITS; do
+      CURRENT_COMMIT_NUMBER=$((CURRENT_COMMIT_NUMBER + 1))
+      # echo -ne "[$CURRENT_COMMIT_NUMBER/$TOTAL_COMMITS] Cherry-picking commit: $COMMIT"\\r
+      echo "[$CURRENT_COMMIT_NUMBER/$TOTAL_COMMITS] Cherry-picking commit: $COMMIT"
+
+      # set environment variables to preserve author/committer and dates
+      export GIT_AUTHOR_NAME=$(git show -s --format='%an' $COMMIT)
+      export GIT_AUTHOR_EMAIL=$(git show -s --format='%ae' $COMMIT)
+      export GIT_COMMITTER_NAME=$(git show -s --format='%cn' $COMMIT)
+      export GIT_COMMITTER_EMAIL=$(git show -s --format='%ce' $COMMIT)
+      export GIT_AUTHOR_DATE=$(git show -s --format='%ad' $COMMIT)
+      export GIT_COMMITTER_DATE=$(git show -s --format='%cd' $COMMIT)
+
+      # cherry-pick the commit
+      if ! GIT_OUTPUT=$(git cherry-pick -m 1 -X theirs $COMMIT 2>&1); then
+        # check if the failure is because of an empty commit
+        if [[ "$GIT_OUTPUT" == *"The previous cherry-pick is now empty"* ]]; then
+          echo "Empty commit detected. Skipping commit $COMMIT"
+          git cherry-pick --skip
+          # log it was empty to the mapping file
+          echo "$COMMIT EMPTY" >> commit-map.txt
+        else
+          # handle other errors or conflicts
+          echo "Cherry-pick failed. Handling error..."
+          echo "$GIT_OUTPUT"
+          exit 1
+        fi
+      else
+        # capture the new commit hash
+        NEW_COMMIT=$(git rev-parse HEAD)
+
+        # save the old and new commit hashes to the mapping file
+        echo "$COMMIT $NEW_COMMIT" >> commit-map.txt
+
+        # append the old commit ID to the commit message
+        git commit --amend -m "$(git log -1 --pretty=%B)" -m "Former-commit-id: $COMMIT" > /dev/null
+      fi
+
+    # prune every 3000 commits to avoid gc errors
+    if [ $((CURRENT_COMMIT_NUMBER % 3000)) -eq 0 ]; then
+      echo "Pruning repo..."
+      git gc
+    fi
+  done
+
+  echo "Rewriting tags..."
+
+  # remove all old tags
+  git tag -l | xargs git tag -d
+
+  # read each line from the tag-commit-map.txt
+  while IFS=' ' read -r TAG OLD_COMMIT; do
+    # search for the new commit in commit-map.txt corresponding to the old commit
+    NEW_COMMIT=$(grep "^$OLD_COMMIT " "commit-map.txt" | awk '{print $2}')
+
+    # check if this is a rebased commit
+    if [ -z "$NEW_COMMIT" ]; then
+      # if not, then just use old commit hash
+      NEW_COMMIT=$OLD_COMMIT
+    fi
+
+    echo "Rewriting tag $TAG from commit $NEW_COMMIT"
+    git tag -f "$TAG" "$NEW_COMMIT"
+  done < "$DIR/tag-commit-map.txt"
+
+  # uninstall lfs since we don't want to touch (push to) lfs right now
+  # git push will also push lfs, if we don't uninstall (--local so just for this repo)
+  git lfs uninstall --local
+
+  # force push new master
+  git push --force origin main:master
+
+  # force push new tags
+  git push --force --tags
+fi
+
+# REWRITE branches based on master
+if [ ! -f "$SRC_CLONE/rewrite-branches-done" ]; then
+  cd $SRC_CLONE
+  > rewrite-branches-done
+
+  # empty file
+  > $BRANCH_DIFF_LOG
+
+  echo "Rewriting branches based on master..."
+
+  # will store raw diffs here, if exist
+  mkdir -p differences
+
+  # get a list of all branches except master and REWRITE_IGNORE_BRANCHES
+  BRANCHES=$(git branch -r | grep -v ' -> ' | sed 's/.*origin\///' | grep -v '^master$' | grep -v -f <(echo "${REWRITE_IGNORE_BRANCHES[*]}" | tr ' ' '\n'))
+
+  for BRANCH in $BRANCHES; do
+    # check if the branch is based on master history
+    MERGE_BASE=$(git merge-base master origin/$BRANCH) || true
+    if [ -n "$MERGE_BASE" ]; then
+      echo "Rewriting branch: $BRANCH"
+
+      # create a new branch based on the new master
+      NEW_MERGE_BASE=$(grep "^$MERGE_BASE " "commit-map.txt" | awk '{print $2}')
+      if [ -z "$NEW_MERGE_BASE" ]; then
+        echo "Error: could not find new merge base for branch $BRANCH" >> $BRANCH_DIFF_LOG
+        continue
+      fi
+      git checkout -b ${BRANCH}_new $NEW_MERGE_BASE
+
+      # get the range of commits unique to this branch
+      COMMITS=$(git rev-list --reverse $MERGE_BASE..origin/${BRANCH})
+
+      HAS_ERROR=0
+
+      # simple delimiter
+      echo "BRANCH ${BRANCH}" >> commit-map.txt
+
+      for COMMIT in $COMMITS; do
+        # set environment variables to preserve author/committer and dates
+        export GIT_AUTHOR_NAME=$(git show -s --format='%an' $COMMIT)
+        export GIT_AUTHOR_EMAIL=$(git show -s --format='%ae' $COMMIT)
+        export GIT_COMMITTER_NAME=$(git show -s --format='%cn' $COMMIT)
+        export GIT_COMMITTER_EMAIL=$(git show -s --format='%ce' $COMMIT)
+        export GIT_AUTHOR_DATE=$(git show -s --format='%ad' $COMMIT)
+        export GIT_COMMITTER_DATE=$(git show -s --format='%cd' $COMMIT)
+
+        # cherry-pick the commit
+        if ! GIT_OUTPUT=$(git cherry-pick -m 1 -X theirs $COMMIT 2>&1); then
+          # check if the failure is because of an empty commit
+          if [[ "$GIT_OUTPUT" == *"The previous cherry-pick is now empty"* ]]; then
+            echo "Empty commit detected. Skipping commit $COMMIT"
+            git cherry-pick --skip
+            # log it was empty to the mapping file
+            echo "$COMMIT EMPTY" >> commit-map.txt
+          else
+            # handle other errors or conflicts
+            echo "Cherry-pick of ${BRANCH} branch failed. Removing branch upstream..." >> $BRANCH_DIFF_LOG
+            echo "$GIT_OUTPUT" > "$LOGS_DIR/branch-${BRANCH}"
+            git cherry-pick --abort
+            git push --delete origin ${BRANCH}
+            HAS_ERROR=1
+            break
+          fi
+        else
+          # capture the new commit hash
+          NEW_COMMIT=$(git rev-parse HEAD)
+
+          # save the old and new commit hashes to the mapping file
+          echo "$COMMIT $NEW_COMMIT" >> commit-map.txt
+
+          # append the old commit ID to the commit message
+          git commit --amend -m "$(git log -1 --pretty=%B)" -m "Former-commit-id: $COMMIT" > /dev/null
+        fi
+      done
+
+      # force push the new branch
+      if [ $HAS_ERROR -eq 0 ]; then
+        # git lfs goes haywire here, so we need to install and uninstall
+        # git lfs install --skip-smudge --local
+        git lfs uninstall --local > /dev/null
+        git push -f origin ${BRANCH}_new:${BRANCH}
+      fi
+
+      # clean up local branch
+      git checkout master > /dev/null
+      git branch -D ${BRANCH}_new > /dev/null
+    else
+      echo "Deleting branch $BRANCH as it's not based on master history" >> $BRANCH_DIFF_LOG
+      git push --delete origin ${BRANCH}
+    fi
+  done
+fi
+
+# VALIDATE cherry-pick
+if [ ! -f "$SRC_CLONE/validation-done" ]; then
+  cd $SRC_CLONE
+  > validation-done
+
+  TOTAL_COMMITS=$(grep -cve '^\s*$' commit-map.txt)
+  CURRENT_COMMIT_NUMBER=0
+  COUNT_SAME=0
+  COUNT_DIFF=0
+
+  # empty file
+  > $COMMIT_DIFF_LOG
+
+  echo "Validating commits..."
+
+  # will store raw diffs here, if exist
+  mkdir -p differences
+
+  # read each line from commit-map.txt
+  while IFS=' ' read -r OLD_COMMIT NEW_COMMIT; do
+    if [ "$NEW_COMMIT" == "EMPTY" ]; then
+      continue
+    fi
+    if [ "$OLD_COMMIT" == "BRANCH" ]; then
+      echo "Branch ${NEW_COMMIT} below:" >> $COMMIT_DIFF_LOG
+      continue
+    fi
+    CURRENT_COMMIT_NUMBER=$((CURRENT_COMMIT_NUMBER + 1))
+    # retrieve short hashes and dates for the old and new commits
+    OLD_COMMIT_SHORT=$(git rev-parse --short $OLD_COMMIT)
+    NEW_COMMIT_SHORT=$(git rev-parse --short $NEW_COMMIT)
+    OLD_DATE=$(git show -s --format='%cd' $OLD_COMMIT)
+    NEW_DATE=$(git show -s --format='%cd' $NEW_COMMIT)
+
+    # echo -ne "[$CURRENT_COMMIT_NUMBER/$TOTAL_COMMITS] Comparing old commit $OLD_COMMIT_SHORT ($OLD_DATE) with new commit $NEW_COMMIT_SHORT ($NEW_DATE)"\\r
+    echo "[$CURRENT_COMMIT_NUMBER/$TOTAL_COMMITS] Comparing old commit $OLD_COMMIT_SHORT ($OLD_DATE) with new commit $NEW_COMMIT_SHORT ($NEW_DATE)"
+
+    # generate lists of files and their hashes for the old and new commits, excluding ignored files
+    OLD_FILES=$(git ls-tree -r $OLD_COMMIT | grep -vE "$(IFS='|'; echo "${VALIDATE_IGNORE_FILES[*]}")")
+    NEW_FILES=$(git ls-tree -r $NEW_COMMIT | grep -vE "$(IFS='|'; echo "${VALIDATE_IGNORE_FILES[*]}")")
+
+    # Compare the diffs
+    if diff <(echo "$OLD_FILES") <(echo "$NEW_FILES") > /dev/null; then
+      # echo "Old commit $OLD_COMMIT_SHORT and new commit $NEW_COMMIT_SHORT are equivalent."
+      COUNT_SAME=$((COUNT_SAME + 1))
+    else
+      echo "[$CURRENT_COMMIT_NUMBER/$TOTAL_COMMITS] Difference found between old commit $OLD_COMMIT_SHORT and new commit $NEW_COMMIT_SHORT" >> $COMMIT_DIFF_LOG
+      COUNT_DIFF=$((COUNT_DIFF + 1))
+      set +e
+      diff -u <(echo "$OLD_FILES") <(echo "$NEW_FILES") > "$LOGS_DIR/commit-$CURRENT_COMMIT_NUMBER-$OLD_COMMIT_SHORT-$NEW_COMMIT_SHORT"
+      set -e
+    fi
+  done < "commit-map.txt"
+
+  echo "Summary:" >> $COMMIT_DIFF_LOG
+  echo "Equivalent commits: $COUNT_SAME" >> $COMMIT_DIFF_LOG
+  echo "Different commits: $COUNT_DIFF" >> $COMMIT_DIFF_LOG
+fi
+
+if [ ! -d $OUT ]; then
+  cp -r $SRC $OUT
+
+  cd $OUT
+
+  # remove all non-master branches
+  # git branch | grep -v "^  master$" | grep -v "\*" | xargs git branch -D
+
+  # echo "cleaning up refs"
+  # delete pull request refs since we can't alter them anyway (https://docs.github.com/en/pull-requests/collaborating-with-pull-requests/reviewing-changes-in-pull-requests/checking-out-pull-requests-locally#error-failed-to-push-some-refs)
+  # git for-each-ref --format='%(refname)' | grep '^refs/pull/' | xargs -I {} git update-ref -d {}
+
+  echo "importing new lfs files"
+  # import "almost" everything to lfs
+  BRANCHES=$(git for-each-ref --format='%(refname)' refs/heads/ | sed 's%refs/heads/%%g' | grep -v -f <(echo "${REWRITE_IGNORE_BRANCHES[*]}" | tr ' ' '\n') | tr '\n' ' ')
+  git lfs migrate import --include="*.dlc,*.onnx,*.svg,*.png,*.gif,*.ttf,*.wav,selfdrive/car/tests/test_models_segs.txt,system/hardware/tici/updater,selfdrive/ui/qt/spinner_larch64,selfdrive/ui/qt/text_larch64,third_party/**/*.a,third_party/**/*.so,third_party/**/*.so.*,third_party/**/*.dylib,third_party/acados/*/t_renderer,third_party/qt5/larch64/bin/lrelease,third_party/qt5/larch64/bin/lupdate,third_party/catch2/include/catch2/catch.hpp,*.apk,*.apkpatch,*.jar,*.pdf,*.jpg,*.mp3,*.thneed,*.tar.gz,*.npy,*.csv,*.a,*.so*,*.dylib,*.o,*.b64,selfdrive/hardware/tici/updater,selfdrive/boardd/tests/test_boardd,selfdrive/ui/qt/spinner_aarch64,installer/updater/updater,selfdrive/debug/profiling/simpleperf/**/*,selfdrive/hardware/eon/updater,selfdrive/ui/qt/text_aarch64,selfdrive/debug/profiling/pyflame/**/*,installer/installers/installer_openpilot,installer/installers/installer_dashcam,selfdrive/ui/text/text,selfdrive/ui/android/text/text,selfdrive/ui/spinner/spinner,selfdrive/visiond/visiond,selfdrive/loggerd/loggerd,selfdrive/sensord/sensord,selfdrive/sensord/gpsd,selfdrive/ui/android/spinner/spinner,selfdrive/ui/qt/spinner,selfdrive/ui/qt/text,_stringdefs.py,dfu-util-aarch64-linux,dfu-util-aarch64,dfu-util-x86_64-linux,dfu-util-x86_64,stb_image.h,clpeak3,clwaste,apk/**/*,external/**/*,phonelibs/**/*,third_party/boringssl/**/*,flask/**/*,panda/**/*,board/**/*,messaging/**/*,opendbc/**/*,tools/cabana/chartswidget.cc,third_party/nanovg/**/*,selfdrive/controls/lib/lateral_mpc/lib_mpc_export/**/*,selfdrive/ui/paint.cc,werkzeug/**/*,pyextra/**/*,third_party/android_hardware_libhardware/**/*,selfdrive/controls/lib/lead_mpc_lib/lib_mpc_export/**/*,selfdrive/locationd/laikad.py,selfdrive/locationd/test/test_laikad.py,tools/gpstest/test_laikad.py,selfdrive/locationd/laikad_helpers.py,tools/nui/**/*,jsonrpc/**/*,selfdrive/controls/lib/longitudinal_mpc/lib_mpc_export/**/*,selfdrive/controls/lib/lateral_mpc/mpc_export/**/*,selfdrive/camerad/cameras/camera_qcom.cc,selfdrive/manager.py,selfdrive/modeld/models/driving.cc,third_party/curl/**/*,selfdrive/modeld/thneed/debug/**/*,selfdrive/modeld/thneed/include/**/*,third_party/openmax/**/*,selfdrive/controls/lib/longitudinal_mpc/mpc_export/**/*,selfdrive/controls/lib/longitudinal_mpc_model/lib_mpc_export/**/*,Pipfile,Pipfile.lock,gunicorn/**/*,*.qm,jinja2/**/*,click/**/*,dbcs/**/*,websocket/**/*" $BRANCHES
+
+  echo "reflog and gc"
+  # this is needed after lfs import
+  git reflog expire --expire=now --all
+  git gc --prune=now --aggressive
+
+  # check the git-filter-repo analysis again - can be found in the repo root/filter-repo/analysis
+  echo "Analyzing with git-filter-repo..."
+  /tmp/git-filter-repo --force --analyze
+
+  echo "New size is $(du -sh .)"
+fi
+
+cd $OUT
+
+# fetch all lfs files from https://github.com/commaai/openpilot.git
+# some lfs files are missing on gitlab, but they can be found on github
+git config lfs.url https://github.com/commaai/openpilot.git/info/lfs
+git config lfs.pushurl ssh://git@github.com/commaai/openpilot.git
+git lfs fetch --all || true
+
+# also fetch all lfs files from https://gitlab.com/commaai/openpilot-lfs.git
+git config lfs.url https://gitlab.com/commaai/openpilot-lfs.git/info/lfs
+git config lfs.pushurl ssh://git@gitlab.com/commaai/openpilot-lfs.git
+git lfs fetch --all || true
+
+# final push - will also push lfs
+# TODO: switch to git@github.com:commaai/openpilot.git when ready
+# git push --mirror git@github.com:commaai/openpilot-tiny.git
+# using this instead to ignore refs/pull/* - since this is also what --mirror does - https://blog.plataformatec.com.br/2013/05/how-to-properly-mirror-a-git-repository/
+git push --prune git@github.com:commaai/openpilot-tiny.git +refs/heads/*:refs/heads/* +refs/tags/*:refs/tags/*
@@ -0,0 +1,59 @@
+#!/usr/bin/env bash
+set -e
+
+DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" >/dev/null && pwd)"
+cd $DIR
+
+git clone --bare https://github.com/commaai/openpilot
+cp -r openpilot.git openpilot_backup
+cd openpilot.git
+
+# backup old repo
+git push git@github.com:commaai/openpilot-archive.git +refs/heads/master:refs/heads/master
+git push git@github.com:commaai/openpilot-archive.git +refs/heads/*:refs/heads/*
+git push git@github.com:commaai/openpilot-archive.git +refs/tags/*:refs/tags/*
+git push --mirror git@github.com:commaai/openpilot-archive.git
+
+# ignore all release branches
+git for-each-ref --format='delete %(refname)' | grep 'dashcam3\|devel\|master-ci\|nightly\|release2\|release3\|release3-staging' | git update-ref --stdin
+
+# re-tag old releases on master
+declare -A TAGS=( ["f8cb04e4a8b032b72a909f68b808a50936184bee"]="v0.9.7" ["0b4d08fab8e35a264bc7383e878538f8083c33e5"]="v0.9.6" ["3b1e9017c560499786d8a0e46aaaeea65037acac"]="v0.9.5" ["fa310d9e2542cf497d92f007baec8fd751ffa99c"]="v0.9.4" ["8704c1ff952b5c85a44f50143bbd1a4f7b4887e2"]="v0.9.3" ["c7d3b28b93faa6c955fb24bc64031512ee985ee9"]="v0.9.2" ["89f68bf0cbf53a81b0553d3816fdbe522f941fa1"]="v0.9.1" ["58b84fb401a804967aa0dd5ee66fafa90194fd30"]="v0.9.0" ["f41dc62a12cc0f3cb8c5453c0caa0ba21e1bd01e"]="v0.8.16" ["5a7c2f90361e72e9c35e88abd2e11acdc4aba354"]="v0.8.15" ["71901c94dbbaa2f9f156a80c14cc7ea65219fc7c"]="v0.8.14" ["95da47079510afc91665263619e5939126da637c"]="v0.8.13" ["472177e2a8a1d002e56f9096326fd2dff62e54f9"]="v0.8.12" ["08078acbd0b4f7da469c7dff6159000e358974a9"]="v0.8.11" ["687925c775c375495f9827946138a724bde00b9d"]="v0.8.10" ["204e5a090735a059d69c29145a4cee49450da07e"]="v0.8.9" ["4be956f8861ecbb521ef9503a3c87b07c9d36721"]="v0.8.8" ["589f82c76627d634761a31a34b2488403556eb0b"]="v0.8.7" ["507cfc8910f74ddb8810039d68b880b426ff9ff9"]="v0.8.6" ["d47b00b45a866bef088f51d1ff31de5885ab04e9"]="v0.8.5" ["553e7d1cce314e7eb0587186b1764c3ff43bed62"]="v0.8.4" ["9896438d1511602a1ff87f7c4eb3c7172b30104a"]="v0.8.3" ["280192ed1443f112463417c2d815ea8ee2762fbd"]="v0.8.2" ["8039361567e4659eae2a084e6f39f34acadf4cac"]="v0.8.1" ["d56e04c0d960c8d3d4ab88b578dc508a2b4e07dc"]="v0.8" ["3d456e5d0fbf0c9887d0499dee812f2b029edf6d"]="v0.7.10" ["81763a18b5d0e379b749e090ecce36a91fca7c43"]="v0.7.9" ["9bc0b350fd273bbb2deb3dcaef0312944e4f6cfd"]="v0.7.8" ["ede5b632b58c55e4ff003f948efae07fe03c2280"]="v0.7.7" ["775acd11ba2e0a8c2f5a5655338718d796491b36"]="v0.7.6.1" ["302417b4cf0dcf00d45e4995b5410e543ad121d1"]="v0.7.5" ["12ff088b42221dd17d9d97decb1fc61a7cb0a861"]="v0.7.4" ["9563f7730252451fdcba9bc3d9fe36dab9c86a26"]="v0.7.3" ["8321cf283abbc2ca3fda7e0c7a069a77a492fe0c"]="v0.7.2" ["1e1de64a1e59476b7b3d3558b92149246d5c3292"]="v0.7.1" ["a2ae18d1dbd1e59c38ce22fa25ddffbd1d3084e3"]="v0.7" ["d4eb5a6eafdd4803d09e6f3963918216cca5a81f"]="v0.6.6" ["70d17cd69b80e7627dcad8fd5b6438f2309ac307"]="v0.6.5" ["58f376002e0c654fbc2de127765fa297cf694a33"]="v0.6.4" ["d5f9caa82d80cdcc7f1b7748f2cf3ccbf94f82a3"]="v0.6.3" ["095ef5f9f60fca1b269aabcc3cfd322b17b9e674"]="v0.6.2" ["cf5c4aeacb1703d0ffd35bdb5297d3494fee9a22"]="v0.6.1" ["60a20537c5f3fcc7f11946d81aebc8f90c08c117"]="v0.6" ["dd34ccfe288ebda8e2568cf550994ae890379f45"]="v0.5.13" ["3f9059fea886f1fa3b0c19a62a981d891dcc84eb"]="v0.5.12" ["2f92d577f995ff6ae1945ef6b89df3cb69b92999"]="v0.5.11" ["5a9d89ed42ddcd209d001a10d7eb828ef0e6d9de"]="v0.5.10" ["0207a970400ee28d3e366f2e8f5c551281accf02"]="v0.5.9" ["b967da5fc1f7a07e3561db072dd714d325e857b0"]="v0.5.8" ["210db686bb89f8696aa040e6e16de65424b808c9"]="v0.5.7" ["860a48765d1016ba226fb2c64aea35a45fe40e4a"]="v0.5.6" ["8f3539a27b28851153454eb737da9624cccaed2d"]="v0.5.5" ["a422246dc30bce11e970514f13f7c110f4470cc3"]="v0.5.4" ["285c52eb693265a0a530543e9ca0aeb593a2a55e"]="v0.5.3" ["0129a8a4ff8da5314e8e4d4d3336e89667ff6d54"]="v0.5.2" ["6f3d10a4c475c4c4509f0b370805419acd13912d"]="v0.5.1" ["de33bc46452b1046387ee2b3a03191b2c71135fb"]="v0.5" ["ae5cb7a0dab8b1bed9d52292f9b4e8e66a0f8ec9"]="v0.4.7" ["c6df34f55ba8c5a911b60d3f9eb20e3fa45f68c1"]="v0.4.6" ["37285038d3f91fa1b49159c4a35a8383168e644f"]="v0.4.5" ["9a9ff839a9b70cb2601d7696af743f5652395389"]="v0.4.4" ["28c0797d30175043bbfa31307b63aab4197cf996"]="v0.4.2" ["4474b9b3718653aeb0aee26422caefb90460cc0e"]="v0.4.1" ["da52d065a4c4f52d6017a537f3a80326f5af8bdc"]="v0.4.0.2" ["9d3963559ae7b15193057937ff3e72481899f40d"]="v0.3.5"  ["1b8c44b5067525a5d266b6e99799d8097da76a29"]="v0.3.4" ["5cf91d0496688fed4f2a6c7021349b1fc0e057a2"]="v0.3.3" ["7fe46f1e1df5dec08a940451ba0feefd5c039165"]="v0.3.2" ["41e3a0f699f5c39cb61a15c0eb7a4aa816d47c24"]="v0.3.1" ["c5d8aec28b5230d34ae4b677c2091cc3dec7e3e8"]="v0.3.0" ["693bcb0f83478f2651db6bac9be5ca5ad60d03f3"]="v0.2.9" ["95a349abcc050712c50d4d85a1c8a804eee7f6c2"]="v0.2.8" ["c6ba5dc5391d3ca6cda479bf1923b88ce45509a0"]="v0.2.7" ["6c3afeec0fb439070b2912978b8dbb659033b1d9"]="v0.2.6" ["29c58b45882ac79595356caf98580c1d2a626011"]="v0.2.5" ["ecc565aa3fdc4c7e719aadc000e1fdc4d80d4fe0"]="v0.2.4" ["adaa4ed350acda4067fc0b455ad15b54cdf4c768"]="v0.2.3" ["a64b9aa9b8cb5863c917b6926516291a63c02fe5"]="v0.2.2" ["17d9becd3c673091b22f09aa02559a9ed9230f50"]="v0.2.1" ["449b482cc3236ccf31829830b4f6a44b2dcc06c2"]="v0.2" ["e94a30bec07e719c5a7b037ca1f4db8312702cce"]="v0.1" )
+for tag in "${!TAGS[@]}"; do git tag -f "${TAGS[$tag]}" "$tag" ; done
+
+# get master root commit
+ROOT_COMMIT=$(git rev-list --max-parents=0 HEAD | tail -n 1)
+
+# link master and devel
+git replace --graft $ROOT_COMMIT v0.7.1
+git-filter-repo --prune-empty never --force --commit-callback 'h=commit.original_id.decode("utf-8");m=commit.message.decode("utf-8");commit.message=str.encode(m + "\n" + "old-commit-hash: " + h)'
+
+# delete replace refs
+git for-each-ref --format='delete %(refname)' refs/replace | git update-ref --stdin
+
+# machine validation
+tail -n +2 "filter-repo/commit-map" | tr ' ' '\n' | xargs -P $(nproc) -n 2 bash -c 'H1=$(cd ../openpilot_backup && git ls-tree -r $0 | sha1sum) && H2=$(git ls-tree -r $1 | sha1sum) && echo "$H1 $H2" >> /tmp/GIT_HASHES && diff <(echo $H1) <(echo $H2) || exit 255'
+# human validation
+less /tmp/GIT_HASH
+
+# cleanup
+git reflog expire --expire=now --all
+git gc --prune=now --aggressive
+
+# get all lfs files
+set +e
+git config lfs.url https://github.com/commaai/openpilot.git/info/lfs
+git lfs fetch --all
+git config lfs.url https://gitlab.com/commaai/openpilot-lfs.git/info/lfs
+git lfs fetch --all
+set -e
+
+# add new files to lfs
+git lfs migrate import --everything --include="*.ico,*.dlc,*.onnx,*.svg,*.png,*.gif,*.ttf,*.wav,system/hardware/tici/updater,selfdrive/ui/qt/spinner_larch64,selfdrive/ui/qt/text_larch64,third_party/**/*.a,third_party/**/*.so,third_party/**/*.so.*,third_party/**/*.dylib,third_party/acados/*/t_renderer,third_party/qt5/larch64/bin/lrelease,third_party/qt5/larch64/bin/lupdate,third_party/catch2/include/catch2/catch.hpp,*.apk,*.apkpatch,*.jar,*.pdf,*.jpg,*.mp3,*.thneed,*.tar.gz,*.npy,*.csv,*.a,*.so*,*.dylib,*.o,*.b64,selfdrive/hardware/tici/updater,selfdrive/boardd/tests/test_boardd,selfdrive/ui/qt/spinner_aarch64,installer/updater/updater,selfdrive/debug/profiling/simpleperf/**/*,selfdrive/hardware/eon/updater,selfdrive/ui/qt/text_aarch64,selfdrive/debug/profiling/pyflame/**/*,installer/installers/installer_openpilot,installer/installers/installer_dashcam,selfdrive/ui/text/text,selfdrive/ui/android/text/text,selfdrive/ui/spinner/spinner,selfdrive/visiond/visiond,selfdrive/loggerd/loggerd,selfdrive/sensord/sensord,selfdrive/sensord/gpsd,selfdrive/ui/android/spinner/spinner,selfdrive/ui/qt/spinner,selfdrive/ui/qt/text,_stringdefs.py,dfu-util-aarch64-linux,dfu-util-aarch64,dfu-util-x86_64-linux,dfu-util-x86_64,stb_image.h,clpeak3,clwaste,apk/**/*,external/**/*,phonelibs/**/*,third_party/boringssl/**/*,pyextra/**/*,panda/board/**/inc/*.h,panda/board/obj/*.elf,board/inc/*.h,third_party/nanovg/**/*,selfdrive/controls/lib/lateral_mpc/lib_mpc_export/**/*,pyextra/**/*,third_party/android_hardware_libhardware/**/*,selfdrive/controls/lib/lead_mpc_lib/lib_mpc_export/**/*,*.pro,selfdrive/controls/lib/longitudinal_mpc/lib_mpc_export/**/*,selfdrive/controls/lib/lateral_mpc/mpc_export/**/*,third_party/curl/**/*,selfdrive/modeld/thneed/debug/**/*,selfdrive/modeld/thneed/include/**/*,third_party/openmax/**/*,selfdrive/controls/lib/longitudinal_mpc/mpc_export/**/*,selfdrive/controls/lib/longitudinal_mpc_model/lib_mpc_export/**/*,Pipfile,Pipfile.lock,poetry.lock,*.qm"
+
+# set new lfs endpoint
+git config lfs.url https://gitlab.com/commaai/openpilot-lfs.git/info/lfs
+git config lfs.pushurl ssh://git@gitlab.com/commaai/openpilot-lfs.git
+
+# push all branch+tag (scary stuff...)
+git push -f --set-upstream git@github.com:commaai/openpilot.git +refs/heads/*:refs/heads/* +refs/tags/*:refs/tags/*
@@ -0,0 +1,82 @@
+v0.1 e94a30bec07e719c5a7b037ca1f4db8312702cce
+v0.2 449b482cc3236ccf31829830b4f6a44b2dcc06c2
+v0.2.1 17d9becd3c673091b22f09aa02559a9ed9230f50
+v0.2.2 a64b9aa9b8cb5863c917b6926516291a63c02fe5
+v0.2.3 adaa4ed350acda4067fc0b455ad15b54cdf4c768
+v0.2.4 ecc565aa3fdc4c7e719aadc000e1fdc4d80d4fe0
+v0.2.5 29c58b45882ac79595356caf98580c1d2a626011
+v0.2.6 6c3afeec0fb439070b2912978b8dbb659033b1d9
+v0.2.7 c6ba5dc5391d3ca6cda479bf1923b88ce45509a0
+v0.2.8 95a349abcc050712c50d4d85a1c8a804eee7f6c2
+v0.2.9 693bcb0f83478f2651db6bac9be5ca5ad60d03f3
+v0.3.0 c5d8aec28b5230d34ae4b677c2091cc3dec7e3e8
+v0.3.1 41e3a0f699f5c39cb61a15c0eb7a4aa816d47c24
+v0.3.2 7fe46f1e1df5dec08a940451ba0feefd5c039165
+v0.3.3 5cf91d0496688fed4f2a6c7021349b1fc0e057a2
+v0.3.4 1b8c44b5067525a5d266b6e99799d8097da76a29
+v0.3.5 b111277f464cf66fa34b67819a83ea683e0f64df
+v0.4.0.2 da52d065a4c4f52d6017a537f3a80326f5af8bdc
+v0.4.1 4474b9b3718653aeb0aee26422caefb90460cc0e
+v0.4.2 28c0797d30175043bbfa31307b63aab4197cf996
+v0.4.4 9a9ff839a9b70cb2601d7696af743f5652395389
+v0.4.5 37285038d3f91fa1b49159c4a35a8383168e644f
+v0.4.6 c6df34f55ba8c5a911b60d3f9eb20e3fa45f68c1
+v0.4.7 ae5cb7a0dab8b1bed9d52292f9b4e8e66a0f8ec9
+v0.5 de33bc46452b1046387ee2b3a03191b2c71135fb
+v0.5.1 8f22f52235c48eada586795ac57edb22688e4d08
+v0.5.2 0129a8a4ff8da5314e8e4d4d3336e89667ff6d54
+v0.5.3 285c52eb693265a0a530543e9ca0aeb593a2a55e
+v0.5.4 a422246dc30bce11e970514f13f7c110f4470cc3
+v0.5.5 8f3539a27b28851153454eb737da9624cccaed2d
+v0.5.6 860a48765d1016ba226fb2c64aea35a45fe40e4a
+v0.5.7 9ce3045f139ee29bf0eea5ec59dfe7df9c3d2c51
+v0.5.8 2cee2e05ba0f3824fdbb8b957958800fa99071a1
+v0.5.9 ad145da3bcded0fe75306df02061d07a633963c3
+v0.5.10 ff4c1557d8358f158f4358788ff18ef93d2470ef
+v0.5.11 d1866845df423c6855e2b365ff230cf7d89a420b
+v0.5.12 f6e8ef27546e9a406724841e75f8df71cc4c2c97
+v0.5.13 dd34ccfe288ebda8e2568cf550994ae890379f45
+v0.6 60a20537c5f3fcc7f11946d81aebc8f90c08c117
+v0.6.1 cf5c4aeacb1703d0ffd35bdb5297d3494fee9a22
+v0.6.2 095ef5f9f60fca1b269aabcc3cfd322b17b9e674
+v0.6.3 d5f9caa82d80cdcc7f1b7748f2cf3ccbf94f82a3
+v0.6.4 58f376002e0c654fbc2de127765fa297cf694a33
+v0.6.5 70d17cd69b80e7627dcad8fd5b6438f2309ac307
+v0.6.6 d4eb5a6eafdd4803d09e6f3963918216cca5a81f
+v0.7 a2ae18d1dbd1e59c38ce22fa25ddffbd1d3084e3
+v0.7.1 1e1de64a1e59476b7b3d3558b92149246d5c3292
+v0.7.2 59bd58c940673b4c4a6a86f299022614bcf42b22
+v0.7.3 d7acd8b68f8131e0e714400cf124a3e228638643
+v0.7.4 e93649882c5e914eec4a8b8b593dc0587e497033
+v0.7.5 8abc0afe464626a461d2c7e192c912eeebeccc65
+v0.7.6 69aacd9d179fe6dd3110253a099c38b34cff7899
+v0.7.7 f1caed7299cdba5e45635d8377da6cc1e5fd7072
+v0.7.8 2189fe8741b635d8394d55dee28959425cfd5ad0
+v0.7.9 86dc54b836a973f132ed26db9f5a60b29f9b25b2
+v0.7.10 47a42ff432db8a2494e922ca5e767e58020f0446
+v0.7.11 f46ed718ba8d6bb4d42cd7b0f0150c406017c373
+v0.8 d56e04c0d960c8d3d4ab88b578dc508a2b4e07dc
+v0.8.1 cd6f26664cb8d32a13847d6648567c47c580e248
+v0.8.2 7cc0999aebfe63b6bb6dd83c1dff62c3915c4820
+v0.8.3 986500fe2f10870018f1fba1e5465476b8915977
+v0.8.4 f0d0b82b8d6f5f450952113e234d0a5a49e80c48
+v0.8.5 f5d9ddc6c2a2802a61e5ce590c6b6688bf736a69
+v0.8.6 75904ed7452c6cbfb2a70cd379a899d8a75b97c2
+v0.8.7 4f9e568019492126e236da85b5ca0a059f292900
+v0.8.8 a949a49d5efaaf2d881143d23e9fb5ff9e28e88c
+v0.8.9 a034926264cd1025c69d6ceb3fe444965f960b75
+v0.8.10 59accdd814398b884167c0f41dbf46dcccf0c29c
+v0.8.11 d630ec9092f039cb5e51c5dd6d92fc47b91407e4
+v0.8.12 57871c99031cf597ffa0d819057ac1401e129f32
+v0.8.13 e43e6e876513450d235124fcb711f1724ed9814c
+v0.8.14 71901c94dbbaa2f9f156a80c14cc7ea65219fc7c
+v0.8.15 5a7c2f90361e72e9c35e88abd2e11acdc4aba354
+v0.8.16 f41dc62a12cc0f3cb8c5453c0caa0ba21e1bd01e
+v0.9.0 58b84fb401a804967aa0dd5ee66fafa90194fd30
+v0.9.1 89f68bf0cbf53a81b0553d3816fdbe522f941fa1
+v0.9.2 c7d3b28b93faa6c955fb24bc64031512ee985ee9
+v0.9.3 8704c1ff952b5c85a44f50143bbd1a4f7b4887e2
+v0.9.4 fa310d9e2542cf497d92f007baec8fd751ffa99c
+v0.9.5 3b1e9017c560499786d8a0e46aaaeea65037acac
+v0.9.6 0b4d08fab8e35a264bc7383e878538f8083c33e5
+v0.9.7 f8cb04e4a8b032b72a909f68b808a50936184bee
@@ -0,0 +1,14 @@
+#!/usr/bin/env python3
+
+from PyQt5.QtWidgets import QApplication, QLabel
+from openpilot.selfdrive.ui.qt.python_helpers import set_main_window
+
+
+if __name__ == "__main__":
+  app = QApplication([])
+  label = QLabel('Hello World!')
+
+  # Set full screen and rotate
+  set_main_window(label)
+
+  app.exec_()
@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:fa3f1c39a4e82adfb52d43fc0ad6773a70dbaa4fc79109a7d6b6c1f73b298eac
+size 2833
@@ -30,7 +30,7 @@ if __name__ == '__main__':
        elif event.type == 3 and event.code == 57 and event.value == -1:
          fingers[current_slot] = [-1, -1]
        elif event.type == 3 and event.code == 53:
-          fingers[current_slot][1] = event.value
+          fingers[current_slot][1] = h - (h - event.value)
          if fingers[current_slot][0] != -1:
            touch_points.append(fingers[current_slot].copy())
        elif event.type == 3 and event.code == 54:
@@ -26,6 +26,10 @@ for pathdef, fn in {'TRANSFORM': 'transforms/transform.cl', 'LOADYUV': 'transfor
    xenv['CXXFLAGS'].append(f'-D{pathdef}_PATH=\\"{File(fn).abspath}\\"')

 # Compile cython
+snpe_rpath_qcom = "/data/pythonpath/third_party/snpe/larch64"
+snpe_rpath_pc = f"{Dir('#').abspath}/third_party/snpe/x86_64-linux-clang"
+snpe_rpath = lenvCython['RPATH'] + [snpe_rpath_qcom if arch == "larch64" else snpe_rpath_pc]
+
 cython_libs = envCython["LIBS"] + libs
 commonmodel_lib = lenv.Library('commonmodel', common_src)
 lenvCython.Program('models/commonmodel_pyx.so', 'models/commonmodel_pyx.pyx', LIBS=[commonmodel_lib, *cython_libs], FRAMEWORKS=frameworks)
@@ -16,6 +16,7 @@ class ModelConstants:
  MODEL_FREQ = 20
  FEATURE_LEN = 512
  FULL_HISTORY_BUFFER_LEN = 99
+  HISTORY_BUFFER_LEN = 24
  DESIRE_LEN = 8
  TRAFFIC_CONVENTION_LEN = 2
  LAT_PLANNER_STATE_LEN = 4
@@ -72,13 +73,14 @@ class Plan:
 class Meta:
  ENGAGED = slice(0, 1)
  # next 2, 4, 6, 8, 10 seconds
-  GAS_DISENGAGE = slice(1, 36, 7)
-  BRAKE_DISENGAGE = slice(2, 36, 7)
-  STEER_OVERRIDE = slice(3, 36, 7)
-  HARD_BRAKE_3 = slice(4, 36, 7)
-  HARD_BRAKE_4 = slice(5, 36, 7)
-  HARD_BRAKE_5 = slice(6, 36, 7)
-  GAS_PRESS = slice(7, 36, 7)
+  GAS_DISENGAGE = slice(1, 31, 6)
+  BRAKE_DISENGAGE = slice(2, 31, 6)
+  STEER_OVERRIDE = slice(3, 31, 6)
+  HARD_BRAKE_3 = slice(4, 31, 6)
+  HARD_BRAKE_4 = slice(5, 31, 6)
+  HARD_BRAKE_5 = slice(6, 31, 6)
  # next 0, 2, 4, 6, 8, 10 seconds
-  LEFT_BLINKER = slice(36, 48, 2)
-  RIGHT_BLINKER = slice(37, 48, 2)
+  GAS_PRESS = slice(31, 55, 4)
+  BRAKE_PRESS = slice(32, 55, 4)
+  LEFT_BLINKER = slice(33, 55, 4)
+  RIGHT_BLINKER = slice(34, 55, 4)
@@ -3,11 +3,21 @@ import capnp
 import numpy as np
 from cereal import log
 from openpilot.selfdrive.modeld.constants import ModelConstants, Plan, Meta
+from openpilot.selfdrive.controls.lib.drive_helpers import MIN_SPEED

 SEND_RAW_PRED = os.getenv('SEND_RAW_PRED')

 ConfidenceClass = log.ModelDataV2.ConfidenceClass

+def curv_from_psis(psi_target, psi_rate, vego, delay):
+  vego = np.clip(vego, MIN_SPEED, np.inf)
+  curv_from_psi = psi_target / (vego * delay)  # epsilon to prevent divide-by-zero
+  return 2*curv_from_psi - psi_rate / vego
+
+def get_curvature_from_plan(plan, vego, delay):
+  psi_target = np.interp(delay, ModelConstants.T_IDXS, plan[:, Plan.T_FROM_CURRENT_EULER][:, 2])
+  psi_rate = plan[:, Plan.ORIENTATION_RATE][0, 2]
+  return curv_from_psis(psi_target, psi_rate, vego, delay)

 class PublishState:
  def __init__(self):
@@ -65,6 +75,8 @@ def fill_model_msg(base_msg: capnp._DynamicStructBuilder, extended_msg: capnp._D
  extended_msg.valid = valid
  base_msg.valid = valid

+  desired_curv = float(get_curvature_from_plan(net_output_data['plan'][0], v_ego, delay))
+
  driving_model_data = base_msg.drivingModelData

  driving_model_data.frameId = vipc_frame_id
@@ -73,7 +85,7 @@ def fill_model_msg(base_msg: capnp._DynamicStructBuilder, extended_msg: capnp._D
  driving_model_data.modelExecutionTime = model_execution_time

  action = driving_model_data.action
-  action.desiredCurvature = float(net_output_data['desired_curvature'][0,0])
+  action.desiredCurvature = desired_curv

  modelV2 = extended_msg.modelV2
  modelV2.frameId = vipc_frame_id
@@ -108,7 +120,7 @@ def fill_model_msg(base_msg: capnp._DynamicStructBuilder, extended_msg: capnp._D

  # lateral planning
  action = modelV2.action
-  action.desiredCurvature = float(net_output_data['desired_curvature'][0,0])
+  action.desiredCurvature = desired_curv

  # times at X_IDXS according to model plan
  PLAN_T_IDXS = [np.nan] * ModelConstants.IDX_N
@@ -169,8 +181,8 @@ def fill_model_msg(base_msg: capnp._DynamicStructBuilder, extended_msg: capnp._D
  disengage_predictions.brake3MetersPerSecondSquaredProbs = net_output_data['meta'][0,Meta.HARD_BRAKE_3].tolist()
  disengage_predictions.brake4MetersPerSecondSquaredProbs = net_output_data['meta'][0,Meta.HARD_BRAKE_4].tolist()
  disengage_predictions.brake5MetersPerSecondSquaredProbs = net_output_data['meta'][0,Meta.HARD_BRAKE_5].tolist()
-  #disengage_predictions.gasPressProbs = net_output_data['meta'][0,Meta.GAS_PRESS].tolist()
-  #disengage_predictions.brakePressProbs = net_output_data['meta'][0,Meta.BRAKE_PRESS].tolist()
+  disengage_predictions.gasPressProbs = net_output_data['meta'][0,Meta.GAS_PRESS].tolist()
+  disengage_predictions.brakePressProbs = net_output_data['meta'][0,Meta.BRAKE_PRESS].tolist()

  publish_state.prev_brake_5ms2_probs[:-1] = publish_state.prev_brake_5ms2_probs[1:]
  publish_state.prev_brake_5ms2_probs[-1] = net_output_data['meta'][0,Meta.HARD_BRAKE_5][0]
@@ -1,15 +1,21 @@
 #!/usr/bin/env python3
+import os
 from openpilot.system.hardware import TICI

-from openpilot.selfdrive.modeld.runners.model_runner import ONNXRunner, TinygradRunner
-from openpilot.sunnypilot.models.helpers import is_active_model_20hz
-
 #
-import os
+if TICI:
+  from tinygrad.tensor import Tensor
+  from tinygrad.dtype import dtypes
+  from openpilot.selfdrive.modeld.runners.tinygrad_helpers import qcom_tensor_from_opencl_address
+  os.environ['QCOM'] = '1'
+else:
+  from openpilot.selfdrive.modeld.runners.ort_helpers import make_onnx_cpu_runner
 import time
+import pickle
 import numpy as np
 import cereal.messaging as messaging
 from cereal import car, log
+from pathlib import Path
 from setproctitle import setproctitle
 from cereal.messaging import PubMaster, SubMaster
 from msgq.visionipc import VisionIpcClient, VisionStreamType, VisionBuf
@@ -21,16 +27,19 @@ from openpilot.common.realtime import config_realtime_process
 from openpilot.common.transformations.camera import DEVICE_CAMERAS
 from openpilot.common.transformations.model import get_warp_matrix
 from openpilot.system import sentry
-from openpilot.system.hardware import PC
 from openpilot.selfdrive.controls.lib.desire_helper import DesireHelper
 from openpilot.selfdrive.modeld.parse_model_outputs import Parser
 from openpilot.selfdrive.modeld.fill_model_msg import fill_model_msg, fill_pose_msg, PublishState
 from openpilot.selfdrive.modeld.constants import ModelConstants
 from openpilot.selfdrive.modeld.models.commonmodel_pyx import DrivingModelFrame, CLContext

+
 PROCESS_NAME = "selfdrive.modeld.modeld"
-USE_ONNX = bool(os.getenv('USE_ONNX', PC))
-IS_20HZ_MODEL_DEFAULT = False
+SEND_RAW_PRED = os.getenv('SEND_RAW_PRED')
+
+MODEL_PATH = Path(__file__).parent / 'models/supercombo.onnx'
+MODEL_PKL_PATH = Path(__file__).parent / 'models/supercombo_tinygrad.pkl'
+METADATA_PATH = Path(__file__).parent / 'models/supercombo_metadata.pkl'

 class FrameMeta:
  frame_id: int = 0
@@ -48,86 +57,79 @@ class ModelState:
  prev_desire: np.ndarray  # for tracking the rising edge of the pulse

  def __init__(self, context: CLContext):
-    self.is_20hz = IS_20HZ_MODEL_DEFAULT
-    if (active_20hz := is_active_model_20hz(None)) is not None:
-      self.is_20hz = active_20hz
-
-    self.frames = {'input_imgs': DrivingModelFrame(context, self.is_20hz), 'big_input_imgs': DrivingModelFrame(context, self.is_20hz)}
+    self.frames = {'input_imgs': DrivingModelFrame(context), 'big_input_imgs': DrivingModelFrame(context)}
    self.prev_desire = np.zeros(ModelConstants.DESIRE_LEN, dtype=np.float32)
    self.full_features_20Hz = np.zeros((ModelConstants.FULL_HISTORY_BUFFER_LEN, ModelConstants.FEATURE_LEN), dtype=np.float32)
    self.desire_20Hz =  np.zeros((ModelConstants.FULL_HISTORY_BUFFER_LEN + 1, ModelConstants.DESIRE_LEN), dtype=np.float32)
-    # Initialize model runner
-    self.model_runner = ONNXRunner(self.frames) if (not TICI) and USE_ONNX else TinygradRunner(self.frames)

    # img buffers are managed in openCL transform code
-    self.numpy_inputs = {}
+    self.numpy_inputs = {
+      'desire': np.zeros((1, (ModelConstants.HISTORY_BUFFER_LEN+1), ModelConstants.DESIRE_LEN), dtype=np.float32),
+      'traffic_convention': np.zeros((1, ModelConstants.TRAFFIC_CONVENTION_LEN), dtype=np.float32),
+      'features_buffer': np.zeros((1, ModelConstants.HISTORY_BUFFER_LEN,  ModelConstants.FEATURE_LEN), dtype=np.float32),
+    }

-    for key, shape in self.model_runner.input_shapes.items():
-      if key not in self.frames:  # Managed by opencl
-        self.numpy_inputs[key] = np.zeros(shape, dtype=np.float32)
+    with open(METADATA_PATH, 'rb') as f:
+      model_metadata = pickle.load(f)
+    self.input_shapes =  model_metadata['input_shapes']

+    self.output_slices = model_metadata['output_slices']
+    net_output_size = model_metadata['output_shapes']['outputs'][1]
+    self.output = np.zeros(net_output_size, dtype=np.float32)
    self.parser = Parser()

-    net_output_size = self.model_runner.model_metadata['output_shapes']['outputs'][1]
-    self.output = np.zeros(net_output_size, dtype=np.float32)
+    if TICI:
+      self.tensor_inputs = {k: Tensor(v, device='NPY').realize() for k,v in self.numpy_inputs.items()}
+      with open(MODEL_PKL_PATH, "rb") as f:
+        self.model_run = pickle.load(f)
+    else:
+      self.onnx_cpu_runner = make_onnx_cpu_runner(MODEL_PATH)

-    num_elements = self.numpy_inputs['features_buffer'].shape[1]
-    step_size = int(-100 / num_elements)
-    self.full_features_20Hz_idxs = np.arange(step_size, step_size * (num_elements + 1), step_size)[::-1]
-    self.desire_reshape_dims = (self.numpy_inputs['desire'].shape[0], self.numpy_inputs['desire'].shape[1], -1, self.numpy_inputs['desire'].shape[2])
+  def slice_outputs(self, model_outputs: np.ndarray) -> dict[str, np.ndarray]:
+    parsed_model_outputs = {k: model_outputs[np.newaxis, v] for k,v in self.output_slices.items()}
+    if SEND_RAW_PRED:
+      parsed_model_outputs['raw_pred'] = model_outputs.copy()
+    return parsed_model_outputs

  def run(self, buf: VisionBuf, wbuf: VisionBuf, transform: np.ndarray, transform_wide: np.ndarray,
-          inputs: dict[str, np.ndarray], prepare_only: bool) -> dict[str, np.ndarray] | None:
+                inputs: dict[str, np.ndarray], prepare_only: bool) -> dict[str, np.ndarray] | None:
    # Model decides when action is completed, so desire input is just a pulse triggered on rising edge
    inputs['desire'][0] = 0
    new_desire = np.where(inputs['desire'] - self.prev_desire > .99, inputs['desire'], 0)
    self.prev_desire[:] = inputs['desire']

-    if self.is_20hz:
-      self.desire_20Hz[:-1] = self.desire_20Hz[1:]
-      self.desire_20Hz[-1] = new_desire
-      self.numpy_inputs['desire'][:] = self.desire_20Hz.reshape(self.desire_reshape_dims).max(axis=2)
-    else:
-      self.numpy_inputs['desire'][0,:-1] = self.numpy_inputs['desire'][0,1:]
-      self.numpy_inputs['desire'][0,-1] = new_desire
-
-    for key in self.numpy_inputs:
-      if key in inputs and key not in ['desire']:
-        self.numpy_inputs[key][:] = inputs[key]
+    self.desire_20Hz[:-1] = self.desire_20Hz[1:]
+    self.desire_20Hz[-1] = new_desire
+    self.numpy_inputs['desire'][:] = self.desire_20Hz.reshape((1,25,4,-1)).max(axis=2)

+    self.numpy_inputs['traffic_convention'][:] = inputs['traffic_convention']
    imgs_cl = {'input_imgs': self.frames['input_imgs'].prepare(buf, transform.flatten()),
               'big_input_imgs': self.frames['big_input_imgs'].prepare(wbuf, transform_wide.flatten())}

-    # Prepare inputs using the model runner
-    self.model_runner.prepare_inputs(imgs_cl, self.numpy_inputs)
+    if TICI:
+      # The imgs tensors are backed by opencl memory, only need init once
+      for key in imgs_cl:
+        if key not in self.tensor_inputs:
+          self.tensor_inputs[key] = qcom_tensor_from_opencl_address(imgs_cl[key].mem_address, self.input_shapes[key], dtype=dtypes.uint8)
+    else:
+      for key in imgs_cl:
+        self.numpy_inputs[key] = self.frames[key].buffer_from_cl(imgs_cl[key]).reshape(self.input_shapes[key])

    if prepare_only:
      return None

-    # Run model inference
-    self.output = self.model_runner.run_model()
-    outputs = self.parser.parse_outputs(self.model_runner.slice_outputs(self.output))
-
-    if self.is_20hz:
-      self.full_features_20Hz[:-1] = self.full_features_20Hz[1:]
-      self.full_features_20Hz[-1] = outputs['hidden_state'][0, :]
-      self.numpy_inputs['features_buffer'][:] = self.full_features_20Hz[self.full_features_20Hz_idxs]
+    if TICI:
+      self.output = self.model_run(**self.tensor_inputs).numpy().flatten()
    else:
-      self.numpy_inputs['features_buffer'][0,:-1] = self.numpy_inputs['features_buffer'][0,1:]
-      self.numpy_inputs['features_buffer'][0,-1] = outputs['hidden_state'][0, :]
+      self.output = self.onnx_cpu_runner.run(None, self.numpy_inputs)[0].flatten()

-    if "desired_curvature" in outputs:
-      input_name_prev = None
+    outputs = self.parser.parse_outputs(self.slice_outputs(self.output))

-      if "prev_desired_curvs" in self.numpy_inputs.keys():
-        input_name_prev = 'prev_desired_curvs'
-      elif "prev_desired_curv" in self.numpy_inputs.keys():
-        input_name_prev = 'prev_desired_curv'
+    self.full_features_20Hz[:-1] = self.full_features_20Hz[1:]
+    self.full_features_20Hz[-1] = outputs['hidden_state'][0, :]

-      if input_name_prev is not None:
-        len = outputs['desired_curvature'][0].size
-        self.numpy_inputs['prev_desired_curv'][0,:-len] = self.numpy_inputs['prev_desired_curv'][0,len:]
-        self.numpy_inputs['prev_desired_curv'][0,-len,:] = outputs['desired_curvature'][0, :]
+    idxs = np.arange(-4,-100,-4)[::-1]
+    self.numpy_inputs['features_buffer'][:] = self.full_features_20Hz[idxs]
    return outputs


@@ -268,10 +270,7 @@ def main(demo=False):
    inputs:dict[str, np.ndarray] = {
      'desire': vec_desire,
      'traffic_convention': traffic_convention,
-    }
-
-    if "lateral_control_params" in model.numpy_inputs.keys():
-      inputs['lateral_control_params'] = np.array([v_ego, steer_delay], dtype=np.float32)
+      }

    mt1 = time.perf_counter()
    model_output = model.run(buf_main, buf_extra, model_transform_main, model_transform_extra, inputs, prepare_only)
@@ -5,11 +5,11 @@

 #include "common/clutil.h"

-DrivingModelFrame::DrivingModelFrame(cl_device_id device_id, cl_context context, bool is_20hz) : ModelFrame(device_id, context), is_20hz(is_20hz) {
+DrivingModelFrame::DrivingModelFrame(cl_device_id device_id, cl_context context) : ModelFrame(device_id, context) {
  input_frames = std::make_unique<uint8_t[]>(buf_size);
  input_frames_cl = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_WRITE, buf_size, NULL, &err));
-  img_buffer_20hz_cl = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_WRITE, buf_len*frame_size_bytes, NULL, &err));
-  region.origin = (buf_len - 1) * frame_size_bytes;
+  img_buffer_20hz_cl = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_WRITE, 5*frame_size_bytes, NULL, &err));
+  region.origin = 4 * frame_size_bytes;
  region.size = frame_size_bytes;
  last_img_cl = CL_CHECK_ERR(clCreateSubBuffer(img_buffer_20hz_cl, CL_MEM_READ_WRITE, CL_BUFFER_CREATE_TYPE_REGION, &region, &err));

@@ -20,7 +20,7 @@ DrivingModelFrame::DrivingModelFrame(cl_device_id device_id, cl_context context,
 cl_mem* DrivingModelFrame::prepare(cl_mem yuv_cl, int frame_width, int frame_height, int frame_stride, int frame_uv_offset, const mat3& projection) {
  run_transform(yuv_cl, MODEL_WIDTH, MODEL_HEIGHT, frame_width, frame_height, frame_stride, frame_uv_offset, projection);

-  for (int i = 0; i < (buf_len - 1); i++) {
+  for (int i = 0; i < 4; i++) {
    CL_CHECK(clEnqueueCopyBuffer(q, img_buffer_20hz_cl, img_buffer_20hz_cl, (i+1)*frame_size_bytes, i*frame_size_bytes, frame_size_bytes, 0, nullptr, nullptr));
  }
  loadyuv_queue(&loadyuv, q, y_cl, u_cl, v_cl, last_img_cl);
@@ -64,7 +64,7 @@ protected:

 class DrivingModelFrame : public ModelFrame {
 public:
-  DrivingModelFrame(cl_device_id device_id, cl_context context, bool is_20hz = false);
+  DrivingModelFrame(cl_device_id device_id, cl_context context);
  ~DrivingModelFrame();
  cl_mem* prepare(cl_mem yuv_cl, int frame_width, int frame_height, int frame_stride, int frame_uv_offset, const mat3& projection);

@@ -74,9 +74,6 @@ public:
  const int buf_size = MODEL_FRAME_SIZE * 2;
  const size_t frame_size_bytes = MODEL_FRAME_SIZE * sizeof(uint8_t);

-  const bool is_20hz;
-  const int buf_len = is_20hz ? 5 : 2;
-
 private:
  LoadYUVState loadyuv;
  cl_mem img_buffer_20hz_cl, last_img_cl, input_frames_cl;
@@ -19,7 +19,7 @@ cdef extern from "selfdrive/modeld/models/commonmodel.h":

  cppclass DrivingModelFrame:
    int buf_size
-    DrivingModelFrame(cl_device_id, cl_context, bint)
+    DrivingModelFrame(cl_device_id, cl_context)

  cppclass MonitoringModelFrame:
    int buf_size
@@ -55,8 +55,8 @@ cdef class ModelFrame:
 cdef class DrivingModelFrame(ModelFrame):
  cdef cppDrivingModelFrame * _frame

-  def __cinit__(self, CLContext context, bint is_20hz=False):
-    self._frame = new cppDrivingModelFrame(context.device_id, context.context, is_20hz)
+  def __cinit__(self, CLContext context):
+    self._frame = new cppDrivingModelFrame(context.device_id, context.context)
    self.frame = <cppModelFrame*>(self._frame)
    self.buf_size = self._frame.buf_size

@@ -1,3 +1,3 @@
 version https://git-lfs.github.com/spec/v1
-oid sha256:39786068cae1ed8c0dc34ef80c281dfcc67ed18a50e06b90765c49bcfdbf7db4
-size 51453312
+oid sha256:72d3d6f8d3c98f5431ec86be77b6350d7d4f43c25075c0106f1d1e7ec7c77668
+size 49096168
@@ -85,7 +85,6 @@ class Parser:
    outs[name + '_stds'] = pred_std_final.reshape(final_shape)

  def parse_outputs(self, outs: dict[str, np.ndarray]) -> dict[str, np.ndarray]:
-    """ Parse the model outputs into a dictionary of numpy arrays. The input_keys are used to determine how the output should be parsed. """
    self.parse_mdn('plan', outs, in_N=ModelConstants.PLAN_MHP_N, out_N=ModelConstants.PLAN_MHP_SELECTION,
                   out_shape=(ModelConstants.IDX_N,ModelConstants.PLAN_WIDTH))
    self.parse_mdn('lane_lines', outs, in_N=0, out_N=0, out_shape=(ModelConstants.NUM_LANE_LINES,ModelConstants.IDX_N,ModelConstants.LANE_LINES_WIDTH))
@@ -97,8 +96,6 @@ class Parser:
                   out_shape=(ModelConstants.LEAD_TRAJ_LEN,ModelConstants.LEAD_WIDTH))
    if 'lat_planner_solution' in outs:
      self.parse_mdn('lat_planner_solution', outs, in_N=0, out_N=0, out_shape=(ModelConstants.IDX_N,ModelConstants.LAT_PLANNER_SOLUTION_WIDTH))
-    if 'desired_curvature' in outs:
-      self.parse_mdn('desired_curvature', outs, in_N=0, out_N=0, out_shape=(ModelConstants.DESIRED_CURV_WIDTH,))
    for k in ['lead_prob', 'lane_lines_prob', 'meta']:
      self.parse_binary_crossentropy(k, outs)
    self.parse_categorical_crossentropy('desire_state', outs, out_shape=(ModelConstants.DESIRE_PRED_WIDTH,))
@@ -1,120 +0,0 @@
-import os
-from openpilot.system.hardware import TICI
-
-from openpilot.sunnypilot.modeld.run_helpers import get_custom_model_paths
-#
-from tinygrad.tensor import Tensor, dtypes
-from openpilot.selfdrive.modeld.runners.tinygrad_helpers import qcom_tensor_from_opencl_address
-from openpilot.selfdrive.modeld.runners.ort_helpers import make_onnx_cpu_runner, ORT_TYPES_TO_NP_TYPES
-import pickle
-import numpy as np
-from pathlib import Path
-from abc import ABC, abstractmethod
-from openpilot.selfdrive.modeld.models.commonmodel_pyx import DrivingModelFrame, CLMem
-from openpilot.system.hardware import PC
-
-if TICI:
-  os.environ['QCOM'] = '1'
-
-SEND_RAW_PRED = os.getenv('SEND_RAW_PRED')
-MODEL_PATH = Path(__file__).parent / '../models/supercombo.onnx'
-MODEL_PKL_PATH = Path(__file__).parent / '../models/supercombo_tinygrad.pkl'
-METADATA_PATH = Path(__file__).parent / '../models/supercombo_metadata.pkl'
-USE_ONNX = os.getenv('USE_ONNX', PC)
-
-
-class ModelRunner(ABC):
-  """Abstract base class for model runners that defines the interface for running ML models."""
-
-  def __init__(self):
-    """Initialize the model runner with paths to model and metadata files."""
-    self.model_paths = ({"model": MODEL_PATH, "metadata": METADATA_PATH} if USE_ONNX else
-                        get_custom_model_paths() or {"model": MODEL_PKL_PATH, "metadata": METADATA_PATH})
-
-    with open(self.model_paths["metadata"], 'rb') as f:
-      self.model_metadata = pickle.load(f)
-
-    self.input_shapes = self.model_metadata['input_shapes']
-    self.output_slices = self.model_metadata['output_slices']
-    self.inputs: dict = {}
-
-  @abstractmethod
-  def prepare_inputs(self, imgs_cl: dict[str, CLMem], numpy_inputs: dict[str, np.ndarray]) -> dict:
-    """Prepare inputs for model inference."""
-
-  @abstractmethod
-  def run_model(self):
-    """Run model inference with prepared inputs."""
-
-  def slice_outputs(self, model_outputs: np.ndarray) -> dict:
-    """Slice model outputs according to metadata configuration."""
-    parsed_outputs = {k: model_outputs[np.newaxis, v] for k, v in self.output_slices.items()}
-    if SEND_RAW_PRED:
-      parsed_outputs['raw_pred'] = model_outputs.copy()
-    return parsed_outputs
-
-
-class TinygradRunner(ModelRunner):
-  """Tinygrad implementation of model runner for TICI hardware."""
-
-  def __init__(self, frames: dict[str, DrivingModelFrame] | None = None):
-    super().__init__()
-    if not str(self.model_paths["model"]).endswith("_tinygrad.pkl"):
-      raise ValueError(f"Tinygrad model must be a _tinygrad.pkl file, we got {self.model_paths['model']}")
-
-    # Load Tinygrad model
-    with open(self.model_paths["model"], "rb") as f:
-      self.model_run = pickle.load(f)
-
-    self.input_to_dtype = {}
-    self.input_to_device = {}
-
-    for idx, name in enumerate(self.model_run.captured.expected_names):
-      self.input_to_dtype[name] = self.model_run.captured.expected_st_vars_dtype_device[idx][2]  # 2 is the dtype
-      self.input_to_device[name] = self.model_run.captured.expected_st_vars_dtype_device[idx][3]  # 3 is the device
-
-    assert TICI or frames is not None, "TinygradRunner requires frames for non-TICI hardware"
-    self.frames = frames
-    self.is_memory_model = None  # Use None to indicate that it hasn't been determined yet
-
-  def prepare_inputs(self, imgs_cl: dict[str, CLMem], numpy_inputs: dict[str, np.ndarray]) -> dict:
-    # Initialize image tensors if not already done
-    for key in imgs_cl:
-      if TICI and key not in self.inputs:
-        self.inputs[key] = qcom_tensor_from_opencl_address(imgs_cl[key].mem_address, self.input_shapes[key], dtype=dtypes.uint8)
-      elif not TICI:
-        shape = self.frames[key].buffer_from_cl(imgs_cl[key]).reshape(self.input_shapes[key])
-        self.inputs[key] = Tensor(shape, device=self.input_to_device[key], dtype=self.input_to_dtype[key]).realize()
-
-    # Update numpy inputs
-    for key, value in numpy_inputs.items():
-      if key not in imgs_cl:
-        self.inputs[key] = Tensor(value, device=self.input_to_device[key], dtype=self.input_to_dtype[key]).realize()
-
-    return self.inputs
-
-  def run_model(self):
-    return self.model_run(**self.inputs).numpy().flatten()
-
-
-class ONNXRunner(ModelRunner):
-  """ONNX implementation of model runner for non-TICI hardware."""
-
-  def __init__(self, frames: dict[str, DrivingModelFrame]):
-    super().__init__()
-    self.runner = make_onnx_cpu_runner(self.model_paths["model"])
-    self.frames = frames
-
-    self.input_to_nptype = {
-      model_input.name: ORT_TYPES_TO_NP_TYPES[model_input.type]
-      for model_input in self.runner.get_inputs()
-    }
-
-  def prepare_inputs(self, imgs_cl: dict[str, CLMem], numpy_inputs: dict[str, np.ndarray]) -> dict:
-    self.inputs = numpy_inputs.copy()
-    for key in imgs_cl:
-      self.inputs[key] = self.frames[key].buffer_from_cl(imgs_cl[key]).reshape(self.input_shapes[key]).astype(dtype=np.float32)
-    return self.inputs
-
-  def run_model(self):
-    return self.runner.run(None, self.inputs)[0].flatten()
@@ -0,0 +1 @@
+benchmark
@@ -0,0 +1,192 @@
+#include <SNPE/SNPE.hpp>
+#include <SNPE/SNPEBuilder.hpp>
+#include <SNPE/SNPEFactory.hpp>
+#include <DlContainer/IDlContainer.hpp>
+#include <DlSystem/DlError.hpp>
+#include <DlSystem/ITensor.hpp>
+#include <DlSystem/ITensorFactory.hpp>
+#include <iostream>
+#include <fstream>
+#include <sstream>
+
+using namespace std;
+
+int64_t timespecDiff(struct timespec *timeA_p, struct timespec *timeB_p) {
+  return ((timeA_p->tv_sec * 1000000000) + timeA_p->tv_nsec) - ((timeB_p->tv_sec * 1000000000) + timeB_p->tv_nsec);
+}
+
+void PrintErrorStringAndExit() {
+  cout << "ERROR!" << endl;
+  const char* const errStr = zdl::DlSystem::getLastErrorString();
+  std::cerr << errStr << std::endl;
+  std::exit(EXIT_FAILURE);
+}
+
+
+zdl::DlSystem::Runtime_t checkRuntime() {
+  static zdl::DlSystem::Version_t Version = zdl::SNPE::SNPEFactory::getLibraryVersion();
+  static zdl::DlSystem::Runtime_t Runtime;
+  std::cout << "SNPE Version: " << Version.asString().c_str() << std::endl; //Print Version number
+  if (zdl::SNPE::SNPEFactory::isRuntimeAvailable(zdl::DlSystem::Runtime_t::DSP)) {
+    std::cout << "Using DSP runtime" << std::endl;
+    Runtime = zdl::DlSystem::Runtime_t::DSP;
+  } else if (zdl::SNPE::SNPEFactory::isRuntimeAvailable(zdl::DlSystem::Runtime_t::GPU)) {
+    std::cout << "Using GPU runtime" << std::endl;
+    Runtime = zdl::DlSystem::Runtime_t::GPU;
+  } else {
+    std::cout << "Using cpu runtime" << std::endl;
+    Runtime = zdl::DlSystem::Runtime_t::CPU;
+  }
+  return Runtime;
+}
+
+void test(char *filename) {
+  static zdl::DlSystem::Runtime_t runtime = checkRuntime();
+  std::unique_ptr<zdl::DlContainer::IDlContainer> container;
+  container = zdl::DlContainer::IDlContainer::open(filename);
+
+  if (!container) { PrintErrorStringAndExit(); }
+  cout << "start build" << endl;
+  std::unique_ptr<zdl::SNPE::SNPE> snpe;
+  {
+    snpe = NULL;
+    zdl::SNPE::SNPEBuilder snpeBuilder(container.get());
+    snpe = snpeBuilder.setOutputLayers({})
+      .setRuntimeProcessor(runtime)
+      .setUseUserSuppliedBuffers(false)
+      //.setDebugMode(true)
+      .build();
+    if (!snpe) {
+      cout << "ERROR!" << endl;
+      const char* const errStr = zdl::DlSystem::getLastErrorString();
+      std::cerr << errStr << std::endl;
+    }
+    cout << "ran snpeBuilder" << endl;
+  }
+
+  const auto &strList_opt = snpe->getInputTensorNames();
+  if (!strList_opt) throw std::runtime_error("Error obtaining input tensor names");
+
+  cout << "get input tensor names done" << endl;
+  const auto &strList = *strList_opt;
+  static zdl::DlSystem::TensorMap inputTensorMap;
+  static zdl::DlSystem::TensorMap outputTensorMap;
+  vector<std::unique_ptr<zdl::DlSystem::ITensor> > inputs;
+  for (int i = 0; i < strList.size(); i++) {
+    cout << "input name: " << strList.at(i) << endl;
+
+    const auto &inputDims_opt = snpe->getInputDimensions(strList.at(i));
+    const auto &inputShape = *inputDims_opt;
+    inputs.push_back(zdl::SNPE::SNPEFactory::getTensorFactory().createTensor(inputShape));
+    inputTensorMap.add(strList.at(i), inputs[i].get());
+  }
+
+  struct timespec start, end;
+  cout << "**** starting benchmark ****" << endl;
+  for (int i = 0; i < 50; i++) {
+    clock_gettime(CLOCK_MONOTONIC, &start);
+    int err = snpe->execute(inputTensorMap, outputTensorMap);
+    assert(err == true);
+    clock_gettime(CLOCK_MONOTONIC, &end);
+    uint64_t timeElapsed = timespecDiff(&end, &start);
+    printf("time: %f ms\n", timeElapsed*1.0/1e6);
+  }
+}
+
+void get_testframe(int index, std::unique_ptr<zdl::DlSystem::ITensor> &input) {
+  FILE * pFile;
+  string filepath="/data/ipt/quantize_samples/sample_input_"+std::to_string(index);
+  pFile = fopen(filepath.c_str(), "rb");
+  int length = 1*6*160*320*4;
+  float * frame_buffer = new float[length/4]; // 32/8
+  fread(frame_buffer, length, 1, pFile);
+  // std::cout << *(frame_buffer+length/4-1) << std::endl;
+  std::copy(frame_buffer, frame_buffer+(length/4), input->begin());
+  fclose(pFile);
+}
+
+void SaveITensor(const std::string& path, const zdl::DlSystem::ITensor* tensor)
+{
+   std::ofstream os(path, std::ofstream::binary);
+   if (!os)
+   {
+      std::cerr << "Failed to open output file for writing: " << path << "\n";
+      std::exit(EXIT_FAILURE);
+   }
+   for ( auto it = tensor->cbegin(); it != tensor->cend(); ++it )
+   {
+      float f = *it;
+      if (!os.write(reinterpret_cast<char*>(&f), sizeof(float)))
+      {
+         std::cerr << "Failed to write data to: " << path << "\n";
+         std::exit(EXIT_FAILURE);
+      }
+   }
+}
+
+void testrun(char* modelfile) {
+  static zdl::DlSystem::Runtime_t runtime = checkRuntime();
+  std::unique_ptr<zdl::DlContainer::IDlContainer> container;
+  container = zdl::DlContainer::IDlContainer::open(modelfile);
+
+  if (!container) { PrintErrorStringAndExit(); }
+  cout << "start build" << endl;
+  std::unique_ptr<zdl::SNPE::SNPE> snpe;
+  {
+    snpe = NULL;
+    zdl::SNPE::SNPEBuilder snpeBuilder(container.get());
+    snpe = snpeBuilder.setOutputLayers({})
+      .setRuntimeProcessor(runtime)
+      .setUseUserSuppliedBuffers(false)
+      //.setDebugMode(true)
+      .build();
+    if (!snpe) {
+      cout << "ERROR!" << endl;
+      const char* const errStr = zdl::DlSystem::getLastErrorString();
+      std::cerr << errStr << std::endl;
+    }
+    cout << "ran snpeBuilder" << endl;
+  }
+
+  const auto &strList_opt = snpe->getInputTensorNames();
+  if (!strList_opt) throw std::runtime_error("Error obtaining input tensor names");
+  cout << "get input tensor names done" << endl;
+
+  const auto &strList = *strList_opt;
+  static zdl::DlSystem::TensorMap inputTensorMap;
+  static zdl::DlSystem::TensorMap outputTensorMap;
+
+  assert(strList.size() == 1);
+  const auto &inputDims_opt = snpe->getInputDimensions(strList.at(0));
+  const auto &inputShape = *inputDims_opt;
+  std::cout << "winkwink" << std::endl;
+
+  for (int i=0; i<10000; i++) {
+    std::unique_ptr<zdl::DlSystem::ITensor> input;
+    input = zdl::SNPE::SNPEFactory::getTensorFactory().createTensor(inputShape);
+    get_testframe(i, input);
+    snpe->execute(input.get(), outputTensorMap);
+    zdl::DlSystem::StringList tensorNames = outputTensorMap.getTensorNames();
+    std::for_each(tensorNames.begin(), tensorNames.end(), [&](const char* name) {
+      std::ostringstream path;
+      path << "/data/opt/Result_" << std::to_string(i) << ".raw";
+      auto tensorPtr = outputTensorMap.getTensor(name);
+      SaveITensor(path.str(), tensorPtr);
+    });
+  }
+}
+
+int main(int argc, char* argv[]) {
+  if (argc < 2) {
+    printf("usage: %s <filename>\n", argv[0]);
+    return -1;
+  }
+
+  if (argc == 2) {
+    while (true) test(argv[1]);
+  } else if (argc == 3) {
+    testrun(argv[1]);
+  }
+  return 0;
+}
+
@@ -0,0 +1,4 @@
+#!/bin/sh -e
+clang++ -I /data/openpilot/third_party/snpe/include/ -L/data/pythonpath/third_party/snpe/aarch64 -lSNPE benchmark.cc -o benchmark
+export LD_LIBRARY_PATH="/data/pythonpath/third_party/snpe/aarch64/:$HOME/openpilot/third_party/snpe/x86_64/:$LD_LIBRARY_PATH"
+exec ./benchmark $1
@@ -0,0 +1,58 @@
+#!/usr/bin/env python3
+import signal
+import subprocess
+
+signal.signal(signal.SIGINT, signal.SIG_DFL)
+signal.signal(signal.SIGTERM, signal.SIG_DFL)
+
+from PyQt5.QtCore import QTimer
+from PyQt5.QtWidgets import QApplication, QWidget, QVBoxLayout, QLabel
+from openpilot.selfdrive.ui.qt.python_helpers import set_main_window
+
+class Window(QWidget):
+  def __init__(self, parent=None):
+    super().__init__(parent)
+
+    layout = QVBoxLayout()
+    self.setLayout(layout)
+
+    self.l = QLabel("jenkins runner")
+    layout.addWidget(self.l)
+    layout.addStretch(1)
+    layout.setContentsMargins(20, 20, 20, 20)
+
+    cmds = [
+      "cat /etc/hostname",
+      "echo AGNOS v$(cat /VERSION)",
+      "uptime -p",
+    ]
+    self.labels = {}
+    for c in cmds:
+      self.labels[c] = QLabel(c)
+      layout.addWidget(self.labels[c])
+
+    self.setStyleSheet("""
+      * {
+        color: white;
+        font-size: 55px;
+        background-color: black;
+        font-family: "JetBrains Mono";
+      }
+    """)
+
+    self.timer = QTimer()
+    self.timer.timeout.connect(self.update)
+    self.timer.start(10 * 1000)
+    self.update()
+
+  def update(self):
+    for cmd, label in self.labels.items():
+      out = subprocess.run(cmd, capture_output=True,
+                           shell=True, check=False, encoding='utf8').stdout
+      label.setText(out.strip())
+
+if __name__ == "__main__":
+  app = QApplication([])
+  w = Window()
+  set_main_window(w)
+  app.exec_()
@@ -0,0 +1,25 @@
+#!/usr/bin/env bash
+set -e
+
+if [ $# -lt 2 ]; then
+  echo "Usage: $0 <base|docs|sim|prebuilt|cl> <arch1> <arch2> ..."
+  exit 1
+fi
+
+SCRIPT_DIR=$(dirname "$0")
+ARCHS=("${@:2}")
+
+source $SCRIPT_DIR/docker_common.sh $1
+
+MANIFEST_AMENDS=""
+for ARCH in ${ARCHS[@]}; do
+  MANIFEST_AMENDS="$MANIFEST_AMENDS --amend $REMOTE_TAG-$ARCH:$COMMIT_SHA"
+done
+
+docker manifest create $REMOTE_TAG $MANIFEST_AMENDS
+docker manifest create $REMOTE_SHA_TAG $MANIFEST_AMENDS
+
+if [[ -n "$PUSH_IMAGE" ]]; then
+  docker manifest push $REMOTE_TAG
+  docker manifest push $REMOTE_SHA_TAG
+fi
@@ -0,0 +1,8 @@
+#!/usr/bin/env bash
+set -e
+
+# Loop something forever until it fails, for verifying new tests
+
+while true; do
+  $@
+done
@@ -0,0 +1 @@
+707434c540e685bbe2886b3ff7c82fd61939d362
@@ -57,7 +57,7 @@ def generate_report(proposed, master, tmp, commit):
                     (lambda x: x.action.desiredCurvature, "desiredCurvature"),
                     (lambda x: x.leadsV3[0].x[0], "leadsV3.x"),
                     (lambda x: x.laneLines[1].y[0], "laneLines.y"),
-                     #(lambda x: x.meta.disengagePredictions.gasPressProbs[1], "gasPressProbs")
+                     (lambda x: x.meta.disengagePredictions.gasPressProbs[1], "gasPressProbs")
                    ], "modelV2")
  DriverStateV2_Plots = zl([
                     (lambda x: x.wheelOnRightProb, "wheelOnRightProb"),
@@ -0,0 +1,98 @@
+import os
+import numpy as np
+import hashlib
+
+import pyopencl as cl  # install with `PYOPENCL_CL_PRETEND_VERSION=2.0 pip install pyopencl`
+
+from openpilot.system.hardware import PC, TICI
+from openpilot.common.basedir import BASEDIR
+from openpilot.common.transformations.camera import DEVICE_CAMERAS
+from openpilot.system.camerad.snapshot.snapshot import yuv_to_rgb
+from openpilot.tools.lib.logreader import LogReader
+
+# TODO: check all sensors
+TEST_ROUTE = "8345e3b82948d454|2022-05-04--13-45-33/0"
+
+cam = DEVICE_CAMERAS[("tici", "ar0231")]
+FRAME_WIDTH, FRAME_HEIGHT = (cam.dcam.width, cam.dcam.height)
+FRAME_STRIDE = FRAME_WIDTH * 12 // 8 + 4
+
+UV_WIDTH = FRAME_WIDTH // 2
+UV_HEIGHT = FRAME_HEIGHT // 2
+UV_SIZE = UV_WIDTH * UV_HEIGHT
+
+
+def init_kernels(frame_offset=0):
+  ctx = cl.create_some_context(interactive=False)
+
+  with open(os.path.join(BASEDIR, 'system/camerad/cameras/process_raw.cl')) as f:
+    build_args = f' -cl-fast-relaxed-math -cl-denorms-are-zero -cl-single-precision-constant -I{BASEDIR}/system/camerad/sensors ' + \
+      f' -DFRAME_WIDTH={FRAME_WIDTH} -DFRAME_HEIGHT={FRAME_WIDTH} -DFRAME_STRIDE={FRAME_STRIDE} -DFRAME_OFFSET={frame_offset} ' + \
+      f' -DRGB_WIDTH={FRAME_WIDTH} -DRGB_HEIGHT={FRAME_HEIGHT} -DYUV_STRIDE={FRAME_WIDTH} -DUV_OFFSET={FRAME_WIDTH*FRAME_HEIGHT}' + \
+      ' -DSENSOR_ID=1 -DVIGNETTING=0 '
+    if PC:
+      build_args += ' -DHALF_AS_FLOAT=1 -cl-std=CL2.0'
+    imgproc_prg = cl.Program(ctx, f.read()).build(options=build_args)
+
+  return ctx, imgproc_prg
+
+def proc_frame(ctx, imgproc_prg, data, rgb=False):
+  q = cl.CommandQueue(ctx)
+
+  yuv_buff = np.empty(FRAME_WIDTH * FRAME_HEIGHT + UV_SIZE * 2, dtype=np.uint8)
+
+  cam_g = cl.Buffer(ctx, cl.mem_flags.READ_ONLY | cl.mem_flags.COPY_HOST_PTR, hostbuf=data)
+  yuv_g = cl.Buffer(ctx, cl.mem_flags.WRITE_ONLY, FRAME_WIDTH * FRAME_HEIGHT + UV_SIZE * 2)
+
+  krn = imgproc_prg.process_raw
+  krn.set_scalar_arg_dtypes([None, None, np.int32])
+  local_worksize = (20, 20) if TICI else (4, 4)
+
+  ev1 = krn(q, (FRAME_WIDTH//2, FRAME_HEIGHT//2), local_worksize, cam_g, yuv_g, 1)
+  cl.enqueue_copy(q, yuv_buff, yuv_g, wait_for=[ev1]).wait()
+  cl.enqueue_barrier(q)
+
+  y = yuv_buff[:FRAME_WIDTH*FRAME_HEIGHT].reshape((FRAME_HEIGHT, FRAME_WIDTH))
+  u = yuv_buff[FRAME_WIDTH*FRAME_HEIGHT::2].reshape((UV_HEIGHT, UV_WIDTH))
+  v = yuv_buff[FRAME_WIDTH*FRAME_HEIGHT+1::2].reshape((UV_HEIGHT, UV_WIDTH))
+
+  if rgb:
+    return yuv_to_rgb(y, u, v)
+  else:
+    return y, u, v
+
+
+def imgproc_replay(lr):
+  ctx, imgproc_prg = init_kernels()
+
+  frames = []
+  for m in lr:
+    if m.which() == 'roadCameraState':
+      cs = m.roadCameraState
+      if cs.image:
+        data = np.frombuffer(cs.image, dtype=np.uint8)
+        img = proc_frame(ctx, imgproc_prg, data)
+
+        frames.append(img)
+
+  return frames
+
+
+if __name__ == "__main__":
+  # load logs
+  lr = list(LogReader(TEST_ROUTE))
+  # run replay
+  out_frames = imgproc_replay(lr)
+
+  all_pix = np.concatenate([np.concatenate([d.flatten() for d in f]) for f in out_frames])
+  pix_hash = hashlib.sha1(all_pix).hexdigest()
+
+  with open('imgproc_replay_ref_hash') as f:
+    ref_hash = f.read()
+
+  if pix_hash != ref_hash:
+    print("result changed! please check kernel")
+    print(f"ref: {ref_hash}")
+    print(f"new: {pix_hash}")
+  else:
+    print("test passed")
@@ -0,0 +1,2 @@
+cachegrind.out.*
+*.prof
@@ -0,0 +1,91 @@
+from collections import defaultdict
+from cereal.services import SERVICE_LIST
+import cereal.messaging as messaging
+import capnp
+
+
+class ReplayDone(Exception):
+  pass
+
+
+class SubSocket:
+  def __init__(self, msgs, trigger):
+    self.i = 0
+    self.trigger = trigger
+    self.msgs = [m.as_builder().to_bytes() for m in msgs if m.which() == trigger]
+    self.max_i = len(self.msgs) - 1
+
+  def receive(self, non_blocking=False):
+    if non_blocking:
+      return None
+
+    if self.i == self.max_i:
+      raise ReplayDone
+
+    while True:
+      msg = self.msgs[self.i]
+      self.i += 1
+      return msg
+
+
+class PubSocket:
+  def send(self, data):
+    pass
+
+
+class SubMaster(messaging.SubMaster):
+  def __init__(self, msgs, trigger, services, check_averag_freq=False):
+    self.frame = 0
+    self.data = {}
+    self.ignore_alive = []
+
+    self.alive = {s: True for s in services}
+    self.updated = {s: False for s in services}
+    self.rcv_time = {s: 0. for s in services}
+    self.rcv_frame = {s: 0 for s in services}
+    self.valid = {s: True for s in services}
+    self.freq_ok = {s: True for s in services}
+    self.freq_tracker = {s: messaging.FrequencyTracker(SERVICE_LIST[s].frequency, SERVICE_LIST[s].frequency, False) for s in services}
+    self.logMonoTime = {}
+    self.sock = {}
+    self.freq = {}
+    self.check_average_freq = check_averag_freq
+    self.non_polled_services = []
+    self.ignore_average_freq = []
+
+    # TODO: specify multiple triggers for service like plannerd that poll on more than one service
+    cur_msgs = []
+    self.msgs = []
+    msgs = [m for m in msgs if m.which() in services]
+
+    for msg in msgs:
+      cur_msgs.append(msg)
+      if msg.which() == trigger:
+        self.msgs.append(cur_msgs)
+        cur_msgs = []
+
+    self.msgs = list(reversed(self.msgs))
+
+    for s in services:
+      self.freq[s] = SERVICE_LIST[s].frequency
+      try:
+        data = messaging.new_message(s)
+      except capnp.lib.capnp.KjException:
+        # lists
+        data = messaging.new_message(s, 0)
+
+      self.data[s] = getattr(data, s)
+      self.logMonoTime[s] = 0
+      self.sock[s] = SubSocket(msgs, s)
+
+  def update(self, timeout=None):
+    if not len(self.msgs):
+      raise ReplayDone
+
+    cur_msgs = self.msgs.pop()
+    self.update_msgs(cur_msgs[0].logMonoTime, self.msgs.pop())
+
+
+class PubMaster(messaging.PubMaster):
+  def __init__(self):
+    self.sock = defaultdict(PubSocket)
@@ -0,0 +1,97 @@
+#!/usr/bin/env python3
+import os
+import sys
+import cProfile
+import pprofile
+import pyprof2calltree
+
+from opendbc.car.toyota.values import CAR as TOYOTA
+from opendbc.car.honda.values import CAR as HONDA
+from opendbc.car.volkswagen.values import CAR as VW
+from openpilot.common.params import Params
+from openpilot.tools.lib.logreader import LogReader
+from openpilot.selfdrive.test.profiling.lib import SubMaster, PubMaster, SubSocket, ReplayDone
+from openpilot.selfdrive.test.process_replay.process_replay import CONFIGS
+
+BASE_URL = "https://commadataci.blob.core.windows.net/openpilotci/"
+
+CARS = {
+  'toyota': ("0982d79ebb0de295|2021-01-03--20-03-36/6", TOYOTA.TOYOTA_RAV4),
+  'honda': ("0982d79ebb0de295|2021-01-08--10-13-10/6", HONDA.HONDA_CIVIC),
+  "vw": ("ef895f46af5fd73f|2021-05-22--14-06-35/6", VW.AUDI_A3_MK3),
+}
+
+
+def get_inputs(msgs, process, fingerprint):
+  for config in CONFIGS:
+    if config.proc_name == process:
+      sub_socks = list(config.pubs)
+      trigger = sub_socks[0]
+      break
+
+  # some procs block on CarParams
+  for msg in msgs:
+    if msg.which() == 'carParams':
+      m = msg.as_builder()
+      m.carParams.carFingerprint = fingerprint
+      Params().put("CarParams", m.carParams.copy().to_bytes())
+      break
+
+  sm = SubMaster(msgs, trigger, sub_socks)
+  pm = PubMaster()
+  if 'can' in sub_socks:
+    can_sock = SubSocket(msgs, 'can')
+  else:
+    can_sock = None
+  return sm, pm, can_sock
+
+
+def profile(proc, func, car='toyota'):
+  segment, fingerprint = CARS[car]
+  segment = segment.replace('|', '/')
+  rlog_url = f"{BASE_URL}{segment}/rlog.bz2"
+  msgs = list(LogReader(rlog_url)) * int(os.getenv("LOOP", "1"))
+
+  os.environ['FINGERPRINT'] = fingerprint
+  os.environ['SKIP_FW_QUERY'] = "1"
+  os.environ['REPLAY'] = "1"
+
+  def run(sm, pm, can_sock):
+    try:
+      if can_sock is not None:
+        func(sm, pm, can_sock)
+      else:
+        func(sm, pm)
+    except ReplayDone:
+      pass
+
+  # Statistical
+  sm, pm, can_sock = get_inputs(msgs, proc, fingerprint)
+  with pprofile.StatisticalProfile()(period=0.00001) as pr:
+    run(sm, pm, can_sock)
+  pr.dump_stats(f'cachegrind.out.{proc}_statistical')
+
+  # Deterministic
+  sm, pm, can_sock = get_inputs(msgs, proc, fingerprint)
+  with cProfile.Profile() as pr:
+    run(sm, pm, can_sock)
+  pyprof2calltree.convert(pr.getstats(), f'cachegrind.out.{proc}_deterministic')
+
+
+if __name__ == '__main__':
+  from openpilot.selfdrive.controls.controlsd import main as controlsd_thread
+  from openpilot.selfdrive.locationd.paramsd import main as paramsd_thread
+  from openpilot.selfdrive.controls.plannerd import main as plannerd_thread
+
+  procs = {
+    'controlsd': controlsd_thread,
+    'paramsd': paramsd_thread,
+    'plannerd': plannerd_thread,
+  }
+
+  proc = sys.argv[1]
+  if proc not in procs:
+    print(f"{proc} not available")
+    sys.exit(0)
+  else:
+    profile(proc, procs[proc])
@@ -126,6 +126,7 @@ void SoftwarePanelSP::handleCurrentModelLblBtnClicked() {
    bundleNames.append(index_to_bundle[index]);
  }

+  currentModelLblBtn->setEnabled(!is_onroad);
  currentModelLblBtn->setValue(GetActiveModelName());

  const QString selectedBundleName = MultiOptionDialog::getSelection(
@@ -0,0 +1 @@
+SConscript(['modeld/SConscript'])
@@ -0,0 +1 @@
+*_pyx.cpp
@@ -0,0 +1,58 @@
+import glob
+
+Import('env', 'envCython', 'arch', 'cereal', 'messaging', 'common', 'gpucommon', 'visionipc', 'transformations')
+lenv = env.Clone()
+lenvCython = envCython.Clone()
+
+libs = [cereal, messaging, visionipc, gpucommon, common, 'capnp', 'kj', 'pthread']
+frameworks = []
+
+common_src = [
+  "models/commonmodel.cc",
+  "transforms/loadyuv.cc",
+  "transforms/transform.cc",
+]
+
+thneed_src_common = [
+  "thneed/thneed_common.cc",
+  "thneed/serialize.cc",
+]
+
+thneed_src_qcom = thneed_src_common + ["thneed/thneed_qcom2.cc"]
+thneed_src_pc = thneed_src_common + ["thneed/thneed_pc.cc"]
+thneed_src = thneed_src_qcom if arch == "larch64" else thneed_src_pc
+
+# SNPE except on Mac and ARM Linux
+snpe_lib = []
+if arch != "Darwin" and arch != "aarch64":
+  common_src += ['runners/snpemodel.cc']
+  snpe_lib += ['SNPE']
+
+# OpenCL is a framework on Mac
+if arch == "Darwin":
+  frameworks += ['OpenCL']
+else:
+  libs += ['OpenCL']
+
+# Set path definitions
+for pathdef, fn in {'TRANSFORM': 'transforms/transform.cl', 'LOADYUV': 'transforms/loadyuv.cl'}.items():
+  for xenv in (lenv, lenvCython):
+    xenv['CXXFLAGS'].append(f'-D{pathdef}_PATH=\\"{File(fn).abspath}\\"')
+
+# Compile cython
+snpe_rpath_qcom = "/data/pythonpath/third_party/snpe/larch64"
+snpe_rpath_pc = f"{Dir('#').abspath}/third_party/snpe/x86_64-linux-clang"
+snpe_rpath = lenvCython['RPATH'] + [snpe_rpath_qcom if arch == "larch64" else snpe_rpath_pc]
+
+cython_libs = envCython["LIBS"] + libs
+snpemodel_lib = lenv.Library('snpemodel', ['runners/snpemodel.cc'])
+commonmodel_lib = lenv.Library('commonmodel', common_src)
+
+lenvCython.Program('runners/runmodel_pyx.so', 'runners/runmodel_pyx.pyx', LIBS=cython_libs, FRAMEWORKS=frameworks)
+lenvCython.Program('runners/snpemodel_pyx.so', 'runners/snpemodel_pyx.pyx', LIBS=[snpemodel_lib, snpe_lib, *cython_libs], FRAMEWORKS=frameworks, RPATH=snpe_rpath)
+lenvCython.Program('models/commonmodel_pyx.so', 'models/commonmodel_pyx.pyx', LIBS=[commonmodel_lib, *cython_libs], FRAMEWORKS=frameworks)
+
+if arch == 'larch64' or GetOption('pc_thneed'):
+  thneed_lib = env.SharedLibrary('thneed', thneed_src, LIBS=[gpucommon, common, 'OpenCL', 'dl'])
+  thneedmodel_lib = env.Library('thneedmodel', ['runners/thneedmodel.cc'])
+  lenvCython.Program('runners/thneedmodel_pyx.so', 'runners/thneedmodel_pyx.pyx', LIBS=envCython["LIBS"]+[thneedmodel_lib, thneed_lib, gpucommon, common, 'dl', 'OpenCL'])
@@ -0,0 +1,86 @@
+import numpy as np
+
+def index_function(idx, max_val=192, max_idx=32):
+  return (max_val) * ((idx/max_idx)**2)
+
+class ModelConstants:
+  # time and distance indices
+  IDX_N = 33
+  T_IDXS = [index_function(idx, max_val=10.0) for idx in range(IDX_N)]
+  X_IDXS = [index_function(idx, max_val=192.0) for idx in range(IDX_N)]
+  LEAD_T_IDXS = [0., 2., 4., 6., 8., 10.]
+  LEAD_T_OFFSETS = [0., 2., 4.]
+  META_T_IDXS = [2., 4., 6., 8., 10.]
+
+  # model inputs constants
+  MODEL_FREQ = 20
+  FEATURE_LEN = 512
+  FULL_HISTORY_BUFFER_LEN = 99
+  HISTORY_BUFFER_LEN = 24
+  DESIRE_LEN = 8
+  TRAFFIC_CONVENTION_LEN = 2
+  LAT_PLANNER_STATE_LEN = 4
+  LATERAL_CONTROL_PARAMS_LEN = 2
+  PREV_DESIRED_CURV_LEN = 1
+
+  # model outputs constants
+  FCW_THRESHOLDS_5MS2 = np.array([.05, .05, .15, .15, .15], dtype=np.float32)
+  FCW_THRESHOLDS_3MS2 = np.array([.7, .7], dtype=np.float32)
+  FCW_5MS2_PROBS_WIDTH = 5
+  FCW_3MS2_PROBS_WIDTH = 2
+
+  DISENGAGE_WIDTH = 5
+  POSE_WIDTH = 6
+  WIDE_FROM_DEVICE_WIDTH = 3
+  SIM_POSE_WIDTH = 6
+  LEAD_WIDTH = 4
+  LANE_LINES_WIDTH = 2
+  ROAD_EDGES_WIDTH = 2
+  PLAN_WIDTH = 15
+  DESIRE_PRED_WIDTH = 8
+  LAT_PLANNER_SOLUTION_WIDTH = 4
+  DESIRED_CURV_WIDTH = 1
+
+  NUM_LANE_LINES = 4
+  NUM_ROAD_EDGES = 2
+
+  LEAD_TRAJ_LEN = 6
+  DESIRE_PRED_LEN = 4
+
+  PLAN_MHP_N = 5
+  LEAD_MHP_N = 2
+  PLAN_MHP_SELECTION = 1
+  LEAD_MHP_SELECTION = 3
+
+  FCW_THRESHOLD_5MS2_HIGH = 0.15
+  FCW_THRESHOLD_5MS2_LOW = 0.05
+  FCW_THRESHOLD_3MS2 = 0.7
+
+  CONFIDENCE_BUFFER_LEN = 5
+  RYG_GREEN = 0.01165
+  RYG_YELLOW = 0.06157
+
+  POLY_PATH_DEGREE = 4
+
+# model outputs slices
+class Plan:
+  POSITION = slice(0, 3)
+  VELOCITY = slice(3, 6)
+  ACCELERATION = slice(6, 9)
+  T_FROM_CURRENT_EULER = slice(9, 12)
+  ORIENTATION_RATE = slice(12, 15)
+
+class Meta:
+  ENGAGED = slice(0, 1)
+  # next 2, 4, 6, 8, 10 seconds
+  GAS_DISENGAGE = slice(1, 31, 6)
+  BRAKE_DISENGAGE = slice(2, 31, 6)
+  STEER_OVERRIDE = slice(3, 31, 6)
+  HARD_BRAKE_3 = slice(4, 31, 6)
+  HARD_BRAKE_4 = slice(5, 31, 6)
+  HARD_BRAKE_5 = slice(6, 31, 6)
+  # next 0, 2, 4, 6, 8, 10 seconds
+  GAS_PRESS = slice(31, 55, 4)
+  BRAKE_PRESS = slice(32, 55, 4)
+  LEFT_BLINKER = slice(33, 55, 4)
+  RIGHT_BLINKER = slice(34, 55, 4)
@@ -0,0 +1,237 @@
+import os
+import capnp
+import numpy as np
+from cereal import log
+from openpilot.sunnypilot.modeld.constants import ModelConstants, Plan, Meta
+from openpilot.selfdrive.controls.lib.drive_helpers import MIN_SPEED
+
+SEND_RAW_PRED = os.getenv('SEND_RAW_PRED')
+
+ConfidenceClass = log.ModelDataV2.ConfidenceClass
+
+def curv_from_psis(psi_target, psi_rate, vego, delay):
+  vego = np.clip(vego, MIN_SPEED, np.inf)
+  curv_from_psi = psi_target / (vego * delay)  # epsilon to prevent divide-by-zero
+  return 2*curv_from_psi - psi_rate / vego
+
+def get_curvature_from_plan(plan, vego, delay):
+  psi_target = np.interp(delay, ModelConstants.T_IDXS, plan[:, Plan.T_FROM_CURRENT_EULER][:, 2])
+  psi_rate = plan[:, Plan.ORIENTATION_RATE][0, 2]
+  return curv_from_psis(psi_target, psi_rate, vego, delay)
+
+class PublishState:
+  def __init__(self):
+    self.disengage_buffer = np.zeros(ModelConstants.CONFIDENCE_BUFFER_LEN*ModelConstants.DISENGAGE_WIDTH, dtype=np.float32)
+    self.prev_brake_5ms2_probs = np.zeros(ModelConstants.FCW_5MS2_PROBS_WIDTH, dtype=np.float32)
+    self.prev_brake_3ms2_probs = np.zeros(ModelConstants.FCW_3MS2_PROBS_WIDTH, dtype=np.float32)
+
+def fill_xyzt(builder, t, x, y, z, x_std=None, y_std=None, z_std=None):
+  builder.t = t
+  builder.x = x.tolist()
+  builder.y = y.tolist()
+  builder.z = z.tolist()
+  if x_std is not None:
+    builder.xStd = x_std.tolist()
+  if y_std is not None:
+    builder.yStd = y_std.tolist()
+  if z_std is not None:
+    builder.zStd = z_std.tolist()
+
+def fill_xyvat(builder, t, x, y, v, a, x_std=None, y_std=None, v_std=None, a_std=None):
+  builder.t = t
+  builder.x = x.tolist()
+  builder.y = y.tolist()
+  builder.v = v.tolist()
+  builder.a = a.tolist()
+  if x_std is not None:
+    builder.xStd = x_std.tolist()
+  if y_std is not None:
+    builder.yStd = y_std.tolist()
+  if v_std is not None:
+    builder.vStd = v_std.tolist()
+  if a_std is not None:
+    builder.aStd = a_std.tolist()
+
+def fill_xyz_poly(builder, degree, x, y, z):
+  xyz = np.stack([x, y, z], axis=1)
+  coeffs = np.polynomial.polynomial.polyfit(ModelConstants.T_IDXS, xyz, deg=degree)
+  builder.xCoefficients = coeffs[:, 0].tolist()
+  builder.yCoefficients = coeffs[:, 1].tolist()
+  builder.zCoefficients = coeffs[:, 2].tolist()
+
+def fill_lane_line_meta(builder, lane_lines, lane_line_probs):
+  builder.leftY = lane_lines[1].y[0]
+  builder.leftProb = lane_line_probs[1]
+  builder.rightY = lane_lines[2].y[0]
+  builder.rightProb = lane_line_probs[2]
+
+def fill_model_msg(base_msg: capnp._DynamicStructBuilder, extended_msg: capnp._DynamicStructBuilder,
+                   net_output_data: dict[str, np.ndarray], v_ego: float, delay: float,
+                   publish_state: PublishState, vipc_frame_id: int, vipc_frame_id_extra: int,
+                   frame_id: int, frame_drop: float, timestamp_eof: int, model_execution_time: float,
+                   valid: bool) -> None:
+  frame_age = frame_id - vipc_frame_id if frame_id > vipc_frame_id else 0
+  frame_drop_perc = frame_drop * 100
+  extended_msg.valid = valid
+  base_msg.valid = valid
+
+  desired_curv = float(get_curvature_from_plan(net_output_data['plan'][0], v_ego, delay))
+
+  driving_model_data = base_msg.drivingModelData
+
+  driving_model_data.frameId = vipc_frame_id
+  driving_model_data.frameIdExtra = vipc_frame_id_extra
+  driving_model_data.frameDropPerc = frame_drop_perc
+  driving_model_data.modelExecutionTime = model_execution_time
+
+  action = driving_model_data.action
+  action.desiredCurvature = desired_curv
+
+  modelV2 = extended_msg.modelV2
+  modelV2.frameId = vipc_frame_id
+  modelV2.frameIdExtra = vipc_frame_id_extra
+  modelV2.frameAge = frame_age
+  modelV2.frameDropPerc = frame_drop_perc
+  modelV2.timestampEof = timestamp_eof
+  modelV2.modelExecutionTime = model_execution_time
+
+  # plan
+  position = modelV2.position
+  fill_xyzt(position, ModelConstants.T_IDXS, *net_output_data['plan'][0,:,Plan.POSITION].T, *net_output_data['plan_stds'][0,:,Plan.POSITION].T)
+  velocity = modelV2.velocity
+  fill_xyzt(velocity, ModelConstants.T_IDXS, *net_output_data['plan'][0,:,Plan.VELOCITY].T)
+  acceleration = modelV2.acceleration
+  fill_xyzt(acceleration, ModelConstants.T_IDXS, *net_output_data['plan'][0,:,Plan.ACCELERATION].T)
+  orientation = modelV2.orientation
+  fill_xyzt(orientation, ModelConstants.T_IDXS, *net_output_data['plan'][0,:,Plan.T_FROM_CURRENT_EULER].T)
+  orientation_rate = modelV2.orientationRate
+  fill_xyzt(orientation_rate, ModelConstants.T_IDXS, *net_output_data['plan'][0,:,Plan.ORIENTATION_RATE].T)
+
+  # temporal pose
+  temporal_pose = modelV2.temporalPose
+  temporal_pose.trans = net_output_data['plan'][0,0,Plan.VELOCITY].tolist()
+  temporal_pose.transStd = net_output_data['plan_stds'][0,0,Plan.VELOCITY].tolist()
+  temporal_pose.rot = net_output_data['plan'][0,0,Plan.ORIENTATION_RATE].tolist()
+  temporal_pose.rotStd = net_output_data['plan_stds'][0,0,Plan.ORIENTATION_RATE].tolist()
+
+  # poly path
+  poly_path = driving_model_data.path
+  fill_xyz_poly(poly_path, ModelConstants.POLY_PATH_DEGREE, *net_output_data['plan'][0,:,Plan.POSITION].T)
+
+  # lateral planning
+  action = modelV2.action
+  action.desiredCurvature = desired_curv
+
+  # times at X_IDXS according to model plan
+  PLAN_T_IDXS = [np.nan] * ModelConstants.IDX_N
+  PLAN_T_IDXS[0] = 0.0
+  plan_x = net_output_data['plan'][0,:,Plan.POSITION][:,0].tolist()
+  for xidx in range(1, ModelConstants.IDX_N):
+    tidx = 0
+    # increment tidx until we find an element that's further away than the current xidx
+    while tidx < ModelConstants.IDX_N - 1 and plan_x[tidx+1] < ModelConstants.X_IDXS[xidx]:
+      tidx += 1
+    if tidx == ModelConstants.IDX_N - 1:
+      # if the Plan doesn't extend far enough, set plan_t to the max value (10s), then break
+      PLAN_T_IDXS[xidx] = ModelConstants.T_IDXS[ModelConstants.IDX_N - 1]
+      break
+    # interpolate to find `t` for the current xidx
+    current_x_val = plan_x[tidx]
+    next_x_val = plan_x[tidx+1]
+    p = (ModelConstants.X_IDXS[xidx] - current_x_val) / (next_x_val - current_x_val) if abs(next_x_val - current_x_val) > 1e-9 else float('nan')
+    PLAN_T_IDXS[xidx] = p * ModelConstants.T_IDXS[tidx+1] + (1 - p) * ModelConstants.T_IDXS[tidx]
+
+  # lane lines
+  modelV2.init('laneLines', 4)
+  for i in range(4):
+    lane_line = modelV2.laneLines[i]
+    fill_xyzt(lane_line, PLAN_T_IDXS, np.array(ModelConstants.X_IDXS), net_output_data['lane_lines'][0,i,:,0], net_output_data['lane_lines'][0,i,:,1])
+  modelV2.laneLineStds = net_output_data['lane_lines_stds'][0,:,0,0].tolist()
+  modelV2.laneLineProbs = net_output_data['lane_lines_prob'][0,1::2].tolist()
+
+  lane_line_meta = driving_model_data.laneLineMeta
+  fill_lane_line_meta(lane_line_meta, modelV2.laneLines, modelV2.laneLineProbs)
+
+  # road edges
+  modelV2.init('roadEdges', 2)
+  for i in range(2):
+    road_edge = modelV2.roadEdges[i]
+    fill_xyzt(road_edge, PLAN_T_IDXS, np.array(ModelConstants.X_IDXS), net_output_data['road_edges'][0,i,:,0], net_output_data['road_edges'][0,i,:,1])
+  modelV2.roadEdgeStds = net_output_data['road_edges_stds'][0,:,0,0].tolist()
+
+  # leads
+  modelV2.init('leadsV3', 3)
+  for i in range(3):
+    lead = modelV2.leadsV3[i]
+    fill_xyvat(lead, ModelConstants.LEAD_T_IDXS, *net_output_data['lead'][0,i].T, *net_output_data['lead_stds'][0,i].T)
+    lead.prob = net_output_data['lead_prob'][0,i].tolist()
+    lead.probTime = ModelConstants.LEAD_T_OFFSETS[i]
+
+  # meta
+  meta = modelV2.meta
+  meta.desireState = net_output_data['desire_state'][0].reshape(-1).tolist()
+  meta.desirePrediction = net_output_data['desire_pred'][0].reshape(-1).tolist()
+  meta.engagedProb = net_output_data['meta'][0,Meta.ENGAGED].item()
+  meta.init('disengagePredictions')
+  disengage_predictions = meta.disengagePredictions
+  disengage_predictions.t = ModelConstants.META_T_IDXS
+  disengage_predictions.brakeDisengageProbs = net_output_data['meta'][0,Meta.BRAKE_DISENGAGE].tolist()
+  disengage_predictions.gasDisengageProbs = net_output_data['meta'][0,Meta.GAS_DISENGAGE].tolist()
+  disengage_predictions.steerOverrideProbs = net_output_data['meta'][0,Meta.STEER_OVERRIDE].tolist()
+  disengage_predictions.brake3MetersPerSecondSquaredProbs = net_output_data['meta'][0,Meta.HARD_BRAKE_3].tolist()
+  disengage_predictions.brake4MetersPerSecondSquaredProbs = net_output_data['meta'][0,Meta.HARD_BRAKE_4].tolist()
+  disengage_predictions.brake5MetersPerSecondSquaredProbs = net_output_data['meta'][0,Meta.HARD_BRAKE_5].tolist()
+  disengage_predictions.gasPressProbs = net_output_data['meta'][0,Meta.GAS_PRESS].tolist()
+  disengage_predictions.brakePressProbs = net_output_data['meta'][0,Meta.BRAKE_PRESS].tolist()
+
+  publish_state.prev_brake_5ms2_probs[:-1] = publish_state.prev_brake_5ms2_probs[1:]
+  publish_state.prev_brake_5ms2_probs[-1] = net_output_data['meta'][0,Meta.HARD_BRAKE_5][0]
+  publish_state.prev_brake_3ms2_probs[:-1] = publish_state.prev_brake_3ms2_probs[1:]
+  publish_state.prev_brake_3ms2_probs[-1] = net_output_data['meta'][0,Meta.HARD_BRAKE_3][0]
+  hard_brake_predicted = (publish_state.prev_brake_5ms2_probs > ModelConstants.FCW_THRESHOLDS_5MS2).all() and \
+    (publish_state.prev_brake_3ms2_probs > ModelConstants.FCW_THRESHOLDS_3MS2).all()
+  meta.hardBrakePredicted = hard_brake_predicted.item()
+
+  # confidence
+  if vipc_frame_id % (2*ModelConstants.MODEL_FREQ) == 0:
+    # any disengage prob
+    brake_disengage_probs = net_output_data['meta'][0,Meta.BRAKE_DISENGAGE]
+    gas_disengage_probs = net_output_data['meta'][0,Meta.GAS_DISENGAGE]
+    steer_override_probs = net_output_data['meta'][0,Meta.STEER_OVERRIDE]
+    any_disengage_probs = 1-((1-brake_disengage_probs)*(1-gas_disengage_probs)*(1-steer_override_probs))
+    # independent disengage prob for each 2s slice
+    ind_disengage_probs = np.r_[any_disengage_probs[0], np.diff(any_disengage_probs) / (1 - any_disengage_probs[:-1])]
+    # rolling buf for 2, 4, 6, 8, 10s
+    publish_state.disengage_buffer[:-ModelConstants.DISENGAGE_WIDTH] = publish_state.disengage_buffer[ModelConstants.DISENGAGE_WIDTH:]
+    publish_state.disengage_buffer[-ModelConstants.DISENGAGE_WIDTH:] = ind_disengage_probs
+
+  score = 0.
+  for i in range(ModelConstants.DISENGAGE_WIDTH):
+    score += publish_state.disengage_buffer[i*ModelConstants.DISENGAGE_WIDTH+ModelConstants.DISENGAGE_WIDTH-1-i].item() / ModelConstants.DISENGAGE_WIDTH
+  if score < ModelConstants.RYG_GREEN:
+    modelV2.confidence = ConfidenceClass.green
+  elif score < ModelConstants.RYG_YELLOW:
+    modelV2.confidence = ConfidenceClass.yellow
+  else:
+    modelV2.confidence = ConfidenceClass.red
+
+  # raw prediction if enabled
+  if SEND_RAW_PRED:
+    modelV2.rawPredictions = net_output_data['raw_pred'].tobytes()
+
+def fill_pose_msg(msg: capnp._DynamicStructBuilder, net_output_data: dict[str, np.ndarray],
+                  vipc_frame_id: int, vipc_dropped_frames: int, timestamp_eof: int, live_calib_seen: bool) -> None:
+  msg.valid = live_calib_seen & (vipc_dropped_frames < 1)
+  cameraOdometry = msg.cameraOdometry
+
+  cameraOdometry.frameId = vipc_frame_id
+  cameraOdometry.timestampEof = timestamp_eof
+
+  cameraOdometry.trans = net_output_data['pose'][0,:3].tolist()
+  cameraOdometry.rot = net_output_data['pose'][0,3:].tolist()
+  cameraOdometry.wideFromDeviceEuler = net_output_data['wide_from_device_euler'][0,:].tolist()
+  cameraOdometry.roadTransformTrans = net_output_data['road_transform'][0,:3].tolist()
+  cameraOdometry.transStd = net_output_data['pose_stds'][0,:3].tolist()
+  cameraOdometry.rotStd = net_output_data['pose_stds'][0,3:].tolist()
+  cameraOdometry.wideFromDeviceEulerStd = net_output_data['wide_from_device_euler_stds'][0,:].tolist()
+  cameraOdometry.roadTransformTransStd = net_output_data['road_transform_stds'][0,:3].tolist()
@@ -0,0 +1,28 @@
+#!/usr/bin/env python3
+import sys
+import pathlib
+import onnx
+import codecs
+import pickle
+
+def get_name_and_shape(value_info:onnx.ValueInfoProto) -> tuple[str, tuple[int,...]]:
+  shape = tuple([int(dim.dim_value) for dim in value_info.type.tensor_type.shape.dim])
+  name = value_info.name
+  return name, shape
+
+if __name__ == "__main__":
+  model_path = pathlib.Path(sys.argv[1])
+  model = onnx.load(str(model_path))
+  i = [x.key for x in model.metadata_props].index('output_slices')
+  output_slices = model.metadata_props[i].value
+
+  metadata = {}
+  metadata['output_slices'] = pickle.loads(codecs.decode(output_slices.encode(), "base64"))
+  metadata['input_shapes'] = dict([get_name_and_shape(x) for x in model.graph.input])
+  metadata['output_shapes'] = dict([get_name_and_shape(x) for x in model.graph.output])
+
+  metadata_path = model_path.parent / (model_path.stem + '_metadata.pkl')
+  with open(metadata_path, 'wb') as f:
+    pickle.dump(metadata, f)
+
+  print(f'saved metadata to {metadata_path}')
@@ -0,0 +1,10 @@
+#!/usr/bin/env bash
+
+DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" >/dev/null && pwd)"
+cd "$DIR/../../"
+
+if [ -f "$DIR/libthneed.so" ]; then
+  export LD_PRELOAD="$DIR/libthneed.so"
+fi
+
+exec "$DIR/modeld.py" "$@"
@@ -0,0 +1,299 @@
+#!/usr/bin/env python3
+import os
+import time
+import pickle
+import numpy as np
+import cereal.messaging as messaging
+from cereal import car, log
+from pathlib import Path
+from setproctitle import setproctitle
+from cereal.messaging import PubMaster, SubMaster
+from msgq.visionipc import VisionIpcClient, VisionStreamType, VisionBuf
+from opendbc.car.car_helpers import get_demo_car_params
+from openpilot.common.swaglog import cloudlog
+from openpilot.common.params import Params
+from openpilot.common.filter_simple import FirstOrderFilter
+from openpilot.common.realtime import config_realtime_process
+from openpilot.common.transformations.camera import DEVICE_CAMERAS
+from openpilot.common.transformations.model import get_warp_matrix
+from openpilot.system import sentry
+from openpilot.selfdrive.controls.lib.desire_helper import DesireHelper
+from openpilot.sunnypilot.modeld.runners import ModelRunner, Runtime
+from openpilot.sunnypilot.modeld.parse_model_outputs import Parser
+from openpilot.sunnypilot.modeld.fill_model_msg import fill_model_msg, fill_pose_msg, PublishState
+from openpilot.sunnypilot.modeld.constants import ModelConstants
+from openpilot.sunnypilot.modeld.models.commonmodel_pyx import DrivingModelFrame, CLContext
+
+PROCESS_NAME = "sunnypilot.modeld.modeld"
+SEND_RAW_PRED = os.getenv('SEND_RAW_PRED')
+
+MODEL_PATHS = {
+  ModelRunner.THNEED: Path(__file__).parent / 'models/supercombo.thneed',
+  ModelRunner.ONNX: Path(__file__).parent / 'models/supercombo.onnx'}
+
+METADATA_PATH = Path(__file__).parent / 'models/supercombo_metadata.pkl'
+
+
+class FrameMeta:
+  frame_id: int = 0
+  timestamp_sof: int = 0
+  timestamp_eof: int = 0
+
+  def __init__(self, vipc=None):
+    if vipc is not None:
+      self.frame_id, self.timestamp_sof, self.timestamp_eof = vipc.frame_id, vipc.timestamp_sof, vipc.timestamp_eof
+
+class ModelState:
+  frame: DrivingModelFrame
+  wide_frame: DrivingModelFrame
+  inputs: dict[str, np.ndarray]
+  output: np.ndarray
+  prev_desire: np.ndarray  # for tracking the rising edge of the pulse
+  model: ModelRunner
+
+  def __init__(self, context: CLContext):
+    self.frame = DrivingModelFrame(context)
+    self.wide_frame = DrivingModelFrame(context)
+    self.prev_desire = np.zeros(ModelConstants.DESIRE_LEN, dtype=np.float32)
+    self.full_features_20Hz = np.zeros((ModelConstants.FULL_HISTORY_BUFFER_LEN, ModelConstants.FEATURE_LEN), dtype=np.float32)
+    self.desire_20Hz =  np.zeros((ModelConstants.FULL_HISTORY_BUFFER_LEN + 1, ModelConstants.DESIRE_LEN), dtype=np.float32)
+
+    # img buffers are managed in openCL transform code
+    self.inputs = {
+      'desire': np.zeros(ModelConstants.DESIRE_LEN * (ModelConstants.HISTORY_BUFFER_LEN+1), dtype=np.float32),
+      'traffic_convention': np.zeros(ModelConstants.TRAFFIC_CONVENTION_LEN, dtype=np.float32),
+      'features_buffer': np.zeros(ModelConstants.HISTORY_BUFFER_LEN * ModelConstants.FEATURE_LEN, dtype=np.float32),
+    }
+
+    with open(METADATA_PATH, 'rb') as f:
+      model_metadata = pickle.load(f)
+
+    self.output_slices = model_metadata['output_slices']
+    net_output_size = model_metadata['output_shapes']['outputs'][1]
+    self.output = np.zeros(net_output_size, dtype=np.float32)
+    self.parser = Parser()
+
+    self.model = ModelRunner(MODEL_PATHS, self.output, Runtime.GPU, False, context)
+    self.model.addInput("input_imgs", None)
+    self.model.addInput("big_input_imgs", None)
+    for k,v in self.inputs.items():
+      self.model.addInput(k, v)
+
+  def slice_outputs(self, model_outputs: np.ndarray) -> dict[str, np.ndarray]:
+    parsed_model_outputs = {k: model_outputs[np.newaxis, v] for k,v in self.output_slices.items()}
+    if SEND_RAW_PRED:
+      parsed_model_outputs['raw_pred'] = model_outputs.copy()
+    return parsed_model_outputs
+
+  def run(self, buf: VisionBuf, wbuf: VisionBuf, transform: np.ndarray, transform_wide: np.ndarray,
+                inputs: dict[str, np.ndarray], prepare_only: bool) -> dict[str, np.ndarray] | None:
+    # Model decides when action is completed, so desire input is just a pulse triggered on rising edge
+    inputs['desire'][0] = 0
+    new_desire = np.where(inputs['desire'] - self.prev_desire > .99, inputs['desire'], 0)
+    self.prev_desire[:] = inputs['desire']
+
+    self.desire_20Hz[:-1] = self.desire_20Hz[1:]
+    self.desire_20Hz[-1] = new_desire
+    self.inputs['desire'][:] = self.desire_20Hz.reshape((25,4,-1)).max(axis=1).flatten()
+
+    self.inputs['traffic_convention'][:] = inputs['traffic_convention']
+
+    self.model.setInputBuffer("input_imgs", self.frame.prepare(buf, transform.flatten(), self.model.getCLBuffer("input_imgs")))
+    self.model.setInputBuffer("big_input_imgs", self.wide_frame.prepare(wbuf, transform_wide.flatten(), self.model.getCLBuffer("big_input_imgs")))
+
+    if prepare_only:
+      return None
+
+    self.model.execute()
+    outputs = self.parser.parse_outputs(self.slice_outputs(self.output))
+
+    self.full_features_20Hz[:-1] = self.full_features_20Hz[1:]
+    self.full_features_20Hz[-1] = outputs['hidden_state'][0, :]
+
+    idxs = np.arange(-4,-100,-4)[::-1]
+    self.inputs['features_buffer'][:] = self.full_features_20Hz[idxs].flatten()
+    return outputs
+
+
+def main(demo=False):
+  cloudlog.warning("modeld init")
+
+  sentry.set_tag("daemon", PROCESS_NAME)
+  cloudlog.bind(daemon=PROCESS_NAME)
+  setproctitle(PROCESS_NAME)
+  config_realtime_process(7, 54)
+
+  cloudlog.warning("setting up CL context")
+  cl_context = CLContext()
+  cloudlog.warning("CL context ready; loading model")
+  model = ModelState(cl_context)
+  cloudlog.warning("models loaded, modeld starting")
+
+  # visionipc clients
+  while True:
+    available_streams = VisionIpcClient.available_streams("camerad", block=False)
+    if available_streams:
+      use_extra_client = VisionStreamType.VISION_STREAM_WIDE_ROAD in available_streams and VisionStreamType.VISION_STREAM_ROAD in available_streams
+      main_wide_camera = VisionStreamType.VISION_STREAM_ROAD not in available_streams
+      break
+    time.sleep(.1)
+
+  vipc_client_main_stream = VisionStreamType.VISION_STREAM_WIDE_ROAD if main_wide_camera else VisionStreamType.VISION_STREAM_ROAD
+  vipc_client_main = VisionIpcClient("camerad", vipc_client_main_stream, True, cl_context)
+  vipc_client_extra = VisionIpcClient("camerad", VisionStreamType.VISION_STREAM_WIDE_ROAD, False, cl_context)
+  cloudlog.warning(f"vision stream set up, main_wide_camera: {main_wide_camera}, use_extra_client: {use_extra_client}")
+
+  while not vipc_client_main.connect(False):
+    time.sleep(0.1)
+  while use_extra_client and not vipc_client_extra.connect(False):
+    time.sleep(0.1)
+
+  cloudlog.warning(f"connected main cam with buffer size: {vipc_client_main.buffer_len} ({vipc_client_main.width} x {vipc_client_main.height})")
+  if use_extra_client:
+    cloudlog.warning(f"connected extra cam with buffer size: {vipc_client_extra.buffer_len} ({vipc_client_extra.width} x {vipc_client_extra.height})")
+
+  # messaging
+  pm = PubMaster(["modelV2", "drivingModelData", "cameraOdometry"])
+  sm = SubMaster(["deviceState", "carState", "roadCameraState", "liveCalibration", "driverMonitoringState", "carControl"])
+
+  publish_state = PublishState()
+  params = Params()
+
+  # setup filter to track dropped frames
+  frame_dropped_filter = FirstOrderFilter(0., 10., 1. / ModelConstants.MODEL_FREQ)
+  frame_id = 0
+  last_vipc_frame_id = 0
+  run_count = 0
+
+  model_transform_main = np.zeros((3, 3), dtype=np.float32)
+  model_transform_extra = np.zeros((3, 3), dtype=np.float32)
+  live_calib_seen = False
+  buf_main, buf_extra = None, None
+  meta_main = FrameMeta()
+  meta_extra = FrameMeta()
+
+
+  if demo:
+    CP = get_demo_car_params()
+  else:
+    CP = messaging.log_from_bytes(params.get("CarParams", block=True), car.CarParams)
+  cloudlog.info("modeld got CarParams: %s", CP.carName)
+
+  # TODO this needs more thought, use .2s extra for now to estimate other delays
+  steer_delay = CP.steerActuatorDelay + .2
+
+  DH = DesireHelper()
+
+  while True:
+    # Keep receiving frames until we are at least 1 frame ahead of previous extra frame
+    while meta_main.timestamp_sof < meta_extra.timestamp_sof + 25000000:
+      buf_main = vipc_client_main.recv()
+      meta_main = FrameMeta(vipc_client_main)
+      if buf_main is None:
+        break
+
+    if buf_main is None:
+      cloudlog.debug("vipc_client_main no frame")
+      continue
+
+    if use_extra_client:
+      # Keep receiving extra frames until frame id matches main camera
+      while True:
+        buf_extra = vipc_client_extra.recv()
+        meta_extra = FrameMeta(vipc_client_extra)
+        if buf_extra is None or meta_main.timestamp_sof < meta_extra.timestamp_sof + 25000000:
+          break
+
+      if buf_extra is None:
+        cloudlog.debug("vipc_client_extra no frame")
+        continue
+
+      if abs(meta_main.timestamp_sof - meta_extra.timestamp_sof) > 10000000:
+        cloudlog.error(f"frames out of sync! main: {meta_main.frame_id} ({meta_main.timestamp_sof / 1e9:.5f}),\
+                         extra: {meta_extra.frame_id} ({meta_extra.timestamp_sof / 1e9:.5f})")
+
+    else:
+      # Use single camera
+      buf_extra = buf_main
+      meta_extra = meta_main
+
+    sm.update(0)
+    desire = DH.desire
+    is_rhd = sm["driverMonitoringState"].isRHD
+    frame_id = sm["roadCameraState"].frameId
+    v_ego = max(sm["carState"].vEgo, 0.)
+    if sm.updated["liveCalibration"] and sm.seen['roadCameraState'] and sm.seen['deviceState']:
+      device_from_calib_euler = np.array(sm["liveCalibration"].rpyCalib, dtype=np.float32)
+      dc = DEVICE_CAMERAS[(str(sm['deviceState'].deviceType), str(sm['roadCameraState'].sensor))]
+      model_transform_main = get_warp_matrix(device_from_calib_euler, dc.ecam.intrinsics if main_wide_camera else dc.fcam.intrinsics, False).astype(np.float32)
+      model_transform_extra = get_warp_matrix(device_from_calib_euler, dc.ecam.intrinsics, True).astype(np.float32)
+      live_calib_seen = True
+
+    traffic_convention = np.zeros(2)
+    traffic_convention[int(is_rhd)] = 1
+
+    vec_desire = np.zeros(ModelConstants.DESIRE_LEN, dtype=np.float32)
+    if desire >= 0 and desire < ModelConstants.DESIRE_LEN:
+      vec_desire[desire] = 1
+
+    # tracked dropped frames
+    vipc_dropped_frames = max(0, meta_main.frame_id - last_vipc_frame_id - 1)
+    frames_dropped = frame_dropped_filter.update(min(vipc_dropped_frames, 10))
+    if run_count < 10: # let frame drops warm up
+      frame_dropped_filter.x = 0.
+      frames_dropped = 0.
+    run_count = run_count + 1
+
+    frame_drop_ratio = frames_dropped / (1 + frames_dropped)
+    prepare_only = vipc_dropped_frames > 0
+    if prepare_only:
+      cloudlog.error(f"skipping model eval. Dropped {vipc_dropped_frames} frames")
+
+    inputs:dict[str, np.ndarray] = {
+      'desire': vec_desire,
+      'traffic_convention': traffic_convention,
+      }
+
+    mt1 = time.perf_counter()
+    model_output = model.run(buf_main, buf_extra, model_transform_main, model_transform_extra, inputs, prepare_only)
+    mt2 = time.perf_counter()
+    model_execution_time = mt2 - mt1
+
+    if model_output is not None:
+      modelv2_send = messaging.new_message('modelV2')
+      drivingdata_send = messaging.new_message('drivingModelData')
+      posenet_send = messaging.new_message('cameraOdometry')
+      fill_model_msg(drivingdata_send, modelv2_send, model_output, v_ego, steer_delay,
+                     publish_state, meta_main.frame_id, meta_extra.frame_id, frame_id,
+                     frame_drop_ratio, meta_main.timestamp_eof, model_execution_time, live_calib_seen)
+
+      desire_state = modelv2_send.modelV2.meta.desireState
+      l_lane_change_prob = desire_state[log.Desire.laneChangeLeft]
+      r_lane_change_prob = desire_state[log.Desire.laneChangeRight]
+      lane_change_prob = l_lane_change_prob + r_lane_change_prob
+      DH.update(sm['carState'], sm['carControl'].latActive, lane_change_prob)
+      modelv2_send.modelV2.meta.laneChangeState = DH.lane_change_state
+      modelv2_send.modelV2.meta.laneChangeDirection = DH.lane_change_direction
+      drivingdata_send.drivingModelData.meta.laneChangeState = DH.lane_change_state
+      drivingdata_send.drivingModelData.meta.laneChangeDirection = DH.lane_change_direction
+
+      fill_pose_msg(posenet_send, model_output, meta_main.frame_id, vipc_dropped_frames, meta_main.timestamp_eof, live_calib_seen)
+      pm.send('modelV2', modelv2_send)
+      pm.send('drivingModelData', drivingdata_send)
+      pm.send('cameraOdometry', posenet_send)
+
+    last_vipc_frame_id = meta_main.frame_id
+
+
+if __name__ == "__main__":
+  try:
+    import argparse
+    parser = argparse.ArgumentParser()
+    parser.add_argument('--demo', action='store_true', help='A boolean for demo mode.')
+    args = parser.parse_args()
+    main(demo=args.demo)
+  except KeyboardInterrupt:
+    cloudlog.warning(f"child {PROCESS_NAME} got SIGINT")
+  except Exception:
+    sentry.capture_exception()
+    raise
@@ -0,0 +1,62 @@
+## Neural networks in openpilot
+To view the architecture of the ONNX networks, you can use [netron](https://netron.app/)
+
+## Supercombo
+### Supercombo input format (Full size: 799906 x float32)
+* **image stream**
+  * Two consecutive images (256 * 512 * 3 in RGB) recorded at 20 Hz : 393216 = 2 * 6 * 128 * 256
+    * Each 256 * 512 image is represented in YUV420 with 6 channels : 6 * 128 * 256
+      * Channels 0,1,2,3 represent the full-res Y channel and are represented in numpy as Y[::2, ::2], Y[::2, 1::2], Y[1::2, ::2], and Y[1::2, 1::2]
+      * Channel 4 represents the half-res U channel
+      * Channel 5 represents the half-res V channel
+* **wide image stream**
+  * Two consecutive images (256 * 512 * 3 in RGB) recorded at 20 Hz : 393216 = 2 * 6 * 128 * 256
+    * Each 256 * 512 image is represented in YUV420 with 6 channels : 6 * 128 * 256
+      * Channels 0,1,2,3 represent the full-res Y channel and are represented in numpy as Y[::2, ::2], Y[::2, 1::2], Y[1::2, ::2], and Y[1::2, 1::2]
+      * Channel 4 represents the half-res U channel
+      * Channel 5 represents the half-res V channel
+* **desire**
+  * one-hot encoded buffer to command model to execute certain actions, bit needs to be sent for the past 5 seconds (at 20FPS) : 100 * 8
+* **traffic convention**
+  * one-hot encoded vector to tell model whether traffic is right-hand or left-hand traffic : 2
+* **feature buffer**
+  * A buffer of intermediate features that gets appended to the current feature to form a 5 seconds temporal context (at 20FPS) : 99 * 512
+
+
+### Supercombo output format (Full size: XXX x float32)
+Read [here](https://github.com/commaai/openpilot/blob/90af436a121164a51da9fa48d093c29f738adf6a/selfdrive/modeld/models/driving.h#L236) for more.
+
+
+## Driver Monitoring Model
+* .onnx model can be run with onnx runtimes
+* .dlc file is a pre-quantized model and only runs on qualcomm DSPs
+
+### input format
+* single image W = 1440 H = 960 luminance channel (Y) from the planar YUV420 format:
+  * full input size is 1440 * 960 = 1382400
+  * normalized ranging from 0.0 to 1.0 in float32 (onnx runner) or ranging from 0 to 255 in uint8 (snpe runner)
+* camera calibration angles (roll, pitch, yaw) from liveCalibration: 3 x float32 inputs
+
+### output format
+* 84 x float32 outputs = 2 + 41 * 2 ([parsing example](https://github.com/commaai/openpilot/blob/22ce4e17ba0d3bfcf37f8255a4dd1dc683fe0c38/selfdrive/modeld/models/dmonitoring.cc#L33))
+  * for each person in the front seats (2 * 41)
+    * face pose: 12 = 6 + 6
+      * face orientation [pitch, yaw, roll] in camera frame: 3
+      * face position [dx, dy] relative to image center: 2
+      * normalized face size: 1
+      * standard deviations for above outputs: 6
+    * face visible probability: 1
+    * eyes: 20 = (8 + 1) + (8 + 1) + 1 + 1
+      * eye position and size, and their standard deviations: 8
+      * eye visible probability: 1
+      * eye closed probability: 1
+    * wearing sunglasses probability: 1
+    * face occluded probability: 1
+    * touching wheel probability: 1
+    * paying attention probability: 1
+    * (deprecated) distracted probabilities: 2
+    * using phone probability: 1
+    * distracted probability: 1
+  * common outputs 2
+    * poor camera vision probability: 1
+    * left hand drive probability: 1
@@ -0,0 +1,69 @@
+#include "sunnypilot/modeld/models/commonmodel.h"
+
+#include <cmath>
+#include <cstring>
+
+#include "common/clutil.h"
+
+DrivingModelFrame::DrivingModelFrame(cl_device_id device_id, cl_context context) : ModelFrame(device_id, context) {
+  input_frames = std::make_unique<uint8_t[]>(buf_size);
+  //input_frames_cl = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_WRITE, buf_size, NULL, &err));
+  img_buffer_20hz_cl = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_WRITE, 5*frame_size_bytes, NULL, &err));
+  region.origin = 4 * frame_size_bytes;
+  region.size = frame_size_bytes;
+  last_img_cl = CL_CHECK_ERR(clCreateSubBuffer(img_buffer_20hz_cl, CL_MEM_READ_WRITE, CL_BUFFER_CREATE_TYPE_REGION, &region, &err));
+
+  loadyuv_init(&loadyuv, context, device_id, MODEL_WIDTH, MODEL_HEIGHT);
+  init_transform(device_id, context, MODEL_WIDTH, MODEL_HEIGHT);
+}
+
+uint8_t* DrivingModelFrame::prepare(cl_mem yuv_cl, int frame_width, int frame_height, int frame_stride, int frame_uv_offset, const mat3& projection, cl_mem* output) {
+  run_transform(yuv_cl, MODEL_WIDTH, MODEL_HEIGHT, frame_width, frame_height, frame_stride, frame_uv_offset, projection);
+
+  for (int i = 0; i < 4; i++) {
+    CL_CHECK(clEnqueueCopyBuffer(q, img_buffer_20hz_cl, img_buffer_20hz_cl, (i+1)*frame_size_bytes, i*frame_size_bytes, frame_size_bytes, 0, nullptr, nullptr));
+  }
+  loadyuv_queue(&loadyuv, q, y_cl, u_cl, v_cl, last_img_cl);
+
+  if (output == NULL) {
+    CL_CHECK(clEnqueueReadBuffer(q, img_buffer_20hz_cl, CL_TRUE, 0, frame_size_bytes, &input_frames[0], 0, nullptr, nullptr));
+    CL_CHECK(clEnqueueReadBuffer(q, last_img_cl, CL_TRUE, 0, frame_size_bytes, &input_frames[MODEL_FRAME_SIZE], 0, nullptr, nullptr));
+    clFinish(q);
+    return &input_frames[0];
+  } else {
+    copy_queue(&loadyuv, q, img_buffer_20hz_cl, *output, 0, 0, frame_size_bytes);
+    copy_queue(&loadyuv, q, last_img_cl, *output, 0, frame_size_bytes, frame_size_bytes);
+
+    // NOTE: Since thneed is using a different command queue, this clFinish is needed to ensure the image is ready.
+    clFinish(q);
+    return NULL;
+  }
+}
+
+DrivingModelFrame::~DrivingModelFrame() {
+  deinit_transform();
+  loadyuv_destroy(&loadyuv);
+  CL_CHECK(clReleaseMemObject(img_buffer_20hz_cl));
+  CL_CHECK(clReleaseMemObject(last_img_cl));
+  CL_CHECK(clReleaseCommandQueue(q));
+}
+
+
+MonitoringModelFrame::MonitoringModelFrame(cl_device_id device_id, cl_context context) : ModelFrame(device_id, context) {
+  input_frames = std::make_unique<uint8_t[]>(buf_size);
+  //input_frame_cl = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_WRITE, buf_size, NULL, &err));
+
+  init_transform(device_id, context, MODEL_WIDTH, MODEL_HEIGHT);
+}
+uint8_t* MonitoringModelFrame::prepare(cl_mem yuv_cl, int frame_width, int frame_height, int frame_stride, int frame_uv_offset, const mat3& projection, cl_mem* output) {
+  run_transform(yuv_cl, MODEL_WIDTH, MODEL_HEIGHT, frame_width, frame_height, frame_stride, frame_uv_offset, projection);
+  CL_CHECK(clEnqueueReadBuffer(q, y_cl, CL_TRUE, 0, MODEL_FRAME_SIZE * sizeof(uint8_t), input_frames.get(), 0, nullptr, nullptr));
+  clFinish(q);
+  //return &y_cl;
+  return input_frames.get();
+}
+
+MonitoringModelFrame::~MonitoringModelFrame() {
+  deinit_transform();
+  CL_CHECK(clReleaseCommandQueue(q));
+}
@@ -0,0 +1,98 @@
+#pragma once
+
+#include <cfloat>
+#include <cstdlib>
+#include <cassert>
+
+#include <memory>
+
+#define CL_USE_DEPRECATED_OPENCL_1_2_APIS
+#ifdef __APPLE__
+#include <OpenCL/cl.h>
+#else
+#include <CL/cl.h>
+#endif
+
+#include "common/mat.h"
+#include "selfdrive/modeld/transforms/loadyuv.h"
+#include "selfdrive/modeld/transforms/transform.h"
+
+class ModelFrame {
+public:
+  ModelFrame(cl_device_id device_id, cl_context context) {
+    q = CL_CHECK_ERR(clCreateCommandQueue(context, device_id, 0, &err));
+  }
+  virtual ~ModelFrame() {}
+  virtual uint8_t* prepare(cl_mem yuv_cl, int frame_width, int frame_height, int frame_stride, int frame_uv_offset, const mat3& projection, cl_mem* output) { return NULL; }
+  /*
+  uint8_t* buffer_from_cl(cl_mem *in_frames, int buffer_size) {
+    CL_CHECK(clEnqueueReadBuffer(q, *in_frames, CL_TRUE, 0, buffer_size, input_frames.get(), 0, nullptr, nullptr));
+    clFinish(q);
+    return &input_frames[0];
+  }
+  */
+
+  int MODEL_WIDTH;
+  int MODEL_HEIGHT;
+  int MODEL_FRAME_SIZE;
+  int buf_size;
+
+protected:
+  cl_mem y_cl, u_cl, v_cl;
+  Transform transform;
+  cl_command_queue q;
+  std::unique_ptr<uint8_t[]> input_frames;
+
+  void init_transform(cl_device_id device_id, cl_context context, int model_width, int model_height) {
+    y_cl = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_WRITE, model_width * model_height, NULL, &err));
+    u_cl = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_WRITE, (model_width / 2) * (model_height / 2), NULL, &err));
+    v_cl = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_WRITE, (model_width / 2) * (model_height / 2), NULL, &err));
+    transform_init(&transform, context, device_id);
+  }
+
+  void deinit_transform() {
+    transform_destroy(&transform);
+    CL_CHECK(clReleaseMemObject(v_cl));
+    CL_CHECK(clReleaseMemObject(u_cl));
+    CL_CHECK(clReleaseMemObject(y_cl));
+  }
+
+  void run_transform(cl_mem yuv_cl, int model_width, int model_height, int frame_width, int frame_height, int frame_stride, int frame_uv_offset, const mat3& projection) {
+    transform_queue(&transform, q,
+        yuv_cl, frame_width, frame_height, frame_stride, frame_uv_offset,
+        y_cl, u_cl, v_cl, model_width, model_height, projection);
+  }
+};
+
+class DrivingModelFrame : public ModelFrame {
+public:
+  DrivingModelFrame(cl_device_id device_id, cl_context context);
+  ~DrivingModelFrame();
+  uint8_t* prepare(cl_mem yuv_cl, int frame_width, int frame_height, int frame_stride, int frame_uv_offset, const mat3& projection, cl_mem* output);
+
+  const int MODEL_WIDTH = 512;
+  const int MODEL_HEIGHT = 256;
+  const int MODEL_FRAME_SIZE = MODEL_WIDTH * MODEL_HEIGHT * 3 / 2;
+  const int buf_size = MODEL_FRAME_SIZE * 2;
+  const size_t frame_size_bytes = MODEL_FRAME_SIZE * sizeof(uint8_t);
+
+private:
+  LoadYUVState loadyuv;
+  cl_mem img_buffer_20hz_cl, last_img_cl;//, input_frames_cl;
+  cl_buffer_region region;
+};
+
+class MonitoringModelFrame : public ModelFrame {
+public:
+  MonitoringModelFrame(cl_device_id device_id, cl_context context);
+  ~MonitoringModelFrame();
+  uint8_t* prepare(cl_mem yuv_cl, int frame_width, int frame_height, int frame_stride, int frame_uv_offset, const mat3& projection, cl_mem* output);
+
+  const int MODEL_WIDTH = 1440;
+  const int MODEL_HEIGHT = 960;
+  const int MODEL_FRAME_SIZE = MODEL_WIDTH * MODEL_HEIGHT;
+  const int buf_size = MODEL_FRAME_SIZE;
+
+private:
+  // cl_mem input_frame_cl;
+};
@@ -0,0 +1,26 @@
+# distutils: language = c++
+
+from msgq.visionipc.visionipc cimport cl_device_id, cl_context, cl_mem
+
+cdef extern from "common/mat.h":
+  cdef struct mat3:
+    float v[9]
+
+cdef extern from "common/clutil.h":
+  cdef unsigned long CL_DEVICE_TYPE_DEFAULT
+  cl_device_id cl_get_device_id(unsigned long)
+  cl_context cl_create_context(cl_device_id)
+
+cdef extern from "sunnypilot/modeld/models/commonmodel.h":
+  cppclass ModelFrame:
+    int buf_size
+    # unsigned char * buffer_from_cl(cl_mem*, int);
+    unsigned char * prepare(cl_mem, int, int, int, int, mat3, cl_mem*)
+
+  cppclass DrivingModelFrame:
+    int buf_size
+    DrivingModelFrame(cl_device_id, cl_context)
+
+  cppclass MonitoringModelFrame:
+    int buf_size
+    MonitoringModelFrame(cl_device_id, cl_context)
@@ -0,0 +1,13 @@
+# distutils: language = c++
+
+from msgq.visionipc.visionipc cimport cl_mem
+from msgq.visionipc.visionipc_pyx cimport CLContext as BaseCLContext
+
+cdef class CLContext(BaseCLContext):
+  pass
+
+cdef class CLMem:
+  cdef cl_mem * mem
+
+  @staticmethod
+  cdef create(void*)
@@ -0,0 +1,76 @@
+# distutils: language = c++
+# cython: c_string_encoding=ascii, language_level=3
+
+import numpy as np
+cimport numpy as cnp
+from libc.string cimport memcpy
+from libc.stdint cimport uintptr_t
+
+from msgq.visionipc.visionipc cimport cl_mem
+from msgq.visionipc.visionipc_pyx cimport VisionBuf, CLContext as BaseCLContext
+from sunnypilot.modeld.models.commonmodel cimport CL_DEVICE_TYPE_DEFAULT, cl_get_device_id, cl_create_context
+from sunnypilot.modeld.models.commonmodel cimport mat3, ModelFrame as cppModelFrame, DrivingModelFrame as cppDrivingModelFrame, MonitoringModelFrame as cppMonitoringModelFrame
+
+
+cdef class CLContext(BaseCLContext):
+  def __cinit__(self):
+    self.device_id = cl_get_device_id(CL_DEVICE_TYPE_DEFAULT)
+    self.context = cl_create_context(self.device_id)
+
+cdef class CLMem:
+  @staticmethod
+  cdef create(void * cmem):
+    mem = CLMem()
+    mem.mem = <cl_mem*> cmem
+    return mem
+
+  @property
+  def mem_address(self):
+    return <uintptr_t>(self.mem)
+
+def cl_from_visionbuf(VisionBuf buf):
+  return CLMem.create(<void*>&buf.buf.buf_cl)
+
+
+cdef class ModelFrame:
+  cdef cppModelFrame * frame
+  cdef int buf_size
+
+  def __dealloc__(self):
+    del self.frame
+
+  def prepare(self, VisionBuf buf, float[:] projection, CLMem output):
+    cdef mat3 cprojection
+    memcpy(cprojection.v, &projection[0], 9*sizeof(float))
+    cdef unsigned char * data
+    if output is None:
+      data = self.frame.prepare(buf.buf.buf_cl, buf.width, buf.height, buf.stride, buf.uv_offset, cprojection, NULL)
+    else:
+      data = self.frame.prepare(buf.buf.buf_cl, buf.width, buf.height, buf.stride, buf.uv_offset, cprojection, output.mem)
+    if not data:
+      return None
+
+    return np.asarray(<cnp.uint8_t[:self.buf_size]> data)
+    # return CLMem.create(data)
+
+  # def buffer_from_cl(self, CLMem in_frames):
+  #   cdef unsigned char * data2
+  #   data2 = self.frame.buffer_from_cl(in_frames.mem, self.buf_size)
+  #   return np.asarray(<cnp.uint8_t[:self.buf_size]> data2)
+
+
+cdef class DrivingModelFrame(ModelFrame):
+  cdef cppDrivingModelFrame * _frame
+
+  def __cinit__(self, CLContext context):
+    self._frame = new cppDrivingModelFrame(context.device_id, context.context)
+    self.frame = <cppModelFrame*>(self._frame)
+    self.buf_size = self._frame.buf_size
+
+cdef class MonitoringModelFrame(ModelFrame):
+  cdef cppMonitoringModelFrame * _frame
+
+  def __cinit__(self, CLContext context):
+    self._frame = new cppMonitoringModelFrame(context.device_id, context.context)
+    self.frame = <cppModelFrame*>(self._frame)
+    self.buf_size = self._frame.buf_size
@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:0c896681fd6851de3968433e12f37834429eba265e938cf383200be3e5835cec
+size 49096168
@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:af2cb689ec9e31292f759b561e70e4558a38f778558dff39ccff460ccafc0d52
+size 49849624
@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:441f2865017c07ee0dfb2488c5d86aab00df7ff5c5ec163959f35c33d74b65e6
+size 594
@@ -0,0 +1,103 @@
+import numpy as np
+from openpilot.sunnypilot.modeld.constants import ModelConstants
+
+def safe_exp(x, out=None):
+  # -11 is around 10**14, more causes float16 overflow
+  return np.exp(np.clip(x, -np.inf, 11), out=out)
+
+def sigmoid(x):
+  return 1. / (1. + safe_exp(-x))
+
+def softmax(x, axis=-1):
+  x -= np.max(x, axis=axis, keepdims=True)
+  if x.dtype == np.float32 or x.dtype == np.float64:
+    safe_exp(x, out=x)
+  else:
+    x = safe_exp(x)
+  x /= np.sum(x, axis=axis, keepdims=True)
+  return x
+
+class Parser:
+  def __init__(self, ignore_missing=False):
+    self.ignore_missing = ignore_missing
+
+  def check_missing(self, outs, name):
+    if name not in outs and not self.ignore_missing:
+      raise ValueError(f"Missing output {name}")
+    return name not in outs
+
+  def parse_categorical_crossentropy(self, name, outs, out_shape=None):
+    if self.check_missing(outs, name):
+      return
+    raw = outs[name]
+    if out_shape is not None:
+      raw = raw.reshape((raw.shape[0],) + out_shape)
+    outs[name] = softmax(raw, axis=-1)
+
+  def parse_binary_crossentropy(self, name, outs):
+    if self.check_missing(outs, name):
+      return
+    raw = outs[name]
+    outs[name] = sigmoid(raw)
+
+  def parse_mdn(self, name, outs, in_N=0, out_N=1, out_shape=None):
+    if self.check_missing(outs, name):
+      return
+    raw = outs[name]
+    raw = raw.reshape((raw.shape[0], max(in_N, 1), -1))
+
+    n_values = (raw.shape[2] - out_N)//2
+    pred_mu = raw[:,:,:n_values]
+    pred_std = safe_exp(raw[:,:,n_values: 2*n_values])
+
+    if in_N > 1:
+      weights = np.zeros((raw.shape[0], in_N, out_N), dtype=raw.dtype)
+      for i in range(out_N):
+        weights[:,:,i - out_N] = softmax(raw[:,:,i - out_N], axis=-1)
+
+      if out_N == 1:
+        for fidx in range(weights.shape[0]):
+          idxs = np.argsort(weights[fidx][:,0])[::-1]
+          weights[fidx] = weights[fidx][idxs]
+          pred_mu[fidx] = pred_mu[fidx][idxs]
+          pred_std[fidx] = pred_std[fidx][idxs]
+      full_shape = tuple([raw.shape[0], in_N] + list(out_shape))
+      outs[name + '_weights'] = weights
+      outs[name + '_hypotheses'] = pred_mu.reshape(full_shape)
+      outs[name + '_stds_hypotheses'] = pred_std.reshape(full_shape)
+
+      pred_mu_final = np.zeros((raw.shape[0], out_N, n_values), dtype=raw.dtype)
+      pred_std_final = np.zeros((raw.shape[0], out_N, n_values), dtype=raw.dtype)
+      for fidx in range(weights.shape[0]):
+        for hidx in range(out_N):
+          idxs = np.argsort(weights[fidx,:,hidx])[::-1]
+          pred_mu_final[fidx, hidx] = pred_mu[fidx, idxs[0]]
+          pred_std_final[fidx, hidx] = pred_std[fidx, idxs[0]]
+    else:
+      pred_mu_final = pred_mu
+      pred_std_final = pred_std
+
+    if out_N > 1:
+      final_shape = tuple([raw.shape[0], out_N] + list(out_shape))
+    else:
+      final_shape = tuple([raw.shape[0],] + list(out_shape))
+    outs[name] = pred_mu_final.reshape(final_shape)
+    outs[name + '_stds'] = pred_std_final.reshape(final_shape)
+
+  def parse_outputs(self, outs: dict[str, np.ndarray]) -> dict[str, np.ndarray]:
+    self.parse_mdn('plan', outs, in_N=ModelConstants.PLAN_MHP_N, out_N=ModelConstants.PLAN_MHP_SELECTION,
+                   out_shape=(ModelConstants.IDX_N,ModelConstants.PLAN_WIDTH))
+    self.parse_mdn('lane_lines', outs, in_N=0, out_N=0, out_shape=(ModelConstants.NUM_LANE_LINES,ModelConstants.IDX_N,ModelConstants.LANE_LINES_WIDTH))
+    self.parse_mdn('road_edges', outs, in_N=0, out_N=0, out_shape=(ModelConstants.NUM_ROAD_EDGES,ModelConstants.IDX_N,ModelConstants.LANE_LINES_WIDTH))
+    self.parse_mdn('pose', outs, in_N=0, out_N=0, out_shape=(ModelConstants.POSE_WIDTH,))
+    self.parse_mdn('road_transform', outs, in_N=0, out_N=0, out_shape=(ModelConstants.POSE_WIDTH,))
+    self.parse_mdn('wide_from_device_euler', outs, in_N=0, out_N=0, out_shape=(ModelConstants.WIDE_FROM_DEVICE_WIDTH,))
+    self.parse_mdn('lead', outs, in_N=ModelConstants.LEAD_MHP_N, out_N=ModelConstants.LEAD_MHP_SELECTION,
+                   out_shape=(ModelConstants.LEAD_TRAJ_LEN,ModelConstants.LEAD_WIDTH))
+    if 'lat_planner_solution' in outs:
+      self.parse_mdn('lat_planner_solution', outs, in_N=0, out_N=0, out_shape=(ModelConstants.IDX_N,ModelConstants.LAT_PLANNER_SOLUTION_WIDTH))
+    for k in ['lead_prob', 'lane_lines_prob', 'meta']:
+      self.parse_binary_crossentropy(k, outs)
+    self.parse_categorical_crossentropy('desire_state', outs, out_shape=(ModelConstants.DESIRE_PRED_WIDTH,))
+    self.parse_categorical_crossentropy('desire_pred', outs, out_shape=(ModelConstants.DESIRE_PRED_LEN,ModelConstants.DESIRE_PRED_WIDTH))
+    return outs
@@ -1,48 +0,0 @@
-# Copyright (c) 2021-, Haibin Wen, sunnypilot, and a number of other contributors.
-#
-# This file is part of sunnypilot and is licensed under the MIT License.
-# See the LICENSE.md file in the root directory for more details.
-
-import pickle
-import numpy as np
-from pathlib import Path
-from cereal import custom
-from openpilot.sunnypilot.models.helpers import get_active_bundle
-from openpilot.system.hardware.hw import Paths
-
-CUSTOM_MODEL_PATH = Paths.model_root()
-METADATA_PATH = Path(__file__).parent / '../models/supercombo_metadata.pkl'
-
-ModelManager = custom.ModelManagerSP
-
-
-def get_custom_model_paths():
-  bundle = get_active_bundle(None)
-  if bundle:
-    drive_model = next((model for model in bundle.models if model.type == ModelManager.Type.drive), None)
-    metadata_model = next(model for model in bundle.models if model.type == ModelManager.Type.metadata)
-    if drive_model and metadata_model:
-      return {"model": f"{CUSTOM_MODEL_PATH}/{drive_model.fileName}", "metadata": f"{CUSTOM_MODEL_PATH}/{metadata_model.fileName}"}
-
-  return None
-
-
-def load_custom_metadata():
-  if not (bundle := get_active_bundle(None)):
-    return None
-
-  metadata_model = next(model for model in bundle.models if model.type == ModelManager.Type.metadata)
-  metadata_path = f"{CUSTOM_MODEL_PATH}/{metadata_model.fileName}"
-  with open(metadata_path, 'rb') as f:
-    return pickle.load(f)
-
-
-def prepare_inputs(model_metadata) -> dict[str, np.ndarray]:
-  # img buffers are managed in openCL transform code so we don't pass them as inputs
-  inputs: dict[str, np.ndarray] = {
-    key: np.zeros(shape, dtype=np.float32).flatten()  # Inputs were defined flattened back then
-    for key, shape in model_metadata['input_shapes'].items()
-    if key not in ['input_imgs', 'big_input_imgs']
-  }
-
-  return inputs
@@ -0,0 +1,27 @@
+import os
+from openpilot.system.hardware import TICI
+from openpilot.sunnypilot.modeld.runners.runmodel_pyx import RunModel, Runtime
+assert Runtime
+
+USE_THNEED = int(os.getenv('USE_THNEED', str(int(TICI))))
+USE_SNPE = int(os.getenv('USE_SNPE', str(int(TICI))))
+
+class ModelRunner(RunModel):
+  THNEED = 'THNEED'
+  SNPE = 'SNPE'
+  ONNX = 'ONNX'
+
+  def __new__(cls, paths, *args, **kwargs):
+    if ModelRunner.THNEED in paths and USE_THNEED:
+      from openpilot.sunnypilot.modeld.runners.thneedmodel_pyx import ThneedModel as Runner
+      runner_type = ModelRunner.THNEED
+    elif ModelRunner.SNPE in paths and USE_SNPE:
+      from openpilot.sunnypilot.modeld.runners.snpemodel_pyx import SNPEModel as Runner
+      runner_type = ModelRunner.SNPE
+    elif ModelRunner.ONNX in paths:
+      from openpilot.sunnypilot.modeld.runners.onnxmodel import ONNXModel as Runner
+      runner_type = ModelRunner.ONNX
+    else:
+      raise Exception("Couldn't select a model runner, make sure to pass at least one valid model path")
+
+    return Runner(str(paths[runner_type]), *args, **kwargs)
@@ -0,0 +1,71 @@
+import os
+import onnx
+import sys
+import numpy as np
+from typing import Any
+
+from openpilot.sunnypilot.modeld.runners.runmodel_pyx import RunModel
+from openpilot.sunnypilot.modeld.runners.ort_helpers import convert_fp16_to_fp32, ORT_TYPES_TO_NP_TYPES
+
+
+def create_ort_session(path, fp16_to_fp32):
+  os.environ["OMP_NUM_THREADS"] = "4"
+  os.environ["OMP_WAIT_POLICY"] = "PASSIVE"
+
+  import onnxruntime as ort
+  print("Onnx available providers: ", ort.get_available_providers(), file=sys.stderr)
+  options = ort.SessionOptions()
+  options.graph_optimization_level = ort.GraphOptimizationLevel.ORT_DISABLE_ALL
+
+  provider: str | tuple[str, dict[Any, Any]]
+  if 'OpenVINOExecutionProvider' in ort.get_available_providers() and 'ONNXCPU' not in os.environ:
+    provider = 'OpenVINOExecutionProvider'
+  elif 'CUDAExecutionProvider' in ort.get_available_providers() and 'ONNXCPU' not in os.environ:
+    options.intra_op_num_threads = 2
+    provider = ('CUDAExecutionProvider', {'cudnn_conv_algo_search': 'EXHAUSTIVE'})
+  else:
+    options.intra_op_num_threads = 2
+    options.execution_mode = ort.ExecutionMode.ORT_SEQUENTIAL
+    options.graph_optimization_level = ort.GraphOptimizationLevel.ORT_ENABLE_ALL
+    provider = 'CPUExecutionProvider'
+
+  model_data = convert_fp16_to_fp32(onnx.load(path)) if fp16_to_fp32 else path
+  print("Onnx selected provider: ", [provider], file=sys.stderr)
+  ort_session = ort.InferenceSession(model_data, options, providers=[provider])
+  print("Onnx using ", ort_session.get_providers(), file=sys.stderr)
+  return ort_session
+
+
+class ONNXModel(RunModel):
+  def __init__(self, path, output, runtime, use_tf8, cl_context):
+    self.inputs = {}
+    self.output = output
+
+    self.session = create_ort_session(path, fp16_to_fp32=True)
+    self.input_names = [x.name for x in self.session.get_inputs()]
+    self.input_shapes = {x.name: [1, *x.shape[1:]] for x in self.session.get_inputs()}
+    self.input_dtypes = {x.name: ORT_TYPES_TO_NP_TYPES[x.type] for x in self.session.get_inputs()}
+
+    # run once to initialize CUDA provider
+    if "CUDAExecutionProvider" in self.session.get_providers():
+      self.session.run(None, {k: np.zeros(self.input_shapes[k], dtype=self.input_dtypes[k]) for k in self.input_names})
+    print("ready to run onnx model", self.input_shapes, file=sys.stderr)
+
+  def addInput(self, name, buffer):
+    assert name in self.input_names
+    self.inputs[name] = buffer
+
+  def setInputBuffer(self, name, buffer):
+    assert name in self.inputs
+    self.inputs[name] = buffer
+
+  def getCLBuffer(self, name):
+    return None
+
+  def execute(self):
+    inputs = {k: v.view(self.input_dtypes[k]) for k,v in self.inputs.items()}
+    inputs = {k: v.reshape(self.input_shapes[k]).astype(self.input_dtypes[k]) for k,v in inputs.items()}
+    outputs = self.session.run(None, inputs)
+    assert len(outputs) == 1, "Only single model outputs are supported"
+    self.output[:] = outputs[0]
+    return self.output
@@ -0,0 +1,36 @@
+import onnx
+import onnxruntime as ort
+import numpy as np
+import itertools
+
+ORT_TYPES_TO_NP_TYPES = {'tensor(float16)': np.float16, 'tensor(float)': np.float32, 'tensor(uint8)': np.uint8}
+
+def attributeproto_fp16_to_fp32(attr):
+  float32_list = np.frombuffer(attr.raw_data, dtype=np.float16)
+  attr.data_type = 1
+  attr.raw_data = float32_list.astype(np.float32).tobytes()
+
+def convert_fp16_to_fp32(model):
+  for i in model.graph.initializer:
+    if i.data_type == 10:
+      attributeproto_fp16_to_fp32(i)
+  for i in itertools.chain(model.graph.input, model.graph.output):
+    if i.type.tensor_type.elem_type == 10:
+      i.type.tensor_type.elem_type = 1
+  for i in model.graph.node:
+    if i.op_type == 'Cast' and i.attribute[0].i == 10:
+      i.attribute[0].i = 1
+    for a in i.attribute:
+      if hasattr(a, 't'):
+        if a.t.data_type == 10:
+          attributeproto_fp16_to_fp32(a.t)
+  return model.SerializeToString()
+
+
+def make_onnx_cpu_runner(model_path):
+  options = ort.SessionOptions()
+  options.intra_op_num_threads = 4
+  options.execution_mode = ort.ExecutionMode.ORT_SEQUENTIAL
+  options.graph_optimization_level = ort.GraphOptimizationLevel.ORT_ENABLE_ALL
+  model_data = convert_fp16_to_fp32(onnx.load(model_path))
+  return ort.InferenceSession(model_data,  options, providers=['CPUExecutionProvider'])
@@ -0,0 +1,4 @@
+#pragma once
+
+#include "sunnypilot/modeld/runners/runmodel.h"
+#include "sunnypilot/modeld/runners/snpemodel.h"
@@ -0,0 +1,49 @@
+#pragma once
+
+#include <string>
+#include <vector>
+#include <memory>
+#include <cassert>
+
+#include "common/clutil.h"
+#include "common/swaglog.h"
+
+#define USE_CPU_RUNTIME 0
+#define USE_GPU_RUNTIME 1
+#define USE_DSP_RUNTIME 2
+
+struct ModelInput {
+  const std::string name;
+  float *buffer;
+  int size;
+
+  ModelInput(const std::string _name, float *_buffer, int _size) : name(_name), buffer(_buffer), size(_size) {}
+  virtual void setBuffer(float *_buffer, int _size) {
+    assert(size == _size || size == 0);
+    buffer = _buffer;
+    size = _size;
+  }
+};
+
+class RunModel {
+public:
+  std::vector<std::unique_ptr<ModelInput>> inputs;
+
+  virtual ~RunModel() {}
+  virtual void execute() {}
+  virtual void* getCLBuffer(const std::string name) { return nullptr; }
+
+  virtual void addInput(const std::string name, float *buffer, int size) {
+    inputs.push_back(std::unique_ptr<ModelInput>(new ModelInput(name, buffer, size)));
+  }
+  virtual void setInputBuffer(const std::string name, float *buffer, int size) {
+    for (auto &input : inputs) {
+      if (name == input->name) {
+        input->setBuffer(buffer, size);
+        return;
+      }
+    }
+    LOGE("Tried to update input `%s` but no input with this name exists", name.c_str());
+    assert(false);
+  }
+};
@@ -0,0 +1,14 @@
+# distutils: language = c++
+
+from libcpp.string cimport string
+
+cdef extern from "sunnypilot/modeld/runners/runmodel.h":
+  cdef int USE_CPU_RUNTIME
+  cdef int USE_GPU_RUNTIME
+  cdef int USE_DSP_RUNTIME
+
+  cdef cppclass RunModel:
+    void addInput(string, float*, int)
+    void setInputBuffer(string, float*, int)
+    void * getCLBuffer(string)
+    void execute()
@@ -0,0 +1,6 @@
+# distutils: language = c++
+
+from .runmodel cimport RunModel as cppRunModel
+
+cdef class RunModel:
+  cdef cppRunModel * model
@@ -0,0 +1,37 @@
+# distutils: language = c++
+# cython: c_string_encoding=ascii, language_level=3
+
+from libcpp.string cimport string
+
+from .runmodel cimport USE_CPU_RUNTIME, USE_GPU_RUNTIME, USE_DSP_RUNTIME
+from sunnypilot.modeld.models.commonmodel_pyx cimport CLMem
+
+class Runtime:
+  CPU = USE_CPU_RUNTIME
+  GPU = USE_GPU_RUNTIME
+  DSP = USE_DSP_RUNTIME
+
+cdef class RunModel:
+  def __dealloc__(self):
+    del self.model
+
+  def addInput(self, string name, float[:] buffer):
+    if buffer is not None:
+      self.model.addInput(name, &buffer[0], len(buffer))
+    else:
+      self.model.addInput(name, NULL, 0)
+
+  def setInputBuffer(self, string name, float[:] buffer):
+    if buffer is not None:
+      self.model.setInputBuffer(name, &buffer[0], len(buffer))
+    else:
+      self.model.setInputBuffer(name, NULL, 0)
+
+  def getCLBuffer(self, string name):
+    cdef void * cl_buf = self.model.getCLBuffer(name)
+    if not cl_buf:
+      return None
+    return CLMem.create(cl_buf)
+
+  def execute(self):
+    self.model.execute()
@@ -0,0 +1,116 @@
+#pragma clang diagnostic ignored "-Wexceptions"
+
+#include "sunnypilot/modeld/runners/snpemodel.h"
+
+#include <cstring>
+#include <memory>
+#include <string>
+#include <utility>
+#include <vector>
+
+#include "common/util.h"
+#include "common/timing.h"
+
+void PrintErrorStringAndExit() {
+  std::cerr << zdl::DlSystem::getLastErrorString() << std::endl;
+  std::exit(EXIT_FAILURE);
+}
+
+SNPEModel::SNPEModel(const std::string path, float *_output, size_t _output_size, int runtime, bool _use_tf8, cl_context context) {
+  output = _output;
+  output_size = _output_size;
+  use_tf8 = _use_tf8;
+
+#ifdef QCOM2
+  if (runtime == USE_GPU_RUNTIME) {
+    snpe_runtime = zdl::DlSystem::Runtime_t::GPU;
+  } else if (runtime == USE_DSP_RUNTIME) {
+    snpe_runtime = zdl::DlSystem::Runtime_t::DSP;
+  } else {
+    snpe_runtime = zdl::DlSystem::Runtime_t::CPU;
+  }
+  assert(zdl::SNPE::SNPEFactory::isRuntimeAvailable(snpe_runtime));
+#endif
+  model_data = util::read_file(path);
+  assert(model_data.size() > 0);
+
+  // load model
+  std::unique_ptr<zdl::DlContainer::IDlContainer> container = zdl::DlContainer::IDlContainer::open((uint8_t*)model_data.data(), model_data.size());
+  if (!container) { PrintErrorStringAndExit(); }
+  LOGW("loaded model with size: %lu", model_data.size());
+
+  // create model runner
+  zdl::SNPE::SNPEBuilder snpe_builder(container.get());
+  while (!snpe) {
+#ifdef QCOM2
+    snpe = snpe_builder.setOutputLayers({})
+                       .setRuntimeProcessor(snpe_runtime)
+                       .setUseUserSuppliedBuffers(true)
+                       .setPerformanceProfile(zdl::DlSystem::PerformanceProfile_t::HIGH_PERFORMANCE)
+                       .build();
+#else
+    snpe = snpe_builder.setOutputLayers({})
+                       .setUseUserSuppliedBuffers(true)
+                       .setPerformanceProfile(zdl::DlSystem::PerformanceProfile_t::HIGH_PERFORMANCE)
+                       .build();
+#endif
+    if (!snpe) std::cerr << zdl::DlSystem::getLastErrorString() << std::endl;
+  }
+
+  // create output buffer
+  zdl::DlSystem::UserBufferEncodingFloat ub_encoding_float;
+  zdl::DlSystem::IUserBufferFactory &ub_factory = zdl::SNPE::SNPEFactory::getUserBufferFactory();
+
+  const auto &output_tensor_names_opt = snpe->getOutputTensorNames();
+  if (!output_tensor_names_opt) throw std::runtime_error("Error obtaining output tensor names");
+  const auto &output_tensor_names = *output_tensor_names_opt;
+  assert(output_tensor_names.size() == 1);
+  const char *output_tensor_name = output_tensor_names.at(0);
+  const zdl::DlSystem::TensorShape &buffer_shape = snpe->getInputOutputBufferAttributes(output_tensor_name)->getDims();
+  if (output_size != 0) {
+    assert(output_size == buffer_shape[1]);
+  } else {
+    output_size = buffer_shape[1];
+  }
+  std::vector<size_t> output_strides = {output_size * sizeof(float), sizeof(float)};
+  output_buffer = ub_factory.createUserBuffer(output, output_size * sizeof(float), output_strides, &ub_encoding_float);
+  output_map.add(output_tensor_name, output_buffer.get());
+}
+
+void SNPEModel::addInput(const std::string name, float *buffer, int size) {
+  const int idx = inputs.size();
+  const auto &input_tensor_names_opt = snpe->getInputTensorNames();
+  if (!input_tensor_names_opt) throw std::runtime_error("Error obtaining input tensor names");
+  const auto &input_tensor_names = *input_tensor_names_opt;
+  const char *input_tensor_name = input_tensor_names.at(idx);
+  const bool input_tf8 = use_tf8 && strcmp(input_tensor_name, "input_img") == 0;  // TODO: This is a terrible hack, get rid of this name check both here and in onnx_runner.py
+  LOGW("adding index %d: %s", idx, input_tensor_name);
+
+  zdl::DlSystem::UserBufferEncodingFloat ub_encoding_float;
+  zdl::DlSystem::UserBufferEncodingTf8 ub_encoding_tf8(0, 1./255); // network takes 0-1
+  zdl::DlSystem::IUserBufferFactory &ub_factory = zdl::SNPE::SNPEFactory::getUserBufferFactory();
+  zdl::DlSystem::UserBufferEncoding *input_encoding = input_tf8 ? (zdl::DlSystem::UserBufferEncoding*)&ub_encoding_tf8 : (zdl::DlSystem::UserBufferEncoding*)&ub_encoding_float;
+
+  const auto &buffer_shape_opt = snpe->getInputDimensions(input_tensor_name);
+  const zdl::DlSystem::TensorShape &buffer_shape = *buffer_shape_opt;
+  size_t size_of_input = input_tf8 ? sizeof(uint8_t) : sizeof(float);
+  std::vector<size_t> strides(buffer_shape.rank());
+  strides[strides.size() - 1] = size_of_input;
+  size_t product = 1;
+  for (size_t i = 0; i < buffer_shape.rank(); i++) product *= buffer_shape[i];
+  size_t stride = strides[strides.size() - 1];
+  for (size_t i = buffer_shape.rank() - 1; i > 0; i--) {
+    stride *= buffer_shape[i];
+    strides[i-1] = stride;
+  }
+
+  auto input_buffer = ub_factory.createUserBuffer(buffer, product*size_of_input, strides, input_encoding);
+  input_map.add(input_tensor_name, input_buffer.get());
+  inputs.push_back(std::unique_ptr<SNPEModelInput>(new SNPEModelInput(name, buffer, size, std::move(input_buffer))));
+}
+
+void SNPEModel::execute() {
+  if (!snpe->execute(input_map, output_map)) {
+    PrintErrorStringAndExit();
+  }
+}
@@ -0,0 +1,52 @@
+#pragma once
+#pragma clang diagnostic ignored "-Wdeprecated-declarations"
+
+#include <memory>
+#include <string>
+#include <utility>
+
+#include <DlContainer/IDlContainer.hpp>
+#include <DlSystem/DlError.hpp>
+#include <DlSystem/ITensor.hpp>
+#include <DlSystem/ITensorFactory.hpp>
+#include <DlSystem/IUserBuffer.hpp>
+#include <DlSystem/IUserBufferFactory.hpp>
+#include <SNPE/SNPE.hpp>
+#include <SNPE/SNPEBuilder.hpp>
+#include <SNPE/SNPEFactory.hpp>
+
+#include "sunnypilot/modeld/runners/runmodel.h"
+
+struct SNPEModelInput : public ModelInput {
+  std::unique_ptr<zdl::DlSystem::IUserBuffer> snpe_buffer;
+
+  SNPEModelInput(const std::string _name, float *_buffer, int _size, std::unique_ptr<zdl::DlSystem::IUserBuffer> _snpe_buffer) : ModelInput(_name, _buffer, _size), snpe_buffer(std::move(_snpe_buffer)) {}
+  void setBuffer(float *_buffer, int _size) {
+    ModelInput::setBuffer(_buffer, _size);
+    assert(snpe_buffer->setBufferAddress(_buffer) == true);
+  }
+};
+
+class SNPEModel : public RunModel {
+public:
+  SNPEModel(const std::string path, float *_output, size_t _output_size, int runtime, bool use_tf8 = false, cl_context context = NULL);
+  void addInput(const std::string name, float *buffer, int size);
+  void execute();
+
+private:
+  std::string model_data;
+
+#ifdef QCOM2
+  zdl::DlSystem::Runtime_t snpe_runtime;
+#endif
+
+  // snpe model stuff
+  std::unique_ptr<zdl::SNPE::SNPE> snpe;
+  zdl::DlSystem::UserBufferMap input_map;
+  zdl::DlSystem::UserBufferMap output_map;
+  std::unique_ptr<zdl::DlSystem::IUserBuffer> output_buffer;
+
+  bool use_tf8;
+  float *output;
+  size_t output_size;
+};
@@ -0,0 +1,9 @@
+# distutils: language = c++
+
+from libcpp.string cimport string
+
+from msgq.visionipc.visionipc cimport cl_context
+
+cdef extern from "sunnypilot/modeld/runners/snpemodel.h":
+  cdef cppclass SNPEModel:
+    SNPEModel(string, float*, size_t, int, bool, cl_context)
@@ -0,0 +1,17 @@
+# distutils: language = c++
+# cython: c_string_encoding=ascii, language_level=3
+
+import os
+from libcpp cimport bool
+from libcpp.string cimport string
+
+from .snpemodel cimport SNPEModel as cppSNPEModel
+from sunnypilot.modeld.models.commonmodel_pyx cimport CLContext
+from sunnypilot.modeld.runners.runmodel_pyx cimport RunModel
+from sunnypilot.modeld.runners.runmodel cimport RunModel as cppRunModel
+
+os.environ['ADSP_LIBRARY_PATH'] = "/data/pythonpath/third_party/snpe/dsp/"
+
+cdef class SNPEModel(RunModel):
+  def __cinit__(self, string path, float[:] output, int runtime, bool use_tf8, CLContext context):
+    self.model = <cppRunModel *> new cppSNPEModel(path, &output[0], len(output), runtime, use_tf8, context.context)
@@ -0,0 +1,58 @@
+#include "sunnypilot/modeld/runners/thneedmodel.h"
+
+#include <string>
+
+#include "common/swaglog.h"
+
+ThneedModel::ThneedModel(const std::string path, float *_output, size_t _output_size, int runtime, bool luse_tf8, cl_context context) {
+  thneed = new Thneed(true, context);
+  thneed->load(path.c_str());
+  thneed->clexec();
+
+  recorded = false;
+  output = _output;
+}
+
+void* ThneedModel::getCLBuffer(const std::string name) {
+  int index = -1;
+  for (int i = 0; i < inputs.size(); i++) {
+    if (name == inputs[i]->name) {
+      index = i;
+      break;
+    }
+  }
+
+  if (index == -1) {
+    LOGE("Tried to get CL buffer for input `%s` but no input with this name exists", name.c_str());
+    assert(false);
+  }
+
+  if (thneed->input_clmem.size() >= inputs.size()) {
+    return &thneed->input_clmem[inputs.size() - index - 1];
+  } else {
+    return nullptr;
+  }
+}
+
+void ThneedModel::execute() {
+  if (!recorded) {
+    thneed->record = true;
+    float *input_buffers[inputs.size()];
+    for (int i = 0; i < inputs.size(); i++) {
+      input_buffers[inputs.size() - i - 1] = inputs[i]->buffer;
+    }
+
+    thneed->copy_inputs(input_buffers);
+    thneed->clexec();
+    thneed->copy_output(output);
+    thneed->stop();
+
+    recorded = true;
+  } else {
+    float *input_buffers[inputs.size()];
+    for (int i = 0; i < inputs.size(); i++) {
+      input_buffers[inputs.size() - i - 1] = inputs[i]->buffer;
+    }
+    thneed->execute(input_buffers, output);
+  }
+}
@@ -0,0 +1,17 @@
+#pragma once
+
+#include <string>
+
+#include "sunnypilot/modeld/runners/runmodel.h"
+#include "sunnypilot/modeld/thneed/thneed.h"
+
+class ThneedModel : public RunModel {
+public:
+  ThneedModel(const std::string path, float *_output, size_t _output_size, int runtime, bool use_tf8 = false, cl_context context = NULL);
+  void *getCLBuffer(const std::string name);
+  void execute();
+private:
+  Thneed *thneed = NULL;
+  bool recorded;
+  float *output;
+};
@@ -0,0 +1,9 @@
+# distutils: language = c++
+
+from libcpp.string cimport string
+
+from msgq.visionipc.visionipc cimport cl_context
+
+cdef extern from "sunnypilot/modeld/runners/thneedmodel.h":
+  cdef cppclass ThneedModel:
+    ThneedModel(string, float*, size_t, int, bool, cl_context)
@@ -0,0 +1,14 @@
+# distutils: language = c++
+# cython: c_string_encoding=ascii, language_level=3
+
+from libcpp cimport bool
+from libcpp.string cimport string
+
+from .thneedmodel cimport ThneedModel as cppThneedModel
+from sunnypilot.modeld.models.commonmodel_pyx cimport CLContext
+from sunnypilot.modeld.runners.runmodel_pyx cimport RunModel
+from sunnypilot.modeld.runners.runmodel cimport RunModel as cppRunModel
+
+cdef class ThneedModel(RunModel):
+  def __cinit__(self, string path, float[:] output, int runtime, bool use_tf8, CLContext context):
+    self.model = <cppRunModel *> new cppThneedModel(path, &output[0], len(output), runtime, use_tf8, context.context)
@@ -0,0 +1,8 @@
+thneed is an SNPE accelerator. I know SNPE is already an accelerator, but sometimes things need to go even faster..
+
+It runs on the local device, and caches a single model run. Then it replays it, but fast.
+
+thneed slices through abstraction layers like a fish.
+
+You need a thneed.
+
@@ -0,0 +1,154 @@
+#include <cassert>
+#include <set>
+
+#include "third_party/json11/json11.hpp"
+#include "common/util.h"
+#include "common/clutil.h"
+#include "common/swaglog.h"
+#include "sunnypilot/modeld/thneed/thneed.h"
+using namespace json11;
+
+extern map<cl_program, string> g_program_source;
+
+void Thneed::load(const char *filename) {
+  LOGD("Thneed::load: loading from %s\n", filename);
+
+  string buf = util::read_file(filename);
+  int jsz = *(int *)buf.data();
+  string jsonerr;
+  string jj(buf.data() + sizeof(int), jsz);
+  Json jdat = Json::parse(jj, jsonerr);
+
+  map<cl_mem, cl_mem> real_mem;
+  real_mem[NULL] = NULL;
+
+  int ptr = sizeof(int)+jsz;
+  for (auto &obj : jdat["objects"].array_items()) {
+    auto mobj = obj.object_items();
+    int sz = mobj["size"].int_value();
+    cl_mem clbuf = NULL;
+
+    if (mobj["buffer_id"].string_value().size() > 0) {
+      // image buffer must already be allocated
+      clbuf = real_mem[*(cl_mem*)(mobj["buffer_id"].string_value().data())];
+      assert(mobj["needs_load"].bool_value() == false);
+    } else {
+      if (mobj["needs_load"].bool_value()) {
+        clbuf = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR | CL_MEM_READ_WRITE, sz, &buf[ptr], NULL);
+        if (debug >= 1) printf("loading %p %d @ 0x%X\n", clbuf, sz, ptr);
+        ptr += sz;
+      } else {
+        // TODO: is there a faster way to init zeroed out buffers?
+        void *host_zeros = calloc(sz, 1);
+        clbuf = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR | CL_MEM_READ_WRITE, sz, host_zeros, NULL);
+        free(host_zeros);
+      }
+    }
+    assert(clbuf != NULL);
+
+    if (mobj["arg_type"] == "image2d_t" || mobj["arg_type"] == "image1d_t") {
+      cl_image_desc desc = {0};
+      desc.image_type = (mobj["arg_type"] == "image2d_t") ? CL_MEM_OBJECT_IMAGE2D : CL_MEM_OBJECT_IMAGE1D_BUFFER;
+      desc.image_width = mobj["width"].int_value();
+      desc.image_height = mobj["height"].int_value();
+      desc.image_row_pitch = mobj["row_pitch"].int_value();
+      assert(sz == desc.image_height*desc.image_row_pitch);
+#ifdef QCOM2
+      desc.buffer = clbuf;
+#else
+      // TODO: we are creating unused buffers on PC
+      clReleaseMemObject(clbuf);
+#endif
+      cl_image_format format = {0};
+      format.image_channel_order = CL_RGBA;
+      format.image_channel_data_type = mobj["float32"].bool_value() ? CL_FLOAT : CL_HALF_FLOAT;
+
+      cl_int errcode;
+
+#ifndef QCOM2
+      if (mobj["needs_load"].bool_value()) {
+        clbuf = clCreateImage(context, CL_MEM_COPY_HOST_PTR | CL_MEM_READ_WRITE, &format, &desc, &buf[ptr-sz], &errcode);
+      } else {
+        clbuf = clCreateImage(context, CL_MEM_READ_WRITE, &format, &desc, NULL, &errcode);
+      }
+#else
+      clbuf = clCreateImage(context, CL_MEM_READ_WRITE, &format, &desc, NULL, &errcode);
+#endif
+      if (clbuf == NULL) {
+        LOGE("clError: %s create image %zux%zu rp %zu with buffer %p\n", cl_get_error_string(errcode),
+             desc.image_width, desc.image_height, desc.image_row_pitch, desc.buffer);
+      }
+      assert(clbuf != NULL);
+    }
+
+    real_mem[*(cl_mem*)(mobj["id"].string_value().data())] = clbuf;
+  }
+
+  map<string, cl_program> g_programs;
+  for (const auto &[name, source] : jdat["programs"].object_items()) {
+    if (debug >= 1) printf("building %s with size %zu\n", name.c_str(), source.string_value().size());
+    g_programs[name] = cl_program_from_source(context, device_id, source.string_value());
+  }
+
+  for (auto &obj : jdat["inputs"].array_items()) {
+    auto mobj = obj.object_items();
+    int sz = mobj["size"].int_value();
+    cl_mem aa = real_mem[*(cl_mem*)(mobj["buffer_id"].string_value().data())];
+    input_clmem.push_back(aa);
+    input_sizes.push_back(sz);
+    LOGD("Thneed::load: adding input %s with size %d\n", mobj["name"].string_value().data(), sz);
+
+    cl_int cl_err;
+    void *ret = clEnqueueMapBuffer(command_queue, aa, CL_TRUE, CL_MAP_WRITE, 0, sz, 0, NULL, NULL, &cl_err);
+    if (cl_err != CL_SUCCESS) LOGE("clError: %s map %p %d\n", cl_get_error_string(cl_err), aa, sz);
+    assert(cl_err == CL_SUCCESS);
+    inputs.push_back(ret);
+  }
+
+  for (auto &obj : jdat["outputs"].array_items()) {
+    auto mobj = obj.object_items();
+    int sz = mobj["size"].int_value();
+    LOGD("Thneed::save: adding output with size %d\n", sz);
+    // TODO: support multiple outputs
+    output = real_mem[*(cl_mem*)(mobj["buffer_id"].string_value().data())];
+    assert(output != NULL);
+  }
+
+  for (auto &obj : jdat["binaries"].array_items()) {
+    string name = obj["name"].string_value();
+    size_t length = obj["length"].int_value();
+    if (debug >= 1) printf("binary %s with size %zu\n", name.c_str(), length);
+    g_programs[name] = cl_program_from_binary(context, device_id, (const uint8_t*)&buf[ptr], length);
+    ptr += length;
+  }
+
+  for (auto &obj : jdat["kernels"].array_items()) {
+    auto gws = obj["global_work_size"];
+    auto lws = obj["local_work_size"];
+    auto kk = shared_ptr<CLQueuedKernel>(new CLQueuedKernel(this));
+
+    kk->name = obj["name"].string_value();
+    kk->program = g_programs[kk->name];
+    kk->work_dim = obj["work_dim"].int_value();
+    for (int i = 0; i < kk->work_dim; i++) {
+      kk->global_work_size[i] = gws[i].int_value();
+      kk->local_work_size[i] = lws[i].int_value();
+    }
+    kk->num_args = obj["num_args"].int_value();
+    for (int i = 0; i < kk->num_args; i++) {
+      string arg = obj["args"].array_items()[i].string_value();
+      int arg_size = obj["args_size"].array_items()[i].int_value();
+      kk->args_size.push_back(arg_size);
+      if (arg_size == 8) {
+        cl_mem val = *(cl_mem*)(arg.data());
+        val = real_mem[val];
+        kk->args.push_back(string((char*)&val, sizeof(val)));
+      } else {
+        kk->args.push_back(arg);
+      }
+    }
+    kq.push_back(kk);
+  }
+
+  clFinish(command_queue);
+}
@@ -0,0 +1,133 @@
+#pragma once
+
+#ifndef __user
+#define __user __attribute__(())
+#endif
+
+#include <cstdint>
+#include <cstdlib>
+#include <memory>
+#include <string>
+#include <vector>
+
+#include <CL/cl.h>
+
+#include "third_party/linux/include/msm_kgsl.h"
+
+using namespace std;
+
+cl_int thneed_clSetKernelArg(cl_kernel kernel, cl_uint arg_index, size_t arg_size, const void *arg_value);
+
+namespace json11 {
+  class Json;
+}
+class Thneed;
+
+class GPUMalloc {
+  public:
+    GPUMalloc(int size, int fd);
+    ~GPUMalloc();
+    void *alloc(int size);
+  private:
+    uint64_t base;
+    int remaining;
+};
+
+class CLQueuedKernel {
+  public:
+    CLQueuedKernel(Thneed *lthneed) { thneed = lthneed; }
+    CLQueuedKernel(Thneed *lthneed,
+                   cl_kernel _kernel,
+                   cl_uint _work_dim,
+                   const size_t *_global_work_size,
+                   const size_t *_local_work_size);
+    cl_int exec();
+    void debug_print(bool verbose);
+    int get_arg_num(const char *search_arg_name);
+    cl_program program;
+    string name;
+    cl_uint num_args;
+    vector<string> arg_names;
+    vector<string> arg_types;
+    vector<string> args;
+    vector<int> args_size;
+    cl_kernel kernel = NULL;
+    json11::Json to_json() const;
+
+    cl_uint work_dim;
+    size_t global_work_size[3] = {0};
+    size_t local_work_size[3] = {0};
+  private:
+    Thneed *thneed;
+};
+
+class CachedIoctl {
+  public:
+    virtual void exec() {}
+};
+
+class CachedSync: public CachedIoctl {
+  public:
+    CachedSync(Thneed *lthneed, string ldata) { thneed = lthneed; data = ldata; }
+    void exec();
+  private:
+    Thneed *thneed;
+    string data;
+};
+
+class CachedCommand: public CachedIoctl {
+  public:
+    CachedCommand(Thneed *lthneed, struct kgsl_gpu_command *cmd);
+    void exec();
+  private:
+    void disassemble(int cmd_index);
+    struct kgsl_gpu_command cache;
+    unique_ptr<kgsl_command_object[]> cmds;
+    unique_ptr<kgsl_command_object[]> objs;
+    Thneed *thneed;
+    vector<shared_ptr<CLQueuedKernel> > kq;
+};
+
+class Thneed {
+  public:
+    Thneed(bool do_clinit=false, cl_context _context = NULL);
+    void stop();
+    void execute(float **finputs, float *foutput, bool slow=false);
+    void wait();
+
+    vector<cl_mem> input_clmem;
+    vector<void *> inputs;
+    vector<size_t> input_sizes;
+    cl_mem output = NULL;
+
+    cl_context context = NULL;
+    cl_command_queue command_queue;
+    cl_device_id device_id;
+    int context_id;
+
+    // protected?
+    bool record = false;
+    int debug;
+    int timestamp;
+
+#ifdef QCOM2
+    unique_ptr<GPUMalloc> ram;
+    vector<unique_ptr<CachedIoctl> > cmds;
+    int fd;
+#endif
+
+    // all CL kernels
+    void copy_inputs(float **finputs, bool internal=false);
+    void copy_output(float *foutput);
+    cl_int clexec();
+    vector<shared_ptr<CLQueuedKernel> > kq;
+
+    // pending CL kernels
+    vector<shared_ptr<CLQueuedKernel> > ckq;
+
+    // loading
+    void load(const char *filename);
+  private:
+    void clinit();
+};
+
@@ -0,0 +1,216 @@
+#include "sunnypilot/modeld/thneed/thneed.h"
+
+#include <cassert>
+#include <cstring>
+#include <map>
+
+#include "common/clutil.h"
+#include "common/timing.h"
+
+map<pair<cl_kernel, int>, string> g_args;
+map<pair<cl_kernel, int>, int> g_args_size;
+map<cl_program, string> g_program_source;
+
+void Thneed::stop() {
+  //printf("Thneed::stop: recorded %lu commands\n", cmds.size());
+  record = false;
+}
+
+void Thneed::clinit() {
+  device_id = cl_get_device_id(CL_DEVICE_TYPE_DEFAULT);
+  if (context == NULL) context = CL_CHECK_ERR(clCreateContext(NULL, 1, &device_id, NULL, NULL, &err));
+  //cl_command_queue_properties props[3] = {CL_QUEUE_PROPERTIES, CL_QUEUE_PROFILING_ENABLE, 0};
+  cl_command_queue_properties props[3] = {CL_QUEUE_PROPERTIES, 0, 0};
+  command_queue = CL_CHECK_ERR(clCreateCommandQueueWithProperties(context, device_id, props, &err));
+  printf("Thneed::clinit done\n");
+}
+
+cl_int Thneed::clexec() {
+  if (debug >= 1) printf("Thneed::clexec: running %lu queued kernels\n", kq.size());
+  for (auto &k : kq) {
+    if (record) ckq.push_back(k);
+    cl_int ret = k->exec();
+    assert(ret == CL_SUCCESS);
+  }
+  return clFinish(command_queue);
+}
+
+void Thneed::copy_inputs(float **finputs, bool internal) {
+  for (int idx = 0; idx < inputs.size(); ++idx) {
+    if (debug >= 1) printf("copying %lu -- %p -> %p (cl %p)\n", input_sizes[idx], finputs[idx], inputs[idx], input_clmem[idx]);
+
+    if (internal) {
+      // if it's internal, using memcpy is fine since the buffer sync is cached in the ioctl layer
+      if (finputs[idx] != NULL) memcpy(inputs[idx], finputs[idx], input_sizes[idx]);
+    } else {
+      if (finputs[idx] != NULL) CL_CHECK(clEnqueueWriteBuffer(command_queue, input_clmem[idx], CL_TRUE, 0, input_sizes[idx], finputs[idx], 0, NULL, NULL));
+    }
+  }
+}
+
+void Thneed::copy_output(float *foutput) {
+  if (output != NULL) {
+    size_t sz;
+    clGetMemObjectInfo(output, CL_MEM_SIZE, sizeof(sz), &sz, NULL);
+    if (debug >= 1) printf("copying %lu for output %p -> %p\n", sz, output, foutput);
+    CL_CHECK(clEnqueueReadBuffer(command_queue, output, CL_TRUE, 0, sz, foutput, 0, NULL, NULL));
+  } else {
+    printf("CAUTION: model output is NULL, does it have no outputs?\n");
+  }
+}
+
+// *********** CLQueuedKernel ***********
+
+CLQueuedKernel::CLQueuedKernel(Thneed *lthneed,
+                               cl_kernel _kernel,
+                               cl_uint _work_dim,
+                               const size_t *_global_work_size,
+                               const size_t *_local_work_size) {
+  thneed = lthneed;
+  kernel = _kernel;
+  work_dim = _work_dim;
+  assert(work_dim <= 3);
+  for (int i = 0; i < work_dim; i++) {
+    global_work_size[i] = _global_work_size[i];
+    local_work_size[i] = _local_work_size[i];
+  }
+
+  char _name[0x100];
+  clGetKernelInfo(kernel, CL_KERNEL_FUNCTION_NAME, sizeof(_name), _name, NULL);
+  name = string(_name);
+  clGetKernelInfo(kernel, CL_KERNEL_NUM_ARGS, sizeof(num_args), &num_args, NULL);
+
+  // get args
+  for (int i = 0; i < num_args; i++) {
+    char arg_name[0x100] = {0};
+    clGetKernelArgInfo(kernel, i, CL_KERNEL_ARG_NAME, sizeof(arg_name), arg_name, NULL);
+    arg_names.push_back(string(arg_name));
+    clGetKernelArgInfo(kernel, i, CL_KERNEL_ARG_TYPE_NAME, sizeof(arg_name), arg_name, NULL);
+    arg_types.push_back(string(arg_name));
+
+    args.push_back(g_args[make_pair(kernel, i)]);
+    args_size.push_back(g_args_size[make_pair(kernel, i)]);
+  }
+
+  // get program
+  clGetKernelInfo(kernel, CL_KERNEL_PROGRAM, sizeof(program), &program, NULL);
+}
+
+int CLQueuedKernel::get_arg_num(const char *search_arg_name) {
+  for (int i = 0; i < num_args; i++) {
+    if (arg_names[i] == search_arg_name) return i;
+  }
+  printf("failed to find %s in %s\n", search_arg_name, name.c_str());
+  assert(false);
+}
+
+cl_int CLQueuedKernel::exec() {
+  if (kernel == NULL) {
+    kernel = clCreateKernel(program, name.c_str(), NULL);
+    arg_names.clear();
+    arg_types.clear();
+
+    for (int j = 0; j < num_args; j++) {
+      char arg_name[0x100] = {0};
+      clGetKernelArgInfo(kernel, j, CL_KERNEL_ARG_NAME, sizeof(arg_name), arg_name, NULL);
+      arg_names.push_back(string(arg_name));
+      clGetKernelArgInfo(kernel, j, CL_KERNEL_ARG_TYPE_NAME, sizeof(arg_name), arg_name, NULL);
+      arg_types.push_back(string(arg_name));
+
+      cl_int ret;
+      if (args[j].size() != 0) {
+        assert(args[j].size() == args_size[j]);
+        ret = thneed_clSetKernelArg(kernel, j, args[j].size(), args[j].data());
+      } else {
+        ret = thneed_clSetKernelArg(kernel, j, args_size[j], NULL);
+      }
+      assert(ret == CL_SUCCESS);
+    }
+  }
+
+  if (thneed->debug >= 1) {
+    debug_print(thneed->debug >= 2);
+  }
+
+  return clEnqueueNDRangeKernel(thneed->command_queue,
+    kernel, work_dim, NULL, global_work_size, local_work_size, 0, NULL, NULL);
+}
+
+void CLQueuedKernel::debug_print(bool verbose) {
+  printf("%p %56s -- ", kernel, name.c_str());
+  for (int i = 0; i < work_dim; i++) {
+    printf("%4zu ", global_work_size[i]);
+  }
+  printf(" -- ");
+  for (int i = 0; i < work_dim; i++) {
+    printf("%4zu ", local_work_size[i]);
+  }
+  printf("\n");
+
+  if (verbose) {
+    for (int i = 0; i < num_args; i++) {
+      string arg = args[i];
+      printf("  %s %s", arg_types[i].c_str(), arg_names[i].c_str());
+      void *arg_value = (void*)arg.data();
+      int arg_size = arg.size();
+      if (arg_size == 0) {
+        printf(" (size) %d", args_size[i]);
+      } else if (arg_size == 1) {
+        printf(" = %d", *((char*)arg_value));
+      } else if (arg_size == 2) {
+        printf(" = %d", *((short*)arg_value));
+      } else if (arg_size == 4) {
+        if (arg_types[i] == "float") {
+          printf(" = %f", *((float*)arg_value));
+        } else {
+          printf(" = %d", *((int*)arg_value));
+        }
+      } else if (arg_size == 8) {
+        cl_mem val = (cl_mem)(*((uintptr_t*)arg_value));
+        printf(" = %p", val);
+        if (val != NULL) {
+          cl_mem_object_type obj_type;
+          clGetMemObjectInfo(val, CL_MEM_TYPE, sizeof(obj_type), &obj_type, NULL);
+          if (arg_types[i] == "image2d_t" || arg_types[i] == "image1d_t" || obj_type == CL_MEM_OBJECT_IMAGE2D) {
+            cl_image_format format;
+            size_t width, height, depth, array_size, row_pitch, slice_pitch;
+            cl_mem buf;
+            clGetImageInfo(val, CL_IMAGE_FORMAT, sizeof(format), &format, NULL);
+            assert(format.image_channel_order == CL_RGBA);
+            assert(format.image_channel_data_type == CL_HALF_FLOAT || format.image_channel_data_type == CL_FLOAT);
+            clGetImageInfo(val, CL_IMAGE_WIDTH, sizeof(width), &width, NULL);
+            clGetImageInfo(val, CL_IMAGE_HEIGHT, sizeof(height), &height, NULL);
+            clGetImageInfo(val, CL_IMAGE_ROW_PITCH, sizeof(row_pitch), &row_pitch, NULL);
+            clGetImageInfo(val, CL_IMAGE_DEPTH, sizeof(depth), &depth, NULL);
+            clGetImageInfo(val, CL_IMAGE_ARRAY_SIZE, sizeof(array_size), &array_size, NULL);
+            clGetImageInfo(val, CL_IMAGE_SLICE_PITCH, sizeof(slice_pitch), &slice_pitch, NULL);
+            assert(depth == 0);
+            assert(array_size == 0);
+            assert(slice_pitch == 0);
+
+            clGetImageInfo(val, CL_IMAGE_BUFFER, sizeof(buf), &buf, NULL);
+            size_t sz = 0;
+            if (buf != NULL) clGetMemObjectInfo(buf, CL_MEM_SIZE, sizeof(sz), &sz, NULL);
+            printf(" image %zu x %zu rp %zu @ %p buffer %zu", width, height, row_pitch, buf, sz);
+          } else {
+            size_t sz;
+            clGetMemObjectInfo(val, CL_MEM_SIZE, sizeof(sz), &sz, NULL);
+            printf(" buffer %zu", sz);
+          }
+        }
+      }
+      printf("\n");
+    }
+  }
+}
+
+cl_int thneed_clSetKernelArg(cl_kernel kernel, cl_uint arg_index, size_t arg_size, const void *arg_value) {
+  g_args_size[make_pair(kernel, arg_index)] = arg_size;
+  if (arg_value != NULL) {
+    g_args[make_pair(kernel, arg_index)] = string((char*)arg_value, arg_size);
+  } else {
+    g_args[make_pair(kernel, arg_index)] = string("");
+  }
+  cl_int ret = clSetKernelArg(kernel, arg_index, arg_size, arg_value);
+  return ret;
+}
--- a/Show More
+++ b/Show More
Author	SHA1	Message	Date
DevTekVE	8120930372	Merge branch 'refs/heads/model-manager-improvements' into model-selector-multi-runner-debug-model-download	2025-01-05 22:55:45 +01:00
DevTekVE	adc9f2cde8	Enable model label button only when conditions are met Previously, the button's state update was misplaced, leading to potential issues with its interactive availability. The logic has been adjusted to ensure it is properly enabled or disabled based on onroad status and download progress. This change improves UX consistency and prevents unintended actions.	2025-01-05 22:54:29 +01:00
DevTekVE	0c89a58e91	Using is_onroad softwarePanel	2025-01-05 22:50:05 +01:00
DevTekVE	c467f8ea08	Merge branch 'refs/heads/model-manager-improvements' into model-selector-multi-runner-debug-model-download	2025-01-05 22:28:53 +01:00
DevTekVE	aa9f830bf4	Update model manager logic and handle offroad transitions Added is_onroad state tracking in SoftwarePanelSP to handle offroad transitions. Updated model manager conditions for improved bundle validation. Removed unnecessary clear operation for ModelManager_DownloadIndex during offroad transitions to optimize behavior.	2025-01-05 22:28:26 +01:00
DevTekVE	30ed809a60	fix prints	2025-01-05 21:24:42 +01:00
DevTekVE	df7d1ef256	Add detailed debug logging to model download process Enhanced logging provides better traceability during the download process. New debug logs include information such as URLs, file paths, response statuses, and model details. This facilitates easier debugging and monitoring of the model download workflow.	2025-01-05 20:12:06 +01:00
Jason Wen	491373bbd6	need this	2025-01-05 10:09:31 -05:00
Jason Wen	0f1c952a3e	must add this back	2025-01-05 10:09:00 -05:00
Jason Wen	6f57822ba1	bring onnx back for sim	2025-01-05 09:46:29 -05:00
Jason Wen	2e7ab9ce85	add to lfs	2025-01-05 09:41:34 -05:00
Jason Wen	0d47532773	need it for default snpe model	2025-01-05 09:07:31 -05:00
Jason Wen	acbcdded4f	don't even compile anymore	2025-01-05 08:45:43 -05:00
Jason Wen	bd3b4dd2e7	Reapply "remove our own" This reverts commit `b1996377b3`.	2025-01-05 08:37:46 -05:00
Jason Wen	a2e30cc7d1	Revert "try using compile2.py again" This reverts commit `914117d2e1`.	2025-01-05 08:37:34 -05:00
Jason Wen	b52347e7b4	Revert "add back symlink" This reverts commit `9f71ad0b8a`.	2025-01-05 08:37:33 -05:00
Jason Wen	36576ad5ad	Revert "fix path" This reverts commit `75d338f2bd`.	2025-01-05 08:37:32 -05:00
Jason Wen	5afa0174c5	Revert "more fix" This reverts commit `23dd423e78`.	2025-01-05 08:37:32 -05:00
Jason Wen	74126eaef8	Revert "wrong path again" This reverts commit `f5301c19d5`.	2025-01-05 08:37:31 -05:00
Jason Wen	b78f14bff3	Reapply "wrong path again" This reverts commit `309639aeb3`.	2025-01-05 08:37:30 -05:00
Jason Wen	c409ac546a	Revert "update" This reverts commit `fb313bd7fb`.	2025-01-05 08:37:30 -05:00
Jason Wen	42af2fbbc2	Revert "hardcode path to our submodule" This reverts commit `5ee1950b6f`.	2025-01-05 08:37:29 -05:00
Jason Wen	8642689c6d	Revert "force path" This reverts commit `5c3b408937`.	2025-01-05 08:37:28 -05:00
Jason Wen	8e6fb8547a	Revert "try this" This reverts commit `41fef87680`.	2025-01-05 08:37:28 -05:00
Jason Wen	0dbb46aa12	Revert "fix file name" This reverts commit `485eef68da`.	2025-01-05 08:37:27 -05:00
Jason Wen	b930a83b8d	Revert "try this" This reverts commit `767f78bbcf`.	2025-01-05 08:37:27 -05:00
Jason Wen	878cec45ad	Revert "again" This reverts commit `17c8cd7376`.	2025-01-05 08:37:26 -05:00
Jason Wen	17c8cd7376	again	2025-01-05 08:34:14 -05:00
Jason Wen	767f78bbcf	try this	2025-01-05 08:30:28 -05:00
Jason Wen	485eef68da	fix file name	2025-01-05 08:27:23 -05:00
Jason Wen	41fef87680	try this	2025-01-05 08:26:23 -05:00
Jason Wen	5c3b408937	force path	2025-01-05 08:21:13 -05:00
Jason Wen	5ee1950b6f	hardcode path to our submodule	2025-01-05 08:12:02 -05:00
Jason Wen	fb313bd7fb	update	2025-01-05 08:09:30 -05:00
Jason Wen	309639aeb3	Revert "wrong path again" This reverts commit `f5301c19d5`.	2025-01-05 08:06:49 -05:00
Jason Wen	f5301c19d5	wrong path again	2025-01-05 08:04:49 -05:00
Jason Wen	23dd423e78	more fix	2025-01-05 07:59:18 -05:00
Jason Wen	75d338f2bd	fix path	2025-01-05 07:56:48 -05:00
Jason Wen	9f71ad0b8a	add back symlink	2025-01-05 07:55:30 -05:00
Jason Wen	914117d2e1	try using compile2.py again	2025-01-05 07:54:34 -05:00
Jason Wen	b1996377b3	Revert "remove our own" This reverts commit `1cf4f57502`.	2025-01-05 07:52:11 -05:00
Jason Wen	158a76289e	try this	2025-01-05 07:35:41 -05:00
Jason Wen	5c125f5fa4	fix thneed	2025-01-05 07:21:11 -05:00
Jason Wen	130ba6b905	use upstream compile3	2025-01-05 06:59:29 -05:00
Jason Wen	1cf4f57502	remove our own	2025-01-05 06:59:19 -05:00
Jason Wen	f9ca110410	mypy	2025-01-05 06:37:08 -05:00
Jason Wen	4bdecdec11	fix thneed paths	2025-01-05 06:32:58 -05:00
Jason Wen	4b6c94e794	Merge branch 'master-new' into model-selector-multi-runner	2025-01-05 06:31:54 -05:00
Jason Wen	59c551ac77	ruff	2025-01-05 06:28:46 -05:00
Jason Wen	c54cc074e2	fix process name	2025-01-05 06:25:18 -05:00
Jason Wen	07391c72b4	ignore tg	2025-01-05 06:20:49 -05:00
DevTekVE	e46aaf0263	Refactor modeld process function checks. Introduce `is_stock_model` to clarify logic and replace direct uses of `is_snpe_model` where the stock model condition is needed. Additionally, rename the duplicate "modeld" process in sunnyPilot to "modeld_snpe" for clarity and consistency.	2025-01-05 12:16:24 +01:00
DevTekVE	f3db1254c3	Adjust modeld execution logic based on active model runner Introduced a check to conditionally execute `modeld` based on the active model runner. Added support for distinguishing between SNPE and TinyGrad runners using new helper functions and updated `custom.capnp` definitions. This change optimizes process management by ensuring compatibility with the selected model runner.	2025-01-05 11:56:31 +01:00
Jason Wen	2c3d776a52	fix more paths	2025-01-05 05:49:31 -05:00
Jason Wen	8516026c74	fix path	2025-01-05 05:35:48 -05:00
Jason Wen	b916e9c655	force with snpe to validate	2025-01-05 05:30:37 -05:00
Jason Wen	15d127889b	tinygrad with snpe	2025-01-05 05:21:25 -05:00