mirror of
https://github.com/sunnypilot/sunnypilot.git
synced 2026-02-18 21:14:01 +08:00
Sync: commaai/openpilot:master into sunnypilot/sunnypilot:master-new (#488)
This commit is contained in:
Jason Wen
@@ -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. |
|
||||
|
||||
Submodule msgq_repo updated: 434ed2312c...5bb86f8bc7
2
panda
2
panda
Submodule panda updated: 45301bf15c...0d4b79a3c7
@@ -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",
|
||||
|
||||
@@ -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",
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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()
|
||||
|
||||
@@ -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'])
|
||||
|
||||
@@ -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" "$@"
|
||||
|
||||
@@ -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):
|
||||
|
||||
@@ -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" "$@"
|
||||
|
||||
@@ -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
|
||||
|
||||
|
||||
|
||||
@@ -7,7 +7,7 @@
|
||||
|
||||
DrivingModelFrame::DrivingModelFrame(cl_device_id device_id, cl_context context) : ModelFrame(device_id, context) {
|
||||
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));
|
||||
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<uint8_t[]>(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() {
|
||||
|
||||
@@ -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;
|
||||
};
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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(<cnp.uint8_t[:self.buf_size]> data)
|
||||
# return CLMem.create(data)
|
||||
|
||||
# def buffer_from_cl(self, CLMem in_frames):
|
||||
# cdef unsigned char * data2
|
||||
# data2 = self.frame.buffer_from_cl(in_frames.mem, self.buf_size)
|
||||
# return np.asarray(<cnp.uint8_t[:self.buf_size]> data2)
|
||||
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):
|
||||
@@ -74,3 +67,4 @@ cdef class MonitoringModelFrame(ModelFrame):
|
||||
self._frame = new cppMonitoringModelFrame(context.device_id, context.context)
|
||||
self.frame = <cppModelFrame*>(self._frame)
|
||||
self.buf_size = self._frame.buf_size
|
||||
|
||||
|
||||
@@ -1,3 +1,3 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:0c896681fd6851de3968433e12f37834429eba265e938cf383200be3e5835cec
|
||||
oid sha256:72d3d6f8d3c98f5431ec86be77b6350d7d4f43c25075c0106f1d1e7ec7c77668
|
||||
size 49096168
|
||||
|
||||
@@ -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)
|
||||
@@ -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
|
||||
@@ -1,4 +0,0 @@
|
||||
#pragma once
|
||||
|
||||
#include "selfdrive/modeld/runners/runmodel.h"
|
||||
#include "selfdrive/modeld/runners/snpemodel.h"
|
||||
@@ -1,49 +0,0 @@
|
||||
#pragma once
|
||||
|
||||
#include <string>
|
||||
#include <vector>
|
||||
#include <memory>
|
||||
#include <cassert>
|
||||
|
||||
#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<std::unique_ptr<ModelInput>> 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<ModelInput>(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);
|
||||
}
|
||||
};
|
||||
@@ -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()
|
||||
@@ -1,6 +0,0 @@
|
||||
# distutils: language = c++
|
||||
|
||||
from .runmodel cimport RunModel as cppRunModel
|
||||
|
||||
cdef class RunModel:
|
||||
cdef cppRunModel * model
|
||||
@@ -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()
|
||||
@@ -1,116 +0,0 @@
|
||||
#pragma clang diagnostic ignored "-Wexceptions"
|
||||
|
||||
#include "selfdrive/modeld/runners/snpemodel.h"
|
||||
|
||||
#include <cstring>
|
||||
#include <memory>
|
||||
#include <string>
|
||||
#include <utility>
|
||||
#include <vector>
|
||||
|
||||
#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<zdl::DlContainer::IDlContainer> 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<size_t> 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<size_t> 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<SNPEModelInput>(new SNPEModelInput(name, buffer, size, std::move(input_buffer))));
|
||||
}
|
||||
|
||||
void SNPEModel::execute() {
|
||||
if (!snpe->execute(input_map, output_map)) {
|
||||
PrintErrorStringAndExit();
|
||||
}
|
||||
}
|
||||
@@ -1,52 +0,0 @@
|
||||
#pragma once
|
||||
#pragma clang diagnostic ignored "-Wdeprecated-declarations"
|
||||
|
||||
#include <memory>
|
||||
#include <string>
|
||||
#include <utility>
|
||||
|
||||
#include <DlContainer/IDlContainer.hpp>
|
||||
#include <DlSystem/DlError.hpp>
|
||||
#include <DlSystem/ITensor.hpp>
|
||||
#include <DlSystem/ITensorFactory.hpp>
|
||||
#include <DlSystem/IUserBuffer.hpp>
|
||||
#include <DlSystem/IUserBufferFactory.hpp>
|
||||
#include <SNPE/SNPE.hpp>
|
||||
#include <SNPE/SNPEBuilder.hpp>
|
||||
#include <SNPE/SNPEFactory.hpp>
|
||||
|
||||
#include "selfdrive/modeld/runners/runmodel.h"
|
||||
|
||||
struct SNPEModelInput : public ModelInput {
|
||||
std::unique_ptr<zdl::DlSystem::IUserBuffer> snpe_buffer;
|
||||
|
||||
SNPEModelInput(const std::string _name, float *_buffer, int _size, std::unique_ptr<zdl::DlSystem::IUserBuffer> _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<zdl::SNPE::SNPE> snpe;
|
||||
zdl::DlSystem::UserBufferMap input_map;
|
||||
zdl::DlSystem::UserBufferMap output_map;
|
||||
std::unique_ptr<zdl::DlSystem::IUserBuffer> output_buffer;
|
||||
|
||||
bool use_tf8;
|
||||
float *output;
|
||||
size_t output_size;
|
||||
};
|
||||
@@ -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)
|
||||
@@ -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 = <cppRunModel *> new cppSNPEModel(path, &output[0], len(output), runtime, use_tf8, context.context)
|
||||
@@ -1,58 +0,0 @@
|
||||
#include "selfdrive/modeld/runners/thneedmodel.h"
|
||||
|
||||
#include <string>
|
||||
|
||||
#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);
|
||||
}
|
||||
}
|
||||
@@ -1,17 +0,0 @@
|
||||
#pragma once
|
||||
|
||||
#include <string>
|
||||
|
||||
#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;
|
||||
};
|
||||
@@ -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)
|
||||
@@ -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 = <cppRunModel *> new cppThneedModel(path, &output[0], len(output), runtime, use_tf8, context.context)
|
||||
8
selfdrive/modeld/runners/tinygrad_helpers.py
Normal file
8
selfdrive/modeld/runners/tinygrad_helpers.py
Normal file
@@ -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')
|
||||
@@ -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.
|
||||
|
||||
@@ -1,154 +0,0 @@
|
||||
#include <cassert>
|
||||
#include <set>
|
||||
|
||||
#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<cl_program, string> 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<cl_mem, cl_mem> 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<string, cl_program> 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<CLQueuedKernel>(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);
|
||||
}
|
||||
@@ -1,133 +0,0 @@
|
||||
#pragma once
|
||||
|
||||
#ifndef __user
|
||||
#define __user __attribute__(())
|
||||
#endif
|
||||
|
||||
#include <cstdint>
|
||||
#include <cstdlib>
|
||||
#include <memory>
|
||||
#include <string>
|
||||
#include <vector>
|
||||
|
||||
#include <CL/cl.h>
|
||||
|
||||
#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<string> arg_names;
|
||||
vector<string> arg_types;
|
||||
vector<string> args;
|
||||
vector<int> 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<kgsl_command_object[]> cmds;
|
||||
unique_ptr<kgsl_command_object[]> objs;
|
||||
Thneed *thneed;
|
||||
vector<shared_ptr<CLQueuedKernel> > 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<cl_mem> input_clmem;
|
||||
vector<void *> inputs;
|
||||
vector<size_t> 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<GPUMalloc> ram;
|
||||
vector<unique_ptr<CachedIoctl> > cmds;
|
||||
int fd;
|
||||
#endif
|
||||
|
||||
// all CL kernels
|
||||
void copy_inputs(float **finputs, bool internal=false);
|
||||
void copy_output(float *foutput);
|
||||
cl_int clexec();
|
||||
vector<shared_ptr<CLQueuedKernel> > kq;
|
||||
|
||||
// pending CL kernels
|
||||
vector<shared_ptr<CLQueuedKernel> > ckq;
|
||||
|
||||
// loading
|
||||
void load(const char *filename);
|
||||
private:
|
||||
void clinit();
|
||||
};
|
||||
|
||||
@@ -1,216 +0,0 @@
|
||||
#include "selfdrive/modeld/thneed/thneed.h"
|
||||
|
||||
#include <cassert>
|
||||
#include <cstring>
|
||||
#include <map>
|
||||
|
||||
#include "common/clutil.h"
|
||||
#include "common/timing.h"
|
||||
|
||||
map<pair<cl_kernel, int>, string> g_args;
|
||||
map<pair<cl_kernel, int>, int> g_args_size;
|
||||
map<cl_program, string> 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;
|
||||
}
|
||||
@@ -1,32 +0,0 @@
|
||||
#include "selfdrive/modeld/thneed/thneed.h"
|
||||
|
||||
#include <cassert>
|
||||
|
||||
#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);
|
||||
}
|
||||
}
|
||||
@@ -1,258 +0,0 @@
|
||||
#include "selfdrive/modeld/thneed/thneed.h"
|
||||
|
||||
#include <dlfcn.h>
|
||||
#include <sys/mman.h>
|
||||
|
||||
#include <cassert>
|
||||
#include <cerrno>
|
||||
#include <cstring>
|
||||
#include <map>
|
||||
#include <string>
|
||||
|
||||
#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<decltype(my_ioctl)>(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<CachedCommand>(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<CachedSync>(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<struct kgsl_command_object[]>(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<struct kgsl_command_object[]>(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<GPUMalloc>(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);
|
||||
}
|
||||
}
|
||||
@@ -36,7 +36,7 @@ PandaUsbHandle::PandaUsbHandle(std::string serial) : PandaCommsHandle(serial) {
|
||||
for (size_t i = 0; i < num_devices; ++i) {
|
||||
libusb_device_descriptor desc;
|
||||
libusb_get_device_descriptor(dev_list[i], &desc);
|
||||
if (desc.idVendor == 0xbbaa && desc.idProduct == 0xddcc) {
|
||||
if (desc.idVendor == 0x3801 && desc.idProduct == 0xddcc) {
|
||||
int ret = libusb_open(dev_list[i], &dev_handle);
|
||||
if (dev_handle == NULL || ret < 0) { goto fail; }
|
||||
|
||||
@@ -110,7 +110,7 @@ std::vector<std::string> PandaUsbHandle::list() {
|
||||
libusb_device *device = dev_list[i];
|
||||
libusb_device_descriptor desc;
|
||||
libusb_get_device_descriptor(device, &desc);
|
||||
if (desc.idVendor == 0xbbaa && desc.idProduct == 0xddcc) {
|
||||
if (desc.idVendor == 0x3801 && desc.idProduct == 0xddcc) {
|
||||
libusb_device_handle *handle = NULL;
|
||||
int ret = libusb_open(device, &handle);
|
||||
if (ret < 0) { goto finish; }
|
||||
|
||||
@@ -0,0 +1,3 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:737f95d34912db53a303ba6499e6f697b510fa5872b8c71f701a4fe924b5466e
|
||||
size 356169
|
||||
@@ -1,3 +0,0 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:6267aadf44c81fb003130aff6f8902014a1add7f18c5c563239943ab015974c9
|
||||
size 356177
|
||||
@@ -0,0 +1,3 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:e359e8f6b5a22b6f3f89b54989dac2110ee3a4463de2d785be83e20cda4f1cb6
|
||||
size 254248
|
||||
@@ -1,3 +0,0 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:433633d2042cfe1191e5b71ec4b34edca21877377f2ace8aaf6ae43f7b805171
|
||||
size 256336
|
||||
@@ -0,0 +1,3 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:a8e4d044812a714ebdf0b15e73d4466e9ddaafa374368f308803c6b68dcd79ab
|
||||
size 332433
|
||||
@@ -1,3 +0,0 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:64aa7c86bf420adfcf78888bfce7233d68ec30ed77f3f9d3f923f26c59f85c74
|
||||
size 332371
|
||||
@@ -0,0 +1,3 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:53b80c3c99a6897cafe7d408872c287ab0bbaac2751e344cf4623b312d2e4866
|
||||
size 268928
|
||||
@@ -1,3 +0,0 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:aa2451a6697da9fb1cf305081337782eb944a623b7662fd0a58ebbed5d0395e7
|
||||
size 268892
|
||||
@@ -0,0 +1,3 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:a4fa8a841c964e90b9c14b5aede8f30de949d319ae072cc291f0afb1c4c24baa
|
||||
size 437808
|
||||
@@ -1,3 +0,0 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:75766b120bc372bad7b2a80be2c195de9b85cf11142f2a9da348637f15e245db
|
||||
size 437801
|
||||
@@ -0,0 +1,3 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:ce370994de01d6240fa753846ddc7e1b852acac3f649a4c29ea16acae126f974
|
||||
size 308578
|
||||
@@ -1,3 +0,0 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:c692e18506a36eeed22d23d1739af8f26fb3f39f34584fdbcb1ebf3bc6430957
|
||||
size 308611
|
||||
@@ -0,0 +1,3 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:82333cfc026735eac81defae9e01b82b81ac30fd01ad265dbdd311dd97322198
|
||||
size 393106
|
||||
@@ -1,3 +0,0 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:8d62d5774bb4be70ef337179c4bab2347f153fc959155bef917527cc6868bcd3
|
||||
size 393161
|
||||
@@ -0,0 +1,3 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:e4ebae1742e3bb7d89f8cd452928c5c6a1b8679155562d23c6303ca4ec2e3b02
|
||||
size 334350
|
||||
@@ -1,3 +0,0 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:d57c998320672e771712bd0b3f35c07bd159dff472368b1bf1abc9b6a3541dad
|
||||
size 334258
|
||||
@@ -0,0 +1,3 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:03326fb4486c1ffdd3609a0e0200e65d1f30dca5d9b047501d98b1b2d5321e75
|
||||
size 470495
|
||||
@@ -1,3 +0,0 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:fe9370fb0f7cad039d12ddc2f23dd4ee1c94d5d75faf9b6143ec390c06b28c9d
|
||||
size 470471
|
||||
@@ -0,0 +1,3 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:bffa062a5eace4d499c2fe12f2b0eb8c71207f4e7bea0e88456614f00e2859b1
|
||||
size 258989
|
||||
@@ -1,3 +0,0 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:917142016b6ed8c800a2c00aecbfdba7eecda3b1a3fcb15a968fa726b5b9241c
|
||||
size 260324
|
||||
@@ -0,0 +1,3 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:fe97d060fb6e60652b26b26406ca8b03890f6bff4b4bc5cb73fca267068a04c9
|
||||
size 217460
|
||||
@@ -1,3 +0,0 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:d4b0e05a8a1b21e008ed5c1a41f267d596cc35a6d49ec8eff496856b34866360
|
||||
size 217481
|
||||
@@ -0,0 +1,3 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:60e33d209f4f600acf344feca8ab6494a1bf3aafe3c7121e2110a82ea06aa61e
|
||||
size 293084
|
||||
@@ -1,3 +0,0 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:4d8e2889e70c310a84eff142ee481cfc92888bbe257809be3694032c67277675
|
||||
size 293076
|
||||
@@ -0,0 +1,3 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:754d601883db2377739f4112a8dff787b4b8181a8bb8bf4d7dbcf77c5cb6ac93
|
||||
size 256838
|
||||
@@ -1,3 +0,0 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:d60f784a226d144e1ac2c9ceec53b30509cbcceffb5e28f9c1b3a88163958352
|
||||
size 260706
|
||||
@@ -0,0 +1,3 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:378d6fc5c674e62b51a12b994cfcd03c435a9731c6e2f4c0d14a39010540cc77
|
||||
size 100067
|
||||
@@ -1,3 +0,0 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:dc8216930cf4b36dd080f32b6f52172e08c2c4fe946450162950e799345d1829
|
||||
size 100106
|
||||
@@ -0,0 +1,3 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:037825cd9b443d038ab3bb370eab802612d1ce528c5de1350fd99965eaa0d131
|
||||
size 262910
|
||||
@@ -1,3 +0,0 @@
|
||||
version https://git-lfs.github.com/spec/v1
|
||||
oid sha256:cb0195a176d9a917ef87698ef6487547ef19f35e86b4ba12bbecc7b5e6a26295
|
||||
size 262860
|
||||
Some files were not shown because too many files have changed in this diff Show More
Reference in New Issue
Block a user