mirror of
https://github.com/infiniteCable2/openpilot.git
synced 2026-08-05 08:16:03 +08:00
modeld_v2: tinygrad transformation warp (#1698)
* chore: sync tinygrad Runs great in sim. now we need to rebuild some models * oops forgot to unblock this after testing * helpers * oh yeah * latest tg * this wont do anything empriically * reduce complexity * okay lint * Update tinygrad_runner.py * Update modeld.py * Update build-all-tinygrad-models.yaml * tinygrad bump * Update modeld.py * Update tinygrad_runner.py * bump * Update SConscript * Update SConscript * com * Update fetcher.py * Update helpers.py * life is froughtless, when you're thoughtless * lint * ozdust ballroom * shiz * Update tinygrad_runner.py * Update tinygrad_runner.py * slough it off as i do * try old support one last time * support mixed input dtypes * use internal * dont need that shiz * Update fill_model_msg.py * Update onnx_runner.py * Update onnx_runner.py * Update model_runner.py * see if this speeds up execution if not, revert me * no * Update helpers.py * rebase * more * planplus --------- Co-authored-by: Jason Wen <haibin.wen3@gmail.com>
This commit is contained in:
James Vecellio-Grant
committed by
GitHub
GitHub
parent
59a16b9cdc
commit
ddf05d7126
@@ -173,7 +173,7 @@ jobs:
|
||||
|
||||
echo "Compiling: $onnx_file -> $output_file"
|
||||
QCOM=1 python3 "${{ env.TINYGRAD_PATH }}/examples/openpilot/compile3.py" "$onnx_file" "$output_file"
|
||||
QCOM=1 python3 "${{ env.MODELS_DIR }}/../get_model_metadata.py" "$onnx_file" || true
|
||||
DEV=QCOM FLOAT16=1 NOLOCALS=1 JIT_BATCH_SIZE=0 python3 "${{ env.MODELS_DIR }}/../get_model_metadata.py" "$onnx_file" || true
|
||||
done
|
||||
|
||||
- name: Prepare Output
|
||||
|
||||
@@ -3,7 +3,7 @@ import time
|
||||
|
||||
from cereal import car, log, messaging
|
||||
from openpilot.common.params import Params
|
||||
from openpilot.system.manager.process_config import managed_processes, is_snpe_model, is_tinygrad_model, is_stock_model
|
||||
from openpilot.system.manager.process_config import managed_processes, is_tinygrad_model, is_stock_model
|
||||
from openpilot.system.hardware import HARDWARE
|
||||
|
||||
if __name__ == "__main__":
|
||||
@@ -11,8 +11,6 @@ if __name__ == "__main__":
|
||||
params = Params()
|
||||
params.put("CarParams", CP.to_bytes())
|
||||
|
||||
if use_snpe_modeld := is_snpe_model(False, params, CP):
|
||||
print("Using SNPE modeld")
|
||||
if use_tinygrad_modeld := is_tinygrad_model(False, params, CP):
|
||||
print("Using TinyGrad modeld")
|
||||
if use_stock_modeld := is_stock_model(False, params, CP):
|
||||
@@ -21,7 +19,7 @@ if __name__ == "__main__":
|
||||
HARDWARE.set_power_save(False)
|
||||
|
||||
procs = ['camerad', 'ui', 'calibrationd', 'plannerd', 'dmonitoringmodeld', 'dmonitoringd']
|
||||
procs += ["modeld_snpe" if use_snpe_modeld else "modeld_tinygrad" if use_tinygrad_modeld else "modeld"]
|
||||
procs += ["modeld_tinygrad" if use_tinygrad_modeld else "modeld"]
|
||||
for p in procs:
|
||||
managed_processes[p].start()
|
||||
|
||||
|
||||
@@ -1,4 +1,3 @@
|
||||
SConscript(['common/transformations/SConscript'])
|
||||
SConscript(['modeld/SConscript'])
|
||||
SConscript(['modeld_v2/SConscript'])
|
||||
SConscript(['selfdrive/locationd/SConscript'])
|
||||
|
||||
@@ -1,34 +1,8 @@
|
||||
import os
|
||||
import glob
|
||||
|
||||
Import('env', 'envCython', 'arch', 'cereal', 'messaging', 'common', 'visionipc')
|
||||
Import('env', 'arch')
|
||||
lenv = env.Clone()
|
||||
lenvCython = envCython.Clone()
|
||||
|
||||
libs = [cereal, messaging, visionipc, common, 'capnp', 'kj', 'pthread']
|
||||
frameworks = []
|
||||
|
||||
common_src = [
|
||||
"models/commonmodel.cc",
|
||||
"transforms/loadyuv.cc",
|
||||
"transforms/transform.cc",
|
||||
]
|
||||
|
||||
# OpenCL is a framework on Mac
|
||||
if arch == "Darwin":
|
||||
frameworks += ['OpenCL']
|
||||
else:
|
||||
libs += ['OpenCL']
|
||||
|
||||
# Set path definitions
|
||||
for pathdef, fn in {'TRANSFORM': 'transforms/transform.cl', 'LOADYUV': 'transforms/loadyuv.cl'}.items():
|
||||
for xenv in (lenv, lenvCython):
|
||||
xenv['CXXFLAGS'].append(f'-D{pathdef}_PATH=\\"{File(fn).abspath}\\"')
|
||||
|
||||
# Compile cython
|
||||
cython_libs = envCython["LIBS"] + libs
|
||||
commonmodel_lib = lenv.Library('commonmodel', common_src)
|
||||
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) if 'pycache' not in x]
|
||||
|
||||
# Get model metadata
|
||||
@@ -47,20 +21,39 @@ if PC:
|
||||
if outputs:
|
||||
lenv.Command(outputs, inputs, cmd)
|
||||
|
||||
tg_flags = {
|
||||
'larch64': 'DEV=QCOM FLOAT16=1 NOLOCALS=1 JIT_BATCH_SIZE=0',
|
||||
'Darwin': f'DEV=CPU THREADS=0 HOME={os.path.expanduser("~")}',
|
||||
}.get(arch, 'DEV=CPU CPU_LLVM=1 THREADS=0')
|
||||
|
||||
image_flag = {
|
||||
'larch64': 'IMAGE=2',
|
||||
}.get(arch, 'IMAGE=0')
|
||||
|
||||
def tg_compile(flags, model_name):
|
||||
pythonpath_string = 'PYTHONPATH="${PYTHONPATH}:' + env.Dir("#tinygrad_repo").abspath + '"'
|
||||
fn = File(f"models/{model_name}").abspath
|
||||
out = fn + "_tinygrad.pkl"
|
||||
|
||||
return lenv.Command(
|
||||
fn + "_tinygrad.pkl",
|
||||
out,
|
||||
[fn + ".onnx"] + tinygrad_files,
|
||||
f'{pythonpath_string} {flags} python3 {Dir("#tinygrad_repo").abspath}/examples/openpilot/compile3.py {fn}.onnx {fn}_tinygrad.pkl'
|
||||
f'{pythonpath_string} {flags} {image_flag} python3 {Dir("#tinygrad_repo").abspath}/examples/openpilot/compile3.py {fn}.onnx {out}'
|
||||
)
|
||||
|
||||
# Compile small models
|
||||
# Compile models
|
||||
for model_name in ['supercombo', 'driving_vision', 'driving_off_policy', 'driving_policy']:
|
||||
if File(f"models/{model_name}.onnx").exists():
|
||||
flags = {
|
||||
'larch64': 'DEV=QCOM FLOAT16=1 NOLOCALS=1 IMAGE=2 JIT_BATCH_SIZE=0',
|
||||
'Darwin': f'DEV=CPU HOME={os.path.expanduser("~")} IMAGE=0', # tinygrad calls brew which needs a $HOME in the env
|
||||
}.get(arch, 'DEV=CPU CPU_LLVM=1 IMAGE=0')
|
||||
tg_compile(flags, model_name)
|
||||
tg_compile(tg_flags, model_name)
|
||||
|
||||
script_files = [File("warp.py"), File(Dir("#selfdrive/modeld").File("compile_warp.py").abspath)]
|
||||
pythonpath_string = 'PYTHONPATH="${PYTHONPATH}:' + env.Dir("#tinygrad_repo").abspath + ':' + env.Dir("#").abspath + '"'
|
||||
compile_warp_cmd = f'{pythonpath_string} {tg_flags} python3 -m sunnypilot.modeld_v2.warp'
|
||||
|
||||
from openpilot.common.transformations.camera import _ar_ox_fisheye, _os_fisheye
|
||||
warp_targets = []
|
||||
for cam in [_ar_ox_fisheye, _os_fisheye]:
|
||||
w, h = cam.width, cam.height
|
||||
for bl in [2, 5]:
|
||||
warp_targets.append(File(f"models/warp_{w}x{h}_b{bl}_tinygrad.pkl").abspath)
|
||||
lenv.Command(warp_targets, tinygrad_files + script_files, compile_warp_cmd)
|
||||
|
||||
@@ -1,36 +1,51 @@
|
||||
#!/usr/bin/env python3
|
||||
import sys
|
||||
import pathlib
|
||||
import onnx
|
||||
import codecs
|
||||
import pickle
|
||||
from typing import Any
|
||||
|
||||
from tinygrad.nn.onnx import OnnxPBParser
|
||||
|
||||
def get_name_and_shape(value_info:onnx.ValueInfoProto) -> tuple[str, tuple[int,...]]:
|
||||
shape = tuple([int(dim.dim_value) for dim in value_info.type.tensor_type.shape.dim])
|
||||
name = value_info.name
|
||||
|
||||
class MetadataOnnxPBParser(OnnxPBParser):
|
||||
def _parse_ModelProto(self) -> dict:
|
||||
obj: dict[str, Any] = {"graph": {"input": [], "output": []}, "metadata_props": []}
|
||||
for fid, wire_type in self._parse_message(self.reader.len):
|
||||
match fid:
|
||||
case 7:
|
||||
obj["graph"] = self._parse_GraphProto()
|
||||
case 14:
|
||||
obj["metadata_props"].append(self._parse_StringStringEntryProto())
|
||||
case _:
|
||||
self.reader.skip_field(wire_type)
|
||||
return obj
|
||||
|
||||
|
||||
def get_name_and_shape(value_info: dict[str, Any]) -> tuple[str, tuple[int, ...]]:
|
||||
shape = tuple(int(dim) if isinstance(dim, int) else 0 for dim in value_info["parsed_type"].shape)
|
||||
name = value_info["name"]
|
||||
return name, shape
|
||||
|
||||
|
||||
def get_metadata_value_by_name(model:onnx.ModelProto, name:str) -> str | Any:
|
||||
for prop in model.metadata_props:
|
||||
if prop.key == name:
|
||||
return prop.value
|
||||
def get_metadata_value_by_name(model: dict[str, Any], name: str) -> str | Any:
|
||||
for prop in model["metadata_props"]:
|
||||
if prop["key"] == name:
|
||||
return prop["value"]
|
||||
return None
|
||||
|
||||
|
||||
if __name__ == "__main__":
|
||||
model_path = pathlib.Path(sys.argv[1])
|
||||
model = onnx.load(str(model_path))
|
||||
model = MetadataOnnxPBParser(model_path).parse()
|
||||
output_slices = get_metadata_value_by_name(model, 'output_slices')
|
||||
assert output_slices is not None, 'output_slices not found in metadata'
|
||||
|
||||
metadata = {
|
||||
'model_checkpoint': get_metadata_value_by_name(model, 'model_checkpoint'),
|
||||
'output_slices': pickle.loads(codecs.decode(output_slices.encode(), "base64")),
|
||||
'input_shapes': dict([get_name_and_shape(x) for x in model.graph.input]),
|
||||
'output_shapes': dict([get_name_and_shape(x) for x in model.graph.output])
|
||||
'input_shapes': dict(get_name_and_shape(x) for x in model["graph"]["input"]),
|
||||
'output_shapes': dict(get_name_and_shape(x) for x in model["graph"]["output"]),
|
||||
}
|
||||
|
||||
metadata_path = model_path.parent / (model_path.stem + '_metadata.pkl')
|
||||
|
||||
@@ -3,41 +3,27 @@ import sys
|
||||
import shutil
|
||||
import pickle
|
||||
import codecs
|
||||
import onnx
|
||||
from pathlib import Path
|
||||
|
||||
from openpilot.system.hardware.hw import Paths
|
||||
|
||||
|
||||
def get_name_and_shape(value_info):
|
||||
shape = tuple([int(dim.dim_value) for dim in value_info.type.tensor_type.shape.dim])
|
||||
return value_info.name, shape
|
||||
|
||||
|
||||
def get_metadata_value_by_name(model, name):
|
||||
for prop in model.metadata_props:
|
||||
if prop.key == name:
|
||||
return prop.value
|
||||
return None
|
||||
from sunnypilot.modeld_v2.get_model_metadata import MetadataOnnxPBParser, get_name_and_shape, get_metadata_value_by_name
|
||||
|
||||
|
||||
def generate_metadata_pkl(model_path, output_path):
|
||||
try:
|
||||
model = onnx.load(str(model_path))
|
||||
model = MetadataOnnxPBParser(model_path).parse()
|
||||
output_slices = get_metadata_value_by_name(model, 'output_slices')
|
||||
|
||||
if output_slices:
|
||||
metadata = {
|
||||
'model_checkpoint': get_metadata_value_by_name(model, 'model_checkpoint'),
|
||||
'output_slices': pickle.loads(codecs.decode(output_slices.encode(), "base64")),
|
||||
'input_shapes': dict([get_name_and_shape(x) for x in model.graph.input]),
|
||||
'output_shapes': dict([get_name_and_shape(x) for x in model.graph.output])
|
||||
}
|
||||
with open(output_path, 'wb') as f:
|
||||
pickle.dump(metadata, f)
|
||||
return True
|
||||
else:
|
||||
if not output_slices:
|
||||
return False
|
||||
metadata = {
|
||||
'model_checkpoint': get_metadata_value_by_name(model, 'model_checkpoint'),
|
||||
'output_slices': pickle.loads(codecs.decode(output_slices.encode(), "base64")),
|
||||
'input_shapes': dict(get_name_and_shape(x) for x in model["graph"]["input"]),
|
||||
'output_shapes': dict(get_name_and_shape(x) for x in model["graph"]["output"]),
|
||||
}
|
||||
with open(output_path, 'wb') as f:
|
||||
pickle.dump(metadata, f)
|
||||
return True
|
||||
except Exception:
|
||||
return False
|
||||
|
||||
|
||||
@@ -26,7 +26,7 @@ from openpilot.selfdrive.controls.lib.drive_helpers import get_accel_from_plan,
|
||||
|
||||
from openpilot.sunnypilot.modeld_v2.fill_model_msg import fill_model_msg, fill_pose_msg, PublishState, get_curvature_from_output
|
||||
from openpilot.sunnypilot.modeld_v2.constants import Plan
|
||||
from openpilot.sunnypilot.modeld_v2.models.commonmodel_pyx import DrivingModelFrame, CLContext
|
||||
from openpilot.sunnypilot.modeld_v2.warp import Warp
|
||||
from openpilot.sunnypilot.modeld_v2.meta_helper import load_meta_constants
|
||||
from openpilot.sunnypilot.modeld_v2.camera_offset_helper import CameraOffsetHelper
|
||||
|
||||
@@ -49,12 +49,12 @@ class FrameMeta:
|
||||
|
||||
|
||||
class ModelState(ModelStateBase):
|
||||
frames: dict[str, DrivingModelFrame]
|
||||
frames: dict[str, Warp]
|
||||
inputs: dict[str, np.ndarray]
|
||||
prev_desire: np.ndarray # for tracking the rising edge of the pulse
|
||||
temporal_idxs: slice | np.ndarray
|
||||
|
||||
def __init__(self, context: CLContext):
|
||||
def __init__(self):
|
||||
ModelStateBase.__init__(self)
|
||||
try:
|
||||
self.model_runner = get_model_runner()
|
||||
@@ -73,17 +73,16 @@ class ModelState(ModelStateBase):
|
||||
self.PLANPLUS_CONTROL: float = 1.0
|
||||
|
||||
buffer_length = 5 if self.model_runner.is_20hz else 2
|
||||
self.frames = {name: DrivingModelFrame(context, buffer_length) for name in self.model_runner.vision_input_names}
|
||||
self.warp = Warp(buffer_length)
|
||||
self.prev_desire = np.zeros(self.constants.DESIRE_LEN, dtype=np.float32)
|
||||
|
||||
# img buffers are managed in openCL transform code
|
||||
self.numpy_inputs = {}
|
||||
self.temporal_buffers = {}
|
||||
self.temporal_idxs_map = {}
|
||||
|
||||
for key, shape in self.model_runner.input_shapes.items():
|
||||
if key not in self.frames: # Managed by opencl
|
||||
if key not in self.model_runner.vision_input_names: # Policy inputs
|
||||
self.numpy_inputs[key] = np.zeros(shape, dtype=np.float32)
|
||||
|
||||
# Temporal input: shape is [batch, history, features]
|
||||
if len(shape) == 3 and shape[1] > 1:
|
||||
buffer_history_len = shape[1] * 4 if shape[1] < 99 else shape[1] # Allow for higher history buffers in the future
|
||||
@@ -129,10 +128,10 @@ class ModelState(ModelStateBase):
|
||||
if key in inputs and key not in [self.desire_key]:
|
||||
self.numpy_inputs[key][:] = inputs[key]
|
||||
|
||||
imgs_cl = {name: self.frames[name].prepare(bufs[name], transforms[name].flatten()) for name in self.model_runner.vision_input_names}
|
||||
|
||||
# Prepare inputs using the model runner
|
||||
self.model_runner.prepare_inputs(imgs_cl, self.numpy_inputs, self.frames)
|
||||
imgs_tensors = self.warp.process(bufs, transforms)
|
||||
for name, tensor in imgs_tensors.items():
|
||||
self.model_runner.inputs[name] = tensor
|
||||
self.model_runner.prepare_inputs(self.numpy_inputs)
|
||||
|
||||
if prepare_only:
|
||||
return None
|
||||
@@ -147,12 +146,11 @@ class ModelState(ModelStateBase):
|
||||
|
||||
if "desired_curvature" in outputs:
|
||||
input_name_prev = None
|
||||
if "prev_desired_curvs" in self.numpy_inputs.keys():
|
||||
input_name_prev = 'prev_desired_curvs'
|
||||
elif "prev_desired_curv" in self.numpy_inputs.keys():
|
||||
if "prev_desired_curv" in self.numpy_inputs.keys():
|
||||
input_name_prev = 'prev_desired_curv'
|
||||
if input_name_prev and input_name_prev in self.temporal_buffers:
|
||||
self.process_desired_curvature(outputs, input_name_prev)
|
||||
|
||||
return outputs
|
||||
|
||||
def process_desired_curvature(self, outputs, input_name_prev):
|
||||
@@ -165,9 +163,8 @@ class ModelState(ModelStateBase):
|
||||
def get_action_from_model(self, model_output: dict[str, np.ndarray], prev_action: log.ModelDataV2.Action,
|
||||
lat_action_t: float, long_action_t: float, v_ego: float) -> log.ModelDataV2.Action:
|
||||
plan = model_output['plan'][0]
|
||||
if 'planplus' in model_output:
|
||||
recovery_power = self.PLANPLUS_CONTROL * (0.75 if v_ego > 20.0 else 1.0)
|
||||
plan = plan + recovery_power * model_output['planplus'][0]
|
||||
if 'planplus' in model_output and self.PLANPLUS_CONTROL != 1.0:
|
||||
plan = plan + (self.PLANPLUS_CONTROL - 1.0) * model_output['planplus'][0]
|
||||
desired_accel, should_stop = get_accel_from_plan(plan[:, Plan.VELOCITY][:, 0], plan[:, Plan.ACCELERATION][:, 0], self.constants.T_IDXS,
|
||||
action_t=long_action_t)
|
||||
desired_accel = smooth_value(desired_accel, prev_action.desiredAcceleration, self.LONG_SMOOTH_SECONDS)
|
||||
@@ -190,10 +187,8 @@ def main(demo=False):
|
||||
setproctitle(PROCESS_NAME)
|
||||
config_realtime_process(7, 54)
|
||||
|
||||
cloudlog.warning("setting up CL context")
|
||||
cl_context = CLContext()
|
||||
cloudlog.warning("CL context ready; loading model")
|
||||
model = ModelState(cl_context)
|
||||
cloudlog.warning("loading model")
|
||||
model = ModelState()
|
||||
cloudlog.warning("models loaded, modeld starting")
|
||||
|
||||
# visionipc clients
|
||||
@@ -206,8 +201,8 @@ def main(demo=False):
|
||||
time.sleep(.1)
|
||||
|
||||
vipc_client_main_stream = VisionStreamType.VISION_STREAM_WIDE_ROAD if main_wide_camera else VisionStreamType.VISION_STREAM_ROAD
|
||||
vipc_client_main = VisionIpcClient("camerad", vipc_client_main_stream, True, cl_context)
|
||||
vipc_client_extra = VisionIpcClient("camerad", VisionStreamType.VISION_STREAM_WIDE_ROAD, False, cl_context)
|
||||
vipc_client_main = VisionIpcClient("camerad", vipc_client_main_stream, True)
|
||||
vipc_client_extra = VisionIpcClient("camerad", VisionStreamType.VISION_STREAM_WIDE_ROAD, False)
|
||||
cloudlog.warning(f"vision stream set up, main_wide_camera: {main_wide_camera}, use_extra_client: {use_extra_client}")
|
||||
|
||||
while not vipc_client_main.connect(False):
|
||||
|
||||
@@ -1,62 +0,0 @@
|
||||
#include "sunnypilot/modeld_v2/models/commonmodel.h"
|
||||
|
||||
#include <cmath>
|
||||
#include <cstring>
|
||||
|
||||
#include "common/clutil.h"
|
||||
|
||||
DrivingModelFrame::DrivingModelFrame(cl_device_id device_id, cl_context context, uint8_t buffer_length) : ModelFrame(device_id, context), buffer_length(buffer_length) {
|
||||
input_frames = std::make_unique<uint8_t[]>(buf_size);
|
||||
input_frames_cl = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_WRITE, buf_size, NULL, &err));
|
||||
img_buffer_20hz_cl = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_WRITE, buffer_length*frame_size_bytes, NULL, &err));
|
||||
region.origin = (buffer_length - 1) * frame_size_bytes;
|
||||
region.size = frame_size_bytes;
|
||||
last_img_cl = CL_CHECK_ERR(clCreateSubBuffer(img_buffer_20hz_cl, CL_MEM_READ_WRITE, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err));
|
||||
// printf("Buffer length: %d, region origin: %lu, region size: %lu\n", buffer_length, region.origin, region.size);
|
||||
|
||||
loadyuv_init(&loadyuv, context, device_id, MODEL_WIDTH, MODEL_HEIGHT);
|
||||
init_transform(device_id, context, MODEL_WIDTH, MODEL_HEIGHT);
|
||||
}
|
||||
|
||||
cl_mem* DrivingModelFrame::prepare(cl_mem yuv_cl, int frame_width, int frame_height, int frame_stride, int frame_uv_offset, const mat3& projection) {
|
||||
run_transform(yuv_cl, MODEL_WIDTH, MODEL_HEIGHT, frame_width, frame_height, frame_stride, frame_uv_offset, projection);
|
||||
|
||||
for (int i = 0; i < (buffer_length - 1); i++) {
|
||||
CL_CHECK(clEnqueueCopyBuffer(q, img_buffer_20hz_cl, img_buffer_20hz_cl, (i+1)*frame_size_bytes, i*frame_size_bytes, frame_size_bytes, 0, nullptr, nullptr));
|
||||
}
|
||||
loadyuv_queue(&loadyuv, q, y_cl, u_cl, v_cl, last_img_cl);
|
||||
|
||||
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 &input_frames_cl;
|
||||
}
|
||||
|
||||
DrivingModelFrame::~DrivingModelFrame() {
|
||||
deinit_transform();
|
||||
loadyuv_destroy(&loadyuv);
|
||||
CL_CHECK(clReleaseMemObject(img_buffer_20hz_cl));
|
||||
CL_CHECK(clReleaseMemObject(last_img_cl));
|
||||
CL_CHECK(clReleaseCommandQueue(q));
|
||||
}
|
||||
|
||||
|
||||
MonitoringModelFrame::MonitoringModelFrame(cl_device_id device_id, cl_context context) : ModelFrame(device_id, context) {
|
||||
input_frames = std::make_unique<uint8_t[]>(buf_size);
|
||||
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);
|
||||
}
|
||||
|
||||
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);
|
||||
clFinish(q);
|
||||
return &y_cl;
|
||||
}
|
||||
|
||||
MonitoringModelFrame::~MonitoringModelFrame() {
|
||||
deinit_transform();
|
||||
CL_CHECK(clReleaseCommandQueue(q));
|
||||
}
|
||||
@@ -1,97 +0,0 @@
|
||||
#pragma once
|
||||
|
||||
#include <cfloat>
|
||||
#include <cstdlib>
|
||||
#include <cassert>
|
||||
|
||||
#include <memory>
|
||||
|
||||
#define CL_USE_DEPRECATED_OPENCL_1_2_APIS
|
||||
#ifdef __APPLE__
|
||||
#include <OpenCL/cl.h>
|
||||
#else
|
||||
#include <CL/cl.h>
|
||||
#endif
|
||||
|
||||
#include "common/mat.h"
|
||||
#include "sunnypilot/modeld_v2/transforms/loadyuv.h"
|
||||
#include "sunnypilot/modeld_v2/transforms/transform.h"
|
||||
|
||||
class ModelFrame {
|
||||
public:
|
||||
ModelFrame(cl_device_id device_id, cl_context context) {
|
||||
q = CL_CHECK_ERR(clCreateCommandQueue(context, device_id, 0, &err));
|
||||
}
|
||||
virtual ~ModelFrame() {}
|
||||
virtual cl_mem* prepare(cl_mem yuv_cl, int frame_width, int frame_height, int frame_stride, int frame_uv_offset, const mat3& projection) { return NULL; }
|
||||
uint8_t* buffer_from_cl(cl_mem *in_frames, int buffer_size) {
|
||||
CL_CHECK(clEnqueueReadBuffer(q, *in_frames, CL_TRUE, 0, buffer_size, input_frames.get(), 0, nullptr, nullptr));
|
||||
clFinish(q);
|
||||
return &input_frames[0];
|
||||
}
|
||||
|
||||
int MODEL_WIDTH;
|
||||
int MODEL_HEIGHT;
|
||||
int MODEL_FRAME_SIZE;
|
||||
int buf_size;
|
||||
|
||||
protected:
|
||||
cl_mem y_cl, u_cl, v_cl;
|
||||
Transform transform;
|
||||
cl_command_queue q;
|
||||
std::unique_ptr<uint8_t[]> input_frames;
|
||||
|
||||
void init_transform(cl_device_id device_id, cl_context context, int model_width, int model_height) {
|
||||
y_cl = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_WRITE, model_width * model_height, NULL, &err));
|
||||
u_cl = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_WRITE, (model_width / 2) * (model_height / 2), NULL, &err));
|
||||
v_cl = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_WRITE, (model_width / 2) * (model_height / 2), NULL, &err));
|
||||
transform_init(&transform, context, device_id);
|
||||
}
|
||||
|
||||
void deinit_transform() {
|
||||
transform_destroy(&transform);
|
||||
CL_CHECK(clReleaseMemObject(v_cl));
|
||||
CL_CHECK(clReleaseMemObject(u_cl));
|
||||
CL_CHECK(clReleaseMemObject(y_cl));
|
||||
}
|
||||
|
||||
void run_transform(cl_mem yuv_cl, int model_width, int model_height, int frame_width, int frame_height, int frame_stride, int frame_uv_offset, const mat3& projection) {
|
||||
transform_queue(&transform, q,
|
||||
yuv_cl, frame_width, frame_height, frame_stride, frame_uv_offset,
|
||||
y_cl, u_cl, v_cl, model_width, model_height, projection);
|
||||
}
|
||||
};
|
||||
|
||||
class DrivingModelFrame : public ModelFrame {
|
||||
public:
|
||||
DrivingModelFrame(cl_device_id device_id, cl_context context, uint8_t buffer_length);
|
||||
~DrivingModelFrame();
|
||||
cl_mem* prepare(cl_mem yuv_cl, int frame_width, int frame_height, int frame_stride, int frame_uv_offset, const mat3& projection);
|
||||
|
||||
const int MODEL_WIDTH = 512;
|
||||
const int MODEL_HEIGHT = 256;
|
||||
const int MODEL_FRAME_SIZE = MODEL_WIDTH * MODEL_HEIGHT * 3 / 2;
|
||||
const int buf_size = MODEL_FRAME_SIZE * 2;
|
||||
const size_t frame_size_bytes = MODEL_FRAME_SIZE * sizeof(uint8_t);
|
||||
const uint8_t buffer_length;
|
||||
|
||||
private:
|
||||
LoadYUVState loadyuv;
|
||||
cl_mem img_buffer_20hz_cl, last_img_cl, input_frames_cl;
|
||||
cl_buffer_region region;
|
||||
};
|
||||
|
||||
class MonitoringModelFrame : public ModelFrame {
|
||||
public:
|
||||
MonitoringModelFrame(cl_device_id device_id, cl_context context);
|
||||
~MonitoringModelFrame();
|
||||
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;
|
||||
const int MODEL_FRAME_SIZE = MODEL_WIDTH * MODEL_HEIGHT;
|
||||
const int buf_size = MODEL_FRAME_SIZE;
|
||||
|
||||
private:
|
||||
cl_mem input_frame_cl;
|
||||
};
|
||||
@@ -1,27 +0,0 @@
|
||||
# distutils: language = c++
|
||||
|
||||
from msgq.visionipc.visionipc cimport cl_device_id, cl_context, cl_mem
|
||||
|
||||
cdef extern from "common/mat.h":
|
||||
cdef struct mat3:
|
||||
float v[9]
|
||||
|
||||
cdef extern from "common/clutil.h":
|
||||
cdef unsigned long CL_DEVICE_TYPE_DEFAULT
|
||||
cl_device_id cl_get_device_id(unsigned long)
|
||||
cl_context cl_create_context(cl_device_id)
|
||||
void cl_release_context(cl_context)
|
||||
|
||||
cdef extern from "sunnypilot/modeld_v2/models/commonmodel.h":
|
||||
cppclass ModelFrame:
|
||||
int buf_size
|
||||
unsigned char * buffer_from_cl(cl_mem*, int);
|
||||
cl_mem * prepare(cl_mem, int, int, int, int, mat3)
|
||||
|
||||
cppclass DrivingModelFrame:
|
||||
int buf_size
|
||||
DrivingModelFrame(cl_device_id, cl_context, unsigned char)
|
||||
|
||||
cppclass MonitoringModelFrame:
|
||||
int buf_size
|
||||
MonitoringModelFrame(cl_device_id, cl_context)
|
||||
@@ -1,13 +0,0 @@
|
||||
# distutils: language = c++
|
||||
|
||||
from msgq.visionipc.visionipc cimport cl_mem
|
||||
from msgq.visionipc.visionipc_pyx cimport CLContext as BaseCLContext
|
||||
|
||||
cdef class CLContext(BaseCLContext):
|
||||
pass
|
||||
|
||||
cdef class CLMem:
|
||||
cdef cl_mem * mem
|
||||
|
||||
@staticmethod
|
||||
cdef create(void*)
|
||||
@@ -1,74 +0,0 @@
|
||||
# distutils: language = c++
|
||||
# cython: c_string_encoding=ascii, language_level=3
|
||||
|
||||
import numpy as np
|
||||
cimport numpy as cnp
|
||||
from libc.string cimport memcpy
|
||||
from libc.stdint cimport uintptr_t, uint8_t
|
||||
|
||||
from msgq.visionipc.visionipc cimport cl_mem
|
||||
from msgq.visionipc.visionipc_pyx cimport VisionBuf, CLContext as BaseCLContext
|
||||
from .commonmodel cimport CL_DEVICE_TYPE_DEFAULT, cl_get_device_id, cl_create_context, cl_release_context
|
||||
from .commonmodel cimport mat3, ModelFrame as cppModelFrame, DrivingModelFrame as cppDrivingModelFrame, MonitoringModelFrame as cppMonitoringModelFrame
|
||||
|
||||
|
||||
cdef class CLContext(BaseCLContext):
|
||||
def __cinit__(self):
|
||||
self.device_id = cl_get_device_id(CL_DEVICE_TYPE_DEFAULT)
|
||||
self.context = cl_create_context(self.device_id)
|
||||
|
||||
def __dealloc__(self):
|
||||
if self.context:
|
||||
cl_release_context(self.context)
|
||||
|
||||
cdef class CLMem:
|
||||
@staticmethod
|
||||
cdef create(void * cmem):
|
||||
mem = CLMem()
|
||||
mem.mem = <cl_mem*> cmem
|
||||
return mem
|
||||
|
||||
@property
|
||||
def mem_address(self):
|
||||
return <uintptr_t>(self.mem)
|
||||
|
||||
def cl_from_visionbuf(VisionBuf buf):
|
||||
return CLMem.create(<void*>&buf.buf.buf_cl)
|
||||
|
||||
|
||||
cdef class ModelFrame:
|
||||
cdef cppModelFrame * frame
|
||||
cdef int buf_size
|
||||
|
||||
def __dealloc__(self):
|
||||
del self.frame
|
||||
|
||||
def prepare(self, VisionBuf buf, float[:] projection):
|
||||
cdef mat3 cprojection
|
||||
memcpy(cprojection.v, &projection[0], 9*sizeof(float))
|
||||
cdef cl_mem * data
|
||||
data = self.frame.prepare(buf.buf.buf_cl, buf.width, buf.height, buf.stride, buf.uv_offset, cprojection)
|
||||
return CLMem.create(data)
|
||||
|
||||
def buffer_from_cl(self, CLMem in_frames):
|
||||
cdef unsigned char * data2
|
||||
data2 = self.frame.buffer_from_cl(in_frames.mem, self.buf_size)
|
||||
return np.asarray(<cnp.uint8_t[:self.buf_size]> data2)
|
||||
|
||||
|
||||
cdef class DrivingModelFrame(ModelFrame):
|
||||
cdef cppDrivingModelFrame * _frame
|
||||
|
||||
def __cinit__(self, CLContext context, int buffer_length=2):
|
||||
self._frame = new cppDrivingModelFrame(context.device_id, context.context, buffer_length)
|
||||
self.frame = <cppModelFrame*>(self._frame)
|
||||
self.buf_size = self._frame.buf_size
|
||||
|
||||
cdef class MonitoringModelFrame(ModelFrame):
|
||||
cdef cppMonitoringModelFrame * _frame
|
||||
|
||||
def __cinit__(self, CLContext context):
|
||||
self._frame = new cppMonitoringModelFrame(context.device_id, context.context)
|
||||
self.frame = <cppModelFrame*>(self._frame)
|
||||
self.buf_size = self._frame.buf_size
|
||||
|
||||
@@ -108,8 +108,8 @@ class Parser:
|
||||
plan_in_N, plan_out_N = (SplitModelConstants.PLAN_MHP_N, SplitModelConstants.PLAN_MHP_SELECTION) if plan_mhp else (0, 0)
|
||||
self.parse_mdn('plan', outs, in_N=plan_in_N, out_N=plan_out_N,
|
||||
out_shape=(SplitModelConstants.IDX_N, SplitModelConstants.PLAN_WIDTH))
|
||||
if 'planplus' in outs:
|
||||
self.parse_mdn('planplus', outs, in_N=plan_in_N, out_N=plan_out_N, out_shape=(SplitModelConstants.IDX_N, SplitModelConstants.PLAN_WIDTH))
|
||||
if 'planplus' in outs:
|
||||
self.parse_mdn('planplus', outs, in_N=0, out_N=0, out_shape=(SplitModelConstants.IDX_N, SplitModelConstants.PLAN_WIDTH))
|
||||
|
||||
def split_outputs(self, outs: dict[str, np.ndarray]) -> None:
|
||||
if 'desired_curvature' in outs:
|
||||
|
||||
@@ -1,36 +0,0 @@
|
||||
import onnx
|
||||
import onnxruntime as ort
|
||||
import numpy as np
|
||||
import itertools
|
||||
|
||||
ORT_TYPES_TO_NP_TYPES = {'tensor(float16)': np.float16, 'tensor(float)': np.float32, 'tensor(uint8)': np.uint8}
|
||||
|
||||
def attributeproto_fp16_to_fp32(attr):
|
||||
float32_list = np.frombuffer(attr.raw_data, dtype=np.float16)
|
||||
attr.data_type = 1
|
||||
attr.raw_data = float32_list.astype(np.float32).tobytes()
|
||||
|
||||
def convert_fp16_to_fp32(model):
|
||||
for i in model.graph.initializer:
|
||||
if i.data_type == 10:
|
||||
attributeproto_fp16_to_fp32(i)
|
||||
for i in itertools.chain(model.graph.input, model.graph.output):
|
||||
if i.type.tensor_type.elem_type == 10:
|
||||
i.type.tensor_type.elem_type = 1
|
||||
for i in model.graph.node:
|
||||
if i.op_type == 'Cast' and i.attribute[0].i == 10:
|
||||
i.attribute[0].i = 1
|
||||
for a in i.attribute:
|
||||
if hasattr(a, 't'):
|
||||
if a.t.data_type == 10:
|
||||
attributeproto_fp16_to_fp32(a.t)
|
||||
return model.SerializeToString()
|
||||
|
||||
|
||||
def make_onnx_cpu_runner(model_path):
|
||||
options = ort.SessionOptions()
|
||||
options.intra_op_num_threads = 4
|
||||
options.execution_mode = ort.ExecutionMode.ORT_SEQUENTIAL
|
||||
options.graph_optimization_level = ort.GraphOptimizationLevel.ORT_ENABLE_ALL
|
||||
model_data = convert_fp16_to_fp32(onnx.load(model_path))
|
||||
return ort.InferenceSession(model_data, options, providers=['CPUExecutionProvider'])
|
||||
@@ -1,8 +0,0 @@
|
||||
|
||||
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')
|
||||
@@ -53,7 +53,7 @@ class DummyModelRunner:
|
||||
self.is_20hz = False
|
||||
|
||||
# Minimal prepare/run methods so ModelState can be run without actually running the model
|
||||
def prepare_inputs(self, imgs_cl, numpy_inputs, frames):
|
||||
def prepare_inputs(self, numpy_inputs):
|
||||
return None
|
||||
|
||||
def run_model(self):
|
||||
@@ -105,7 +105,7 @@ def get_expected_indices(shape, constants, mode, key=None):
|
||||
|
||||
@pytest.mark.parametrize("shapes,mode", SHAPE_MODE_PARAMS, indirect=["shapes"])
|
||||
def test_buffer_shapes_and_indices(shapes, mode, apply_patches):
|
||||
state = ModelState(None)
|
||||
state = ModelState()
|
||||
constants = DummyModelRunner(shapes).constants
|
||||
for key in shapes:
|
||||
buf = state.temporal_buffers.get(key, None)
|
||||
@@ -236,7 +236,7 @@ def dynamic_buffer_update(state, key, new_val, mode):
|
||||
@pytest.mark.parametrize("shapes,mode", SHAPE_MODE_PARAMS, indirect=["shapes"])
|
||||
@pytest.mark.parametrize("key", ["desire", "features_buffer", "prev_desired_curv"])
|
||||
def test_buffer_update_equivalence(shapes, mode, key, apply_patches):
|
||||
state = ModelState(None)
|
||||
state = ModelState()
|
||||
if key == "desire":
|
||||
desire_keys = [k for k in shapes.keys() if k.startswith('desire')]
|
||||
if desire_keys:
|
||||
|
||||
@@ -0,0 +1,102 @@
|
||||
import os
|
||||
os.environ['DEV'] = 'CPU'
|
||||
import pytest
|
||||
import numpy as np
|
||||
from openpilot.selfdrive.modeld.compile_warp import get_nv12_info, CAMERA_CONFIGS
|
||||
from openpilot.sunnypilot.modeld_v2.warp import Warp, MODEL_W, MODEL_H
|
||||
|
||||
VISION_NAME_PAIRS = [ # needed to account for supercombos input_imgs
|
||||
('img', 'big_img'),
|
||||
('input_imgs', 'big_input_imgs'),
|
||||
]
|
||||
|
||||
|
||||
class MockVisionBuf:
|
||||
def __init__(self, w, h):
|
||||
self.width = w
|
||||
self.height = h
|
||||
_, _, _, yuv_size = get_nv12_info(w, h)
|
||||
self.data = np.zeros(yuv_size, dtype=np.uint8)
|
||||
|
||||
|
||||
@pytest.mark.parametrize("buffer_length", [2, 5])
|
||||
def test_warp_initialization(buffer_length):
|
||||
warp = Warp(buffer_length)
|
||||
assert warp.buffer_length == buffer_length
|
||||
assert warp.img_buffer_shape == (buffer_length * 6, MODEL_H // 2, MODEL_W // 2)
|
||||
|
||||
|
||||
@pytest.mark.parametrize("buffer_length", [2, 5])
|
||||
@pytest.mark.parametrize("cam_w, cam_h", CAMERA_CONFIGS)
|
||||
@pytest.mark.parametrize("road, wide", VISION_NAME_PAIRS)
|
||||
def test_warp_process(buffer_length, cam_w, cam_h, road, wide):
|
||||
warp = Warp(buffer_length)
|
||||
mock_buf = MockVisionBuf(cam_w, cam_h)
|
||||
transform = np.eye(3, dtype=np.float32).flatten()
|
||||
bufs = {road: mock_buf, wide: mock_buf}
|
||||
transforms = {road: transform, wide: transform}
|
||||
|
||||
out = warp.process(bufs, transforms)
|
||||
assert isinstance(out, dict)
|
||||
assert road in out and wide in out
|
||||
assert out[road].shape == (1, 12, MODEL_H // 2, MODEL_W // 2)
|
||||
assert out[wide].shape == (1, 12, MODEL_H // 2, MODEL_W // 2)
|
||||
|
||||
key = (cam_w, cam_h)
|
||||
assert key in warp.jit_cache
|
||||
|
||||
out2 = warp.process(bufs, transforms)
|
||||
assert out2[road].shape == out[road].shape
|
||||
|
||||
|
||||
@pytest.mark.parametrize("road, wide", VISION_NAME_PAIRS)
|
||||
def test_warp_buffer_shift(road, wide):
|
||||
warp = Warp(2)
|
||||
cam_w, cam_h = CAMERA_CONFIGS[1]
|
||||
transform = np.eye(3, dtype=np.float32).flatten()
|
||||
|
||||
buf1 = MockVisionBuf(cam_w, cam_h)
|
||||
buf1.data[0] = 255
|
||||
bufs1 = {road: buf1, wide: buf1}
|
||||
transforms = {road: transform, wide: transform}
|
||||
out1 = warp.process(bufs1, transforms)
|
||||
road1 = out1[road].numpy().copy()
|
||||
|
||||
buf2 = MockVisionBuf(cam_w, cam_h)
|
||||
buf2.data[0] = 128
|
||||
bufs2 = {road: buf2, wide: buf2}
|
||||
out2 = warp.process(bufs2, transforms)
|
||||
assert not np.array_equal(road1, out2[road].numpy())
|
||||
|
||||
|
||||
@pytest.mark.parametrize("buffer_length", [2, 5])
|
||||
@pytest.mark.parametrize("road, wide", VISION_NAME_PAIRS)
|
||||
def test_warp_buffer_accumulation(buffer_length, road, wide):
|
||||
warp = Warp(buffer_length)
|
||||
cam_w, cam_h = CAMERA_CONFIGS[0]
|
||||
transform = np.eye(3, dtype=np.float32).flatten()
|
||||
transforms = {road: transform, wide: transform}
|
||||
outputs = []
|
||||
|
||||
for i in range(buffer_length + 1):
|
||||
buf = MockVisionBuf(cam_w, cam_h)
|
||||
buf.data[:] = i * 10
|
||||
out = warp.process({road: buf, wide: buf}, transforms)
|
||||
outputs.append(out[road].numpy().copy())
|
||||
|
||||
assert warp.full_buffers['img'].shape == (buffer_length * 6, MODEL_H // 2, MODEL_W // 2)
|
||||
for i in range(1, len(outputs)):
|
||||
assert not np.array_equal(outputs[i - 1], outputs[i])
|
||||
|
||||
|
||||
def test_warp_different_cameras_same_instance():
|
||||
warp = Warp(2)
|
||||
transform = np.eye(3, dtype=np.float32).flatten()
|
||||
|
||||
buf1 = MockVisionBuf(*CAMERA_CONFIGS[0])
|
||||
warp.process({'img': buf1, 'big_img': buf1}, {'img': transform, 'big_img': transform})
|
||||
assert len(warp.jit_cache) == 1
|
||||
|
||||
buf2 = MockVisionBuf(*CAMERA_CONFIGS[1])
|
||||
warp.process({'img': buf2, 'big_img': buf2}, {'img': transform, 'big_img': transform})
|
||||
assert len(warp.jit_cache) == 2
|
||||
@@ -1,76 +0,0 @@
|
||||
#include "sunnypilot/modeld_v2/transforms/loadyuv.h"
|
||||
|
||||
#include <cassert>
|
||||
#include <cstdio>
|
||||
#include <cstring>
|
||||
|
||||
void loadyuv_init(LoadYUVState* s, cl_context ctx, cl_device_id device_id, int width, int height) {
|
||||
memset(s, 0, sizeof(*s));
|
||||
|
||||
s->width = width;
|
||||
s->height = height;
|
||||
|
||||
char args[1024];
|
||||
snprintf(args, sizeof(args),
|
||||
"-cl-fast-relaxed-math -cl-denorms-are-zero "
|
||||
"-DTRANSFORMED_WIDTH=%d -DTRANSFORMED_HEIGHT=%d",
|
||||
width, height);
|
||||
cl_program prg = cl_program_from_file(ctx, device_id, LOADYUV_PATH, args);
|
||||
|
||||
s->loadys_krnl = CL_CHECK_ERR(clCreateKernel(prg, "loadys", &err));
|
||||
s->loaduv_krnl = CL_CHECK_ERR(clCreateKernel(prg, "loaduv", &err));
|
||||
s->copy_krnl = CL_CHECK_ERR(clCreateKernel(prg, "copy", &err));
|
||||
|
||||
// done with this
|
||||
CL_CHECK(clReleaseProgram(prg));
|
||||
}
|
||||
|
||||
void loadyuv_destroy(LoadYUVState* s) {
|
||||
CL_CHECK(clReleaseKernel(s->loadys_krnl));
|
||||
CL_CHECK(clReleaseKernel(s->loaduv_krnl));
|
||||
CL_CHECK(clReleaseKernel(s->copy_krnl));
|
||||
}
|
||||
|
||||
void loadyuv_queue(LoadYUVState* s, cl_command_queue q,
|
||||
cl_mem y_cl, cl_mem u_cl, cl_mem v_cl,
|
||||
cl_mem out_cl) {
|
||||
cl_int global_out_off = 0;
|
||||
|
||||
CL_CHECK(clSetKernelArg(s->loadys_krnl, 0, sizeof(cl_mem), &y_cl));
|
||||
CL_CHECK(clSetKernelArg(s->loadys_krnl, 1, sizeof(cl_mem), &out_cl));
|
||||
CL_CHECK(clSetKernelArg(s->loadys_krnl, 2, sizeof(cl_int), &global_out_off));
|
||||
|
||||
const size_t loadys_work_size = (s->width*s->height)/8;
|
||||
CL_CHECK(clEnqueueNDRangeKernel(q, s->loadys_krnl, 1, NULL,
|
||||
&loadys_work_size, NULL, 0, 0, NULL));
|
||||
|
||||
const size_t loaduv_work_size = ((s->width/2)*(s->height/2))/8;
|
||||
global_out_off += (s->width*s->height);
|
||||
|
||||
CL_CHECK(clSetKernelArg(s->loaduv_krnl, 0, sizeof(cl_mem), &u_cl));
|
||||
CL_CHECK(clSetKernelArg(s->loaduv_krnl, 1, sizeof(cl_mem), &out_cl));
|
||||
CL_CHECK(clSetKernelArg(s->loaduv_krnl, 2, sizeof(cl_int), &global_out_off));
|
||||
|
||||
CL_CHECK(clEnqueueNDRangeKernel(q, s->loaduv_krnl, 1, NULL,
|
||||
&loaduv_work_size, NULL, 0, 0, NULL));
|
||||
|
||||
global_out_off += (s->width/2)*(s->height/2);
|
||||
|
||||
CL_CHECK(clSetKernelArg(s->loaduv_krnl, 0, sizeof(cl_mem), &v_cl));
|
||||
CL_CHECK(clSetKernelArg(s->loaduv_krnl, 1, sizeof(cl_mem), &out_cl));
|
||||
CL_CHECK(clSetKernelArg(s->loaduv_krnl, 2, sizeof(cl_int), &global_out_off));
|
||||
|
||||
CL_CHECK(clEnqueueNDRangeKernel(q, s->loaduv_krnl, 1, NULL,
|
||||
&loaduv_work_size, NULL, 0, 0, NULL));
|
||||
}
|
||||
|
||||
void copy_queue(LoadYUVState* s, cl_command_queue q, cl_mem src, cl_mem dst,
|
||||
size_t src_offset, size_t dst_offset, size_t size) {
|
||||
CL_CHECK(clSetKernelArg(s->copy_krnl, 0, sizeof(cl_mem), &src));
|
||||
CL_CHECK(clSetKernelArg(s->copy_krnl, 1, sizeof(cl_mem), &dst));
|
||||
CL_CHECK(clSetKernelArg(s->copy_krnl, 2, sizeof(cl_int), &src_offset));
|
||||
CL_CHECK(clSetKernelArg(s->copy_krnl, 3, sizeof(cl_int), &dst_offset));
|
||||
const size_t copy_work_size = size/8;
|
||||
CL_CHECK(clEnqueueNDRangeKernel(q, s->copy_krnl, 1, NULL,
|
||||
©_work_size, NULL, 0, 0, NULL));
|
||||
}
|
||||
@@ -1,47 +0,0 @@
|
||||
#define UV_SIZE ((TRANSFORMED_WIDTH/2)*(TRANSFORMED_HEIGHT/2))
|
||||
|
||||
__kernel void loadys(__global uchar8 const * const Y,
|
||||
__global uchar * out,
|
||||
int out_offset)
|
||||
{
|
||||
const int gid = get_global_id(0);
|
||||
const int ois = gid * 8;
|
||||
const int oy = ois / TRANSFORMED_WIDTH;
|
||||
const int ox = ois % TRANSFORMED_WIDTH;
|
||||
|
||||
const uchar8 ys = Y[gid];
|
||||
|
||||
// 02
|
||||
// 13
|
||||
|
||||
__global uchar* outy0;
|
||||
__global uchar* outy1;
|
||||
if ((oy & 1) == 0) {
|
||||
outy0 = out + out_offset; //y0
|
||||
outy1 = out + out_offset + UV_SIZE*2; //y2
|
||||
} else {
|
||||
outy0 = out + out_offset + UV_SIZE; //y1
|
||||
outy1 = out + out_offset + UV_SIZE*3; //y3
|
||||
}
|
||||
|
||||
vstore4(ys.s0246, 0, outy0 + (oy/2) * (TRANSFORMED_WIDTH/2) + ox/2);
|
||||
vstore4(ys.s1357, 0, outy1 + (oy/2) * (TRANSFORMED_WIDTH/2) + ox/2);
|
||||
}
|
||||
|
||||
__kernel void loaduv(__global uchar8 const * const in,
|
||||
__global uchar8 * out,
|
||||
int out_offset)
|
||||
{
|
||||
const int gid = get_global_id(0);
|
||||
const uchar8 inv = in[gid];
|
||||
out[gid + out_offset / 8] = inv;
|
||||
}
|
||||
|
||||
__kernel void copy(__global uchar8 * in,
|
||||
__global uchar8 * out,
|
||||
int in_offset,
|
||||
int out_offset)
|
||||
{
|
||||
const int gid = get_global_id(0);
|
||||
out[gid + out_offset / 8] = in[gid + in_offset / 8];
|
||||
}
|
||||
@@ -1,20 +0,0 @@
|
||||
#pragma once
|
||||
|
||||
#include "common/clutil.h"
|
||||
|
||||
typedef struct {
|
||||
int width, height;
|
||||
cl_kernel loadys_krnl, loaduv_krnl, copy_krnl;
|
||||
} LoadYUVState;
|
||||
|
||||
void loadyuv_init(LoadYUVState* s, cl_context ctx, cl_device_id device_id, int width, int height);
|
||||
|
||||
void loadyuv_destroy(LoadYUVState* s);
|
||||
|
||||
void loadyuv_queue(LoadYUVState* s, cl_command_queue q,
|
||||
cl_mem y_cl, cl_mem u_cl, cl_mem v_cl,
|
||||
cl_mem out_cl);
|
||||
|
||||
|
||||
void copy_queue(LoadYUVState* s, cl_command_queue q, cl_mem src, cl_mem dst,
|
||||
size_t src_offset, size_t dst_offset, size_t size);
|
||||
@@ -1,97 +0,0 @@
|
||||
#include "sunnypilot/modeld_v2/transforms/transform.h"
|
||||
|
||||
#include <cassert>
|
||||
#include <cstring>
|
||||
|
||||
#include "common/clutil.h"
|
||||
|
||||
void transform_init(Transform* s, cl_context ctx, cl_device_id device_id) {
|
||||
memset(s, 0, sizeof(*s));
|
||||
|
||||
cl_program prg = cl_program_from_file(ctx, device_id, TRANSFORM_PATH, "");
|
||||
s->krnl = CL_CHECK_ERR(clCreateKernel(prg, "warpPerspective", &err));
|
||||
// done with this
|
||||
CL_CHECK(clReleaseProgram(prg));
|
||||
|
||||
s->m_y_cl = CL_CHECK_ERR(clCreateBuffer(ctx, CL_MEM_READ_WRITE, 3*3*sizeof(float), NULL, &err));
|
||||
s->m_uv_cl = CL_CHECK_ERR(clCreateBuffer(ctx, CL_MEM_READ_WRITE, 3*3*sizeof(float), NULL, &err));
|
||||
}
|
||||
|
||||
void transform_destroy(Transform* s) {
|
||||
CL_CHECK(clReleaseMemObject(s->m_y_cl));
|
||||
CL_CHECK(clReleaseMemObject(s->m_uv_cl));
|
||||
CL_CHECK(clReleaseKernel(s->krnl));
|
||||
}
|
||||
|
||||
void transform_queue(Transform* s,
|
||||
cl_command_queue q,
|
||||
cl_mem in_yuv, int in_width, int in_height, int in_stride, int in_uv_offset,
|
||||
cl_mem out_y, cl_mem out_u, cl_mem out_v,
|
||||
int out_width, int out_height,
|
||||
const mat3& projection) {
|
||||
const int zero = 0;
|
||||
|
||||
// sampled using pixel center origin
|
||||
// (because that's how fastcv and opencv does it)
|
||||
|
||||
mat3 projection_y = projection;
|
||||
|
||||
// in and out uv is half the size of y.
|
||||
mat3 projection_uv = transform_scale_buffer(projection, 0.5);
|
||||
|
||||
CL_CHECK(clEnqueueWriteBuffer(q, s->m_y_cl, CL_TRUE, 0, 3*3*sizeof(float), (void*)projection_y.v, 0, NULL, NULL));
|
||||
CL_CHECK(clEnqueueWriteBuffer(q, s->m_uv_cl, CL_TRUE, 0, 3*3*sizeof(float), (void*)projection_uv.v, 0, NULL, NULL));
|
||||
|
||||
const int in_y_width = in_width;
|
||||
const int in_y_height = in_height;
|
||||
const int in_y_px_stride = 1;
|
||||
const int in_uv_width = in_width/2;
|
||||
const int in_uv_height = in_height/2;
|
||||
const int in_uv_px_stride = 2;
|
||||
const int in_u_offset = in_uv_offset;
|
||||
const int in_v_offset = in_uv_offset + 1;
|
||||
|
||||
const int out_y_width = out_width;
|
||||
const int out_y_height = out_height;
|
||||
const int out_uv_width = out_width/2;
|
||||
const int out_uv_height = out_height/2;
|
||||
|
||||
CL_CHECK(clSetKernelArg(s->krnl, 0, sizeof(cl_mem), &in_yuv)); // src
|
||||
CL_CHECK(clSetKernelArg(s->krnl, 1, sizeof(cl_int), &in_stride)); // src_row_stride
|
||||
CL_CHECK(clSetKernelArg(s->krnl, 2, sizeof(cl_int), &in_y_px_stride)); // src_px_stride
|
||||
CL_CHECK(clSetKernelArg(s->krnl, 3, sizeof(cl_int), &zero)); // src_offset
|
||||
CL_CHECK(clSetKernelArg(s->krnl, 4, sizeof(cl_int), &in_y_height)); // src_rows
|
||||
CL_CHECK(clSetKernelArg(s->krnl, 5, sizeof(cl_int), &in_y_width)); // src_cols
|
||||
CL_CHECK(clSetKernelArg(s->krnl, 6, sizeof(cl_mem), &out_y)); // dst
|
||||
CL_CHECK(clSetKernelArg(s->krnl, 7, sizeof(cl_int), &out_y_width)); // dst_row_stride
|
||||
CL_CHECK(clSetKernelArg(s->krnl, 8, sizeof(cl_int), &zero)); // dst_offset
|
||||
CL_CHECK(clSetKernelArg(s->krnl, 9, sizeof(cl_int), &out_y_height)); // dst_rows
|
||||
CL_CHECK(clSetKernelArg(s->krnl, 10, sizeof(cl_int), &out_y_width)); // dst_cols
|
||||
CL_CHECK(clSetKernelArg(s->krnl, 11, sizeof(cl_mem), &s->m_y_cl)); // M
|
||||
|
||||
const size_t work_size_y[2] = {(size_t)out_y_width, (size_t)out_y_height};
|
||||
|
||||
CL_CHECK(clEnqueueNDRangeKernel(q, s->krnl, 2, NULL,
|
||||
(const size_t*)&work_size_y, NULL, 0, 0, NULL));
|
||||
|
||||
const size_t work_size_uv[2] = {(size_t)out_uv_width, (size_t)out_uv_height};
|
||||
|
||||
CL_CHECK(clSetKernelArg(s->krnl, 2, sizeof(cl_int), &in_uv_px_stride)); // src_px_stride
|
||||
CL_CHECK(clSetKernelArg(s->krnl, 3, sizeof(cl_int), &in_u_offset)); // src_offset
|
||||
CL_CHECK(clSetKernelArg(s->krnl, 4, sizeof(cl_int), &in_uv_height)); // src_rows
|
||||
CL_CHECK(clSetKernelArg(s->krnl, 5, sizeof(cl_int), &in_uv_width)); // src_cols
|
||||
CL_CHECK(clSetKernelArg(s->krnl, 6, sizeof(cl_mem), &out_u)); // dst
|
||||
CL_CHECK(clSetKernelArg(s->krnl, 7, sizeof(cl_int), &out_uv_width)); // dst_row_stride
|
||||
CL_CHECK(clSetKernelArg(s->krnl, 8, sizeof(cl_int), &zero)); // dst_offset
|
||||
CL_CHECK(clSetKernelArg(s->krnl, 9, sizeof(cl_int), &out_uv_height)); // dst_rows
|
||||
CL_CHECK(clSetKernelArg(s->krnl, 10, sizeof(cl_int), &out_uv_width)); // dst_cols
|
||||
CL_CHECK(clSetKernelArg(s->krnl, 11, sizeof(cl_mem), &s->m_uv_cl)); // M
|
||||
|
||||
CL_CHECK(clEnqueueNDRangeKernel(q, s->krnl, 2, NULL,
|
||||
(const size_t*)&work_size_uv, NULL, 0, 0, NULL));
|
||||
CL_CHECK(clSetKernelArg(s->krnl, 3, sizeof(cl_int), &in_v_offset)); // src_ofset
|
||||
CL_CHECK(clSetKernelArg(s->krnl, 6, sizeof(cl_mem), &out_v)); // dst
|
||||
|
||||
CL_CHECK(clEnqueueNDRangeKernel(q, s->krnl, 2, NULL,
|
||||
(const size_t*)&work_size_uv, NULL, 0, 0, NULL));
|
||||
}
|
||||
@@ -1,54 +0,0 @@
|
||||
#define INTER_BITS 5
|
||||
#define INTER_TAB_SIZE (1 << INTER_BITS)
|
||||
#define INTER_SCALE 1.f / INTER_TAB_SIZE
|
||||
|
||||
#define INTER_REMAP_COEF_BITS 15
|
||||
#define INTER_REMAP_COEF_SCALE (1 << INTER_REMAP_COEF_BITS)
|
||||
|
||||
__kernel void warpPerspective(__global const uchar * src,
|
||||
int src_row_stride, int src_px_stride, int src_offset, int src_rows, int src_cols,
|
||||
__global uchar * dst,
|
||||
int dst_row_stride, int dst_offset, int dst_rows, int dst_cols,
|
||||
__constant float * M)
|
||||
{
|
||||
int dx = get_global_id(0);
|
||||
int dy = get_global_id(1);
|
||||
|
||||
if (dx < dst_cols && dy < dst_rows)
|
||||
{
|
||||
float X0 = M[0] * dx + M[1] * dy + M[2];
|
||||
float Y0 = M[3] * dx + M[4] * dy + M[5];
|
||||
float W = M[6] * dx + M[7] * dy + M[8];
|
||||
W = W != 0.0f ? INTER_TAB_SIZE / W : 0.0f;
|
||||
int X = rint(X0 * W), Y = rint(Y0 * W);
|
||||
|
||||
int sx = convert_short_sat(X >> INTER_BITS);
|
||||
int sy = convert_short_sat(Y >> INTER_BITS);
|
||||
|
||||
short sx_clamp = clamp(sx, 0, src_cols - 1);
|
||||
short sx_p1_clamp = clamp(sx + 1, 0, src_cols - 1);
|
||||
short sy_clamp = clamp(sy, 0, src_rows - 1);
|
||||
short sy_p1_clamp = clamp(sy + 1, 0, src_rows - 1);
|
||||
int v0 = convert_int(src[mad24(sy_clamp, src_row_stride, src_offset + sx_clamp*src_px_stride)]);
|
||||
int v1 = convert_int(src[mad24(sy_clamp, src_row_stride, src_offset + sx_p1_clamp*src_px_stride)]);
|
||||
int v2 = convert_int(src[mad24(sy_p1_clamp, src_row_stride, src_offset + sx_clamp*src_px_stride)]);
|
||||
int v3 = convert_int(src[mad24(sy_p1_clamp, src_row_stride, src_offset + sx_p1_clamp*src_px_stride)]);
|
||||
|
||||
short ay = (short)(Y & (INTER_TAB_SIZE - 1));
|
||||
short ax = (short)(X & (INTER_TAB_SIZE - 1));
|
||||
float taby = 1.f/INTER_TAB_SIZE*ay;
|
||||
float tabx = 1.f/INTER_TAB_SIZE*ax;
|
||||
|
||||
int dst_index = mad24(dy, dst_row_stride, dst_offset + dx);
|
||||
|
||||
int itab0 = convert_short_sat_rte( (1.0f-taby)*(1.0f-tabx) * INTER_REMAP_COEF_SCALE );
|
||||
int itab1 = convert_short_sat_rte( (1.0f-taby)*tabx * INTER_REMAP_COEF_SCALE );
|
||||
int itab2 = convert_short_sat_rte( taby*(1.0f-tabx) * INTER_REMAP_COEF_SCALE );
|
||||
int itab3 = convert_short_sat_rte( taby*tabx * INTER_REMAP_COEF_SCALE );
|
||||
|
||||
int val = v0 * itab0 + v1 * itab1 + v2 * itab2 + v3 * itab3;
|
||||
|
||||
uchar pix = convert_uchar_sat((val + (1 << (INTER_REMAP_COEF_BITS-1))) >> INTER_REMAP_COEF_BITS);
|
||||
dst[dst_index] = pix;
|
||||
}
|
||||
}
|
||||
@@ -1,25 +0,0 @@
|
||||
#pragma once
|
||||
|
||||
#define CL_USE_DEPRECATED_OPENCL_1_2_APIS
|
||||
#ifdef __APPLE__
|
||||
#include <OpenCL/cl.h>
|
||||
#else
|
||||
#include <CL/cl.h>
|
||||
#endif
|
||||
|
||||
#include "common/mat.h"
|
||||
|
||||
typedef struct {
|
||||
cl_kernel krnl;
|
||||
cl_mem m_y_cl, m_uv_cl;
|
||||
} Transform;
|
||||
|
||||
void transform_init(Transform* s, cl_context ctx, cl_device_id device_id);
|
||||
|
||||
void transform_destroy(Transform* transform);
|
||||
|
||||
void transform_queue(Transform* s, cl_command_queue q,
|
||||
cl_mem yuv, int in_width, int in_height, int in_stride, int in_uv_offset,
|
||||
cl_mem out_y, cl_mem out_u, cl_mem out_v,
|
||||
int out_width, int out_height,
|
||||
const mat3& projection);
|
||||
@@ -0,0 +1,137 @@
|
||||
import pickle
|
||||
import time
|
||||
import numpy as np
|
||||
from pathlib import Path
|
||||
from tinygrad.tensor import Tensor
|
||||
from tinygrad.engine.jit import TinyJit
|
||||
from tinygrad.device import Device
|
||||
|
||||
from openpilot.system.camerad.cameras.nv12_info import get_nv12_info
|
||||
from openpilot.selfdrive.modeld.compile_warp import (
|
||||
CAMERA_CONFIGS, MEDMODEL_INPUT_SIZE, make_frame_prepare, make_update_both_imgs,
|
||||
warp_pkl_path,
|
||||
)
|
||||
|
||||
MODELS_DIR = Path(__file__).parent / 'models'
|
||||
MODEL_W, MODEL_H = MEDMODEL_INPUT_SIZE
|
||||
UPSTREAM_BUFFER_LENGTH = 5
|
||||
|
||||
|
||||
def v2_warp_pkl_path(cam_w, cam_h, buffer_length):
|
||||
return MODELS_DIR / f'warp_{cam_w}x{cam_h}_b{buffer_length}_tinygrad.pkl'
|
||||
|
||||
|
||||
def compile_v2_warp(cam_w, cam_h, buffer_length):
|
||||
_, _, _, yuv_size = get_nv12_info(cam_w, cam_h)
|
||||
img_buffer_shape = (buffer_length * 6, MODEL_H // 2, MODEL_W // 2)
|
||||
|
||||
print(f"Compiling v2 warp for {cam_w}x{cam_h} buffer_length={buffer_length}...")
|
||||
|
||||
frame_prepare = make_frame_prepare(cam_w, cam_h, MODEL_W, MODEL_H)
|
||||
update_both_imgs = make_update_both_imgs(frame_prepare, MODEL_W, MODEL_H)
|
||||
update_img_jit = TinyJit(update_both_imgs, prune=True)
|
||||
|
||||
full_buffer = Tensor.zeros(img_buffer_shape, dtype='uint8').contiguous().realize()
|
||||
big_full_buffer = Tensor.zeros(img_buffer_shape, dtype='uint8').contiguous().realize()
|
||||
full_buffer_np = np.zeros(img_buffer_shape, dtype=np.uint8)
|
||||
big_full_buffer_np = np.zeros(img_buffer_shape, dtype=np.uint8)
|
||||
|
||||
for i in range(10):
|
||||
new_frame_np = (32 * np.random.randn(yuv_size).astype(np.float32) + 128).clip(0, 255).astype(np.uint8)
|
||||
img_inputs = [full_buffer,
|
||||
Tensor.from_blob(new_frame_np.ctypes.data, (yuv_size,), dtype='uint8').realize(),
|
||||
Tensor(Tensor.randn(3, 3).mul(8).realize().numpy(), device='NPY')]
|
||||
new_big_frame_np = (32 * np.random.randn(yuv_size).astype(np.float32) + 128).clip(0, 255).astype(np.uint8)
|
||||
big_img_inputs = [big_full_buffer,
|
||||
Tensor.from_blob(new_big_frame_np.ctypes.data, (yuv_size,), dtype='uint8').realize(),
|
||||
Tensor(Tensor.randn(3, 3).mul(8).realize().numpy(), device='NPY')]
|
||||
inputs = img_inputs + big_img_inputs
|
||||
Device.default.synchronize()
|
||||
|
||||
inputs_np = [x.numpy() for x in inputs]
|
||||
inputs_np[0] = full_buffer_np
|
||||
inputs_np[3] = big_full_buffer_np
|
||||
|
||||
st = time.perf_counter()
|
||||
out = update_img_jit(*inputs)
|
||||
full_buffer = out[0].contiguous().realize().clone()
|
||||
big_full_buffer = out[2].contiguous().realize().clone()
|
||||
mt = time.perf_counter()
|
||||
Device.default.synchronize()
|
||||
et = time.perf_counter()
|
||||
print(f" [{i+1}/10] enqueue {(mt-st)*1e3:6.2f} ms -- total {(et-st)*1e3:6.2f} ms")
|
||||
|
||||
pkl_path = v2_warp_pkl_path(cam_w, cam_h, buffer_length)
|
||||
with open(pkl_path, "wb") as f:
|
||||
pickle.dump(update_img_jit, f)
|
||||
print(f" Saved to {pkl_path}")
|
||||
|
||||
jit = pickle.load(open(pkl_path, "rb"))
|
||||
jit(*inputs)
|
||||
|
||||
|
||||
class Warp:
|
||||
def __init__(self, buffer_length=2):
|
||||
self.buffer_length = buffer_length
|
||||
self.img_buffer_shape = (buffer_length * 6, MODEL_H // 2, MODEL_W // 2)
|
||||
|
||||
self.jit_cache = {}
|
||||
self.full_buffers = {k: Tensor.zeros(self.img_buffer_shape, dtype='uint8').contiguous().realize() for k in ['img', 'big_img']}
|
||||
self._blob_cache: dict[int, Tensor] = {}
|
||||
self._nv12_cache: dict[tuple[int, int], int] = {}
|
||||
self.transforms_np = {k: np.zeros((3, 3), dtype=np.float32) for k in ['img', 'big_img']}
|
||||
self.transforms = {k: Tensor(v, device='NPY').realize() for k, v in self.transforms_np.items()}
|
||||
|
||||
def process(self, bufs, transforms):
|
||||
if not bufs:
|
||||
return {}
|
||||
road = next(n for n in bufs if 'big' not in n)
|
||||
wide = next(n for n in bufs if 'big' in n)
|
||||
cam_w, cam_h = bufs[road].width, bufs[road].height
|
||||
key = (cam_w, cam_h)
|
||||
|
||||
if key not in self.jit_cache:
|
||||
v2_pkl = v2_warp_pkl_path(cam_w, cam_h, self.buffer_length)
|
||||
if v2_pkl.exists():
|
||||
with open(v2_pkl, 'rb') as f:
|
||||
self.jit_cache[key] = pickle.load(f)
|
||||
elif self.buffer_length == UPSTREAM_BUFFER_LENGTH:
|
||||
upstream_pkl = warp_pkl_path(cam_w, cam_h)
|
||||
if upstream_pkl.exists():
|
||||
with open(upstream_pkl, 'rb') as f:
|
||||
self.jit_cache[key] = pickle.load(f)
|
||||
if key not in self.jit_cache:
|
||||
frame_prepare = make_frame_prepare(cam_w, cam_h, MODEL_W, MODEL_H)
|
||||
update_both_imgs = make_update_both_imgs(frame_prepare, MODEL_W, MODEL_H)
|
||||
self.jit_cache[key] = TinyJit(update_both_imgs, prune=True)
|
||||
|
||||
if key not in self._nv12_cache:
|
||||
self._nv12_cache[key] = get_nv12_info(cam_w, cam_h)[3]
|
||||
yuv_size = self._nv12_cache[key]
|
||||
|
||||
road_ptr = bufs[road].data.ctypes.data
|
||||
wide_ptr = bufs[wide].data.ctypes.data
|
||||
if road_ptr not in self._blob_cache:
|
||||
self._blob_cache[road_ptr] = Tensor.from_blob(road_ptr, (yuv_size,), dtype='uint8')
|
||||
if wide_ptr not in self._blob_cache:
|
||||
self._blob_cache[wide_ptr] = Tensor.from_blob(wide_ptr, (yuv_size,), dtype='uint8')
|
||||
road_blob = self._blob_cache[road_ptr]
|
||||
wide_blob = self._blob_cache[wide_ptr] if wide_ptr != road_ptr else Tensor.from_blob(wide_ptr, (yuv_size,), dtype='uint8')
|
||||
np.copyto(self.transforms_np['img'], transforms[road].reshape(3, 3))
|
||||
np.copyto(self.transforms_np['big_img'], transforms[wide].reshape(3, 3))
|
||||
|
||||
Device.default.synchronize()
|
||||
res = self.jit_cache[key](
|
||||
self.full_buffers['img'], road_blob, self.transforms['img'],
|
||||
self.full_buffers['big_img'], wide_blob, self.transforms['big_img'],
|
||||
)
|
||||
self.full_buffers['img'], out_road = res[0].realize(), res[1].realize()
|
||||
self.full_buffers['big_img'], out_wide = res[2].realize(), res[3].realize()
|
||||
|
||||
return {road: out_road, wide: out_wide}
|
||||
|
||||
|
||||
if __name__ == "__main__":
|
||||
for cam_w, cam_h in CAMERA_CONFIGS:
|
||||
for bl in [2, 5]:
|
||||
compile_v2_warp(cam_w, cam_h, bl)
|
||||
@@ -116,7 +116,7 @@ class ModelCache:
|
||||
|
||||
class ModelFetcher:
|
||||
"""Handles fetching and caching of model data from remote source"""
|
||||
MODEL_URL = "https://raw.githubusercontent.com/sunnypilot/sunnypilot-docs/refs/heads/gh-pages/docs/driving_models_v14.json"
|
||||
MODEL_URL = "https://raw.githubusercontent.com/sunnypilot/sunnypilot-docs/refs/heads/gh-pages/docs/driving_models_v15.json"
|
||||
|
||||
def __init__(self, params: Params):
|
||||
self.params = params
|
||||
|
||||
@@ -13,16 +13,13 @@ import numpy as np
|
||||
from openpilot.common.params import Params
|
||||
from cereal import custom
|
||||
from openpilot.sunnypilot.modeld.constants import Meta, MetaTombRaider, MetaSimPose
|
||||
from openpilot.sunnypilot.modeld.runners import ModelRunner
|
||||
from openpilot.system.hardware import PC
|
||||
from openpilot.system.hardware.hw import Paths
|
||||
from pathlib import Path
|
||||
|
||||
# see the README.md for more details on the model selector versioning
|
||||
CURRENT_SELECTOR_VERSION = 14
|
||||
CURRENT_SELECTOR_VERSION = 15
|
||||
REQUIRED_MIN_SELECTOR_VERSION = 14
|
||||
|
||||
USE_ONNX = os.getenv('USE_ONNX', PC)
|
||||
|
||||
CUSTOM_MODEL_PATH = Paths.model_root()
|
||||
METADATA_PATH = Path(__file__).parent / '../models/supercombo_metadata.pkl'
|
||||
@@ -122,16 +119,6 @@ def _get_model():
|
||||
|
||||
return None
|
||||
|
||||
def get_model_path():
|
||||
if USE_ONNX:
|
||||
return {ModelRunner.ONNX: Path(__file__).parent / '../models/supercombo.onnx'}
|
||||
|
||||
if model := _get_model():
|
||||
return {ModelRunner.THNEED: f"{CUSTOM_MODEL_PATH}/{model.artifact.fileName}"}
|
||||
|
||||
return {ModelRunner.THNEED: Path(__file__).parent / '../models/supercombo.thneed'}
|
||||
|
||||
|
||||
def load_metadata():
|
||||
metadata_path = METADATA_PATH
|
||||
|
||||
|
||||
@@ -1,6 +1,5 @@
|
||||
import os
|
||||
import numpy as np
|
||||
from openpilot.sunnypilot.modeld_v2.models.commonmodel_pyx import DrivingModelFrame, CLMem
|
||||
from openpilot.system.hardware.hw import Paths
|
||||
from cereal import custom
|
||||
|
||||
@@ -8,8 +7,6 @@ from cereal import custom
|
||||
NumpyDict = dict[str, np.ndarray]
|
||||
ShapeDict = dict[str, tuple[int, ...]]
|
||||
SliceDict = dict[str, slice]
|
||||
CLMemDict = dict[str, CLMem]
|
||||
FrameDict = dict[str, DrivingModelFrame]
|
||||
|
||||
ModelType = custom.ModelManagerSP.Model.Type
|
||||
Model = custom.ModelManagerSP.Model
|
||||
|
||||
@@ -2,7 +2,7 @@ from abc import abstractmethod, ABC
|
||||
|
||||
import numpy as np
|
||||
from openpilot.sunnypilot.models.helpers import get_active_bundle
|
||||
from openpilot.sunnypilot.models.runners.constants import NumpyDict, ShapeDict, CLMemDict, FrameDict, Model, SliceDict, SEND_RAW_PRED
|
||||
from openpilot.sunnypilot.models.runners.constants import NumpyDict, ShapeDict, Model, SliceDict, SEND_RAW_PRED
|
||||
from openpilot.system.hardware.hw import Paths
|
||||
import pickle
|
||||
|
||||
@@ -133,13 +133,11 @@ class ModelRunner(ModularRunner):
|
||||
raise ValueError("Model data is not available. Ensure the model is loaded correctly.")
|
||||
|
||||
@abstractmethod
|
||||
def prepare_inputs(self, imgs_cl: CLMemDict, numpy_inputs: NumpyDict, frames: FrameDict) -> dict:
|
||||
def prepare_inputs(self, numpy_inputs: NumpyDict) -> dict:
|
||||
"""
|
||||
Abstract method to prepare inputs for model inference.
|
||||
|
||||
:param imgs_cl: Dictionary of OpenCL memory objects for image inputs.
|
||||
:param numpy_inputs: Dictionary of numpy arrays for non-image inputs.
|
||||
:param frames: Dictionary of DrivingModelFrame objects for context.
|
||||
:return: Dictionary of prepared inputs ready for the model.
|
||||
"""
|
||||
raise NotImplementedError
|
||||
|
||||
@@ -1,62 +0,0 @@
|
||||
import numpy as np
|
||||
|
||||
from openpilot.sunnypilot.modeld_v2 import MODEL_PATH
|
||||
from openpilot.sunnypilot.modeld_v2.runners.ort_helpers import make_onnx_cpu_runner, ORT_TYPES_TO_NP_TYPES
|
||||
from openpilot.sunnypilot.models.runners.constants import ModelType, ShapeDict, CLMemDict, NumpyDict, FrameDict
|
||||
from openpilot.sunnypilot.models.runners.model_runner import ModelRunner
|
||||
from openpilot.sunnypilot.modeld_v2.constants import ModelConstants
|
||||
|
||||
|
||||
class ONNXRunner(ModelRunner):
|
||||
"""
|
||||
A ModelRunner implementation for executing ONNX models using ONNX Runtime CPU.
|
||||
|
||||
Handles loading the ONNX model, preparing inputs as numpy arrays, running
|
||||
inference, and parsing outputs. This runner is typically used on non-TICI platforms.
|
||||
"""
|
||||
def __init__(self):
|
||||
super().__init__()
|
||||
# Initialize ONNX Runtime session for the model at MODEL_PATH
|
||||
self.runner = make_onnx_cpu_runner(MODEL_PATH)
|
||||
# Map expected input names to numpy dtypes
|
||||
self.input_to_nptype = {
|
||||
model_input.name: ORT_TYPES_TO_NP_TYPES[model_input.type]
|
||||
for model_input in self.runner.get_inputs()
|
||||
}
|
||||
# For ONNX, _model_data isn't strictly necessary as shapes/types come from the runner
|
||||
# However, we might still need output_slices if custom models define them.
|
||||
# We assume supercombo type for potentially loading output_slices metadata if available.
|
||||
self._model_data = self.models.get(ModelType.supercombo)
|
||||
self._constants = ModelConstants # Constants for ONNX models, if needed
|
||||
|
||||
@property
|
||||
def input_shapes(self) -> ShapeDict:
|
||||
"""Returns the input shapes defined in the ONNX model."""
|
||||
# ONNX shapes are derived directly from the model definition via the runner
|
||||
return {runner_input.name: runner_input.shape for runner_input in self.runner.get_inputs()}
|
||||
|
||||
def prepare_inputs(self, imgs_cl: CLMemDict, numpy_inputs: NumpyDict, frames: FrameDict) -> dict:
|
||||
"""Prepares inputs for the ONNX model as numpy arrays."""
|
||||
self.inputs = numpy_inputs # Start with non-image numpy inputs
|
||||
# Convert image inputs from OpenCL buffers to numpy arrays
|
||||
for key in imgs_cl:
|
||||
buffer = frames[key].buffer_from_cl(imgs_cl[key])
|
||||
reshaped_buffer = buffer.reshape(self.input_shapes[key])
|
||||
self.inputs[key] = reshaped_buffer.astype(dtype=self.input_to_nptype[key])
|
||||
return self.inputs
|
||||
|
||||
def _parse_outputs(self, model_outputs: np.ndarray) -> NumpyDict:
|
||||
"""Parses the raw ONNX model outputs using the standard Parser."""
|
||||
# Use slicing if metadata is available, otherwise pass raw outputs
|
||||
if self._model_data is None:
|
||||
raise ValueError("Model data is not available. Ensure the model is loaded correctly.")
|
||||
|
||||
outputs_to_parse = self._slice_outputs(model_outputs) if self._model_data else {'raw_pred': model_outputs}
|
||||
result: NumpyDict = self.parser_method_dict[self._model_data.model.type.raw](outputs_to_parse)
|
||||
return result
|
||||
|
||||
def _run_model(self) -> NumpyDict:
|
||||
"""Runs the ONNX model inference and parses the outputs."""
|
||||
# Execute the ONNX Runtime session
|
||||
outputs = self.runner.run(None, self.inputs)[0].flatten()
|
||||
return self._parse_outputs(outputs)
|
||||
@@ -1,11 +1,9 @@
|
||||
import pickle
|
||||
|
||||
import numpy as np
|
||||
from openpilot.sunnypilot.modeld_v2.runners.tinygrad_helpers import qcom_tensor_from_opencl_address
|
||||
from openpilot.sunnypilot.models.runners.constants import CLMemDict, FrameDict, NumpyDict, ModelType, ShapeDict, CUSTOM_MODEL_PATH, SliceDict
|
||||
from openpilot.sunnypilot.models.runners.constants import NumpyDict, ModelType, ShapeDict, CUSTOM_MODEL_PATH, SliceDict
|
||||
from openpilot.sunnypilot.models.runners.model_runner import ModelRunner
|
||||
from openpilot.sunnypilot.models.runners.tinygrad.model_types import PolicyTinygrad, VisionTinygrad, SupercomboTinygrad, OffPolicyTinygrad
|
||||
from openpilot.system.hardware import TICI
|
||||
from openpilot.sunnypilot.models.split_model_constants import SplitModelConstants
|
||||
from openpilot.sunnypilot.modeld_v2.constants import ModelConstants
|
||||
|
||||
@@ -54,37 +52,31 @@ class TinygradRunner(ModelRunner, SupercomboTinygrad, PolicyTinygrad, VisionTiny
|
||||
info = self.model_run.captured.expected_input_info[idx]
|
||||
self.input_to_dtype[name] = info[2] # dtype
|
||||
self.input_to_device[name] = info[3] # device
|
||||
self._policy_cached = False
|
||||
|
||||
@property
|
||||
def vision_input_names(self) -> list[str]:
|
||||
"""Returns the list of vision input names from the input shapes."""
|
||||
return [name for name in self.input_shapes.keys() if 'img' in name]
|
||||
|
||||
def prepare_vision_inputs(self, imgs_cl: CLMemDict, frames: FrameDict):
|
||||
"""Prepares vision (image) inputs as Tinygrad Tensors."""
|
||||
for key in imgs_cl:
|
||||
if TICI and key not in self.inputs:
|
||||
# On TICI, directly use OpenCL memory address for efficiency via QCOM extensions
|
||||
self.inputs[key] = qcom_tensor_from_opencl_address(imgs_cl[key].mem_address, self.input_shapes[key], dtype=self.input_to_dtype[key])
|
||||
elif not TICI:
|
||||
# On other platforms, copy data from CL buffer to a numpy array first
|
||||
shape = frames[key].buffer_from_cl(imgs_cl[key]).reshape(self.input_shapes[key])
|
||||
self.inputs[key] = Tensor(shape, device=self.input_to_device[key], dtype=self.input_to_dtype[key]).realize()
|
||||
|
||||
def prepare_policy_inputs(self, numpy_inputs: NumpyDict):
|
||||
"""Prepares non-image (policy) inputs as Tinygrad Tensors."""
|
||||
for key, value in numpy_inputs.items():
|
||||
self.inputs[key] = Tensor(value, device=self.input_to_device[key], dtype=self.input_to_dtype[key]).realize()
|
||||
if not self._policy_cached:
|
||||
for key, value in numpy_inputs.items():
|
||||
self.inputs[key] = Tensor(value, device='NPY').realize()
|
||||
self._policy_cached = True
|
||||
|
||||
def prepare_inputs(self, imgs_cl: CLMemDict, numpy_inputs: NumpyDict, frames: FrameDict) -> dict:
|
||||
def prepare_inputs(self, numpy_inputs: NumpyDict) -> dict:
|
||||
"""Prepares all vision and policy inputs for the model."""
|
||||
self.prepare_vision_inputs(imgs_cl, frames)
|
||||
self.prepare_policy_inputs(numpy_inputs)
|
||||
for key in self.vision_input_names:
|
||||
if key in self.inputs:
|
||||
self.inputs[key] = self.inputs[key].cast(self.input_to_dtype[key])
|
||||
return self.inputs
|
||||
|
||||
def _run_model(self) -> NumpyDict:
|
||||
"""Runs the Tinygrad model inference and parses the outputs."""
|
||||
outputs = self.model_run(**self.inputs).numpy().flatten()
|
||||
outputs = self.model_run(**self.inputs).contiguous().realize().uop.base.buffer.numpy().flatten()
|
||||
return self._parse_outputs(outputs)
|
||||
|
||||
def _parse_outputs(self, model_outputs: np.ndarray) -> NumpyDict:
|
||||
@@ -120,6 +112,9 @@ class TinygradSplitRunner(ModelRunner):
|
||||
off_policy_output = self.off_policy_runner.run_model()
|
||||
outputs.update(off_policy_output)
|
||||
|
||||
if 'planplus' in outputs and 'plan' in outputs:
|
||||
outputs['plan'] = outputs['plan'] + outputs['planplus']
|
||||
|
||||
return outputs
|
||||
|
||||
@property
|
||||
@@ -143,17 +138,18 @@ class TinygradSplitRunner(ModelRunner):
|
||||
slices.update(self.off_policy_runner.output_slices)
|
||||
return slices
|
||||
|
||||
def prepare_inputs(self, imgs_cl: CLMemDict, numpy_inputs: NumpyDict, frames: FrameDict) -> dict:
|
||||
def prepare_inputs(self, numpy_inputs: NumpyDict) -> dict:
|
||||
"""Prepares inputs for both vision and policy models."""
|
||||
# Policy inputs only depend on numpy_inputs
|
||||
self.policy_runner.prepare_policy_inputs(numpy_inputs)
|
||||
# Vision inputs depend on imgs_cl and frames
|
||||
self.vision_runner.prepare_vision_inputs(imgs_cl, frames)
|
||||
|
||||
for key in self.vision_input_names:
|
||||
if key in self.inputs:
|
||||
self.vision_runner.inputs[key] = self.inputs[key].cast(self.vision_runner.input_to_dtype[key])
|
||||
|
||||
inputs = {**self.policy_runner.inputs, **self.vision_runner.inputs}
|
||||
|
||||
if self.off_policy_runner:
|
||||
self.off_policy_runner.prepare_policy_inputs(numpy_inputs)
|
||||
inputs.update(self.off_policy_runner.inputs)
|
||||
|
||||
# Return combined inputs (though they are stored within respective runners)
|
||||
return inputs
|
||||
|
||||
@@ -80,10 +80,6 @@ def use_sunnylink_uploader_shim(started, params, CP: car.CarParams) -> bool:
|
||||
"""Shim for use_sunnylink_uploader to match the process manager signature."""
|
||||
return use_sunnylink_uploader(params)
|
||||
|
||||
def is_snpe_model(started, params, CP: car.CarParams) -> bool:
|
||||
"""Check if the active model runner is SNPE."""
|
||||
return bool(get_active_model_runner(params, not started) == custom.ModelManagerSP.Runner.snpe)
|
||||
|
||||
def is_tinygrad_model(started, params, CP: car.CarParams) -> bool:
|
||||
"""Check if the active model runner is SNPE."""
|
||||
return bool(get_active_model_runner(params, not started) == custom.ModelManagerSP.Runner.tinygrad)
|
||||
@@ -170,7 +166,6 @@ procs = [
|
||||
procs += [
|
||||
# Models
|
||||
PythonProcess("models_manager", "sunnypilot.models.manager", only_offroad),
|
||||
NativeProcess("modeld_snpe", "sunnypilot/modeld", ["./modeld"], and_(only_onroad, is_snpe_model)),
|
||||
NativeProcess("modeld_tinygrad", "sunnypilot/modeld_v2", ["./modeld"], and_(only_onroad, is_tinygrad_model)),
|
||||
|
||||
# Backup
|
||||
|
||||
Reference in New Issue
Block a user