Refactor modeld for flexible input rates and output parsing.

Add support for both 20Hz and variable input rates in `modeld.py` by introducing conditional data updates based on a new `is_20hz` flag. Additionally, enhance `parse_outputs` in `parse_model_outputs.py` to handle `desired_curvature` parsing when relevant input keys are present.

.codespellignore

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

.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

@@ -74,6 +74,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 +178,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,72 @@ $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;
+  }
 }
 
 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/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,36 @@ 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 +146,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 +158,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 +174,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 +200,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
 };
 

pyproject.toml

@@ -163,6 +163,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()

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

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/modeld/SConscript

@@ -21,7 +21,7 @@ else:
   libs += ['OpenCL']
 
 # Set path definitions
-for pathdef, fn in {'TRANSFORM': 'transforms/transform.cl', 'LOADYUV': 'transforms/loadyuv.cl'}.items():
+for pathdef, fn in {'TRANSFORM': 'transforms/transform.cl', 'LOADYUV': 'transforms/loadyuv.cl', 'LOADYUV_FLOAT': 'transforms/loadyuv_float.cl'}.items():
   for xenv in (lenv, lenvCython):
     xenv['CXXFLAGS'].append(f'-D{pathdef}_PATH=\\"{File(fn).abspath}\\"')
 

selfdrive/modeld/modeld.py

@@ -1,21 +1,13 @@
 #!/usr/bin/env python3
-import os
 from openpilot.system.hardware import TICI
 
+from openpilot.selfdrive.modeld.runners.model_runner import ONNXRunner, TinygradRunner
+
 #
-if TICI:
-  from tinygrad.tensor import Tensor
-  from tinygrad.dtype import dtypes
-  from openpilot.selfdrive.modeld.runners.tinygrad_helpers import qcom_tensor_from_opencl_address
-  os.environ['QCOM'] = '1'
-else:
-  from openpilot.selfdrive.modeld.runners.ort_helpers import make_onnx_cpu_runner
 import time
-import pickle
 import numpy as np
 import cereal.messaging as messaging
 from cereal import car, log
-from pathlib import Path
 from setproctitle import setproctitle
 from cereal.messaging import PubMaster, SubMaster
 from msgq.visionipc import VisionIpcClient, VisionStreamType, VisionBuf
@@ -31,15 +23,12 @@ from openpilot.selfdrive.controls.lib.desire_helper import DesireHelper
 from openpilot.selfdrive.modeld.parse_model_outputs import Parser
 from openpilot.selfdrive.modeld.fill_model_msg import fill_model_msg, fill_pose_msg, PublishState
 from openpilot.selfdrive.modeld.constants import ModelConstants
-from openpilot.selfdrive.modeld.models.commonmodel_pyx import DrivingModelFrame, CLContext
+# from openpilot.selfdrive.modeld.models.commonmodel_pyx import DrivingModelFrame, CLContext
+from openpilot.selfdrive.modeld.models.commonmodel_pyx import DrivingModelFrameLegacy as DrivingModelFrame, CLContext
 
 
 PROCESS_NAME = "selfdrive.modeld.modeld"
-SEND_RAW_PRED = os.getenv('SEND_RAW_PRED')
 
-MODEL_PATH = Path(__file__).parent / 'models/supercombo.onnx'
-MODEL_PKL_PATH = Path(__file__).parent / 'models/supercombo_tinygrad.pkl'
-METADATA_PATH = Path(__file__).parent / 'models/supercombo_metadata.pkl'
 
 class FrameMeta:
   frame_id: int = 0
@@ -61,35 +50,26 @@ class ModelState:
     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)
+    self.is_20hz = False
+    # Initialize model runner
+    self.model_runner = TinygradRunner(self.frames) if TICI else ONNXRunner(self.frames)
 
     # img buffers are managed in openCL transform code
-    self.numpy_inputs = {
-      'desire': np.zeros((1, (ModelConstants.HISTORY_BUFFER_LEN+1), ModelConstants.DESIRE_LEN), dtype=np.float32),
-      'traffic_convention': np.zeros((1, ModelConstants.TRAFFIC_CONVENTION_LEN), dtype=np.float32),
-      'features_buffer': np.zeros((1, ModelConstants.HISTORY_BUFFER_LEN,  ModelConstants.FEATURE_LEN), dtype=np.float32),
-    }
+    self.numpy_inputs = {}
 
-    with open(METADATA_PATH, 'rb') as f:
-      model_metadata = pickle.load(f)
-    self.input_shapes =  model_metadata['input_shapes']
+    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)
 
-    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()
 
-    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)
+    net_output_size = self.model_runner.model_metadata['output_shapes']['outputs'][1]
+    self.output = np.zeros(net_output_size, dtype=np.float32)
 
-  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
+    num_elements = self.numpy_inputs['features_buffer'].shape[1]
+    step_size = int(-100 / num_elements)
+    self.full_features_20Hz_idxs = np.arange(step_size, step_size * (num_elements + 1), step_size)[::-1]
+    self.desire_reshape_dims = (self.numpy_inputs['desire'].shape[0], self.numpy_inputs['desire'].shape[1], -1, self.numpy_inputs['desire'].shape[2])
 
   def 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:
@@ -98,38 +78,54 @@ class ModelState:
     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.numpy_inputs['desire'][:] = self.desire_20Hz.reshape((1,25,4,-1)).max(axis=2)
+    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:
+      len = inputs['desire'].shape[0]
+      self.numpy_inputs['desire'][0, :-1] = self.numpy_inputs['desire'][0, 1:]
+      self.numpy_inputs['desire'][0, -1, :len] = new_desire[:len]
+
+    for key in self.numpy_inputs:
+      if key in inputs and key not in ['desire']:
+        self.numpy_inputs[key][:] = inputs[key]
 
-    self.numpy_inputs['traffic_convention'][:] = inputs['traffic_convention']
     imgs_cl = {'input_imgs': self.frames['input_imgs'].prepare(buf, transform.flatten()),
                'big_input_imgs': self.frames['big_input_imgs'].prepare(wbuf, transform_wide.flatten())}
 
-    if TICI:
-      # The imgs tensors are backed by opencl memory, only need init once
-      for key in imgs_cl:
-        if key not in self.tensor_inputs:
-          self.tensor_inputs[key] = qcom_tensor_from_opencl_address(imgs_cl[key].mem_address, self.input_shapes[key], dtype=dtypes.uint8)
-    else:
-      for key in imgs_cl:
-        self.numpy_inputs[key] = self.frames[key].buffer_from_cl(imgs_cl[key]).reshape(self.input_shapes[key])
+    # Prepare inputs using the model runner
+    self.model_runner.prepare_inputs(imgs_cl, self.numpy_inputs)
 
     if prepare_only:
       return None
 
-    if TICI:
-      self.output = self.model_run(**self.tensor_inputs).numpy().flatten()
+    # Run model inference
+    self.output = self.model_runner.run_model()
+    outputs = self.parser.parse_outputs(self.model_runner.slice_outputs(self.output), self.numpy_inputs.keys())
+
+    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.output = self.onnx_cpu_runner.run(None, self.numpy_inputs)[0].flatten()
+      feature_len = outputs['hidden_state'].shape[1]
+      self.numpy_inputs['features_buffer'][0, :-1] = self.numpy_inputs['features_buffer'][0, 1:]
+      self.numpy_inputs['features_buffer'][0, -1, :feature_len] = outputs['hidden_state'][0, :feature_len]
 
-    outputs = self.parser.parse_outputs(self.slice_outputs(self.output))
 
-    self.full_features_20Hz[:-1] = self.full_features_20Hz[1:]
-    self.full_features_20Hz[-1] = outputs['hidden_state'][0, :]
+    if "desired_curvature" in outputs:
+      input_name_prev = None
 
-    idxs = np.arange(-4,-100,-4)[::-1]
-    self.numpy_inputs['features_buffer'][:] = self.full_features_20Hz[idxs]
+      if "prev_desired_curvs" in self.numpy_inputs.keys():
+        input_name_prev = 'prev_desired_curvs'
+      elif "prev_desired_curv" in self.numpy_inputs.keys():
+        input_name_prev = 'prev_desired_curv'
+
+      if input_name_prev is not None:
+        len = outputs['desired_curvature'][0].size
+        self.numpy_inputs[input_name_prev][0, :-len, 0] = self.numpy_inputs[input_name_prev][0, len:, 0]
+        self.numpy_inputs[input_name_prev][0, -len:, 0] = outputs['desired_curvature'][0]
     return outputs
 
 
@@ -190,7 +186,6 @@ def main(demo=False):
   meta_main = FrameMeta()
   meta_extra = FrameMeta()
 
-
   if demo:
     CP = get_demo_car_params()
   else:
@@ -272,6 +267,9 @@ def main(demo=False):
       'traffic_convention': traffic_convention,
       }
 
+    if "lateral_control_params" in model.numpy_inputs.keys():
+      inputs['lateral_control_params'] = np.array([sm["carState"].vEgo, steer_delay], dtype=np.float32)
+
     mt1 = time.perf_counter()
     model_output = model.run(buf_main, buf_extra, model_transform_main, model_transform_extra, inputs, prepare_only)
     mt2 = time.perf_counter()

selfdrive/modeld/models/commonmodel.cc

@@ -8,25 +8,30 @@
 DrivingModelFrame::DrivingModelFrame(cl_device_id device_id, cl_context context) : ModelFrame(device_id, context) {
   input_frames = std::make_unique<uint8_t[]>(buf_size);
   input_frames_cl = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_WRITE, buf_size, NULL, &err));
-  img_buffer_20hz_cl = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_WRITE, 5*frame_size_bytes, NULL, &err));
-  region.origin = 4 * frame_size_bytes;
+  img_buffer_20hz_cl = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_WRITE, buffer_length*frame_size_bytes, NULL, &err));
+  region.origin = (buffer_length - 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));
+  printf("Buffer length: %d, region origin: %lu, region size: %lu\n", buffer_length, region.origin, region.size);
 
-  loadyuv_init(&loadyuv, context, device_id, MODEL_WIDTH, MODEL_HEIGHT);
+  loadyuv_init(&loadyuv, context, device_id, MODEL_WIDTH, MODEL_HEIGHT, is_float);
   init_transform(device_id, context, MODEL_WIDTH, MODEL_HEIGHT);
 }
 
 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 < (buffer_length - 1); i++) {
+    printf("Moving %d to %d from src_offset %lu to dst_offset %lu with size %lu\n", i+1, i, (i+1)*frame_size_bytes, i*frame_size_bytes, frame_size_bytes);
     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));
   }
+  printf("Loop done\n");
   loadyuv_queue(&loadyuv, q, y_cl, u_cl, v_cl, last_img_cl);
 
   copy_queue(&loadyuv, q, img_buffer_20hz_cl, input_frames_cl, 0, 0, frame_size_bytes);
+  printf("Copying from %p to %p with size %lu\n", img_buffer_20hz_cl, input_frames_cl, frame_size_bytes);
   copy_queue(&loadyuv, q, last_img_cl, input_frames_cl, 0, frame_size_bytes, frame_size_bytes);
+  printf("Copying from %p to %p with size %lu\n", last_img_cl, input_frames_cl, frame_size_bytes);
 
   // NOTE: Since thneed is using a different command queue, this clFinish is needed to ensure the image is ready.
   clFinish(q);
@@ -41,6 +46,45 @@ DrivingModelFrame::~DrivingModelFrame() {
   CL_CHECK(clReleaseCommandQueue(q));
 }
 
+DrivingModelFrameLegacy::DrivingModelFrameLegacy(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 * sizeof(float), NULL, &err));
+  // Reduce buffer to 2 frames due to float size (4x bigger than uint8_t)
+  img_buffer_20hz_cl = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_WRITE, 2 * frame_size_bytes * sizeof(float), NULL, &err));
+    
+  // Adjust region for 2-frame buffer
+  region.origin = frame_size_bytes * sizeof(float);  // Point to second frame
+  region.size = frame_size_bytes * sizeof(float);
+  last_img_cl = CL_CHECK_ERR(clCreateSubBuffer(img_buffer_20hz_cl, CL_MEM_READ_WRITE, CL_BUFFER_CREATE_TYPE_REGION, &region, &err));
+
+  loadyuv_init(&loadyuv, context, device_id, MODEL_WIDTH, MODEL_HEIGHT, true);
+  init_transform(device_id, context, MODEL_WIDTH, MODEL_HEIGHT);
+}
+
+cl_mem* DrivingModelFrameLegacy::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);
+
+  // Rolling buffer with just 2 frames
+  CL_CHECK(clEnqueueCopyBuffer(q, img_buffer_20hz_cl, img_buffer_20hz_cl, frame_size_bytes * sizeof(float), 0, frame_size_bytes * sizeof(float), 0, nullptr, nullptr));
+    
+  loadyuv_queue(&loadyuv, q, y_cl, u_cl, v_cl, last_img_cl);
+
+  // Copy both frames to input buffer
+  copy_queue(&loadyuv, q, img_buffer_20hz_cl, input_frames_cl, 0, 0, frame_size_bytes * sizeof(float));
+  copy_queue(&loadyuv, q, last_img_cl, input_frames_cl, 0, frame_size_bytes * sizeof(float), frame_size_bytes * sizeof(float));
+
+  clFinish(q);
+  return &input_frames_cl;
+}
+
+DrivingModelFrameLegacy::~DrivingModelFrameLegacy() {
+  deinit_transform();
+  loadyuv_destroy(&loadyuv);
+  CL_CHECK(clReleaseMemObject(img_buffer_20hz_cl));
+  CL_CHECK(clReleaseMemObject(input_frames_cl));
+  CL_CHECK(clReleaseCommandQueue(q));
+}
 
 MonitoringModelFrame::MonitoringModelFrame(cl_device_id device_id, cl_context context) : ModelFrame(device_id, context) {
   input_frames = std::make_unique<uint8_t[]>(buf_size);

selfdrive/modeld/models/commonmodel.h

@@ -17,6 +17,7 @@
 #include "selfdrive/modeld/transforms/loadyuv.h"
 #include "selfdrive/modeld/transforms/transform.h"
 
+template <typename T = uint8_t>
 class ModelFrame {
 public:
   ModelFrame(cl_device_id device_id, cl_context context) {
@@ -24,7 +25,7 @@ public:
   }
   virtual ~ModelFrame() {}
   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) {
+  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];
@@ -39,7 +40,7 @@ protected:
   cl_mem y_cl, u_cl, v_cl;
   Transform transform;
   cl_command_queue q;
-  std::unique_ptr<uint8_t[]> input_frames;
+  std::unique_ptr<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));
@@ -62,7 +63,11 @@ protected:
   }
 };
 
-class DrivingModelFrame : public ModelFrame {
+class DrivingModelFrame : public ModelFrame<> {
+  using ModelFrame::q, ModelFrame::y_cl, ModelFrame::u_cl, ModelFrame::v_cl;
+  using ModelFrame::init_transform, ModelFrame::deinit_transform, ModelFrame::run_transform;
+  using ModelFrame::input_frames;
+  
 public:
   DrivingModelFrame(cl_device_id device_id, cl_context context);
   ~DrivingModelFrame();
@@ -73,6 +78,8 @@ public:
   const int MODEL_FRAME_SIZE = MODEL_WIDTH * MODEL_HEIGHT * 3 / 2;
   const int buf_size = MODEL_FRAME_SIZE * 2;
   const size_t frame_size_bytes = MODEL_FRAME_SIZE * sizeof(uint8_t);
+  const bool is_float = false;
+  const uint8_t buffer_length = is_float ? 2 : 5;
 
 private:
   LoadYUVState loadyuv;
@@ -80,7 +87,7 @@ private:
   cl_buffer_region region;
 };
 
-class MonitoringModelFrame : public ModelFrame {
+class MonitoringModelFrame : public ModelFrame<> {
 public:
   MonitoringModelFrame(cl_device_id device_id, cl_context context);
   ~MonitoringModelFrame();
@@ -94,3 +101,23 @@ public:
 private:
   cl_mem input_frame_cl;
 };
+
+class DrivingModelFrameLegacy : public ModelFrame<float> {
+public:
+  DrivingModelFrameLegacy(cl_device_id device_id, cl_context context);
+  ~DrivingModelFrameLegacy();
+  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;
+
+private:
+  cl_mem input_frames_cl;
+  cl_mem img_buffer_20hz_cl;
+  cl_mem last_img_cl;
+  LoadYUVState loadyuv;
+  cl_buffer_region region;
+  const size_t frame_size_bytes = MODEL_FRAME_SIZE * sizeof(float);  // This is the only change
+};

selfdrive/modeld/models/commonmodel.pxd

@@ -12,15 +12,19 @@ cdef extern from "common/clutil.h":
   cl_context cl_create_context(cl_device_id)
 
 cdef extern from "selfdrive/modeld/models/commonmodel.h":
-  cppclass ModelFrame:
+  cppclass ModelFrame[T]:
     int buf_size
-    unsigned char * buffer_from_cl(cl_mem*, int);
+    T * 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)
 
+  cppclass DrivingModelFrameLegacy:
+    int buf_size
+    DrivingModelFrameLegacy(cl_device_id, cl_context)
+
   cppclass MonitoringModelFrame:
     int buf_size
     MonitoringModelFrame(cl_device_id, cl_context)

selfdrive/modeld/models/commonmodel_pyx.pyx

@@ -4,12 +4,12 @@
 import numpy as np
 cimport numpy as cnp
 from libc.string cimport memcpy
-from libc.stdint cimport uintptr_t
+from libc.stdint cimport uintptr_t, uint8_t
 
 from msgq.visionipc.visionipc cimport cl_mem
 from msgq.visionipc.visionipc_pyx cimport VisionBuf, CLContext as BaseCLContext
 from .commonmodel cimport CL_DEVICE_TYPE_DEFAULT, cl_get_device_id, cl_create_context
-from .commonmodel cimport mat3, ModelFrame as cppModelFrame, DrivingModelFrame as cppDrivingModelFrame, MonitoringModelFrame as cppMonitoringModelFrame
+from .commonmodel cimport mat3, ModelFrame as cppModelFrame, DrivingModelFrame as cppDrivingModelFrame, DrivingModelFrameLegacy as cppDrivingModelFrameLegacy, MonitoringModelFrame as cppMonitoringModelFrame
 
 
 cdef class CLContext(BaseCLContext):
@@ -33,7 +33,7 @@ def cl_from_visionbuf(VisionBuf buf):
 
 
 cdef class ModelFrame:
-  cdef cppModelFrame * frame
+  cdef cppModelFrame[uint8_t] * frame
   cdef int buf_size
 
   def __dealloc__(self):
@@ -52,12 +52,40 @@ cdef class ModelFrame:
     return np.asarray(<cnp.uint8_t[:self.buf_size]> data2)
 
 
+cdef class ModelFrame_float:
+  cdef cppModelFrame[float] * frame
+  cdef int buf_size
+
+  def __dealloc__(self):
+    del self.frame
+
+  def prepare(self, VisionBuf buf, float[:] projection):
+    cdef mat3 cprojection
+    memcpy(cprojection.v, &projection[0], 9*sizeof(float))
+    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)
+
+  def buffer_from_cl(self, CLMem in_frames):
+    cdef float * data2
+    data2 = self.frame.buffer_from_cl(in_frames.mem, self.buf_size)
+    return np.asarray(<cnp.float32_t[:self.buf_size]> data2)
+
+
 cdef class DrivingModelFrame(ModelFrame):
   cdef cppDrivingModelFrame * _frame
 
   def __cinit__(self, CLContext context):
     self._frame = new cppDrivingModelFrame(context.device_id, context.context)
-    self.frame = <cppModelFrame*>(self._frame)
+    self.frame = <cppModelFrame[uint8_t]*>(self._frame)
+    self.buf_size = self._frame.buf_size
+
+cdef class DrivingModelFrameLegacy(ModelFrame_float):
+  cdef cppDrivingModelFrameLegacy * _frame
+
+  def __cinit__(self, CLContext context):
+    self._frame = new cppDrivingModelFrameLegacy(context.device_id, context.context)
+    self.frame = <cppModelFrame[float]*>(self._frame)
     self.buf_size = self._frame.buf_size
 
 cdef class MonitoringModelFrame(ModelFrame):
@@ -65,6 +93,6 @@ cdef class MonitoringModelFrame(ModelFrame):
 
   def __cinit__(self, CLContext context):
     self._frame = new cppMonitoringModelFrame(context.device_id, context.context)
-    self.frame = <cppModelFrame*>(self._frame)
+    self.frame = <cppModelFrame[uint8_t]*>(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:72d3d6f8d3c98f5431ec86be77b6350d7d4f43c25075c0106f1d1e7ec7c77668
-size 49096168
+oid sha256:b31b504bc0b440d3bc72967507a00eb4f112285626fbfb3135011500325ee6d6
+size 51452435

selfdrive/modeld/parse_model_outputs.py

@@ -84,7 +84,8 @@ class Parser:
     outs[name] = pred_mu_final.reshape(final_shape)
     outs[name + '_stds'] = pred_std_final.reshape(final_shape)
 
-  def parse_outputs(self, outs: dict[str, np.ndarray]) -> dict[str, np.ndarray]:
+  def parse_outputs(self, outs: dict[str, np.ndarray], input_keys: [str]) -> 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 and "prev_desired_curv" in input_keys:
+      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/model_runner.py

@@ -0,0 +1,115 @@
+import os
+from openpilot.system.hardware import TICI
+
+#
+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.selfdrive.modeld.models.commonmodel_pyx import DrivingModelFrameLegacy as DrivingModelFrame, CLMem
+
+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'
+
+
+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."""
+    with open(METADATA_PATH, '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__()
+    # Load Tinygrad model
+    with open(MODEL_PKL_PATH, "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:
+    if self.is_memory_model is None:
+      self.is_memory_model = any(self.input_to_dtype[key] == dtypes.uint8 for key in imgs_cl)
+      print(f"Memory model: {self.is_memory_model}")
+
+    # Initialize image tensors if not already done
+    for key in imgs_cl:
+      if TICI and self.is_memory_model 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 or not self.is_memory_model:
+        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(MODEL_PATH)
+    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])
+    return self.inputs
+
+  def run_model(self):
+    return self.runner.run(None, self.inputs)[0].flatten()

selfdrive/modeld/transforms/loadyuv.cc

@@ -4,7 +4,7 @@
 #include <cstdio>
 #include <cstring>
 
-void loadyuv_init(LoadYUVState* s, cl_context ctx, cl_device_id device_id, int width, int height) {
+void loadyuv_init(LoadYUVState* s, cl_context ctx, cl_device_id device_id, int width, int height, bool use_float) {
   memset(s, 0, sizeof(*s));
 
   s->width = width;
@@ -15,7 +15,9 @@ void loadyuv_init(LoadYUVState* s, cl_context ctx, cl_device_id device_id, int w
            "-cl-fast-relaxed-math -cl-denorms-are-zero "
            "-DTRANSFORMED_WIDTH=%d -DTRANSFORMED_HEIGHT=%d",
            width, height);
-  cl_program prg = cl_program_from_file(ctx, device_id, LOADYUV_PATH, args);
+  const char * loadyuv_path = use_float ? LOADYUV_FLOAT_PATH : LOADYUV_PATH;
+  printf(" Use float: %d\n Using loadyuv_path: %s\n", use_float, loadyuv_path);
+  cl_program prg = cl_program_from_file(ctx, device_id, loadyuv_path, args);
 
   s->loadys_krnl = CL_CHECK_ERR(clCreateKernel(prg, "loadys", &err));
   s->loaduv_krnl = CL_CHECK_ERR(clCreateKernel(prg, "loaduv", &err));

selfdrive/modeld/transforms/loadyuv.h

@@ -7,7 +7,8 @@ typedef struct {
   cl_kernel loadys_krnl, loaduv_krnl, copy_krnl;
 } LoadYUVState;
 
-void loadyuv_init(LoadYUVState* s, cl_context ctx, cl_device_id device_id, int width, int height);
+void loadyuv_init(LoadYUVState* s, cl_context ctx, cl_device_id device_id, int width, int height, bool use_float);
+inline void loadyuv_init(LoadYUVState* s, cl_context ctx, cl_device_id device_id, int width, int height) { loadyuv_init(s, ctx, device_id, width, height, false); };
 
 void loadyuv_destroy(LoadYUVState* s);
 

selfdrive/modeld/transforms/loadyuv_float.cl

@@ -0,0 +1,49 @@
+#define UV_SIZE ((TRANSFORMED_WIDTH/2)*(TRANSFORMED_HEIGHT/2))
+
+__kernel void loadys(__global uchar8 const * const Y,
+                     __global float * out,
+                     int out_offset)
+{
+    const int gid = get_global_id(0);
+    const int ois = gid * 8;
+    const int oy = ois / TRANSFORMED_WIDTH;
+    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;
+    if ((oy & 1) == 0) {
+      outy0 = out + out_offset; //y0
+      outy1 = out + out_offset + UV_SIZE*2; //y2
+    } else {
+      outy0 = out + out_offset + UV_SIZE; //y1
+      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);
+}
+
+__kernel void loaduv(__global uchar8 const * const in,
+                     __global float8 * 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;
+}
+
+__kernel void copy(__global float8 * in,
+                   __global float8 * out,
+                   int in_offset,
+                   int out_offset)
+{
+  const int gid = get_global_id(0);
+  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/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

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

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

selfdrive/test/process_replay/fakedata/regen306779F6870|2024-10-03--04-03-23--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst

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

selfdrive/test/process_replay/fakedata/regen306779F6870|2024-10-03--04-03-23--0_card_efa9c321c2f985d948b5d9a5e708f5d2d31b7737.zst

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

selfdrive/test/process_replay/fakedata/regen4B38A7428CD|2024-08-30--02-56-31--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst

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

selfdrive/test/process_replay/fakedata/regen4B38A7428CD|2024-08-30--02-56-31--0_card_efa9c321c2f985d948b5d9a5e708f5d2d31b7737.zst

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

selfdrive/test/process_replay/fakedata/regen4CE950B0267|2024-08-30--02-51-30--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst

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

selfdrive/test/process_replay/fakedata/regen4CE950B0267|2024-08-30--02-51-30--0_card_efa9c321c2f985d948b5d9a5e708f5d2d31b7737.zst

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

selfdrive/test/process_replay/fakedata/regen58464878D07|2024-08-30--03-15-31--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst

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

selfdrive/test/process_replay/fakedata/regen58464878D07|2024-08-30--03-15-31--0_card_efa9c321c2f985d948b5d9a5e708f5d2d31b7737.zst

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

selfdrive/test/process_replay/fakedata/regen6E484EDAB96|2024-08-30--02-47-37--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst

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

selfdrive/test/process_replay/fakedata/regen6E484EDAB96|2024-08-30--02-47-37--0_card_efa9c321c2f985d948b5d9a5e708f5d2d31b7737.zst

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

selfdrive/test/process_replay/fakedata/regen720F2BA4CF6|2024-08-30--03-09-15--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst

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

selfdrive/test/process_replay/fakedata/regen720F2BA4CF6|2024-08-30--03-09-15--0_card_efa9c321c2f985d948b5d9a5e708f5d2d31b7737.zst

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

selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst

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

selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_card_efa9c321c2f985d948b5d9a5e708f5d2d31b7737.zst

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

selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_plannerd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst

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

selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_plannerd_efa9c321c2f985d948b5d9a5e708f5d2d31b7737.zst

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

selfdrive/test/process_replay/fakedata/regen9ADBECBCD1C|2024-08-30--03-13-04--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst

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

selfdrive/test/process_replay/fakedata/regen9ADBECBCD1C|2024-08-30--03-13-04--0_card_efa9c321c2f985d948b5d9a5e708f5d2d31b7737.zst

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

selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst

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

selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_card_efa9c321c2f985d948b5d9a5e708f5d2d31b7737.zst

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

selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_plannerd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst

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

selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_plannerd_efa9c321c2f985d948b5d9a5e708f5d2d31b7737.zst

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

selfdrive/test/process_replay/fakedata/regenA67A128BCD8|2024-08-30--02-36-22--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst

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

selfdrive/test/process_replay/fakedata/regenA67A128BCD8|2024-08-30--02-36-22--0_card_efa9c321c2f985d948b5d9a5e708f5d2d31b7737.zst

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

selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst

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

selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_card_efa9c321c2f985d948b5d9a5e708f5d2d31b7737.zst

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

selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_plannerd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst

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

selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_plannerd_efa9c321c2f985d948b5d9a5e708f5d2d31b7737.zst

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