Switch model inputs and buffers from uint8_t to float.

Updated data types across the model pipeline, including ONNX model, buffers, and transformations, to utilize `float` instead of `uint8_t`.

.devcontainer/.gitignore

@@ -1,3 +0,0 @@
-.Xauthority
-.env
-.host/

.devcontainer/Dockerfile

@@ -1,18 +0,0 @@
-FROM ghcr.io/commaai/openpilot-base:latest
-
-RUN apt update && apt install -y vim net-tools usbutils htop ripgrep tmux wget mesa-utils xvfb libxtst6 libxv1 libglu1-mesa gdb bash-completion
-RUN python3 -m ensurepip --upgrade
-RUN pip3 install ipython jupyter jupyterlab
-
-RUN cd /tmp && \
-    ARCH=$(arch | sed s/aarch64/arm64/ | sed s/x86_64/amd64/) && \
-    curl -L -o virtualgl.deb "https://github.com/VirtualGL/virtualgl/releases/download/3.1.1/virtualgl_3.1.1_$ARCH.deb" && \
-    dpkg -i virtualgl.deb
-
-RUN usermod -aG video batman
-
-USER batman
-
-RUN cd $HOME && \
-    curl -O https://raw.githubusercontent.com/commaai/agnos-builder/master/userspace/home/.tmux.conf && \
-    curl -O https://raw.githubusercontent.com/commaai/agnos-builder/master/userspace/home/.vimrc

.devcontainer/container_post_create.sh

@@ -1,38 +0,0 @@
-#!/usr/bin/env bash
-
-TARGET_USER=batman
-source .devcontainer/.host/.env
-
-# override display flag for mac hosts
-if [[ $HOST_OS == darwin ]]; then
-  echo "Setting up DISPLAY override for macOS..."
-  cat <<EOF >> /home/$TARGET_USER/.bashrc
-source .devcontainer/.host/.env
-if [ -n "\$HOST_DISPLAY" ]; then
-  DISPLAY_NUM=\$(echo "\$HOST_DISPLAY" | awk -F: '{print \$NF}')
-  export DISPLAY=host.docker.internal:\$DISPLAY_NUM
-fi
-EOF
-fi
-
-# setup virtualgl for mac hosts
-if [[ $HOST_OS == darwin ]]; then
-  echo "Setting up virtualgl for macOS..."
-  cat <<EOF >> /home/$TARGET_USER/.bashrc
-if [ -n "\$HOST_DISPLAY" ]; then
-  export VGL_PORT=10000
-  export VGL_CLIENT=host.docker.internal
-  export VGL_COMPRESS=rgb
-  export VGL_DISPLAY=:99
-  export VGL_FPS=60
-  # prevent vglrun from running exec
-  alias exec=:; source vglrun :; unalias exec
-fi
-EOF
-fi
-
-# These lines are temporary, to remain backwards compatible with old devcontainers
-# that were running as root and therefore had their caches written as root
-sudo chown -R $TARGET_USER: /tmp/scons_cache
-sudo chown -R $TARGET_USER: /tmp/comma_download_cache
-sudo chown -R $TARGET_USER: /home/batman/.comma

.devcontainer/container_post_start.sh

@@ -1,15 +0,0 @@
-#!/usr/bin/env bash
-
-source .devcontainer/.host/.env
-
-# setup safe directories for submodules
-SUBMODULE_DIRS=$(git config --file .gitmodules --get-regexp path | awk '{ print $2 }')
-for DIR in $SUBMODULE_DIRS; do 
-  git config --global --add safe.directory "$PWD/$DIR"
-done
-
-# virtual display for virtualgl
-if [[ "$HOST_OS" == "darwin" ]] && [[ -n "$HOST_DISPLAY" ]]; then
-  echo "Starting virtual display at :99 ..."
-  tmux new-session -d -s fakedisplay Xvfb :99 -screen 0 1920x1080x24
-fi

.devcontainer/devcontainer.json

@@ -1,53 +0,0 @@
-{
-  "name": "openpilot devcontainer",
-  "build": {
-    "dockerfile": "Dockerfile"
-  },
-  "postCreateCommand": ".devcontainer/container_post_create.sh",
-  "postStartCommand": ".devcontainer/container_post_start.sh",
-  "initializeCommand": [".devcontainer/host_setup"],
-  "privileged": true,
-  "containerEnv": {
-    "DISPLAY": "${localEnv:DISPLAY}",
-    "PYTHONPATH": "${containerWorkspaceFolder}",
-    "TERM": "xterm-256color",
-    "force_color_prompt": "1"
-  },
-  "runArgs": [
-    "--volume=/dev:/dev",
-    "--volume=/tmp/.X11-unix:/tmp/.X11-unix",
-    "--volume=${localWorkspaceFolder}/.devcontainer/.host/.Xauthority:/home/batman/.Xauthority",
-    "--volume=${localEnv:HOME}/.comma:/home/batman/.comma",
-    "--volume=${localEnv:HOME}/.azure:/home/batman/.azure",
-    "--volume=/tmp/comma_download_cache:/tmp/comma_download_cache",
-    "--shm-size=1G",
-    "--add-host=host.docker.internal:host-gateway", // required to use host.docker.internal on linux
-    "--publish=0.0.0.0:8070-8079:8070-8079"  // body ZMQ services
-  ],
-  "features": {
-    "ghcr.io/devcontainers/features/common-utils:2": {
-      "installZsh": false,
-      "installOhMyZsh": false,
-      "upgradePackages": false,
-      "username": "batman"
-    },
-    "ghcr.io/devcontainers-contrib/features/gh-cli:1": {},
-    "ghcr.io/devcontainers/features/azure-cli:1": {}
-  },
-  "containerUser": "batman",
-  "remoteUser": "batman",
-  "customizations": {
-    "vscode": {
-      "extensions": [
-        "ms-python.python",
-        "ms-vscode.cpptools",
-        "ms-toolsai.jupyter",
-        "guyskk.language-cython",
-        "lharri73.dbc"
-      ]
-    }
-  },
-  "mounts": [
-    "type=volume,source=scons_cache,target=/tmp/scons_cache"
-  ]
-}

.devcontainer/host_setup

@@ -1,47 +0,0 @@
-#!/usr/bin/env bash
-
-# pull base image
-if [[ -z $USE_LOCAL_IMAGE ]]; then
-  echo "Updating openpilot_base image if needed..."
-  docker pull ghcr.io/commaai/openpilot-base:latest
-fi
-
-# setup .host dir
-mkdir -p .devcontainer/.host
-
-# setup links to Xauthority
-XAUTHORITY_LINK=".devcontainer/.host/.Xauthority"
-rm -f $XAUTHORITY_LINK
-if [[ -z $XAUTHORITY ]]; then
-  echo "XAUTHORITY not set. Fallback to ~/.Xauthority ..."
-  if ! [[ -f $HOME/.Xauthority ]]; then
-    echo "~/.XAuthority file does not exist. GUI tools may not work properly."
-    touch $XAUTHORITY_LINK # dummy file to satisfy container volume mount
-  else
-    ln -sf $HOME/.Xauthority $XAUTHORITY_LINK
-  fi
-else
-    ln -sf $XAUTHORITY $XAUTHORITY_LINK
-fi
-
-# setup host env file
-HOST_INFO_FILE=".devcontainer/.host/.env"
-SYSTEM=$(uname -s | tr '[:upper:]' '[:lower:]')
-echo "HOST_OS=\"$SYSTEM\"" > $HOST_INFO_FILE
-echo "HOST_DISPLAY=\"$DISPLAY\"" >> $HOST_INFO_FILE
-
-# run virtualgl if macos
-if [[ $SYSTEM == "darwin" ]]; then
-  echo
-  if [[ -f /opt/VirtualGL/bin/vglclient ]]; then
-    echo "Starting VirtualGL client at port 10000..."
-    VGL_LOG_FILE=".devcontainer/.host/.vgl/vglclient.log"
-    mkdir -p "$(dirname $VGL_LOG_FILE)"
-    /opt/VirtualGL/bin/vglclient -l "$VGL_LOG_FILE" -display "$DISPLAY" -port 10000 -detach
-  else
-    echo "VirtualGL not found. GUI tools may not work properly. Some GUI tools require OpenGL to work properly. To use them with XQuartz on mac, VirtualGL needs to be installed. To install it run:"
-    echo
-    echo "  brew install --cask virtualgl"
-    echo
-  fi
-fi

.devcontainer/host_setup.cmd

@@ -1,10 +0,0 @@
-:: pull base image
-IF NOT DEFINED USE_LOCAL_IMAGE ^
-echo "Updating openpilot_base image if needed..." && ^
-docker pull ghcr.io/commaai/openpilot-base:latest
-
-:: setup .host dir
-mkdir .devcontainer\.host
-
-:: setup host env file
-echo "" > .devcontainer\.host\.env

.github/PULL_REQUEST_TEMPLATE/car_port.md

@@ -8,7 +8,7 @@ assignees: ''
 
 **Checklist**
 
-- [ ] added entry to CAR in selfdrive/car/*/values.py and ran `selfdrive/opcar/docs.py` to generate new docs
+- [ ] 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:

.github/pull_request_template.md

@@ -44,7 +44,7 @@ Explain how you tested this bug fix.
 
 **Checklist**
 
-- [ ] added entry to CAR in selfdrive/car/*/values.py and ran `selfdrive/opcar/docs.py` to generate new docs
+- [ ] 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:

.github/workflows/auto-cache/action.yaml

@@ -14,17 +14,25 @@ inputs:
     description: 'whether to save the cache'
     default: 'false'
     required: false
+outputs:
+  cache-hit:
+    description: 'cache hit occurred'
+    value: ${{ (contains(runner.name, 'nsc') && steps.ns-cache.outputs.cache-hit) ||
+               (!contains(runner.name, 'nsc') && inputs.save != 'false' && steps.gha-cache.outputs.cache-hit) ||
+               (!contains(runner.name, 'nsc') && inputs.save == 'false' && steps.gha-cache-ro.outputs.cache-hit) }}
 
 runs:
   using: "composite"
   steps:
     - name: setup namespace cache
+      id: ns-cache
       if: ${{ contains(runner.name, 'nsc') }}
       uses: namespacelabs/nscloud-cache-action@v1
       with:
         path: ${{ inputs.path }}
 
     - name: setup github cache
+      id: gha-cache
       if: ${{ !contains(runner.name, 'nsc') && inputs.save != 'false' }}
       uses: 'actions/cache@v4'
       with:
@@ -33,6 +41,7 @@ runs:
         restore-keys: ${{ inputs.restore-keys }}
 
     - name: setup github cache
+      id: gha-cache-ro
       if: ${{ !contains(runner.name, 'nsc') && inputs.save == 'false' }}
       uses: 'actions/cache/restore@v4'
       with:

.github/workflows/ci_weekly_run.yaml

@@ -15,7 +15,3 @@ jobs:
     uses: commaai/openpilot/.github/workflows/selfdrive_tests.yaml@master
     with:
       run_number: ${{ inputs.run_number }}
-  tools_tests:
-    uses: commaai/openpilot/.github/workflows/tools_tests.yaml@master
-    with:
-      run_number: ${{ inputs.run_number }}

.github/workflows/docs.yaml

@@ -18,9 +18,10 @@ concurrency:
 jobs:
   docs:
     name: build docs
-    runs-on: ubuntu-latest
-    timeout-minutes: 1
+    runs-on: ubuntu-24.04
     steps:
+    - uses: commaai/timeout@v1
+
     - uses: actions/checkout@v4
       with:
         submodules: true

.github/workflows/repo-maintenance.yaml

@@ -28,7 +28,7 @@ jobs:
         git add .
     - name: update car docs
       run: |
-        scons -j$(nproc) --minimal opendbc
+        scons -j$(nproc) --minimal opendbc_repo
         PYTHONPATH=. python selfdrive/car/docs.py
         git add docs/CARS.md
     - name: Create Pull Request

.github/workflows/selfdrive_tests.yaml

@@ -32,9 +32,9 @@ env:
 jobs:
   build_release:
     name: build release
-    runs-on: ${{ ((github.repository == 'commaai/openpilot') &&
-                   ((github.event_name != 'pull_request') ||
-                    (github.event.pull_request.head.repo.full_name == 'commaai/openpilot'))) && 'namespace-profile-amd64-8x16' || 'ubuntu-latest' }}
+    runs-on:
+      - ${{ ((github.repository == 'commaai/openpilot') && ((github.event_name != 'pull_request') || (github.event.pull_request.head.repo.full_name == 'commaai/openpilot'))) && 'namespace-profile-amd64-8x16' || 'ubuntu-24.04' }}
+      - ${{ ((github.repository == 'commaai/openpilot') && ((github.event_name != 'pull_request') || (github.event.pull_request.head.repo.full_name == 'commaai/openpilot'))) && 'namespace-experiments:docker.builds.local-cache=separate' || 'ubuntu-24.04' }}
     env:
       STRIPPED_DIR: /tmp/releasepilot
     steps:
@@ -75,15 +75,9 @@ jobs:
         ${{ env.RUN }} "scripts/lint/lint.sh --skip check_added_large_files"
 
   build:
-    strategy:
-      matrix:
-        arch: ${{ fromJson(
-           ((github.repository == 'commaai/openpilot') &&
-              ((github.event_name != 'pull_request') ||
-               (github.event.pull_request.head.repo.full_name == 'commaai/openpilot'))) && '["x86_64", "aarch64"]' || '["x86_64"]' ) }}
-    runs-on: ${{ ((matrix.arch == 'aarch64') && 'namespace-profile-arm64-2x8') ||
-                 ((matrix.arch == 'x86_64') && ((github.repository == 'commaai/openpilot') && ((github.event_name != 'pull_request') || (github.event.pull_request.head.repo.full_name == 'commaai/openpilot'))) && 'namespace-profile-amd64-8x16') ||
-                   'ubuntu-latest'}}
+    runs-on:
+      - ${{ ((github.repository == 'commaai/openpilot') && ((github.event_name != 'pull_request') || (github.event.pull_request.head.repo.full_name == 'commaai/openpilot'))) && 'namespace-profile-amd64-8x16' || 'ubuntu-24.04' }}
+      - ${{ ((github.repository == 'commaai/openpilot') && ((github.event_name != 'pull_request') || (github.event.pull_request.head.repo.full_name == 'commaai/openpilot'))) && 'namespace-experiments:docker.builds.local-cache=separate' || 'ubuntu-24.04' }}
     steps:
     - uses: actions/checkout@v4
       with:
@@ -92,17 +86,14 @@ jobs:
       if: github.ref == 'refs/heads/master' && github.event_name != 'pull_request' && github.repository == 'commaai/openpilot'
       run: |
         echo "PUSH_IMAGE=true" >> "$GITHUB_ENV"
-        echo "TARGET_ARCHITECTURE=${{ matrix.arch }}" >> "$GITHUB_ENV"
         $DOCKER_LOGIN
     - uses: ./.github/workflows/setup-with-retry
-      with:
-        docker_hub_pat: ${{ secrets.DOCKER_HUB_PAT }}
     - uses: ./.github/workflows/compile-openpilot
-      timeout-minutes: ${{ ((steps.restore-scons-cache.outputs.cache-hit == 'true') && 15 || 30) }} # allow more time when we missed the scons cache
+      timeout-minutes: 30
 
   build_mac:
     name: build macOS
-    runs-on: namespace-profile-macos-8x14
+    runs-on: ${{ github.repository == 'commaai/openpilot' && 'namespace-profile-macos-8x14' || 'macos-latest' }}
     steps:
     - uses: actions/checkout@v4
       with:
@@ -125,28 +116,11 @@ jobs:
     - name: Building openpilot
       run: . .venv/bin/activate && scons -j$(nproc)
 
-  docker_push_multiarch:
-    name: docker push multiarch tag
-    runs-on: ubuntu-latest
-    if: github.ref == 'refs/heads/master' && github.event_name != 'pull_request' && github.repository == 'commaai/openpilot'
-    needs: [build]
-    steps:
-    - uses: actions/checkout@v4
-      with:
-        submodules: false
-    - name: Setup docker
-      run: |
-        $DOCKER_LOGIN
-    - name: Merge x64 and arm64 tags
-      run: |
-        export PUSH_IMAGE=true
-        scripts/retry.sh selfdrive/test/docker_tag_multiarch.sh base x86_64 aarch64
-
   static_analysis:
     name: static analysis
-    runs-on: ${{ ((github.repository == 'commaai/openpilot') &&
-                   ((github.event_name != 'pull_request') ||
-                    (github.event.pull_request.head.repo.full_name == 'commaai/openpilot'))) && 'namespace-profile-amd64-8x16' || 'ubuntu-24.04' }}
+    runs-on:
+      - ${{ ((github.repository == 'commaai/openpilot') && ((github.event_name != 'pull_request') || (github.event.pull_request.head.repo.full_name == 'commaai/openpilot'))) && 'namespace-profile-amd64-8x16' || 'ubuntu-24.04' }}
+      - ${{ ((github.repository == 'commaai/openpilot') && ((github.event_name != 'pull_request') || (github.event.pull_request.head.repo.full_name == 'commaai/openpilot'))) && 'namespace-experiments:docker.builds.local-cache=separate' || 'ubuntu-24.04' }}
     env:
       PYTHONWARNINGS: default
     steps:
@@ -160,24 +134,16 @@ jobs:
 
   unit_tests:
     name: unit tests
-    runs-on: ${{ ((github.repository == 'commaai/openpilot') &&
-                   ((github.event_name != 'pull_request') ||
-                    (github.event.pull_request.head.repo.full_name == 'commaai/openpilot'))) && 'namespace-profile-amd64-8x16' || 'ubuntu-latest' }}
+    runs-on:
+      - ${{ ((github.repository == 'commaai/openpilot') && ((github.event_name != 'pull_request') || (github.event.pull_request.head.repo.full_name == 'commaai/openpilot'))) && 'namespace-profile-amd64-8x16' || 'ubuntu-24.04' }}
+      - ${{ ((github.repository == 'commaai/openpilot') && ((github.event_name != 'pull_request') || (github.event.pull_request.head.repo.full_name == 'commaai/openpilot'))) && 'namespace-experiments:docker.builds.local-cache=separate' || 'ubuntu-24.04' }}
     steps:
     - uses: actions/checkout@v4
       with:
         submodules: true
     - uses: ./.github/workflows/setup-with-retry
-      with:
-        docker_hub_pat: ${{ secrets.DOCKER_HUB_PAT }}
     - name: Build openpilot
-      timeout-minutes: ${{ ((steps.restore-scons-cache.outputs.cache-hit == 'true') && 10 || 30) }} # allow more time when we missed the scons cache
       run: ${{ env.RUN }} "scons -j$(nproc)"
-    - name: Setup cache
-      uses: ./.github/workflows/auto-cache
-      with:
-        path: .ci_cache/comma_download_cache
-        key: unit_tests_${{ hashFiles('.github/workflows/selfdrive_tests.yaml') }}
     - name: Run unit tests
       timeout-minutes: ${{ contains(runner.name, 'nsc') && 1 || 20 }}
       run: |
@@ -195,27 +161,25 @@ jobs:
 
   process_replay:
     name: process replay
-    runs-on: ${{ ((github.repository == 'commaai/openpilot') &&
-                   ((github.event_name != 'pull_request') ||
-                    (github.event.pull_request.head.repo.full_name == 'commaai/openpilot'))) && 'namespace-profile-amd64-8x16' || 'ubuntu-latest' }}
+    runs-on:
+      - ${{ ((github.repository == 'commaai/openpilot') && ((github.event_name != 'pull_request') || (github.event.pull_request.head.repo.full_name == 'commaai/openpilot'))) && 'namespace-profile-amd64-8x16' || 'ubuntu-24.04' }}
+      - ${{ ((github.repository == 'commaai/openpilot') && ((github.event_name != 'pull_request') || (github.event.pull_request.head.repo.full_name == 'commaai/openpilot'))) && 'namespace-experiments:docker.builds.local-cache=separate' || 'ubuntu-24.04' }}
     steps:
     - uses: actions/checkout@v4
       with:
         submodules: true
     - uses: ./.github/workflows/setup-with-retry
-      with:
-        docker_hub_pat: ${{ secrets.DOCKER_HUB_PAT }}
     - name: Cache test routes
       id: dependency-cache
       uses: actions/cache@v4
       with:
         path: .ci_cache/comma_download_cache
-        key: proc-replay-${{ hashFiles('.github/workflows/selfdrive_tests.yaml', 'selfdrive/test/process_replay/ref_commit', 'selfdrive/test/process_replay/test_regen.py') }}
+        key: proc-replay-${{ hashFiles('selfdrive/test/process_replay/ref_commit', 'selfdrive/test/process_replay/test_processes.py') }}
     - name: Build openpilot
       run: |
         ${{ env.RUN }} "scons -j$(nproc)"
     - name: Run replay
-      timeout-minutes: ${{ contains(runner.name, 'nsc') && 1 || 20 }}
+      timeout-minutes: ${{ contains(runner.name, 'nsc') && (steps.dependency-cache.outputs.cache-hit == 'true') && 1 || 20 }}
       run: |
         ${{ env.RUN }} "coverage run selfdrive/test/process_replay/test_processes.py -j$(nproc) && \
                         chmod -R 777 /tmp/comma_download_cache && \
@@ -250,9 +214,9 @@ jobs:
 
   test_cars:
     name: cars
-    runs-on: ${{ ((github.repository == 'commaai/openpilot') &&
-                   ((github.event_name != 'pull_request') ||
-                    (github.event.pull_request.head.repo.full_name == 'commaai/openpilot'))) && 'namespace-profile-amd64-8x16' || 'ubuntu-latest' }}
+    runs-on:
+      - ${{ ((github.repository == 'commaai/openpilot') && ((github.event_name != 'pull_request') || (github.event.pull_request.head.repo.full_name == 'commaai/openpilot'))) && 'namespace-profile-amd64-8x16' || 'ubuntu-24.04' }}
+      - ${{ ((github.repository == 'commaai/openpilot') && ((github.event_name != 'pull_request') || (github.event.pull_request.head.repo.full_name == 'commaai/openpilot'))) && 'namespace-experiments:docker.builds.local-cache=separate' || 'ubuntu-24.04' }}
     strategy:
       fail-fast: false
       matrix:
@@ -263,17 +227,17 @@ jobs:
         submodules: true
     - uses: ./.github/workflows/setup-with-retry
     - name: Cache test routes
-      id: dependency-cache
-      uses: ./.github/workflows/auto-cache
+      id: routes-cache
+      uses: actions/cache@v4
       with:
         path: .ci_cache/comma_download_cache
-        key: car_models-${{ hashFiles('selfdrive/car/tests/test_models.py', 'selfdrive/car/tests/routes.py') }}-${{ matrix.job }}
+        key: car_models-${{ hashFiles('selfdrive/car/tests/test_models.py', 'opendbc/car/tests/routes.py') }}-${{ matrix.job }}
     - name: Build openpilot
       run: ${{ env.RUN }} "scons -j$(nproc)"
     - name: Test car models
-      timeout-minutes: ${{ contains(runner.name, 'nsc') && 1 || 20 }}
+      timeout-minutes: ${{ contains(runner.name, 'nsc') && (steps.routes-cache.outputs.cache-hit == 'true') && 1 || 20 }}
       run: |
-        ${{ env.RUN }} "FILEREADER_CACHE=1 MAX_EXAMPLES=1 $PYTEST selfdrive/car/tests/test_models.py && \
+        ${{ env.RUN }} "MAX_EXAMPLES=1 $PYTEST selfdrive/car/tests/test_models.py && \
                         chmod -R 777 /tmp/comma_download_cache"
       env:
         NUM_JOBS: 4
@@ -342,9 +306,10 @@ jobs:
 
   simulator_driving:
     name: simulator driving
-    runs-on: ${{ ((github.repository == 'commaai/openpilot') &&
-                   ((github.event_name != 'pull_request') ||
-                    (github.event.pull_request.head.repo.full_name == 'commaai/openpilot'))) && 'namespace-profile-amd64-8x16' || 'ubuntu-24.04' }}
+    runs-on:
+      - ${{ ((github.repository == 'commaai/openpilot') && ((github.event_name != 'pull_request') || (github.event.pull_request.head.repo.full_name == 'commaai/openpilot'))) && 'namespace-profile-amd64-8x16' || 'ubuntu-24.04' }}
+      - ${{ ((github.repository == 'commaai/openpilot') && ((github.event_name != 'pull_request') || (github.event.pull_request.head.repo.full_name == 'commaai/openpilot'))) && 'namespace-experiments:docker.builds.local-cache=separate' || 'ubuntu-24.04' }}
+    if: (github.repository == 'commaai/openpilot') && ((github.event_name != 'pull_request') || (github.event.pull_request.head.repo.full_name == 'commaai/openpilot'))
     steps:
     - uses: actions/checkout@v4
       with:
@@ -363,9 +328,9 @@ jobs:
   create_ui_report:
     # This job name needs to be the same as UI_JOB_NAME in ui_preview.yaml
     name: Create UI Report
-    runs-on: ${{ ((github.repository == 'commaai/openpilot') &&
-                   ((github.event_name != 'pull_request') ||
-                    (github.event.pull_request.head.repo.full_name == 'commaai/openpilot'))) && 'namespace-profile-amd64-8x16' || 'ubuntu-24.04' }}
+    runs-on:
+      - ${{ ((github.repository == 'commaai/openpilot') && ((github.event_name != 'pull_request') || (github.event.pull_request.head.repo.full_name == 'commaai/openpilot'))) && 'namespace-profile-amd64-8x16' || 'ubuntu-24.04' }}
+      - ${{ ((github.repository == 'commaai/openpilot') && ((github.event_name != 'pull_request') || (github.event.pull_request.head.repo.full_name == 'commaai/openpilot'))) && 'namespace-experiments:docker.builds.local-cache=separate' || 'ubuntu-24.04' }}
     steps:
       - uses: actions/checkout@v4
         with:

.github/workflows/setup-with-retry/action.yaml

@@ -17,7 +17,6 @@ runs:
       uses: ./.github/workflows/setup
       continue-on-error: true
       with:
-        docker_hub_pat: ${{ inputs.docker_hub_pat }}
         is_retried: true
     - if: steps.setup1.outcome == 'failure'
       shell: bash
@@ -27,7 +26,6 @@ runs:
       uses: ./.github/workflows/setup
       continue-on-error: true
       with:
-        docker_hub_pat: ${{ inputs.docker_hub_pat }}
         is_retried: true
     - if: steps.setup2.outcome == 'failure'
       shell: bash
@@ -36,5 +34,4 @@ runs:
       if: steps.setup2.outcome == 'failure'
       uses: ./.github/workflows/setup
       with:
-        docker_hub_pat: ${{ inputs.docker_hub_pat }}
         is_retried: true

.github/workflows/setup/action.yaml

@@ -1,10 +1,6 @@
 name: 'openpilot env setup'
 
 inputs:
-  docker_hub_pat:
-    description: 'Auth token for Docker Hub, required for BuildJet jobs'
-    required: true
-    default: ''
   is_retried:
     description: 'A mock param that asserts that we use the setup-with-retry instead of this action directly'
     required: false
@@ -24,11 +20,9 @@ runs:
       name: No retries!
       run: |
         if [ "${{ github.run_attempt }}" -gt 1 ]; then
-          echo -e "\033[31m"
-          echo "##################################################"
-          echo "    Retries not allowed! Fix the flaky test!      "
-          echo "##################################################"
-          echo -e "\033[0m"
+          echo -e "\033[0;31m##################################################"
+          echo -e "\033[0;31m    Retries not allowed! Fix the flaky test!      "
+          echo -e "\033[0;31m##################################################\033[0m"
           exit 1
         fi
 
@@ -36,19 +30,6 @@ runs:
     - shell: bash
       run: git lfs pull
 
-    # on BuildJet runners, must be logged into DockerHub to avoid rate limiting
-    # https://buildjet.com/for-github-actions/docs/guides/docker
-    - shell: bash
-      if: ${{ contains(runner.name, 'buildjet') && inputs.docker_hub_pat == '' }}
-      run: |
-        echo "Need to set the Docker Hub PAT secret as an input to this action"
-        exit 1
-    - name: Login to Docker Hub
-      if: contains(runner.name, 'buildjet')
-      shell: bash
-      run: |
-        docker login -u adeebshihadeh -p ${{ inputs.docker_hub_pat }}
-
     # build cache
     - id: date
       shell: bash

.github/workflows/tools_tests.yaml

@@ -1,54 +0,0 @@
-name: tools
-
-on:
-  push:
-    branches:
-      - master
-  pull_request:
-  workflow_call:
-    inputs:
-      run_number:
-        default: '1'
-        required: true
-        type: string
-concurrency:
-  group: tools-tests-ci-run-${{ inputs.run_number }}-${{ github.event_name == 'push' && github.ref == 'refs/heads/master' && github.run_id || github.head_ref || github.ref }}-${{ github.workflow }}-${{ github.event_name }}
-  cancel-in-progress: true
-
-env:
-  BASE_IMAGE: openpilot-base
-  DOCKER_LOGIN: docker login ghcr.io -u ${{ github.actor }} -p ${{ secrets.GITHUB_TOKEN }}
-
-  BUILD: selfdrive/test/docker_build.sh base
-
-  RUN: docker run --shm-size 2G -v $GITHUB_WORKSPACE:/tmp/openpilot -w /tmp/openpilot -e FILEREADER_CACHE=1 -e PYTHONPATH=/tmp/openpilot -e NUM_JOBS -e JOB_ID -e GITHUB_ACTION -e GITHUB_REF -e GITHUB_HEAD_REF -e GITHUB_SHA -e GITHUB_REPOSITORY -e GITHUB_RUN_ID -v $GITHUB_WORKSPACE/.ci_cache/scons_cache:/tmp/scons_cache -v $GITHUB_WORKSPACE/.ci_cache/comma_download_cache:/tmp/comma_download_cache -v $GITHUB_WORKSPACE/.ci_cache/openpilot_cache:/tmp/openpilot_cache $BASE_IMAGE /bin/bash -c
-
-
-jobs:
-  devcontainer:
-    name: devcontainer
-    runs-on: ubuntu-latest
-    if: false  # we can re-enable once this is faster
-    steps:
-    - uses: actions/checkout@v4
-      with:
-        submodules: true
-    - uses: ./.github/workflows/setup-with-retry
-    - name: Use local image for testing devcontainer with latest base image
-      run: |
-        echo "USE_LOCAL_IMAGE=true" >> "$GITHUB_ENV"
-    - name: Setup Dev Container CLI
-      run: npm install -g @devcontainers/cli
-    - name: Build dev container image
-      run: ./scripts/retry.sh devcontainer build --workspace-folder .
-    - name: Run dev container
-      run: |
-        mkdir -p /tmp/devcontainer_scons_cache/
-        cp -r $GITHUB_WORKSPACE/.ci_cache/scons_cache/. /tmp/devcontainer_scons_cache/
-        devcontainer up --workspace-folder .
-    - name: Test environment
-      run: |
-        devcontainer exec --workspace-folder . scons -j$(nproc) cereal/ common/
-        devcontainer exec --workspace-folder . pip3 install pip-install-test
-        devcontainer exec --workspace-folder . touch /home/batman/.comma/auth.json
-        devcontainer exec --workspace-folder . sudo touch /root/test.txt

Dockerfile.openpilot

@@ -1,9 +1,9 @@
 FROM ghcr.io/commaai/openpilot-base:latest
 
-ENV PYTHONUNBUFFERED 1
+ENV PYTHONUNBUFFERED=1
 
-ENV OPENPILOT_PATH /home/batman/openpilot
-ENV PYTHONPATH ${OPENPILOT_PATH}:${PYTHONPATH}
+ENV OPENPILOT_PATH=/home/batman/openpilot
+ENV PYTHONPATH=${OPENPILOT_PATH}:${PYTHONPATH}
 
 RUN mkdir -p ${OPENPILOT_PATH}
 WORKDIR ${OPENPILOT_PATH}

Dockerfile.openpilot_base

@@ -1,6 +1,6 @@
 FROM ubuntu:24.04
 
-ENV PYTHONUNBUFFERED 1
+ENV PYTHONUNBUFFERED=1
 
 ENV DEBIAN_FRONTEND=noninteractive
 RUN apt-get update && \
@@ -8,9 +8,9 @@ RUN apt-get update && \
     rm -rf /var/lib/apt/lists/*
 
 RUN sed -i -e 's/# en_US.UTF-8 UTF-8/en_US.UTF-8 UTF-8/' /etc/locale.gen && locale-gen
-ENV LANG en_US.UTF-8
-ENV LANGUAGE en_US:en
-ENV LC_ALL en_US.UTF-8
+ENV LANG=en_US.UTF-8
+ENV LANGUAGE=en_US:en
+ENV LC_ALL=en_US.UTF-8
 
 COPY tools/install_ubuntu_dependencies.sh /tmp/tools/
 RUN /tmp/tools/install_ubuntu_dependencies.sh && \
@@ -55,9 +55,9 @@ RUN mkdir -p /tmp/opencl-driver-intel && \
     cd / && \
     rm -rf /tmp/opencl-driver-intel
 
-ENV NVIDIA_VISIBLE_DEVICES all
-ENV NVIDIA_DRIVER_CAPABILITIES graphics,utility,compute
-ENV QTWEBENGINE_DISABLE_SANDBOX 1
+ENV NVIDIA_VISIBLE_DEVICES=all
+ENV NVIDIA_DRIVER_CAPABILITIES=graphics,utility,compute
+ENV QTWEBENGINE_DISABLE_SANDBOX=1
 
 RUN dbus-uuidgen > /etc/machine-id
 
@@ -68,17 +68,14 @@ RUN usermod -aG sudo $USER
 RUN echo '%sudo ALL=(ALL) NOPASSWD:ALL' >> /etc/sudoers
 USER $USER
 
-COPY --chown=$USER pyproject.toml uv.lock /tmp/
-COPY --chown=$USER tools/install_python_dependencies.sh /tmp/tools/
+COPY --chown=$USER pyproject.toml uv.lock /home/$USER
+COPY --chown=$USER tools/install_python_dependencies.sh /home/$USER/tools/
 
 ENV VIRTUAL_ENV=/home/$USER/.venv
 ENV PATH="$VIRTUAL_ENV/bin:$PATH"
-RUN cd /tmp && \
+RUN cd /home/$USER && \
     tools/install_python_dependencies.sh && \
-    mkdir -p $VIRTUAL_ENV && \
-    cp -r /tmp/.venv/* $VIRTUAL_ENV && \
-    rm -rf /tmp/* && \
-    rm -rf /home/$USER/.cache
+    rm -rf tools/ pyproject.toml uv.lock .cache
 
 USER root
 RUN sudo git config --global --add safe.directory /tmp/openpilot

Jenkinsfile

@@ -79,6 +79,10 @@ def deviceStage(String stageName, String deviceType, List extra_env, def steps)
         return
     }
 
+    if (isReplay()) {
+      error("REPLAYING TESTS IS NOT ALLOWED. FIX THEM INSTEAD.")
+    }
+
     def extra = extra_env.collect { "export ${it}" }.join('\n');
     def branch = env.BRANCH_NAME ?: 'master';
     def gitDiff = sh returnStdout: true, script: 'curl -s -H "Authorization: Bearer ${GITHUB_COMMENTS_TOKEN}" https://api.github.com/repos/commaai/openpilot/compare/master...${GIT_BRANCH} | jq .files[].filename || echo "/"', label: 'Getting changes'
@@ -123,6 +127,11 @@ def hasPathChanged(String gitDiff, List<String> paths) {
   return false
 }
 
+def isReplay() {
+  def replayClass = "org.jenkinsci.plugins.workflow.cps.replay.ReplayCause"
+  return currentBuild.rawBuild.getCauses().any{ cause -> cause.toString().contains(replayClass) }
+}
+
 def setupCredentials() {
   withCredentials([
     string(credentialsId: 'azure_token', variable: 'AZURE_TOKEN'),
@@ -207,8 +216,8 @@ node {
           step("build", "cd system/manager && ./build.py"),
           step("test pandad", "pytest selfdrive/pandad/tests/test_pandad.py", [diffPaths: ["panda", "selfdrive/pandad/"]]),
           step("test power draw", "pytest -s system/hardware/tici/tests/test_power_draw.py"),
-          step("test encoder", "LD_LIBRARY_PATH=/usr/local/lib pytest system/loggerd/tests/test_encoder.py"),
-          step("test pigeond", "pytest system/ubloxd/tests/test_pigeond.py"),
+          step("test encoder", "LD_LIBRARY_PATH=/usr/local/lib pytest system/loggerd/tests/test_encoder.py", [diffPaths: ["system/loggerd/"]]),
+          step("test pigeond", "pytest system/ubloxd/tests/test_pigeond.py", [diffPaths: ["system/ubloxd/"]]),
           step("test manager", "pytest system/manager/test/test_manager.py"),
         ])
       },
@@ -243,7 +252,7 @@ node {
       'replay': {
         deviceStage("model-replay", "tici-replay", ["UNSAFE=1"], [
           step("build", "cd system/manager && ./build.py", [diffPaths: ["selfdrive/modeld/", "tinygrad_repo", "selfdrive/test/process_replay/model_replay.py"]]),
-          step("model replay", "selfdrive/test/process_replay/model_replay.py", [diffPaths: ["selfdrive/modeld/"]]),
+          step("model replay", "selfdrive/test/process_replay/model_replay.py", [diffPaths: ["selfdrive/modeld/", "tinygrad_repo", "selfdrive/test/process_replay/model_replay.py"]]),
         ])
       },
       'tizi': {
@@ -253,7 +262,7 @@ node {
           step("test pandad spi", "pytest selfdrive/pandad/tests/test_pandad_spi.py"),
           step("test pandad", "pytest selfdrive/pandad/tests/test_pandad.py", [diffPaths: ["panda", "selfdrive/pandad/"]]),
           step("test amp", "pytest system/hardware/tici/tests/test_amplifier.py"),
-          step("test qcomgpsd", "pytest system/qcomgpsd/tests/test_qcomgpsd.py"),
+          step("test qcomgpsd", "pytest system/qcomgpsd/tests/test_qcomgpsd.py", [diffPaths: ["system/qcomgpsd/"]]),
         ])
       },
 

README.md

@@ -38,7 +38,8 @@ Quick start: `bash <(curl -fsSL openpilot.comma.ai)`
   </tr>
 </table>
 
-To start using openpilot in a car
+
+Using openpilot in a car
 ------
 
 To use openpilot in a car, you need four things:
@@ -49,6 +50,14 @@ To use openpilot in a car, you need four things:
 
 We have detailed instructions for [how to install the harness and device in a car](https://comma.ai/setup). Note that it's possible to run openpilot on [other hardware](https://blog.comma.ai/self-driving-car-for-free/), although it's not plug-and-play.
 
+### Branches
+| branch           | URL                                    | description                                                                         |
+|------------------|----------------------------------------|-------------------------------------------------------------------------------------|
+| `release3`         | openpilot.comma.ai                      | This is openpilot's release branch.                                                 |
+| `release3-staging` | openpilot-test.comma.ai                | This is the staging branch for releases. Use it to get new releases slightly early. |
+| `nightly`          | openpilot-nightly.comma.ai             | This is the bleeding edge development branch. Do not expect this to be stable.      |
+| `nightly-dev`      | installer.comma.ai/commaai/nightly-dev | Same as nightly, but includes experimental development features for some cars.      |
+
 To start developing openpilot
 ------
 

SConstruct

@@ -349,7 +349,7 @@ Export('common', 'gpucommon')
 env_swaglog = env.Clone()
 env_swaglog['CXXFLAGS'].append('-DSWAGLOG="\\"common/swaglog.h\\""')
 SConscript(['msgq_repo/SConscript'], exports={'env': env_swaglog})
-SConscript(['opendbc/can/SConscript'], exports={'env': env_swaglog})
+SConscript(['opendbc_repo/SConscript'], exports={'env': env_swaglog})
 
 SConscript(['cereal/SConscript'])
 
@@ -366,6 +366,7 @@ SConscript(['rednose/SConscript'])
 
 # Build system services
 SConscript([
+  'system/ui/SConscript',
   'system/proclogd/SConscript',
   'system/ubloxd/SConscript',
   'system/loggerd/SConscript',

cereal/log.capnp

@@ -486,6 +486,9 @@ struct DeviceState @0xa4d8b5af2aa492eb {
   nvmeTempC @35 :List(Float32);
   modemTempC @36 :List(Float32);
   pmicTempC @39 :List(Float32);
+  intakeTempC @46 :Float32;
+  exhaustTempC @47 :Float32;
+  caseTempC @48 :Float32;
   maxTempC @44 :Float32;  # max of other temps, used to control fan
   thermalZones @38 :List(ThermalZone);
   thermalStatus @14 :ThermalStatus;
@@ -2437,6 +2440,14 @@ struct Microphone {
   filteredSoundPressureWeightedDb @2 :Float32;
 }
 
+struct Touch {
+  sec @0 :Int64;
+  usec @1 :Int64;
+  type @2 :UInt8;
+  code @3 :Int32;
+  value @4 :Int32;
+}
+
 struct Event {
   logMonoTime @0 :UInt64;  # nanoseconds
   valid @67 :Bool = true;
@@ -2517,6 +2528,9 @@ struct Event {
     logMessage @18 :Text;
     errorLogMessage @85 :Text;
 
+    # touch frame
+    touch @135 :List(Touch);
+
     # navigation
     navInstruction @82 :NavInstruction;
     navRoute @83 :NavRoute;

cereal/messaging/tests/test_services.py

@@ -17,5 +17,5 @@ class TestServices:
 
   def test_generated_header(self):
     with tempfile.NamedTemporaryFile(suffix=".h") as f:
-      ret = os.system(f"python3 {services.__file__} > {f.name} && clang++ {f.name}")
+      ret = os.system(f"python3 {services.__file__} > {f.name} && clang++ {f.name} -std=c++11")
       assert ret == 0, "generated services header is not valid C"

cereal/services.py

@@ -22,6 +22,7 @@ _services: dict[str, tuple] = {
   "temperatureSensor2": (True, 2., 200),
   "gpsNMEA": (True, 9.),
   "deviceState": (True, 2., 1),
+  "touch": (True, 20., 1),
   "can": (True, 100., 2053),  # decimation gives ~3 msgs in a full segment
   "controlsState": (True, 100., 10),
   "selfdriveState": (True, 100., 10),

common/params.cc

@@ -112,7 +112,6 @@ std::unordered_map<std::string, uint32_t> keys = {
     {"DisablePowerDown", PERSISTENT},
     {"DisableUpdates", PERSISTENT},
     {"DisengageOnAccelerator", PERSISTENT},
-    {"DmModelInitialized", CLEAR_ON_ONROAD_TRANSITION},
     {"DongleId", PERSISTENT},
     {"DoReboot", CLEAR_ON_MANAGER_START},
     {"DoShutdown", CLEAR_ON_MANAGER_START},

common/prefix.h

@@ -13,7 +13,7 @@ public:
     if (prefix.empty()) {
       prefix = util::random_string(15);
     }
-    msgq_path = "/dev/shm/" + prefix;
+    msgq_path = Path::shm_path() + "/" + prefix;
     bool ret = util::create_directories(msgq_path, 0777);
     assert(ret);
     setenv("OPENPILOT_PREFIX", prefix.c_str(), 1);

common/prefix.py

@@ -11,7 +11,7 @@ from openpilot.system.hardware.hw import DEFAULT_DOWNLOAD_CACHE_ROOT
 class OpenpilotPrefix:
   def __init__(self, prefix: str = None, clean_dirs_on_exit: bool = True, shared_download_cache: bool = False):
     self.prefix = prefix if prefix else str(uuid.uuid4().hex[0:15])
-    self.msgq_path = os.path.join('/dev/shm', self.prefix)
+    self.msgq_path = os.path.join(Paths.shm_path(), self.prefix)
     self.clean_dirs_on_exit = clean_dirs_on_exit
     self.shared_download_cache = shared_download_cache
 

common/realtime.py

@@ -48,13 +48,13 @@ class Ratekeeper:
   def __init__(self, rate: float, print_delay_threshold: float | None = 0.0) -> None:
     """Rate in Hz for ratekeeping. print_delay_threshold must be nonnegative."""
     self._interval = 1. / rate
-    self._next_frame_time = time.monotonic() + self._interval
     self._print_delay_threshold = print_delay_threshold
     self._frame = 0
     self._remaining = 0.0
     self._process_name = getproctitle()
     self._dts = deque([self._interval], maxlen=100)
-    self._last_monitor_time = time.monotonic()
+    self._last_monitor_time = -1.
+    self._next_frame_time = -1.
 
   @property
   def frame(self) -> int:
@@ -79,6 +79,10 @@ class Ratekeeper:
 
   # Monitors the cumulative lag, but does not enforce a rate
   def monitor_time(self) -> bool:
+    if self._last_monitor_time < 0:
+      self._next_frame_time = time.monotonic() + self._interval
+      self._last_monitor_time = time.monotonic()
+
     prev = self._last_monitor_time
     self._last_monitor_time = time.monotonic()
     self._dts.append(self._last_monitor_time - prev)

common/transformations/model.py

@@ -1,7 +1,7 @@
 import numpy as np
 
 from openpilot.common.transformations.orientation import rot_from_euler
-from openpilot.common.transformations.camera import get_view_frame_from_calib_frame, view_frame_from_device_frame
+from openpilot.common.transformations.camera import get_view_frame_from_calib_frame, view_frame_from_device_frame, _ar_ox_fisheye
 
 # segnet
 SEGNET_SIZE = (512, 384)
@@ -39,6 +39,13 @@ sbigmodel_intrinsics = np.array([
   [0.0,  sbigmodel_fl,      0.5 * (256 + MEDMODEL_CY)],
   [0.0,  0.0,                                     1.0]])
 
+DM_INPUT_SIZE = (1440, 960)
+dmonitoringmodel_fl = _ar_ox_fisheye.focal_length
+dmonitoringmodel_intrinsics = np.array([
+  [dmonitoringmodel_fl,  0.0, DM_INPUT_SIZE[0]/2],
+  [0.0, dmonitoringmodel_fl, DM_INPUT_SIZE[1]/2 - (_ar_ox_fisheye.height - DM_INPUT_SIZE[1])/2],
+  [0.0,  0.0, 1.0]])
+
 bigmodel_frame_from_calib_frame = np.dot(bigmodel_intrinsics,
   get_view_frame_from_calib_frame(0, 0, 0, 0))
 

common/watchdog.cc

@@ -2,8 +2,9 @@
 
 #include "common/watchdog.h"
 #include "common/util.h"
+#include "system/hardware/hw.h"
 
-const std::string watchdog_fn_prefix = "/dev/shm/wd_";  // + <pid>
+const std::string watchdog_fn_prefix = Path::shm_path() + "/wd_";  // + <pid>
 
 bool watchdog_kick(uint64_t ts) {
   static std::string fn = watchdog_fn_prefix + std::to_string(getpid());

docs/CARS.md

@@ -103,7 +103,7 @@ A supported vehicle is one that just works when you install a comma device. All
 |Hyundai|Ioniq Plug-in Hybrid 2020-22|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 H 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 Plug-in Hybrid 2020-22">Buy Here</a></sub></details>||
 |Hyundai|Kona 2020|Smart Cruise Control (SCC)|openpilot available[<sup>1</sup>](#footnotes)|6 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-empty.svg)](##)|<details><summary>Parts</summary><sub>- 1 Hyundai B 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=Kona 2020">Buy Here</a></sub></details>||
 |Hyundai|Kona Electric 2018-21|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 G 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=Kona Electric 2018-21">Buy Here</a></sub></details>||
-|Hyundai|Kona Electric 2022-23|Smart Cruise Control (SCC)|Stock|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 Hyundai O 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=Kona Electric 2022-23">Buy Here</a></sub></details>||
+|Hyundai|Kona Electric 2022-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 O 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=Kona Electric 2022-23">Buy Here</a></sub></details>||
 |Hyundai|Kona Electric (with HDA II, Korea only) 2023[<sup>5</sup>](#footnotes)|Smart Cruise Control (SCC)|Stock|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 Hyundai R 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=Kona Electric (with HDA II, Korea only) 2023">Buy Here</a></sub></details>|<a href="https://www.youtube.com/watch?v=U2fOCmcQ8hw" target="_blank"><img height="18px" src="assets/icon-youtube.svg"></img></a>|
 |Hyundai|Kona Hybrid 2020|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 I 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=Kona Hybrid 2020">Buy Here</a></sub></details>||
 |Hyundai|Palisade 2020-22|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 H 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=Palisade 2020-22">Buy Here</a></sub></details>|<a href="https://youtu.be/TAnDqjF4fDY?t=456" target="_blank"><img height="18px" src="assets/icon-youtube.svg"></img></a>|
@@ -167,7 +167,7 @@ A supported vehicle is one that just works when you install a comma device. All
 |Lexus|ES 2017-18|All|openpilot available[<sup>2</sup>](#footnotes)|19 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-empty.svg)](##)|<details><summary>Parts</summary><sub>- 1 RJ45 cable (7 ft)<br>- 1 Toyota A connector<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=Lexus&model=ES 2017-18">Buy Here</a></sub></details>||
 |Lexus|ES 2019-24|All|openpilot|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 RJ45 cable (7 ft)<br>- 1 Toyota A connector<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=Lexus&model=ES 2019-24">Buy Here</a></sub></details>||
 |Lexus|ES Hybrid 2017-18|All|openpilot available[<sup>2</sup>](#footnotes)|19 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-empty.svg)](##)|<details><summary>Parts</summary><sub>- 1 RJ45 cable (7 ft)<br>- 1 Toyota A connector<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=Lexus&model=ES Hybrid 2017-18">Buy Here</a></sub></details>||
-|Lexus|ES Hybrid 2019-24|All|openpilot|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 RJ45 cable (7 ft)<br>- 1 Toyota A connector<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=Lexus&model=ES Hybrid 2019-24">Buy Here</a></sub></details>|<a href="https://youtu.be/BZ29osRVJeg?t=12" target="_blank"><img height="18px" src="assets/icon-youtube.svg"></img></a>|
+|Lexus|ES Hybrid 2019-25|All|openpilot|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 RJ45 cable (7 ft)<br>- 1 Toyota A connector<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=Lexus&model=ES Hybrid 2019-25">Buy Here</a></sub></details>|<a href="https://youtu.be/BZ29osRVJeg?t=12" target="_blank"><img height="18px" src="assets/icon-youtube.svg"></img></a>|
 |Lexus|GS F 2016|All|Stock|19 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-empty.svg)](##)|<details><summary>Parts</summary><sub>- 1 RJ45 cable (7 ft)<br>- 1 Toyota A connector<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=Lexus&model=GS F 2016">Buy Here</a></sub></details>||
 |Lexus|IS 2017-19|All|Stock|19 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-empty.svg)](##)|<details><summary>Parts</summary><sub>- 1 RJ45 cable (7 ft)<br>- 1 Toyota A connector<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=Lexus&model=IS 2017-19">Buy Here</a></sub></details>||
 |Lexus|IS 2022-23|All|openpilot|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 RJ45 cable (7 ft)<br>- 1 Toyota A connector<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=Lexus&model=IS 2022-23">Buy Here</a></sub></details>||
@@ -188,7 +188,7 @@ A supported vehicle is one that just works when you install a comma device. All
 |Lincoln|Aviator Plug-in Hybrid 2020-24|Co-Pilot360 Plus|openpilot|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 Ford Q3 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=Lincoln&model=Aviator Plug-in Hybrid 2020-24">Buy Here</a></sub></details>||
 |MAN|eTGE 2020-24|Adaptive Cruise Control (ACC) & Lane Assist|openpilot available[<sup>1,12</sup>](#footnotes)|0 mph|31 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 angled mount (8 degrees)<br>- 1 comma 3X<br>- 1 harness box<br>- 1 long OBD-C cable<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=MAN&model=eTGE 2020-24">Buy Here</a></sub></details>|<a href="https://youtu.be/4100gLeabmo" target="_blank"><img height="18px" src="assets/icon-youtube.svg"></img></a>|
 |MAN|TGE 2017-24|Adaptive Cruise Control (ACC) & Lane Assist|openpilot available[<sup>1,12</sup>](#footnotes)|0 mph|31 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 angled mount (8 degrees)<br>- 1 comma 3X<br>- 1 harness box<br>- 1 long OBD-C cable<br>- 1 right angle OBD-C cable (1.5 ft)<br><a href="https://comma.ai/shop/comma-3x.html?make=MAN&model=TGE 2017-24">Buy Here</a></sub></details>|<a href="https://youtu.be/4100gLeabmo" target="_blank"><img height="18px" src="assets/icon-youtube.svg"></img></a>|
-|Mazda|CX-5 2022-24|All|Stock|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 Mazda 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=Mazda&model=CX-5 2022-24">Buy Here</a></sub></details>||
+|Mazda|CX-5 2022-25|All|Stock|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 Mazda 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=Mazda&model=CX-5 2022-25">Buy Here</a></sub></details>||
 |Mazda|CX-9 2021-23|All|Stock|0 mph|28 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-full.svg)](##)|<details><summary>Parts</summary><sub>- 1 Mazda 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=Mazda&model=CX-9 2021-23">Buy Here</a></sub></details>|<a href="https://youtu.be/dA3duO4a0O4" target="_blank"><img height="18px" src="assets/icon-youtube.svg"></img></a>|
 |Nissan|Altima 2019-20|ProPILOT Assist|Stock|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-empty.svg)](##)|<details><summary>Parts</summary><sub>- 1 Nissan B connector<br>- 1 RJ45 cable (7 ft)<br>- 1 USB-C coupler<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=Nissan&model=Altima 2019-20">Buy Here</a></sub></details>||
 |Nissan|Leaf 2018-23|ProPILOT Assist|Stock|0 mph|0 mph|[![star](assets/icon-star-full.svg)](##)|[![star](assets/icon-star-empty.svg)](##)|<details><summary>Parts</summary><sub>- 1 Nissan A connector<br>- 1 RJ45 cable (7 ft)<br>- 1 USB-C coupler<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=Nissan&model=Leaf 2018-23">Buy Here</a></sub></details>|<a href="https://youtu.be/vaMbtAh_0cY" target="_blank"><img height="18px" src="assets/icon-youtube.svg"></img></a>|

docs/css/tooltip.css

@@ -0,0 +1,44 @@
+[data-tooltip] {
+  position: relative;
+  display: inline-block;
+  border-bottom: 1px dotted black;
+}
+
+[data-tooltip] .tooltip-content {
+  width: max-content;
+  max-width: 25em;
+  position: absolute;
+  top: 100%;
+  left: 50%;
+  transform: translateX(-50%);
+  background-color: white;
+  color: #404040;
+  box-shadow: 0 4px 14px 0 rgba(0,0,0,.2), 0 0 0 1px rgba(0,0,0,.05);
+  padding: 10px;
+  font: 14px/1.5 Lato, proxima-nova, Helvetica Neue, Arial, sans-serif;
+  text-decoration: none;
+  opacity: 0;
+  visibility: hidden;
+  transition: opacity 0.1s, visibility 0s;
+  z-index: 1000;
+  pointer-events: none; /* Prevent accidental interaction */
+}
+
+[data-tooltip]:hover .tooltip-content {
+  opacity: 1;
+  visibility: visible;
+  pointer-events: auto; /* Allow interaction when visible */
+}
+
+.tooltip-content .tooltip-glossary-link {
+  display: inline-block;
+  margin-top: 8px;
+  font-size: 12px;
+  color: #007bff;
+  text-decoration: none;
+}
+
+.tooltip-content .tooltip-glossary-link:hover {
+  color: #0056b3;
+  text-decoration: underline;
+}

docs/hooks/glossary.py

@@ -0,0 +1,68 @@
+import re
+import tomllib
+
+def load_glossary(file_path="docs/glossary.toml"):
+  with open(file_path, "rb") as f:
+    glossary_data = tomllib.load(f)
+  return glossary_data.get("glossary", {})
+
+def generate_anchor_id(name):
+  return name.replace(" ", "-").replace("_", "-").lower()
+
+def format_markdown_term(name, definition):
+  anchor_id = generate_anchor_id(name)
+  markdown = f"* [**{name.replace('_', ' ').title()}**](#{anchor_id})"
+  if definition.get("abbreviation"):
+    markdown += f" *({definition['abbreviation']})*"
+  if definition.get("description"):
+    markdown += f": {definition['description']}\n"
+  return markdown
+
+def glossary_markdown(vocabulary):
+  markdown = ""
+  for category, terms in vocabulary.items():
+    markdown += f"## {category.replace('_', ' ').title()}\n\n"
+    for name, definition in terms.items():
+      markdown += format_markdown_term(name, definition)
+  return markdown
+
+def format_tooltip_html(term_key, definition, html):
+  display_term = term_key.replace("_", " ").title()
+  clean_description = re.sub(r"\[(.+)]\(.+\)", r"\1", definition["description"])
+  glossary_link = (
+    f"<a href='/concepts/glossary#{term_key}' class='tooltip-glossary-link' title='View in glossary'>Glossary🔗</a>"
+  )
+  return re.sub(
+    re.escape(display_term),
+    lambda
+    match: f"<span data-tooltip>{match.group(0)}<span class='tooltip-content'>{clean_description} {glossary_link}</span></span>",
+    html,
+    flags=re.IGNORECASE,
+  )
+
+def apply_tooltip(_term_key, _definition, pattern, html):
+  return re.sub(
+    pattern,
+    lambda match: format_tooltip_html(_term_key, _definition, match.group(0)),
+    html,
+    flags=re.IGNORECASE,
+  )
+
+def tooltip_html(vocabulary, html):
+  for _category, terms in vocabulary.items():
+    for term_key, definition in terms.items():
+      if definition.get("description"):
+        pattern = rf"(?<!\w){re.escape(term_key.replace('_', ' ').title())}(?![^<]*<\/a>)(?!\([^)]*\))"
+        html = apply_tooltip(term_key, definition, pattern, html)
+  return html
+
+# Page Hooks
+def on_page_markdown(markdown, **kwargs):
+  glossary = load_glossary()
+  return markdown.replace("{{GLOSSARY_DEFINITIONS}}", glossary_markdown(glossary))
+
+def on_page_content(html, **kwargs):
+  if kwargs.get("page").title == "Glossary":
+    return html
+  glossary = load_glossary()
+  return tooltip_html(glossary, html)

launch_env.sh

@@ -7,7 +7,7 @@ export OPENBLAS_NUM_THREADS=1
 export VECLIB_MAXIMUM_THREADS=1
 
 if [ -z "$AGNOS_VERSION" ]; then
-  export AGNOS_VERSION="11.2"
+  export AGNOS_VERSION="11.4"
 fi
 
 export STAGING_ROOT="/data/safe_staging"

mkdocs.yml

@@ -8,6 +8,10 @@ strict: true
 docs_dir: docs
 site_dir: docs_site/
 
+hooks:
+  - docs/hooks/glossary.py
+extra_css:
+  - css/tooltip.css
 theme:
   name: readthedocs
   navigation_depth: 3

pyproject.toml

@@ -31,6 +31,9 @@ dependencies = [
   # body / webrtcd
   "aiohttp",
   "aiortc",
+  # aiortc does not put an upper bound on pyopenssl and is now incompatible
+  # with the latest release
+  "pyopenssl < 24.3.0",
   "pyaudio",
 
   # panda
@@ -98,7 +101,6 @@ dev = [
   "azure-identity",
   "azure-storage-blob",
   "dictdiffer",
-  "flaky",
   "lru-dict",
   "matplotlib",
   "parameterized >=0.8, <0.9",
@@ -239,6 +241,7 @@ exclude = [
   "cereal",
   "panda",
   "opendbc",
+  "opendbc_repo",
   "rednose_repo",
   "tinygrad_repo",
   "teleoprtc",

release/release_files.py

@@ -32,7 +32,6 @@ blacklist = [
 
   ".git/",
   ".github/",
-  ".devcontainer/",
   "Darwin/",
   ".vscode",
 

scripts/code_stats.py

@@ -37,7 +37,7 @@ for f in sorted(pyf):
   lns = len(src.split("\n"))
   tree = ast.parse(src)
   Analyzer().visit(tree)
-  print("%5d %s %s" % (lns, f, xbit))
+  print(f"{lns:5d} {f} {xbit}")
   if 'test' in f:
     testlns += lns
   elif f.startswith(('tools/', 'scripts/', 'selfdrive/debug')):
@@ -47,8 +47,8 @@ for f in sorted(pyf):
   else:
     tlns += lns
 
-print("%d lines of openpilot python" % tlns)
-print("%d lines of car ports" % carlns)
-print("%d lines of tools/scripts/debug" % scriptlns)
-print("%d lines of tests" % testlns)
+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/lint/check_raylib_includes.sh

@@ -0,0 +1,10 @@
+#!/usr/bin/env bash
+
+FAIL=0
+
+if grep -n '#include "third_party/raylib/include/raylib\.h"' $@ | grep -v '^system/ui/raylib/raylib\.h'; then
+  echo -e "Bad raylib include found! Use '#include \"system/ui/raylib/raylib.h\"' instead\n"
+  FAIL=1
+fi
+
+exit $FAIL

scripts/lint/lint.sh

@@ -53,6 +53,7 @@ function run_tests() {
   run "check_shebang_scripts_are_executable" python3 -m pre_commit_hooks.check_shebang_scripts_are_executable $ALL_FILES
   run "check_shebang_format" $DIR/check_shebang_format.sh $ALL_FILES
   run "check_nomerge_comments" $DIR/check_nomerge_comments.sh $ALL_FILES
+  run "check_raylib_includes" $DIR/check_raylib_includes.sh $ALL_FILES
 
   if [[ -z "$FAST" ]]; then
     run "mypy" mypy $PYTHON_FILES

scripts/waste.py

@@ -16,9 +16,9 @@ def waste(core):
   j = 0
   while 1:
     if (i % 100) == 0:
-      setproctitle("%3d: %8d" % (core, i))
+      setproctitle(f"{core:3d}: {i:8d}")
       lt = time.monotonic()
-      print("%3d: %8d %f  %.2f" % (core, i, lt-st, j))
+      print(f"{core:3d}: {i:8d} {lt-st:f}  {j:.2f}")
       st = lt
     i += 1
     j = np.sum(np.matmul(m1, m2))

selfdrive/assets/img_driver_face_static.png

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

selfdrive/car/car_specific.py

@@ -148,7 +148,8 @@ class CarSpecificEvents:
       # To avoid re-engaging when openpilot cancels, check user engagement intention via buttons
       # Main button also can trigger an engagement on these cars
       self.cruise_buttons.append(any(ev.type in HYUNDAI_ENABLE_BUTTONS for ev in CS.buttonEvents))
-      events = self.create_common_events(CS, CS_prev, pcm_enable=self.CP.pcmCruise, allow_enable=any(self.cruise_buttons))
+      events = self.create_common_events(CS, CS_prev, extra_gears=(GearShifter.sport, GearShifter.manumatic),
+                                         pcm_enable=self.CP.pcmCruise, allow_enable=any(self.cruise_buttons))
 
       # low speed steer alert hysteresis logic (only for cars with steer cut off above 10 m/s)
       if CS.vEgo < (self.CP.minSteerSpeed + 2.) and self.CP.minSteerSpeed > 10.:

selfdrive/car/tests/test_docs.py

@@ -17,7 +17,7 @@ class TestCarDocs:
     with open(CARS_MD_OUT) as f:
       current_cars_md = f.read()
 
-    assert generated_cars_md == current_cars_md, "Run selfdrive/opcar/docs.py to update the compatibility documentation"
+    assert generated_cars_md == current_cars_md, "Run selfdrive/car/docs.py to update the compatibility documentation"
 
   def test_docs_diff(self):
     dump_path = os.path.join(BASEDIR, "selfdrive", "car", "tests", "cars_dump")

selfdrive/car/tests/test_models.py

@@ -394,7 +394,7 @@ class TestCarModelBase(unittest.TestCase):
       for msg in filter(lambda m: m.src in range(64), can.can):
         to_send = libpanda_py.make_CANPacket(msg.address, msg.src % 4, msg.dat)
         ret = self.safety.safety_rx_hook(to_send)
-        self.assertEqual(1, ret, f"safety rx failed ({ret=}): {to_send}")
+        self.assertEqual(1, ret, f"safety rx failed ({ret=}): {(msg.address, msg.src % 4)}")
 
       # Skip first frame so CS_prev is properly initialized
       if idx == 0:

selfdrive/controls/lib/drive_helpers.py

@@ -5,12 +5,15 @@ from openpilot.common.realtime import DT_CTRL
 MIN_SPEED = 1.0
 CONTROL_N = 17
 CAR_ROTATION_RADIUS = 0.0
+# This is a turn radius smaller than most cars can achieve
+MAX_CURVATURE = 0.2
 
 # EU guidelines
 MAX_LATERAL_JERK = 5.0
 MAX_VEL_ERR = 5.0
 
 def clip_curvature(v_ego, prev_curvature, new_curvature):
+  new_curvature = clip(new_curvature, -MAX_CURVATURE, MAX_CURVATURE)
   v_ego = max(MIN_SPEED, v_ego)
   max_curvature_rate = MAX_LATERAL_JERK / (v_ego**2) # inexact calculation, check https://github.com/commaai/openpilot/pull/24755
   safe_desired_curvature = clip(new_curvature,

selfdrive/debug/can_printer.py

@@ -28,7 +28,7 @@ def can_printer(bus, max_msg, addr, ascii_decode):
         x = binascii.hexlify(msgs[_addr][-1]).decode('ascii')
         freq = len(msgs[_addr]) / (time.monotonic() - start)
         if max_msg is None or _addr < max_msg:
-          dd += "%04X(%4d)(%6d)(%3dHz) %s %s\n" % (_addr, _addr, len(msgs[_addr]), freq, x.ljust(20), a)
+          dd += f"{_addr:04X}({_addr:4d})({len(msgs[_addr]):6d})({freq:3}dHz) {x.ljust(20)} {a}\n"
       print(dd)
       lp = time.monotonic()
 

selfdrive/debug/check_can_parser_performance.py

@@ -29,7 +29,7 @@ if __name__ == '__main__':
     start_t = time.process_time_ns()
     for msg in msgs:
       can_list = can_capnp_to_list([msg])
-      for cp in tm.CI.can_parsers:
+      for cp in tm.CI.can_parsers.values():
         if cp is not None:
           cp.update_strings(can_list)
     ets.append((time.process_time_ns() - start_t) * 1e-6)

selfdrive/debug/count_events.py

@@ -8,6 +8,7 @@ from typing import cast
 
 from cereal.services import SERVICE_LIST
 from openpilot.tools.lib.logreader import LogReader, ReadMode
+from openpilot.selfdrive.test.process_replay.migration import migrate_all
 
 if __name__ == "__main__":
   cnt_events: Counter = Counter()
@@ -20,7 +21,7 @@ if __name__ == "__main__":
   start_time = math.inf
   end_time = -math.inf
   ignition_off = None
-  for msg in LogReader(sys.argv[1], ReadMode.QLOG):
+  for msg in migrate_all(LogReader(sys.argv[1], ReadMode.QLOG)):
     t = (msg.logMonoTime - start_time) / 1e9
     end_time = max(end_time, msg.logMonoTime)
     start_time = min(start_time, msg.logMonoTime)

selfdrive/debug/fingerprint_from_route.py

@@ -22,7 +22,7 @@ def get_fingerprint(lr):
           msgs[c.address] = len(c.dat)
 
   # show CAN fingerprint
-  fingerprint = ', '.join("%d: %d" % v for v in sorted(msgs.items()))
+  fingerprint = ', '.join(f"{v[0]}: {v[1]}" for v in sorted(msgs.items()))
   print(f"\nfound {len(msgs)} messages. CAN fingerprint:\n")
   print(fingerprint)
 

selfdrive/debug/get_fingerprint.py

@@ -25,7 +25,7 @@ while True:
     if c.src % 0x80 == 0 and c.address < 0x800 and c.address not in (0x7df, 0x7e0, 0x7e8):
       msgs[c.address] = len(c.dat)
 
-  fingerprint = ', '.join("%d: %d" % v for v in sorted(msgs.items()))
+  fingerprint = ', '.join(f"{v[0]}: {v[1]}" for v in sorted(msgs.items()))
 
   print(f"number of messages {len(msgs)}:")
   print(f"fingerprint {fingerprint}")

selfdrive/debug/touch_replay.py

@@ -0,0 +1,54 @@
+#!/usr/bin/env python3
+import argparse
+
+import numpy as np
+import matplotlib.pyplot as plt
+
+from openpilot.tools.lib.logreader import LogReader
+
+if __name__ == '__main__':
+  parser = argparse.ArgumentParser()
+  parser.add_argument('--width', default=2160, type=int)
+  parser.add_argument('--height', default=1080, type=int)
+  parser.add_argument('--route', default='rlog', type=str)
+  args = parser.parse_args()
+
+  w = args.width
+  h = args.height
+  route = args.route
+
+  fingers = [[-1, -1]] * 5
+  touch_points = []
+  current_slot = 0
+
+  lr = list(LogReader(route))
+  for msg in lr:
+    if msg.which() == 'touch':
+      for event in msg.touch:
+        if event.type == 3 and event.code == 47:
+          current_slot = event.value
+        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)
+          if fingers[current_slot][0] != -1:
+            touch_points.append(fingers[current_slot].copy())
+        elif event.type == 3 and event.code == 54:
+          fingers[current_slot][0] = w - event.value
+          if fingers[current_slot][1] != -1:
+            touch_points.append(fingers[current_slot].copy())
+
+  if not touch_points:
+    print(f'No touch events found for {route}')
+    quit()
+
+  unique_points, counts = np.unique(touch_points, axis=0, return_counts=True)
+
+  plt.figure(figsize=(10, 3))
+  plt.scatter(unique_points[:, 0], unique_points[:, 1], c=counts, s=counts * 20, edgecolors='red')
+  plt.colorbar()
+  plt.title(f'Touches for {route}')
+  plt.xlim(0, w)
+  plt.ylim(0, h)
+  plt.grid(True)
+  plt.show()

selfdrive/locationd/locationd.py

@@ -8,6 +8,7 @@ from enum import Enum
 from collections import defaultdict
 
 from cereal import log, messaging
+from cereal.services import SERVICE_LIST
 from openpilot.common.transformations.orientation import rot_from_euler
 from openpilot.common.realtime import config_realtime_process
 from openpilot.common.params import Params
@@ -23,8 +24,10 @@ MIN_STD_SANITY_CHECK = 1e-5  # m or rad
 MAX_FILTER_REWIND_TIME = 0.8  # s
 MAX_SENSOR_TIME_DIFF = 0.1  # s
 YAWRATE_CROSS_ERR_CHECK_FACTOR = 30
-INPUT_INVALID_THRESHOLD = 0.5
-INPUT_INVALID_DECAY = 0.9993  # ~10 secs to resume after a bad input
+INPUT_INVALID_THRESHOLD = 0.5  # 0 bad inputs ignored
+TIMING_INVALID_THRESHOLD = 2.5  # 2 bad timings ignored
+INPUT_INVALID_DECAY = 0.9993  # ~10 secs to resume after exceeding allowed bad inputs by one (at 100hz)
+TIMING_INVALID_DECAY = 0.9990  # ~2 secs to resume after exceeding allowed bad timings by one (at 100hz)
 POSENET_STD_INITIAL_VALUE = 10.0
 POSENET_STD_HIST_HALF = 20
 
@@ -265,10 +268,13 @@ def main():
   estimator = LocationEstimator(DEBUG)
 
   filter_initialized = False
-  critcal_services = ["accelerometer", "gyroscope", "liveCalibration", "cameraOdometry"]
-  observation_timing_invalid = False
+  critcal_services = ["accelerometer", "gyroscope", "cameraOdometry"]
+  observation_timing_invalid = defaultdict(int)
   observation_input_invalid = defaultdict(int)
 
+  input_invalid_decay = {s: INPUT_INVALID_DECAY ** (100. / SERVICE_LIST[s].frequency) for s in critcal_services}
+  timing_invalid_decay = {s: TIMING_INVALID_DECAY ** (100. / SERVICE_LIST[s].frequency) for s in critcal_services}
+
   initial_pose = params.get("LocationFilterInitialState")
   if initial_pose is not None:
     initial_pose = json.loads(initial_pose)
@@ -282,8 +288,6 @@ def main():
     acc_msgs, gyro_msgs = (messaging.drain_sock(sock) for sock in sensor_sockets)
 
     if filter_initialized:
-      observation_timing_invalid = False
-
       msgs = []
       for msg in acc_msgs + gyro_msgs:
         t, valid, which, data = msg.logMonoTime, msg.valid, msg.which(), getattr(msg, msg.which())
@@ -298,18 +302,23 @@ def main():
         if valid:
           t = log_mono_time * 1e-9
           res = estimator.handle_log(t, which, msg)
+          if which not in critcal_services:
+            continue
+
           if res == HandleLogResult.TIMING_INVALID:
-            observation_timing_invalid = True
+            observation_timing_invalid[which] += 1
           elif res == HandleLogResult.INPUT_INVALID:
             observation_input_invalid[which] += 1
           else:
-            observation_input_invalid[which] *= INPUT_INVALID_DECAY
+            observation_input_invalid[which] *= input_invalid_decay[which]
+            observation_timing_invalid[which] *= timing_invalid_decay[which]
     else:
       filter_initialized = sm.all_checks() and sensor_all_checks(acc_msgs, gyro_msgs, sensor_valid, sensor_recv_time, sensor_alive, SIMULATION)
 
     if sm.updated["cameraOdometry"]:
       critical_service_inputs_valid = all(observation_input_invalid[s] < INPUT_INVALID_THRESHOLD for s in critcal_services)
-      inputs_valid = sm.all_valid() and critical_service_inputs_valid and not observation_timing_invalid
+      critical_service_timing_valid = all(observation_timing_invalid[s] < TIMING_INVALID_THRESHOLD for s in critcal_services)
+      inputs_valid = sm.all_valid() and critical_service_inputs_valid and critical_service_timing_valid
       sensors_valid = sensor_all_checks(acc_msgs, gyro_msgs, sensor_valid, sensor_recv_time, sensor_alive, SIMULATION)
 
       msg = estimator.get_msg(sensors_valid, inputs_valid, filter_initialized)

selfdrive/locationd/test/test_locationd_scenarios.py

@@ -1,4 +1,3 @@
-import pytest
 import numpy as np
 from collections import defaultdict
 from enum import Enum
@@ -17,6 +16,7 @@ SELECT_COMPARE_FIELDS = {
   'sensors_flag': ['sensorsOK'],
 }
 JUNK_IDX = 100
+CONSISTENT_SPIKES_COUNT = 10
 
 
 class Scenario(Enum):
@@ -25,6 +25,8 @@ class Scenario(Enum):
   GYRO_SPIKE_MIDWAY = 'gyro_spike_midway'
   ACCEL_OFF = 'accel_off'
   ACCEL_SPIKE_MIDWAY = 'accel_spike_midway'
+  SENSOR_TIMING_SPIKE_MIDWAY = 'timing_spikes'
+  SENSOR_TIMING_CONSISTENT_SPIKES = 'timing_consistent_spikes'
 
 
 def get_select_fields_data(logs):
@@ -43,6 +45,17 @@ def get_select_fields_data(logs):
   return data
 
 
+def modify_logs_midway(logs, which, count, fn):
+  non_which = [x for x in logs if x.which() != which]
+  which = [x for x in logs if x.which() == which]
+  temps = which[len(which) // 2:len(which) // 2 + count]
+  for i, temp in enumerate(temps):
+    temp = temp.as_builder()
+    fn(temp)
+    which[len(which) // 2 + i] = temp.as_reader()
+  return sorted(non_which + which, key=lambda x: x.logMonoTime)
+
+
 def run_scenarios(scenario, logs):
   if scenario == Scenario.BASE:
     pass
@@ -51,30 +64,28 @@ def run_scenarios(scenario, logs):
     logs = sorted([x for x in logs if x.which() != 'gyroscope'], key=lambda x: x.logMonoTime)
 
   elif scenario == Scenario.GYRO_SPIKE_MIDWAY:
-    non_gyro = [x for x in logs if x.which() not in 'gyroscope']
-    gyro = [x for x in logs if x.which() in 'gyroscope']
-    temp = gyro[len(gyro) // 2].as_builder()
-    temp.gyroscope.gyroUncalibrated.v[0] += 3.0
-    gyro[len(gyro) // 2] = temp.as_reader()
-    logs = sorted(non_gyro + gyro, key=lambda x: x.logMonoTime)
+    def gyro_spike(msg):
+      msg.gyroscope.gyroUncalibrated.v[0] += 3.0
+    logs = modify_logs_midway(logs, 'gyroscope', 1, gyro_spike)
 
   elif scenario == Scenario.ACCEL_OFF:
     logs = sorted([x for x in logs if x.which() != 'accelerometer'], key=lambda x: x.logMonoTime)
 
   elif scenario == Scenario.ACCEL_SPIKE_MIDWAY:
-    non_accel = [x for x in logs if x.which() not in 'accelerometer']
-    accel = [x for x in logs if x.which() in 'accelerometer']
-    temp = accel[len(accel) // 2].as_builder()
-    temp.accelerometer.acceleration.v[0] += 10.0
-    accel[len(accel) // 2] = temp.as_reader()
-    logs = sorted(non_accel + accel, key=lambda x: x.logMonoTime)
+    def acc_spike(msg):
+      msg.accelerometer.acceleration.v[0] += 10.0
+    logs = modify_logs_midway(logs, 'accelerometer', 1, acc_spike)
+
+  elif scenario == Scenario.SENSOR_TIMING_SPIKE_MIDWAY or scenario == Scenario.SENSOR_TIMING_CONSISTENT_SPIKES:
+    def timing_spike(msg):
+      msg.accelerometer.timestamp -= int(0.150 * 1e9)
+    count = 1 if scenario == Scenario.SENSOR_TIMING_SPIKE_MIDWAY else CONSISTENT_SPIKES_COUNT
+    logs = modify_logs_midway(logs, 'accelerometer', count, timing_spike)
 
   replayed_logs = replay_process_with_name(name='locationd', lr=logs)
   return get_select_fields_data(logs), get_select_fields_data(replayed_logs)
 
 
-@pytest.mark.xdist_group("test_locationd_scenarios")
-@pytest.mark.shared_download_cache
 class TestLocationdScenarios:
   """
   Test locationd with different scenarios. In all these scenarios, we expect the following:
@@ -122,7 +133,7 @@ class TestLocationdScenarios:
     assert np.allclose(orig_data['yaw_rate'], replayed_data['yaw_rate'], atol=np.radians(0.35))
     assert np.allclose(orig_data['roll'], replayed_data['roll'], atol=np.radians(0.55))
     assert np.diff(replayed_data['inputs_flag'])[499] == -1.0
-    assert np.diff(replayed_data['inputs_flag'])[696] == 1.0
+    assert np.diff(replayed_data['inputs_flag'])[704] == 1.0
 
   def test_accel_off(self):
     """
@@ -146,3 +157,21 @@ class TestLocationdScenarios:
     orig_data, replayed_data = run_scenarios(Scenario.ACCEL_SPIKE_MIDWAY, self.logs)
     assert np.allclose(orig_data['yaw_rate'], replayed_data['yaw_rate'], atol=np.radians(0.35))
     assert np.allclose(orig_data['roll'], replayed_data['roll'], atol=np.radians(0.55))
+
+  def test_single_timing_spike(self):
+    """
+    Test: timing of 150ms off for the single accelerometer message in the middle of the segment
+    Expected Result: the message is ignored, and inputsOK is False for that time
+    """
+    orig_data, replayed_data = run_scenarios(Scenario.SENSOR_TIMING_SPIKE_MIDWAY, self.logs)
+    assert np.all(replayed_data['inputs_flag'] == orig_data['inputs_flag'])
+    assert np.all(replayed_data['sensors_flag'] == orig_data['sensors_flag'])
+
+  def test_consistent_timing_spikes(self):
+    """
+    Test: consistent timing spikes for N accelerometer messages in the middle of the segment
+    Expected Result: inputsOK becomes False after N of bad measurements
+    """
+    orig_data, replayed_data = run_scenarios(Scenario.SENSOR_TIMING_CONSISTENT_SPIKES, self.logs)
+    assert np.diff(replayed_data['inputs_flag'])[500] == -1.0
+    assert np.diff(replayed_data['inputs_flag'])[787] == 1.0

selfdrive/modeld/dmonitoringmodeld.py

@@ -12,15 +12,14 @@ from cereal import messaging
 from cereal.messaging import PubMaster, SubMaster
 from msgq.visionipc import VisionIpcClient, VisionStreamType, VisionBuf
 from openpilot.common.swaglog import cloudlog
-from openpilot.common.params import Params
 from openpilot.common.realtime import set_realtime_priority
+from openpilot.common.transformations.model import dmonitoringmodel_intrinsics
+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.models.commonmodel_pyx import CLContext
 from openpilot.selfdrive.modeld.parse_model_outputs import sigmoid
 
 CALIB_LEN = 3
-MODEL_WIDTH = 1440
-MODEL_HEIGHT = 960
 FEATURE_LEN = 512
 OUTPUT_SIZE = 84 + FEATURE_LEN
 
@@ -63,25 +62,21 @@ class ModelState:
 
   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 = {
-      'input_img': np.zeros(MODEL_HEIGHT * MODEL_WIDTH, dtype=np.uint8),
       'calib': np.zeros(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'])
 
-  def run(self, buf:VisionBuf, calib:np.ndarray) -> tuple[np.ndarray, float]:
+  def run(self, buf:VisionBuf, calib:np.ndarray, transform:np.ndarray) -> tuple[np.ndarray, float]:
     self.inputs['calib'][:] = calib
 
-    v_offset = buf.height - MODEL_HEIGHT
-    h_offset = (buf.width - MODEL_WIDTH) // 2
-    buf_data = buf.data.reshape(-1, buf.stride)
-    input_data = self.inputs['input_img'].reshape(MODEL_HEIGHT, MODEL_WIDTH)
-    input_data[:] = buf_data[v_offset:v_offset+MODEL_HEIGHT, h_offset:h_offset+MODEL_WIDTH]
+    self.model.setInputBuffer("input_img", self.frame.prepare(buf, transform.flatten(), None).view(np.float32))
 
-    self.model.setInputBuffer("input_img", self.inputs['input_img'].view(np.float32))
     t1 = time.perf_counter()
     self.model.execute()
     t2 = time.perf_counter()
@@ -126,7 +121,6 @@ def main():
   cl_context = CLContext()
   model = ModelState(cl_context)
   cloudlog.warning("models loaded, dmonitoringmodeld starting")
-  Params().put_bool("DmModelInitialized", True)
 
   cloudlog.warning("connecting to driver stream")
   vipc_client = VisionIpcClient("camerad", VisionStreamType.VISION_STREAM_DRIVER, True, cl_context)
@@ -139,18 +133,23 @@ def main():
   pm = PubMaster(["driverStateV2"])
 
   calib = np.zeros(CALIB_LEN, dtype=np.float32)
+  model_transform = None
 
   while True:
     buf = vipc_client.recv()
     if buf is None:
       continue
 
+    if model_transform is None:
+      cam = _os_fisheye if buf.width == _os_fisheye.width else _ar_ox_fisheye
+      model_transform = np.linalg.inv(np.dot(dmonitoringmodel_intrinsics, np.linalg.inv(cam.intrinsics))).astype(np.float32)
+
     sm.update(0)
     if sm.updated["liveCalibration"]:
       calib[:] = np.array(sm["liveCalibration"].rpyCalib)
 
     t1 = time.perf_counter()
-    model_output, gpu_execution_time = model.run(buf, calib)
+    model_output, gpu_execution_time = model.run(buf, calib, model_transform)
     t2 = time.perf_counter()
 
     pm.send("driverStateV2", get_driverstate_packet(model_output, vipc_client.frame_id, vipc_client.timestamp_sof, t2 - t1, gpu_execution_time))

selfdrive/modeld/fill_model_msg.py

@@ -3,11 +3,22 @@ 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):
     self.disengage_buffer = np.zeros(ModelConstants.CONFIDENCE_BUFFER_LEN*ModelConstants.DISENGAGE_WIDTH, dtype=np.float32)
@@ -55,14 +66,17 @@ def fill_lane_line_meta(builder, lane_lines, lane_line_probs):
   builder.rightProb = lane_line_probs[2]
 
 def fill_model_msg(base_msg: capnp._DynamicStructBuilder, extended_msg: capnp._DynamicStructBuilder,
-                   net_output_data: dict[str, np.ndarray], publish_state: PublishState,
-                   vipc_frame_id: int, vipc_frame_id_extra: int, frame_id: int, frame_drop: float,
-                   timestamp_eof: int, model_execution_time: float, valid: bool) -> None:
+                   net_output_data: dict[str, np.ndarray], v_ego: float, delay: float,
+                   publish_state: PublishState, vipc_frame_id: int, vipc_frame_id_extra: int,
+                   frame_id: int, frame_drop: float, timestamp_eof: int, model_execution_time: float,
+                   valid: bool) -> None:
   frame_age = frame_id - vipc_frame_id if frame_id > vipc_frame_id else 0
   frame_drop_perc = frame_drop * 100
   extended_msg.valid = valid
   base_msg.valid = valid
 
+  desired_curv = float(get_curvature_from_plan(net_output_data['plan'][0], v_ego, delay))
+
   driving_model_data = base_msg.drivingModelData
 
   driving_model_data.frameId = vipc_frame_id
@@ -71,7 +85,7 @@ def fill_model_msg(base_msg: capnp._DynamicStructBuilder, extended_msg: capnp._D
   driving_model_data.modelExecutionTime = model_execution_time
 
   action = driving_model_data.action
-  action.desiredCurvature = float(net_output_data['desired_curvature'][0,0])
+  action.desiredCurvature = desired_curv
 
   modelV2 = extended_msg.modelV2
   modelV2.frameId = vipc_frame_id
@@ -106,7 +120,7 @@ def fill_model_msg(base_msg: capnp._DynamicStructBuilder, extended_msg: capnp._D
 
   # lateral planning
   action = modelV2.action
-  action.desiredCurvature = float(net_output_data['desired_curvature'][0,0])
+  action.desiredCurvature = desired_curv
 
   # times at X_IDXS according to model plan
   PLAN_T_IDXS = [np.nan] * ModelConstants.IDX_N

selfdrive/modeld/modeld.py

@@ -12,6 +12,8 @@ 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
@@ -22,7 +24,7 @@ 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 ModelFrame, CLContext
+from openpilot.selfdrive.modeld.models.commonmodel_pyx import DrivingModelFrame, CLContext
 
 PROCESS_NAME = "selfdrive.modeld.modeld"
 SEND_RAW_PRED = os.getenv('SEND_RAW_PRED')
@@ -44,33 +46,30 @@ class FrameMeta:
       self.frame_id, self.timestamp_sof, self.timestamp_eof = vipc.frame_id, vipc.timestamp_sof, vipc.timestamp_eof
 
 class ModelState:
-  frame: ModelFrame
-  wide_frame: ModelFrame
+  frame: DrivingModelFrame
+  wide_frame: DrivingModelFrame
   inputs: dict[str, np.ndarray]
   output: np.ndarray
   prev_desire: np.ndarray  # for tracking the rising edge of the pulse
   model: ModelRunner
 
   def __init__(self, context: CLContext):
-    self.frame = ModelFrame(context)
-    self.wide_frame = ModelFrame(context)
+    self.frame = DrivingModelFrame(context)
+    self.wide_frame = DrivingModelFrame(context)
     self.prev_desire = np.zeros(ModelConstants.DESIRE_LEN, dtype=np.float32)
     self.full_features_20Hz = np.zeros((ModelConstants.FULL_HISTORY_BUFFER_LEN, ModelConstants.FEATURE_LEN), dtype=np.float32)
     self.desire_20Hz =  np.zeros((ModelConstants.FULL_HISTORY_BUFFER_LEN + 1, ModelConstants.DESIRE_LEN), dtype=np.float32)
-    self.prev_desired_curv_20hz = np.zeros((ModelConstants.FULL_HISTORY_BUFFER_LEN + 1, ModelConstants.PREV_DESIRED_CURV_LEN), dtype=np.float32)
 
     # img buffers are managed in openCL transform code
-    self.inputs = {
-      'desire': np.zeros(ModelConstants.DESIRE_LEN * (ModelConstants.HISTORY_BUFFER_LEN+1), dtype=np.float32),
-      'traffic_convention': np.zeros(ModelConstants.TRAFFIC_CONVENTION_LEN, dtype=np.float32),
-      'lateral_control_params': np.zeros(ModelConstants.LATERAL_CONTROL_PARAMS_LEN, dtype=np.float32),
-      'prev_desired_curv': np.zeros(ModelConstants.PREV_DESIRED_CURV_LEN * (ModelConstants.HISTORY_BUFFER_LEN+1), dtype=np.float32),
-      'features_buffer': np.zeros(ModelConstants.HISTORY_BUFFER_LEN * ModelConstants.FEATURE_LEN, dtype=np.float32),
-    }
+    self.inputs = {}
 
     with open(METADATA_PATH, 'rb') as f:
       model_metadata = pickle.load(f)
 
+    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)
@@ -82,6 +81,14 @@ class ModelState:
     for k,v in self.inputs.items():
       self.model.addInput(k, v)
 
+    num_elements = model_metadata['input_shapes']['features_buffer'][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:
@@ -97,10 +104,9 @@ class ModelState:
 
     self.desire_20Hz[:-1] = self.desire_20Hz[1:]
     self.desire_20Hz[-1] = new_desire
-    self.inputs['desire'][:] = self.desire_20Hz.reshape((25,4,-1)).max(axis=1).flatten()
+    self.inputs['desire'][:] = self.desire_20Hz.reshape(self.desire_reshape_dims).max(axis=1).flatten()
 
     self.inputs['traffic_convention'][:] = inputs['traffic_convention']
-    self.inputs['lateral_control_params'][:] = inputs['lateral_control_params']
 
     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")))
@@ -114,13 +120,25 @@ class ModelState:
     self.full_features_20Hz[:-1] = self.full_features_20Hz[1:]
     self.full_features_20Hz[-1] = outputs['hidden_state'][0, :]
 
-    self.prev_desired_curv_20hz[:-1] = self.prev_desired_curv_20hz[1:]
-    self.prev_desired_curv_20hz[-1] = outputs['desired_curvature'][0, :]
+    # idxs = np.arange(-4,-100,-4)[::-1]
+    self.inputs['features_buffer'][:] = self.full_features_20Hz[self.full_features_20Hz_idxs].flatten()
 
-    idxs = np.arange(-4,-100,-4)[::-1]
-    self.inputs['features_buffer'][:] = self.full_features_20Hz[idxs].flatten()
-    # TODO model only uses last value now, once that changes we need to input strided action history buffer
-    self.inputs['prev_desired_curv'][-ModelConstants.PREV_DESIRED_CURV_LEN:] = 0. * self.prev_desired_curv_20hz[-4, :]
+    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 "desired_curvature" in outputs:
+      input_name_prev = None
+      if "prev_desired_curvs" in self.inputs.keys():
+        input_name_prev = 'prev_desired_curvs'
+      elif "prev_desired_curv" in self.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]
     return outputs
 
 
@@ -231,7 +249,6 @@ def main(demo=False):
     is_rhd = sm["driverMonitoringState"].isRHD
     frame_id = sm["roadCameraState"].frameId
     v_ego = max(sm["carState"].vEgo, 0.)
-    lateral_control_params = np.array([v_ego, steer_delay], dtype=np.float32)
     if sm.updated["liveCalibration"] and sm.seen['roadCameraState'] and sm.seen['deviceState']:
       device_from_calib_euler = np.array(sm["liveCalibration"].rpyCalib, dtype=np.float32)
       dc = DEVICE_CAMERAS[(str(sm['deviceState'].deviceType), str(sm['roadCameraState'].sensor))]
@@ -262,9 +279,20 @@ def main(demo=False):
     inputs:dict[str, np.ndarray] = {
       'desire': vec_desire,
       'traffic_convention': traffic_convention,
-      'lateral_control_params': lateral_control_params,
       }
 
+    if "lateral_control_params" in model.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()
@@ -274,7 +302,8 @@ def main(demo=False):
       modelv2_send = messaging.new_message('modelV2')
       drivingdata_send = messaging.new_message('drivingModelData')
       posenet_send = messaging.new_message('cameraOdometry')
-      fill_model_msg(drivingdata_send, modelv2_send, model_output, publish_state, meta_main.frame_id, meta_extra.frame_id, frame_id,
+      fill_model_msg(drivingdata_send, modelv2_send, model_output, v_ego, steer_delay,
+                     publish_state, meta_main.frame_id, meta_extra.frame_id, frame_id,
                      frame_drop_ratio, meta_main.timestamp_eof, model_execution_time, live_calib_seen)
 
       desire_state = modelv2_send.modelV2.meta.desireState

selfdrive/modeld/models/commonmodel.cc

@@ -1,36 +1,30 @@
 #include "selfdrive/modeld/models/commonmodel.h"
 
-#include <cassert>
 #include <cmath>
 #include <cstring>
 
 #include "common/clutil.h"
 
-ModelFrame::ModelFrame(cl_device_id device_id, cl_context context) {
-  input_frames = std::make_unique<uint8_t[]>(buf_size);
-
-  q = CL_CHECK_ERR(clCreateCommandQueue(context, device_id, 0, &err));
-  y_cl = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_WRITE, MODEL_WIDTH * MODEL_HEIGHT, NULL, &err));
-  u_cl = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_WRITE, (MODEL_WIDTH / 2) * (MODEL_HEIGHT / 2), NULL, &err));
-  v_cl = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_WRITE, (MODEL_WIDTH / 2) * (MODEL_HEIGHT / 2), NULL, &err));
+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;
   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));
 
-  transform_init(&transform, context, device_id);
   loadyuv_init(&loadyuv, context, device_id, MODEL_WIDTH, MODEL_HEIGHT);
+  init_transform(device_id, context, MODEL_WIDTH, MODEL_HEIGHT);
 }
 
-uint8_t* ModelFrame::prepare(cl_mem yuv_cl, int frame_width, int frame_height, int frame_stride, int frame_uv_offset, const mat3 &projection, cl_mem *output) {
-  transform_queue(&this->transform, q,
-                yuv_cl, frame_width, frame_height, frame_stride, frame_uv_offset,
-                y_cl, u_cl, v_cl, MODEL_WIDTH, MODEL_HEIGHT, projection);
+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) {
+  run_transform(yuv_cl, MODEL_WIDTH, MODEL_HEIGHT, frame_width, frame_height, frame_stride, frame_uv_offset, projection);
 
   for (int i = 0; i < 4; i++) {
     CL_CHECK(clEnqueueCopyBuffer(q, img_buffer_20hz_cl, img_buffer_20hz_cl, (i+1)*frame_size_bytes, i*frame_size_bytes, frame_size_bytes, 0, nullptr, nullptr));
   }
   loadyuv_queue(&loadyuv, q, y_cl, u_cl, v_cl, last_img_cl);
+
   if (output == NULL) {
     CL_CHECK(clEnqueueReadBuffer(q, img_buffer_20hz_cl, CL_TRUE, 0, frame_size_bytes, &input_frames[0], 0, nullptr, nullptr));
     CL_CHECK(clEnqueueReadBuffer(q, last_img_cl, CL_TRUE, 0, frame_size_bytes, &input_frames[MODEL_FRAME_SIZE], 0, nullptr, nullptr));
@@ -46,13 +40,30 @@ uint8_t* ModelFrame::prepare(cl_mem yuv_cl, int frame_width, int frame_height, i
   }
 }
 
-ModelFrame::~ModelFrame() {
-  transform_destroy(&transform);
+DrivingModelFrame::~DrivingModelFrame() {
+  deinit_transform();
   loadyuv_destroy(&loadyuv);
   CL_CHECK(clReleaseMemObject(img_buffer_20hz_cl));
   CL_CHECK(clReleaseMemObject(last_img_cl));
-  CL_CHECK(clReleaseMemObject(v_cl));
-  CL_CHECK(clReleaseMemObject(u_cl));
-  CL_CHECK(clReleaseMemObject(y_cl));
   CL_CHECK(clReleaseCommandQueue(q));
-}
+}
+
+
+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));
+
+  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) {
+  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();
+}
+
+MonitoringModelFrame::~MonitoringModelFrame() {
+  deinit_transform();
+  CL_CHECK(clReleaseCommandQueue(q));
+}

selfdrive/modeld/models/commonmodel.h

@@ -2,6 +2,7 @@
 
 #include <cfloat>
 #include <cstdlib>
+#include <cassert>
 
 #include <memory>
 
@@ -18,21 +19,80 @@
 
 class ModelFrame {
 public:
-  ModelFrame(cl_device_id device_id, cl_context context);
-  ~ModelFrame();
-  uint8_t* prepare(cl_mem yuv_cl, int width, int height, int frame_stride, int frame_uv_offset, const mat3& transform, cl_mem *output);
+  ModelFrame(cl_device_id device_id, cl_context context) {
+    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; }
+  /*
+  uint8_t* buffer_from_cl(cl_mem *in_frames, int buffer_size) {
+    CL_CHECK(clEnqueueReadBuffer(q, *in_frames, CL_TRUE, 0, buffer_size, input_frames.get(), 0, nullptr, nullptr));
+    clFinish(q);
+    return &input_frames[0];
+  }
+  */
+
+  int MODEL_WIDTH;
+  int MODEL_HEIGHT;
+  int MODEL_FRAME_SIZE;
+  int buf_size;
+
+protected:
+  cl_mem y_cl, u_cl, v_cl;
+  Transform transform;
+  cl_command_queue q;
+  std::unique_ptr<float[]> input_frames;
+
+  void init_transform(cl_device_id device_id, cl_context context, int model_width, int model_height) {
+    y_cl = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_WRITE, model_width * model_height, NULL, &err));
+    u_cl = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_WRITE, (model_width / 2) * (model_height / 2), NULL, &err));
+    v_cl = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_WRITE, (model_width / 2) * (model_height / 2), NULL, &err));
+    transform_init(&transform, context, device_id);
+  }
+
+  void deinit_transform() {
+    transform_destroy(&transform);
+    CL_CHECK(clReleaseMemObject(v_cl));
+    CL_CHECK(clReleaseMemObject(u_cl));
+    CL_CHECK(clReleaseMemObject(y_cl));
+  }
+
+  void run_transform(cl_mem yuv_cl, int model_width, int model_height, int frame_width, int frame_height, int frame_stride, int frame_uv_offset, const mat3& projection) {
+    transform_queue(&transform, q,
+        yuv_cl, frame_width, frame_height, frame_stride, frame_uv_offset,
+        y_cl, u_cl, v_cl, model_width, model_height, projection);
+  }
+};
+
+class DrivingModelFrame : public ModelFrame {
+public:
+  DrivingModelFrame(cl_device_id device_id, cl_context context);
+  ~DrivingModelFrame();
+  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);
 
   const int MODEL_WIDTH = 512;
   const int MODEL_HEIGHT = 256;
   const int MODEL_FRAME_SIZE = MODEL_WIDTH * MODEL_HEIGHT * 3 / 2;
   const int buf_size = MODEL_FRAME_SIZE * 2;
-  const size_t frame_size_bytes = MODEL_FRAME_SIZE * sizeof(uint8_t);
+  const size_t frame_size_bytes = MODEL_FRAME_SIZE * sizeof(float);
 
 private:
-  Transform transform;
   LoadYUVState loadyuv;
-  cl_command_queue q;
-  cl_mem y_cl, u_cl, v_cl, img_buffer_20hz_cl, last_img_cl;
+  cl_mem img_buffer_20hz_cl, last_img_cl;//, input_frames_cl;
   cl_buffer_region region;
-  std::unique_ptr<uint8_t[]> input_frames;
-};
+};
+
+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);
+
+  const int MODEL_WIDTH = 1440;
+  const int MODEL_HEIGHT = 960;
+  const int MODEL_FRAME_SIZE = MODEL_WIDTH * MODEL_HEIGHT;
+  const int buf_size = MODEL_FRAME_SIZE;
+
+private:
+  // cl_mem input_frame_cl;
+};

selfdrive/modeld/models/commonmodel.pxd

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

selfdrive/modeld/models/commonmodel_pyx.pyx

@@ -4,11 +4,12 @@
 import numpy as np
 cimport numpy as cnp
 from libc.string cimport memcpy
+from libc.stdint cimport uintptr_t
 
 from msgq.visionipc.visionipc cimport cl_mem
 from msgq.visionipc.visionipc_pyx cimport VisionBuf, CLContext as BaseCLContext
 from .commonmodel cimport CL_DEVICE_TYPE_DEFAULT, cl_get_device_id, cl_create_context
-from .commonmodel cimport mat3, ModelFrame as cppModelFrame
+from .commonmodel cimport mat3, ModelFrame as cppModelFrame, DrivingModelFrame as cppDrivingModelFrame, MonitoringModelFrame as cppMonitoringModelFrame
 
 
 cdef class CLContext(BaseCLContext):
@@ -23,11 +24,17 @@ cdef class CLMem:
     mem.mem = <cl_mem*> cmem
     return mem
 
+  @property
+  def mem_address(self):
+    return <uintptr_t>(self.mem)
+
+def cl_from_visionbuf(VisionBuf buf):
+  return CLMem.create(<void*>&buf.buf.buf_cl)
+
+
 cdef class ModelFrame:
   cdef cppModelFrame * frame
-
-  def __cinit__(self, CLContext context):
-    self.frame = new cppModelFrame(context.device_id, context.context)
+  cdef int buf_size
 
   def __dealloc__(self):
     del self.frame
@@ -35,11 +42,35 @@ cdef class ModelFrame:
   def prepare(self, VisionBuf buf, float[:] projection, CLMem output):
     cdef mat3 cprojection
     memcpy(cprojection.v, &projection[0], 9*sizeof(float))
-    cdef unsigned char * data
+    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
-    return np.asarray(<cnp.uint8_t[:self.frame.buf_size]> 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)
+
+
+cdef class DrivingModelFrame(ModelFrame):
+  cdef cppDrivingModelFrame * _frame
+
+  def __cinit__(self, CLContext context):
+    self._frame = new cppDrivingModelFrame(context.device_id, context.context)
+    self.frame = <cppModelFrame*>(self._frame)
+    self.buf_size = self._frame.buf_size
+
+cdef class MonitoringModelFrame(ModelFrame):
+  cdef cppMonitoringModelFrame * _frame
+
+  def __cinit__(self, CLContext context):
+    self._frame = new cppMonitoringModelFrame(context.device_id, context.context)
+    self.frame = <cppModelFrame*>(self._frame)
+    self.buf_size = self._frame.buf_size

selfdrive/modeld/models/supercombo.onnx

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

selfdrive/modeld/parse_model_outputs.py

@@ -96,8 +96,6 @@ class Parser:
                    out_shape=(ModelConstants.LEAD_TRAJ_LEN,ModelConstants.LEAD_WIDTH))
     if 'lat_planner_solution' in outs:
       self.parse_mdn('lat_planner_solution', outs, in_N=0, out_N=0, out_shape=(ModelConstants.IDX_N,ModelConstants.LAT_PLANNER_SOLUTION_WIDTH))
-    if 'desired_curvature' in outs:
-      self.parse_mdn('desired_curvature', outs, in_N=0, out_N=0, out_shape=(ModelConstants.DESIRED_CURV_WIDTH,))
     for k in ['lead_prob', 'lane_lines_prob', 'meta']:
       self.parse_binary_crossentropy(k, outs)
     self.parse_categorical_crossentropy('desire_state', outs, out_shape=(ModelConstants.DESIRE_PRED_WIDTH,))

selfdrive/modeld/runners/onnxmodel.py

@@ -1,39 +1,12 @@
-import onnx
-import itertools
 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
 
-ORT_TYPES_TO_NP_TYPES = {'tensor(float16)': np.float16, 'tensor(float)': np.float32, 'tensor(uint8)': np.uint8}
-
-def attributeproto_fp16_to_fp32(attr):
-  float32_list = np.frombuffer(attr.raw_data, dtype=np.float16)
-  attr.data_type = 1
-  attr.raw_data = float32_list.astype(np.float32).tobytes()
-
-def convert_fp16_to_fp32(onnx_path_or_bytes):
-  if isinstance(onnx_path_or_bytes, bytes):
-    model = onnx.load_from_string(onnx_path_or_bytes)
-  elif isinstance(onnx_path_or_bytes, str):
-    model = onnx.load(onnx_path_or_bytes)
-
-  for i in model.graph.initializer:
-    if i.data_type == 10:
-      attributeproto_fp16_to_fp32(i)
-  for i in itertools.chain(model.graph.input, model.graph.output):
-    if i.type.tensor_type.elem_type == 10:
-      i.type.tensor_type.elem_type = 1
-  for i in model.graph.node:
-    if i.op_type == 'Cast' and i.attribute[0].i == 10:
-      i.attribute[0].i = 1
-    for a in i.attribute:
-      if hasattr(a, 't'):
-        if a.t.data_type == 10:
-          attributeproto_fp16_to_fp32(a.t)
-  return model.SerializeToString()
 
 def create_ort_session(path, fp16_to_fp32):
   os.environ["OMP_NUM_THREADS"] = "4"
@@ -49,14 +22,14 @@ def create_ort_session(path, fp16_to_fp32):
     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': 'DEFAULT'})
+    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(path) if fp16_to_fp32 else path
+  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)

selfdrive/modeld/runners/ort_helpers.py

@@ -0,0 +1,36 @@
+import onnx
+import onnxruntime as ort
+import numpy as np
+import itertools
+
+ORT_TYPES_TO_NP_TYPES = {'tensor(float16)': np.float16, 'tensor(float)': np.float32, 'tensor(uint8)': np.uint8}
+
+def attributeproto_fp16_to_fp32(attr):
+  float32_list = np.frombuffer(attr.raw_data, dtype=np.float16)
+  attr.data_type = 1
+  attr.raw_data = float32_list.astype(np.float32).tobytes()
+
+def convert_fp16_to_fp32(model):
+  for i in model.graph.initializer:
+    if i.data_type == 10:
+      attributeproto_fp16_to_fp32(i)
+  for i in itertools.chain(model.graph.input, model.graph.output):
+    if i.type.tensor_type.elem_type == 10:
+      i.type.tensor_type.elem_type = 1
+  for i in model.graph.node:
+    if i.op_type == 'Cast' and i.attribute[0].i == 10:
+      i.attribute[0].i = 1
+    for a in i.attribute:
+      if hasattr(a, 't'):
+        if a.t.data_type == 10:
+          attributeproto_fp16_to_fp32(a.t)
+  return model.SerializeToString()
+
+
+def make_onnx_cpu_runner(model_path):
+  options = ort.SessionOptions()
+  options.intra_op_num_threads = 4
+  options.execution_mode = ort.ExecutionMode.ORT_SEQUENTIAL
+  options.graph_optimization_level = ort.GraphOptimizationLevel.ORT_ENABLE_ALL
+  model_data = convert_fp16_to_fp32(onnx.load(model_path))
+  return ort.InferenceSession(model_data,  options, providers=['CPUExecutionProvider'])

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

selfdrive/pandad/panda.cc

@@ -145,6 +145,10 @@ void Panda::set_can_speed_kbps(uint16_t bus, uint16_t speed) {
   handle->control_write(0xde, bus, (speed * 10));
 }
 
+void Panda::set_can_fd_auto(uint16_t bus, bool enabled) {
+  handle->control_write(0xe8, bus, enabled);
+}
+
 void Panda::set_data_speed_kbps(uint16_t bus, uint16_t speed) {
   handle->control_write(0xf9, bus, (speed * 10));
 }

selfdrive/pandad/panda.h

@@ -77,6 +77,7 @@ public:
   void enable_deepsleep();
   void send_heartbeat(bool engaged);
   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);
   void set_canfd_non_iso(uint16_t bus, bool non_iso);
   void can_send(const capnp::List<cereal::CanData>::Reader &can_data_list);

selfdrive/pandad/panda_comms.cc

@@ -36,7 +36,7 @@ PandaUsbHandle::PandaUsbHandle(std::string serial) : PandaCommsHandle(serial) {
   for (size_t i = 0; i < num_devices; ++i) {
     libusb_device_descriptor desc;
     libusb_get_device_descriptor(dev_list[i], &desc);
-    if (desc.idVendor == 0xbbaa && desc.idProduct == 0xddcc) {
+    if (desc.idVendor == 0x3801 && desc.idProduct == 0xddcc) {
       int ret = libusb_open(dev_list[i], &dev_handle);
       if (dev_handle == NULL || ret < 0) { goto fail; }
 
@@ -110,7 +110,7 @@ std::vector<std::string> PandaUsbHandle::list() {
     libusb_device *device = dev_list[i];
     libusb_device_descriptor desc;
     libusb_get_device_descriptor(device, &desc);
-    if (desc.idVendor == 0xbbaa && desc.idProduct == 0xddcc) {
+    if (desc.idVendor == 0x3801 && desc.idProduct == 0xddcc) {
       libusb_device_handle *handle = NULL;
       int ret = libusb_open(device, &handle);
       if (ret < 0) { goto finish; }

selfdrive/pandad/pandad.cc

@@ -67,6 +67,10 @@ Panda *connect(std::string serial="", uint32_t index=0) {
   }
   //panda->enable_deepsleep();
 
+  for (int i = 0; i < PANDA_BUS_CNT; i++) {
+    panda->set_can_fd_auto(i, true);
+  }
+
   if (!panda->up_to_date() && !getenv("BOARDD_SKIP_FW_CHECK")) {
     throw std::runtime_error("Panda firmware out of date. Run pandad.py to update.");
   }
@@ -412,6 +416,7 @@ void process_peripheral_state(Panda *panda, PubMaster *pm, bool no_fan_control)
 
     if (ir_pwr != prev_ir_pwr || sm.frame % 100 == 0 || ir_pwr >= 50.0) {
       panda->set_ir_pwr(ir_pwr);
+      Hardware::set_ir_power(ir_pwr);
       prev_ir_pwr = ir_pwr;
     }
   }

selfdrive/pandad/spi.cc

@@ -239,7 +239,7 @@ int PandaSpiHandle::spi_transfer_retry(uint8_t endpoint, uint8_t *tx_data, uint1
         // due to full TX buffers
         nack_count += 1;
         if (nack_count > 3) {
-          SPILOG(LOGE, "NACK sleep %d", nack_count);
+          SPILOG(LOGD, "NACK sleep %d", nack_count);
           usleep(std::clamp(nack_count*10, 200, 2000));
         }
       }
@@ -418,7 +418,7 @@ fail:
     }
   }
 
-  if (ret > 0) ret = -1;
+  if (ret >= 0) ret = -1;
   return ret;
 }
 #endif

selfdrive/pandad/tests/test_pandad.py

@@ -6,6 +6,7 @@ import cereal.messaging as messaging
 from cereal import log
 from openpilot.common.gpio import gpio_set, gpio_init
 from panda import Panda, PandaDFU, PandaProtocolMismatch
+from openpilot.common.retry import retry
 from openpilot.system.manager.process_config import managed_processes
 from openpilot.system.hardware import HARDWARE
 from openpilot.system.hardware.tici.pins import GPIO
@@ -51,6 +52,7 @@ class TestPandad:
     assert not Panda.wait_for_dfu(None, 3)
     assert not Panda.wait_for_panda(None, 3)
 
+  @retry(attempts=3)
   def _flash_bootstub_and_test(self, fn, expect_mismatch=False):
     self._go_to_dfu()
     pd = PandaDFU(None)

selfdrive/pandad/tests/test_pandad_loopback.py

@@ -14,7 +14,7 @@ from openpilot.common.params import Params
 from openpilot.common.timeout import Timeout
 from openpilot.selfdrive.pandad import can_list_to_can_capnp
 from openpilot.system.hardware import TICI
-from openpilot.selfdrive.test.helpers import phone_only, with_processes
+from openpilot.selfdrive.test.helpers import with_processes
 
 
 @retry(attempts=3)
@@ -72,7 +72,6 @@ class TestBoarddLoopback:
     os.environ['STARTED'] = '1'
     os.environ['BOARDD_LOOPBACK'] = '1'
 
-  @phone_only
   @with_processes(['pandad'])
   def test_loopback(self):
     num_pandas = 2 if TICI and "SINGLE_PANDA" not in os.environ else 1

selfdrive/pandad/tests/test_pandad_spi.py

@@ -7,7 +7,7 @@ import random
 import cereal.messaging as messaging
 from cereal.services import SERVICE_LIST
 from openpilot.system.hardware import HARDWARE
-from openpilot.selfdrive.test.helpers import phone_only, with_processes
+from openpilot.selfdrive.test.helpers import with_processes
 from openpilot.selfdrive.pandad.tests.test_pandad_loopback import setup_pandad, send_random_can_messages
 
 JUNGLE_SPAM = "JUNGLE_SPAM" in os.environ
@@ -23,7 +23,6 @@ class TestBoarddSpi:
     if not JUNGLE_SPAM:
       os.environ['BOARDD_LOOPBACK'] = '1'
 
-  @phone_only
   @with_processes(['pandad'])
   def test_spi_corruption(self, subtests):
     setup_pandad(1)

selfdrive/selfdrived/tests/test_alerts.py

@@ -42,7 +42,7 @@ class TestAlerts:
 
     for name, e in events.items():
       if not name.endswith("DEPRECATED"):
-        fail_msg = "%s @%d not in EVENTS" % (name, e)
+        fail_msg = f"{name} @{e} not in EVENTS"
         assert e in EVENTS.keys(), fail_msg
 
   # ensure alert text doesn't exceed allowed width

selfdrive/test/helpers.py

@@ -10,7 +10,6 @@ from functools import wraps
 import cereal.messaging as messaging
 from openpilot.common.params import Params
 from openpilot.system.manager.process_config import managed_processes
-from openpilot.system.hardware import PC
 from openpilot.system.version import training_version, terms_version
 
 
@@ -29,14 +28,6 @@ def set_params_enabled():
   msg.liveCalibration.rpyCalib = [0.0, 0.0, 0.0]
   params.put("CalibrationParams", msg.to_bytes())
 
-def phone_only(f):
-  @wraps(f)
-  def wrap(self, *args, **kwargs):
-    if PC:
-      pytest.skip("This test is not meant to run on PC")
-    return f(self, *args, **kwargs)
-  return wrap
-
 def release_only(f):
   @wraps(f)
   def wrap(self, *args, **kwargs):

selfdrive/test/process_replay/model_replay.py

@@ -7,19 +7,20 @@ import tempfile
 from itertools import zip_longest
 
 import matplotlib.pyplot as plt
+import numpy as np
 
 from openpilot.common.git import get_commit
 from openpilot.system.hardware import PC
 from openpilot.tools.lib.openpilotci import get_url
 from openpilot.selfdrive.test.process_replay.compare_logs import compare_logs, format_diff
 from openpilot.selfdrive.test.process_replay.process_replay import get_process_config, replay_process
-from openpilot.tools.lib.framereader import FrameReader
+from openpilot.tools.lib.framereader import FrameReader, NumpyFrameReader
 from openpilot.tools.lib.logreader import LogReader, save_log
 from openpilot.tools.lib.github_utils import GithubUtils
 
 TEST_ROUTE = "2f4452b03ccb98f0|2022-12-03--13-45-30"
 SEGMENT = 6
-MAX_FRAMES = 100 if PC else 600
+MAX_FRAMES = 100 if PC else 400
 
 NO_MODEL = "NO_MODEL" in os.environ
 SEND_EXTRA_INPUTS = bool(int(os.getenv("SEND_EXTRA_INPUTS", "0")))
@@ -31,14 +32,14 @@ GITHUB = GithubUtils(API_TOKEN, DATA_TOKEN)
 
 
 def get_log_fn(test_route, ref="master"):
-  return f"{test_route}_model_tici_{ref}.bz2"
+  return f"{test_route}_model_tici_{ref}.zst"
 
 def plot(proposed, master, title, tmp):
   proposed = list(proposed)
   master = list(master)
   fig, ax = plt.subplots()
-  ax.plot(proposed, label='PROPOSED')
   ax.plot(master, label='MASTER')
+  ax.plot(proposed, label='PROPOSED')
   plt.legend(loc='best')
   plt.title(title)
   plt.savefig(f'{tmp}/{title}.png')
@@ -151,21 +152,44 @@ def model_replay(lr, frs):
   dmonitoringmodeld = get_process_config("dmonitoringmodeld")
 
   modeld_msgs = replay_process(modeld, modeld_logs, frs)
+  if isinstance(frs['roadCameraState'], NumpyFrameReader):
+    del frs['roadCameraState'].frames
+    del frs['wideRoadCameraState'].frames
   dmonitoringmodeld_msgs = replay_process(dmonitoringmodeld, dmodeld_logs, frs)
   return modeld_msgs + dmonitoringmodeld_msgs
 
 
+def get_frames():
+  regen_cache = "--regen-cache" in sys.argv
+  cache = "--cache" in sys.argv or not PC or regen_cache
+  videos = ('fcamera.hevc', 'dcamera.hevc', 'ecamera.hevc')
+  cams = ('roadCameraState', 'driverCameraState', 'wideRoadCameraState')
+
+  if cache:
+    frames_cache = '/tmp/model_replay_cache' if PC else '/data/model_replay_cache'
+    os.makedirs(frames_cache, exist_ok=True)
+
+    cache_size = 200
+    for v in videos:
+      if not all(os.path.isfile(f'{frames_cache}/{TEST_ROUTE}_{v}_{i}.npy') for i in range(MAX_FRAMES//cache_size)) or regen_cache:
+        f = FrameReader(get_url(TEST_ROUTE, SEGMENT, v)).get(0, MAX_FRAMES + 1, pix_fmt="nv12")
+        print(f'Caching {v}...')
+        for i in range(MAX_FRAMES//cache_size):
+          np.save(f'{frames_cache}/{TEST_ROUTE}_{v}_{i}', f[(i * cache_size) + 1:((i + 1) * cache_size) + 1])
+        del f
+
+    return {c : NumpyFrameReader(f"{frames_cache}/{TEST_ROUTE}_{v}", 1928, 1208, cache_size) for c,v in zip(cams, videos, strict=True)}
+  else:
+    return {c : FrameReader(get_url(TEST_ROUTE, SEGMENT, v), readahead=True) for c,v in zip(cams, videos, strict=True)}
+
+
 if __name__ == "__main__":
   update = "--update" in sys.argv or (os.getenv("GIT_BRANCH", "") == 'master')
   replay_dir = os.path.dirname(os.path.abspath(__file__))
 
   # load logs
-  lr = list(LogReader(get_url(TEST_ROUTE, SEGMENT, "rlog.bz2")))
-  frs = {
-    'roadCameraState': FrameReader(get_url(TEST_ROUTE, SEGMENT, "fcamera.hevc"), readahead=True),
-    'driverCameraState': FrameReader(get_url(TEST_ROUTE, SEGMENT, "dcamera.hevc"), readahead=True),
-    'wideRoadCameraState': FrameReader(get_url(TEST_ROUTE, SEGMENT, "ecamera.hevc"), readahead=True)
-  }
+  lr = list(LogReader(get_url(TEST_ROUTE, SEGMENT, "rlog.zst")))
+  frs = get_frames()
 
   log_msgs = []
   # run replays

selfdrive/test/process_replay/process_replay.py

@@ -5,13 +5,13 @@ import copy
 import json
 import heapq
 import signal
-import platform
 from collections import Counter, OrderedDict
 from dataclasses import dataclass, field
 from typing import Any
 from collections.abc import Callable, Iterable
 from tqdm import tqdm
 import capnp
+from openpilot.system.hardware.hw import Paths
 
 import cereal.messaging as messaging
 from cereal import car
@@ -780,8 +780,7 @@ def generate_params_config(lr=None, CP=None, fingerprint=None, custom_params=Non
 
 def generate_environ_config(CP=None, fingerprint=None, log_dir=None) -> dict[str, Any]:
   environ_dict = {}
-  if platform.system() != "Darwin":
-    environ_dict["PARAMS_ROOT"] = "/dev/shm/params"
+  environ_dict["PARAMS_ROOT"] = f"{Paths.shm_path()}/params"
   if log_dir is not None:
     environ_dict["LOG_ROOT"] = log_dir
 

selfdrive/test/process_replay/ref_commit

@@ -1 +1 @@
-2fc2e865ab77fd8145feab86d454f2111c5d9871
+cae12bc0a2960de17104a9e22fafe33d797fbcee

selfdrive/test/test_onroad.py

@@ -72,7 +72,7 @@ PROCS = {
 
 PROCS.update({
   "tici": {
-    "./pandad": 4.0,
+    "./pandad": 5.0,
     "./ubloxd": 1.0,
     "system.ubloxd.pigeond": 6.0,
   },
@@ -138,6 +138,7 @@ class TestOnroad:
     proc = None
     try:
       manager_path = os.path.join(BASEDIR, "system/manager/manager.py")
+      cls.manager_st = time.monotonic()
       proc = subprocess.Popen(["python", manager_path])
 
       sm = messaging.SubMaster(['carState'])
@@ -202,6 +203,10 @@ class TestOnroad:
       with subtests.test(service=s):
         assert len(msgs) >= math.floor(SERVICE_LIST[s].frequency*int(TEST_DURATION*0.8))
 
+  def test_manager_starting_time(self):
+    st = self.msgs['managerState'][0].logMonoTime / 1e9
+    assert (st - self.manager_st) < 10, f"manager.py took {st - self.manager_st}s to publish the first 'managerState' msg"
+
   def test_cloudlog_size(self):
     msgs = self.msgs['logMessage']
 
@@ -295,13 +300,18 @@ class TestOnroad:
     assert cpu_ok
 
   def test_memory_usage(self):
+    print("\n------------------------------------------------")
+    print("--------------- Memory Usage -------------------")
+    print("------------------------------------------------")
     offset = int(SERVICE_LIST['deviceState'].frequency * LOG_OFFSET)
     mems = [m.deviceState.memoryUsagePercent for m in self.msgs['deviceState'][offset:]]
     print("Memory usage: ", mems)
 
     # check for big leaks. note that memory usage is
     # expected to go up while the MSGQ buffers fill up
-    assert max(mems) - min(mems) <= 3.0
+    assert np.average(mems) <= 65, "Average memory usage above 65%"
+    assert np.max(np.diff(mems)) <= 4, "Max memory increase too high"
+    assert np.average(np.diff(mems)) <= 1, "Average memory increase too high"
 
   def test_gpu_usage(self):
     assert self.gpu_procs == {"weston", "ui", "camerad", "selfdrive.modeld.modeld", "selfdrive.modeld.dmonitoringmodeld"}
@@ -368,6 +378,8 @@ class TestOnroad:
     ]
     for (s, instant_max, avg_max) in cfgs:
       ts = [getattr(m, s).modelExecutionTime for m in self.msgs[s]]
+      # TODO some init can happen in first iteration
+      ts = ts[1:]
       assert max(ts) < instant_max, f"high '{s}' execution time: {max(ts)}"
       assert np.mean(ts) < avg_max, f"high avg '{s}' execution time: {np.mean(ts)}"
       result += f"'{s}' execution time: min  {min(ts):.5f}s\n"

selfdrive/ui/SConscript

@@ -67,14 +67,13 @@ if GetOption('extras'):
   qt_src.remove("main.cc")  # replaced by test_runner
   qt_env.Program('tests/test_translations', [asset_obj, 'tests/test_runner.cc', 'tests/test_translations.cc'] + qt_src, LIBS=qt_libs)
 
-if GetOption('extras') and arch != "Darwin":
+if GetOption('extras'):
   qt_env.SharedLibrary("qt/python_helpers", ["qt/qt_window.cc"], LIBS=qt_libs)
 
   # spinner and text window
   qt_env.Program("_text", ["qt/text.cc"], LIBS=qt_libs)
   qt_env.Program("_spinner", ["qt/spinner.cc"], LIBS=qt_libs)
 
-
   # setup and factory resetter
   qt_env.Program("qt/setup/reset", ["qt/setup/reset.cc"], LIBS=qt_libs)
   qt_env.Program("qt/setup/setup", ["qt/setup/setup.cc", asset_obj],
@@ -83,32 +82,31 @@ if GetOption('extras') and arch != "Darwin":
   # build updater UI
   qt_env.Program("qt/setup/updater", ["qt/setup/updater.cc", asset_obj], LIBS=qt_libs)
 
-  # build installers
-  senv = qt_env.Clone()
-  senv['LINKFLAGS'].append('-Wl,-strip-debug')
+  if arch != "Darwin":
+    # build installers
+    senv = qt_env.Clone()
+    senv['LINKFLAGS'].append('-Wl,-strip-debug')
 
-  release = "release3"
-  installers = [
-    ("openpilot", release),
-    ("openpilot_test", f"{release}-staging"),
-    ("openpilot_nightly", "nightly"),
-    ("openpilot_internal", "nightly-dev"),
-  ]
+    release = "release3"
+    installers = [
+      ("openpilot", release),
+      ("openpilot_test", f"{release}-staging"),
+      ("openpilot_nightly", "nightly"),
+      ("openpilot_internal", "nightly-dev"),
+    ]
 
-  cont = senv.Command(f"installer/continue_openpilot.o", f"installer/continue_openpilot.sh",
-                      "ld -r -b binary -o $TARGET $SOURCE")
-  for name, branch in installers:
-    d = {'BRANCH': f"'\"{branch}\"'"}
-    if "internal" in name:
-      d['INTERNAL'] = "1"
+    cont = senv.Command(f"installer/continue_openpilot.o", f"installer/continue_openpilot.sh",
+                        "ld -r -b binary -o $TARGET $SOURCE")
+    for name, branch in installers:
+      d = {'BRANCH': f"'\"{branch}\"'"}
+      if "internal" in name:
+        d['INTERNAL'] = "1"
 
-    obj = senv.Object(f"installer/installers/installer_{name}.o", ["installer/installer.cc"], CPPDEFINES=d)
-    f = senv.Program(f"installer/installers/installer_{name}", [obj, cont], LIBS=qt_libs)
-    # keep installers small
-    assert f[0].get_size() < 370*1e3
+      obj = senv.Object(f"installer/installers/installer_{name}.o", ["installer/installer.cc"], CPPDEFINES=d)
+      f = senv.Program(f"installer/installers/installer_{name}", [obj, cont], LIBS=qt_libs)
+      # keep installers small
+      assert f[0].get_size() < 370*1e3
 
 # build watch3
 if arch in ['x86_64', 'aarch64', 'Darwin'] or GetOption('extras'):
   qt_env.Program("watch3", ["watch3.cc"], LIBS=qt_libs + ['common', 'msgq', 'visionipc'])
-
-SConscript(['raylib/SConscript'])

selfdrive/ui/qt/home.cc

@@ -124,6 +124,7 @@ OffroadHome::OffroadHome(QWidget* parent) : QFrame(parent) {
     // left: PrimeAdWidget
     QStackedWidget *left_widget = new QStackedWidget(this);
     QVBoxLayout *left_prime_layout = new QVBoxLayout();
+    left_prime_layout->setContentsMargins(0, 0, 0, 0);
     QWidget *prime_user = new PrimeUserWidget();
     prime_user->setStyleSheet(R"(
     border-radius: 10px;

selfdrive/ui/qt/network/wifi_manager.cc

@@ -203,7 +203,7 @@ void WifiManager::connect(const Network &n, const bool is_hidden, const QString
   connection["ipv4"]["dns-priority"] = 600;
   connection["ipv6"]["method"] = "ignore";
 
-  call(NM_DBUS_PATH_SETTINGS, NM_DBUS_INTERFACE_SETTINGS, "AddConnection", QVariant::fromValue(connection));
+  asyncCall(NM_DBUS_PATH_SETTINGS, NM_DBUS_INTERFACE_SETTINGS, "AddConnection", QVariant::fromValue(connection));
 }
 
 void WifiManager::deactivateConnectionBySsid(const QString &ssid) {
@@ -330,6 +330,10 @@ void WifiManager::initConnections() {
       lteConnectionPath = path;
     }
   }
+
+  if (!isKnownConnection(tethering_ssid)) {
+    addTetheringConnection();
+  }
 }
 
 std::optional<QDBusPendingCall> WifiManager::activateWifiConnection(const QString &ssid) {
@@ -399,9 +403,13 @@ void WifiManager::updateGsmSettings(bool roaming, QString apn, bool metered) {
     }
 
     if (changes) {
-      call(lteConnectionPath.path(), NM_DBUS_INTERFACE_SETTINGS_CONNECTION, "UpdateUnsaved", QVariant::fromValue(settings));  // update is temporary
-      deactivateConnection(lteConnectionPath);
-      activateModemConnection(lteConnectionPath);
+      QDBusPendingCall pending_call = asyncCall(lteConnectionPath.path(), NM_DBUS_INTERFACE_SETTINGS_CONNECTION, "UpdateUnsaved", QVariant::fromValue(settings));  // update is temporary
+      QDBusPendingCallWatcher *watcher = new QDBusPendingCallWatcher(pending_call);
+      QObject::connect(watcher, &QDBusPendingCallWatcher::finished, this, [this, watcher]() {
+        deactivateConnection(lteConnectionPath);
+        activateModemConnection(lteConnectionPath);
+        watcher->deleteLater();
+      });
     }
   }
 }
@@ -434,10 +442,7 @@ void WifiManager::addTetheringConnection() {
   connection["ipv4"]["route-metric"] = 1100;
   connection["ipv6"]["method"] = "ignore";
 
-  auto path = call<QDBusObjectPath>(NM_DBUS_PATH_SETTINGS, NM_DBUS_INTERFACE_SETTINGS, "AddConnection", QVariant::fromValue(connection));
-  if (!path.path().isEmpty()) {
-    knownConnections[path] = tethering_ssid;
-  }
+  asyncCall(NM_DBUS_PATH_SETTINGS, NM_DBUS_INTERFACE_SETTINGS, "AddConnection", QVariant::fromValue(connection));
 }
 
 void WifiManager::tetheringActivated(QDBusPendingCallWatcher *call) {
@@ -453,10 +458,6 @@ void WifiManager::tetheringActivated(QDBusPendingCallWatcher *call) {
 
 void WifiManager::setTetheringEnabled(bool enabled) {
   if (enabled) {
-    if (!isKnownConnection(tethering_ssid)) {
-      addTetheringConnection();
-    }
-
     auto pending_call = activateWifiConnection(tethering_ssid);
 
     if (pending_call) {
@@ -478,9 +479,6 @@ bool WifiManager::isTetheringEnabled() {
 }
 
 QString WifiManager::getTetheringPassword() {
-  if (!isKnownConnection(tethering_ssid)) {
-    addTetheringConnection();
-  }
   const QDBusObjectPath &path = getConnectionPath(tethering_ssid);
   if (!path.path().isEmpty()) {
     QDBusReply<QMap<QString, QVariantMap>> response = call(path.path(), NM_DBUS_INTERFACE_SETTINGS_CONNECTION, "GetSecrets", "802-11-wireless-security");

selfdrive/ui/qt/setup/setup.cc

@@ -408,7 +408,7 @@ Setup::Setup(QWidget *parent) : QStackedWidget(parent) {
   std::stringstream buffer;
   buffer << std::ifstream("/sys/class/hwmon/hwmon1/in1_input").rdbuf();
   float voltage = (float)std::atoi(buffer.str().c_str()) / 1000.;
-  if (voltage > 0 && voltage < 7) {
+  if (voltage < 7) {
     addWidget(low_voltage());
   }
 

selfdrive/ui/raylib/.gitignore

@@ -1 +0,0 @@
-_spinner

selfdrive/ui/raylib/SConscript

@@ -1,17 +0,0 @@
-Import('env', 'arch', 'common')
-
-raylib_env = env.Clone()
-raylib_util_lib = env.Library("raylib_util_lib", ['util.cc'], LIBS='raylib')
-linked_libs = ['raylib', raylib_util_lib, common]
-raylib_env['LIBPATH'] += [f'#third_party/raylib/{arch}/']
-
-mac_frameworks = []
-if arch == "Darwin":
-  mac_frameworks += ['OpenCL', 'CoreVideo', 'Cocoa', 'GLUT', 'CoreFoundation', 'OpenGL', 'IOKit']
-elif arch == 'larch64':
-  linked_libs += []
-else:
-  linked_libs += ['OpenCL', 'dl', 'pthread']
-
-if arch != 'aarch64':
-  raylib_env.Program("_spinner", ["spinner.cc"], LIBS=linked_libs, FRAMEWORKS=mac_frameworks)

system/athena/athenad.py

@@ -54,6 +54,7 @@ RETRY_DELAY = 10  # seconds
 MAX_RETRY_COUNT = 30  # Try for at most 5 minutes if upload fails immediately
 MAX_AGE = 31 * 24 * 3600  # seconds
 WS_FRAME_SIZE = 4096
+DEVICE_STATE_UPDATE_INTERVAL = 1.0  # in seconds
 
 NetworkType = log.DeviceState.NetworkType
 
@@ -211,10 +212,11 @@ def retry_upload(tid: int, end_event: threading.Event, increase_count: bool = Tr
 def cb(sm, item, tid, end_event: threading.Event, sz: int, cur: int) -> None:
   # Abort transfer if connection changed to metered after starting upload
   # or if athenad is shutting down to re-connect the websocket
-  sm.update(0)
-  metered = sm['deviceState'].networkMetered
-  if metered and (not item.allow_cellular):
-    raise AbortTransferException
+  if not item.allow_cellular:
+    if (time.monotonic() - sm.recv_time['deviceState']) > DEVICE_STATE_UPDATE_INTERVAL:
+      sm.update(0)
+      if sm['deviceState'].networkMetered:
+        raise AbortTransferException
 
   if end_event.is_set():
     raise AbortTransferException
@@ -258,13 +260,13 @@ def upload_handler(end_event: threading.Event) -> None:
           sz = -1
 
         cloudlog.event("athena.upload_handler.upload_start", fn=fn, sz=sz, network_type=network_type, metered=metered, retry_count=item.retry_count)
-        response = _do_upload(item, partial(cb, sm, item, tid, end_event))
 
-        if response.status_code not in (200, 201, 401, 403, 412):
-          cloudlog.event("athena.upload_handler.retry", status_code=response.status_code, fn=fn, sz=sz, network_type=network_type, metered=metered)
-          retry_upload(tid, end_event)
-        else:
-          cloudlog.event("athena.upload_handler.success", fn=fn, sz=sz, network_type=network_type, metered=metered)
+        with _do_upload(item, partial(cb, sm, item, tid, end_event)) as response:
+          if response.status_code not in (200, 201, 401, 403, 412):
+            cloudlog.event("athena.upload_handler.retry", status_code=response.status_code, fn=fn, sz=sz, network_type=network_type, metered=metered)
+            retry_upload(tid, end_event)
+          else:
+            cloudlog.event("athena.upload_handler.success", fn=fn, sz=sz, network_type=network_type, metered=metered)
 
         UploadQueueCache.cache(upload_queue)
       except (requests.exceptions.Timeout, requests.exceptions.ConnectionError, requests.exceptions.SSLError):
@@ -309,13 +311,16 @@ def getMessage(service: str, timeout: int = 1000) -> dict:
     raise Exception("invalid service")
 
   socket = messaging.sub_sock(service, timeout=timeout)
-  ret = messaging.recv_one(socket)
+  try:
+    ret = messaging.recv_one(socket)
 
-  if ret is None:
-    raise TimeoutError
+    if ret is None:
+      raise TimeoutError
 
-  # this is because capnp._DynamicStructReader doesn't have typing information
-  return cast(dict, ret.to_dict())
+    # this is because capnp._DynamicStructReader doesn't have typing information
+    return cast(dict, ret.to_dict())
+  finally:
+    del socket
 
 
 @dispatcher.add_method
@@ -626,8 +631,9 @@ def log_handler(end_event: threading.Event) -> None:
 
 def stat_handler(end_event: threading.Event) -> None:
   STATS_DIR = Paths.stats_root()
+  last_scan = 0.0
+
   while not end_event.is_set():
-    last_scan = 0.
     curr_scan = time.monotonic()
     try:
       if curr_scan - last_scan > 10:
@@ -801,6 +807,8 @@ def main(exit_event: threading.Event = None):
       cur_upload_items.clear()
 
       handle_long_poll(ws, exit_event)
+
+      ws.close()
     except (KeyboardInterrupt, SystemExit):
       break
     except (ConnectionError, TimeoutError, WebSocketException):

system/athena/tests/test_athenad.py

@@ -240,7 +240,7 @@ class TestAthenadMethods:
   @with_upload_handler
   def test_upload_handler_retry(self, mocker, host, status, retry):
     mock_put = mocker.patch('requests.put')
-    mock_put.return_value.status_code = status
+    mock_put.return_value.__enter__.return_value.status_code = status
     fn = self._create_file('qlog.zst')
     item = athenad.UploadItem(path=fn, url=f"{host}/qlog.zst", headers={}, created_at=int(time.time()*1000), id='', allow_cellular=True)
 

system/camerad/cameras/camera_common.cc

@@ -65,16 +65,19 @@ void CameraBuf::init(cl_device_id device_id, cl_context context, SpectraCamera *
   const SensorInfo *sensor = cam->sensor.get();
 
   is_raw = cam->is_raw;
-  camera_bufs_raw = std::make_unique<VisionBuf[]>(frame_buf_count);
   frame_metadata = std::make_unique<FrameMetadata[]>(frame_buf_count);
 
-  // RAW + final frames from ISP
-  const int raw_frame_size = (sensor->frame_height + sensor->extra_height) * sensor->frame_stride;
-  for (int i = 0; i < frame_buf_count; i++) {
-    camera_bufs_raw[i].allocate(raw_frame_size);
-    camera_bufs_raw[i].init_cl(device_id, context);
+  // RAW frames from ISP
+  if (is_raw) {
+    camera_bufs_raw = std::make_unique<VisionBuf[]>(frame_buf_count);
+
+    const int raw_frame_size = (sensor->frame_height + sensor->extra_height) * sensor->frame_stride;
+    for (int i = 0; i < frame_buf_count; i++) {
+      camera_bufs_raw[i].allocate(raw_frame_size);
+      camera_bufs_raw[i].init_cl(device_id, context);
+    }
+    LOGD("allocated %d CL buffers", frame_buf_count);
   }
-  LOGD("allocated %d CL buffers", frame_buf_count);
 
   out_img_width = sensor->frame_width;
   out_img_height = sensor->hdr_offset > 0 ? (sensor->frame_height - sensor->hdr_offset) / 2 : sensor->frame_height;
@@ -83,15 +86,17 @@ void CameraBuf::init(cl_device_id device_id, cl_context context, SpectraCamera *
   // TODO: VENUS_BUFFER_SIZE should give the size, but it's too small. dependent on encoder settings?
   size_t nv12_size = (out_img_width <= 1344 ? 2900 : 2346)*cam->stride;
 
-  vipc_server->create_buffers_with_sizes(stream_type, YUV_BUFFER_COUNT, out_img_width, out_img_height, nv12_size, cam->stride, cam->uv_offset);
-  LOGD("created %d YUV vipc buffers with size %dx%d", YUV_BUFFER_COUNT, cam->stride, cam->y_height);
+  vipc_server->create_buffers_with_sizes(stream_type, VIPC_BUFFER_COUNT, out_img_width, out_img_height, nv12_size, cam->stride, cam->uv_offset);
+  LOGD("created %d YUV vipc buffers with size %dx%d", VIPC_BUFFER_COUNT, cam->stride, cam->y_height);
 
-  imgproc = new ImgProc(device_id, context, this, sensor, cam->cc.camera_num, cam->stride, cam->uv_offset);
+  if (is_raw) imgproc = new ImgProc(device_id, context, this, sensor, cam->cc.camera_num, cam->stride, cam->uv_offset);
 }
 
 CameraBuf::~CameraBuf() {
-  for (int i = 0; i < frame_buf_count; i++) {
-    camera_bufs_raw[i].free();
+  if (camera_bufs_raw != nullptr) {
+    for (int i = 0; i < frame_buf_count; i++) {
+      camera_bufs_raw[i].free();
+    }
   }
   if (imgproc) delete imgproc;
 }

system/camerad/cameras/camera_common.h

@@ -8,7 +8,7 @@
 #include "common/util.h"
 
 
-const int YUV_BUFFER_COUNT = 20;
+const int VIPC_BUFFER_COUNT = 18;
 
 typedef struct FrameMetadata {
   uint32_t frame_id;

system/camerad/cameras/camera_qcom2.cc

@@ -55,7 +55,7 @@ public:
 
   float fl_pix = 0;
 
-  CameraState(SpectraMaster *master, const CameraConfig &config) : camera(master, config, true /*config.stream_type == VISION_STREAM_ROAD*/) {};
+  CameraState(SpectraMaster *master, const CameraConfig &config) : camera(master, config, config.stream_type == VISION_STREAM_ROAD) {};
   ~CameraState();
   void init(VisionIpcServer *v, cl_device_id device_id, cl_context ctx);
   void update_exposure_score(float desired_ev, int exp_t, int exp_g_idx, float exp_gain);
@@ -69,10 +69,10 @@ public:
 };
 
 void CameraState::init(VisionIpcServer *v, cl_device_id device_id, cl_context ctx) {
-  if (!camera.enabled) return;
-
   camera.camera_open(v, device_id, ctx);
 
+  if (!camera.enabled) return;
+
   fl_pix = camera.cc.focal_len / camera.sensor->pixel_size_mm;
   set_exposure_rect();
 
@@ -140,14 +140,14 @@ void CameraState::set_camera_exposure(float grey_frac) {
   // Therefore we use the target EV from 3 frames ago, the grey fraction that was just measured was the result of that control action.
   // TODO: Lower latency to 2 frames, by using the histogram outputted by the sensor we can do AE before the debayering is complete
 
-  const float cur_ev_ = cur_ev[camera.buf.cur_frame_data.frame_id % 3];
   const auto &sensor = camera.sensor;
+  const float cur_ev_ = cur_ev[camera.buf.cur_frame_data.frame_id % 3] * sensor->ev_scale;
 
   // Scale target grey between 0.1 and 0.4 depending on lighting conditions
   float new_target_grey = std::clamp(0.4 - 0.3 * log2(1.0 + sensor->target_grey_factor*cur_ev_) / log2(6000.0), 0.1, 0.4);
   float target_grey = (1.0 - k_grey) * target_grey_fraction + k_grey * new_target_grey;
 
-  float desired_ev = std::clamp(cur_ev_ * target_grey / grey_frac, sensor->min_ev, sensor->max_ev);
+  float desired_ev = std::clamp(cur_ev_ / sensor->ev_scale * target_grey / grey_frac, sensor->min_ev, sensor->max_ev);
   float k = (1.0 - k_ev) / 3.0;
   desired_ev = (k * cur_ev[0]) + (k * cur_ev[1]) + (k * cur_ev[2]) + (k_ev * desired_ev);
 

system/camerad/cameras/cdm.cc

@@ -14,7 +14,7 @@ int write_dmi(uint8_t *dst, uint64_t *addr, uint32_t length, uint32_t dmi_addr,
   return sizeof(struct cdm_dmi_cmd);
 }
 
-int write_cont(uint8_t *dst, uint32_t reg, std::vector<uint32_t> vals) {
+int write_cont(uint8_t *dst, uint32_t reg, const std::vector<uint32_t> &vals) {
   struct cdm_regcontinuous_cmd *cmd = (struct cdm_regcontinuous_cmd*)dst;
   cmd->cmd = CAM_CDM_CMD_REG_CONT;
   cmd->count = vals.size();
@@ -31,7 +31,7 @@ int write_cont(uint8_t *dst, uint32_t reg, std::vector<uint32_t> vals) {
   return sizeof(struct cdm_regcontinuous_cmd) + vals.size()*sizeof(uint32_t);
 }
 
-int write_random(uint8_t *dst, std::vector<uint32_t> vals) {
+int write_random(uint8_t *dst, const std::vector<uint32_t> &vals) {
   struct cdm_regrandom_cmd *cmd = (struct cdm_regrandom_cmd*)dst;
   cmd->cmd = CAM_CDM_CMD_REG_RANDOM;
   cmd->count = vals.size() / 2;

system/camerad/cameras/cdm.h

@@ -7,8 +7,8 @@
 #include <memory>
 
 // our helpers
-int write_random(uint8_t *dst, std::vector<uint32_t> vals);
-int write_cont(uint8_t *dst, uint32_t reg, std::vector<uint32_t> vals);
+int write_random(uint8_t *dst, const std::vector<uint32_t> &vals);
+int write_cont(uint8_t *dst, uint32_t reg, const std::vector<uint32_t> &vals);
 int write_dmi(uint8_t *dst, uint64_t *addr, uint32_t length, uint32_t dmi_addr, uint8_t sel);
 
 // from drivers/media/platform/msm/camera/cam_cdm/cam_cdm_util.{c,h}

system/camerad/cameras/process_raw.cl

@@ -19,6 +19,17 @@
 #endif
 
 float get_vignetting_s(float r) {
+#if defined(VIGNETTE_PROFILE_4DT6MM)
+  if (r < 100000) {
+    return 1.0f + 0.0000013f*r;
+  } else if (r < 250000) {
+    return 1.02f + 0.0000011f*r;
+  } else if (r < 400000) {
+    return 0.92f + 0.0000015f*r;
+  } else {
+    return 0.44f + 0.0000027f*r;
+  }
+#elif defined(VIGNETTE_PROFILE_8DT0MM)
   if (r < 62500) {
     return (1.0f + 0.0000008f*r);
   } else if (r < 490000) {
@@ -28,6 +39,9 @@ float get_vignetting_s(float r) {
   } else {
     return (0.53503625f + 0.0000000000022f*r*r);
   }
+#else
+  return 1.0f;
+#endif
 }
 
 int4 parse_12bit(uchar8 pvs) {
@@ -65,7 +79,7 @@ __kernel void process_raw(const __global uchar * in, __global uchar * out, int e
   #if VIGNETTING
     int gx = (gid_x*2 - RGB_WIDTH/2);
     int gy = (gid_y*2 - RGB_HEIGHT/2);
-    const float vignette_factor = get_vignetting_s((gx*gx + gy*gy) / VIGNETTE_RSZ);
+    const float vignette_factor = get_vignetting_s(gx*gx + gy*gy);
   #else
     const float vignette_factor = 1.0;
   #endif

system/camerad/cameras/spectra.cc

@@ -51,12 +51,12 @@ int do_sync_control(int fd, uint32_t id, void *handle, uint32_t size) {
   };
   int ret = HANDLE_EINTR(ioctl(fd, CAM_PRIVATE_IOCTL_CMD, &arg));
 
-  int32_t ioctl_result = (int32_t)arg.result;
+  int32_t ioctl_result = static_cast<int32_t>(arg.result);
   if (ret < 0) {
     LOGE("CAM_SYNC error: id %u - errno %d - ret %d - ioctl_result %d", id, errno, ret, ioctl_result);
     return ret;
   }
-  if (ioctl_result < 0) {
+  if (ioctl_result != 0) {
     LOGE("CAM_SYNC error: id %u - errno %d - ret %d - ioctl_result %d", id, errno, ret, ioctl_result);
     return ioctl_result;
   }
@@ -196,9 +196,9 @@ void SpectraMaster::init() {
   assert(isp_fd >= 0);
   LOGD("opened isp");
 
-  //icp_fd = open_v4l_by_name_and_index("cam-icp");
-  //assert(icp_fd >= 0);
-  //LOGD("opened icp");
+  icp_fd = open_v4l_by_name_and_index("cam-icp");
+  assert(icp_fd >= 0);
+  LOGD("opened icp");
 
   // query ISP for MMU handles
   LOG("-- Query for MMU handles");
@@ -215,7 +215,6 @@ void SpectraMaster::init() {
   cdm_iommu = isp_query_cap_cmd.cdm_iommu.non_secure;
 
   // query ICP for MMU handles
-  /*
   struct cam_icp_query_cap_cmd icp_query_cap_cmd = {0};
   query_cap_cmd.caps_handle = (uint64_t)&icp_query_cap_cmd;
   query_cap_cmd.size = sizeof(icp_query_cap_cmd);
@@ -223,7 +222,6 @@ void SpectraMaster::init() {
   assert(ret == 0);
   LOGD("using ICP MMU handle: %x", icp_query_cap_cmd.dev_iommu_handle.non_secure);
   icp_device_iommu = icp_query_cap_cmd.dev_iommu_handle.non_secure;
-  */
 
   // subscribe
   LOG("-- Subscribing");
@@ -242,6 +240,9 @@ SpectraCamera::SpectraCamera(SpectraMaster *master, const CameraConfig &config,
     cc(config),
     is_raw(raw) {
   mm.init(m->video0_fd);
+
+  ife_buf_depth = is_raw ? 4 : VIPC_BUFFER_COUNT;
+  assert(ife_buf_depth < MAX_IFE_BUFS);
 }
 
 SpectraCamera::~SpectraCamera() {
@@ -261,12 +262,12 @@ int SpectraCamera::clear_req_queue() {
 }
 
 void SpectraCamera::camera_open(VisionIpcServer *v, cl_device_id device_id, cl_context ctx) {
-  if (!enabled) return;
-
   if (!openSensor()) {
     return;
   }
 
+  if (!enabled) return;
+
   // size is driven by all the HW that handles frames,
   // the video encoder has certain alignment requirements in this case
   stride = VENUS_Y_STRIDE(COLOR_FMT_NV12, sensor->frame_width);
@@ -288,14 +289,15 @@ void SpectraCamera::camera_open(VisionIpcServer *v, cl_device_id device_id, cl_c
   linkDevices();
 
   LOGD("camera init %d", cc.camera_num);
-  buf.init(device_id, ctx, this, v, FRAME_BUF_COUNT, cc.stream_type);
+  buf.init(device_id, ctx, this, v, ife_buf_depth, cc.stream_type);
   camera_map_bufs();
+  enqueue_req_multi(1, ife_buf_depth, 0);
 }
 
 void SpectraCamera::enqueue_req_multi(uint64_t start, int n, bool dp) {
   for (uint64_t i = start; i < start + n; ++i) {
-    request_ids[(i - 1) % FRAME_BUF_COUNT] = i;
-    enqueue_buffer((i - 1) % FRAME_BUF_COUNT, dp);
+    request_ids[(i - 1) % ife_buf_depth] = i;
+    enqueue_buffer((i - 1) % ife_buf_depth, dp);
   }
 }
 
@@ -670,19 +672,22 @@ void SpectraCamera::enqueue_buffer(int i, bool dp) {
   int ret;
   uint64_t request_id = request_ids[i];
 
-  if (buf_handle_raw[i] && sync_objs[i]) {
-    // wait
+  if (sync_objs[i]) {
+    // SOF has come in, wait until readout is complete
     struct cam_sync_wait sync_wait = {0};
     sync_wait.sync_obj = sync_objs[i];
-    sync_wait.timeout_ms = 50; // max dt tolerance, typical should be 23
+    sync_wait.timeout_ms = 100;
     ret = do_sync_control(m->cam_sync_fd, CAM_SYNC_WAIT, &sync_wait, sizeof(sync_wait));
     if (ret != 0) {
-      LOGE("failed to wait for sync: %d %d", ret, sync_wait.sync_obj);
       // TODO: handle frame drop cleanly
+      // when this happens, it messes up future frames
+      LOGE("failed to wait for sync: %d %d", ret, sync_wait.sync_obj);
     }
     buf.frame_metadata[i].timestamp_end_of_isp = (uint64_t)nanos_since_boot();
     buf.frame_metadata[i].timestamp_eof = buf.frame_metadata[i].timestamp_sof + sensor->readout_time_ns;
-    if (dp) buf.queue(i);
+    if (dp) {
+      buf.queue(i);
+    }
 
     // destroy old output fence
     for (auto so : {sync_objs, sync_objs_bps_out}) {
@@ -705,13 +710,13 @@ void SpectraCamera::enqueue_buffer(int i, bool dp) {
   }
   sync_objs[i] = sync_create.sync_obj;
 
-  /*
-  ret = do_cam_control(m->cam_sync_fd, CAM_SYNC_CREATE, &sync_create, sizeof(sync_create));
-  if (ret != 0) {
-    LOGE("failed to create fence: %d %d", ret, sync_create.sync_obj);
+  if (icp_dev_handle > 0) {
+    ret = do_cam_control(m->cam_sync_fd, CAM_SYNC_CREATE, &sync_create, sizeof(sync_create));
+    if (ret != 0) {
+      LOGE("failed to create fence: %d %d", ret, sync_create.sync_obj);
+    }
+    sync_objs_bps_out[i] = sync_create.sync_obj;
   }
-  sync_objs_bps_out[i] = sync_create.sync_obj;
-  */
 
   // schedule request with camera request manager
   struct cam_req_mgr_sched_request req_mgr_sched_request = {0};
@@ -733,31 +738,33 @@ void SpectraCamera::enqueue_buffer(int i, bool dp) {
 
 void SpectraCamera::camera_map_bufs() {
   int ret;
-  for (int i = 0; i < FRAME_BUF_COUNT; i++) {
+  for (int i = 0; i < ife_buf_depth; i++) {
     // configure ISP to put the image in place
     struct cam_mem_mgr_map_cmd mem_mgr_map_cmd = {0};
-    mem_mgr_map_cmd.mmu_hdls[0] = m->device_iommu;
-    //mem_mgr_map_cmd.mmu_hdls[1] = m->icp_device_iommu;
-    //mem_mgr_map_cmd.num_hdl = 2;
-    mem_mgr_map_cmd.num_hdl = 1;
     mem_mgr_map_cmd.flags = CAM_MEM_FLAG_HW_READ_WRITE;
+    mem_mgr_map_cmd.mmu_hdls[0] = m->device_iommu;
+    mem_mgr_map_cmd.num_hdl = 1;
+    if (icp_dev_handle > 0) {
+      mem_mgr_map_cmd.num_hdl = 2;
+      mem_mgr_map_cmd.mmu_hdls[1] = m->icp_device_iommu;
+    }
 
-    // RAW bayer images
-    mem_mgr_map_cmd.fd = buf.camera_bufs_raw[i].fd;
-    ret = do_cam_control(m->video0_fd, CAM_REQ_MGR_MAP_BUF, &mem_mgr_map_cmd, sizeof(mem_mgr_map_cmd));
-    assert(ret == 0);
-    LOGD("map buf req: (fd: %d) 0x%x %d", buf.camera_bufs_raw[i].fd, mem_mgr_map_cmd.out.buf_handle, ret);
-    buf_handle_raw[i] = mem_mgr_map_cmd.out.buf_handle;
-
-    // TODO: this needs to match camera bufs length
-    // final processed images
-    VisionBuf *vb = buf.vipc_server->get_buffer(buf.stream_type, i);
-    mem_mgr_map_cmd.fd = vb->fd;
-    ret = do_cam_control(m->video0_fd, CAM_REQ_MGR_MAP_BUF, &mem_mgr_map_cmd, sizeof(mem_mgr_map_cmd));
-    LOGD("map buf req: (fd: %d) 0x%x %d", vb->fd, mem_mgr_map_cmd.out.buf_handle, ret);
-    buf_handle_yuv[i] = mem_mgr_map_cmd.out.buf_handle;
+    if (is_raw) {
+      // RAW bayer images
+      mem_mgr_map_cmd.fd = buf.camera_bufs_raw[i].fd;
+      ret = do_cam_control(m->video0_fd, CAM_REQ_MGR_MAP_BUF, &mem_mgr_map_cmd, sizeof(mem_mgr_map_cmd));
+      assert(ret == 0);
+      LOGD("map buf req: (fd: %d) 0x%x %d", buf.camera_bufs_raw[i].fd, mem_mgr_map_cmd.out.buf_handle, ret);
+      buf_handle_raw[i] = mem_mgr_map_cmd.out.buf_handle;
+    } else {
+      // final processed images
+      VisionBuf *vb = buf.vipc_server->get_buffer(buf.stream_type, i);
+      mem_mgr_map_cmd.fd = vb->fd;
+      ret = do_cam_control(m->video0_fd, CAM_REQ_MGR_MAP_BUF, &mem_mgr_map_cmd, sizeof(mem_mgr_map_cmd));
+      LOGD("map buf req: (fd: %d) 0x%x %d", vb->fd, mem_mgr_map_cmd.out.buf_handle, ret);
+      buf_handle_yuv[i] = mem_mgr_map_cmd.out.buf_handle;
+    }
   }
-  enqueue_req_multi(1, FRAME_BUF_COUNT, 0);
 }
 
 bool SpectraCamera::openSensor() {
@@ -870,7 +877,7 @@ void SpectraCamera::configISP() {
   // allocate IFE memory, then configure it
   ife_cmd.init(m, 67984, 0x20,
                CAM_MEM_FLAG_HW_READ_WRITE | CAM_MEM_FLAG_KMD_ACCESS | CAM_MEM_FLAG_UMD_ACCESS | CAM_MEM_FLAG_CMD_BUF_TYPE,
-               m->device_iommu, m->cdm_iommu, FRAME_BUF_COUNT);
+               m->device_iommu, m->cdm_iommu, ife_buf_depth);
   if (!is_raw) {
     ife_gamma_lut.init(m, 64*sizeof(uint32_t), 0x20,
                        CAM_MEM_FLAG_HW_READ_WRITE | CAM_MEM_FLAG_KMD_ACCESS | CAM_MEM_FLAG_UMD_ACCESS | CAM_MEM_FLAG_CMD_BUF_TYPE,
@@ -892,8 +899,6 @@ void SpectraCamera::configISP() {
 }
 
 void SpectraCamera::configICP() {
-  if (!enabled) return;
-
   /*
     Configures both the ICP and BPS.
   */
@@ -925,7 +930,7 @@ void SpectraCamera::configICP() {
 
   // BPS CMD buffer
   unsigned char striping_out[] = "\x00";
-  bps_cmd.init(m, FRAME_BUF_COUNT*ALIGNED_SIZE(464, 0x20), 0x20,
+  bps_cmd.init(m, ife_buf_depth*ALIGNED_SIZE(464, 0x20), 0x20,
                CAM_MEM_FLAG_HW_READ_WRITE | CAM_MEM_FLAG_KMD_ACCESS | CAM_MEM_FLAG_UMD_ACCESS | CAM_MEM_FLAG_CMD_BUF_TYPE | CAM_MEM_FLAG_HW_SHARED_ACCESS,
                m->icp_device_iommu);
 
@@ -1041,14 +1046,17 @@ void SpectraCamera::camera_close() {
 
     // release devices
     LOGD("-- Release devices");
+    if (icp_dev_handle > 0) {
+      ret = device_control(m->icp_fd, CAM_RELEASE_DEV, session_handle, icp_dev_handle);
+      LOGD("release icp: %d", ret);
+    }
     ret = device_control(m->isp_fd, CAM_RELEASE_DEV, session_handle, isp_dev_handle);
     LOGD("release isp: %d", ret);
     ret = device_control(csiphy_fd, CAM_RELEASE_DEV, session_handle, csiphy_dev_handle);
     LOGD("release csiphy: %d", ret);
 
-    for (int i = 0; i < FRAME_BUF_COUNT; i++) {
-      release(m->video0_fd, buf_handle_yuv[i]);
-      release(m->video0_fd, buf_handle_raw[i]);
+    for (int i = 0; i < ife_buf_depth; i++) {
+      release(m->video0_fd, is_raw ? buf_handle_raw[i] : buf_handle_yuv[i]);
     }
     LOGD("released buffers");
   }
@@ -1071,13 +1079,13 @@ void SpectraCamera::handle_camera_event(const cam_req_mgr_message *event_data) {
 
   if (real_id != 0) { // next ready
     if (real_id == 1) {idx_offset = main_id;}
-    int buf_idx = (real_id - 1) % FRAME_BUF_COUNT;
+    int buf_idx = (real_id - 1) % ife_buf_depth;
 
     // check for skipped frames
     if (main_id > frame_id_last + 1 && !skipped) {
       LOGE("camera %d realign", cc.camera_num);
       clear_req_queue();
-      enqueue_req_multi(real_id + 1, FRAME_BUF_COUNT - 1, 0);
+      enqueue_req_multi(real_id + 1, ife_buf_depth - 1, 0);
       skipped = true;
     } else if (main_id == frame_id_last + 1) {
       skipped = false;
@@ -1086,7 +1094,7 @@ void SpectraCamera::handle_camera_event(const cam_req_mgr_message *event_data) {
     // check for dropped requests
     if (real_id > request_id_last + 1) {
       LOGE("camera %d dropped requests %ld %ld", cc.camera_num, real_id, request_id_last);
-      enqueue_req_multi(request_id_last + 1 + FRAME_BUF_COUNT, real_id - (request_id_last + 1), 0);
+      enqueue_req_multi(request_id_last + 1 + ife_buf_depth, real_id - (request_id_last + 1), 0);
     }
 
     // metas
@@ -1099,12 +1107,12 @@ void SpectraCamera::handle_camera_event(const cam_req_mgr_message *event_data) {
     meta_data.timestamp_sof = timestamp; // this is timestamped in the kernel's SOF IRQ callback
 
     // dispatch
-    enqueue_req_multi(real_id + FRAME_BUF_COUNT, 1, 1);
+    enqueue_req_multi(real_id + ife_buf_depth, 1, 1);
   } else { // not ready
     if (main_id > frame_id_last + 10) {
       LOGE("camera %d reset after half second of no response", cc.camera_num);
       clear_req_queue();
-      enqueue_req_multi(request_id_last + 1, FRAME_BUF_COUNT, 0);
+      enqueue_req_multi(request_id_last + 1, ife_buf_depth, 0);
       frame_id_last = main_id;
       skipped = true;
     }