mawei/sunnypilot
1
0
Fork 0
mirror of https://github.com/sunnypilot/sunnypilot.git synced 2026-07-15 19:42:04 +08:00
Code Issues Packages Projects Releases Wiki Activity

Merge branch 'master' into mads-new

Browse Source
This commit is contained in:
Jason Wen
2024-09-30 16:31:52 -04:00
parent 738484e628 c6cd23210c
commit 42e3061748
17 changed files with 72 additions and 47 deletions
Download Patch File Download Diff File Expand all files Collapse all files
.github/workflows/stale.yaml
+1 -1
View File
@@ -21,7 +21,7 @@ jobs:
close-pr-message: 'This PR has been automatically closed due to inactivity. Feel free to re-open once activity resumes.'
stale-pr-label: stale
delete-branch: ${{ github.event.pull_request.head.repo.full_name == 'commaai/openpilot' }} # only delete branches on the main repo
exempt-pr-labels: "ignore stale,needs testing,car port,car" # if wip or it needs testing from the community, don't mark as stale
exempt-pr-labels: "ignore stale,needs testing" # if wip or it needs testing from the community, don't mark as stale
days-before-pr-stale: ${{ env.DAYS_BEFORE_PR_STALE }}
days-before-pr-close: ${{ env.DAYS_BEFORE_PR_CLOSE }}
.github/workflows/ui_preview.yaml
+2
View File
@@ -135,11 +135,13 @@ jobs:
TABLE="${TABLE}</table>"
TABLE="${TABLE}</details>"
if $open; then
TABLE="<details open>${TABLE}"
else
TABLE="<details>${TABLE}"
fi
echo "DIFF=$DIFF$TABLE" >> "$GITHUB_OUTPUT"
- name: Saving proposed ui
cereal/messaging/__init__.py
+17 -4
View File
@@ -109,12 +109,26 @@ class FrequencyTracker:
self.min_freq = min_freq * 0.8
self.max_freq = max_freq * 1.2
self.recv_dts: Deque[float] = deque(maxlen=int(10 * freq))
self.recv_dts_sum = 0.0
self.recent_recv_dts: Deque[float] = deque(maxlen=int(freq))
self.recent_recv_dts_sum = 0.0
self.prev_time = 0.0
def record_recv_time(self, cur_time: float) -> None:
# TODO: Handle case where cur_time is less than prev_time
if self.prev_time > 1e-5:
self.recv_dts.append(cur_time - self.prev_time)
dt = cur_time - self.prev_time
if len(self.recv_dts) == self.recv_dts.maxlen:
self.recv_dts_sum -= self.recv_dts[0]
self.recv_dts.append(dt)
self.recv_dts_sum += dt
if len(self.recent_recv_dts) == self.recent_recv_dts.maxlen:
self.recent_recv_dts_sum -= self.recent_recv_dts[0]
self.recent_recv_dts.append(dt)
self.recent_recv_dts_sum += dt
self.prev_time = cur_time
@property
@@ -122,12 +136,11 @@ class FrequencyTracker:
if not self.recv_dts:
return False
avg_freq = len(self.recv_dts) / sum(self.recv_dts)
avg_freq = len(self.recv_dts) / self.recv_dts_sum
if self.min_freq <= avg_freq <= self.max_freq:
return True
recent_dts = list(self.recv_dts)[-int(self.recv_dts.maxlen / 10):]
avg_freq_recent = len(recent_dts) / sum(recent_dts)
avg_freq_recent = len(self.recent_recv_dts) / self.recent_recv_dts_sum
return self.min_freq <= avg_freq_recent <= self.max_freq
selfdrive/assets/offroad/icon_openpilot.png
-3
View File
@@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:a3cd3b91673eded1282e2082be0efa8e54ed477b5feb3580e521d08078e18ed1
size 42640
selfdrive/modeld/modeld.py
+1 -2
View File
@@ -96,8 +96,7 @@ class ModelState:
# if getCLBuffer is not None, frame will be None
self.model.setInputBuffer("input_imgs", self.frame.prepare(buf, transform.flatten(), self.model.getCLBuffer("input_imgs")))
if wbuf is not None:
self.model.setInputBuffer("big_input_imgs", self.wide_frame.prepare(wbuf, transform_wide.flatten(), self.model.getCLBuffer("big_input_imgs")))
self.model.setInputBuffer("big_input_imgs", self.wide_frame.prepare(wbuf, transform_wide.flatten(), self.model.getCLBuffer("big_input_imgs")))
if prepare_only:
return None
selfdrive/modeld/models/commonmodel.cc
+7 -6
View File
@@ -13,7 +13,7 @@ ModelFrame::ModelFrame(cl_device_id device_id, cl_context context) {
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));
net_input_cl = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_WRITE, MODEL_FRAME_SIZE * sizeof(uint8_t), NULL, &err));
net_input_cl = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_WRITE, frame_size_bytes, NULL, &err));
transform_init(&transform, context, device_id);
loadyuv_init(&loadyuv, context, device_id, MODEL_WIDTH, MODEL_HEIGHT);
@@ -24,15 +24,16 @@ uint8_t* ModelFrame::prepare(cl_mem yuv_cl, int frame_width, int frame_height, i
yuv_cl, frame_width, frame_height, frame_stride, frame_uv_offset,
y_cl, u_cl, v_cl, MODEL_WIDTH, MODEL_HEIGHT, projection);
loadyuv_queue(&loadyuv, q, y_cl, u_cl, v_cl, net_input_cl);
if (output == NULL) {
loadyuv_queue(&loadyuv, q, y_cl, u_cl, v_cl, net_input_cl);
std::memmove(&input_frames[0], &input_frames[MODEL_FRAME_SIZE], sizeof(uint8_t) * MODEL_FRAME_SIZE);
CL_CHECK(clEnqueueReadBuffer(q, net_input_cl, CL_TRUE, 0, MODEL_FRAME_SIZE * sizeof(uint8_t), &input_frames[MODEL_FRAME_SIZE], 0, nullptr, nullptr));
std::memmove(&input_frames[0], &input_frames[MODEL_FRAME_SIZE], frame_size_bytes);
CL_CHECK(clEnqueueReadBuffer(q, net_input_cl, CL_TRUE, 0, frame_size_bytes, &input_frames[MODEL_FRAME_SIZE], 0, nullptr, nullptr));
clFinish(q);
return &input_frames[0];
} else {
loadyuv_queue(&loadyuv, q, y_cl, u_cl, v_cl, *output, true);
copy_queue(&loadyuv, q, *output, *output, frame_size_bytes, 0, frame_size_bytes);
copy_queue(&loadyuv, q, net_input_cl, *output, 0, frame_size_bytes, frame_size_bytes);
// NOTE: Since thneed is using a different command queue, this clFinish is needed to ensure the image is ready.
clFinish(q);
return NULL;
selfdrive/modeld/models/commonmodel.h
+1
View File
@@ -26,6 +26,7 @@ public:
const int MODEL_HEIGHT = 256;
const int MODEL_FRAME_SIZE = MODEL_WIDTH * MODEL_HEIGHT * 3 / 2;
const int buf_size = MODEL_FRAME_SIZE * 2;
const size_t frame_size_bytes = MODEL_FRAME_SIZE * sizeof(uint8_t);
private:
Transform transform;
selfdrive/modeld/models/supercombo.onnx
+1 -1
View File
@@ -1,3 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:47417a13d8a8d6af6a1562834eee538e3b43242c6e277ab6f1978a78a6785b7d
oid sha256:21c88ca8a3de59fb11dc3e5888476f6bb627f3647eb0d199680d598e8bf31c0c
size 62486469
selfdrive/modeld/transforms/loadyuv.cc
+12 -10
View File
@@ -33,17 +33,8 @@ void loadyuv_destroy(LoadYUVState* s) {
void loadyuv_queue(LoadYUVState* s, cl_command_queue q,
cl_mem y_cl, cl_mem u_cl, cl_mem v_cl,
cl_mem out_cl, bool do_shift) {
cl_mem out_cl) {
cl_int global_out_off = 0;
if (do_shift) {
// shift the image in slot 1 to slot 0, then place the new image in slot 1
global_out_off += (s->width*s->height) + (s->width/2)*(s->height/2)*2;
CL_CHECK(clSetKernelArg(s->copy_krnl, 0, sizeof(cl_mem), &out_cl));
CL_CHECK(clSetKernelArg(s->copy_krnl, 1, sizeof(cl_int), &global_out_off));
const size_t copy_work_size = global_out_off/8;
CL_CHECK(clEnqueueNDRangeKernel(q, s->copy_krnl, 1, NULL,
&copy_work_size, NULL, 0, 0, NULL));
}
CL_CHECK(clSetKernelArg(s->loadys_krnl, 0, sizeof(cl_mem), &y_cl));
CL_CHECK(clSetKernelArg(s->loadys_krnl, 1, sizeof(cl_mem), &out_cl));
@@ -72,3 +63,14 @@ void loadyuv_queue(LoadYUVState* s, cl_command_queue q,
CL_CHECK(clEnqueueNDRangeKernel(q, s->loaduv_krnl, 1, NULL,
&loaduv_work_size, NULL, 0, 0, NULL));
}
void copy_queue(LoadYUVState* s, cl_command_queue q, cl_mem src, cl_mem dst,
size_t src_offset, size_t dst_offset, size_t size) {
CL_CHECK(clSetKernelArg(s->copy_krnl, 0, sizeof(cl_mem), &src));
CL_CHECK(clSetKernelArg(s->copy_krnl, 1, sizeof(cl_mem), &dst));
CL_CHECK(clSetKernelArg(s->copy_krnl, 2, sizeof(cl_int), &src_offset));
CL_CHECK(clSetKernelArg(s->copy_krnl, 3, sizeof(cl_int), &dst_offset));
const size_t copy_work_size = size/8;
CL_CHECK(clEnqueueNDRangeKernel(q, s->copy_krnl, 1, NULL,
&copy_work_size, NULL, 0, 0, NULL));
}
selfdrive/modeld/transforms/loadyuv.cl
+5 -3
View File
@@ -37,9 +37,11 @@ __kernel void loaduv(__global uchar8 const * const in,
out[gid + out_offset / 8] = inv;
}
__kernel void copy(__global uchar8 * inout,
int in_offset)
__kernel void copy(__global uchar8 * in,
__global uchar8 * out,
int in_offset,
int out_offset)
{
const int gid = get_global_id(0);
inout[gid] = inout[gid + in_offset / 8];
out[gid + out_offset / 8] = in[gid + in_offset / 8];
}
selfdrive/modeld/transforms/loadyuv.h
+5 -1
View File
@@ -13,4 +13,8 @@ void loadyuv_destroy(LoadYUVState* s);
void loadyuv_queue(LoadYUVState* s, cl_command_queue q,
cl_mem y_cl, cl_mem u_cl, cl_mem v_cl,
cl_mem out_cl, bool do_shift = false);
cl_mem out_cl);
void copy_queue(LoadYUVState* s, cl_command_queue q, cl_mem src, cl_mem dst,
size_t src_offset, size_t dst_offset, size_t size);
selfdrive/test/process_replay/model_replay.py
+3
View File
@@ -129,6 +129,9 @@ if __name__ == "__main__":
for i in range(4):
for field in ('x', 'y', 'z', 't'):
ignore.append(f'modelV2.laneLines.{i}.{field}')
for i in range(2):
for field in ('x', 'y', 'z', 't'):
ignore.append(f'modelV2.roadEdges.{i}.{field}')
tolerance = .3 if PC else None
results: Any = {TEST_ROUTE: {}}
log_paths: Any = {TEST_ROUTE: {"models": {'ref': BASE_URL + log_fn, 'new': log_fn}}}
selfdrive/test/process_replay/model_replay_ref_commit
+1 -1
View File
@@ -1 +1 @@
ff10a803005f5116a8f3d88629503bc52b15b93f
76b8121429ad69d6548f744a39680ade324947e2
selfdrive/test/test_onroad.py
+4 -4
View File
@@ -36,19 +36,19 @@ CPU usage budget
MAX_TOTAL_CPU = 265. # total for all 8 cores
PROCS = {
# Baseline CPU usage by process
"selfdrive.controls.controlsd": 18.0,
"selfdrive.selfdrived.selfdrived": 21.0,
"selfdrive.controls.controlsd": 16.0,
"selfdrive.selfdrived.selfdrived": 16.0,
"selfdrive.car.card": 30.0,
"./loggerd": 14.0,
"./encoderd": 17.0,
"./camerad": 14.5,
"selfdrive.controls.plannerd": 11.0,
"selfdrive.controls.plannerd": 9.0,
"./ui": 18.0,
"selfdrive.locationd.paramsd": 9.0,
"./sensord": 7.0,
"selfdrive.controls.radard": 2.0,
"selfdrive.modeld.modeld": 17.0,
"selfdrive.modeld.dmonitoringmodeld": 8.0,
"selfdrive.modeld.dmonitoringmodeld": 11.0,
"system.hardware.hardwared": 3.87,
"selfdrive.locationd.calibrationd": 2.0,
"selfdrive.locationd.torqued": 5.0,
selfdrive/ui/qt/offroad/settings.cc
+1 -1
View File
@@ -22,7 +22,7 @@ TogglesPanel::TogglesPanel(SettingsWindow *parent) : ListWidget(parent) {
"OpenpilotEnabledToggle",
tr("Enable openpilot"),
tr("Use the openpilot system for adaptive cruise control and lane keep driver assistance. Your attention is required at all times to use this feature. Changing this setting takes effect when the car is powered off."),
"../assets/offroad/icon_openpilot.png",
"../assets/img_chffr_wheel.png",
},
{
"ExperimentalLongitudinalEnabled",
selfdrive/ui/qt/widgets/prime.cc
+7 -7
View File
@@ -180,20 +180,20 @@ SetupWidget::SetupWidget(QWidget* parent) : QFrame(parent) {
QFrame* finishRegistration = new QFrame;
finishRegistration->setObjectName("primeWidget");
QVBoxLayout* finishRegistationLayout = new QVBoxLayout(finishRegistration);
finishRegistationLayout->setSpacing(38);
finishRegistationLayout->setContentsMargins(64, 48, 64, 48);
QVBoxLayout* finishRegistrationLayout = new QVBoxLayout(finishRegistration);
finishRegistrationLayout->setSpacing(38);
finishRegistrationLayout->setContentsMargins(64, 48, 64, 48);
QLabel* registrationTitle = new QLabel(tr("Finish Setup"));
registrationTitle->setStyleSheet("font-size: 75px; font-weight: bold;");
finishRegistationLayout->addWidget(registrationTitle);
finishRegistrationLayout->addWidget(registrationTitle);
QLabel* registrationDescription = new QLabel(tr("Pair your device with comma connect (connect.comma.ai) and claim your comma prime offer."));
registrationDescription->setWordWrap(true);
registrationDescription->setStyleSheet("font-size: 50px; font-weight: light;");
finishRegistationLayout->addWidget(registrationDescription);
finishRegistrationLayout->addWidget(registrationDescription);
finishRegistationLayout->addStretch();
finishRegistrationLayout->addStretch();
QPushButton* pair = new QPushButton(tr("Pair device"));
pair->setStyleSheet(R"(
@@ -208,7 +208,7 @@ SetupWidget::SetupWidget(QWidget* parent) : QFrame(parent) {
background-color: #3049F4;
}
)");
finishRegistationLayout->addWidget(pair);
finishRegistrationLayout->addWidget(pair);
popup = new PairingPopup(this);
QObject::connect(pair, &QPushButton::clicked, popup, &PairingPopup::exec);
system/camerad/cameras/camera_qcom2.cc
+4 -3
View File
@@ -178,9 +178,10 @@ void CameraState::set_camera_exposure(float grey_frac) {
new_exp_t = exposure_time;
enable_dc_gain = false;
} else {
// Simple brute force optimizer to choose sensor parameters
// to reach desired EV
for (int g = std::max((int)sensor->analog_gain_min_idx, gain_idx - 1); g <= std::min((int)sensor->analog_gain_max_idx, gain_idx + 1); g++) {
// Simple brute force optimizer to choose sensor parameters to reach desired EV
int min_g = std::max(gain_idx - 1, sensor->analog_gain_min_idx);
int max_g = std::min(gain_idx + 1, sensor->analog_gain_max_idx);
for (int g = min_g; g <= max_g; g++) {
float gain = sensor->sensor_analog_gains[g] * get_gain_factor();
// Compute optimal time for given gain