diff --git a/README.md b/README.md index b0959eee5d..77002481d3 100644 --- a/README.md +++ b/README.md @@ -53,7 +53,7 @@ We have detailed instructions for [how to install the harness and device in a ca ### Branches | branch | URL | description | |------------------|----------------------------------------|-------------------------------------------------------------------------------------| -| `release3` | openilot.comma.ai | This is openpilot's release branch. | +| `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. | diff --git a/msgq_repo b/msgq_repo index 434ed2312c..5bb86f8bc7 160000 --- a/msgq_repo +++ b/msgq_repo @@ -1 +1 @@ -Subproject commit 434ed2312c980b5504a4a75001d04c3ecbddf93f +Subproject commit 5bb86f8bc7434048ab2d5ce0243a97d9848b34de diff --git a/panda b/panda index 45301bf15c..0d4b79a3c7 160000 --- a/panda +++ b/panda @@ -1 +1 @@ -Subproject commit 45301bf15c8236e3ae76cfe6ff6938540a1bafd5 +Subproject commit 0d4b79a3c7b9c1012832ec414989e56e3a2a8d46 diff --git a/pyproject.toml b/pyproject.toml index 2da580b6fa..9cd28e6baf 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -42,8 +42,7 @@ dependencies = [ # modeld "onnx >= 1.14.0", - "onnxruntime >=1.16.3; platform_system == 'Linux' and platform_machine == 'aarch64'", - "onnxruntime-gpu >=1.16.3; platform_system == 'Linux' and platform_machine == 'x86_64'", + "onnxruntime >=1.16.3", # logging "pyzmq", diff --git a/release/release_files.py b/release/release_files.py index 52974ba711..0e1ed852a2 100755 --- a/release/release_files.py +++ b/release/release_files.py @@ -54,7 +54,7 @@ whitelist = [ "tools/joystick/", "tools/longitudinal_maneuvers/", - "tinygrad_repo/openpilot/compile2.py", + "tinygrad_repo/examples/openpilot/compile3.py", "tinygrad_repo/extra/onnx.py", "tinygrad_repo/extra/onnx_ops.py", "tinygrad_repo/extra/thneed.py", diff --git a/selfdrive/controls/lib/longitudinal_planner.py b/selfdrive/controls/lib/longitudinal_planner.py index f1637d960c..eba8019117 100755 --- a/selfdrive/controls/lib/longitudinal_planner.py +++ b/selfdrive/controls/lib/longitudinal_planner.py @@ -50,24 +50,20 @@ def limit_accel_in_turns(v_ego, angle_steers, a_target, CP): return [a_target[0], min(a_target[1], a_x_allowed)] -def get_accel_from_plan(CP, speeds, accels): +def get_accel_from_plan(speeds, accels, action_t=DT_MDL, vEgoStopping=0.05): if len(speeds) == CONTROL_N: - v_target_now = interp(DT_MDL, CONTROL_N_T_IDX, speeds) - a_target_now = interp(DT_MDL, CONTROL_N_T_IDX, accels) + v_now = speeds[0] + a_now = accels[0] - v_target = interp(CP.longitudinalActuatorDelay + DT_MDL, CONTROL_N_T_IDX, speeds) - if v_target != v_target_now: - a_target = 2 * (v_target - v_target_now) / CP.longitudinalActuatorDelay - a_target_now - else: - a_target = a_target_now - - v_target_1sec = interp(CP.longitudinalActuatorDelay + DT_MDL + 1.0, CONTROL_N_T_IDX, speeds) + v_target = interp(action_t, CONTROL_N_T_IDX, speeds) + a_target = 2 * (v_target - v_now) / (action_t) - a_now + v_target_1sec = interp(action_t + 1.0, CONTROL_N_T_IDX, speeds) else: v_target = 0.0 v_target_1sec = 0.0 a_target = 0.0 - should_stop = (v_target < CP.vEgoStopping and - v_target_1sec < CP.vEgoStopping) + should_stop = (v_target < vEgoStopping and + v_target_1sec < vEgoStopping) return a_target, should_stop @@ -201,7 +197,9 @@ class LongitudinalPlanner: longitudinalPlan.longitudinalPlanSource = self.mpc.source longitudinalPlan.fcw = self.fcw - a_target, should_stop = get_accel_from_plan(self.CP, longitudinalPlan.speeds, longitudinalPlan.accels) + action_t = self.CP.longitudinalActuatorDelay + DT_MDL + a_target, should_stop = get_accel_from_plan(longitudinalPlan.speeds, longitudinalPlan.accels, + action_t=action_t, vEgoStopping=self.CP.vEgoStopping) longitudinalPlan.aTarget = a_target longitudinalPlan.shouldStop = should_stop longitudinalPlan.allowBrake = True diff --git a/selfdrive/debug/touch_replay.py b/selfdrive/debug/touch_replay.py index c397520918..1e1596d264 100755 --- a/selfdrive/debug/touch_replay.py +++ b/selfdrive/debug/touch_replay.py @@ -38,11 +38,17 @@ if __name__ == '__main__': 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() diff --git a/selfdrive/modeld/SConscript b/selfdrive/modeld/SConscript index d472998416..ef0fd52f33 100644 --- a/selfdrive/modeld/SConscript +++ b/selfdrive/modeld/SConscript @@ -13,20 +13,6 @@ common_src = [ "transforms/transform.cc", ] -thneed_src_common = [ - "thneed/thneed_common.cc", - "thneed/serialize.cc", -] - -thneed_src_qcom = thneed_src_common + ["thneed/thneed_qcom2.cc"] -thneed_src_pc = thneed_src_common + ["thneed/thneed_pc.cc"] -thneed_src = thneed_src_qcom if arch == "larch64" else thneed_src_pc - -# SNPE except on Mac and ARM Linux -snpe_lib = [] -if arch != "Darwin" and arch != "aarch64": - common_src += ['runners/snpemodel.cc'] - snpe_lib += ['SNPE'] # OpenCL is a framework on Mac if arch == "Darwin": @@ -45,34 +31,24 @@ snpe_rpath_pc = f"{Dir('#').abspath}/third_party/snpe/x86_64-linux-clang" snpe_rpath = lenvCython['RPATH'] + [snpe_rpath_qcom if arch == "larch64" else snpe_rpath_pc] cython_libs = envCython["LIBS"] + libs -snpemodel_lib = lenv.Library('snpemodel', ['runners/snpemodel.cc']) commonmodel_lib = lenv.Library('commonmodel', common_src) - -lenvCython.Program('runners/runmodel_pyx.so', 'runners/runmodel_pyx.pyx', LIBS=cython_libs, FRAMEWORKS=frameworks) -lenvCython.Program('runners/snpemodel_pyx.so', 'runners/snpemodel_pyx.pyx', LIBS=[snpemodel_lib, snpe_lib, *cython_libs], FRAMEWORKS=frameworks, RPATH=snpe_rpath) lenvCython.Program('models/commonmodel_pyx.so', 'models/commonmodel_pyx.pyx', LIBS=[commonmodel_lib, *cython_libs], FRAMEWORKS=frameworks) - -tinygrad_files = ["#"+x for x in glob.glob(env.Dir("#tinygrad_repo").relpath + "/**", recursive=True, root_dir=env.Dir("#").abspath)] +tinygrad_files = ["#"+x for x in glob.glob(env.Dir("#tinygrad_repo").relpath + "/**", recursive=True, root_dir=env.Dir("#").abspath) if 'pycache' not in x] # Get model metadata fn = File("models/supercombo").abspath cmd = f'python3 {Dir("#selfdrive/modeld").abspath}/get_model_metadata.py {fn}.onnx' lenv.Command(fn + "_metadata.pkl", [fn + ".onnx"] + tinygrad_files, cmd) -# Build thneed model -if arch == "larch64" or GetOption('pc_thneed'): - tinygrad_opts = [] - if not GetOption('pc_thneed'): - # use FLOAT16 on device for speed + don't cache the CL kernels for space - tinygrad_opts += ["FLOAT16=1", "PYOPENCL_NO_CACHE=1"] - cmd = f"cd {Dir('#').abspath}/tinygrad_repo && " + ' '.join(tinygrad_opts) + f" python3 openpilot/compile2.py {fn}.onnx {fn}.thneed" +# Compile tinygrad model +pythonpath_string = 'PYTHONPATH="${PYTHONPATH}:' + env.Dir("#tinygrad_repo").abspath + '"' +if arch == 'larch64': + device_string = 'QCOM=1' +else: + device_string = 'CLANG=1 IMAGE=0' - lenv.Command(fn + ".thneed", [fn + ".onnx"] + tinygrad_files, cmd) +for model_name in ['supercombo', 'dmonitoring_model']: + fn = File(f"models/{model_name}").abspath + cmd = f'{pythonpath_string} {device_string} python3 {Dir("#tinygrad_repo").abspath}/examples/openpilot/compile3.py {fn}.onnx {fn}_tinygrad.pkl' + lenv.Command(fn + "_tinygrad.pkl", [fn + ".onnx"] + tinygrad_files, cmd) - fn_dm = File("models/dmonitoring_model").abspath - cmd = f"cd {Dir('#').abspath}/tinygrad_repo && " + ' '.join(tinygrad_opts) + f" python3 openpilot/compile2.py {fn_dm}.onnx {fn_dm}.thneed" - lenv.Command(fn_dm + ".thneed", [fn_dm + ".onnx"] + tinygrad_files, cmd) - - thneed_lib = env.SharedLibrary('thneed', thneed_src, LIBS=[gpucommon, common, 'OpenCL', 'dl']) - thneedmodel_lib = env.Library('thneedmodel', ['runners/thneedmodel.cc']) - lenvCython.Program('runners/thneedmodel_pyx.so', 'runners/thneedmodel_pyx.pyx', LIBS=envCython["LIBS"]+[thneedmodel_lib, thneed_lib, gpucommon, common, 'dl', 'OpenCL']) diff --git a/selfdrive/modeld/dmonitoringmodeld b/selfdrive/modeld/dmonitoringmodeld index 80157e1751..90b43800fe 100755 --- a/selfdrive/modeld/dmonitoringmodeld +++ b/selfdrive/modeld/dmonitoringmodeld @@ -1,10 +1,4 @@ #!/usr/bin/env bash DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" >/dev/null && pwd)" -cd "$DIR/../../" - -if [ -f "$DIR/libthneed.so" ]; then - export LD_PRELOAD="$DIR/libthneed.so" -fi - exec "$DIR/dmonitoringmodeld.py" "$@" diff --git a/selfdrive/modeld/dmonitoringmodeld.py b/selfdrive/modeld/dmonitoringmodeld.py index a29dc081e1..ebc24ad024 100755 --- a/selfdrive/modeld/dmonitoringmodeld.py +++ b/selfdrive/modeld/dmonitoringmodeld.py @@ -1,8 +1,17 @@ #!/usr/bin/env python3 import os +from openpilot.system.hardware import TICI +if TICI: + from tinygrad.tensor import Tensor + from tinygrad.dtype import dtypes + from openpilot.selfdrive.modeld.runners.tinygrad_helpers import qcom_tensor_from_opencl_address + os.environ['QCOM'] = '1' +else: + from openpilot.selfdrive.modeld.runners.ort_helpers import make_onnx_cpu_runner import gc import math import time +import pickle import ctypes import numpy as np from pathlib import Path @@ -13,21 +22,20 @@ from cereal.messaging import PubMaster, SubMaster from msgq.visionipc import VisionIpcClient, VisionStreamType, VisionBuf from openpilot.common.swaglog import cloudlog from openpilot.common.realtime import set_realtime_priority -from openpilot.common.transformations.model import dmonitoringmodel_intrinsics +from openpilot.common.transformations.model import dmonitoringmodel_intrinsics, DM_INPUT_SIZE from openpilot.common.transformations.camera import _ar_ox_fisheye, _os_fisheye from openpilot.selfdrive.modeld.models.commonmodel_pyx import CLContext, MonitoringModelFrame -from openpilot.selfdrive.modeld.runners import ModelRunner, Runtime from openpilot.selfdrive.modeld.parse_model_outputs import sigmoid +MODEL_WIDTH, MODEL_HEIGHT = DM_INPUT_SIZE CALIB_LEN = 3 FEATURE_LEN = 512 OUTPUT_SIZE = 84 + FEATURE_LEN PROCESS_NAME = "selfdrive.modeld.dmonitoringmodeld" SEND_RAW_PRED = os.getenv('SEND_RAW_PRED') -MODEL_PATHS = { - ModelRunner.THNEED: Path(__file__).parent / 'models/dmonitoring_model.thneed', - ModelRunner.ONNX: Path(__file__).parent / 'models/dmonitoring_model.onnx'} +MODEL_PATH = Path(__file__).parent / 'models/dmonitoring_model.onnx' +MODEL_PKL_PATH = Path(__file__).parent / 'models/dmonitoring_model_tinygrad.pkl' class DriverStateResult(ctypes.Structure): _fields_ = [ @@ -58,29 +66,42 @@ class DMonitoringModelResult(ctypes.Structure): class ModelState: inputs: dict[str, np.ndarray] output: np.ndarray - model: ModelRunner def __init__(self, cl_ctx): assert ctypes.sizeof(DMonitoringModelResult) == OUTPUT_SIZE * ctypes.sizeof(ctypes.c_float) self.frame = MonitoringModelFrame(cl_ctx) - self.output = np.zeros(OUTPUT_SIZE, dtype=np.float32) - self.inputs = { - 'calib': np.zeros(CALIB_LEN, dtype=np.float32)} + self.numpy_inputs = { + 'calib': np.zeros((1, CALIB_LEN), dtype=np.float32), + } - self.model = ModelRunner(MODEL_PATHS, self.output, Runtime.GPU, False, cl_ctx) - self.model.addInput("input_img", None) - self.model.addInput("calib", self.inputs['calib']) + if TICI: + self.tensor_inputs = {k: Tensor(v, device='NPY').realize() for k,v in self.numpy_inputs.items()} + with open(MODEL_PKL_PATH, "rb") as f: + self.model_run = pickle.load(f) + else: + self.onnx_cpu_runner = make_onnx_cpu_runner(MODEL_PATH) def run(self, buf:VisionBuf, calib:np.ndarray, transform:np.ndarray) -> tuple[np.ndarray, float]: - self.inputs['calib'][:] = calib - - self.model.setInputBuffer("input_img", self.frame.prepare(buf, transform.flatten(), None).view(np.float32)) + self.numpy_inputs['calib'][0,:] = calib t1 = time.perf_counter() - self.model.execute() + + input_img_cl = self.frame.prepare(buf, transform.flatten()) + if TICI: + # The imgs tensors are backed by opencl memory, only need init once + if 'input_img' not in self.tensor_inputs: + self.tensor_inputs['input_img'] = qcom_tensor_from_opencl_address(input_img_cl.mem_address, (1, MODEL_WIDTH*MODEL_HEIGHT), dtype=dtypes.uint8) + else: + self.numpy_inputs['input_img'] = self.frame.buffer_from_cl(input_img_cl).reshape((1, MODEL_WIDTH*MODEL_HEIGHT)) + + if TICI: + output = self.model_run(**self.tensor_inputs).numpy().flatten() + else: + output = self.onnx_cpu_runner.run(None, self.numpy_inputs)[0].flatten() + t2 = time.perf_counter() - return self.output, t2 - t1 + return output, t2 - t1 def fill_driver_state(msg, ds_result: DriverStateResult): diff --git a/selfdrive/modeld/modeld b/selfdrive/modeld/modeld index e1cef4dcc3..5ba4688554 100755 --- a/selfdrive/modeld/modeld +++ b/selfdrive/modeld/modeld @@ -1,10 +1,4 @@ #!/usr/bin/env bash DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" >/dev/null && pwd)" -cd "$DIR/../../" - -if [ -f "$DIR/libthneed.so" ]; then - export LD_PRELOAD="$DIR/libthneed.so" -fi - exec "$DIR/modeld.py" "$@" diff --git a/selfdrive/modeld/modeld.py b/selfdrive/modeld/modeld.py index e0f78c5aa0..49f8c483f0 100755 --- a/selfdrive/modeld/modeld.py +++ b/selfdrive/modeld/modeld.py @@ -1,5 +1,15 @@ #!/usr/bin/env python3 import os +from openpilot.system.hardware import TICI + +# +if TICI: + from tinygrad.tensor import Tensor + from tinygrad.dtype import dtypes + from openpilot.selfdrive.modeld.runners.tinygrad_helpers import qcom_tensor_from_opencl_address + os.environ['QCOM'] = '1' +else: + from openpilot.selfdrive.modeld.runners.ort_helpers import make_onnx_cpu_runner import time import pickle import numpy as np @@ -18,22 +28,19 @@ from openpilot.common.transformations.camera import DEVICE_CAMERAS from openpilot.common.transformations.model import get_warp_matrix from openpilot.system import sentry from openpilot.selfdrive.controls.lib.desire_helper import DesireHelper -from openpilot.selfdrive.modeld.runners import ModelRunner, Runtime from openpilot.selfdrive.modeld.parse_model_outputs import Parser from openpilot.selfdrive.modeld.fill_model_msg import fill_model_msg, fill_pose_msg, PublishState from openpilot.selfdrive.modeld.constants import ModelConstants from openpilot.selfdrive.modeld.models.commonmodel_pyx import DrivingModelFrame, CLContext + PROCESS_NAME = "selfdrive.modeld.modeld" SEND_RAW_PRED = os.getenv('SEND_RAW_PRED') -MODEL_PATHS = { - ModelRunner.THNEED: Path(__file__).parent / 'models/supercombo.thneed', - ModelRunner.ONNX: Path(__file__).parent / 'models/supercombo.onnx'} - +MODEL_PATH = Path(__file__).parent / 'models/supercombo.onnx' +MODEL_PKL_PATH = Path(__file__).parent / 'models/supercombo_tinygrad.pkl' METADATA_PATH = Path(__file__).parent / 'models/supercombo_metadata.pkl' - class FrameMeta: frame_id: int = 0 timestamp_sof: int = 0 @@ -44,40 +51,39 @@ class FrameMeta: self.frame_id, self.timestamp_sof, self.timestamp_eof = vipc.frame_id, vipc.timestamp_sof, vipc.timestamp_eof class ModelState: - frame: DrivingModelFrame - wide_frame: DrivingModelFrame + frames: dict[str, DrivingModelFrame] inputs: dict[str, np.ndarray] output: np.ndarray prev_desire: np.ndarray # for tracking the rising edge of the pulse - model: ModelRunner def __init__(self, context: CLContext): - self.frame = DrivingModelFrame(context) - self.wide_frame = DrivingModelFrame(context) + self.frames = {'input_imgs': DrivingModelFrame(context), 'big_input_imgs': 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) # 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), - 'features_buffer': np.zeros(ModelConstants.HISTORY_BUFFER_LEN * ModelConstants.FEATURE_LEN, dtype=np.float32), + self.numpy_inputs = { + 'desire': np.zeros((1, (ModelConstants.HISTORY_BUFFER_LEN+1), ModelConstants.DESIRE_LEN), dtype=np.float32), + 'traffic_convention': np.zeros((1, ModelConstants.TRAFFIC_CONVENTION_LEN), dtype=np.float32), + 'features_buffer': np.zeros((1, ModelConstants.HISTORY_BUFFER_LEN, ModelConstants.FEATURE_LEN), dtype=np.float32), } with open(METADATA_PATH, 'rb') as f: model_metadata = pickle.load(f) + self.input_shapes = model_metadata['input_shapes'] self.output_slices = model_metadata['output_slices'] net_output_size = model_metadata['output_shapes']['outputs'][1] self.output = np.zeros(net_output_size, dtype=np.float32) self.parser = Parser() - self.model = ModelRunner(MODEL_PATHS, self.output, Runtime.GPU, False, context) - self.model.addInput("input_imgs", None) - self.model.addInput("big_input_imgs", None) - for k,v in self.inputs.items(): - self.model.addInput(k, v) + if TICI: + self.tensor_inputs = {k: Tensor(v, device='NPY').realize() for k,v in self.numpy_inputs.items()} + with open(MODEL_PKL_PATH, "rb") as f: + self.model_run = pickle.load(f) + else: + self.onnx_cpu_runner = make_onnx_cpu_runner(MODEL_PATH) def 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()} @@ -94,24 +100,36 @@ 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.numpy_inputs['desire'][:] = self.desire_20Hz.reshape((1,25,4,-1)).max(axis=2) - self.inputs['traffic_convention'][:] = inputs['traffic_convention'] + self.numpy_inputs['traffic_convention'][:] = inputs['traffic_convention'] + imgs_cl = {'input_imgs': self.frames['input_imgs'].prepare(buf, transform.flatten()), + 'big_input_imgs': self.frames['big_input_imgs'].prepare(wbuf, transform_wide.flatten())} - 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"))) + if TICI: + # The imgs tensors are backed by opencl memory, only need init once + for key in imgs_cl: + if key not in self.tensor_inputs: + self.tensor_inputs[key] = qcom_tensor_from_opencl_address(imgs_cl[key].mem_address, self.input_shapes[key], dtype=dtypes.uint8) + else: + for key in imgs_cl: + self.numpy_inputs[key] = self.frames[key].buffer_from_cl(imgs_cl[key]).reshape(self.input_shapes[key]) if prepare_only: return None - self.model.execute() + if TICI: + self.output = self.model_run(**self.tensor_inputs).numpy().flatten() + else: + self.output = self.onnx_cpu_runner.run(None, self.numpy_inputs)[0].flatten() + outputs = self.parser.parse_outputs(self.slice_outputs(self.output)) self.full_features_20Hz[:-1] = self.full_features_20Hz[1:] self.full_features_20Hz[-1] = outputs['hidden_state'][0, :] idxs = np.arange(-4,-100,-4)[::-1] - self.inputs['features_buffer'][:] = self.full_features_20Hz[idxs].flatten() + self.numpy_inputs['features_buffer'][:] = self.full_features_20Hz[idxs] return outputs @@ -281,7 +299,6 @@ def main(demo=False): pm.send('modelV2', modelv2_send) pm.send('drivingModelData', drivingdata_send) pm.send('cameraOdometry', posenet_send) - last_vipc_frame_id = meta_main.frame_id diff --git a/selfdrive/modeld/models/commonmodel.cc b/selfdrive/modeld/models/commonmodel.cc index cd77903680..ad2620c7b4 100644 --- a/selfdrive/modeld/models/commonmodel.cc +++ b/selfdrive/modeld/models/commonmodel.cc @@ -7,7 +7,7 @@ DrivingModelFrame::DrivingModelFrame(cl_device_id device_id, cl_context context) : ModelFrame(device_id, context) { input_frames = std::make_unique(buf_size); - //input_frames_cl = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_WRITE, buf_size, NULL, &err)); + 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; @@ -17,7 +17,7 @@ DrivingModelFrame::DrivingModelFrame(cl_device_id device_id, cl_context context) init_transform(device_id, context, MODEL_WIDTH, MODEL_HEIGHT); } -uint8_t* DrivingModelFrame::prepare(cl_mem yuv_cl, int frame_width, int frame_height, int frame_stride, int frame_uv_offset, const mat3& projection, cl_mem* output) { +cl_mem* DrivingModelFrame::prepare(cl_mem yuv_cl, int frame_width, int frame_height, int frame_stride, int frame_uv_offset, const mat3& projection) { run_transform(yuv_cl, MODEL_WIDTH, MODEL_HEIGHT, frame_width, frame_height, frame_stride, frame_uv_offset, projection); for (int i = 0; i < 4; i++) { @@ -25,19 +25,12 @@ uint8_t* DrivingModelFrame::prepare(cl_mem yuv_cl, int frame_width, int frame_he } loadyuv_queue(&loadyuv, q, y_cl, u_cl, v_cl, last_img_cl); - if (output == NULL) { - CL_CHECK(clEnqueueReadBuffer(q, img_buffer_20hz_cl, CL_TRUE, 0, frame_size_bytes, &input_frames[0], 0, nullptr, nullptr)); - CL_CHECK(clEnqueueReadBuffer(q, last_img_cl, CL_TRUE, 0, frame_size_bytes, &input_frames[MODEL_FRAME_SIZE], 0, nullptr, nullptr)); - clFinish(q); - return &input_frames[0]; - } else { - copy_queue(&loadyuv, q, img_buffer_20hz_cl, *output, 0, 0, frame_size_bytes); - copy_queue(&loadyuv, q, last_img_cl, *output, 0, frame_size_bytes, frame_size_bytes); + copy_queue(&loadyuv, q, img_buffer_20hz_cl, input_frames_cl, 0, 0, frame_size_bytes); + copy_queue(&loadyuv, q, last_img_cl, input_frames_cl, 0, frame_size_bytes, frame_size_bytes); - // NOTE: Since thneed is using a different command queue, this clFinish is needed to ensure the image is ready. - clFinish(q); - return NULL; - } + // NOTE: Since thneed is using a different command queue, this clFinish is needed to ensure the image is ready. + clFinish(q); + return &input_frames_cl; } DrivingModelFrame::~DrivingModelFrame() { @@ -51,16 +44,15 @@ DrivingModelFrame::~DrivingModelFrame() { MonitoringModelFrame::MonitoringModelFrame(cl_device_id device_id, cl_context context) : ModelFrame(device_id, context) { input_frames = std::make_unique(buf_size); - //input_frame_cl = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_WRITE, buf_size, NULL, &err)); + 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); } -uint8_t* MonitoringModelFrame::prepare(cl_mem yuv_cl, int frame_width, int frame_height, int frame_stride, int frame_uv_offset, const mat3& projection, cl_mem* output) { + +cl_mem* MonitoringModelFrame::prepare(cl_mem yuv_cl, int frame_width, int frame_height, int frame_stride, int frame_uv_offset, const mat3& projection) { run_transform(yuv_cl, MODEL_WIDTH, MODEL_HEIGHT, frame_width, frame_height, frame_stride, frame_uv_offset, projection); - CL_CHECK(clEnqueueReadBuffer(q, y_cl, CL_TRUE, 0, MODEL_FRAME_SIZE * sizeof(uint8_t), input_frames.get(), 0, nullptr, nullptr)); clFinish(q); - //return &y_cl; - return input_frames.get(); + return &y_cl; } MonitoringModelFrame::~MonitoringModelFrame() { diff --git a/selfdrive/modeld/models/commonmodel.h b/selfdrive/modeld/models/commonmodel.h index e55f78e248..14409943e4 100644 --- a/selfdrive/modeld/models/commonmodel.h +++ b/selfdrive/modeld/models/commonmodel.h @@ -23,14 +23,12 @@ public: q = CL_CHECK_ERR(clCreateCommandQueue(context, device_id, 0, &err)); } virtual ~ModelFrame() {} - virtual uint8_t* prepare(cl_mem yuv_cl, int frame_width, int frame_height, int frame_stride, int frame_uv_offset, const mat3& projection, cl_mem* output) { return NULL; } - /* + virtual cl_mem* prepare(cl_mem yuv_cl, int frame_width, int frame_height, int frame_stride, int frame_uv_offset, const mat3& projection) { return NULL; } uint8_t* buffer_from_cl(cl_mem *in_frames, int buffer_size) { CL_CHECK(clEnqueueReadBuffer(q, *in_frames, CL_TRUE, 0, buffer_size, input_frames.get(), 0, nullptr, nullptr)); clFinish(q); return &input_frames[0]; } - */ int MODEL_WIDTH; int MODEL_HEIGHT; @@ -68,7 +66,7 @@ class DrivingModelFrame : public ModelFrame { public: DrivingModelFrame(cl_device_id device_id, cl_context context); ~DrivingModelFrame(); - uint8_t* prepare(cl_mem yuv_cl, int frame_width, int frame_height, int frame_stride, int frame_uv_offset, const mat3& projection, cl_mem* output); + cl_mem* prepare(cl_mem yuv_cl, int frame_width, int frame_height, int frame_stride, int frame_uv_offset, const mat3& projection); const int MODEL_WIDTH = 512; const int MODEL_HEIGHT = 256; @@ -78,7 +76,7 @@ public: private: LoadYUVState loadyuv; - cl_mem img_buffer_20hz_cl, last_img_cl;//, input_frames_cl; + cl_mem img_buffer_20hz_cl, last_img_cl, input_frames_cl; cl_buffer_region region; }; @@ -86,7 +84,7 @@ class MonitoringModelFrame : public ModelFrame { public: MonitoringModelFrame(cl_device_id device_id, cl_context context); ~MonitoringModelFrame(); - uint8_t* prepare(cl_mem yuv_cl, int frame_width, int frame_height, int frame_stride, int frame_uv_offset, const mat3& projection, cl_mem* output); + cl_mem* prepare(cl_mem yuv_cl, int frame_width, int frame_height, int frame_stride, int frame_uv_offset, const mat3& projection); const int MODEL_WIDTH = 1440; const int MODEL_HEIGHT = 960; @@ -94,5 +92,5 @@ public: const int buf_size = MODEL_FRAME_SIZE; private: - // cl_mem input_frame_cl; + cl_mem input_frame_cl; }; diff --git a/selfdrive/modeld/models/commonmodel.pxd b/selfdrive/modeld/models/commonmodel.pxd index 61acb2bb2f..d2a8fb4dcd 100644 --- a/selfdrive/modeld/models/commonmodel.pxd +++ b/selfdrive/modeld/models/commonmodel.pxd @@ -14,8 +14,8 @@ cdef extern from "common/clutil.h": cdef extern from "selfdrive/modeld/models/commonmodel.h": cppclass ModelFrame: int buf_size - # unsigned char * buffer_from_cl(cl_mem*, int); - unsigned char * prepare(cl_mem, int, int, int, int, mat3, cl_mem*) + unsigned char * buffer_from_cl(cl_mem*, int); + cl_mem * prepare(cl_mem, int, int, int, int, mat3) cppclass DrivingModelFrame: int buf_size diff --git a/selfdrive/modeld/models/commonmodel_pyx.pyx b/selfdrive/modeld/models/commonmodel_pyx.pyx index c7ea5ecac2..b754086544 100644 --- a/selfdrive/modeld/models/commonmodel_pyx.pyx +++ b/selfdrive/modeld/models/commonmodel_pyx.pyx @@ -39,24 +39,17 @@ cdef class ModelFrame: def __dealloc__(self): del self.frame - def prepare(self, VisionBuf buf, float[:] projection, CLMem output): + def prepare(self, VisionBuf buf, float[:] projection): cdef mat3 cprojection memcpy(cprojection.v, &projection[0], 9*sizeof(float)) - cdef unsigned char * data - if output is None: - data = self.frame.prepare(buf.buf.buf_cl, buf.width, buf.height, buf.stride, buf.uv_offset, cprojection, NULL) - else: - data = self.frame.prepare(buf.buf.buf_cl, buf.width, buf.height, buf.stride, buf.uv_offset, cprojection, output.mem) - if not data: - return None + cdef cl_mem * data + data = self.frame.prepare(buf.buf.buf_cl, buf.width, buf.height, buf.stride, buf.uv_offset, cprojection) + return CLMem.create(data) - return np.asarray( 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( data2) + def buffer_from_cl(self, CLMem in_frames): + cdef unsigned char * data2 + data2 = self.frame.buffer_from_cl(in_frames.mem, self.buf_size) + return np.asarray( data2) cdef class DrivingModelFrame(ModelFrame): @@ -74,3 +67,4 @@ cdef class MonitoringModelFrame(ModelFrame): self._frame = new cppMonitoringModelFrame(context.device_id, context.context) self.frame = (self._frame) self.buf_size = self._frame.buf_size + diff --git a/selfdrive/modeld/models/supercombo.onnx b/selfdrive/modeld/models/supercombo.onnx index 384072f426..777aa8aa7e 100644 --- a/selfdrive/modeld/models/supercombo.onnx +++ b/selfdrive/modeld/models/supercombo.onnx @@ -1,3 +1,3 @@ version https://git-lfs.github.com/spec/v1 -oid sha256:0c896681fd6851de3968433e12f37834429eba265e938cf383200be3e5835cec +oid sha256:72d3d6f8d3c98f5431ec86be77b6350d7d4f43c25075c0106f1d1e7ec7c77668 size 49096168 diff --git a/selfdrive/modeld/runners/__init__.py b/selfdrive/modeld/runners/__init__.py deleted file mode 100644 index 4c29bf3f1c..0000000000 --- a/selfdrive/modeld/runners/__init__.py +++ /dev/null @@ -1,27 +0,0 @@ -import os -from openpilot.system.hardware import TICI -from openpilot.selfdrive.modeld.runners.runmodel_pyx import RunModel, Runtime -assert Runtime - -USE_THNEED = int(os.getenv('USE_THNEED', str(int(TICI)))) -USE_SNPE = int(os.getenv('USE_SNPE', str(int(TICI)))) - -class ModelRunner(RunModel): - THNEED = 'THNEED' - SNPE = 'SNPE' - ONNX = 'ONNX' - - def __new__(cls, paths, *args, **kwargs): - if ModelRunner.THNEED in paths and USE_THNEED: - from openpilot.selfdrive.modeld.runners.thneedmodel_pyx import ThneedModel as Runner - runner_type = ModelRunner.THNEED - elif ModelRunner.SNPE in paths and USE_SNPE: - from openpilot.selfdrive.modeld.runners.snpemodel_pyx import SNPEModel as Runner - runner_type = ModelRunner.SNPE - elif ModelRunner.ONNX in paths: - from openpilot.selfdrive.modeld.runners.onnxmodel import ONNXModel as Runner - runner_type = ModelRunner.ONNX - else: - raise Exception("Couldn't select a model runner, make sure to pass at least one valid model path") - - return Runner(str(paths[runner_type]), *args, **kwargs) diff --git a/selfdrive/modeld/runners/onnxmodel.py b/selfdrive/modeld/runners/onnxmodel.py deleted file mode 100644 index b2827a8619..0000000000 --- a/selfdrive/modeld/runners/onnxmodel.py +++ /dev/null @@ -1,71 +0,0 @@ -import os -import onnx -import sys -import numpy as np -from typing import Any - -from openpilot.selfdrive.modeld.runners.runmodel_pyx import RunModel -from openpilot.selfdrive.modeld.runners.ort_helpers import convert_fp16_to_fp32, ORT_TYPES_TO_NP_TYPES - - -def create_ort_session(path, fp16_to_fp32): - os.environ["OMP_NUM_THREADS"] = "4" - os.environ["OMP_WAIT_POLICY"] = "PASSIVE" - - import onnxruntime as ort - print("Onnx available providers: ", ort.get_available_providers(), file=sys.stderr) - options = ort.SessionOptions() - options.graph_optimization_level = ort.GraphOptimizationLevel.ORT_DISABLE_ALL - - provider: str | tuple[str, dict[Any, Any]] - if 'OpenVINOExecutionProvider' in ort.get_available_providers() and 'ONNXCPU' not in os.environ: - provider = 'OpenVINOExecutionProvider' - elif 'CUDAExecutionProvider' in ort.get_available_providers() and 'ONNXCPU' not in os.environ: - options.intra_op_num_threads = 2 - provider = ('CUDAExecutionProvider', {'cudnn_conv_algo_search': 'EXHAUSTIVE'}) - else: - options.intra_op_num_threads = 2 - options.execution_mode = ort.ExecutionMode.ORT_SEQUENTIAL - options.graph_optimization_level = ort.GraphOptimizationLevel.ORT_ENABLE_ALL - provider = 'CPUExecutionProvider' - - model_data = convert_fp16_to_fp32(onnx.load(path)) if fp16_to_fp32 else path - print("Onnx selected provider: ", [provider], file=sys.stderr) - ort_session = ort.InferenceSession(model_data, options, providers=[provider]) - print("Onnx using ", ort_session.get_providers(), file=sys.stderr) - return ort_session - - -class ONNXModel(RunModel): - def __init__(self, path, output, runtime, use_tf8, cl_context): - self.inputs = {} - self.output = output - - self.session = create_ort_session(path, fp16_to_fp32=True) - self.input_names = [x.name for x in self.session.get_inputs()] - self.input_shapes = {x.name: [1, *x.shape[1:]] for x in self.session.get_inputs()} - self.input_dtypes = {x.name: ORT_TYPES_TO_NP_TYPES[x.type] for x in self.session.get_inputs()} - - # run once to initialize CUDA provider - if "CUDAExecutionProvider" in self.session.get_providers(): - self.session.run(None, {k: np.zeros(self.input_shapes[k], dtype=self.input_dtypes[k]) for k in self.input_names}) - print("ready to run onnx model", self.input_shapes, file=sys.stderr) - - def addInput(self, name, buffer): - assert name in self.input_names - self.inputs[name] = buffer - - def setInputBuffer(self, name, buffer): - assert name in self.inputs - self.inputs[name] = buffer - - def getCLBuffer(self, name): - return None - - def execute(self): - inputs = {k: v.view(self.input_dtypes[k]) for k,v in self.inputs.items()} - inputs = {k: v.reshape(self.input_shapes[k]).astype(self.input_dtypes[k]) for k,v in inputs.items()} - outputs = self.session.run(None, inputs) - assert len(outputs) == 1, "Only single model outputs are supported" - self.output[:] = outputs[0] - return self.output diff --git a/selfdrive/modeld/runners/run.h b/selfdrive/modeld/runners/run.h deleted file mode 100644 index 36ad262a5b..0000000000 --- a/selfdrive/modeld/runners/run.h +++ /dev/null @@ -1,4 +0,0 @@ -#pragma once - -#include "selfdrive/modeld/runners/runmodel.h" -#include "selfdrive/modeld/runners/snpemodel.h" diff --git a/selfdrive/modeld/runners/runmodel.h b/selfdrive/modeld/runners/runmodel.h deleted file mode 100644 index 18cc180cb7..0000000000 --- a/selfdrive/modeld/runners/runmodel.h +++ /dev/null @@ -1,49 +0,0 @@ -#pragma once - -#include -#include -#include -#include - -#include "common/clutil.h" -#include "common/swaglog.h" - -#define USE_CPU_RUNTIME 0 -#define USE_GPU_RUNTIME 1 -#define USE_DSP_RUNTIME 2 - -struct ModelInput { - const std::string name; - float *buffer; - int size; - - ModelInput(const std::string _name, float *_buffer, int _size) : name(_name), buffer(_buffer), size(_size) {} - virtual void setBuffer(float *_buffer, int _size) { - assert(size == _size || size == 0); - buffer = _buffer; - size = _size; - } -}; - -class RunModel { -public: - std::vector> inputs; - - virtual ~RunModel() {} - virtual void execute() {} - virtual void* getCLBuffer(const std::string name) { return nullptr; } - - virtual void addInput(const std::string name, float *buffer, int size) { - inputs.push_back(std::unique_ptr(new ModelInput(name, buffer, size))); - } - virtual void setInputBuffer(const std::string name, float *buffer, int size) { - for (auto &input : inputs) { - if (name == input->name) { - input->setBuffer(buffer, size); - return; - } - } - LOGE("Tried to update input `%s` but no input with this name exists", name.c_str()); - assert(false); - } -}; diff --git a/selfdrive/modeld/runners/runmodel.pxd b/selfdrive/modeld/runners/runmodel.pxd deleted file mode 100644 index 01b2a9cf2c..0000000000 --- a/selfdrive/modeld/runners/runmodel.pxd +++ /dev/null @@ -1,14 +0,0 @@ -# distutils: language = c++ - -from libcpp.string cimport string - -cdef extern from "selfdrive/modeld/runners/runmodel.h": - cdef int USE_CPU_RUNTIME - cdef int USE_GPU_RUNTIME - cdef int USE_DSP_RUNTIME - - cdef cppclass RunModel: - void addInput(string, float*, int) - void setInputBuffer(string, float*, int) - void * getCLBuffer(string) - void execute() diff --git a/selfdrive/modeld/runners/runmodel_pyx.pxd b/selfdrive/modeld/runners/runmodel_pyx.pxd deleted file mode 100644 index b6ede7cf37..0000000000 --- a/selfdrive/modeld/runners/runmodel_pyx.pxd +++ /dev/null @@ -1,6 +0,0 @@ -# distutils: language = c++ - -from .runmodel cimport RunModel as cppRunModel - -cdef class RunModel: - cdef cppRunModel * model diff --git a/selfdrive/modeld/runners/runmodel_pyx.pyx b/selfdrive/modeld/runners/runmodel_pyx.pyx deleted file mode 100644 index 12b8ec10ff..0000000000 --- a/selfdrive/modeld/runners/runmodel_pyx.pyx +++ /dev/null @@ -1,37 +0,0 @@ -# distutils: language = c++ -# cython: c_string_encoding=ascii, language_level=3 - -from libcpp.string cimport string - -from .runmodel cimport USE_CPU_RUNTIME, USE_GPU_RUNTIME, USE_DSP_RUNTIME -from selfdrive.modeld.models.commonmodel_pyx cimport CLMem - -class Runtime: - CPU = USE_CPU_RUNTIME - GPU = USE_GPU_RUNTIME - DSP = USE_DSP_RUNTIME - -cdef class RunModel: - def __dealloc__(self): - del self.model - - def addInput(self, string name, float[:] buffer): - if buffer is not None: - self.model.addInput(name, &buffer[0], len(buffer)) - else: - self.model.addInput(name, NULL, 0) - - def setInputBuffer(self, string name, float[:] buffer): - if buffer is not None: - self.model.setInputBuffer(name, &buffer[0], len(buffer)) - else: - self.model.setInputBuffer(name, NULL, 0) - - def getCLBuffer(self, string name): - cdef void * cl_buf = self.model.getCLBuffer(name) - if not cl_buf: - return None - return CLMem.create(cl_buf) - - def execute(self): - self.model.execute() diff --git a/selfdrive/modeld/runners/snpemodel.cc b/selfdrive/modeld/runners/snpemodel.cc deleted file mode 100644 index 15c1db0086..0000000000 --- a/selfdrive/modeld/runners/snpemodel.cc +++ /dev/null @@ -1,116 +0,0 @@ -#pragma clang diagnostic ignored "-Wexceptions" - -#include "selfdrive/modeld/runners/snpemodel.h" - -#include -#include -#include -#include -#include - -#include "common/util.h" -#include "common/timing.h" - -void PrintErrorStringAndExit() { - std::cerr << zdl::DlSystem::getLastErrorString() << std::endl; - std::exit(EXIT_FAILURE); -} - -SNPEModel::SNPEModel(const std::string path, float *_output, size_t _output_size, int runtime, bool _use_tf8, cl_context context) { - output = _output; - output_size = _output_size; - use_tf8 = _use_tf8; - -#ifdef QCOM2 - if (runtime == USE_GPU_RUNTIME) { - snpe_runtime = zdl::DlSystem::Runtime_t::GPU; - } else if (runtime == USE_DSP_RUNTIME) { - snpe_runtime = zdl::DlSystem::Runtime_t::DSP; - } else { - snpe_runtime = zdl::DlSystem::Runtime_t::CPU; - } - assert(zdl::SNPE::SNPEFactory::isRuntimeAvailable(snpe_runtime)); -#endif - model_data = util::read_file(path); - assert(model_data.size() > 0); - - // load model - std::unique_ptr container = zdl::DlContainer::IDlContainer::open((uint8_t*)model_data.data(), model_data.size()); - if (!container) { PrintErrorStringAndExit(); } - LOGW("loaded model with size: %lu", model_data.size()); - - // create model runner - zdl::SNPE::SNPEBuilder snpe_builder(container.get()); - while (!snpe) { -#ifdef QCOM2 - snpe = snpe_builder.setOutputLayers({}) - .setRuntimeProcessor(snpe_runtime) - .setUseUserSuppliedBuffers(true) - .setPerformanceProfile(zdl::DlSystem::PerformanceProfile_t::HIGH_PERFORMANCE) - .build(); -#else - snpe = snpe_builder.setOutputLayers({}) - .setUseUserSuppliedBuffers(true) - .setPerformanceProfile(zdl::DlSystem::PerformanceProfile_t::HIGH_PERFORMANCE) - .build(); -#endif - if (!snpe) std::cerr << zdl::DlSystem::getLastErrorString() << std::endl; - } - - // create output buffer - zdl::DlSystem::UserBufferEncodingFloat ub_encoding_float; - zdl::DlSystem::IUserBufferFactory &ub_factory = zdl::SNPE::SNPEFactory::getUserBufferFactory(); - - const auto &output_tensor_names_opt = snpe->getOutputTensorNames(); - if (!output_tensor_names_opt) throw std::runtime_error("Error obtaining output tensor names"); - const auto &output_tensor_names = *output_tensor_names_opt; - assert(output_tensor_names.size() == 1); - const char *output_tensor_name = output_tensor_names.at(0); - const zdl::DlSystem::TensorShape &buffer_shape = snpe->getInputOutputBufferAttributes(output_tensor_name)->getDims(); - if (output_size != 0) { - assert(output_size == buffer_shape[1]); - } else { - output_size = buffer_shape[1]; - } - std::vector output_strides = {output_size * sizeof(float), sizeof(float)}; - output_buffer = ub_factory.createUserBuffer(output, output_size * sizeof(float), output_strides, &ub_encoding_float); - output_map.add(output_tensor_name, output_buffer.get()); -} - -void SNPEModel::addInput(const std::string name, float *buffer, int size) { - const int idx = inputs.size(); - const auto &input_tensor_names_opt = snpe->getInputTensorNames(); - if (!input_tensor_names_opt) throw std::runtime_error("Error obtaining input tensor names"); - const auto &input_tensor_names = *input_tensor_names_opt; - const char *input_tensor_name = input_tensor_names.at(idx); - const bool input_tf8 = use_tf8 && strcmp(input_tensor_name, "input_img") == 0; // TODO: This is a terrible hack, get rid of this name check both here and in onnx_runner.py - LOGW("adding index %d: %s", idx, input_tensor_name); - - zdl::DlSystem::UserBufferEncodingFloat ub_encoding_float; - zdl::DlSystem::UserBufferEncodingTf8 ub_encoding_tf8(0, 1./255); // network takes 0-1 - zdl::DlSystem::IUserBufferFactory &ub_factory = zdl::SNPE::SNPEFactory::getUserBufferFactory(); - zdl::DlSystem::UserBufferEncoding *input_encoding = input_tf8 ? (zdl::DlSystem::UserBufferEncoding*)&ub_encoding_tf8 : (zdl::DlSystem::UserBufferEncoding*)&ub_encoding_float; - - const auto &buffer_shape_opt = snpe->getInputDimensions(input_tensor_name); - const zdl::DlSystem::TensorShape &buffer_shape = *buffer_shape_opt; - size_t size_of_input = input_tf8 ? sizeof(uint8_t) : sizeof(float); - std::vector strides(buffer_shape.rank()); - strides[strides.size() - 1] = size_of_input; - size_t product = 1; - for (size_t i = 0; i < buffer_shape.rank(); i++) product *= buffer_shape[i]; - size_t stride = strides[strides.size() - 1]; - for (size_t i = buffer_shape.rank() - 1; i > 0; i--) { - stride *= buffer_shape[i]; - strides[i-1] = stride; - } - - auto input_buffer = ub_factory.createUserBuffer(buffer, product*size_of_input, strides, input_encoding); - input_map.add(input_tensor_name, input_buffer.get()); - inputs.push_back(std::unique_ptr(new SNPEModelInput(name, buffer, size, std::move(input_buffer)))); -} - -void SNPEModel::execute() { - if (!snpe->execute(input_map, output_map)) { - PrintErrorStringAndExit(); - } -} diff --git a/selfdrive/modeld/runners/snpemodel.h b/selfdrive/modeld/runners/snpemodel.h deleted file mode 100644 index 86b2c86084..0000000000 --- a/selfdrive/modeld/runners/snpemodel.h +++ /dev/null @@ -1,52 +0,0 @@ -#pragma once -#pragma clang diagnostic ignored "-Wdeprecated-declarations" - -#include -#include -#include - -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include "selfdrive/modeld/runners/runmodel.h" - -struct SNPEModelInput : public ModelInput { - std::unique_ptr snpe_buffer; - - SNPEModelInput(const std::string _name, float *_buffer, int _size, std::unique_ptr _snpe_buffer) : ModelInput(_name, _buffer, _size), snpe_buffer(std::move(_snpe_buffer)) {} - void setBuffer(float *_buffer, int _size) { - ModelInput::setBuffer(_buffer, _size); - assert(snpe_buffer->setBufferAddress(_buffer) == true); - } -}; - -class SNPEModel : public RunModel { -public: - SNPEModel(const std::string path, float *_output, size_t _output_size, int runtime, bool use_tf8 = false, cl_context context = NULL); - void addInput(const std::string name, float *buffer, int size); - void execute(); - -private: - std::string model_data; - -#ifdef QCOM2 - zdl::DlSystem::Runtime_t snpe_runtime; -#endif - - // snpe model stuff - std::unique_ptr snpe; - zdl::DlSystem::UserBufferMap input_map; - zdl::DlSystem::UserBufferMap output_map; - std::unique_ptr output_buffer; - - bool use_tf8; - float *output; - size_t output_size; -}; diff --git a/selfdrive/modeld/runners/snpemodel.pxd b/selfdrive/modeld/runners/snpemodel.pxd deleted file mode 100644 index a911b43584..0000000000 --- a/selfdrive/modeld/runners/snpemodel.pxd +++ /dev/null @@ -1,9 +0,0 @@ -# distutils: language = c++ - -from libcpp.string cimport string - -from msgq.visionipc.visionipc cimport cl_context - -cdef extern from "selfdrive/modeld/runners/snpemodel.h": - cdef cppclass SNPEModel: - SNPEModel(string, float*, size_t, int, bool, cl_context) diff --git a/selfdrive/modeld/runners/snpemodel_pyx.pyx b/selfdrive/modeld/runners/snpemodel_pyx.pyx deleted file mode 100644 index f83b7c8cff..0000000000 --- a/selfdrive/modeld/runners/snpemodel_pyx.pyx +++ /dev/null @@ -1,17 +0,0 @@ -# distutils: language = c++ -# cython: c_string_encoding=ascii, language_level=3 - -import os -from libcpp cimport bool -from libcpp.string cimport string - -from .snpemodel cimport SNPEModel as cppSNPEModel -from selfdrive.modeld.models.commonmodel_pyx cimport CLContext -from selfdrive.modeld.runners.runmodel_pyx cimport RunModel -from selfdrive.modeld.runners.runmodel cimport RunModel as cppRunModel - -os.environ['ADSP_LIBRARY_PATH'] = "/data/pythonpath/third_party/snpe/dsp/" - -cdef class SNPEModel(RunModel): - def __cinit__(self, string path, float[:] output, int runtime, bool use_tf8, CLContext context): - self.model = new cppSNPEModel(path, &output[0], len(output), runtime, use_tf8, context.context) diff --git a/selfdrive/modeld/runners/thneedmodel.cc b/selfdrive/modeld/runners/thneedmodel.cc deleted file mode 100644 index a16d8b42aa..0000000000 --- a/selfdrive/modeld/runners/thneedmodel.cc +++ /dev/null @@ -1,58 +0,0 @@ -#include "selfdrive/modeld/runners/thneedmodel.h" - -#include - -#include "common/swaglog.h" - -ThneedModel::ThneedModel(const std::string path, float *_output, size_t _output_size, int runtime, bool luse_tf8, cl_context context) { - thneed = new Thneed(true, context); - thneed->load(path.c_str()); - thneed->clexec(); - - recorded = false; - output = _output; -} - -void* ThneedModel::getCLBuffer(const std::string name) { - int index = -1; - for (int i = 0; i < inputs.size(); i++) { - if (name == inputs[i]->name) { - index = i; - break; - } - } - - if (index == -1) { - LOGE("Tried to get CL buffer for input `%s` but no input with this name exists", name.c_str()); - assert(false); - } - - if (thneed->input_clmem.size() >= inputs.size()) { - return &thneed->input_clmem[inputs.size() - index - 1]; - } else { - return nullptr; - } -} - -void ThneedModel::execute() { - if (!recorded) { - thneed->record = true; - float *input_buffers[inputs.size()]; - for (int i = 0; i < inputs.size(); i++) { - input_buffers[inputs.size() - i - 1] = inputs[i]->buffer; - } - - thneed->copy_inputs(input_buffers); - thneed->clexec(); - thneed->copy_output(output); - thneed->stop(); - - recorded = true; - } else { - float *input_buffers[inputs.size()]; - for (int i = 0; i < inputs.size(); i++) { - input_buffers[inputs.size() - i - 1] = inputs[i]->buffer; - } - thneed->execute(input_buffers, output); - } -} diff --git a/selfdrive/modeld/runners/thneedmodel.h b/selfdrive/modeld/runners/thneedmodel.h deleted file mode 100644 index 6ed479c081..0000000000 --- a/selfdrive/modeld/runners/thneedmodel.h +++ /dev/null @@ -1,17 +0,0 @@ -#pragma once - -#include - -#include "selfdrive/modeld/runners/runmodel.h" -#include "selfdrive/modeld/thneed/thneed.h" - -class ThneedModel : public RunModel { -public: - ThneedModel(const std::string path, float *_output, size_t _output_size, int runtime, bool use_tf8 = false, cl_context context = NULL); - void *getCLBuffer(const std::string name); - void execute(); -private: - Thneed *thneed = NULL; - bool recorded; - float *output; -}; diff --git a/selfdrive/modeld/runners/thneedmodel.pxd b/selfdrive/modeld/runners/thneedmodel.pxd deleted file mode 100644 index 79e24dbdd6..0000000000 --- a/selfdrive/modeld/runners/thneedmodel.pxd +++ /dev/null @@ -1,9 +0,0 @@ -# distutils: language = c++ - -from libcpp.string cimport string - -from msgq.visionipc.visionipc cimport cl_context - -cdef extern from "selfdrive/modeld/runners/thneedmodel.h": - cdef cppclass ThneedModel: - ThneedModel(string, float*, size_t, int, bool, cl_context) diff --git a/selfdrive/modeld/runners/thneedmodel_pyx.pyx b/selfdrive/modeld/runners/thneedmodel_pyx.pyx deleted file mode 100644 index 6f8fdd255f..0000000000 --- a/selfdrive/modeld/runners/thneedmodel_pyx.pyx +++ /dev/null @@ -1,14 +0,0 @@ -# distutils: language = c++ -# cython: c_string_encoding=ascii, language_level=3 - -from libcpp cimport bool -from libcpp.string cimport string - -from .thneedmodel cimport ThneedModel as cppThneedModel -from selfdrive.modeld.models.commonmodel_pyx cimport CLContext -from selfdrive.modeld.runners.runmodel_pyx cimport RunModel -from selfdrive.modeld.runners.runmodel cimport RunModel as cppRunModel - -cdef class ThneedModel(RunModel): - def __cinit__(self, string path, float[:] output, int runtime, bool use_tf8, CLContext context): - self.model = new cppThneedModel(path, &output[0], len(output), runtime, use_tf8, context.context) diff --git a/selfdrive/modeld/runners/tinygrad_helpers.py b/selfdrive/modeld/runners/tinygrad_helpers.py new file mode 100644 index 0000000000..776381341c --- /dev/null +++ b/selfdrive/modeld/runners/tinygrad_helpers.py @@ -0,0 +1,8 @@ + +from tinygrad.tensor import Tensor +from tinygrad.helpers import to_mv + +def qcom_tensor_from_opencl_address(opencl_address, shape, dtype): + cl_buf_desc_ptr = to_mv(opencl_address, 8).cast('Q')[0] + rawbuf_ptr = to_mv(cl_buf_desc_ptr, 0x100).cast('Q')[20] # offset 0xA0 is a raw gpu pointer. + return Tensor.from_blob(rawbuf_ptr, shape, dtype=dtype, device='QCOM') diff --git a/selfdrive/modeld/thneed/README b/selfdrive/modeld/thneed/README deleted file mode 100644 index f3bc66d8fc..0000000000 --- a/selfdrive/modeld/thneed/README +++ /dev/null @@ -1,8 +0,0 @@ -thneed is an SNPE accelerator. I know SNPE is already an accelerator, but sometimes things need to go even faster.. - -It runs on the local device, and caches a single model run. Then it replays it, but fast. - -thneed slices through abstraction layers like a fish. - -You need a thneed. - diff --git a/selfdrive/modeld/thneed/__init__.py b/selfdrive/modeld/thneed/__init__.py deleted file mode 100644 index e69de29bb2..0000000000 diff --git a/selfdrive/modeld/thneed/serialize.cc b/selfdrive/modeld/thneed/serialize.cc deleted file mode 100644 index 3dc2bef414..0000000000 --- a/selfdrive/modeld/thneed/serialize.cc +++ /dev/null @@ -1,154 +0,0 @@ -#include -#include - -#include "third_party/json11/json11.hpp" -#include "common/util.h" -#include "common/clutil.h" -#include "common/swaglog.h" -#include "selfdrive/modeld/thneed/thneed.h" -using namespace json11; - -extern map g_program_source; - -void Thneed::load(const char *filename) { - LOGD("Thneed::load: loading from %s\n", filename); - - string buf = util::read_file(filename); - int jsz = *(int *)buf.data(); - string jsonerr; - string jj(buf.data() + sizeof(int), jsz); - Json jdat = Json::parse(jj, jsonerr); - - map real_mem; - real_mem[NULL] = NULL; - - int ptr = sizeof(int)+jsz; - for (auto &obj : jdat["objects"].array_items()) { - auto mobj = obj.object_items(); - int sz = mobj["size"].int_value(); - cl_mem clbuf = NULL; - - if (mobj["buffer_id"].string_value().size() > 0) { - // image buffer must already be allocated - clbuf = real_mem[*(cl_mem*)(mobj["buffer_id"].string_value().data())]; - assert(mobj["needs_load"].bool_value() == false); - } else { - if (mobj["needs_load"].bool_value()) { - clbuf = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR | CL_MEM_READ_WRITE, sz, &buf[ptr], NULL); - if (debug >= 1) printf("loading %p %d @ 0x%X\n", clbuf, sz, ptr); - ptr += sz; - } else { - // TODO: is there a faster way to init zeroed out buffers? - void *host_zeros = calloc(sz, 1); - clbuf = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR | CL_MEM_READ_WRITE, sz, host_zeros, NULL); - free(host_zeros); - } - } - assert(clbuf != NULL); - - if (mobj["arg_type"] == "image2d_t" || mobj["arg_type"] == "image1d_t") { - cl_image_desc desc = {0}; - desc.image_type = (mobj["arg_type"] == "image2d_t") ? CL_MEM_OBJECT_IMAGE2D : CL_MEM_OBJECT_IMAGE1D_BUFFER; - desc.image_width = mobj["width"].int_value(); - desc.image_height = mobj["height"].int_value(); - desc.image_row_pitch = mobj["row_pitch"].int_value(); - assert(sz == desc.image_height*desc.image_row_pitch); -#ifdef QCOM2 - desc.buffer = clbuf; -#else - // TODO: we are creating unused buffers on PC - clReleaseMemObject(clbuf); -#endif - cl_image_format format = {0}; - format.image_channel_order = CL_RGBA; - format.image_channel_data_type = mobj["float32"].bool_value() ? CL_FLOAT : CL_HALF_FLOAT; - - cl_int errcode; - -#ifndef QCOM2 - if (mobj["needs_load"].bool_value()) { - clbuf = clCreateImage(context, CL_MEM_COPY_HOST_PTR | CL_MEM_READ_WRITE, &format, &desc, &buf[ptr-sz], &errcode); - } else { - clbuf = clCreateImage(context, CL_MEM_READ_WRITE, &format, &desc, NULL, &errcode); - } -#else - clbuf = clCreateImage(context, CL_MEM_READ_WRITE, &format, &desc, NULL, &errcode); -#endif - if (clbuf == NULL) { - LOGE("clError: %s create image %zux%zu rp %zu with buffer %p\n", cl_get_error_string(errcode), - desc.image_width, desc.image_height, desc.image_row_pitch, desc.buffer); - } - assert(clbuf != NULL); - } - - real_mem[*(cl_mem*)(mobj["id"].string_value().data())] = clbuf; - } - - map g_programs; - for (const auto &[name, source] : jdat["programs"].object_items()) { - if (debug >= 1) printf("building %s with size %zu\n", name.c_str(), source.string_value().size()); - g_programs[name] = cl_program_from_source(context, device_id, source.string_value()); - } - - for (auto &obj : jdat["inputs"].array_items()) { - auto mobj = obj.object_items(); - int sz = mobj["size"].int_value(); - cl_mem aa = real_mem[*(cl_mem*)(mobj["buffer_id"].string_value().data())]; - input_clmem.push_back(aa); - input_sizes.push_back(sz); - LOGD("Thneed::load: adding input %s with size %d\n", mobj["name"].string_value().data(), sz); - - cl_int cl_err; - void *ret = clEnqueueMapBuffer(command_queue, aa, CL_TRUE, CL_MAP_WRITE, 0, sz, 0, NULL, NULL, &cl_err); - if (cl_err != CL_SUCCESS) LOGE("clError: %s map %p %d\n", cl_get_error_string(cl_err), aa, sz); - assert(cl_err == CL_SUCCESS); - inputs.push_back(ret); - } - - for (auto &obj : jdat["outputs"].array_items()) { - auto mobj = obj.object_items(); - int sz = mobj["size"].int_value(); - LOGD("Thneed::save: adding output with size %d\n", sz); - // TODO: support multiple outputs - output = real_mem[*(cl_mem*)(mobj["buffer_id"].string_value().data())]; - assert(output != NULL); - } - - for (auto &obj : jdat["binaries"].array_items()) { - string name = obj["name"].string_value(); - size_t length = obj["length"].int_value(); - if (debug >= 1) printf("binary %s with size %zu\n", name.c_str(), length); - g_programs[name] = cl_program_from_binary(context, device_id, (const uint8_t*)&buf[ptr], length); - ptr += length; - } - - for (auto &obj : jdat["kernels"].array_items()) { - auto gws = obj["global_work_size"]; - auto lws = obj["local_work_size"]; - auto kk = shared_ptr(new CLQueuedKernel(this)); - - kk->name = obj["name"].string_value(); - kk->program = g_programs[kk->name]; - kk->work_dim = obj["work_dim"].int_value(); - for (int i = 0; i < kk->work_dim; i++) { - kk->global_work_size[i] = gws[i].int_value(); - kk->local_work_size[i] = lws[i].int_value(); - } - kk->num_args = obj["num_args"].int_value(); - for (int i = 0; i < kk->num_args; i++) { - string arg = obj["args"].array_items()[i].string_value(); - int arg_size = obj["args_size"].array_items()[i].int_value(); - kk->args_size.push_back(arg_size); - if (arg_size == 8) { - cl_mem val = *(cl_mem*)(arg.data()); - val = real_mem[val]; - kk->args.push_back(string((char*)&val, sizeof(val))); - } else { - kk->args.push_back(arg); - } - } - kq.push_back(kk); - } - - clFinish(command_queue); -} diff --git a/selfdrive/modeld/thneed/thneed.h b/selfdrive/modeld/thneed/thneed.h deleted file mode 100644 index 47e18e0be3..0000000000 --- a/selfdrive/modeld/thneed/thneed.h +++ /dev/null @@ -1,133 +0,0 @@ -#pragma once - -#ifndef __user -#define __user __attribute__(()) -#endif - -#include -#include -#include -#include -#include - -#include - -#include "third_party/linux/include/msm_kgsl.h" - -using namespace std; - -cl_int thneed_clSetKernelArg(cl_kernel kernel, cl_uint arg_index, size_t arg_size, const void *arg_value); - -namespace json11 { - class Json; -} -class Thneed; - -class GPUMalloc { - public: - GPUMalloc(int size, int fd); - ~GPUMalloc(); - void *alloc(int size); - private: - uint64_t base; - int remaining; -}; - -class CLQueuedKernel { - public: - CLQueuedKernel(Thneed *lthneed) { thneed = lthneed; } - CLQueuedKernel(Thneed *lthneed, - cl_kernel _kernel, - cl_uint _work_dim, - const size_t *_global_work_size, - const size_t *_local_work_size); - cl_int exec(); - void debug_print(bool verbose); - int get_arg_num(const char *search_arg_name); - cl_program program; - string name; - cl_uint num_args; - vector arg_names; - vector arg_types; - vector args; - vector args_size; - cl_kernel kernel = NULL; - json11::Json to_json() const; - - cl_uint work_dim; - size_t global_work_size[3] = {0}; - size_t local_work_size[3] = {0}; - private: - Thneed *thneed; -}; - -class CachedIoctl { - public: - virtual void exec() {} -}; - -class CachedSync: public CachedIoctl { - public: - CachedSync(Thneed *lthneed, string ldata) { thneed = lthneed; data = ldata; } - void exec(); - private: - Thneed *thneed; - string data; -}; - -class CachedCommand: public CachedIoctl { - public: - CachedCommand(Thneed *lthneed, struct kgsl_gpu_command *cmd); - void exec(); - private: - void disassemble(int cmd_index); - struct kgsl_gpu_command cache; - unique_ptr cmds; - unique_ptr objs; - Thneed *thneed; - vector > kq; -}; - -class Thneed { - public: - Thneed(bool do_clinit=false, cl_context _context = NULL); - void stop(); - void execute(float **finputs, float *foutput, bool slow=false); - void wait(); - - vector input_clmem; - vector inputs; - vector input_sizes; - cl_mem output = NULL; - - cl_context context = NULL; - cl_command_queue command_queue; - cl_device_id device_id; - int context_id; - - // protected? - bool record = false; - int debug; - int timestamp; - -#ifdef QCOM2 - unique_ptr ram; - vector > cmds; - int fd; -#endif - - // all CL kernels - void copy_inputs(float **finputs, bool internal=false); - void copy_output(float *foutput); - cl_int clexec(); - vector > kq; - - // pending CL kernels - vector > ckq; - - // loading - void load(const char *filename); - private: - void clinit(); -}; - diff --git a/selfdrive/modeld/thneed/thneed_common.cc b/selfdrive/modeld/thneed/thneed_common.cc deleted file mode 100644 index ecdf1237e3..0000000000 --- a/selfdrive/modeld/thneed/thneed_common.cc +++ /dev/null @@ -1,216 +0,0 @@ -#include "selfdrive/modeld/thneed/thneed.h" - -#include -#include -#include - -#include "common/clutil.h" -#include "common/timing.h" - -map, string> g_args; -map, int> g_args_size; -map g_program_source; - -void Thneed::stop() { - //printf("Thneed::stop: recorded %lu commands\n", cmds.size()); - record = false; -} - -void Thneed::clinit() { - device_id = cl_get_device_id(CL_DEVICE_TYPE_DEFAULT); - if (context == NULL) context = CL_CHECK_ERR(clCreateContext(NULL, 1, &device_id, NULL, NULL, &err)); - //cl_command_queue_properties props[3] = {CL_QUEUE_PROPERTIES, CL_QUEUE_PROFILING_ENABLE, 0}; - cl_command_queue_properties props[3] = {CL_QUEUE_PROPERTIES, 0, 0}; - command_queue = CL_CHECK_ERR(clCreateCommandQueueWithProperties(context, device_id, props, &err)); - printf("Thneed::clinit done\n"); -} - -cl_int Thneed::clexec() { - if (debug >= 1) printf("Thneed::clexec: running %lu queued kernels\n", kq.size()); - for (auto &k : kq) { - if (record) ckq.push_back(k); - cl_int ret = k->exec(); - assert(ret == CL_SUCCESS); - } - return clFinish(command_queue); -} - -void Thneed::copy_inputs(float **finputs, bool internal) { - for (int idx = 0; idx < inputs.size(); ++idx) { - if (debug >= 1) printf("copying %lu -- %p -> %p (cl %p)\n", input_sizes[idx], finputs[idx], inputs[idx], input_clmem[idx]); - - if (internal) { - // if it's internal, using memcpy is fine since the buffer sync is cached in the ioctl layer - if (finputs[idx] != NULL) memcpy(inputs[idx], finputs[idx], input_sizes[idx]); - } else { - if (finputs[idx] != NULL) CL_CHECK(clEnqueueWriteBuffer(command_queue, input_clmem[idx], CL_TRUE, 0, input_sizes[idx], finputs[idx], 0, NULL, NULL)); - } - } -} - -void Thneed::copy_output(float *foutput) { - if (output != NULL) { - size_t sz; - clGetMemObjectInfo(output, CL_MEM_SIZE, sizeof(sz), &sz, NULL); - if (debug >= 1) printf("copying %lu for output %p -> %p\n", sz, output, foutput); - CL_CHECK(clEnqueueReadBuffer(command_queue, output, CL_TRUE, 0, sz, foutput, 0, NULL, NULL)); - } else { - printf("CAUTION: model output is NULL, does it have no outputs?\n"); - } -} - -// *********** CLQueuedKernel *********** - -CLQueuedKernel::CLQueuedKernel(Thneed *lthneed, - cl_kernel _kernel, - cl_uint _work_dim, - const size_t *_global_work_size, - const size_t *_local_work_size) { - thneed = lthneed; - kernel = _kernel; - work_dim = _work_dim; - assert(work_dim <= 3); - for (int i = 0; i < work_dim; i++) { - global_work_size[i] = _global_work_size[i]; - local_work_size[i] = _local_work_size[i]; - } - - char _name[0x100]; - clGetKernelInfo(kernel, CL_KERNEL_FUNCTION_NAME, sizeof(_name), _name, NULL); - name = string(_name); - clGetKernelInfo(kernel, CL_KERNEL_NUM_ARGS, sizeof(num_args), &num_args, NULL); - - // get args - for (int i = 0; i < num_args; i++) { - char arg_name[0x100] = {0}; - clGetKernelArgInfo(kernel, i, CL_KERNEL_ARG_NAME, sizeof(arg_name), arg_name, NULL); - arg_names.push_back(string(arg_name)); - clGetKernelArgInfo(kernel, i, CL_KERNEL_ARG_TYPE_NAME, sizeof(arg_name), arg_name, NULL); - arg_types.push_back(string(arg_name)); - - args.push_back(g_args[make_pair(kernel, i)]); - args_size.push_back(g_args_size[make_pair(kernel, i)]); - } - - // get program - clGetKernelInfo(kernel, CL_KERNEL_PROGRAM, sizeof(program), &program, NULL); -} - -int CLQueuedKernel::get_arg_num(const char *search_arg_name) { - for (int i = 0; i < num_args; i++) { - if (arg_names[i] == search_arg_name) return i; - } - printf("failed to find %s in %s\n", search_arg_name, name.c_str()); - assert(false); -} - -cl_int CLQueuedKernel::exec() { - if (kernel == NULL) { - kernel = clCreateKernel(program, name.c_str(), NULL); - arg_names.clear(); - arg_types.clear(); - - for (int j = 0; j < num_args; j++) { - char arg_name[0x100] = {0}; - clGetKernelArgInfo(kernel, j, CL_KERNEL_ARG_NAME, sizeof(arg_name), arg_name, NULL); - arg_names.push_back(string(arg_name)); - clGetKernelArgInfo(kernel, j, CL_KERNEL_ARG_TYPE_NAME, sizeof(arg_name), arg_name, NULL); - arg_types.push_back(string(arg_name)); - - cl_int ret; - if (args[j].size() != 0) { - assert(args[j].size() == args_size[j]); - ret = thneed_clSetKernelArg(kernel, j, args[j].size(), args[j].data()); - } else { - ret = thneed_clSetKernelArg(kernel, j, args_size[j], NULL); - } - assert(ret == CL_SUCCESS); - } - } - - if (thneed->debug >= 1) { - debug_print(thneed->debug >= 2); - } - - return clEnqueueNDRangeKernel(thneed->command_queue, - kernel, work_dim, NULL, global_work_size, local_work_size, 0, NULL, NULL); -} - -void CLQueuedKernel::debug_print(bool verbose) { - printf("%p %56s -- ", kernel, name.c_str()); - for (int i = 0; i < work_dim; i++) { - printf("%4zu ", global_work_size[i]); - } - printf(" -- "); - for (int i = 0; i < work_dim; i++) { - printf("%4zu ", local_work_size[i]); - } - printf("\n"); - - if (verbose) { - for (int i = 0; i < num_args; i++) { - string arg = args[i]; - printf(" %s %s", arg_types[i].c_str(), arg_names[i].c_str()); - void *arg_value = (void*)arg.data(); - int arg_size = arg.size(); - if (arg_size == 0) { - printf(" (size) %d", args_size[i]); - } else if (arg_size == 1) { - printf(" = %d", *((char*)arg_value)); - } else if (arg_size == 2) { - printf(" = %d", *((short*)arg_value)); - } else if (arg_size == 4) { - if (arg_types[i] == "float") { - printf(" = %f", *((float*)arg_value)); - } else { - printf(" = %d", *((int*)arg_value)); - } - } else if (arg_size == 8) { - cl_mem val = (cl_mem)(*((uintptr_t*)arg_value)); - printf(" = %p", val); - if (val != NULL) { - cl_mem_object_type obj_type; - clGetMemObjectInfo(val, CL_MEM_TYPE, sizeof(obj_type), &obj_type, NULL); - if (arg_types[i] == "image2d_t" || arg_types[i] == "image1d_t" || obj_type == CL_MEM_OBJECT_IMAGE2D) { - cl_image_format format; - size_t width, height, depth, array_size, row_pitch, slice_pitch; - cl_mem buf; - clGetImageInfo(val, CL_IMAGE_FORMAT, sizeof(format), &format, NULL); - assert(format.image_channel_order == CL_RGBA); - assert(format.image_channel_data_type == CL_HALF_FLOAT || format.image_channel_data_type == CL_FLOAT); - clGetImageInfo(val, CL_IMAGE_WIDTH, sizeof(width), &width, NULL); - clGetImageInfo(val, CL_IMAGE_HEIGHT, sizeof(height), &height, NULL); - clGetImageInfo(val, CL_IMAGE_ROW_PITCH, sizeof(row_pitch), &row_pitch, NULL); - clGetImageInfo(val, CL_IMAGE_DEPTH, sizeof(depth), &depth, NULL); - clGetImageInfo(val, CL_IMAGE_ARRAY_SIZE, sizeof(array_size), &array_size, NULL); - clGetImageInfo(val, CL_IMAGE_SLICE_PITCH, sizeof(slice_pitch), &slice_pitch, NULL); - assert(depth == 0); - assert(array_size == 0); - assert(slice_pitch == 0); - - clGetImageInfo(val, CL_IMAGE_BUFFER, sizeof(buf), &buf, NULL); - size_t sz = 0; - if (buf != NULL) clGetMemObjectInfo(buf, CL_MEM_SIZE, sizeof(sz), &sz, NULL); - printf(" image %zu x %zu rp %zu @ %p buffer %zu", width, height, row_pitch, buf, sz); - } else { - size_t sz; - clGetMemObjectInfo(val, CL_MEM_SIZE, sizeof(sz), &sz, NULL); - printf(" buffer %zu", sz); - } - } - } - printf("\n"); - } - } -} - -cl_int thneed_clSetKernelArg(cl_kernel kernel, cl_uint arg_index, size_t arg_size, const void *arg_value) { - g_args_size[make_pair(kernel, arg_index)] = arg_size; - if (arg_value != NULL) { - g_args[make_pair(kernel, arg_index)] = string((char*)arg_value, arg_size); - } else { - g_args[make_pair(kernel, arg_index)] = string(""); - } - cl_int ret = clSetKernelArg(kernel, arg_index, arg_size, arg_value); - return ret; -} diff --git a/selfdrive/modeld/thneed/thneed_pc.cc b/selfdrive/modeld/thneed/thneed_pc.cc deleted file mode 100644 index 8d0037628e..0000000000 --- a/selfdrive/modeld/thneed/thneed_pc.cc +++ /dev/null @@ -1,32 +0,0 @@ -#include "selfdrive/modeld/thneed/thneed.h" - -#include - -#include "common/clutil.h" -#include "common/timing.h" - -Thneed::Thneed(bool do_clinit, cl_context _context) { - context = _context; - if (do_clinit) clinit(); - char *thneed_debug_env = getenv("THNEED_DEBUG"); - debug = (thneed_debug_env != NULL) ? atoi(thneed_debug_env) : 0; -} - -void Thneed::execute(float **finputs, float *foutput, bool slow) { - uint64_t tb, te; - if (debug >= 1) tb = nanos_since_boot(); - - // ****** copy inputs - copy_inputs(finputs); - - // ****** run commands - clexec(); - - // ****** copy outputs - copy_output(foutput); - - if (debug >= 1) { - te = nanos_since_boot(); - printf("model exec in %lu us\n", (te-tb)/1000); - } -} diff --git a/selfdrive/modeld/thneed/thneed_qcom2.cc b/selfdrive/modeld/thneed/thneed_qcom2.cc deleted file mode 100644 index 21de15d17c..0000000000 --- a/selfdrive/modeld/thneed/thneed_qcom2.cc +++ /dev/null @@ -1,258 +0,0 @@ -#include "selfdrive/modeld/thneed/thneed.h" - -#include -#include - -#include -#include -#include -#include -#include - -#include "common/clutil.h" -#include "common/timing.h" - -Thneed *g_thneed = NULL; -int g_fd = -1; - -void hexdump(uint8_t *d, int len) { - assert((len%4) == 0); - printf(" dumping %p len 0x%x\n", d, len); - for (int i = 0; i < len/4; i++) { - if (i != 0 && (i%0x10) == 0) printf("\n"); - printf("%8x ", d[i]); - } - printf("\n"); -} - -// *********** ioctl interceptor *********** - -extern "C" { - -int (*my_ioctl)(int filedes, unsigned long request, void *argp) = NULL; -#undef ioctl -int ioctl(int filedes, unsigned long request, void *argp) { - request &= 0xFFFFFFFF; // needed on QCOM2 - if (my_ioctl == NULL) my_ioctl = reinterpret_cast(dlsym(RTLD_NEXT, "ioctl")); - Thneed *thneed = g_thneed; - - // save the fd - if (request == IOCTL_KGSL_GPUOBJ_ALLOC) g_fd = filedes; - - // note that this runs always, even without a thneed object - if (request == IOCTL_KGSL_DRAWCTXT_CREATE) { - struct kgsl_drawctxt_create *create = (struct kgsl_drawctxt_create *)argp; - create->flags &= ~KGSL_CONTEXT_PRIORITY_MASK; - create->flags |= 6 << KGSL_CONTEXT_PRIORITY_SHIFT; // priority from 1-15, 1 is max priority - printf("IOCTL_KGSL_DRAWCTXT_CREATE: creating context with flags 0x%x\n", create->flags); - } - - if (thneed != NULL) { - if (request == IOCTL_KGSL_GPU_COMMAND) { - struct kgsl_gpu_command *cmd = (struct kgsl_gpu_command *)argp; - if (thneed->record) { - thneed->timestamp = cmd->timestamp; - thneed->context_id = cmd->context_id; - thneed->cmds.push_back(unique_ptr(new CachedCommand(thneed, cmd))); - } - if (thneed->debug >= 1) { - printf("IOCTL_KGSL_GPU_COMMAND(%2zu): flags: 0x%lx context_id: %u timestamp: %u numcmds: %d numobjs: %d\n", - thneed->cmds.size(), - cmd->flags, - cmd->context_id, cmd->timestamp, cmd->numcmds, cmd->numobjs); - } - } else if (request == IOCTL_KGSL_GPUOBJ_SYNC) { - struct kgsl_gpuobj_sync *cmd = (struct kgsl_gpuobj_sync *)argp; - struct kgsl_gpuobj_sync_obj *objs = (struct kgsl_gpuobj_sync_obj *)(cmd->objs); - - if (thneed->debug >= 2) { - printf("IOCTL_KGSL_GPUOBJ_SYNC count:%d ", cmd->count); - for (int i = 0; i < cmd->count; i++) { - printf(" -- offset:0x%lx len:0x%lx id:%d op:%d ", objs[i].offset, objs[i].length, objs[i].id, objs[i].op); - } - printf("\n"); - } - - if (thneed->record) { - thneed->cmds.push_back(unique_ptr(new - CachedSync(thneed, string((char *)objs, sizeof(struct kgsl_gpuobj_sync_obj)*cmd->count)))); - } - } else if (request == IOCTL_KGSL_DEVICE_WAITTIMESTAMP_CTXTID) { - struct kgsl_device_waittimestamp_ctxtid *cmd = (struct kgsl_device_waittimestamp_ctxtid *)argp; - if (thneed->debug >= 1) { - printf("IOCTL_KGSL_DEVICE_WAITTIMESTAMP_CTXTID: context_id: %d timestamp: %d timeout: %d\n", - cmd->context_id, cmd->timestamp, cmd->timeout); - } - } else if (request == IOCTL_KGSL_SETPROPERTY) { - if (thneed->debug >= 1) { - struct kgsl_device_getproperty *prop = (struct kgsl_device_getproperty *)argp; - printf("IOCTL_KGSL_SETPROPERTY: 0x%x sizebytes:%zu\n", prop->type, prop->sizebytes); - if (thneed->debug >= 2) { - hexdump((uint8_t *)prop->value, prop->sizebytes); - if (prop->type == KGSL_PROP_PWR_CONSTRAINT) { - struct kgsl_device_constraint *constraint = (struct kgsl_device_constraint *)prop->value; - hexdump((uint8_t *)constraint->data, constraint->size); - } - } - } - } else if (request == IOCTL_KGSL_DRAWCTXT_CREATE || request == IOCTL_KGSL_DRAWCTXT_DESTROY) { - // this happens - } else if (request == IOCTL_KGSL_GPUOBJ_ALLOC || request == IOCTL_KGSL_GPUOBJ_FREE) { - // this happens - } else { - if (thneed->debug >= 1) { - printf("other ioctl %lx\n", request); - } - } - } - - int ret = my_ioctl(filedes, request, argp); - // NOTE: This error message goes into stdout and messes up pyenv - // if (ret != 0) printf("ioctl returned %d with errno %d\n", ret, errno); - return ret; -} - -} - -// *********** GPUMalloc *********** - -GPUMalloc::GPUMalloc(int size, int fd) { - struct kgsl_gpuobj_alloc alloc; - memset(&alloc, 0, sizeof(alloc)); - alloc.size = size; - alloc.flags = 0x10000a00; - ioctl(fd, IOCTL_KGSL_GPUOBJ_ALLOC, &alloc); - void *addr = mmap64(NULL, alloc.mmapsize, 0x3, 0x1, fd, alloc.id*0x1000); - assert(addr != MAP_FAILED); - - base = (uint64_t)addr; - remaining = size; -} - -GPUMalloc::~GPUMalloc() { - // TODO: free the GPU malloced area -} - -void *GPUMalloc::alloc(int size) { - void *ret = (void*)base; - size = (size+0xff) & (~0xFF); - assert(size <= remaining); - remaining -= size; - base += size; - return ret; -} - -// *********** CachedSync, at the ioctl layer *********** - -void CachedSync::exec() { - struct kgsl_gpuobj_sync cmd; - - cmd.objs = (uint64_t)data.data(); - cmd.obj_len = data.length(); - cmd.count = data.length() / sizeof(struct kgsl_gpuobj_sync_obj); - - int ret = ioctl(thneed->fd, IOCTL_KGSL_GPUOBJ_SYNC, &cmd); - assert(ret == 0); -} - -// *********** CachedCommand, at the ioctl layer *********** - -CachedCommand::CachedCommand(Thneed *lthneed, struct kgsl_gpu_command *cmd) { - thneed = lthneed; - assert(cmd->numsyncs == 0); - - memcpy(&cache, cmd, sizeof(cache)); - - if (cmd->numcmds > 0) { - cmds = make_unique(cmd->numcmds); - memcpy(cmds.get(), (void *)cmd->cmdlist, sizeof(struct kgsl_command_object)*cmd->numcmds); - cache.cmdlist = (uint64_t)cmds.get(); - for (int i = 0; i < cmd->numcmds; i++) { - void *nn = thneed->ram->alloc(cmds[i].size); - memcpy(nn, (void*)cmds[i].gpuaddr, cmds[i].size); - cmds[i].gpuaddr = (uint64_t)nn; - } - } - - if (cmd->numobjs > 0) { - objs = make_unique(cmd->numobjs); - memcpy(objs.get(), (void *)cmd->objlist, sizeof(struct kgsl_command_object)*cmd->numobjs); - cache.objlist = (uint64_t)objs.get(); - for (int i = 0; i < cmd->numobjs; i++) { - void *nn = thneed->ram->alloc(objs[i].size); - memset(nn, 0, objs[i].size); - objs[i].gpuaddr = (uint64_t)nn; - } - } - - kq = thneed->ckq; - thneed->ckq.clear(); -} - -void CachedCommand::exec() { - cache.timestamp = ++thneed->timestamp; - int ret = ioctl(thneed->fd, IOCTL_KGSL_GPU_COMMAND, &cache); - - if (thneed->debug >= 1) printf("CachedCommand::exec got %d\n", ret); - - if (thneed->debug >= 2) { - for (auto &it : kq) { - it->debug_print(false); - } - } - - assert(ret == 0); -} - -// *********** Thneed *********** - -Thneed::Thneed(bool do_clinit, cl_context _context) { - // TODO: QCOM2 actually requires a different context - //context = _context; - if (do_clinit) clinit(); - assert(g_fd != -1); - fd = g_fd; - ram = make_unique(0x80000, fd); - timestamp = -1; - g_thneed = this; - char *thneed_debug_env = getenv("THNEED_DEBUG"); - debug = (thneed_debug_env != NULL) ? atoi(thneed_debug_env) : 0; -} - -void Thneed::wait() { - struct kgsl_device_waittimestamp_ctxtid wait; - wait.context_id = context_id; - wait.timestamp = timestamp; - wait.timeout = -1; - - uint64_t tb = nanos_since_boot(); - int wret = ioctl(fd, IOCTL_KGSL_DEVICE_WAITTIMESTAMP_CTXTID, &wait); - uint64_t te = nanos_since_boot(); - - if (debug >= 1) printf("wait %d after %lu us\n", wret, (te-tb)/1000); -} - -void Thneed::execute(float **finputs, float *foutput, bool slow) { - uint64_t tb, te; - if (debug >= 1) tb = nanos_since_boot(); - - // ****** copy inputs - copy_inputs(finputs, true); - - // ****** run commands - int i = 0; - for (auto &it : cmds) { - ++i; - if (debug >= 1) printf("run %2d @ %7lu us: ", i, (nanos_since_boot()-tb)/1000); - it->exec(); - if ((i == cmds.size()) || slow) wait(); - } - - // ****** copy outputs - copy_output(foutput); - - if (debug >= 1) { - te = nanos_since_boot(); - printf("model exec in %lu us\n", (te-tb)/1000); - } -} diff --git a/selfdrive/pandad/panda_comms.cc b/selfdrive/pandad/panda_comms.cc index 59908f0cde..8a20f397d3 100644 --- a/selfdrive/pandad/panda_comms.cc +++ b/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 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; } diff --git a/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_calibrationd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_calibrationd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_calibrationd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_calibrationd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst b/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst new file mode 100644 index 0000000000..cf5d085a89 --- /dev/null +++ b/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:737f95d34912db53a303ba6499e6f697b510fa5872b8c71f701a4fe924b5466e +size 356169 diff --git a/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst deleted file mode 100644 index 9f6f2a34b0..0000000000 --- a/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst +++ /dev/null @@ -1,3 +0,0 @@ -version https://git-lfs.github.com/spec/v1 -oid sha256:6267aadf44c81fb003130aff6f8902014a1add7f18c5c563239943ab015974c9 -size 356177 diff --git a/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_controlsd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_controlsd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_controlsd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_controlsd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_dmonitoringd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_dmonitoringd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_dmonitoringd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_dmonitoringd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_locationd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_locationd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_locationd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_locationd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_paramsd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_paramsd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_paramsd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_paramsd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_plannerd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst b/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_plannerd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst new file mode 100644 index 0000000000..23421b51ab --- /dev/null +++ b/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_plannerd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:e359e8f6b5a22b6f3f89b54989dac2110ee3a4463de2d785be83e20cda4f1cb6 +size 254248 diff --git a/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_plannerd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_plannerd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst deleted file mode 100644 index 1e5248fd79..0000000000 --- a/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_plannerd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst +++ /dev/null @@ -1,3 +0,0 @@ -version https://git-lfs.github.com/spec/v1 -oid sha256:433633d2042cfe1191e5b71ec4b34edca21877377f2ace8aaf6ae43f7b805171 -size 256336 diff --git a/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_radard_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_radard_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_radard_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_radard_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_selfdrived_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_selfdrived_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_selfdrived_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_selfdrived_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_torqued_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_torqued_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_torqued_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_torqued_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_ubloxd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_ubloxd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_ubloxd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regen1CA7A48E6F7|2024-08-30--02-45-08--0_ubloxd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regen306779F6870|2024-10-03--04-03-23--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst b/selfdrive/test/process_replay/fakedata/regen306779F6870|2024-10-03--04-03-23--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst new file mode 100644 index 0000000000..e396d23e11 --- /dev/null +++ b/selfdrive/test/process_replay/fakedata/regen306779F6870|2024-10-03--04-03-23--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:a8e4d044812a714ebdf0b15e73d4466e9ddaafa374368f308803c6b68dcd79ab +size 332433 diff --git a/selfdrive/test/process_replay/fakedata/regen306779F6870|2024-10-03--04-03-23--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen306779F6870|2024-10-03--04-03-23--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst deleted file mode 100644 index 7990a4bfc6..0000000000 --- a/selfdrive/test/process_replay/fakedata/regen306779F6870|2024-10-03--04-03-23--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst +++ /dev/null @@ -1,3 +0,0 @@ -version https://git-lfs.github.com/spec/v1 -oid sha256:64aa7c86bf420adfcf78888bfce7233d68ec30ed77f3f9d3f923f26c59f85c74 -size 332371 diff --git a/selfdrive/test/process_replay/fakedata/regen4B38A7428CD|2024-08-30--02-56-31--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst b/selfdrive/test/process_replay/fakedata/regen4B38A7428CD|2024-08-30--02-56-31--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst new file mode 100644 index 0000000000..e44ef1aeed --- /dev/null +++ b/selfdrive/test/process_replay/fakedata/regen4B38A7428CD|2024-08-30--02-56-31--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:53b80c3c99a6897cafe7d408872c287ab0bbaac2751e344cf4623b312d2e4866 +size 268928 diff --git a/selfdrive/test/process_replay/fakedata/regen4B38A7428CD|2024-08-30--02-56-31--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen4B38A7428CD|2024-08-30--02-56-31--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst deleted file mode 100644 index d80463af60..0000000000 --- a/selfdrive/test/process_replay/fakedata/regen4B38A7428CD|2024-08-30--02-56-31--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst +++ /dev/null @@ -1,3 +0,0 @@ -version https://git-lfs.github.com/spec/v1 -oid sha256:aa2451a6697da9fb1cf305081337782eb944a623b7662fd0a58ebbed5d0395e7 -size 268892 diff --git a/selfdrive/test/process_replay/fakedata/regen4CE950B0267|2024-08-30--02-51-30--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst b/selfdrive/test/process_replay/fakedata/regen4CE950B0267|2024-08-30--02-51-30--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst new file mode 100644 index 0000000000..af4d4d784c --- /dev/null +++ b/selfdrive/test/process_replay/fakedata/regen4CE950B0267|2024-08-30--02-51-30--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:a4fa8a841c964e90b9c14b5aede8f30de949d319ae072cc291f0afb1c4c24baa +size 437808 diff --git a/selfdrive/test/process_replay/fakedata/regen4CE950B0267|2024-08-30--02-51-30--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen4CE950B0267|2024-08-30--02-51-30--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst deleted file mode 100644 index 532267f974..0000000000 --- a/selfdrive/test/process_replay/fakedata/regen4CE950B0267|2024-08-30--02-51-30--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst +++ /dev/null @@ -1,3 +0,0 @@ -version https://git-lfs.github.com/spec/v1 -oid sha256:75766b120bc372bad7b2a80be2c195de9b85cf11142f2a9da348637f15e245db -size 437801 diff --git a/selfdrive/test/process_replay/fakedata/regen58464878D07|2024-08-30--03-15-31--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst b/selfdrive/test/process_replay/fakedata/regen58464878D07|2024-08-30--03-15-31--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst new file mode 100644 index 0000000000..3abf95e8ac --- /dev/null +++ b/selfdrive/test/process_replay/fakedata/regen58464878D07|2024-08-30--03-15-31--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:ce370994de01d6240fa753846ddc7e1b852acac3f649a4c29ea16acae126f974 +size 308578 diff --git a/selfdrive/test/process_replay/fakedata/regen58464878D07|2024-08-30--03-15-31--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen58464878D07|2024-08-30--03-15-31--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst deleted file mode 100644 index e4b083c7e1..0000000000 --- a/selfdrive/test/process_replay/fakedata/regen58464878D07|2024-08-30--03-15-31--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst +++ /dev/null @@ -1,3 +0,0 @@ -version https://git-lfs.github.com/spec/v1 -oid sha256:c692e18506a36eeed22d23d1739af8f26fb3f39f34584fdbcb1ebf3bc6430957 -size 308611 diff --git a/selfdrive/test/process_replay/fakedata/regen6E484EDAB96|2024-08-30--02-47-37--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst b/selfdrive/test/process_replay/fakedata/regen6E484EDAB96|2024-08-30--02-47-37--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst new file mode 100644 index 0000000000..f6f348be88 --- /dev/null +++ b/selfdrive/test/process_replay/fakedata/regen6E484EDAB96|2024-08-30--02-47-37--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:82333cfc026735eac81defae9e01b82b81ac30fd01ad265dbdd311dd97322198 +size 393106 diff --git a/selfdrive/test/process_replay/fakedata/regen6E484EDAB96|2024-08-30--02-47-37--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen6E484EDAB96|2024-08-30--02-47-37--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst deleted file mode 100644 index dcd535b13d..0000000000 --- a/selfdrive/test/process_replay/fakedata/regen6E484EDAB96|2024-08-30--02-47-37--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst +++ /dev/null @@ -1,3 +0,0 @@ -version https://git-lfs.github.com/spec/v1 -oid sha256:8d62d5774bb4be70ef337179c4bab2347f153fc959155bef917527cc6868bcd3 -size 393161 diff --git a/selfdrive/test/process_replay/fakedata/regen720F2BA4CF6|2024-08-30--03-09-15--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst b/selfdrive/test/process_replay/fakedata/regen720F2BA4CF6|2024-08-30--03-09-15--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst new file mode 100644 index 0000000000..fb3b707b85 --- /dev/null +++ b/selfdrive/test/process_replay/fakedata/regen720F2BA4CF6|2024-08-30--03-09-15--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:e4ebae1742e3bb7d89f8cd452928c5c6a1b8679155562d23c6303ca4ec2e3b02 +size 334350 diff --git a/selfdrive/test/process_replay/fakedata/regen720F2BA4CF6|2024-08-30--03-09-15--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen720F2BA4CF6|2024-08-30--03-09-15--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst deleted file mode 100644 index 0987d4e6b3..0000000000 --- a/selfdrive/test/process_replay/fakedata/regen720F2BA4CF6|2024-08-30--03-09-15--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst +++ /dev/null @@ -1,3 +0,0 @@ -version https://git-lfs.github.com/spec/v1 -oid sha256:d57c998320672e771712bd0b3f35c07bd159dff472368b1bf1abc9b6a3541dad -size 334258 diff --git a/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_calibrationd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_calibrationd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_calibrationd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_calibrationd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst b/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst new file mode 100644 index 0000000000..cef3e0150c --- /dev/null +++ b/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:03326fb4486c1ffdd3609a0e0200e65d1f30dca5d9b047501d98b1b2d5321e75 +size 470495 diff --git a/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst deleted file mode 100644 index 2321b17bd1..0000000000 --- a/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst +++ /dev/null @@ -1,3 +0,0 @@ -version https://git-lfs.github.com/spec/v1 -oid sha256:fe9370fb0f7cad039d12ddc2f23dd4ee1c94d5d75faf9b6143ec390c06b28c9d -size 470471 diff --git a/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_controlsd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_controlsd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_controlsd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_controlsd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_dmonitoringd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_dmonitoringd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_dmonitoringd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_dmonitoringd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_locationd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_locationd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_locationd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_locationd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_paramsd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_paramsd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_paramsd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_paramsd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_plannerd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst b/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_plannerd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst new file mode 100644 index 0000000000..3a13d7e7e2 --- /dev/null +++ b/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_plannerd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:bffa062a5eace4d499c2fe12f2b0eb8c71207f4e7bea0e88456614f00e2859b1 +size 258989 diff --git a/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_plannerd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_plannerd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst deleted file mode 100644 index a2bf931178..0000000000 --- a/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_plannerd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst +++ /dev/null @@ -1,3 +0,0 @@ -version https://git-lfs.github.com/spec/v1 -oid sha256:917142016b6ed8c800a2c00aecbfdba7eecda3b1a3fcb15a968fa726b5b9241c -size 260324 diff --git a/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_radard_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_radard_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_radard_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_radard_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_selfdrived_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_selfdrived_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_selfdrived_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_selfdrived_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_torqued_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_torqued_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_torqued_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_torqued_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_ubloxd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_ubloxd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_ubloxd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regen756F8230C21|2024-11-07--00-08-24--0_ubloxd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regen9ADBECBCD1C|2024-08-30--03-13-04--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst b/selfdrive/test/process_replay/fakedata/regen9ADBECBCD1C|2024-08-30--03-13-04--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst new file mode 100644 index 0000000000..12c275e0b7 --- /dev/null +++ b/selfdrive/test/process_replay/fakedata/regen9ADBECBCD1C|2024-08-30--03-13-04--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:fe97d060fb6e60652b26b26406ca8b03890f6bff4b4bc5cb73fca267068a04c9 +size 217460 diff --git a/selfdrive/test/process_replay/fakedata/regen9ADBECBCD1C|2024-08-30--03-13-04--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen9ADBECBCD1C|2024-08-30--03-13-04--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst deleted file mode 100644 index 83abf15339..0000000000 --- a/selfdrive/test/process_replay/fakedata/regen9ADBECBCD1C|2024-08-30--03-13-04--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst +++ /dev/null @@ -1,3 +0,0 @@ -version https://git-lfs.github.com/spec/v1 -oid sha256:d4b0e05a8a1b21e008ed5c1a41f267d596cc35a6d49ec8eff496856b34866360 -size 217481 diff --git a/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_calibrationd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_calibrationd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_calibrationd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_calibrationd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst b/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst new file mode 100644 index 0000000000..7714c7931f --- /dev/null +++ b/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:60e33d209f4f600acf344feca8ab6494a1bf3aafe3c7121e2110a82ea06aa61e +size 293084 diff --git a/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst deleted file mode 100644 index 228b01854a..0000000000 --- a/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst +++ /dev/null @@ -1,3 +0,0 @@ -version https://git-lfs.github.com/spec/v1 -oid sha256:4d8e2889e70c310a84eff142ee481cfc92888bbe257809be3694032c67277675 -size 293076 diff --git a/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_controlsd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_controlsd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_controlsd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_controlsd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_dmonitoringd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_dmonitoringd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_dmonitoringd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_dmonitoringd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_locationd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_locationd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_locationd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_locationd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_paramsd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_paramsd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_paramsd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_paramsd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_plannerd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst b/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_plannerd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst new file mode 100644 index 0000000000..22350b93c7 --- /dev/null +++ b/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_plannerd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:754d601883db2377739f4112a8dff787b4b8181a8bb8bf4d7dbcf77c5cb6ac93 +size 256838 diff --git a/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_plannerd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_plannerd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst deleted file mode 100644 index 6fbe7b3aec..0000000000 --- a/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_plannerd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst +++ /dev/null @@ -1,3 +0,0 @@ -version https://git-lfs.github.com/spec/v1 -oid sha256:d60f784a226d144e1ac2c9ceec53b30509cbcceffb5e28f9c1b3a88163958352 -size 260706 diff --git a/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_radard_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_radard_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_radard_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_radard_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_selfdrived_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_selfdrived_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_selfdrived_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_selfdrived_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_torqued_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_torqued_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_torqued_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_torqued_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_ubloxd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_ubloxd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_ubloxd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regen9CBD921E93E|2024-08-30--02-38-51--0_ubloxd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regenA67A128BCD8|2024-08-30--02-36-22--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst b/selfdrive/test/process_replay/fakedata/regenA67A128BCD8|2024-08-30--02-36-22--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst new file mode 100644 index 0000000000..1bd5a628a9 --- /dev/null +++ b/selfdrive/test/process_replay/fakedata/regenA67A128BCD8|2024-08-30--02-36-22--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:378d6fc5c674e62b51a12b994cfcd03c435a9731c6e2f4c0d14a39010540cc77 +size 100067 diff --git a/selfdrive/test/process_replay/fakedata/regenA67A128BCD8|2024-08-30--02-36-22--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regenA67A128BCD8|2024-08-30--02-36-22--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst deleted file mode 100644 index 811af99c44..0000000000 --- a/selfdrive/test/process_replay/fakedata/regenA67A128BCD8|2024-08-30--02-36-22--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst +++ /dev/null @@ -1,3 +0,0 @@ -version https://git-lfs.github.com/spec/v1 -oid sha256:dc8216930cf4b36dd080f32b6f52172e08c2c4fe946450162950e799345d1829 -size 100106 diff --git a/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_calibrationd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_calibrationd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_calibrationd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_calibrationd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst b/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst new file mode 100644 index 0000000000..467c8614c7 --- /dev/null +++ b/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:037825cd9b443d038ab3bb370eab802612d1ce528c5de1350fd99965eaa0d131 +size 262910 diff --git a/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst deleted file mode 100644 index 658ad8a62e..0000000000 --- a/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst +++ /dev/null @@ -1,3 +0,0 @@ -version https://git-lfs.github.com/spec/v1 -oid sha256:cb0195a176d9a917ef87698ef6487547ef19f35e86b4ba12bbecc7b5e6a26295 -size 262860 diff --git a/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_controlsd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_controlsd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_controlsd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_controlsd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_dmonitoringd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_dmonitoringd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_dmonitoringd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_dmonitoringd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_locationd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_locationd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_locationd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_locationd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_paramsd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_paramsd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_paramsd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_paramsd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_plannerd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst b/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_plannerd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst new file mode 100644 index 0000000000..8c3a0bff12 --- /dev/null +++ b/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_plannerd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:58645abd939aa17cf291a254517865704c74124743a8a77da296c519a29567df +size 256653 diff --git a/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_plannerd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_plannerd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst deleted file mode 100644 index ce87ea8d54..0000000000 --- a/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_plannerd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst +++ /dev/null @@ -1,3 +0,0 @@ -version https://git-lfs.github.com/spec/v1 -oid sha256:3ddaa636dc318b33e356787852aad7559c656ce55469c49b7b97fae093dc7889 -size 259081 diff --git a/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_radard_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_radard_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_radard_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_radard_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_selfdrived_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_selfdrived_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_selfdrived_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_selfdrived_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_torqued_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_torqued_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_torqued_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_torqued_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_ubloxd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_ubloxd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_ubloxd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regenAA1FF48CF1F|2024-08-30--03-06-45--0_ubloxd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regenACF84CCF482|2024-08-30--03-21-55--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst b/selfdrive/test/process_replay/fakedata/regenACF84CCF482|2024-08-30--03-21-55--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst new file mode 100644 index 0000000000..17c59881b2 --- /dev/null +++ b/selfdrive/test/process_replay/fakedata/regenACF84CCF482|2024-08-30--03-21-55--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:fed5fa87ab0dfbf410b5a142e98eae9dacdb1a7c2ec38f91ba5ae18d8ad9a9e0 +size 268234 diff --git a/selfdrive/test/process_replay/fakedata/regenACF84CCF482|2024-08-30--03-21-55--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regenACF84CCF482|2024-08-30--03-21-55--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst deleted file mode 100644 index 951836bcee..0000000000 --- a/selfdrive/test/process_replay/fakedata/regenACF84CCF482|2024-08-30--03-21-55--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst +++ /dev/null @@ -1,3 +0,0 @@ -version https://git-lfs.github.com/spec/v1 -oid sha256:f282a05c43be27ca4b4d522e5dcdddf8597ae92418328e951f7fef74a5a82cbc -size 268145 diff --git a/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_calibrationd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_calibrationd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_calibrationd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_calibrationd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst b/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst new file mode 100644 index 0000000000..4ed688afbd --- /dev/null +++ b/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:f8affdd98bd39d326a7d66e5bca4028ac81595a3f55a6fc094e8b93f535eab8f +size 305144 diff --git a/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst deleted file mode 100644 index b921150f44..0000000000 --- a/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst +++ /dev/null @@ -1,3 +0,0 @@ -version https://git-lfs.github.com/spec/v1 -oid sha256:7ef1a55cad5cede01685360813781d30ee576a9f3d201f5caf1144327688fc0b -size 305127 diff --git a/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_controlsd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_controlsd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_controlsd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_controlsd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_dmonitoringd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_dmonitoringd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_dmonitoringd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_dmonitoringd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_locationd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_locationd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_locationd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_locationd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_paramsd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_paramsd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_paramsd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_paramsd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_plannerd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst b/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_plannerd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst new file mode 100644 index 0000000000..141c4d4d4c --- /dev/null +++ b/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_plannerd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:d82c14c33dc91196a5febc1c45d307a157d5b64b4eb1f8c880fa319af79bc5f5 +size 253450 diff --git a/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_plannerd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_plannerd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst deleted file mode 100644 index 1fe0b3b5a2..0000000000 --- a/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_plannerd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst +++ /dev/null @@ -1,3 +0,0 @@ -version https://git-lfs.github.com/spec/v1 -oid sha256:055486d036c48811a0e0c4e461a7a1bdf66e3bef2fb6afdaf88aefb7b63a4f21 -size 255708 diff --git a/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_radard_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_radard_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_radard_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_radard_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_selfdrived_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_selfdrived_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_selfdrived_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_selfdrived_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_torqued_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_torqued_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_torqued_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_torqued_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_ubloxd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_ubloxd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst similarity index 100% rename from selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_ubloxd_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst rename to selfdrive/test/process_replay/fakedata/regenC8F0D6ADC5C|2024-08-30--02-54-01--0_ubloxd_a20d7a427906c43549cf58b7a859f264aabdf28d.zst diff --git a/selfdrive/test/process_replay/fakedata/regenDB02684E00A|2024-08-30--03-02-54--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst b/selfdrive/test/process_replay/fakedata/regenDB02684E00A|2024-08-30--03-02-54--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst new file mode 100644 index 0000000000..15edf70fa3 --- /dev/null +++ b/selfdrive/test/process_replay/fakedata/regenDB02684E00A|2024-08-30--03-02-54--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:43150b3ef31b8cd8193150e95e7663a5cd0596b2a0fc9c374861d6c7bb761dbd +size 271983 diff --git a/selfdrive/test/process_replay/fakedata/regenDB02684E00A|2024-08-30--03-02-54--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regenDB02684E00A|2024-08-30--03-02-54--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst deleted file mode 100644 index 7ddb8e3c11..0000000000 --- a/selfdrive/test/process_replay/fakedata/regenDB02684E00A|2024-08-30--03-02-54--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst +++ /dev/null @@ -1,3 +0,0 @@ -version https://git-lfs.github.com/spec/v1 -oid sha256:45a953af023b1900aab038c5e8731356c77f7716e36be63cd24b823ce5c34d04 -size 271972 diff --git a/selfdrive/test/process_replay/fakedata/regenED976DEB757|2024-08-30--03-18-02--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst b/selfdrive/test/process_replay/fakedata/regenED976DEB757|2024-08-30--03-18-02--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst new file mode 100644 index 0000000000..df781bb4de --- /dev/null +++ b/selfdrive/test/process_replay/fakedata/regenED976DEB757|2024-08-30--03-18-02--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:aec553e5975a5fe7263b93261f44f909af664e7e9e704fd3d95b41766cbc1a81 +size 332401 diff --git a/selfdrive/test/process_replay/fakedata/regenED976DEB757|2024-08-30--03-18-02--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regenED976DEB757|2024-08-30--03-18-02--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst deleted file mode 100644 index 1bb4bd49e7..0000000000 --- a/selfdrive/test/process_replay/fakedata/regenED976DEB757|2024-08-30--03-18-02--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst +++ /dev/null @@ -1,3 +0,0 @@ -version https://git-lfs.github.com/spec/v1 -oid sha256:82bdd9410a4300cbda41cb59a3038b2799d6dec79691dd51717f09b3e3863db6 -size 332408 diff --git a/selfdrive/test/process_replay/fakedata/regenF3DBBA9E8DF|2024-08-30--02-59-03--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst b/selfdrive/test/process_replay/fakedata/regenF3DBBA9E8DF|2024-08-30--02-59-03--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst new file mode 100644 index 0000000000..c2388f98ea --- /dev/null +++ b/selfdrive/test/process_replay/fakedata/regenF3DBBA9E8DF|2024-08-30--02-59-03--0_card_a20d7a427906c43549cf58b7a859f264aabdf28d.zst @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:b85e323fb1c29b15f488ebb18e8f81f1e70996515d722c9cc0a7a466ea9119ae +size 294529 diff --git a/selfdrive/test/process_replay/fakedata/regenF3DBBA9E8DF|2024-08-30--02-59-03--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst b/selfdrive/test/process_replay/fakedata/regenF3DBBA9E8DF|2024-08-30--02-59-03--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst deleted file mode 100644 index 4d9bee43c6..0000000000 --- a/selfdrive/test/process_replay/fakedata/regenF3DBBA9E8DF|2024-08-30--02-59-03--0_card_fffa98ee8531a24a34a2141c9017193e0a26f50d.zst +++ /dev/null @@ -1,3 +0,0 @@ -version https://git-lfs.github.com/spec/v1 -oid sha256:0e7eb9dcadd6ce9a39d9eb9a2b325944acb55dd791f9f3b9579ac232da21776b -size 294509 diff --git a/selfdrive/test/process_replay/migration.py b/selfdrive/test/process_replay/migration.py index 5376f91aee..d531b3b7ed 100644 --- a/selfdrive/test/process_replay/migration.py +++ b/selfdrive/test/process_replay/migration.py @@ -107,7 +107,7 @@ def migrate_longitudinalPlan(msgs): if msg.which() != 'longitudinalPlan': continue new_msg = msg.as_builder() - new_msg.longitudinalPlan.aTarget, new_msg.longitudinalPlan.shouldStop = get_accel_from_plan(CP, msg.longitudinalPlan.speeds, msg.longitudinalPlan.accels) + new_msg.longitudinalPlan.aTarget, new_msg.longitudinalPlan.shouldStop = get_accel_from_plan(msg.longitudinalPlan.speeds, msg.longitudinalPlan.accels) ops.append((index, new_msg.as_reader())) return ops, [], [] diff --git a/selfdrive/test/process_replay/ref_commit b/selfdrive/test/process_replay/ref_commit index 7a622fab98..99c3356cf9 100644 --- a/selfdrive/test/process_replay/ref_commit +++ b/selfdrive/test/process_replay/ref_commit @@ -1 +1 @@ -fffa98ee8531a24a34a2141c9017193e0a26f50d \ No newline at end of file +a20d7a427906c43549cf58b7a859f264aabdf28d \ No newline at end of file diff --git a/selfdrive/test/test_onroad.py b/selfdrive/test/test_onroad.py index 7b1519a3d3..25aaa95f0f 100644 --- a/selfdrive/test/test_onroad.py +++ b/selfdrive/test/test_onroad.py @@ -36,7 +36,7 @@ CPU usage budget TEST_DURATION = 25 LOG_OFFSET = 8 -MAX_TOTAL_CPU = 265. # total for all 8 cores +MAX_TOTAL_CPU = 275. # total for all 8 cores PROCS = { # Baseline CPU usage by process "selfdrive.controls.controlsd": 16.0, @@ -50,8 +50,8 @@ PROCS = { "selfdrive.locationd.paramsd": 9.0, "./sensord": 7.0, "selfdrive.controls.radard": 2.0, - "selfdrive.modeld.modeld": 17.0, - "selfdrive.modeld.dmonitoringmodeld": 11.0, + "selfdrive.modeld.modeld": 22.0, + "selfdrive.modeld.dmonitoringmodeld": 21.0, "system.hardware.hardwared": 4.0, "selfdrive.locationd.calibrationd": 2.0, "selfdrive.locationd.torqued": 5.0, @@ -371,10 +371,9 @@ class TestOnroad: result += "------------------------------------------------\n" result += "----------------- Model Timing -----------------\n" result += "------------------------------------------------\n" - # TODO: this went up when plannerd cpu usage increased, why? cfgs = [ - ("modelV2", 0.050, 0.036), - ("driverStateV2", 0.050, 0.026), + ("modelV2", 0.045, 0.035), + ("driverStateV2", 0.045, 0.035), ] for (s, instant_max, avg_max) in cfgs: ts = [getattr(m, s).modelExecutionTime for m in self.msgs[s]] diff --git a/selfdrive/ui/qt/network/wifi_manager.cc b/selfdrive/ui/qt/network/wifi_manager.cc index fd9ee17078..1717210634 100644 --- a/selfdrive/ui/qt/network/wifi_manager.cc +++ b/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 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(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> response = call(path.path(), NM_DBUS_INTERFACE_SETTINGS_CONNECTION, "GetSecrets", "802-11-wireless-security"); diff --git a/system/camerad/cameras/process_raw.cl b/system/camerad/cameras/process_raw.cl index 6f6612fab0..ff6060d855 100644 --- a/system/camerad/cameras/process_raw.cl +++ b/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 diff --git a/system/camerad/sensors/ar0231_cl.h b/system/camerad/sensors/ar0231_cl.h index 8e96bbce5b..c79242543b 100644 --- a/system/camerad/sensors/ar0231_cl.h +++ b/system/camerad/sensors/ar0231_cl.h @@ -1,9 +1,10 @@ #if SENSOR_ID == 1 +#define VIGNETTE_PROFILE_8DT0MM + #define BIT_DEPTH 12 #define PV_MAX 4096 #define BLACK_LVL 168 -#define VIGNETTE_RSZ 1.0f float4 normalize_pv(int4 parsed, float vignette_factor) { float4 pv = (convert_float4(parsed) - BLACK_LVL) / (PV_MAX - BLACK_LVL); @@ -30,4 +31,4 @@ float3 apply_gamma(float3 rgb, int expo_time) { ((rk * (rgb-mp) * (gamma_k*mp+gamma_b) * (1+1/(rk*mp)) / (1-rk*(rgb-mp))) + gamma_k*mp + gamma_b); } -#endif \ No newline at end of file +#endif diff --git a/system/camerad/sensors/os04c10_cl.h b/system/camerad/sensors/os04c10_cl.h index 3da3fab8df..3b5cf88839 100644 --- a/system/camerad/sensors/os04c10_cl.h +++ b/system/camerad/sensors/os04c10_cl.h @@ -1,13 +1,13 @@ #if SENSOR_ID == 3 #define BGGR +#define VIGNETTE_PROFILE_4DT6MM #define BIT_DEPTH 12 #define PV_MAX10 1023 #define PV_MAX12 4095 #define PV_MAX16 65536 // gamma curve is calibrated to 16bit #define BLACK_LVL 48 -#define VIGNETTE_RSZ 2.2545f float combine_dual_pvs(float lv, float sv, int expo_time) { float svc = fmax(sv * expo_time, (float)(64 * (PV_MAX10 - BLACK_LVL))); diff --git a/system/camerad/sensors/ox03c10_cl.h b/system/camerad/sensors/ox03c10_cl.h index f1fbd16ccb..c8cec7cf8a 100644 --- a/system/camerad/sensors/ox03c10_cl.h +++ b/system/camerad/sensors/ox03c10_cl.h @@ -1,8 +1,9 @@ #if SENSOR_ID == 2 +#define VIGNETTE_PROFILE_8DT0MM + #define BIT_DEPTH 12 #define BLACK_LVL 64 -#define VIGNETTE_RSZ 1.0f float ox_lut_func(int x) { if (x < 512) { diff --git a/system/hardware/tici/tests/test_power_draw.py b/system/hardware/tici/tests/test_power_draw.py index 22fd70e1a5..e1b9845c4c 100644 --- a/system/hardware/tici/tests/test_power_draw.py +++ b/system/hardware/tici/tests/test_power_draw.py @@ -33,7 +33,7 @@ class Proc: PROCS = [ Proc(['camerad'], 1.75, msgs=['roadCameraState', 'wideRoadCameraState', 'driverCameraState']), Proc(['modeld'], 1.12, atol=0.2, msgs=['modelV2']), - Proc(['dmonitoringmodeld'], 0.65, msgs=['driverStateV2']), + Proc(['dmonitoringmodeld'], 0.6, msgs=['driverStateV2']), Proc(['encoderd'], 0.23, msgs=[]), ] diff --git a/system/manager/process_config.py b/system/manager/process_config.py index d26fc19afe..a25be615f4 100644 --- a/system/manager/process_config.py +++ b/system/manager/process_config.py @@ -70,10 +70,12 @@ procs = [ PythonProcess("micd", "system.micd", iscar), PythonProcess("timed", "system.timed", always_run, enabled=not PC), + # TODO Make python process once TG allows opening QCOM from child proc NativeProcess("dmonitoringmodeld", "selfdrive/modeld", ["./dmonitoringmodeld"], driverview, enabled=(not PC or WEBCAM)), NativeProcess("encoderd", "system/loggerd", ["./encoderd"], only_onroad), NativeProcess("stream_encoderd", "system/loggerd", ["./encoderd", "--stream"], notcar), NativeProcess("loggerd", "system/loggerd", ["./loggerd"], logging), + # TODO Make python process once TG allows opening QCOM from child proc NativeProcess("modeld", "selfdrive/modeld", ["./modeld"], only_onroad), NativeProcess("sensord", "system/sensord", ["./sensord"], only_onroad, enabled=not PC), NativeProcess("ui", "selfdrive/ui", ["./ui"], always_run, watchdog_max_dt=(5 if not PC else None)), diff --git a/tinygrad_repo b/tinygrad_repo index 9dda6d260d..480e5e7a12 160000 --- a/tinygrad_repo +++ b/tinygrad_repo @@ -1 +1 @@ -Subproject commit 9dda6d260db0255750bacff61e3cee1e580567e1 +Subproject commit 480e5e7a1292bf2f84e18edffd06a985c4b48e65 diff --git a/tools/cabana/tools/findsignal.cc b/tools/cabana/tools/findsignal.cc index 8e1749ba8e..ec56fcaac0 100644 --- a/tools/cabana/tools/findsignal.cc +++ b/tools/cabana/tools/findsignal.cc @@ -225,7 +225,7 @@ void FindSignalDlg::setInitialSignals() { model->initial_signals.clear(); for (const auto &[id, m] : can->lastMessages()) { - if (buses.isEmpty() || buses.contains(id.source) && (addresses.isEmpty() || addresses.contains(id.address))) { + if ((buses.isEmpty() || buses.contains(id.source)) && (addresses.isEmpty() || addresses.contains(id.address))) { const auto &events = can->events(id); auto e = std::lower_bound(events.cbegin(), events.cend(), first_time, CompareCanEvent()); if (e != events.cend()) {