From c4712a85c6cbe91c5026bd803eda92a704591d38 Mon Sep 17 00:00:00 2001 From: mitchellgoffpc Date: Tue, 10 Sep 2024 11:43:12 -0700 Subject: [PATCH] Slow tinygrad runner, remove thneed runner --- selfdrive/modeld/modeld.py | 2 +- selfdrive/modeld/runners/__init__.py | 10 +- selfdrive/modeld/runners/thneedmodel.cc | 58 ----- selfdrive/modeld/runners/thneedmodel.h | 17 -- selfdrive/modeld/runners/thneedmodel.pxd | 9 - selfdrive/modeld/runners/thneedmodel_pyx.pyx | 14 - selfdrive/modeld/runners/tinygradmodel.py | 91 +++++++ selfdrive/modeld/thneed/README | 8 - selfdrive/modeld/thneed/__init__.py | 0 selfdrive/modeld/thneed/serialize.cc | 154 ----------- selfdrive/modeld/thneed/thneed.h | 133 ---------- selfdrive/modeld/thneed/thneed_common.cc | 216 ---------------- selfdrive/modeld/thneed/thneed_pc.cc | 32 --- selfdrive/modeld/thneed/thneed_qcom2.cc | 258 ------------------- 14 files changed, 97 insertions(+), 905 deletions(-) delete mode 100644 selfdrive/modeld/runners/thneedmodel.cc delete mode 100644 selfdrive/modeld/runners/thneedmodel.h delete mode 100644 selfdrive/modeld/runners/thneedmodel.pxd delete mode 100644 selfdrive/modeld/runners/thneedmodel_pyx.pyx create mode 100644 selfdrive/modeld/runners/tinygradmodel.py delete mode 100644 selfdrive/modeld/thneed/README delete mode 100644 selfdrive/modeld/thneed/__init__.py delete mode 100644 selfdrive/modeld/thneed/serialize.cc delete mode 100644 selfdrive/modeld/thneed/thneed.h delete mode 100644 selfdrive/modeld/thneed/thneed_common.cc delete mode 100644 selfdrive/modeld/thneed/thneed_pc.cc delete mode 100644 selfdrive/modeld/thneed/thneed_qcom2.cc diff --git a/selfdrive/modeld/modeld.py b/selfdrive/modeld/modeld.py index 86813fcb8ecd66..97258384745d07 100755 --- a/selfdrive/modeld/modeld.py +++ b/selfdrive/modeld/modeld.py @@ -29,7 +29,7 @@ SEND_RAW_PRED = os.getenv('SEND_RAW_PRED') MODEL_PATHS = { - ModelRunner.THNEED: Path(__file__).parent / 'models/supercombo.thneed', + ModelRunner.TINYGRAD: Path(__file__).parent / 'models/supercombo.onnx', ModelRunner.ONNX: Path(__file__).parent / 'models/supercombo.onnx'} METADATA_PATH = Path(__file__).parent / 'models/supercombo_metadata.pkl' diff --git a/selfdrive/modeld/runners/__init__.py b/selfdrive/modeld/runners/__init__.py index 4c29bf3f1cfbf7..afcb97df5203c9 100644 --- a/selfdrive/modeld/runners/__init__.py +++ b/selfdrive/modeld/runners/__init__.py @@ -3,18 +3,18 @@ from openpilot.selfdrive.modeld.runners.runmodel_pyx import RunModel, Runtime assert Runtime -USE_THNEED = int(os.getenv('USE_THNEED', str(int(TICI)))) +USE_TINYGRAD = int(os.getenv('USE_TINYGRAD', str(int(TICI)))) USE_SNPE = int(os.getenv('USE_SNPE', str(int(TICI)))) class ModelRunner(RunModel): - THNEED = 'THNEED' + TINYGRAD = 'TINYGRAD' 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 + if ModelRunner.TINYGRAD in paths and USE_TINYGRAD: + from openpilot.selfdrive.modeld.runners.tinygradmodel import TinygradModel as Runner + runner_type = ModelRunner.TINYGRAD elif ModelRunner.SNPE in paths and USE_SNPE: from openpilot.selfdrive.modeld.runners.snpemodel_pyx import SNPEModel as Runner runner_type = ModelRunner.SNPE diff --git a/selfdrive/modeld/runners/thneedmodel.cc b/selfdrive/modeld/runners/thneedmodel.cc deleted file mode 100644 index a16d8b42aab223..00000000000000 --- a/selfdrive/modeld/runners/thneedmodel.cc +++ /dev/null @@ -1,58 +0,0 @@ -#include "selfdrive/modeld/runners/thneedmodel.h" - -#include - -#include "common/swaglog.h" - -ThneedModel::ThneedModel(const std::string path, float *_output, size_t _output_size, int runtime, bool luse_tf8, cl_context context) { - thneed = new Thneed(true, context); - thneed->load(path.c_str()); - thneed->clexec(); - - recorded = false; - output = _output; -} - -void* ThneedModel::getCLBuffer(const std::string name) { - int index = -1; - for (int i = 0; i < inputs.size(); i++) { - if (name == inputs[i]->name) { - index = i; - break; - } - } - - if (index == -1) { - LOGE("Tried to get CL buffer for input `%s` but no input with this name exists", name.c_str()); - assert(false); - } - - if (thneed->input_clmem.size() >= inputs.size()) { - return &thneed->input_clmem[inputs.size() - index - 1]; - } else { - return nullptr; - } -} - -void ThneedModel::execute() { - if (!recorded) { - thneed->record = true; - float *input_buffers[inputs.size()]; - for (int i = 0; i < inputs.size(); i++) { - input_buffers[inputs.size() - i - 1] = inputs[i]->buffer; - } - - thneed->copy_inputs(input_buffers); - thneed->clexec(); - thneed->copy_output(output); - thneed->stop(); - - recorded = true; - } else { - float *input_buffers[inputs.size()]; - for (int i = 0; i < inputs.size(); i++) { - input_buffers[inputs.size() - i - 1] = inputs[i]->buffer; - } - thneed->execute(input_buffers, output); - } -} diff --git a/selfdrive/modeld/runners/thneedmodel.h b/selfdrive/modeld/runners/thneedmodel.h deleted file mode 100644 index 6ed479c081634d..00000000000000 --- a/selfdrive/modeld/runners/thneedmodel.h +++ /dev/null @@ -1,17 +0,0 @@ -#pragma once - -#include - -#include "selfdrive/modeld/runners/runmodel.h" -#include "selfdrive/modeld/thneed/thneed.h" - -class ThneedModel : public RunModel { -public: - ThneedModel(const std::string path, float *_output, size_t _output_size, int runtime, bool use_tf8 = false, cl_context context = NULL); - void *getCLBuffer(const std::string name); - void execute(); -private: - Thneed *thneed = NULL; - bool recorded; - float *output; -}; diff --git a/selfdrive/modeld/runners/thneedmodel.pxd b/selfdrive/modeld/runners/thneedmodel.pxd deleted file mode 100644 index 79e24dbdd62518..00000000000000 --- a/selfdrive/modeld/runners/thneedmodel.pxd +++ /dev/null @@ -1,9 +0,0 @@ -# distutils: language = c++ - -from libcpp.string cimport string - -from msgq.visionipc.visionipc cimport cl_context - -cdef extern from "selfdrive/modeld/runners/thneedmodel.h": - cdef cppclass ThneedModel: - ThneedModel(string, float*, size_t, int, bool, cl_context) diff --git a/selfdrive/modeld/runners/thneedmodel_pyx.pyx b/selfdrive/modeld/runners/thneedmodel_pyx.pyx deleted file mode 100644 index 6f8fdd255fa5bb..00000000000000 --- a/selfdrive/modeld/runners/thneedmodel_pyx.pyx +++ /dev/null @@ -1,14 +0,0 @@ -# distutils: language = c++ -# cython: c_string_encoding=ascii, language_level=3 - -from libcpp cimport bool -from libcpp.string cimport string - -from .thneedmodel cimport ThneedModel as cppThneedModel -from selfdrive.modeld.models.commonmodel_pyx cimport CLContext -from selfdrive.modeld.runners.runmodel_pyx cimport RunModel -from selfdrive.modeld.runners.runmodel cimport RunModel as cppRunModel - -cdef class ThneedModel(RunModel): - def __cinit__(self, string path, float[:] output, int runtime, bool use_tf8, CLContext context): - self.model = new cppThneedModel(path, &output[0], len(output), runtime, use_tf8, context.context) diff --git a/selfdrive/modeld/runners/tinygradmodel.py b/selfdrive/modeld/runners/tinygradmodel.py new file mode 100644 index 00000000000000..ef00ade211f279 --- /dev/null +++ b/selfdrive/modeld/runners/tinygradmodel.py @@ -0,0 +1,91 @@ +from pathlib import Path +import sys +sys.path.append(str(Path(__file__).parent.parent.parent.parent / 'tinygrad_repo')) + +import onnx +import itertools +import numpy as np + +from tinygrad.jit import TinyJit +from tinygrad.tensor import Tensor +from tinygrad.helpers import dtypes +from extra.onnx import get_run_onnx +from openpilot.selfdrive.modeld.runners.runmodel_pyx import RunModel + +ONNX_TYPES_TO_NP_TYPES: dict[int, np.dtype] = { + i: onnx.helper.tensor_dtype_to_np_dtype(i) + for dtype, i in onnx.TensorProto.DataType.items() + if dtype in ['FLOAT', 'FLOAT16', 'INT64', 'INT32', 'UINT8'] +} + +@TinyJit +def model_exec(run_onnx, **inputs): + return next(iter(run_onnx(inputs).values())).cast(dtypes.float32).realize() + + +class TinygradModel(RunModel): + def __init__(self, path, output, runtime, use_tf8, cl_context): + self.inputs = {} + self.output = output + self.use_tf8 = use_tf8 + + Tensor.manual_seed(1337) + Tensor.no_grad = True + + onnx_model = onnx.load(path) + self.run_onnx = get_run_onnx(onnx_model) + self.input_shapes = {inp.name:tuple(x.dim_value for x in inp.type.tensor_type.shape.dim) for inp in onnx_model.graph.input} + self.input_dtypes = {inp.name: ONNX_TYPES_TO_NP_TYPES[inp.type.tensor_type.elem_type] for inp in onnx_model.graph.input} + + def addInput(self, name, buffer): + assert name in self.input_shapes + 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(np.uint8) / 255. if self.use_tf8 and k == 'input_img' else v) 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()} + inputs = {k: Tensor(v) for k,v in inputs.items()} + outputs = model_exec(self.run_onnx, **inputs) + self.output[:] = outputs.numpy() + return self.output + + +if __name__ == "__main__": + import pickle + from tqdm import trange + from openpilot.selfdrive.modeld.runners import Runtime + from openpilot.selfdrive.modeld.constants import ModelConstants + from openpilot.selfdrive.modeld.models.commonmodel_pyx import CLContext + + MODEL_PATH = Path(__file__).parent.parent / 'models/supercombo.onnx' + METADATA_PATH = Path(__file__).parent.parent / 'models/supercombo_metadata.pkl' + with open(METADATA_PATH, 'rb') as f: + model_metadata = pickle.load(f) + + net_output_size = model_metadata['output_shapes']['outputs'][1] + output = np.zeros(net_output_size, dtype=np.float32) + + model = TinygradModel(MODEL_PATH, output, Runtime.CPU, False, CLContext()) + + inputs = { + 'input_imgs': np.zeros(128 * 256 * 12, dtype=np.uint8), + 'big_input_imgs': np.zeros(128 * 256 * 12, dtype=np.uint8), + '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), + 'lateral_control_params': np.zeros(ModelConstants.LATERAL_CONTROL_PARAMS_LEN, dtype=np.float32), + 'prev_desired_curv': np.zeros(ModelConstants.PREV_DESIRED_CURV_LEN * (ModelConstants.HISTORY_BUFFER_LEN+1), dtype=np.float32), + 'features_buffer': np.zeros(ModelConstants.HISTORY_BUFFER_LEN * ModelConstants.FEATURE_LEN, dtype=np.float32), + } + + for k,v in inputs.items(): + model.addInput(k, v) + + for _ in trange(100): + model.execute() diff --git a/selfdrive/modeld/thneed/README b/selfdrive/modeld/thneed/README deleted file mode 100644 index f3bc66d8fc26ff..00000000000000 --- a/selfdrive/modeld/thneed/README +++ /dev/null @@ -1,8 +0,0 @@ -thneed is an SNPE accelerator. I know SNPE is already an accelerator, but sometimes things need to go even faster.. - -It runs on the local device, and caches a single model run. Then it replays it, but fast. - -thneed slices through abstraction layers like a fish. - -You need a thneed. - diff --git a/selfdrive/modeld/thneed/__init__.py b/selfdrive/modeld/thneed/__init__.py deleted file mode 100644 index e69de29bb2d1d6..00000000000000 diff --git a/selfdrive/modeld/thneed/serialize.cc b/selfdrive/modeld/thneed/serialize.cc deleted file mode 100644 index 3dc2bef41448f8..00000000000000 --- a/selfdrive/modeld/thneed/serialize.cc +++ /dev/null @@ -1,154 +0,0 @@ -#include -#include - -#include "third_party/json11/json11.hpp" -#include "common/util.h" -#include "common/clutil.h" -#include "common/swaglog.h" -#include "selfdrive/modeld/thneed/thneed.h" -using namespace json11; - -extern map g_program_source; - -void Thneed::load(const char *filename) { - LOGD("Thneed::load: loading from %s\n", filename); - - string buf = util::read_file(filename); - int jsz = *(int *)buf.data(); - string jsonerr; - string jj(buf.data() + sizeof(int), jsz); - Json jdat = Json::parse(jj, jsonerr); - - map real_mem; - real_mem[NULL] = NULL; - - int ptr = sizeof(int)+jsz; - for (auto &obj : jdat["objects"].array_items()) { - auto mobj = obj.object_items(); - int sz = mobj["size"].int_value(); - cl_mem clbuf = NULL; - - if (mobj["buffer_id"].string_value().size() > 0) { - // image buffer must already be allocated - clbuf = real_mem[*(cl_mem*)(mobj["buffer_id"].string_value().data())]; - assert(mobj["needs_load"].bool_value() == false); - } else { - if (mobj["needs_load"].bool_value()) { - clbuf = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR | CL_MEM_READ_WRITE, sz, &buf[ptr], NULL); - if (debug >= 1) printf("loading %p %d @ 0x%X\n", clbuf, sz, ptr); - ptr += sz; - } else { - // TODO: is there a faster way to init zeroed out buffers? - void *host_zeros = calloc(sz, 1); - clbuf = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR | CL_MEM_READ_WRITE, sz, host_zeros, NULL); - free(host_zeros); - } - } - assert(clbuf != NULL); - - if (mobj["arg_type"] == "image2d_t" || mobj["arg_type"] == "image1d_t") { - cl_image_desc desc = {0}; - desc.image_type = (mobj["arg_type"] == "image2d_t") ? CL_MEM_OBJECT_IMAGE2D : CL_MEM_OBJECT_IMAGE1D_BUFFER; - desc.image_width = mobj["width"].int_value(); - desc.image_height = mobj["height"].int_value(); - desc.image_row_pitch = mobj["row_pitch"].int_value(); - assert(sz == desc.image_height*desc.image_row_pitch); -#ifdef QCOM2 - desc.buffer = clbuf; -#else - // TODO: we are creating unused buffers on PC - clReleaseMemObject(clbuf); -#endif - cl_image_format format = {0}; - format.image_channel_order = CL_RGBA; - format.image_channel_data_type = mobj["float32"].bool_value() ? CL_FLOAT : CL_HALF_FLOAT; - - cl_int errcode; - -#ifndef QCOM2 - if (mobj["needs_load"].bool_value()) { - clbuf = clCreateImage(context, CL_MEM_COPY_HOST_PTR | CL_MEM_READ_WRITE, &format, &desc, &buf[ptr-sz], &errcode); - } else { - clbuf = clCreateImage(context, CL_MEM_READ_WRITE, &format, &desc, NULL, &errcode); - } -#else - clbuf = clCreateImage(context, CL_MEM_READ_WRITE, &format, &desc, NULL, &errcode); -#endif - if (clbuf == NULL) { - LOGE("clError: %s create image %zux%zu rp %zu with buffer %p\n", cl_get_error_string(errcode), - desc.image_width, desc.image_height, desc.image_row_pitch, desc.buffer); - } - assert(clbuf != NULL); - } - - real_mem[*(cl_mem*)(mobj["id"].string_value().data())] = clbuf; - } - - map g_programs; - for (const auto &[name, source] : jdat["programs"].object_items()) { - if (debug >= 1) printf("building %s with size %zu\n", name.c_str(), source.string_value().size()); - g_programs[name] = cl_program_from_source(context, device_id, source.string_value()); - } - - for (auto &obj : jdat["inputs"].array_items()) { - auto mobj = obj.object_items(); - int sz = mobj["size"].int_value(); - cl_mem aa = real_mem[*(cl_mem*)(mobj["buffer_id"].string_value().data())]; - input_clmem.push_back(aa); - input_sizes.push_back(sz); - LOGD("Thneed::load: adding input %s with size %d\n", mobj["name"].string_value().data(), sz); - - cl_int cl_err; - void *ret = clEnqueueMapBuffer(command_queue, aa, CL_TRUE, CL_MAP_WRITE, 0, sz, 0, NULL, NULL, &cl_err); - if (cl_err != CL_SUCCESS) LOGE("clError: %s map %p %d\n", cl_get_error_string(cl_err), aa, sz); - assert(cl_err == CL_SUCCESS); - inputs.push_back(ret); - } - - for (auto &obj : jdat["outputs"].array_items()) { - auto mobj = obj.object_items(); - int sz = mobj["size"].int_value(); - LOGD("Thneed::save: adding output with size %d\n", sz); - // TODO: support multiple outputs - output = real_mem[*(cl_mem*)(mobj["buffer_id"].string_value().data())]; - assert(output != NULL); - } - - for (auto &obj : jdat["binaries"].array_items()) { - string name = obj["name"].string_value(); - size_t length = obj["length"].int_value(); - if (debug >= 1) printf("binary %s with size %zu\n", name.c_str(), length); - g_programs[name] = cl_program_from_binary(context, device_id, (const uint8_t*)&buf[ptr], length); - ptr += length; - } - - for (auto &obj : jdat["kernels"].array_items()) { - auto gws = obj["global_work_size"]; - auto lws = obj["local_work_size"]; - auto kk = shared_ptr(new CLQueuedKernel(this)); - - kk->name = obj["name"].string_value(); - kk->program = g_programs[kk->name]; - kk->work_dim = obj["work_dim"].int_value(); - for (int i = 0; i < kk->work_dim; i++) { - kk->global_work_size[i] = gws[i].int_value(); - kk->local_work_size[i] = lws[i].int_value(); - } - kk->num_args = obj["num_args"].int_value(); - for (int i = 0; i < kk->num_args; i++) { - string arg = obj["args"].array_items()[i].string_value(); - int arg_size = obj["args_size"].array_items()[i].int_value(); - kk->args_size.push_back(arg_size); - if (arg_size == 8) { - cl_mem val = *(cl_mem*)(arg.data()); - val = real_mem[val]; - kk->args.push_back(string((char*)&val, sizeof(val))); - } else { - kk->args.push_back(arg); - } - } - kq.push_back(kk); - } - - clFinish(command_queue); -} diff --git a/selfdrive/modeld/thneed/thneed.h b/selfdrive/modeld/thneed/thneed.h deleted file mode 100644 index 47e18e0be3bb30..00000000000000 --- a/selfdrive/modeld/thneed/thneed.h +++ /dev/null @@ -1,133 +0,0 @@ -#pragma once - -#ifndef __user -#define __user __attribute__(()) -#endif - -#include -#include -#include -#include -#include - -#include - -#include "third_party/linux/include/msm_kgsl.h" - -using namespace std; - -cl_int thneed_clSetKernelArg(cl_kernel kernel, cl_uint arg_index, size_t arg_size, const void *arg_value); - -namespace json11 { - class Json; -} -class Thneed; - -class GPUMalloc { - public: - GPUMalloc(int size, int fd); - ~GPUMalloc(); - void *alloc(int size); - private: - uint64_t base; - int remaining; -}; - -class CLQueuedKernel { - public: - CLQueuedKernel(Thneed *lthneed) { thneed = lthneed; } - CLQueuedKernel(Thneed *lthneed, - cl_kernel _kernel, - cl_uint _work_dim, - const size_t *_global_work_size, - const size_t *_local_work_size); - cl_int exec(); - void debug_print(bool verbose); - int get_arg_num(const char *search_arg_name); - cl_program program; - string name; - cl_uint num_args; - vector arg_names; - vector arg_types; - vector args; - vector args_size; - cl_kernel kernel = NULL; - json11::Json to_json() const; - - cl_uint work_dim; - size_t global_work_size[3] = {0}; - size_t local_work_size[3] = {0}; - private: - Thneed *thneed; -}; - -class CachedIoctl { - public: - virtual void exec() {} -}; - -class CachedSync: public CachedIoctl { - public: - CachedSync(Thneed *lthneed, string ldata) { thneed = lthneed; data = ldata; } - void exec(); - private: - Thneed *thneed; - string data; -}; - -class CachedCommand: public CachedIoctl { - public: - CachedCommand(Thneed *lthneed, struct kgsl_gpu_command *cmd); - void exec(); - private: - void disassemble(int cmd_index); - struct kgsl_gpu_command cache; - unique_ptr cmds; - unique_ptr objs; - Thneed *thneed; - vector > kq; -}; - -class Thneed { - public: - Thneed(bool do_clinit=false, cl_context _context = NULL); - void stop(); - void execute(float **finputs, float *foutput, bool slow=false); - void wait(); - - vector input_clmem; - vector inputs; - vector input_sizes; - cl_mem output = NULL; - - cl_context context = NULL; - cl_command_queue command_queue; - cl_device_id device_id; - int context_id; - - // protected? - bool record = false; - int debug; - int timestamp; - -#ifdef QCOM2 - unique_ptr ram; - vector > cmds; - int fd; -#endif - - // all CL kernels - void copy_inputs(float **finputs, bool internal=false); - void copy_output(float *foutput); - cl_int clexec(); - vector > kq; - - // pending CL kernels - vector > ckq; - - // loading - void load(const char *filename); - private: - void clinit(); -}; - diff --git a/selfdrive/modeld/thneed/thneed_common.cc b/selfdrive/modeld/thneed/thneed_common.cc deleted file mode 100644 index ecdf1237e384ff..00000000000000 --- a/selfdrive/modeld/thneed/thneed_common.cc +++ /dev/null @@ -1,216 +0,0 @@ -#include "selfdrive/modeld/thneed/thneed.h" - -#include -#include -#include - -#include "common/clutil.h" -#include "common/timing.h" - -map, string> g_args; -map, int> g_args_size; -map g_program_source; - -void Thneed::stop() { - //printf("Thneed::stop: recorded %lu commands\n", cmds.size()); - record = false; -} - -void Thneed::clinit() { - device_id = cl_get_device_id(CL_DEVICE_TYPE_DEFAULT); - if (context == NULL) context = CL_CHECK_ERR(clCreateContext(NULL, 1, &device_id, NULL, NULL, &err)); - //cl_command_queue_properties props[3] = {CL_QUEUE_PROPERTIES, CL_QUEUE_PROFILING_ENABLE, 0}; - cl_command_queue_properties props[3] = {CL_QUEUE_PROPERTIES, 0, 0}; - command_queue = CL_CHECK_ERR(clCreateCommandQueueWithProperties(context, device_id, props, &err)); - printf("Thneed::clinit done\n"); -} - -cl_int Thneed::clexec() { - if (debug >= 1) printf("Thneed::clexec: running %lu queued kernels\n", kq.size()); - for (auto &k : kq) { - if (record) ckq.push_back(k); - cl_int ret = k->exec(); - assert(ret == CL_SUCCESS); - } - return clFinish(command_queue); -} - -void Thneed::copy_inputs(float **finputs, bool internal) { - for (int idx = 0; idx < inputs.size(); ++idx) { - if (debug >= 1) printf("copying %lu -- %p -> %p (cl %p)\n", input_sizes[idx], finputs[idx], inputs[idx], input_clmem[idx]); - - if (internal) { - // if it's internal, using memcpy is fine since the buffer sync is cached in the ioctl layer - if (finputs[idx] != NULL) memcpy(inputs[idx], finputs[idx], input_sizes[idx]); - } else { - if (finputs[idx] != NULL) CL_CHECK(clEnqueueWriteBuffer(command_queue, input_clmem[idx], CL_TRUE, 0, input_sizes[idx], finputs[idx], 0, NULL, NULL)); - } - } -} - -void Thneed::copy_output(float *foutput) { - if (output != NULL) { - size_t sz; - clGetMemObjectInfo(output, CL_MEM_SIZE, sizeof(sz), &sz, NULL); - if (debug >= 1) printf("copying %lu for output %p -> %p\n", sz, output, foutput); - CL_CHECK(clEnqueueReadBuffer(command_queue, output, CL_TRUE, 0, sz, foutput, 0, NULL, NULL)); - } else { - printf("CAUTION: model output is NULL, does it have no outputs?\n"); - } -} - -// *********** CLQueuedKernel *********** - -CLQueuedKernel::CLQueuedKernel(Thneed *lthneed, - cl_kernel _kernel, - cl_uint _work_dim, - const size_t *_global_work_size, - const size_t *_local_work_size) { - thneed = lthneed; - kernel = _kernel; - work_dim = _work_dim; - assert(work_dim <= 3); - for (int i = 0; i < work_dim; i++) { - global_work_size[i] = _global_work_size[i]; - local_work_size[i] = _local_work_size[i]; - } - - char _name[0x100]; - clGetKernelInfo(kernel, CL_KERNEL_FUNCTION_NAME, sizeof(_name), _name, NULL); - name = string(_name); - clGetKernelInfo(kernel, CL_KERNEL_NUM_ARGS, sizeof(num_args), &num_args, NULL); - - // get args - for (int i = 0; i < num_args; i++) { - char arg_name[0x100] = {0}; - clGetKernelArgInfo(kernel, i, CL_KERNEL_ARG_NAME, sizeof(arg_name), arg_name, NULL); - arg_names.push_back(string(arg_name)); - clGetKernelArgInfo(kernel, i, CL_KERNEL_ARG_TYPE_NAME, sizeof(arg_name), arg_name, NULL); - arg_types.push_back(string(arg_name)); - - args.push_back(g_args[make_pair(kernel, i)]); - args_size.push_back(g_args_size[make_pair(kernel, i)]); - } - - // get program - clGetKernelInfo(kernel, CL_KERNEL_PROGRAM, sizeof(program), &program, NULL); -} - -int CLQueuedKernel::get_arg_num(const char *search_arg_name) { - for (int i = 0; i < num_args; i++) { - if (arg_names[i] == search_arg_name) return i; - } - printf("failed to find %s in %s\n", search_arg_name, name.c_str()); - assert(false); -} - -cl_int CLQueuedKernel::exec() { - if (kernel == NULL) { - kernel = clCreateKernel(program, name.c_str(), NULL); - arg_names.clear(); - arg_types.clear(); - - for (int j = 0; j < num_args; j++) { - char arg_name[0x100] = {0}; - clGetKernelArgInfo(kernel, j, CL_KERNEL_ARG_NAME, sizeof(arg_name), arg_name, NULL); - arg_names.push_back(string(arg_name)); - clGetKernelArgInfo(kernel, j, CL_KERNEL_ARG_TYPE_NAME, sizeof(arg_name), arg_name, NULL); - arg_types.push_back(string(arg_name)); - - cl_int ret; - if (args[j].size() != 0) { - assert(args[j].size() == args_size[j]); - ret = thneed_clSetKernelArg(kernel, j, args[j].size(), args[j].data()); - } else { - ret = thneed_clSetKernelArg(kernel, j, args_size[j], NULL); - } - assert(ret == CL_SUCCESS); - } - } - - if (thneed->debug >= 1) { - debug_print(thneed->debug >= 2); - } - - return clEnqueueNDRangeKernel(thneed->command_queue, - kernel, work_dim, NULL, global_work_size, local_work_size, 0, NULL, NULL); -} - -void CLQueuedKernel::debug_print(bool verbose) { - printf("%p %56s -- ", kernel, name.c_str()); - for (int i = 0; i < work_dim; i++) { - printf("%4zu ", global_work_size[i]); - } - printf(" -- "); - for (int i = 0; i < work_dim; i++) { - printf("%4zu ", local_work_size[i]); - } - printf("\n"); - - if (verbose) { - for (int i = 0; i < num_args; i++) { - string arg = args[i]; - printf(" %s %s", arg_types[i].c_str(), arg_names[i].c_str()); - void *arg_value = (void*)arg.data(); - int arg_size = arg.size(); - if (arg_size == 0) { - printf(" (size) %d", args_size[i]); - } else if (arg_size == 1) { - printf(" = %d", *((char*)arg_value)); - } else if (arg_size == 2) { - printf(" = %d", *((short*)arg_value)); - } else if (arg_size == 4) { - if (arg_types[i] == "float") { - printf(" = %f", *((float*)arg_value)); - } else { - printf(" = %d", *((int*)arg_value)); - } - } else if (arg_size == 8) { - cl_mem val = (cl_mem)(*((uintptr_t*)arg_value)); - printf(" = %p", val); - if (val != NULL) { - cl_mem_object_type obj_type; - clGetMemObjectInfo(val, CL_MEM_TYPE, sizeof(obj_type), &obj_type, NULL); - if (arg_types[i] == "image2d_t" || arg_types[i] == "image1d_t" || obj_type == CL_MEM_OBJECT_IMAGE2D) { - cl_image_format format; - size_t width, height, depth, array_size, row_pitch, slice_pitch; - cl_mem buf; - clGetImageInfo(val, CL_IMAGE_FORMAT, sizeof(format), &format, NULL); - assert(format.image_channel_order == CL_RGBA); - assert(format.image_channel_data_type == CL_HALF_FLOAT || format.image_channel_data_type == CL_FLOAT); - clGetImageInfo(val, CL_IMAGE_WIDTH, sizeof(width), &width, NULL); - clGetImageInfo(val, CL_IMAGE_HEIGHT, sizeof(height), &height, NULL); - clGetImageInfo(val, CL_IMAGE_ROW_PITCH, sizeof(row_pitch), &row_pitch, NULL); - clGetImageInfo(val, CL_IMAGE_DEPTH, sizeof(depth), &depth, NULL); - clGetImageInfo(val, CL_IMAGE_ARRAY_SIZE, sizeof(array_size), &array_size, NULL); - clGetImageInfo(val, CL_IMAGE_SLICE_PITCH, sizeof(slice_pitch), &slice_pitch, NULL); - assert(depth == 0); - assert(array_size == 0); - assert(slice_pitch == 0); - - clGetImageInfo(val, CL_IMAGE_BUFFER, sizeof(buf), &buf, NULL); - size_t sz = 0; - if (buf != NULL) clGetMemObjectInfo(buf, CL_MEM_SIZE, sizeof(sz), &sz, NULL); - printf(" image %zu x %zu rp %zu @ %p buffer %zu", width, height, row_pitch, buf, sz); - } else { - size_t sz; - clGetMemObjectInfo(val, CL_MEM_SIZE, sizeof(sz), &sz, NULL); - printf(" buffer %zu", sz); - } - } - } - printf("\n"); - } - } -} - -cl_int thneed_clSetKernelArg(cl_kernel kernel, cl_uint arg_index, size_t arg_size, const void *arg_value) { - g_args_size[make_pair(kernel, arg_index)] = arg_size; - if (arg_value != NULL) { - g_args[make_pair(kernel, arg_index)] = string((char*)arg_value, arg_size); - } else { - g_args[make_pair(kernel, arg_index)] = string(""); - } - cl_int ret = clSetKernelArg(kernel, arg_index, arg_size, arg_value); - return ret; -} diff --git a/selfdrive/modeld/thneed/thneed_pc.cc b/selfdrive/modeld/thneed/thneed_pc.cc deleted file mode 100644 index 8d0037628e2f3d..00000000000000 --- a/selfdrive/modeld/thneed/thneed_pc.cc +++ /dev/null @@ -1,32 +0,0 @@ -#include "selfdrive/modeld/thneed/thneed.h" - -#include - -#include "common/clutil.h" -#include "common/timing.h" - -Thneed::Thneed(bool do_clinit, cl_context _context) { - context = _context; - if (do_clinit) clinit(); - char *thneed_debug_env = getenv("THNEED_DEBUG"); - debug = (thneed_debug_env != NULL) ? atoi(thneed_debug_env) : 0; -} - -void Thneed::execute(float **finputs, float *foutput, bool slow) { - uint64_t tb, te; - if (debug >= 1) tb = nanos_since_boot(); - - // ****** copy inputs - copy_inputs(finputs); - - // ****** run commands - clexec(); - - // ****** copy outputs - copy_output(foutput); - - if (debug >= 1) { - te = nanos_since_boot(); - printf("model exec in %lu us\n", (te-tb)/1000); - } -} diff --git a/selfdrive/modeld/thneed/thneed_qcom2.cc b/selfdrive/modeld/thneed/thneed_qcom2.cc deleted file mode 100644 index 21de15d17c9cfc..00000000000000 --- a/selfdrive/modeld/thneed/thneed_qcom2.cc +++ /dev/null @@ -1,258 +0,0 @@ -#include "selfdrive/modeld/thneed/thneed.h" - -#include -#include - -#include -#include -#include -#include -#include - -#include "common/clutil.h" -#include "common/timing.h" - -Thneed *g_thneed = NULL; -int g_fd = -1; - -void hexdump(uint8_t *d, int len) { - assert((len%4) == 0); - printf(" dumping %p len 0x%x\n", d, len); - for (int i = 0; i < len/4; i++) { - if (i != 0 && (i%0x10) == 0) printf("\n"); - printf("%8x ", d[i]); - } - printf("\n"); -} - -// *********** ioctl interceptor *********** - -extern "C" { - -int (*my_ioctl)(int filedes, unsigned long request, void *argp) = NULL; -#undef ioctl -int ioctl(int filedes, unsigned long request, void *argp) { - request &= 0xFFFFFFFF; // needed on QCOM2 - if (my_ioctl == NULL) my_ioctl = reinterpret_cast(dlsym(RTLD_NEXT, "ioctl")); - Thneed *thneed = g_thneed; - - // save the fd - if (request == IOCTL_KGSL_GPUOBJ_ALLOC) g_fd = filedes; - - // note that this runs always, even without a thneed object - if (request == IOCTL_KGSL_DRAWCTXT_CREATE) { - struct kgsl_drawctxt_create *create = (struct kgsl_drawctxt_create *)argp; - create->flags &= ~KGSL_CONTEXT_PRIORITY_MASK; - create->flags |= 6 << KGSL_CONTEXT_PRIORITY_SHIFT; // priority from 1-15, 1 is max priority - printf("IOCTL_KGSL_DRAWCTXT_CREATE: creating context with flags 0x%x\n", create->flags); - } - - if (thneed != NULL) { - if (request == IOCTL_KGSL_GPU_COMMAND) { - struct kgsl_gpu_command *cmd = (struct kgsl_gpu_command *)argp; - if (thneed->record) { - thneed->timestamp = cmd->timestamp; - thneed->context_id = cmd->context_id; - thneed->cmds.push_back(unique_ptr(new CachedCommand(thneed, cmd))); - } - if (thneed->debug >= 1) { - printf("IOCTL_KGSL_GPU_COMMAND(%2zu): flags: 0x%lx context_id: %u timestamp: %u numcmds: %d numobjs: %d\n", - thneed->cmds.size(), - cmd->flags, - cmd->context_id, cmd->timestamp, cmd->numcmds, cmd->numobjs); - } - } else if (request == IOCTL_KGSL_GPUOBJ_SYNC) { - struct kgsl_gpuobj_sync *cmd = (struct kgsl_gpuobj_sync *)argp; - struct kgsl_gpuobj_sync_obj *objs = (struct kgsl_gpuobj_sync_obj *)(cmd->objs); - - if (thneed->debug >= 2) { - printf("IOCTL_KGSL_GPUOBJ_SYNC count:%d ", cmd->count); - for (int i = 0; i < cmd->count; i++) { - printf(" -- offset:0x%lx len:0x%lx id:%d op:%d ", objs[i].offset, objs[i].length, objs[i].id, objs[i].op); - } - printf("\n"); - } - - if (thneed->record) { - thneed->cmds.push_back(unique_ptr(new - CachedSync(thneed, string((char *)objs, sizeof(struct kgsl_gpuobj_sync_obj)*cmd->count)))); - } - } else if (request == IOCTL_KGSL_DEVICE_WAITTIMESTAMP_CTXTID) { - struct kgsl_device_waittimestamp_ctxtid *cmd = (struct kgsl_device_waittimestamp_ctxtid *)argp; - if (thneed->debug >= 1) { - printf("IOCTL_KGSL_DEVICE_WAITTIMESTAMP_CTXTID: context_id: %d timestamp: %d timeout: %d\n", - cmd->context_id, cmd->timestamp, cmd->timeout); - } - } else if (request == IOCTL_KGSL_SETPROPERTY) { - if (thneed->debug >= 1) { - struct kgsl_device_getproperty *prop = (struct kgsl_device_getproperty *)argp; - printf("IOCTL_KGSL_SETPROPERTY: 0x%x sizebytes:%zu\n", prop->type, prop->sizebytes); - if (thneed->debug >= 2) { - hexdump((uint8_t *)prop->value, prop->sizebytes); - if (prop->type == KGSL_PROP_PWR_CONSTRAINT) { - struct kgsl_device_constraint *constraint = (struct kgsl_device_constraint *)prop->value; - hexdump((uint8_t *)constraint->data, constraint->size); - } - } - } - } else if (request == IOCTL_KGSL_DRAWCTXT_CREATE || request == IOCTL_KGSL_DRAWCTXT_DESTROY) { - // this happens - } else if (request == IOCTL_KGSL_GPUOBJ_ALLOC || request == IOCTL_KGSL_GPUOBJ_FREE) { - // this happens - } else { - if (thneed->debug >= 1) { - printf("other ioctl %lx\n", request); - } - } - } - - int ret = my_ioctl(filedes, request, argp); - // NOTE: This error message goes into stdout and messes up pyenv - // if (ret != 0) printf("ioctl returned %d with errno %d\n", ret, errno); - return ret; -} - -} - -// *********** GPUMalloc *********** - -GPUMalloc::GPUMalloc(int size, int fd) { - struct kgsl_gpuobj_alloc alloc; - memset(&alloc, 0, sizeof(alloc)); - alloc.size = size; - alloc.flags = 0x10000a00; - ioctl(fd, IOCTL_KGSL_GPUOBJ_ALLOC, &alloc); - void *addr = mmap64(NULL, alloc.mmapsize, 0x3, 0x1, fd, alloc.id*0x1000); - assert(addr != MAP_FAILED); - - base = (uint64_t)addr; - remaining = size; -} - -GPUMalloc::~GPUMalloc() { - // TODO: free the GPU malloced area -} - -void *GPUMalloc::alloc(int size) { - void *ret = (void*)base; - size = (size+0xff) & (~0xFF); - assert(size <= remaining); - remaining -= size; - base += size; - return ret; -} - -// *********** CachedSync, at the ioctl layer *********** - -void CachedSync::exec() { - struct kgsl_gpuobj_sync cmd; - - cmd.objs = (uint64_t)data.data(); - cmd.obj_len = data.length(); - cmd.count = data.length() / sizeof(struct kgsl_gpuobj_sync_obj); - - int ret = ioctl(thneed->fd, IOCTL_KGSL_GPUOBJ_SYNC, &cmd); - assert(ret == 0); -} - -// *********** CachedCommand, at the ioctl layer *********** - -CachedCommand::CachedCommand(Thneed *lthneed, struct kgsl_gpu_command *cmd) { - thneed = lthneed; - assert(cmd->numsyncs == 0); - - memcpy(&cache, cmd, sizeof(cache)); - - if (cmd->numcmds > 0) { - cmds = make_unique(cmd->numcmds); - memcpy(cmds.get(), (void *)cmd->cmdlist, sizeof(struct kgsl_command_object)*cmd->numcmds); - cache.cmdlist = (uint64_t)cmds.get(); - for (int i = 0; i < cmd->numcmds; i++) { - void *nn = thneed->ram->alloc(cmds[i].size); - memcpy(nn, (void*)cmds[i].gpuaddr, cmds[i].size); - cmds[i].gpuaddr = (uint64_t)nn; - } - } - - if (cmd->numobjs > 0) { - objs = make_unique(cmd->numobjs); - memcpy(objs.get(), (void *)cmd->objlist, sizeof(struct kgsl_command_object)*cmd->numobjs); - cache.objlist = (uint64_t)objs.get(); - for (int i = 0; i < cmd->numobjs; i++) { - void *nn = thneed->ram->alloc(objs[i].size); - memset(nn, 0, objs[i].size); - objs[i].gpuaddr = (uint64_t)nn; - } - } - - kq = thneed->ckq; - thneed->ckq.clear(); -} - -void CachedCommand::exec() { - cache.timestamp = ++thneed->timestamp; - int ret = ioctl(thneed->fd, IOCTL_KGSL_GPU_COMMAND, &cache); - - if (thneed->debug >= 1) printf("CachedCommand::exec got %d\n", ret); - - if (thneed->debug >= 2) { - for (auto &it : kq) { - it->debug_print(false); - } - } - - assert(ret == 0); -} - -// *********** Thneed *********** - -Thneed::Thneed(bool do_clinit, cl_context _context) { - // TODO: QCOM2 actually requires a different context - //context = _context; - if (do_clinit) clinit(); - assert(g_fd != -1); - fd = g_fd; - ram = make_unique(0x80000, fd); - timestamp = -1; - g_thneed = this; - char *thneed_debug_env = getenv("THNEED_DEBUG"); - debug = (thneed_debug_env != NULL) ? atoi(thneed_debug_env) : 0; -} - -void Thneed::wait() { - struct kgsl_device_waittimestamp_ctxtid wait; - wait.context_id = context_id; - wait.timestamp = timestamp; - wait.timeout = -1; - - uint64_t tb = nanos_since_boot(); - int wret = ioctl(fd, IOCTL_KGSL_DEVICE_WAITTIMESTAMP_CTXTID, &wait); - uint64_t te = nanos_since_boot(); - - if (debug >= 1) printf("wait %d after %lu us\n", wret, (te-tb)/1000); -} - -void Thneed::execute(float **finputs, float *foutput, bool slow) { - uint64_t tb, te; - if (debug >= 1) tb = nanos_since_boot(); - - // ****** copy inputs - copy_inputs(finputs, true); - - // ****** run commands - int i = 0; - for (auto &it : cmds) { - ++i; - if (debug >= 1) printf("run %2d @ %7lu us: ", i, (nanos_since_boot()-tb)/1000); - it->exec(); - if ((i == cmds.size()) || slow) wait(); - } - - // ****** copy outputs - copy_output(foutput); - - if (debug >= 1) { - te = nanos_since_boot(); - printf("model exec in %lu us\n", (te-tb)/1000); - } -}