Condensing a little

This commit introduces functionality to detect if the active model operates at 20Hz. A new helper function `is_active_model_20hz` was added, alongside updates to model initialization and configuration handling. Additionally, a new `is20hz` field was added to the model bundle for improved introspection and performance adaptation.

.codespellignore

@@ -0,0 +1,2 @@
+Wen
+REGIST

.github/PULL_REQUEST_TEMPLATE/bugfix.md

@@ -1,15 +0,0 @@
----
-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. -->

.github/PULL_REQUEST_TEMPLATE/car_bugfix.md

@@ -1,19 +0,0 @@
----
-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]

.github/PULL_REQUEST_TEMPLATE/car_port.md

@@ -1,15 +0,0 @@
----
-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):

.github/PULL_REQUEST_TEMPLATE/fingerprint.md

@@ -1,13 +0,0 @@
----
-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

.github/PULL_REQUEST_TEMPLATE/refactor.md

@@ -1,15 +0,0 @@
----
-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. -->

.github/PULL_REQUEST_TEMPLATE/tuning.md

@@ -1,31 +0,0 @@
----
-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
--->

.github/build.py

@@ -1,30 +0,0 @@
-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()

.github/workflows/ui_preview.yaml

@@ -86,7 +86,7 @@ jobs:
         run: >-
           sudo apt-get install -y imagemagick
 
-          scenes="homescreen settings_device settings_software settings_toggles settings_developer offroad_alert update_available prime onroad onroad_disengaged onroad_override onroad_sidebar onroad_wide onroad_wide_sidebar onroad_alert_small onroad_alert_mid onroad_alert_full driver_camera body keyboard"
+          scenes="homescreen settings_device settings_software settings_sunnylink settings_toggles settings_sunnypilot settings_sunnypilot_mads settings_developer offroad_alert update_available prime onroad onroad_disengaged onroad_override onroad_sidebar onroad_wide onroad_wide_sidebar onroad_alert_small onroad_alert_mid onroad_alert_full driver_camera body keyboard"
           A=($scenes)
 
           DIFF=""

SConstruct

@@ -49,10 +49,6 @@ 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',
@@ -74,6 +70,12 @@ AddOption('--minimal',
           default=os.path.exists(File('#.lfsconfig').abspath), # minimal by default on release branch (where there's no LFS)
           help='the minimum build to run openpilot. no tests, tools, etc.')
 
+AddOption('--stock-ui',
+          action='store_true',
+          dest='stock_ui',
+          default=False,
+          help='Build stock openpilot UI instead of sunnypilot UI')
+
 ## Architecture name breakdown (arch)
 ## - larch64: linux tici aarch64
 ## - aarch64: linux pc aarch64
@@ -172,6 +174,10 @@ else:
 if arch != "Darwin":
   ldflags += ["-Wl,--as-needed", "-Wl,--no-undefined"]
 
+if not GetOption('stock_ui'):
+  cflags += ["-DSUNNYPILOT"]
+  cxxflags += ["-DSUNNYPILOT"]
+
 ccflags_option = GetOption('ccflags')
 if ccflags_option:
   ccflags += ccflags_option.split(' ')

cereal/custom.capnp

@@ -8,10 +8,73 @@ $Cxx.namespace("cereal");
 # cereal, so use these if you want custom events in your fork.
 
 # you can rename the struct, but don't change the identifier
-struct CustomReserved0 @0x81c2f05a394cf4af {
+struct SelfdriveStateSP @0x81c2f05a394cf4af {
+  mads @0 :ModularAssistiveDrivingSystem;
+
+  struct ModularAssistiveDrivingSystem {
+    state @0 :ModularAssistiveDrivingSystemState;
+    enabled @1 :Bool;
+    active @2 :Bool;
+    available @3 :Bool;
+
+    enum ModularAssistiveDrivingSystemState {
+      disabled @0;
+      paused @1;
+      enabled @2;
+      softDisabling @3;
+      overriding @4;
+    }
+  }
 }
 
-struct CustomReserved1 @0xaedffd8f31e7b55d {
+struct ModelManagerSP @0xaedffd8f31e7b55d {
+  activeBundle @0 :ModelBundle;
+  selectedBundle @1 :ModelBundle;
+  availableBundles @2 :List(ModelBundle);
+
+  struct DownloadUri {
+    uri @0 :Text;
+    sha256 @1 :Text;
+  }
+
+  enum Type {
+    drive @0;
+    navigation @1;
+    metadata @2;
+  }
+
+  struct Model {
+    fullName @0 :Text;
+    fileName @1 :Text;
+    downloadUri @2 :DownloadUri;
+    downloadProgress @3 :DownloadProgress;
+    type @4 :Type;
+  }
+
+  enum DownloadStatus {
+    notDownloading @0;
+    downloading @1;
+    downloaded @2;
+    cached @3;
+    failed @4;
+  }
+
+  struct DownloadProgress {
+    status @0 :DownloadStatus;
+    progress @1 :Float32;
+    eta @2 :UInt32;
+  }
+
+  struct ModelBundle {
+    index @0 :UInt32;
+    internalName @1 :Text;
+    displayName @2 :Text;
+    models @3 :List(Model);
+    status @4 :DownloadStatus;
+    generation @5 :UInt32;
+    environment @6 :Text;
+    is20hz @7 :Bool;
+  }
 }
 
 struct CustomReserved2 @0xf35cc4560bbf6ec2 {

cereal/log.capnp

@@ -125,6 +125,79 @@ struct OnroadEvent @0xc4fa6047f024e718 {
     espActive @90;
     personalityChanged @91;
     aeb @92;
+    eventReserved93 @93;
+    eventReserved94 @94;
+    eventReserved95 @95;
+    eventReserved96 @96;
+    eventReserved97 @97;
+    eventReserved98 @98;
+    eventReserved99 @99;
+    eventReserved100 @100;
+    eventReserved101 @101;
+    eventReserved102 @102;
+    eventReserved103 @103;
+    eventReserved104 @104;
+    eventReserved105 @105;
+    eventReserved106 @106;
+    eventReserved107 @107;
+    eventReserved108 @108;
+    eventReserved109 @109;
+    eventReserved110 @110;
+    eventReserved111 @111;
+    eventReserved112 @112;
+    eventReserved113 @113;
+    eventReserved114 @114;
+    eventReserved115 @115;
+    eventReserved116 @116;
+    eventReserved117 @117;
+    eventReserved118 @118;
+    eventReserved119 @119;
+    eventReserved120 @120;
+    eventReserved121 @121;
+    eventReserved122 @122;
+    eventReserved123 @123;
+    eventReserved124 @124;
+    eventReserved125 @125;
+    eventReserved126 @126;
+    eventReserved127 @127;
+    eventReserved128 @128;
+    eventReserved129 @129;
+    eventReserved130 @130;
+    eventReserved131 @131;
+    eventReserved132 @132;
+    eventReserved133 @133;
+    eventReserved134 @134;
+    eventReserved135 @135;
+    eventReserved136 @136;
+    eventReserved137 @137;
+    eventReserved138 @138;
+    eventReserved139 @139;
+    eventReserved140 @140;
+    eventReserved141 @141;
+    eventReserved142 @142;
+    eventReserved143 @143;
+    eventReserved144 @144;
+    eventReserved145 @145;
+    eventReserved146 @146;
+    eventReserved147 @147;
+    eventReserved148 @148;
+    eventReserved149 @149;
+    eventReserved150 @150;
+
+    # sunnypilot
+    lkasEnable @151;
+    lkasDisable @152;
+    manualSteeringRequired @153;
+    manualLongitudinalRequired @154;
+    silentLkasEnable @155;
+    silentLkasDisable @156;
+    silentBrakeHold @157;
+    silentWrongGear @158;
+    silentReverseGear @159;
+    silentDoorOpen @160;
+    silentSeatbeltNotLatched @161;
+    silentParkBrake @162;
+    controlsMismatchLateral @163;
 
     soundsUnavailableDEPRECATED @47;
   }
@@ -589,6 +662,7 @@ struct PandaState @0xa7649e2575e4591e {
 
   # safety stuff
   controlsAllowed @3 :Bool;
+  controlsAllowedLat @5 :Bool;
   safetyRxInvalid @19 :UInt32;
   safetyTxBlocked @24 :UInt32;
   safetyModel @14 :Car.CarParams.SafetyModel;
@@ -696,7 +770,6 @@ struct PandaState @0xa7649e2575e4591e {
   }
 
   gasInterceptorDetectedDEPRECATED @4 :Bool;
-  startedSignalDetectedDEPRECATED @5 :Bool;
   hasGpsDEPRECATED @6 :Bool;
   gmlanSendErrsDEPRECATED @9 :UInt32;
   fanSpeedRpmDEPRECATED @11 :UInt16;
@@ -2558,8 +2631,8 @@ struct Event {
     customReservedRawData2 @126 :Data;
 
     # *********** Custom: reserved for forks ***********
-    customReserved0 @107 :Custom.CustomReserved0;
-    customReserved1 @108 :Custom.CustomReserved1;
+    selfdriveStateSP @107 :Custom.SelfdriveStateSP;
+    modelManagerSP @108 :Custom.ModelManagerSP;
     customReserved2 @109 :Custom.CustomReserved2;
     customReserved3 @110 :Custom.CustomReserved3;
     customReserved4 @111 :Custom.CustomReserved4;

cereal/maptile.capnp

@@ -1,49 +0,0 @@
-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);
-}

cereal/services.py

@@ -74,6 +74,10 @@ _services: dict[str, tuple] = {
   "userFlag": (True, 0., 1),
   "microphone": (True, 10., 10),
 
+  # sunnypilot
+  "modelManagerSP": (False, 1., 1),
+  "selfdriveStateSP": (True, 100., 10),
+
   # debug
   "uiDebug": (True, 0., 1),
   "testJoystick": (True, 0.),

common/api/__init__.py

@@ -1,46 +1,27 @@
-import jwt
-import os
-import requests
-from datetime import datetime, timedelta, UTC
-from openpilot.system.hardware.hw import Paths
-from openpilot.system.version import get_version
+from openpilot.common.api.comma_connect import CommaConnectApi
+from sunnypilot.sunnylink.api import SunnylinkApi
 
-API_HOST = os.getenv('API_HOST', 'https://api.commadotai.com')
 
 class Api:
-  def __init__(self, dongle_id):
-    self.dongle_id = dongle_id
-    with open(Paths.persist_root()+'/comma/id_rsa') as f:
-      self.private_key = f.read()
+  def __init__(self, dongle_id, use_sunnylink=False):
+    if use_sunnylink:
+      self.service = SunnylinkApi(dongle_id)
+    else:
+      self.service = CommaConnectApi(dongle_id)
+
+  def request(self, method, endpoint, **params):
+    return self.service.request(method, endpoint, **params)
 
   def get(self, *args, **kwargs):
-    return self.request('GET', *args, **kwargs)
+    return self.service.get(*args, **kwargs)
 
   def post(self, *args, **kwargs):
-    return self.request('POST', *args, **kwargs)
-
-  def request(self, method, endpoint, timeout=None, access_token=None, **params):
-    return api_get(endpoint, method=method, timeout=timeout, access_token=access_token, **params)
+    return self.service.post(*args, **kwargs)
 
   def get_token(self, expiry_hours=1):
-    now = datetime.now(UTC).replace(tzinfo=None)
-    payload = {
-      'identity': self.dongle_id,
-      'nbf': now,
-      'iat': now,
-      'exp': now + timedelta(hours=expiry_hours)
-    }
-    token = jwt.encode(payload, self.private_key, algorithm='RS256')
-    if isinstance(token, bytes):
-      token = token.decode('utf8')
-    return token
+    return self.service.get_token(expiry_hours)
 
 
-def api_get(endpoint, method='GET', timeout=None, access_token=None, **params):
-  headers = {}
-  if access_token is not None:
-    headers['Authorization'] = "JWT " + access_token
-
-  headers['User-Agent'] = "openpilot-" + get_version()
-
-  return requests.request(method, API_HOST + "/" + endpoint, timeout=timeout, headers=headers, params=params)
+def api_get(endpoint, method='GET', timeout=None, access_token=None, use_sunnylink=False, **params):
+  return SunnylinkApi(None).api_get(endpoint, method, timeout, access_token, **params) if use_sunnylink \
+    else CommaConnectApi(None).api_get(endpoint, method, timeout, access_token, **params)

common/api/base.py

@@ -0,0 +1,56 @@
+import jwt
+import requests
+import unicodedata
+from datetime import datetime, timedelta, UTC
+from openpilot.system.hardware.hw import Paths
+from openpilot.system.version import get_version
+
+
+class BaseApi:
+  def __init__(self, dongle_id, api_host, user_agent="openpilot-"):
+    self.dongle_id = dongle_id
+    self.api_host = api_host
+    self.user_agent = user_agent
+    with open(f'{Paths.persist_root()}/comma/id_rsa') as f:
+      self.private_key = f.read()
+
+  def get(self, *args, **kwargs):
+    return self.request('GET', *args, **kwargs)
+
+  def post(self, *args, **kwargs):
+    return self.request('POST', *args, **kwargs)
+
+  def request(self, method, endpoint, timeout=None, access_token=None, **params):
+    return self.api_get(endpoint, method=method, timeout=timeout, access_token=access_token, **params)
+
+  def _get_token(self, expiry_hours=1, **extra_payload):
+    now = datetime.now(UTC).replace(tzinfo=None)
+    payload = {
+      'identity': self.dongle_id,
+      'nbf': now,
+      'iat': now,
+      'exp': now + timedelta(hours=expiry_hours),
+      **extra_payload
+    }
+    token = jwt.encode(payload, self.private_key, algorithm='RS256')
+    if isinstance(token, bytes):
+      token = token.decode('utf8')
+    return token
+
+  def get_token(self, expiry_hours=1):
+    return self._get_token(expiry_hours)
+
+  def remove_non_ascii_chars(self, text):
+    normalized_text = unicodedata.normalize('NFD', text)
+    ascii_encoded_text = normalized_text.encode('ascii', 'ignore')
+    return ascii_encoded_text.decode()
+
+  def api_get(self, endpoint, method='GET', timeout=None, access_token=None, **params):
+    headers = {}
+    if access_token is not None:
+      headers['Authorization'] = "JWT " + access_token
+
+    version = self.remove_non_ascii_chars(get_version())
+    headers['User-Agent'] = self.user_agent + version
+
+    return requests.request(method, f"{self.api_host}/{endpoint}", timeout=timeout, headers=headers, params=params)

common/api/comma_connect.py

@@ -0,0 +1,11 @@
+import os
+
+from openpilot.common.api.base import BaseApi
+
+API_HOST = os.getenv('API_HOST', 'https://api.commadotai.com')
+
+
+class CommaConnectApi(BaseApi):
+  def __init__(self, dongle_id):
+    super().__init__(dongle_id, API_HOST)
+    self.user_agent = "openpilot-"

common/params.cc

@@ -109,36 +109,35 @@ std::unordered_map<std::string, uint32_t> keys = {
     {"CurrentBootlog", PERSISTENT},
     {"CurrentRoute", CLEAR_ON_MANAGER_START | CLEAR_ON_ONROAD_TRANSITION},
     {"DisableLogging", CLEAR_ON_MANAGER_START | CLEAR_ON_ONROAD_TRANSITION},
-    {"DisablePowerDown", PERSISTENT},
-    {"DisableUpdates", PERSISTENT},
-    {"DisengageOnAccelerator", PERSISTENT},
+    {"DisablePowerDown", PERSISTENT | BACKUP},
+    {"DisableUpdates", PERSISTENT | BACKUP},
+    {"DisengageOnAccelerator", PERSISTENT | BACKUP},
     {"DongleId", PERSISTENT},
     {"DoReboot", CLEAR_ON_MANAGER_START},
     {"DoShutdown", CLEAR_ON_MANAGER_START},
     {"DoUninstall", CLEAR_ON_MANAGER_START},
-    {"ExperimentalLongitudinalEnabled", PERSISTENT | DEVELOPMENT_ONLY},
-    {"ExperimentalMode", PERSISTENT},
-    {"ExperimentalModeConfirmed", PERSISTENT},
+    {"ExperimentalLongitudinalEnabled", PERSISTENT | DEVELOPMENT_ONLY | BACKUP},
+    {"ExperimentalMode", PERSISTENT | BACKUP},
+    {"ExperimentalModeConfirmed", PERSISTENT | BACKUP},
     {"FirmwareQueryDone", CLEAR_ON_MANAGER_START | CLEAR_ON_ONROAD_TRANSITION},
     {"ForcePowerDown", PERSISTENT},
     {"GitBranch", PERSISTENT},
     {"GitCommit", PERSISTENT},
     {"GitCommitDate", PERSISTENT},
     {"GitDiff", PERSISTENT},
-    {"GithubSshKeys", PERSISTENT},
-    {"GithubUsername", PERSISTENT},
+    {"GithubSshKeys", PERSISTENT | BACKUP},
+    {"GithubUsername", PERSISTENT | BACKUP},
     {"GitRemote", PERSISTENT},
-    {"GsmApn", PERSISTENT},
-    {"GsmMetered", PERSISTENT},
-    {"GsmRoaming", PERSISTENT},
+    {"GsmApn", PERSISTENT | BACKUP},
+    {"GsmMetered", PERSISTENT | BACKUP},
+    {"GsmRoaming", PERSISTENT | BACKUP},
     {"HardwareSerial", PERSISTENT},
     {"HasAcceptedTerms", PERSISTENT},
-    {"IMEI", PERSISTENT},
     {"InstallDate", PERSISTENT},
     {"IsDriverViewEnabled", CLEAR_ON_MANAGER_START},
     {"IsEngaged", PERSISTENT},
-    {"IsLdwEnabled", PERSISTENT},
-    {"IsMetric", PERSISTENT},
+    {"IsLdwEnabled", PERSISTENT | BACKUP},
+    {"IsMetric", PERSISTENT | BACKUP},
     {"IsOffroad", CLEAR_ON_MANAGER_START},
     {"IsOnroad", PERSISTENT},
     {"IsRhdDetected", PERSISTENT},
@@ -146,7 +145,7 @@ std::unordered_map<std::string, uint32_t> keys = {
     {"IsTakingSnapshot", CLEAR_ON_MANAGER_START},
     {"IsTestedBranch", CLEAR_ON_MANAGER_START},
     {"JoystickDebugMode", CLEAR_ON_MANAGER_START | CLEAR_ON_OFFROAD_TRANSITION},
-    {"LanguageSetting", PERSISTENT},
+    {"LanguageSetting", PERSISTENT | BACKUP},
     {"LastAthenaPingTime", CLEAR_ON_MANAGER_START},
     {"LastGPSPosition", PERSISTENT},
     {"LastManagerExitReason", CLEAR_ON_MANAGER_START},
@@ -158,7 +157,7 @@ std::unordered_map<std::string, uint32_t> keys = {
     {"LiveTorqueParameters", PERSISTENT | DONT_LOG},
     {"LocationFilterInitialState", PERSISTENT},
     {"LongitudinalManeuverMode", CLEAR_ON_MANAGER_START | CLEAR_ON_OFFROAD_TRANSITION},
-    {"LongitudinalPersonality", PERSISTENT},
+    {"LongitudinalPersonality", PERSISTENT | BACKUP},
     {"NetworkMetered", PERSISTENT},
     {"ObdMultiplexingChanged", CLEAR_ON_MANAGER_START | CLEAR_ON_ONROAD_TRANSITION},
     {"ObdMultiplexingEnabled", CLEAR_ON_MANAGER_START | CLEAR_ON_ONROAD_TRANSITION},
@@ -174,17 +173,17 @@ std::unordered_map<std::string, uint32_t> keys = {
     {"Offroad_TemperatureTooHigh", CLEAR_ON_MANAGER_START},
     {"Offroad_UnofficialHardware", CLEAR_ON_MANAGER_START},
     {"Offroad_UpdateFailed", CLEAR_ON_MANAGER_START},
-    {"OpenpilotEnabledToggle", PERSISTENT},
+    {"OpenpilotEnabledToggle", PERSISTENT | BACKUP},
     {"PandaHeartbeatLost", CLEAR_ON_MANAGER_START | CLEAR_ON_OFFROAD_TRANSITION},
     {"PandaSomResetTriggered", CLEAR_ON_MANAGER_START | CLEAR_ON_OFFROAD_TRANSITION},
     {"PandaSignatures", CLEAR_ON_MANAGER_START},
     {"PrimeType", PERSISTENT},
-    {"RecordFront", PERSISTENT},
+    {"RecordFront", PERSISTENT | BACKUP},
     {"RecordFrontLock", PERSISTENT},  // for the internal fleet
-    {"SecOCKey", PERSISTENT | DONT_LOG},
+    {"SecOCKey", PERSISTENT | DONT_LOG},  // Candidate for | BACKUP
     {"RouteCount", PERSISTENT},
     {"SnoozeUpdate", CLEAR_ON_MANAGER_START | CLEAR_ON_OFFROAD_TRANSITION},
-    {"SshEnabled", PERSISTENT},
+    {"SshEnabled", PERSISTENT | BACKUP},
     {"TermsVersion", PERSISTENT},
     {"TrainingVersion", PERSISTENT},
     {"UbloxAvailable", PERSISTENT},
@@ -200,7 +199,28 @@ std::unordered_map<std::string, uint32_t> keys = {
     {"UpdaterTargetBranch", CLEAR_ON_MANAGER_START},
     {"UpdaterLastFetchTime", PERSISTENT},
     {"Version", PERSISTENT},
-    {"EnableGithubRunner", PERSISTENT},
+
+    // --- sunnypilot params --- //
+    {"EnableGithubRunner", PERSISTENT | BACKUP},
+
+    // MADS params
+    {"Mads", PERSISTENT | BACKUP},
+    {"MadsMainCruiseAllowed", PERSISTENT | BACKUP},
+    {"MadsPauseLateralOnBrake", PERSISTENT | BACKUP},
+    {"MadsUnifiedEngagementMode", PERSISTENT | BACKUP},
+
+    // Model Manager params
+    {"ModelManager_ActiveBundle", PERSISTENT},
+    {"ModelManager_DownloadIndex", CLEAR_ON_MANAGER_START | CLEAR_ON_ONROAD_TRANSITION},
+    {"ModelManager_LastSyncTime", CLEAR_ON_MANAGER_START | CLEAR_ON_OFFROAD_TRANSITION},
+    {"ModelManager_ModelsCache", PERSISTENT | BACKUP},
+
+    // sunnylink params
+    {"EnableSunnylinkUploader", PERSISTENT | BACKUP},
+    {"LastSunnylinkPingTime", CLEAR_ON_MANAGER_START},
+    {"SunnylinkDongleId", PERSISTENT},
+    {"SunnylinkdPid", PERSISTENT},
+    {"SunnylinkEnabled", PERSISTENT},
 };
 
 } // namespace

common/params.h

@@ -16,6 +16,7 @@ enum ParamKeyType {
   CLEAR_ON_OFFROAD_TRANSITION = 0x10,
   DONT_LOG = 0x20,
   DEVELOPMENT_ONLY = 0x40,
+  BACKUP = 0x80,
   ALL = 0xFFFFFFFF
 };
 

common/util.cc

@@ -255,6 +255,29 @@ 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;

common/util.h

@@ -74,6 +74,7 @@ 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);

docs/CARS.md

@@ -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.
 
-# 289 Supported Cars
+# 290 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,6 +56,7 @@ 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>|
@@ -79,7 +80,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-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>||
+|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>||
 |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>||
@@ -91,7 +92,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 (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 (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 (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>||
@@ -294,7 +295,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-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 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 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>||
 

pyproject.toml

@@ -42,8 +42,7 @@ dependencies = [
 
   # modeld
   "onnx >= 1.14.0",
-  "onnxruntime >=1.16.3; platform_system == 'Linux' and platform_machine == 'aarch64'",
-  "onnxruntime-gpu >=1.16.3; platform_system == 'Linux' and platform_machine == 'x86_64'",
+  "onnxruntime >=1.16.3",
 
   # logging
   "pyzmq",
@@ -104,7 +103,6 @@ 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
@@ -113,13 +111,10 @@ 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/metadrive_simulator-0.4.2.3-py3-none-any.whl ; (platform_machine != 'aarch64')",
+  "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')",
   "rerun-sdk >= 0.18",
 ]
 
@@ -164,6 +159,7 @@ testpaths = [
   "tools/replay",
   "tools/cabana",
   "cereal/messaging/tests",
+  "sunnypilot",
 ]
 
 [tool.codespell]

release/ci/squash_and_merge.py

@@ -0,0 +1,361 @@
+#!/usr/bin/env python3
+import argparse
+import subprocess
+import sys
+import shutil
+import signal
+import contextlib
+import tempfile
+import os
+
+
+def run_command(command: str) -> tuple[int, str, str]:
+  """Run a shell command and return exit code, stdout, and stderr."""
+  process = subprocess.Popen(
+    command,
+    shell=True,
+    stdout=subprocess.PIPE,
+    stderr=subprocess.PIPE,
+    text=True
+  )
+  stdout, stderr = process.communicate()
+  return process.returncode, stdout.strip(), stderr.strip()
+
+
+def is_gh_available() -> bool:
+  """Check if GitHub CLI is available."""
+  return shutil.which('gh') is not None
+
+
+def get_current_branch() -> str | None:
+  """Get the name of the current git branch."""
+  code, output, error = run_command("git rev-parse --abbrev-ref HEAD")
+  if code != 0:
+    print(f"Error getting current branch: {error}")
+    return None
+  return output
+
+
+def backup_branch(branch_name: str) -> bool:
+  """Create a backup of the current branch."""
+  backup_name = f"{branch_name}-backup-$(date +%Y%m%d_%H%M%S)"
+  code, _, error = run_command(f"git branch {backup_name}")
+  if code != 0:
+    print(f"Error creating backup branch: {error}")
+    return False
+  print(f"Created backup branch: {backup_name}")
+  return True
+
+
+def get_commit_messages(source_branch: str, target_branch: str) -> list[str] | None:
+  """Get all commit messages between source and target branches."""
+  code, output, error = run_command(f"git log {target_branch}..{source_branch} --format=%B")
+  if code != 0:
+    print(f"Error getting commit messages: {error}")
+    return None
+  return [msg.strip() for msg in output.splitlines() if msg and not msg.startswith('Merge')]
+
+
+def get_pr_info(branch_name: str) -> str | None:
+  """Get PR title using GitHub CLI."""
+  if not is_gh_available():
+    print("Warning: GitHub CLI not found. Install it to auto-fetch PR titles:")
+    print("  https://cli.github.com/")
+    return None
+
+  # Try to get PR info using gh cli
+  code, output, error = run_command(f"gh pr view --json title --jq .title {branch_name}")
+  if code != 0:
+    print(f"No open PR found for branch '{branch_name}'")
+    return None
+
+  return output
+
+
+def create_squash_message(pr_title: str | None, commit_messages: list[str], source_branch: str) -> str:
+  """Create a squash commit message from PR title and commit messages."""
+  parts = []
+
+  # Add PR title if provided
+  if pr_title:
+    parts.append(pr_title)
+  else:
+    parts.append(f"Squashed changes from {source_branch}")
+  parts.append("")  # Empty line after title
+
+  # Add original commits section
+  if commit_messages:
+    parts.append("Original commits:")
+    parts.append("")  # Empty line before list
+    parts.extend(f"* {msg}" for msg in commit_messages)
+
+  return '\n'.join(parts)
+
+
+def prompt_for_title() -> str:
+  """Prompt user for a commit title."""
+  return input("Enter commit title (or press Enter to use default): ").strip()
+
+
+@contextlib.contextmanager
+def workspace_manager(original_branch: str):
+  """Context manager to handle workspace state and cleanup."""
+  stash_created = False
+  stash_restored = False
+  temp_branch: str | None = None
+
+  def cleanup_handler(signum=None, frame=None):
+    """Clean up workspace state."""
+    nonlocal temp_branch, stash_created, stash_restored
+    try:
+      if signum and stash_restored:
+        # If we're handling Ctrl+C but stash was already restored,
+        # just clean up branches and exit
+        current = get_current_branch()
+        if current and current != original_branch:
+          run_command(f"git checkout {original_branch}")
+        if temp_branch:
+          run_command(f"git branch -D {temp_branch}")
+        print("\nOperation interrupted, but changes were already restored.")
+        sys.exit(1)
+
+      # First, switch back to original branch
+      current = get_current_branch()
+      if current and current != original_branch:
+        run_command(f"git checkout {original_branch}")
+
+      # Then clean up temp branch
+      if temp_branch:
+        run_command(f"git branch -D {temp_branch}")
+
+      # Finally, restore stash if needed - AFTER switching branches
+      if stash_created and not stash_restored:
+        print("Restoring your uncommitted changes...")
+        code, stash_list, _ = run_command("git stash list")
+        if code == 0 and "Automatic stash by squash script" in stash_list:
+          run_command("git stash pop")
+          stash_restored = True
+          stash_created = False
+
+      if signum:
+        print("\nOperation interrupted. Cleaned up and restored original state.")
+        sys.exit(1)
+
+    except Exception as e:
+      print(f"Error during cleanup: {e}")
+      if signum:
+        sys.exit(1)
+
+  try:
+    # Set up signal handlers
+    signal.signal(signal.SIGINT, cleanup_handler)
+    signal.signal(signal.SIGTERM, cleanup_handler)
+
+    # Check for changes (including untracked files)
+    code, output, _ = run_command("git status --porcelain")
+    if output:
+      print("Stashing uncommitted changes...")
+      run_command("git stash push -u -m 'Automatic stash by squash script'")
+      stash_created = True
+
+    yield lambda x: setattr(x, 'temp_branch', temp_branch)
+
+  except Exception as e:
+    print(f"\nError occurred: {str(e)}")
+    cleanup_handler()
+    raise
+  finally:
+    cleanup_handler()
+
+
+def create_commit_with_message(message: str) -> bool:
+  """Create a commit with the given message using a temporary file."""
+  try:
+    with tempfile.NamedTemporaryFile(mode='w', delete=False) as f:
+      f.write(message)
+      temp_path = f.name
+
+    # Use the temporary file for the commit message
+    code, _, error = run_command(f"git commit -F {temp_path}")
+    os.unlink(temp_path)  # Clean up the temp file
+
+    if code != 0:
+      print(f"Error creating commit: {error}")
+      return False
+    return True
+  except Exception as e:
+    print(f"Error handling commit message: {e}")
+    if os.path.exists(temp_path):
+      os.unlink(temp_path)
+    return False
+
+
+def squash_and_merge(source_branch: str, target_branch: str, manual_title: str | None, backup: bool = False, push: bool = False) -> bool:
+  """
+  Squash the source branch and merge into target branch.
+  """
+  # Get original branch right away
+  original_branch = get_current_branch()
+  if not original_branch:
+    return False
+
+  class State:
+    temp_branch: str | None = None
+
+  state = State()
+
+  with workspace_manager(original_branch) as set_temp_branch:
+    # Validate source branch exists
+    code, _, error = run_command(f"git rev-parse --verify {source_branch}")
+    if code != 0:
+      print(f"Error: Source branch {source_branch} not found")
+      return False
+
+    if source_branch == target_branch:
+      print(f"Error: Source and target branches cannot be the same ({source_branch})")
+      return False
+
+    # Ensure target branch exists
+    code, _, error = run_command(f"git rev-parse --verify {target_branch}")
+    if code != 0:
+      print(f"Error: Target branch {target_branch} not found")
+      return False
+
+    # Find merge base
+    code, merge_base, error = run_command(f"git merge-base {target_branch} {source_branch}")
+    if code != 0:
+      print(f"Error finding merge base: {error}")
+      return False
+
+    # Create backup unless explicitly skipped
+    if backup and not backup_branch(source_branch):
+      return False
+
+    # Get commit messages
+    commit_messages = get_commit_messages(source_branch, target_branch)
+    if commit_messages is None:
+      return False
+
+    # Get title (priority: manual title > PR title > prompt user)
+    title = manual_title
+    if not title:
+      title = get_pr_info(source_branch)
+    if not title:
+      title = prompt_for_title()
+
+    try:
+      # Create and switch to temporary branch
+      temp_branch = f"temp-squash-{source_branch}"
+      state.temp_branch = temp_branch
+      set_temp_branch(state)
+
+      print(f"\nCreating temporary branch {temp_branch}...")
+      code, _, error = run_command(f"git checkout -b {temp_branch} {source_branch}")
+      if code != 0:
+        print(f"Error creating temp branch: {error}")
+        return False
+
+      print("Preparing squash by resetting temporary branch to merge base...")
+      code, _, error = run_command(f"git reset --soft {merge_base}")
+      if code != 0:
+        print(f"Error resetting for squash: {error}")
+        return False
+
+      # Create commit with message
+      print("Creating squash commit...")
+      squash_message = create_squash_message(title, commit_messages, source_branch)
+      if not create_commit_with_message(squash_message):
+        return False
+
+      # Switch to target and try merge
+      print(f"\nSwitching to target branch {target_branch}...")
+      code, _, error = run_command(f"git checkout {target_branch}")
+      if code != 0:
+        print(f"Error checking out target branch: {error}")
+        return False
+
+      print(f"Attempting to merge changes from {temp_branch}...")
+      code, _, error = run_command(f"git merge {temp_branch}")
+
+      if code != 0:
+        print(f"\nMerge failed with error: {error}")
+        print("\nThe squash was successful, and your changes are preserved in the temporary branch.")
+        print("To complete the merge manually, follow these steps:")
+        print(f"\n1. Your squashed changes are in branch: '{temp_branch}'")
+        print(f"2. The target branch is: '{target_branch}'")
+        print("\nTo resolve the conflicts:")
+        print(f"   git checkout {target_branch}")
+        print(f"   git merge {temp_branch}")
+        print("   # resolve conflicts in your editor")
+        print("   git add <resolved-files>")
+        print("   git commit")
+        print(f"   git push origin {target_branch}  # when ready to push")
+        print("\nTo clean up after successful merge:")
+        print(f"   git branch -D {temp_branch}")
+
+        # Make sure to abort the merge
+        print("\nAborting current merge attempt...")
+        run_command("git merge --abort")
+
+        # Return to original branch, but keep temp branch
+        print(f"Returning to {original_branch}...")
+        run_command(f"git checkout {original_branch}")
+        return False
+
+      # Clean up temp branch on success
+      run_command(f"git branch -D {temp_branch}")
+
+      # Push if requested
+      if push:
+        code, _, error = run_command(f"git push origin {target_branch}")
+        if code != 0:
+          print(f"Error pushing to {target_branch}: {error}")
+          return False
+        print(f"Successfully pushed to {target_branch}")
+      else:
+        print(f"Changes squashed and merged into {target_branch} locally")
+        print(f"To push the changes: git push origin {target_branch}")
+
+      # Return to original branch
+      code, _, error = run_command(f"git checkout {original_branch}")
+      if code != 0:
+        print(f"Warning: Failed to return to original branch: {error}")
+        return False
+
+      return True
+
+    except Exception as e:
+      print(f"Error during squash process: {e}")
+      return False
+
+
+def main():
+  parser = argparse.ArgumentParser(
+    description='Squash branch and merge into target branch'
+  )
+  parser.add_argument('--target', '-t', required=True,
+                      help='Target branch to merge changes into')
+  parser.add_argument('--source', '-s',
+                      help='Source branch to squash (default: current branch)')
+  parser.add_argument('--title', '-m',
+                      help='Optional manual title (overrides PR title)')
+  parser.add_argument('--backup', action='store_true',
+                      help='Creates a backup branch for the source branch')
+  parser.add_argument('--push', action='store_true',
+                      help='Push changes to remote after squashing')
+
+  args = parser.parse_args()
+
+  # Determine source branch early
+  source_branch = args.source
+  if not source_branch:
+    source_branch = get_current_branch()
+    if not source_branch:
+      sys.exit(1)
+
+  if not squash_and_merge(source_branch, args.target, args.title, args.backup, args.push):
+    sys.exit(1)
+
+
+if __name__ == "__main__":
+  main()

release/release_files.py

@@ -91,7 +91,7 @@ whitelist = [
   "tools/joystick/",
   "tools/longitudinal_maneuvers/",
 
-  "tinygrad_repo/openpilot/compile2.py",
+  "tinygrad_repo/examples/openpilot/compile3.py",
   "tinygrad_repo/extra/onnx.py",
   "tinygrad_repo/extra/onnx_ops.py",
   "tinygrad_repo/extra/thneed.py",

scripts/build_small.sh

@@ -1,39 +0,0 @@
-#!/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 .)"

scripts/code_stats.py

@@ -1,54 +0,0 @@
-#!/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)))

scripts/count_cars.py

@@ -1,11 +0,0 @@
-#!/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)

scripts/git_rewrite/rewrite-git-history.sh

@@ -1,391 +0,0 @@
-#!/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/*

scripts/git_rewrite/rewrite.sh

@@ -1,59 +0,0 @@
-#!/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/*

scripts/git_rewrite/tag-commit-map.txt

@@ -1,82 +0,0 @@
-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

scripts/lint/lint.sh

@@ -57,7 +57,7 @@ function run_tests() {
 
   if [[ -z "$FAST" ]]; then
     run "mypy" mypy $PYTHON_FILES
-    run "codespell" codespell $ALL_FILES
+    run "codespell" codespell $ALL_FILES --ignore-words=$ROOT/.codespellignore
   fi
 
   return $FAILED

scripts/pyqt_demo.py

@@ -1,14 +0,0 @@
-#!/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_()

selfdrive/assets/offroad/icon_plus.png

@@ -1,3 +0,0 @@
-version https://git-lfs.github.com/spec/v1
-oid sha256:fa3f1c39a4e82adfb52d43fc0ad6773a70dbaa4fc79109a7d6b6c1f73b298eac
-size 2833

selfdrive/car/card.py

@@ -22,6 +22,8 @@ from openpilot.selfdrive.pandad import can_capnp_to_list, can_list_to_can_capnp
 from openpilot.selfdrive.car.cruise import VCruiseHelper
 from openpilot.selfdrive.car.car_specific import MockCarState
 
+from openpilot.sunnypilot.mads.mads import MadsParams
+
 REPLAY = "REPLAY" in os.environ
 
 EventName = log.OnroadEvent.EventName
@@ -113,6 +115,10 @@ class Car:
     if not disengage_on_accelerator:
       self.CP.alternativeExperience |= ALTERNATIVE_EXPERIENCE.DISABLE_DISENGAGE_ON_GAS
 
+    # mads
+    MadsParams().set_alternative_experience(self.CP)
+    MadsParams().set_car_specific_params(self.CP)
+
     openpilot_enabled_toggle = self.params.get_bool("OpenpilotEnabledToggle")
 
     controller_available = self.CI.CC is not None and openpilot_enabled_toggle and not self.CP.dashcamOnly

selfdrive/controls/controlsd.py

@@ -19,6 +19,7 @@ from openpilot.selfdrive.controls.lib.longcontrol import LongControl
 from openpilot.selfdrive.controls.lib.vehicle_model import VehicleModel
 from openpilot.selfdrive.locationd.helpers import PoseCalibrator, Pose
 
+from opendbc.sunnypilot import SunnypilotParamFlags
 
 State = log.SelfdriveState.OpenpilotState
 LaneChangeState = log.LaneChangeState
@@ -56,6 +57,9 @@ class Controls:
     elif self.CP.lateralTuning.which() == 'torque':
       self.LaC = LatControlTorque(self.CP, self.CI)
 
+    data_services = list(self.sm.data.keys()) + ['selfdriveStateSP']
+    self.sm = messaging.SubMaster(data_services, poll='selfdriveState')
+
   def update(self):
     self.sm.update(15)
     if self.sm.updated["liveCalibration"]:
@@ -88,7 +92,16 @@ class Controls:
 
     # Check which actuators can be enabled
     standstill = abs(CS.vEgo) <= max(self.CP.minSteerSpeed, MIN_LATERAL_CONTROL_SPEED) or CS.standstill
-    CC.latActive = self.sm['selfdriveState'].active and not CS.steerFaultTemporary and not CS.steerFaultPermanent and not standstill
+
+    ss_sp = self.sm['selfdriveStateSP']
+    CC.madsEnabled = ss_sp.mads.enabled
+    if ss_sp.mads.available:
+      CC.sunnypilotParams |= SunnypilotParamFlags.ENABLE_MADS.value
+      _lat_active = ss_sp.mads.active
+    else:
+      _lat_active = self.sm['selfdriveState'].active
+
+    CC.latActive = _lat_active and not CS.steerFaultTemporary and not CS.steerFaultPermanent and not standstill
     CC.longActive = CC.enabled and not any(e.overrideLongitudinal for e in self.sm['onroadEvents']) and self.CP.openpilotLongitudinalControl
 
     actuators = CC.actuators

selfdrive/debug/touch_replay.py

@@ -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] = h - (h - event.value)
+          fingers[current_slot][1] = event.value
           if fingers[current_slot][0] != -1:
             touch_points.append(fingers[current_slot].copy())
         elif event.type == 3 and event.code == 54:

selfdrive/modeld/SConscript

@@ -13,20 +13,6 @@ common_src = [
   "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":
@@ -40,39 +26,25 @@ 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
-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)
-
-tinygrad_files = ["#"+x for x in glob.glob(env.Dir("#tinygrad_repo").relpath + "/**", recursive=True, root_dir=env.Dir("#").abspath)]
+tinygrad_files = ["#"+x for x in glob.glob(env.Dir("#tinygrad_repo").relpath + "/**", recursive=True, root_dir=env.Dir("#").abspath) if 'pycache' not in x]
 
 # Get model metadata
 fn = File("models/supercombo").abspath
 cmd = f'python3 {Dir("#selfdrive/modeld").abspath}/get_model_metadata.py {fn}.onnx'
 lenv.Command(fn + "_metadata.pkl", [fn + ".onnx"] + tinygrad_files, cmd)
 
-# Build thneed model
-if arch == "larch64" or GetOption('pc_thneed'):
-  tinygrad_opts = []
-  if not GetOption('pc_thneed'):
-    # use FLOAT16 on device for speed + don't cache the CL kernels for space
-    tinygrad_opts += ["FLOAT16=1", "PYOPENCL_NO_CACHE=1"]
-  cmd = f"cd {Dir('#').abspath}/tinygrad_repo && " + ' '.join(tinygrad_opts) + f" python3 openpilot/compile2.py {fn}.onnx {fn}.thneed"
+# Compile tinygrad model
+pythonpath_string = 'PYTHONPATH="${PYTHONPATH}:' + env.Dir("#tinygrad_repo").abspath + '"'
+if arch == 'larch64':
+  device_string = 'QCOM=1'
+else:
+  device_string = 'CLANG=1 IMAGE=0'
 
-  lenv.Command(fn + ".thneed", [fn + ".onnx"] + tinygrad_files, cmd)
+for model_name in ['supercombo', 'dmonitoring_model']:
+  fn = File(f"models/{model_name}").abspath
+  cmd = f'{pythonpath_string} {device_string} python3 {Dir("#tinygrad_repo").abspath}/examples/openpilot/compile3.py {fn}.onnx {fn}_tinygrad.pkl'
+  lenv.Command(fn + "_tinygrad.pkl", [fn + ".onnx"] + tinygrad_files, cmd)
 
-  fn_dm = File("models/dmonitoring_model").abspath
-  cmd = f"cd {Dir('#').abspath}/tinygrad_repo && " + ' '.join(tinygrad_opts) + f" python3 openpilot/compile2.py {fn_dm}.onnx {fn_dm}.thneed"
-  lenv.Command(fn_dm + ".thneed", [fn_dm + ".onnx"] + tinygrad_files, cmd)
-
-  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'])

selfdrive/modeld/constants.py

@@ -16,7 +16,6 @@ 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
@@ -73,14 +72,13 @@ class Plan:
 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)
+  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)
   # 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)
+  LEFT_BLINKER = slice(36, 48, 2)
+  RIGHT_BLINKER = slice(37, 48, 2)

selfdrive/modeld/dmonitoringmodeld

@@ -1,10 +1,4 @@
 #!/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/dmonitoringmodeld.py" "$@"

selfdrive/modeld/dmonitoringmodeld.py

@@ -1,8 +1,17 @@
 #!/usr/bin/env python3
 import os
+from openpilot.system.hardware import TICI
+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 gc
 import math
 import time
+import pickle
 import ctypes
 import numpy as np
 from pathlib import Path
@@ -13,21 +22,20 @@ from cereal.messaging import PubMaster, SubMaster
 from msgq.visionipc import VisionIpcClient, VisionStreamType, VisionBuf
 from openpilot.common.swaglog import cloudlog
 from openpilot.common.realtime import set_realtime_priority
-from openpilot.common.transformations.model import dmonitoringmodel_intrinsics
+from openpilot.common.transformations.model import dmonitoringmodel_intrinsics, DM_INPUT_SIZE
 from openpilot.common.transformations.camera import _ar_ox_fisheye, _os_fisheye
 from openpilot.selfdrive.modeld.models.commonmodel_pyx import CLContext, MonitoringModelFrame
-from openpilot.selfdrive.modeld.runners import ModelRunner, Runtime
 from openpilot.selfdrive.modeld.parse_model_outputs import sigmoid
 
+MODEL_WIDTH, MODEL_HEIGHT = DM_INPUT_SIZE
 CALIB_LEN = 3
 FEATURE_LEN = 512
 OUTPUT_SIZE = 84 + FEATURE_LEN
 
 PROCESS_NAME = "selfdrive.modeld.dmonitoringmodeld"
 SEND_RAW_PRED = os.getenv('SEND_RAW_PRED')
-MODEL_PATHS = {
-  ModelRunner.THNEED: Path(__file__).parent / 'models/dmonitoring_model.thneed',
-  ModelRunner.ONNX: Path(__file__).parent / 'models/dmonitoring_model.onnx'}
+MODEL_PATH = Path(__file__).parent / 'models/dmonitoring_model.onnx'
+MODEL_PKL_PATH = Path(__file__).parent / 'models/dmonitoring_model_tinygrad.pkl'
 
 class DriverStateResult(ctypes.Structure):
   _fields_ = [
@@ -58,29 +66,42 @@ class DMonitoringModelResult(ctypes.Structure):
 class ModelState:
   inputs: dict[str, np.ndarray]
   output: np.ndarray
-  model: ModelRunner
 
   def __init__(self, cl_ctx):
     assert ctypes.sizeof(DMonitoringModelResult) == OUTPUT_SIZE * ctypes.sizeof(ctypes.c_float)
 
     self.frame = MonitoringModelFrame(cl_ctx)
-    self.output = np.zeros(OUTPUT_SIZE, dtype=np.float32)
-    self.inputs = {
-      'calib': np.zeros(CALIB_LEN, dtype=np.float32)}
+    self.numpy_inputs = {
+      'calib': np.zeros((1, CALIB_LEN), dtype=np.float32),
+    }
 
-    self.model = ModelRunner(MODEL_PATHS, self.output, Runtime.GPU, False, cl_ctx)
-    self.model.addInput("input_img", None)
-    self.model.addInput("calib", self.inputs['calib'])
+    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)
 
   def run(self, buf:VisionBuf, calib:np.ndarray, transform:np.ndarray) -> tuple[np.ndarray, float]:
-    self.inputs['calib'][:] = calib
-
-    self.model.setInputBuffer("input_img", self.frame.prepare(buf, transform.flatten(), None).view(np.float32))
+    self.numpy_inputs['calib'][0,:] = calib
 
     t1 = time.perf_counter()
-    self.model.execute()
+
+    input_img_cl = self.frame.prepare(buf, transform.flatten())
+    if TICI:
+      # The imgs tensors are backed by opencl memory, only need init once
+      if 'input_img' not in self.tensor_inputs:
+        self.tensor_inputs['input_img'] = qcom_tensor_from_opencl_address(input_img_cl.mem_address, (1, MODEL_WIDTH*MODEL_HEIGHT), dtype=dtypes.uint8)
+    else:
+      self.numpy_inputs['input_img'] = self.frame.buffer_from_cl(input_img_cl).reshape((1, MODEL_WIDTH*MODEL_HEIGHT))
+
+    if TICI:
+      output = self.model_run(**self.tensor_inputs).numpy().flatten()
+    else:
+      output = self.onnx_cpu_runner.run(None, self.numpy_inputs)[0].flatten()
+
     t2 = time.perf_counter()
-    return self.output, t2 - t1
+    return output, t2 - t1
 
 
 def fill_driver_state(msg, ds_result: DriverStateResult):

selfdrive/modeld/fill_model_msg.py

@@ -3,21 +3,11 @@ 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):
@@ -75,8 +65,6 @@ 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
@@ -85,7 +73,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 = desired_curv
+  action.desiredCurvature = float(net_output_data['desired_curvature'][0,0])
 
   modelV2 = extended_msg.modelV2
   modelV2.frameId = vipc_frame_id
@@ -120,7 +108,7 @@ def fill_model_msg(base_msg: capnp._DynamicStructBuilder, extended_msg: capnp._D
 
   # lateral planning
   action = modelV2.action
-  action.desiredCurvature = desired_curv
+  action.desiredCurvature = float(net_output_data['desired_curvature'][0,0])
 
   # times at X_IDXS according to model plan
   PLAN_T_IDXS = [np.nan] * ModelConstants.IDX_N
@@ -181,8 +169,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]

selfdrive/modeld/modeld

@@ -1,10 +1,4 @@
 #!/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" "$@"

selfdrive/modeld/modeld.py

@@ -1,40 +1,36 @@
 #!/usr/bin/env python3
+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
 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.realtime import DT_MDL
-from openpilot.common.numpy_fast import interp
 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.system.hardware import PC
 from openpilot.selfdrive.controls.lib.desire_helper import DesireHelper
-from openpilot.selfdrive.modeld.runners import ModelRunner, Runtime
 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"
-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'
-
+USE_ONNX = bool(os.getenv('USE_ONNX', PC))
+IS_20HZ_MODEL_DEFAULT = False
 
 class FrameMeta:
   frame_id: int = 0
@@ -46,99 +42,92 @@ class FrameMeta:
       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
+  frames: dict[str, 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.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.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.inputs = {}
+    self.numpy_inputs = {}
 
-    with open(METADATA_PATH, 'rb') as f:
-      model_metadata = pickle.load(f)
+    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)
 
-    for key, shape in model_metadata['input_shapes'].items():
-      if key not in ["input_imgs", "big_input_imgs"]:
-        self.inputs[key] = np.zeros(shape, dtype=np.float32).flatten()
-
-    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)
+    net_output_size = self.model_runner.model_metadata['output_shapes']['outputs'][1]
+    self.output = np.zeros(net_output_size, dtype=np.float32)
 
-    num_elements = model_metadata['input_shapes']['features_buffer'][1]
+    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]
-
-    desired_shape = int(self.inputs['desire'].shape[0] / self.desire_20Hz.shape[1])
-    middle_dim = int(self.desire_20Hz.shape[0] / desired_shape)
-    self.desire_reshape_dims = (desired_shape, middle_dim, -1)
-
-  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
+    self.desire_reshape_dims = (self.numpy_inputs['desire'].shape[0], self.numpy_inputs['desire'].shape[1], -1, self.numpy_inputs['desire'].shape[2])
 
   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']
 
-    self.desire_20Hz[:-1] = self.desire_20Hz[1:]
-    self.desire_20Hz[-1] = new_desire
-    self.inputs['desire'][:] = self.desire_20Hz.reshape(self.desire_reshape_dims).max(axis=1).flatten()
+    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
 
-    self.inputs['traffic_convention'][:] = inputs['traffic_convention']
+    for key in self.numpy_inputs:
+      if key in inputs and key not in ['desire']:
+        self.numpy_inputs[key][:] = inputs[key]
 
-    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")))
+    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 prepare_only:
       return None
 
-    self.model.execute()
-    outputs = self.parser.parse_outputs(self.slice_outputs(self.output))
+    # Run model inference
+    self.output = self.model_runner.run_model()
+    outputs = self.parser.parse_outputs(self.model_runner.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[self.full_features_20Hz_idxs].flatten()
-
-    if "lat_planner_solution" in outputs:
-      if "lat_planner_state" in self.inputs.keys():
-        self.inputs['lat_planner_state'][2] = interp(DT_MDL, ModelConstants.T_IDXS, outputs['lat_planner_solution'][0, :, 2])
-        self.inputs['lat_planner_state'][3] = interp(DT_MDL, ModelConstants.T_IDXS, outputs['lat_planner_solution'][0, :, 3])
+    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]
+    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, :]
 
     if "desired_curvature" in outputs:
       input_name_prev = None
-      if "prev_desired_curvs" in self.inputs.keys():
+
+      if "prev_desired_curvs" in self.numpy_inputs.keys():
         input_name_prev = 'prev_desired_curvs'
-      elif "prev_desired_curv" in self.inputs.keys():
+      elif "prev_desired_curv" in self.numpy_inputs.keys():
         input_name_prev = 'prev_desired_curv'
 
       if input_name_prev is not None:
         len = outputs['desired_curvature'][0].size
-        self.inputs[input_name_prev][:-len] = self.inputs[input_name_prev][len:]
-        self.inputs[input_name_prev][-len:] = outputs['desired_curvature'][0, :len]
+        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, :]
     return outputs
 
 
@@ -279,20 +268,11 @@ def main(demo=False):
     inputs:dict[str, np.ndarray] = {
       'desire': vec_desire,
       'traffic_convention': traffic_convention,
-      }
+    }
 
-    if "lateral_control_params" in model.inputs.keys():
+    if "lateral_control_params" in model.numpy_inputs.keys():
       inputs['lateral_control_params'] = np.array([v_ego, steer_delay], dtype=np.float32)
 
-    if "driving_style" in model.inputs.keys():
-      inputs['driving_style'] = np.array([1, 0, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0], dtype=np.float32)
-
-    if "nav_features" in model.inputs.keys():
-      inputs['nav_features'] = np.zeros(ModelConstants.NAV_FEATURE_LEN, dtype=np.float32)  # Get size from shape
-
-    if "nav_instructions" in model.inputs.keys():
-      inputs['nav_instructions'] = np.zeros(ModelConstants.NAV_INSTRUCTION_LEN, dtype=np.float32)  # Get size from shape
-
     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()
@@ -320,7 +300,6 @@ def main(demo=False):
       pm.send('modelV2', modelv2_send)
       pm.send('drivingModelData', drivingdata_send)
       pm.send('cameraOdometry', posenet_send)
-
     last_vipc_frame_id = meta_main.frame_id
 
 

selfdrive/modeld/models/commonmodel.cc

@@ -5,11 +5,11 @@
 
 #include "common/clutil.h"
 
-DrivingModelFrame::DrivingModelFrame(cl_device_id device_id, cl_context context) : ModelFrame(device_id, context) {
-  input_frames = std::make_unique<float[]>(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;
+DrivingModelFrame::DrivingModelFrame(cl_device_id device_id, cl_context context, bool is_20hz) : ModelFrame(device_id, context), is_20hz(is_20hz) {
+  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;
   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));
 
@@ -17,27 +17,20 @@ DrivingModelFrame::DrivingModelFrame(cl_device_id device_id, cl_context context)
   init_transform(device_id, context, MODEL_WIDTH, MODEL_HEIGHT);
 }
 
-float* 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) {
+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 < 4; i++) {
+  for (int i = 0; i < (buf_len - 1); 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);
+  copy_queue(&loadyuv, q, img_buffer_20hz_cl, input_frames_cl, 0, 0, frame_size_bytes);
+  copy_queue(&loadyuv, q, last_img_cl, input_frames_cl, 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;
-  }
+  // NOTE: Since thneed is using a different command queue, this clFinish is needed to ensure the image is ready.
+  clFinish(q);
+  return &input_frames_cl;
 }
 
 DrivingModelFrame::~DrivingModelFrame() {
@@ -50,17 +43,16 @@ DrivingModelFrame::~DrivingModelFrame() {
 
 
 MonitoringModelFrame::MonitoringModelFrame(cl_device_id device_id, cl_context context) : ModelFrame(device_id, context) {
-  input_frames = std::make_unique<float[]>(buf_size);
-  //input_frame_cl = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_WRITE, buf_size, NULL, &err));
+  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);
 }
-float* 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) {
+
+cl_mem* MonitoringModelFrame::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);
-  CL_CHECK(clEnqueueReadBuffer(q, y_cl, CL_TRUE, 0, MODEL_FRAME_SIZE * sizeof(float), input_frames.get(), 0, nullptr, nullptr));
   clFinish(q);
-  //return &y_cl;
-  return input_frames.get();
+  return &y_cl;
 }
 
 MonitoringModelFrame::~MonitoringModelFrame() {

selfdrive/modeld/models/commonmodel.h

@@ -23,14 +23,12 @@ public:
     q = CL_CHECK_ERR(clCreateCommandQueue(context, device_id, 0, &err));
   }
   virtual ~ModelFrame() {}
-  virtual float* 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; }
-  /*
+  virtual cl_mem* prepare(cl_mem yuv_cl, int frame_width, int frame_height, int frame_stride, int frame_uv_offset, const mat3& projection) { 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;
@@ -41,7 +39,7 @@ protected:
   cl_mem y_cl, u_cl, v_cl;
   Transform transform;
   cl_command_queue q;
-  std::unique_ptr<float[]> input_frames;
+  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));
@@ -66,19 +64,22 @@ protected:
 
 class DrivingModelFrame : public ModelFrame {
 public:
-  DrivingModelFrame(cl_device_id device_id, cl_context context);
+  DrivingModelFrame(cl_device_id device_id, cl_context context, bool is_20hz = false);
   ~DrivingModelFrame();
-  float* prepare(cl_mem yuv_cl, int frame_width, int frame_height, int frame_stride, int frame_uv_offset, const mat3& projection, cl_mem* output);
+  cl_mem* prepare(cl_mem yuv_cl, int frame_width, int frame_height, int frame_stride, int frame_uv_offset, const mat3& projection);
 
   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(float);
+  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;
+  cl_mem img_buffer_20hz_cl, last_img_cl, input_frames_cl;
   cl_buffer_region region;
 };
 
@@ -86,7 +87,7 @@ class MonitoringModelFrame : public ModelFrame {
 public:
   MonitoringModelFrame(cl_device_id device_id, cl_context context);
   ~MonitoringModelFrame();
-  float* prepare(cl_mem yuv_cl, int frame_width, int frame_height, int frame_stride, int frame_uv_offset, const mat3& projection, cl_mem* output);
+  cl_mem* prepare(cl_mem yuv_cl, int frame_width, int frame_height, int frame_stride, int frame_uv_offset, const mat3& projection);
 
   const int MODEL_WIDTH = 1440;
   const int MODEL_HEIGHT = 960;
@@ -94,5 +95,5 @@ public:
   const int buf_size = MODEL_FRAME_SIZE;
 
 private:
-  // cl_mem input_frame_cl;
+  cl_mem input_frame_cl;
 };

selfdrive/modeld/models/commonmodel.pxd

@@ -14,12 +14,12 @@ cdef extern from "common/clutil.h":
 cdef extern from "selfdrive/modeld/models/commonmodel.h":
   cppclass ModelFrame:
     int buf_size
-    # unsigned char * buffer_from_cl(cl_mem*, int);
-    float * prepare(cl_mem, int, int, int, int, mat3, cl_mem*)
+    unsigned char * buffer_from_cl(cl_mem*, int);
+    cl_mem * prepare(cl_mem, int, int, int, int, mat3)
 
   cppclass DrivingModelFrame:
     int buf_size
-    DrivingModelFrame(cl_device_id, cl_context)
+    DrivingModelFrame(cl_device_id, cl_context, bint)
 
   cppclass MonitoringModelFrame:
     int buf_size

selfdrive/modeld/models/commonmodel_pyx.pyx

@@ -39,31 +39,24 @@ cdef class ModelFrame:
   def __dealloc__(self):
     del self.frame
 
-  def prepare(self, VisionBuf buf, float[:] projection, CLMem output):
+  def prepare(self, VisionBuf buf, float[:] projection):
     cdef mat3 cprojection
     memcpy(cprojection.v, &projection[0], 9*sizeof(float))
-    cdef float * 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
+    cdef cl_mem * data
+    data = self.frame.prepare(buf.buf.buf_cl, buf.width, buf.height, buf.stride, buf.uv_offset, cprojection)
+    return CLMem.create(data)
 
-    return np.asarray(<cnp.float32_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)
+  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)
+  def __cinit__(self, CLContext context, bint is_20hz=False):
+    self._frame = new cppDrivingModelFrame(context.device_id, context.context, is_20hz)
     self.frame = <cppModelFrame*>(self._frame)
     self.buf_size = self._frame.buf_size
 
@@ -74,3 +67,4 @@ cdef class MonitoringModelFrame(ModelFrame):
     self._frame = new cppMonitoringModelFrame(context.device_id, context.context)
     self.frame = <cppModelFrame*>(self._frame)
     self.buf_size = self._frame.buf_size
+

selfdrive/modeld/models/supercombo.onnx

@@ -1,3 +1,3 @@
 version https://git-lfs.github.com/spec/v1
-oid sha256:b31b504bc0b440d3bc72967507a00eb4f112285626fbfb3135011500325ee6d6
-size 51452435
+oid sha256:39786068cae1ed8c0dc34ef80c281dfcc67ed18a50e06b90765c49bcfdbf7db4
+size 51453312

selfdrive/modeld/parse_model_outputs.py

@@ -85,6 +85,7 @@ 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))
@@ -96,6 +97,8 @@ 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,))

selfdrive/modeld/runners/__init__.py

@@ -1,27 +0,0 @@
-import os
-from openpilot.system.hardware import TICI
-from openpilot.selfdrive.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.selfdrive.modeld.runners.thneedmodel_pyx import ThneedModel as Runner
-      runner_type = ModelRunner.THNEED
-    elif ModelRunner.SNPE in paths and USE_SNPE:
-      from openpilot.selfdrive.modeld.runners.snpemodel_pyx import SNPEModel as Runner
-      runner_type = ModelRunner.SNPE
-    elif ModelRunner.ONNX in paths:
-      from openpilot.selfdrive.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)

selfdrive/modeld/runners/model_runner.py

@@ -0,0 +1,120 @@
+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()

selfdrive/modeld/runners/onnxmodel.py

@@ -1,71 +0,0 @@
-import os
-import onnx
-import sys
-import numpy as np
-from typing import Any
-
-from openpilot.selfdrive.modeld.runners.runmodel_pyx import RunModel
-from openpilot.selfdrive.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

selfdrive/modeld/runners/run.h

@@ -1,4 +0,0 @@
-#pragma once
-
-#include "selfdrive/modeld/runners/runmodel.h"
-#include "selfdrive/modeld/runners/snpemodel.h"

selfdrive/modeld/runners/runmodel.h

@@ -1,49 +0,0 @@
-#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);
-  }
-};

selfdrive/modeld/runners/runmodel.pxd

@@ -1,14 +0,0 @@
-# distutils: language = c++
-
-from libcpp.string cimport string
-
-cdef extern from "selfdrive/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()

selfdrive/modeld/runners/runmodel_pyx.pxd

@@ -1,6 +0,0 @@
-# distutils: language = c++
-
-from .runmodel cimport RunModel as cppRunModel
-
-cdef class RunModel:
-  cdef cppRunModel * model

selfdrive/modeld/runners/runmodel_pyx.pyx

@@ -1,37 +0,0 @@
-# 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 selfdrive.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()

selfdrive/modeld/runners/snpemodel.cc

@@ -1,116 +0,0 @@
-#pragma clang diagnostic ignored "-Wexceptions"
-
-#include "selfdrive/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();
-  }
-}

selfdrive/modeld/runners/snpemodel.h

@@ -1,52 +0,0 @@
-#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 "selfdrive/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;
-};

selfdrive/modeld/runners/snpemodel.pxd

@@ -1,9 +0,0 @@
-# distutils: language = c++
-
-from libcpp.string cimport string
-
-from msgq.visionipc.visionipc cimport cl_context
-
-cdef extern from "selfdrive/modeld/runners/snpemodel.h":
-  cdef cppclass SNPEModel:
-    SNPEModel(string, float*, size_t, int, bool, cl_context)

selfdrive/modeld/runners/snpemodel_pyx.pyx

@@ -1,17 +0,0 @@
-# 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 selfdrive.modeld.models.commonmodel_pyx cimport CLContext
-from selfdrive.modeld.runners.runmodel_pyx cimport RunModel
-from selfdrive.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)

selfdrive/modeld/runners/thneedmodel.cc

@@ -1,58 +0,0 @@
-#include "selfdrive/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);
-  }
-}

selfdrive/modeld/runners/thneedmodel.h

@@ -1,17 +0,0 @@
-#pragma once
-
-#include <string>
-
-#include "selfdrive/modeld/runners/runmodel.h"
-#include "selfdrive/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;
-};

selfdrive/modeld/runners/thneedmodel.pxd

@@ -1,9 +0,0 @@
-# distutils: language = c++
-
-from libcpp.string cimport string
-
-from msgq.visionipc.visionipc cimport cl_context
-
-cdef extern from "selfdrive/modeld/runners/thneedmodel.h":
-  cdef cppclass ThneedModel:
-    ThneedModel(string, float*, size_t, int, bool, cl_context)

selfdrive/modeld/runners/thneedmodel_pyx.pyx

@@ -1,14 +0,0 @@
-# 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 selfdrive.modeld.models.commonmodel_pyx cimport CLContext
-from selfdrive.modeld.runners.runmodel_pyx cimport RunModel
-from selfdrive.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)

selfdrive/modeld/runners/tinygrad_helpers.py

@@ -0,0 +1,8 @@
+
+from tinygrad.tensor import Tensor
+from tinygrad.helpers import to_mv
+
+def qcom_tensor_from_opencl_address(opencl_address, shape, dtype):
+  cl_buf_desc_ptr = to_mv(opencl_address, 8).cast('Q')[0]
+  rawbuf_ptr = to_mv(cl_buf_desc_ptr, 0x100).cast('Q')[20] # offset 0xA0 is a raw gpu pointer.
+  return Tensor.from_blob(rawbuf_ptr, shape, dtype=dtype, device='QCOM')

selfdrive/modeld/tests/snpe_benchmark/.gitignore

@@ -1 +0,0 @@
-benchmark

selfdrive/modeld/tests/snpe_benchmark/benchmark.cc

@@ -1,192 +0,0 @@
-#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;
-}
-

selfdrive/modeld/tests/snpe_benchmark/benchmark.sh

@@ -1,4 +0,0 @@
-#!/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

selfdrive/modeld/thneed/README

@@ -1,8 +0,0 @@
-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.
-

selfdrive/modeld/thneed/serialize.cc

@@ -1,154 +0,0 @@
-#include <cassert>
-#include <set>
-
-#include "third_party/json11/json11.hpp"
-#include "common/util.h"
-#include "common/clutil.h"
-#include "common/swaglog.h"
-#include "selfdrive/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);
-}

selfdrive/modeld/thneed/thneed.h

@@ -1,133 +0,0 @@
-#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();
-};
-

selfdrive/modeld/thneed/thneed_common.cc

@@ -1,216 +0,0 @@
-#include "selfdrive/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;
-}

selfdrive/modeld/thneed/thneed_pc.cc

@@ -1,32 +0,0 @@
-#include "selfdrive/modeld/thneed/thneed.h"
-
-#include <cassert>
-
-#include "common/clutil.h"
-#include "common/timing.h"
-
-Thneed::Thneed(bool do_clinit, cl_context _context) {
-  context = _context;
-  if (do_clinit) clinit();
-  char *thneed_debug_env = getenv("THNEED_DEBUG");
-  debug = (thneed_debug_env != NULL) ? atoi(thneed_debug_env) : 0;
-}
-
-void Thneed::execute(float **finputs, float *foutput, bool slow) {
-  uint64_t tb, te;
-  if (debug >= 1) tb = nanos_since_boot();
-
-  // ****** copy inputs
-  copy_inputs(finputs);
-
-  // ****** run commands
-  clexec();
-
-  // ****** copy outputs
-  copy_output(foutput);
-
-  if (debug >= 1) {
-    te = nanos_since_boot();
-    printf("model exec in %lu us\n", (te-tb)/1000);
-  }
-}

selfdrive/modeld/thneed/thneed_qcom2.cc

@@ -1,258 +0,0 @@
-#include "selfdrive/modeld/thneed/thneed.h"
-
-#include <dlfcn.h>
-#include <sys/mman.h>
-
-#include <cassert>
-#include <cerrno>
-#include <cstring>
-#include <map>
-#include <string>
-
-#include "common/clutil.h"
-#include "common/timing.h"
-
-Thneed *g_thneed = NULL;
-int g_fd = -1;
-
-void hexdump(uint8_t *d, int len) {
-  assert((len%4) == 0);
-  printf("  dumping %p len 0x%x\n", d, len);
-  for (int i = 0; i < len/4; i++) {
-    if (i != 0 && (i%0x10) == 0) printf("\n");
-    printf("%8x ", d[i]);
-  }
-  printf("\n");
-}
-
-// *********** ioctl interceptor ***********
-
-extern "C" {
-
-int (*my_ioctl)(int filedes, unsigned long request, void *argp) = NULL;
-#undef ioctl
-int ioctl(int filedes, unsigned long request, void *argp) {
-  request &= 0xFFFFFFFF;  // needed on QCOM2
-  if (my_ioctl == NULL) my_ioctl = reinterpret_cast<decltype(my_ioctl)>(dlsym(RTLD_NEXT, "ioctl"));
-  Thneed *thneed = g_thneed;
-
-  // save the fd
-  if (request == IOCTL_KGSL_GPUOBJ_ALLOC) g_fd = filedes;
-
-  // note that this runs always, even without a thneed object
-  if (request == IOCTL_KGSL_DRAWCTXT_CREATE) {
-    struct kgsl_drawctxt_create *create = (struct kgsl_drawctxt_create *)argp;
-    create->flags &= ~KGSL_CONTEXT_PRIORITY_MASK;
-    create->flags |= 6 << KGSL_CONTEXT_PRIORITY_SHIFT;   // priority from 1-15, 1 is max priority
-    printf("IOCTL_KGSL_DRAWCTXT_CREATE: creating context with flags 0x%x\n", create->flags);
-  }
-
-  if (thneed != NULL) {
-    if (request == IOCTL_KGSL_GPU_COMMAND) {
-      struct kgsl_gpu_command *cmd = (struct kgsl_gpu_command *)argp;
-      if (thneed->record) {
-        thneed->timestamp = cmd->timestamp;
-        thneed->context_id = cmd->context_id;
-        thneed->cmds.push_back(unique_ptr<CachedCommand>(new CachedCommand(thneed, cmd)));
-      }
-      if (thneed->debug >= 1) {
-        printf("IOCTL_KGSL_GPU_COMMAND(%2zu): flags: 0x%lx    context_id: %u  timestamp: %u  numcmds: %d  numobjs: %d\n",
-            thneed->cmds.size(),
-            cmd->flags,
-            cmd->context_id, cmd->timestamp, cmd->numcmds, cmd->numobjs);
-      }
-    } else if (request == IOCTL_KGSL_GPUOBJ_SYNC) {
-      struct kgsl_gpuobj_sync *cmd = (struct kgsl_gpuobj_sync *)argp;
-      struct kgsl_gpuobj_sync_obj *objs = (struct kgsl_gpuobj_sync_obj *)(cmd->objs);
-
-      if (thneed->debug >= 2) {
-        printf("IOCTL_KGSL_GPUOBJ_SYNC count:%d ", cmd->count);
-        for (int i = 0; i < cmd->count; i++) {
-          printf(" -- offset:0x%lx len:0x%lx id:%d op:%d  ", objs[i].offset, objs[i].length, objs[i].id, objs[i].op);
-        }
-        printf("\n");
-      }
-
-      if (thneed->record) {
-        thneed->cmds.push_back(unique_ptr<CachedSync>(new
-              CachedSync(thneed, string((char *)objs, sizeof(struct kgsl_gpuobj_sync_obj)*cmd->count))));
-      }
-    } else if (request == IOCTL_KGSL_DEVICE_WAITTIMESTAMP_CTXTID) {
-      struct kgsl_device_waittimestamp_ctxtid *cmd = (struct kgsl_device_waittimestamp_ctxtid *)argp;
-      if (thneed->debug >= 1) {
-        printf("IOCTL_KGSL_DEVICE_WAITTIMESTAMP_CTXTID: context_id: %d  timestamp: %d  timeout: %d\n",
-            cmd->context_id, cmd->timestamp, cmd->timeout);
-      }
-    } else if (request == IOCTL_KGSL_SETPROPERTY) {
-      if (thneed->debug >= 1) {
-        struct kgsl_device_getproperty *prop = (struct kgsl_device_getproperty *)argp;
-        printf("IOCTL_KGSL_SETPROPERTY: 0x%x sizebytes:%zu\n", prop->type, prop->sizebytes);
-        if (thneed->debug >= 2) {
-          hexdump((uint8_t *)prop->value, prop->sizebytes);
-          if (prop->type == KGSL_PROP_PWR_CONSTRAINT) {
-            struct kgsl_device_constraint *constraint = (struct kgsl_device_constraint *)prop->value;
-            hexdump((uint8_t *)constraint->data, constraint->size);
-          }
-        }
-      }
-    } else if (request == IOCTL_KGSL_DRAWCTXT_CREATE || request == IOCTL_KGSL_DRAWCTXT_DESTROY) {
-      // this happens
-    } else if (request == IOCTL_KGSL_GPUOBJ_ALLOC || request == IOCTL_KGSL_GPUOBJ_FREE) {
-      // this happens
-    } else {
-      if (thneed->debug >= 1) {
-        printf("other ioctl %lx\n", request);
-      }
-    }
-  }
-
-  int ret = my_ioctl(filedes, request, argp);
-  // NOTE: This error message goes into stdout and messes up pyenv
-  // if (ret != 0) printf("ioctl returned %d with errno %d\n", ret, errno);
-  return ret;
-}
-
-}
-
-// *********** GPUMalloc ***********
-
-GPUMalloc::GPUMalloc(int size, int fd) {
-  struct kgsl_gpuobj_alloc alloc;
-  memset(&alloc, 0, sizeof(alloc));
-  alloc.size = size;
-  alloc.flags = 0x10000a00;
-  ioctl(fd, IOCTL_KGSL_GPUOBJ_ALLOC, &alloc);
-  void *addr = mmap64(NULL, alloc.mmapsize, 0x3, 0x1, fd, alloc.id*0x1000);
-  assert(addr != MAP_FAILED);
-
-  base = (uint64_t)addr;
-  remaining = size;
-}
-
-GPUMalloc::~GPUMalloc() {
-  // TODO: free the GPU malloced area
-}
-
-void *GPUMalloc::alloc(int size) {
-  void *ret = (void*)base;
-  size = (size+0xff) & (~0xFF);
-  assert(size <= remaining);
-  remaining -= size;
-  base += size;
-  return ret;
-}
-
-// *********** CachedSync, at the ioctl layer ***********
-
-void CachedSync::exec() {
-  struct kgsl_gpuobj_sync cmd;
-
-  cmd.objs = (uint64_t)data.data();
-  cmd.obj_len = data.length();
-  cmd.count = data.length() / sizeof(struct kgsl_gpuobj_sync_obj);
-
-  int ret = ioctl(thneed->fd, IOCTL_KGSL_GPUOBJ_SYNC, &cmd);
-  assert(ret == 0);
-}
-
-// *********** CachedCommand, at the ioctl layer ***********
-
-CachedCommand::CachedCommand(Thneed *lthneed, struct kgsl_gpu_command *cmd) {
-  thneed = lthneed;
-  assert(cmd->numsyncs == 0);
-
-  memcpy(&cache, cmd, sizeof(cache));
-
-  if (cmd->numcmds > 0) {
-    cmds = make_unique<struct kgsl_command_object[]>(cmd->numcmds);
-    memcpy(cmds.get(), (void *)cmd->cmdlist, sizeof(struct kgsl_command_object)*cmd->numcmds);
-    cache.cmdlist = (uint64_t)cmds.get();
-    for (int i = 0; i < cmd->numcmds; i++) {
-      void *nn = thneed->ram->alloc(cmds[i].size);
-      memcpy(nn, (void*)cmds[i].gpuaddr, cmds[i].size);
-      cmds[i].gpuaddr = (uint64_t)nn;
-    }
-  }
-
-  if (cmd->numobjs > 0) {
-    objs = make_unique<struct kgsl_command_object[]>(cmd->numobjs);
-    memcpy(objs.get(), (void *)cmd->objlist, sizeof(struct kgsl_command_object)*cmd->numobjs);
-    cache.objlist = (uint64_t)objs.get();
-    for (int i = 0; i < cmd->numobjs; i++) {
-      void *nn = thneed->ram->alloc(objs[i].size);
-      memset(nn, 0, objs[i].size);
-      objs[i].gpuaddr = (uint64_t)nn;
-    }
-  }
-
-  kq = thneed->ckq;
-  thneed->ckq.clear();
-}
-
-void CachedCommand::exec() {
-  cache.timestamp = ++thneed->timestamp;
-  int ret = ioctl(thneed->fd, IOCTL_KGSL_GPU_COMMAND, &cache);
-
-  if (thneed->debug >= 1) printf("CachedCommand::exec got %d\n", ret);
-
-  if (thneed->debug >= 2) {
-    for (auto &it : kq) {
-      it->debug_print(false);
-    }
-  }
-
-  assert(ret == 0);
-}
-
-// *********** Thneed ***********
-
-Thneed::Thneed(bool do_clinit, cl_context _context) {
-  // TODO: QCOM2 actually requires a different context
-  //context = _context;
-  if (do_clinit) clinit();
-  assert(g_fd != -1);
-  fd = g_fd;
-  ram = make_unique<GPUMalloc>(0x80000, fd);
-  timestamp = -1;
-  g_thneed = this;
-  char *thneed_debug_env = getenv("THNEED_DEBUG");
-  debug = (thneed_debug_env != NULL) ? atoi(thneed_debug_env) : 0;
-}
-
-void Thneed::wait() {
-  struct kgsl_device_waittimestamp_ctxtid wait;
-  wait.context_id = context_id;
-  wait.timestamp = timestamp;
-  wait.timeout = -1;
-
-  uint64_t tb = nanos_since_boot();
-  int wret = ioctl(fd, IOCTL_KGSL_DEVICE_WAITTIMESTAMP_CTXTID, &wait);
-  uint64_t te = nanos_since_boot();
-
-  if (debug >= 1) printf("wait %d after %lu us\n", wret, (te-tb)/1000);
-}
-
-void Thneed::execute(float **finputs, float *foutput, bool slow) {
-  uint64_t tb, te;
-  if (debug >= 1) tb = nanos_since_boot();
-
-  // ****** copy inputs
-  copy_inputs(finputs, true);
-
-  // ****** run commands
-  int i = 0;
-  for (auto &it : cmds) {
-    ++i;
-    if (debug >= 1) printf("run %2d @ %7lu us: ", i, (nanos_since_boot()-tb)/1000);
-    it->exec();
-    if ((i == cmds.size()) || slow) wait();
-  }
-
-  // ****** copy outputs
-  copy_output(foutput);
-
-  if (debug >= 1) {
-    te = nanos_since_boot();
-    printf("model exec in %lu us\n", (te-tb)/1000);
-  }
-}

selfdrive/modeld/transforms/loadyuv.cl

@@ -1,7 +1,7 @@
 #define UV_SIZE ((TRANSFORMED_WIDTH/2)*(TRANSFORMED_HEIGHT/2))
 
 __kernel void loadys(__global uchar8 const * const Y,
-                     __global float * out,
+                     __global uchar * out,
                      int out_offset)
 {
     const int gid = get_global_id(0);
@@ -10,13 +10,12 @@ __kernel void loadys(__global uchar8 const * const Y,
     const int ox = ois % TRANSFORMED_WIDTH;
 
     const uchar8 ys = Y[gid];
-    const float8 ysf = convert_float8(ys);
 
     // 02
     // 13
 
-    __global float* outy0;
-    __global float* outy1;
+    __global uchar* outy0;
+    __global uchar* outy1;
     if ((oy & 1) == 0) {
       outy0 = out + out_offset; //y0
       outy1 = out + out_offset + UV_SIZE*2; //y2
@@ -25,23 +24,24 @@ __kernel void loadys(__global uchar8 const * const Y,
       outy1 = out + out_offset + UV_SIZE*3; //y3
     }
 
-    vstore4(ysf.s0246, 0, outy0 + (oy/2) * (TRANSFORMED_WIDTH/2) + ox/2);
-    vstore4(ysf.s1357, 0, outy1 + (oy/2) * (TRANSFORMED_WIDTH/2) + ox/2);
+    vstore4(ys.s0246, 0, outy0 + (oy/2) * (TRANSFORMED_WIDTH/2) + ox/2);
+    vstore4(ys.s1357, 0, outy1 + (oy/2) * (TRANSFORMED_WIDTH/2) + ox/2);
 }
 
 __kernel void loaduv(__global uchar8 const * const in,
-                     __global float8 * out,
+                     __global uchar8 * out,
                      int out_offset)
 {
   const int gid = get_global_id(0);
   const uchar8 inv = in[gid];
-  const float8 outv  = convert_float8(inv);
-  out[gid + out_offset / 8] = outv;
+  out[gid + out_offset / 8] = inv;
 }
 
-__kernel void copy(__global float8 * inout,
-                   int in_offset)
+__kernel void copy(__global uchar8 * in,
+                   __global uchar8 * out,
+                   int in_offset,
+                   int out_offset)
 {
   const int gid = get_global_id(0);
-  inout[gid] = inout[gid + in_offset / 8];
-}
+  out[gid + out_offset / 8] = in[gid + in_offset / 8];
+}

selfdrive/pandad/panda.cc

@@ -137,8 +137,8 @@ void Panda::enable_deepsleep() {
   handle->control_write(0xfb, 0, 0);
 }
 
-void Panda::send_heartbeat(bool engaged) {
-  handle->control_write(0xf3, engaged, 0);
+void Panda::send_heartbeat(bool engaged, bool engaged_mads) {
+  handle->control_write(0xf3, engaged, engaged_mads);
 }
 
 void Panda::set_can_speed_kbps(uint16_t bus, uint16_t speed) {

selfdrive/pandad/panda.h

@@ -75,7 +75,7 @@ public:
   std::optional<std::string> get_serial();
   void set_power_saving(bool power_saving);
   void enable_deepsleep();
-  void send_heartbeat(bool engaged);
+  void send_heartbeat(bool engaged, bool engaged_mads);
   void set_can_speed_kbps(uint16_t bus, uint16_t speed);
   void set_can_fd_auto(uint16_t bus, bool enabled);
   void set_data_speed_kbps(uint16_t bus, uint16_t speed);

selfdrive/pandad/pandad.cc

@@ -41,6 +41,8 @@
 #define CUTOFF_IL 400
 #define SATURATE_IL 1000
 
+#define ALT_EXP_DISENGAGE_LATERAL_ON_BRAKE 2048
+
 ExitHandler do_exit;
 
 bool check_all_connected(const std::vector<Panda *> &pandas) {
@@ -53,6 +55,18 @@ bool check_all_connected(const std::vector<Panda *> &pandas) {
   return true;
 }
 
+bool process_mads_heartbeat(SubMaster *sm) {
+  const int &alt_exp = (*sm)["carParams"].getCarParams().getAlternativeExperience();
+  const bool disengage_lateral_on_brake = (alt_exp & ALT_EXP_DISENGAGE_LATERAL_ON_BRAKE) != 0;
+
+  const auto &mads = (*sm)["selfdriveStateSP"].getSelfdriveStateSP().getMads();
+  const bool heartbeat_type = disengage_lateral_on_brake ? mads.getActive() : mads.getEnabled();
+
+  const bool engaged = sm->allAliveAndValid({"selfdriveStateSP"}) && heartbeat_type;
+
+  return engaged;
+}
+
 Panda *connect(std::string serial="", uint32_t index=0) {
   std::unique_ptr<Panda> panda;
   try {
@@ -144,6 +158,7 @@ void fill_panda_state(cereal::PandaState::Builder &ps, cereal::PandaState::Panda
   ps.setIgnitionLine(health.ignition_line_pkt);
   ps.setIgnitionCan(health.ignition_can_pkt);
   ps.setControlsAllowed(health.controls_allowed_pkt);
+  ps.setControlsAllowedLat(health.controls_allowed_lat_pkt);
   ps.setTxBufferOverflow(health.tx_buffer_overflow_pkt);
   ps.setRxBufferOverflow(health.rx_buffer_overflow_pkt);
   ps.setPandaType(hw_type);
@@ -327,7 +342,7 @@ void send_peripheral_state(Panda *panda, PubMaster *pm) {
 }
 
 void process_panda_state(std::vector<Panda *> &pandas, PubMaster *pm, bool spoofing_started) {
-  static SubMaster sm({"selfdriveState"});
+  static SubMaster sm({"selfdriveState", "selfdriveStateSP", "carParams"});
 
   std::vector<std::string> connected_serials;
   for (Panda *p : pandas) {
@@ -366,8 +381,9 @@ void process_panda_state(std::vector<Panda *> &pandas, PubMaster *pm, bool spoof
 
     sm.update(0);
     const bool engaged = sm.allAliveAndValid({"selfdriveState"}) && sm["selfdriveState"].getSelfdriveState().getEnabled();
+    const bool engaged_mads = process_mads_heartbeat(&sm);
     for (const auto &panda : pandas) {
-      panda->send_heartbeat(engaged);
+      panda->send_heartbeat(engaged, engaged_mads);
     }
   }
 }

selfdrive/selfdrived/events.py

@@ -105,6 +105,24 @@ class Events:
       ret.append(event)
     return ret
 
+  def has(self, event_name: int) -> bool:
+    return event_name in self.events
+
+  def contains_in_list(self, events_list: list[int]) -> bool:
+    return any(event_name in self.events for event_name in events_list)
+
+  def remove(self, event_name: int, static: bool = False) -> None:
+    if static and event_name in self.static_events:
+      self.static_events.remove(event_name)
+
+    if event_name in self.events:
+      self.event_counters[event_name] = self.event_counters[event_name] + 1
+      self.events.remove(event_name)
+
+  def replace(self, prev_event_name: int, cur_event_name: int, static: bool = False) -> None:
+    self.remove(prev_event_name, static)
+    self.add(cur_event_name, static)
+
 
 class Alert:
   def __init__(self,
@@ -951,6 +969,99 @@ EVENTS: dict[int, dict[str, Alert | AlertCallbackType]] = {
     ET.WARNING: personality_changed_alert,
   },
 
+  # sunnypilot
+  EventName.lkasEnable: {
+    ET.ENABLE: EngagementAlert(AudibleAlert.engage),
+  },
+
+  EventName.lkasDisable: {
+    ET.USER_DISABLE: EngagementAlert(AudibleAlert.disengage),
+  },
+
+  EventName.manualSteeringRequired: {
+    ET.USER_DISABLE: Alert(
+      "Automatic Lane Centering is OFF",
+      "Manual Steering Required",
+      AlertStatus.normal, AlertSize.mid,
+      Priority.LOW, VisualAlert.none, AudibleAlert.disengage, 1.),
+  },
+
+  EventName.manualLongitudinalRequired: {
+    ET.WARNING: Alert(
+      "Smart/Adaptive Cruise Control: OFF",
+      "Manual Speed Control Required",
+      AlertStatus.normal, AlertSize.mid,
+      Priority.LOW, VisualAlert.none, AudibleAlert.none, 1.),
+  },
+
+  EventName.silentLkasEnable: {
+    ET.ENABLE: EngagementAlert(AudibleAlert.none),
+  },
+
+  EventName.silentLkasDisable: {
+    ET.USER_DISABLE: EngagementAlert(AudibleAlert.none),
+  },
+
+  EventName.silentBrakeHold: {
+    ET.USER_DISABLE: EngagementAlert(AudibleAlert.none),
+    ET.NO_ENTRY: NoEntryAlert("Brake Hold Active"),
+  },
+
+  EventName.silentWrongGear: {
+    ET.WARNING: Alert(
+      "",
+      "",
+      AlertStatus.normal, AlertSize.none,
+      Priority.LOWEST, VisualAlert.none, AudibleAlert.none, 0.),
+    ET.NO_ENTRY: Alert(
+      "Gear not D",
+      "openpilot Unavailable",
+      AlertStatus.normal, AlertSize.mid,
+      Priority.LOW, VisualAlert.none, AudibleAlert.none, 0.),
+  },
+
+  EventName.silentReverseGear: {
+    ET.PERMANENT: Alert(
+      "Reverse\nGear",
+      "",
+      AlertStatus.normal, AlertSize.full,
+      Priority.LOWEST, VisualAlert.none, AudibleAlert.none, .2, creation_delay=0.5),
+    ET.NO_ENTRY: NoEntryAlert("Reverse Gear"),
+  },
+
+  EventName.silentDoorOpen: {
+    ET.WARNING: Alert(
+      "",
+      "",
+      AlertStatus.normal, AlertSize.none,
+      Priority.LOWEST, VisualAlert.none, AudibleAlert.none, 0.),
+    ET.NO_ENTRY: NoEntryAlert("Door Open"),
+  },
+
+  EventName.silentSeatbeltNotLatched: {
+    ET.WARNING: Alert(
+      "",
+      "",
+      AlertStatus.normal, AlertSize.none,
+      Priority.LOWEST, VisualAlert.none, AudibleAlert.none, 0.),
+    ET.NO_ENTRY: NoEntryAlert("Seatbelt Unlatched"),
+  },
+
+  EventName.silentParkBrake: {
+    ET.WARNING: Alert(
+      "",
+      "",
+      AlertStatus.normal, AlertSize.none,
+      Priority.LOWEST, VisualAlert.none, AudibleAlert.none, 0.),
+    ET.NO_ENTRY: NoEntryAlert("Parking Brake Engaged"),
+  },
+
+  EventName.controlsMismatchLateral: {
+    ET.IMMEDIATE_DISABLE: ImmediateDisableAlert("Controls Mismatch: Lateral"),
+    ET.NO_ENTRY: NoEntryAlert("Controls Mismatch: Lateral"),
+  },
+
+
 }
 
 

selfdrive/selfdrived/selfdrived.py

@@ -23,6 +23,8 @@ from openpilot.selfdrive.controls.lib.latcontrol import MIN_LATERAL_CONTROL_SPEE
 
 from openpilot.system.version import get_build_metadata
 
+from openpilot.sunnypilot.mads.mads import ModularAssistiveDrivingSystem
+
 REPLAY = "REPLAY" in os.environ
 SIMULATION = "SIMULATION" in os.environ
 TESTING_CLOSET = "TESTING_CLOSET" in os.environ
@@ -131,6 +133,10 @@ class SelfdriveD:
     elif self.CP.passive:
       self.events.add(EventName.dashcamMode, static=True)
 
+    self.mads = ModularAssistiveDrivingSystem(self)
+    sock_services = list(self.pm.sock.keys()) + ['selfdriveStateSP']
+    self.pm = messaging.PubMaster(sock_services)
+
   def update_events(self, CS):
     """Compute onroadEvents from carState"""
 
@@ -451,11 +457,25 @@ class SelfdriveD:
       self.pm.send('onroadEvents', ce_send)
     self.events_prev = self.events.names.copy()
 
+    # selfdriveStateSP
+    ss_sp_msg = messaging.new_message('selfdriveStateSP')
+    ss_sp_msg.valid = True
+    ss_sp = ss_sp_msg.selfdriveStateSP
+    mads = ss_sp.mads
+    mads.state = self.mads.state_machine.state
+    mads.enabled = self.mads.enabled
+    mads.active = self.mads.active
+    mads.available = self.mads.enabled_toggle
+
+    self.pm.send('selfdriveStateSP', ss_sp_msg)
+
   def step(self):
     CS = self.data_sample()
     self.update_events(CS)
     if not self.CP.passive and self.initialized:
       self.enabled, self.active = self.state_machine.update(self.events)
+    if not self.CP.notCar:
+      self.mads.update(CS, self.sm)
     self.update_alerts(CS)
 
     self.publish_selfdriveState(CS)
@@ -473,6 +493,8 @@ class SelfdriveD:
       self.is_metric = self.params.get_bool("IsMetric")
       self.experimental_mode = self.params.get_bool("ExperimentalMode") and self.CP.openpilotLongitudinalControl
       self.personality = self.read_personality_param()
+
+      self.mads.read_params()
       time.sleep(0.1)
 
   def run(self):

selfdrive/selfdrived/tests/test_alerts.py

@@ -41,7 +41,7 @@ class TestAlerts:
     events = log.OnroadEvent.EventName.schema.enumerants
 
     for name, e in events.items():
-      if not name.endswith("DEPRECATED"):
+      if not name.endswith("DEPRECATED") and not name.startswith("eventReserved"):
         fail_msg = f"{name} @{e} not in EVENTS"
         assert e in EVENTS.keys(), fail_msg
 

selfdrive/test/ciui.py

@@ -1,58 +0,0 @@
-#!/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_()

selfdrive/test/docker_tag_multiarch.sh

@@ -1,25 +0,0 @@
-#!/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

selfdrive/test/loop_until_fail.sh

@@ -1,8 +0,0 @@
-#!/usr/bin/env bash
-set -e
-
-# Loop something forever until it fails, for verifying new tests
-
-while true; do
-  $@
-done

selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst

@@ -1,3 +0,0 @@
-version https://git-lfs.github.com/spec/v1
-oid sha256:737f95d34912db53a303ba6499e6f697b510fa5872b8c71f701a4fe924b5466e
-size 356169

selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_card_efa9c321c2f985d948b5d9a5e708f5d2d31b7737.zst

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:79ccc3bd2094ba8a55adedf0007b0152eb3a68edc5e2d35aeccbba122d5811c6
+size 356151

selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_plannerd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst

@@ -1,3 +0,0 @@
-version https://git-lfs.github.com/spec/v1
-oid sha256:e359e8f6b5a22b6f3f89b54989dac2110ee3a4463de2d785be83e20cda4f1cb6
-size 254248