Up to Date Pull

* ui: fix onroad transition with active nav (commaai#23757)

* update release date

* bump panda

* nav: handle route responses with duplicate points (commaai#23769)

* nav: set settings stack based on uiState.prime_type (commaai#23770)

* nav: set settings stack based on uiState.prime_type

* also show on success

* Hyundai: move unused cars to dashcam only (commaai#23767)

* Move OptimaH to dashcamOnly

* Elantra GT also has no users for a year

* Update selfdrive/car/hyundai/interface.py

Co-authored-by: Adeeb Shihadeh <adeebshihadeh@gmail.com>

* add thneed optimizer (commaai#23772)

* add thneed optimizer

* local work group opt

* kernels and final mods

* release files

* build system touchups

* fix kernel path, rand inputs for self test

* broken since extra is gone

* update model replay ref

Co-authored-by: Comma Device <device@comma.ai>

* bump to 0.8.14

* UI: fix power buttons disable logic (commaai#23774)

* same check

* statsd: fix crash if there is no git repository (commaai#23775)

* fix crash if there is no git repository

* return default

Co-authored-by: Willem Melching <willem.melching@gmail.com>

* bump cereal

* Longitudinal: add Kf to carParams (commaai#23752)

* sane_kf

* long_kf + cereal

* update ref

Co-authored-by: Willem Melching <willem.melching@gmail.com>

* Snpemodel: Clean up after commaai#23772 (commaai#23780)

* #clean up snpemodel after commaai#23772

* Leave benchmark remove loop

* add branch name to startup master alert (commaai#23782)

* add branch name to startup master alert

* still orange

* update refs

* keep the new alert consistent for process replay

* process replay: fix branch name alert (commaai#23783)

* env is modified after start

* update ref

* fix in window for xmonad with SDL2 (commaai#23786)

* camerad: adjust CL priority on comma three (commaai#23790)

Co-authored-by: Comma Device <device@comma.ai>

* Longitudinal control: interpolate longitudinal plan (commaai#23787)

* interpolate longitudinal actuator delay

rename

* formatting

* interpolate v_target most importantly!

* fix interpolation and rename

* nicer setup

* left in from testing

* update refs

* camerad: fast debayer on c2 cameras (commaai#23795)

* fast debayer on c2 dcam

* add casts

* 128 local worksize on HDR debayer, 8 ms -> 3.5 ms

* width instead of saving rgb_width

Co-authored-by: Comma Device <device@comma.ai>

* camerad: use open_v4l_by_name_and_index on c2 also (commaai#23794)

* use open_v4l_by_name_and_index on c2 also

* remove open_v4l_by_name_and_index from qcom2

* Add missing CAMRYH_TSS2 fwdCamera f/w (commaai#23796)

`@Ty.ler#0303`  2022 Toyota Camry SE Hybrid  DongleID/route 9477a77ffc4e8bd4|2022-02-18--15-40-25

* HKG: Add FW for 2019 Veloster (commaai#23756)

* HKG: Add FW for 2019 Veloster

* Add esp FW

* bigmodel (commaai#23684)

* Added wide cam vipc client and bigmodel transform logic

* Added wide_frame to ModelState, should still work normally

* Refactored image input into addImage method, should still work normally

* Updated thneed/compile.cc

* Bigmodel, untested: 44f83118-b375-4d4c-ae12-2017124f0cf4/200

* Have to initialize extra buffer in SNPEModel

* Default paramater value in the wrong place I think

* Move USE_EXTRA to SConscript

* New model: 6c34d59a-acc3-4877-84bd-904c10745ba6/250

* move use extra check to runtime, not on C2

* this is always true

* more C2 checks

* log if frames are out of sync

* more logging on no frame

* store in pointer

* print sof

* add sync logic

* log based on sof difference as well

* keep both models

* less assumptions

* define above thneed

* typo

* simplify

* no need for second client is main is already wide

* more comments update

* no optional reference

* more logging to debug lags

* add to release files

* both defines

* New model: 6831a77f-2574-4bfb-8077-79b0972a2771/950

* Path offset no longer relevant

* Remove duplicate execute

* Moved bigmodel back to big_supercombo.dlc

* add wide vipc stream

* Tici must be tici

* Needs state too

* add wide cam support to model replay

* handle syncing better

* ugh, c2

* print that

* handle ecam lag

* skip first one

* so close

* update refs

Co-authored-by: mitchellgoffpc <mitchellgoffpc@gmail.com>
Co-authored-by: Harald Schafer <harald.the.engineer@gmail.com>
Co-authored-by: Adeeb Shihadeh <adeebshihadeh@gmail.com>
Co-authored-by: Comma Device <device@comma.ai>

* add to release notes

* paramsd cpu usage has been increasing

* boardd: connect to all pandas if none specified (commaai#23805)

* Refactor convolutions (commaai#23807)

* one conv with defines

* add conv

* building works on C3

* this is num_outputs too, process replay is so useful

Co-authored-by: Comma Device <device@comma.ai>

* cleanup import chaining of T_IDXS (commaai#23806)

* bump panda

* ui.py: fix model lead plotting

* Param to allow upload over cellular (commaai#23811)

* param to allow upload over cellular

* brackets to be consistent

* move outside loop

* fix tests

Co-authored-by: Willem Melching <willem.melching@gmail.com>

* longitudinal planner: disable change cost when stopped. not engaged or gas pressed (commaai#23639)

* disable change cost completely on standstill and gas press

* cleanup

* set accel to zero

* clean up logic around standstill

* update ref

* onroad_test: increase camerad cpu usage after debayer changes

* Update raw_logger/framereader to new new ffmpeg api (commaai#23799)

* Update raw_logger/framereader to new new ffmpeg api

* Can be removed now

* Handled by avcodec_free_context

* handle EAGAIN and clean up

Co-authored-by: Willem Melching <willem.melching@gmail.com>
Co-authored-by: Adeeb Shihadeh <adeebshihadeh@gmail.com>
Co-authored-by: Shane Smiskol <shane@smiskol.com>
Co-authored-by: George Hotz <72895+geohot@users.noreply.github.com>
Co-authored-by: Comma Device <device@comma.ai>
Co-authored-by: Jack Huang <jackhuang1205@gmail.com>
Co-authored-by: ClockeNessMnstr <locke.dftc@gmail.com>
Co-authored-by: redacid95 <jonathanc.olivier@gmail.com>
Co-authored-by: Erich Moraga <33645296+ErichMoraga@users.noreply.github.com>
Co-authored-by: Jason Wen <47793918+sunnyhaibin@users.noreply.github.com>
Co-authored-by: mitchellgoffpc <mitchellgoffpc@gmail.com>
Co-authored-by: Harald Schafer <harald.the.engineer@gmail.com>
Co-authored-by: Robbe Derks <robbe.derks@gmail.com>

RELEASES.md

@@ -1,4 +1,8 @@
-Version 0.8.13 (2022-02-16)
+Version 0.8.14 (2022-0X-XX)
+========================
+ * bigmodel!
+
+Version 0.8.13 (2022-02-18)
 ========================
  * Improved driver monitoring
    * Retuned driver pose learner for relaxed driving positions

common/window.py

@@ -11,11 +11,18 @@ def __init__(self, w, h, caption="window", double=False, halve=False):
     self.double = double
     self.halve = halve
     if self.double:
-      self.screen = pygame.display.set_mode((w*2, h*2))
+      self.rw, self.rh = w*2, h*2
     elif self.halve:
-      self.screen = pygame.display.set_mode((w//2, h//2))
+      self.rw, self.rh = w//2, h//2
     else:
-      self.screen = pygame.display.set_mode((w, h))
+      self.rw, self.rh = w, h
+    self.screen = pygame.display.set_mode((self.rw, self.rh))
+    pygame.display.flip()
+
+    # hack for xmonad, it shrinks the window by 6 pixels after the display.flip
+    if self.screen.get_width() != self.rw:
+      self.screen = pygame.display.set_mode((self.rw+(self.rw-self.screen.get_width()), self.rh+(self.rh-self.screen.get_height())))
+      pygame.display.flip()
 
   def draw(self, out):
     pygame.event.pump()

release/files_common

@@ -58,6 +58,7 @@ common/transformations/transformations.pyx
 common/api/__init__.py
 
 models/supercombo.dlc
+models/big_supercombo.dlc
 models/dmonitoring_model_q.dlc
 
 release/*
@@ -426,7 +427,9 @@ selfdrive/modeld/transforms/transform.cl
 selfdrive/modeld/thneed/thneed.*
 selfdrive/modeld/thneed/serialize.cc
 selfdrive/modeld/thneed/compile.cc
+selfdrive/modeld/thneed/optimizer.cc
 selfdrive/modeld/thneed/include/*
+selfdrive/modeld/thneed/kernels/*.cl
 
 selfdrive/modeld/runners/snpemodel.cc
 selfdrive/modeld/runners/snpemodel.h

selfdrive/athena/athenad.py

@@ -163,6 +163,8 @@ def upload_handler(end_event: threading.Event) -> None:
   sm = messaging.SubMaster(['deviceState'])
   tid = threading.get_ident()
 
+  cellular_unmetered = Params().get_bool("CellularUnmetered")
+
   while not end_event.is_set():
     cur_upload_items[tid] = None
 
@@ -182,7 +184,7 @@ def upload_handler(end_event: threading.Event) -> None:
       # Check if uploading over cell is allowed
       sm.update(0)
       cell = sm['deviceState'].networkType not in [NetworkType.wifi, NetworkType.ethernet]
-      if cell and (not cur_upload_items[tid].allow_cellular):
+      if cell and (not cur_upload_items[tid].allow_cellular) and (not cellular_unmetered):
         retry_upload(tid, end_event, False)
         continue
 

selfdrive/athena/tests/helpers.py

@@ -53,14 +53,18 @@ class MockParams():
   default_params = {
     "DongleId": b"0000000000000000",
     "GithubSshKeys": b"ssh-rsa AAAAB3NzaC1yc2EAAAADAQABAAABAQC307aE+nuHzTAgaJhzSf5v7ZZQW9gaperjhCmyPyl4PzY7T1mDGenTlVTN7yoVFZ9UfO9oMQqo0n1OwDIiqbIFxqnhrHU0cYfj88rI85m5BEKlNu5RdaVTj1tcbaPpQc5kZEolaI1nDDjzV0lwS7jo5VYDHseiJHlik3HH1SgtdtsuamGR2T80q1SyW+5rHoMOJG73IH2553NnWuikKiuikGHUYBd00K1ilVAK2xSiMWJp55tQfZ0ecr9QjEsJ+J/efL4HqGNXhffxvypCXvbUYAFSddOwXUPo5BTKevpxMtH+2YrkpSjocWA04VnTYFiPG6U4ItKmbLOTFZtPzoez private",  # noqa: E501
-    "AthenadUploadQueue": '[]'
+    "AthenadUploadQueue": '[]',
+    "CellularUnmetered": False,
   }
   params = default_params.copy()
 
   @staticmethod
   def restore_defaults():
     MockParams.params = MockParams.default_params.copy()
 
+  def get_bool(self, k):
+    return bool(MockParams.params.get(k))
+
   def get(self, k, encoding=None):
     ret = MockParams.params.get(k)
     if ret is not None and encoding is not None:

selfdrive/boardd/boardd.cc

@@ -610,11 +610,11 @@ void pigeon_thread(Panda *panda) {
 }
 
 void boardd_main_thread(std::vector<std::string> serials) {
-  if (serials.size() == 0) serials.push_back("");
-
   PubMaster pm({"pandaStates", "peripheralState"});
   LOGW("attempting to connect");
 
+  if (serials.size() == 0) serials = Panda::list();
+
   // connect to all provided serials
   std::vector<Panda *> pandas;
   for (int i = 0; i < serials.size() && !do_exit; /**/) {

selfdrive/camerad/cameras/camera_common.cc

@@ -22,6 +22,7 @@
 #include "CL/cl_ext_qcom.h"
 #include "selfdrive/camerad/cameras/camera_qcom.h"
 #elif QCOM2
+#include "CL/cl_ext_qcom.h"
 #include "selfdrive/camerad/cameras/camera_qcom2.h"
 #elif WEBCAM
 #include "selfdrive/camerad/cameras/camera_webcam.h"
@@ -36,6 +37,7 @@ class Debayer {
   Debayer(cl_device_id device_id, cl_context context, const CameraBuf *b, const CameraState *s) {
     char args[4096];
     const CameraInfo *ci = &s->ci;
+    hdr_ = ci->hdr;
     snprintf(args, sizeof(args),
              "-cl-fast-relaxed-math -cl-denorms-are-zero "
              "-DFRAME_WIDTH=%d -DFRAME_HEIGHT=%d -DFRAME_STRIDE=%d "
@@ -62,9 +64,20 @@ class Debayer {
       CL_CHECK(clSetKernelArg(krnl_, 2, localMemSize, 0));
       CL_CHECK(clEnqueueNDRangeKernel(q, krnl_, 2, NULL, globalWorkSize, localWorkSize, 0, 0, debayer_event));
     } else {
-      const size_t debayer_work_size = height;  // doesn't divide evenly, is this okay?
-      CL_CHECK(clSetKernelArg(krnl_, 2, sizeof(float), &gain));
-      CL_CHECK(clEnqueueNDRangeKernel(q, krnl_, 1, NULL, &debayer_work_size, NULL, 0, 0, debayer_event));
+      if (hdr_) {
+        // HDR requires a 1-D kernel due to the DPCM compression
+        const size_t debayer_local_worksize = 128;
+        const size_t debayer_work_size = height;  // doesn't divide evenly, is this okay?
+        CL_CHECK(clSetKernelArg(krnl_, 2, sizeof(float), &gain));
+        CL_CHECK(clEnqueueNDRangeKernel(q, krnl_, 1, NULL, &debayer_work_size, &debayer_local_worksize, 0, 0, debayer_event));
+      } else {
+        const int debayer_local_worksize = 32;
+        assert(width % 2 == 0);
+        const size_t globalWorkSize[] = {size_t(height), size_t(width / 2)};
+        const size_t localWorkSize[] = {debayer_local_worksize, debayer_local_worksize};
+        CL_CHECK(clSetKernelArg(krnl_, 2, sizeof(float), &gain));
+        CL_CHECK(clEnqueueNDRangeKernel(q, krnl_, 2, NULL, globalWorkSize, localWorkSize, 0, 0, debayer_event));
+      }
     }
   }
 
@@ -74,6 +87,7 @@ class Debayer {
 
 private:
   cl_kernel krnl_;
+  bool hdr_;
 };
 
 void CameraBuf::init(cl_device_id device_id, cl_context context, CameraState *s, VisionIpcServer * v, int frame_cnt, VisionStreamType init_rgb_type, VisionStreamType init_yuv_type, release_cb init_release_callback) {
@@ -427,8 +441,7 @@ void common_process_driver_camera(MultiCameraState *s, CameraState *c, int cnt)
 
 void camerad_thread() {
   cl_device_id device_id = cl_get_device_id(CL_DEVICE_TYPE_DEFAULT);
-   // TODO: do this for QCOM2 too
-#if defined(QCOM)
+#if defined(QCOM) || defined(QCOM2)
   const cl_context_properties props[] = {CL_CONTEXT_PRIORITY_HINT_QCOM, CL_PRIORITY_HINT_HIGH_QCOM, 0};
   cl_context context = CL_CHECK_ERR(clCreateContext(props, 1, &device_id, NULL, NULL, &err));
 #else
@@ -447,3 +460,16 @@ void camerad_thread() {
 
   CL_CHECK(clReleaseContext(context));
 }
+
+int open_v4l_by_name_and_index(const char name[], int index, int flags) {
+  for (int v4l_index = 0; /**/; ++v4l_index) {
+    std::string v4l_name = util::read_file(util::string_format("/sys/class/video4linux/v4l-subdev%d/name", v4l_index));
+    if (v4l_name.empty()) return -1;
+    if (v4l_name.find(name) == 0) {
+      if (index == 0) {
+        return HANDLE_EINTR(open(util::string_format("/dev/v4l-subdev%d", v4l_index).c_str(), flags));
+      }
+      index--;
+    }
+  }
+}

selfdrive/camerad/cameras/camera_common.h

@@ -135,3 +135,5 @@ void cameras_run(MultiCameraState *s);
 void cameras_close(MultiCameraState *s);
 void camera_autoexposure(CameraState *s, float grey_frac);
 void camerad_thread();
+
+int open_v4l_by_name_and_index(const char name[], int index = 0, int flags = O_RDWR | O_NONBLOCK);

selfdrive/camerad/cameras/camera_qcom.cc

@@ -399,7 +399,7 @@ static void sensors_init(MultiCameraState *s) {
     .output_format = MSM_SENSOR_BAYER,
   }};
 
-  unique_fd sensorinit_fd = HANDLE_EINTR(open("/dev/v4l-subdev11", O_RDWR | O_NONBLOCK));
+  unique_fd sensorinit_fd = open_v4l_by_name_and_index("msm_sensor_init");
   assert(sensorinit_fd >= 0);
   for (auto &info : slave_infos) {
     info.power_setting_array.power_setting = &info.power_setting_array.power_setting_a[0];
@@ -417,31 +417,22 @@ static void camera_open(CameraState *s, bool is_road_cam) {
   struct msm_actuator_cfg_data actuator_cfg_data = {};
 
   // open devices
-  const char *sensor_dev;
+  s->csid_fd = open_v4l_by_name_and_index("msm_csid", is_road_cam ? 0 : 2);
+  assert(s->csid_fd >= 0);
+  s->csiphy_fd = open_v4l_by_name_and_index("msm_csiphy", is_road_cam ? 0 : 2);
+  assert(s->csiphy_fd >= 0);
+  s->isp_fd = open_v4l_by_name_and_index("vfe", is_road_cam ? 0 : 1);
+  assert(s->isp_fd >= 0);
+
   if (is_road_cam) {
-    s->csid_fd = HANDLE_EINTR(open("/dev/v4l-subdev3", O_RDWR | O_NONBLOCK));
-    assert(s->csid_fd >= 0);
-    s->csiphy_fd = HANDLE_EINTR(open("/dev/v4l-subdev0", O_RDWR | O_NONBLOCK));
-    assert(s->csiphy_fd >= 0);
-    sensor_dev = "/dev/v4l-subdev17";
-    s->isp_fd = HANDLE_EINTR(open("/dev/v4l-subdev13", O_RDWR | O_NONBLOCK));
-    assert(s->isp_fd >= 0);
-    s->actuator_fd = HANDLE_EINTR(open("/dev/v4l-subdev7", O_RDWR | O_NONBLOCK));
+    s->actuator_fd = open_v4l_by_name_and_index("msm_actuator");
     assert(s->actuator_fd >= 0);
-  } else {
-    s->csid_fd = HANDLE_EINTR(open("/dev/v4l-subdev5", O_RDWR | O_NONBLOCK));
-    assert(s->csid_fd >= 0);
-    s->csiphy_fd = HANDLE_EINTR(open("/dev/v4l-subdev2", O_RDWR | O_NONBLOCK));
-    assert(s->csiphy_fd >= 0);
-    sensor_dev = "/dev/v4l-subdev18";
-    s->isp_fd = HANDLE_EINTR(open("/dev/v4l-subdev14", O_RDWR | O_NONBLOCK));
-    assert(s->isp_fd >= 0);
   }
 
   // wait for sensor device
   // on first startup, these devices aren't present yet
   for (int i = 0; i < 10; i++) {
-    s->sensor_fd = HANDLE_EINTR(open(sensor_dev, O_RDWR | O_NONBLOCK));
+    s->sensor_fd = open_v4l_by_name_and_index(is_road_cam ? "imx298" : "ov8865_sunny");
     if (s->sensor_fd >= 0) break;
     LOGW("waiting for sensors...");
     util::sleep_for(1000); // sleep one second
@@ -908,7 +899,7 @@ void cameras_open(MultiCameraState *s) {
   s->v4l_fd = HANDLE_EINTR(open("/dev/video0", O_RDWR | O_NONBLOCK));
   assert(s->v4l_fd >= 0);
 
-  s->ispif_fd = HANDLE_EINTR(open("/dev/v4l-subdev15", O_RDWR | O_NONBLOCK));
+  s->ispif_fd = open_v4l_by_name_and_index("msm_ispif");
   assert(s->ispif_fd >= 0);
 
   // ISPIF: stop

selfdrive/camerad/cameras/camera_qcom2.cc

@@ -578,19 +578,6 @@ static void camera_init(MultiCameraState *multi_cam_state, VisionIpcServer * v,
   s->buf.init(device_id, ctx, s, v, FRAME_BUF_COUNT, rgb_type, yuv_type);
 }
 
-int open_v4l_by_name_and_index(const char name[], int index, int flags = O_RDWR | O_NONBLOCK) {
-  for (int v4l_index = 0; /**/; ++v4l_index) {
-    std::string v4l_name = util::read_file(util::string_format("/sys/class/video4linux/v4l-subdev%d/name", v4l_index));
-    if (v4l_name.empty()) return -1;
-    if (v4l_name.find(name) == 0) {
-      if (index == 0) {
-        return open(util::string_format("/dev/v4l-subdev%d", v4l_index).c_str(), flags);
-      }
-      index--;
-    }
-  }
-}
-
 static void camera_open(CameraState *s) {
   s->sensor_fd = open_v4l_by_name_and_index("cam-sensor-driver", s->camera_num);
   assert(s->sensor_fd >= 0);

selfdrive/camerad/cameras/debayer.cl

@@ -26,6 +26,8 @@ float3 srgb_gamma(float3 p) {
 	return select(ph, pl, islessequal(p, 0.0031308f));
 }
 
+#if HDR
+
 __constant int dpcm_lookup[512] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31, 0, -1, -2, -3, -4, -5, -6, -7, -8, -9, -10, -11, -12, -13, -14, -15, -16, -17, -18, -19, -20, -21, -22, -23, -24, -25, -26, -27, -28, -29, -30, -31, 935, 951, 967, 983, 999, 1015, 1031, 1047, 1063, 1079, 1095, 1111, 1127, 1143, 1159, 1175, 1191, 1207, 1223, 1239, 1255, 1271, 1287, 1303, 1319, 1335, 1351, 1367, 1383, 1399, 1415, 1431, -935, -951, -967, -983, -999, -1015, -1031, -1047, -1063, -1079, -1095, -1111, -1127, -1143, -1159, -1175, -1191, -1207, -1223, -1239, -1255, -1271, -1287, -1303, -1319, -1335, -1351, -1367, -1383, -1399, -1415, -1431, 419, 427, 435, 443, 451, 459, 467, 475, 483, 491, 499, 507, 515, 523, 531, 539, 547, 555, 563, 571, 579, 587, 595, 603, 611, 619, 627, 635, 643, 651, 659, 667, 675, 683, 691, 699, 707, 715, 723, 731, 739, 747, 755, 763, 771, 779, 787, 795, 803, 811, 819, 827, 835, 843, 851, 859, 867, 875, 883, 891, 899, 907, 915, 923, -419, -427, -435, -443, -451, -459, -467, -475, -483, -491, -499, -507, -515, -523, -531, -539, -547, -555, -563, -571, -579, -587, -595, -603, -611, -619, -627, -635, -643, -651, -659, -667, -675, -683, -691, -699, -707, -715, -723, -731, -739, -747, -755, -763, -771, -779, -787, -795, -803, -811, -819, -827, -835, -843, -851, -859, -867, -875, -883, -891, -899, -907, -915, -923, 161, 165, 169, 173, 177, 181, 185, 189, 193, 197, 201, 205, 209, 213, 217, 221, 225, 229, 233, 237, 241, 245, 249, 253, 257, 261, 265, 269, 273, 277, 281, 285, 289, 293, 297, 301, 305, 309, 313, 317, 321, 325, 329, 333, 337, 341, 345, 349, 353, 357, 361, 365, 369, 373, 377, 381, 385, 389, 393, 397, 401, 405, 409, 413, -161, -165, -169, -173, -177, -181, -185, -189, -193, -197, -201, -205, -209, -213, -217, -221, -225, -229, -233, -237, -241, -245, -249, -253, -257, -261, -265, -269, -273, -277, -281, -285, -289, -293, -297, -301, -305, -309, -313, -317, -321, -325, -329, -333, -337, -341, -345, -349, -353, -357, -361, -365, -369, -373, -377, -381, -385, -389, -393, -397, -401, -405, -409, -413, 32, 34, 36, 38, 40, 42, 44, 46, 48, 50, 52, 54, 56, 58, 60, 62, 64, 66, 68, 70, 72, 74, 76, 78, 80, 82, 84, 86, 88, 90, 92, 94, 96, 98, 100, 102, 104, 106, 108, 110, 112, 114, 116, 118, 120, 122, 124, 126, 128, 130, 132, 134, 136, 138, 140, 142, 144, 146, 148, 150, 152, 154, 156, 158, -32, -34, -36, -38, -40, -42, -44, -46, -48, -50, -52, -54, -56, -58, -60, -62, -64, -66, -68, -70, -72, -74, -76, -78, -80, -82, -84, -86, -88, -90, -92, -94, -96, -98, -100, -102, -104, -106, -108, -110, -112, -114, -116, -118, -120, -122, -124, -126, -128, -130, -132, -134, -136, -138, -140, -142, -144, -146, -148, -150, -152, -154, -156, -158};
 
 inline uint4 decompress(uint4 p, uint4 pl) {
@@ -35,15 +37,22 @@ inline uint4 decompress(uint4 p, uint4 pl) {
   return select(r2, r1, p < 0x200);
 }
 
+#endif
+
 __kernel void debayer10(__global uchar const * const in,
                         __global uchar * out, float digital_gain)
 {
   const int oy = get_global_id(0);
   if (oy >= RGB_HEIGHT) return;
   const int iy = oy * 2;
 
+#if HDR
   uint4 pint_last;
   for (int ox = 0; ox < RGB_WIDTH; ox += 2) {
+#else
+  int ox = get_global_id(1) * 2;
+  {
+#endif
     const int ix = (ox/2) * 5;
 
     // TODO: why doesn't this work for the frontview

selfdrive/car/hyundai/interface.py

@@ -30,6 +30,11 @@ def get_params(candidate, fingerprint=gen_empty_fingerprint(), car_fw=[]):  # py
 
     ret.pcmCruise = not ret.openpilotLongitudinalControl
 
+    # These cars have been put into dashcam only due to both a lack of users and test coverage.
+    # These cars likely still work fine. Once a user confirms each car works and a test route is
+    # added to selfdrive/test/test_routes, we can remove it from this list.
+    ret.dashcamOnly = candidate in {CAR.KIA_OPTIMA_H, CAR.ELANTRA_GT_I30}
+
     ret.steerActuatorDelay = 0.1  # Default delay
     ret.steerRateCost = 0.5
     ret.steerLimitTimer = 0.4

selfdrive/car/hyundai/values.py

@@ -647,10 +647,14 @@ class Buttons:
       b'\xf1\x00JS__ SCC H-CUP      1.00 1.02 95650-J3200         ',
       b'\xf1\x00JS__ SCC HNCUP      1.00 1.02 95650-J3100         ',
     ],
-    (Ecu.esp, 0x7d1, None): [b'\xf1\x00\x00\x00\x00\x00\x00\x00', ],
+    (Ecu.esp, 0x7d1, None): [
+      b'\xf1\x00\x00\x00\x00\x00\x00\x00',
+      b'\xf1\x816V8RAC00121.ELF\xf1\x00\x00\x00\x00\x00\x00\x00',
+    ],
     (Ecu.engine, 0x7e0, None): [
       b'\x01TJS-JNU06F200H0A',
       b'\x01TJS-JDK06F200H0A',
+      b'391282BJF5 ',
     ],
     (Ecu.eps, 0x7d4, None): [b'\xf1\x00JSL MDPS C 1.00 1.03 56340-J3000 8308', ],
     (Ecu.fwdCamera, 0x7c4, None): [

selfdrive/car/interfaces.py

@@ -92,10 +92,12 @@ def get_std_params(candidate, fingerprint):
     ret.stoppingControl = True
     ret.longitudinalTuning.deadzoneBP = [0.]
     ret.longitudinalTuning.deadzoneV = [0.]
+    ret.longitudinalTuning.kf = 1.
     ret.longitudinalTuning.kpBP = [0.]
     ret.longitudinalTuning.kpV = [1.]
     ret.longitudinalTuning.kiBP = [0.]
     ret.longitudinalTuning.kiV = [1.]
+    # TODO estimate car specific lag, use .15s for now
     ret.longitudinalActuatorDelayLowerBound = 0.15
     ret.longitudinalActuatorDelayUpperBound = 0.15
     ret.steerLimitTimer = 1.0

selfdrive/car/toyota/values.py

@@ -398,6 +398,7 @@ class CAR:
     (Ecu.fwdCamera, 0x750, 109): [
       b'\x028646F3305200\x00\x00\x00\x008646G5301200\x00\x00\x00\x00',
       b'\x028646F3305300\x00\x00\x00\x008646G5301200\x00\x00\x00\x00',
+      b'\x028646F3305300\x00\x00\x00\x008646G3304000\x00\x00\x00\x00',
     ],
   },
   CAR.CHR: {

selfdrive/common/params.cc

@@ -91,6 +91,7 @@ std::unordered_map<std::string, uint32_t> keys = {
     {"CarParams", CLEAR_ON_MANAGER_START | CLEAR_ON_IGNITION_ON},
     {"CarParamsCache", CLEAR_ON_MANAGER_START},
     {"CarVin", CLEAR_ON_MANAGER_START | CLEAR_ON_IGNITION_ON},
+    {"CellularUnmetered", PERSISTENT},
     {"CompletedTrainingVersion", PERSISTENT},
     {"ControlsReady", CLEAR_ON_MANAGER_START | CLEAR_ON_IGNITION_ON},
     {"CurrentRoute", CLEAR_ON_MANAGER_START | CLEAR_ON_IGNITION_ON},

selfdrive/common/version.h

@@ -1 +1 @@
-#define COMMA_VERSION "0.8.13"
+#define COMMA_VERSION "0.8.14"

selfdrive/controls/controlsd.py

@@ -492,7 +492,8 @@ def state_control(self, CS):
     if not self.joystick_mode:
       # accel PID loop
       pid_accel_limits = self.CI.get_pid_accel_limits(self.CP, CS.vEgo, self.v_cruise_kph * CV.KPH_TO_MS)
-      actuators.accel = self.LoC.update(self.active, CS, self.CP, long_plan, pid_accel_limits)
+      t_since_plan = (self.sm.frame - self.sm.rcv_frame['longitudinalPlan']) * DT_CTRL
+      actuators.accel = self.LoC.update(self.active, CS, self.CP, long_plan, pid_accel_limits, t_since_plan)
 
       # Steering PID loop and lateral MPC
       lat_active = self.active and not CS.steerWarning and not CS.steerError and CS.vEgo > self.CP.minSteerSpeed