diff --git a/selfdrive/modeld/runners/thneedmodel.cc b/selfdrive/modeld/runners/thneedmodel.cc index a85b2ac02..edc091bda 100644 --- a/selfdrive/modeld/runners/thneedmodel.cc +++ b/selfdrive/modeld/runners/thneedmodel.cc @@ -47,7 +47,7 @@ void* ThneedModel::getExtraBuf() { void ThneedModel::execute() { if (!recorded) { - thneed->record = THNEED_RECORD; + thneed->record = true; if (use_extra) { float *inputs[5] = {recurrent, trafficConvention, desire, extra, input}; thneed->copy_inputs(inputs); diff --git a/selfdrive/modeld/thneed/compile.cc b/selfdrive/modeld/thneed/compile.cc index c2a357fcf..c22156d2c 100644 --- a/selfdrive/modeld/thneed/compile.cc +++ b/selfdrive/modeld/thneed/compile.cc @@ -38,7 +38,7 @@ int main(int argc, char* argv[]) { // test model auto thneed = new Thneed(true); - thneed->record &= ~THNEED_RECORD; + thneed->record = false; thneed->load(argv[2]); thneed->clexec(); thneed->find_inputs_outputs(); diff --git a/selfdrive/modeld/thneed/serialize.cc b/selfdrive/modeld/thneed/serialize.cc index cd5584553..89b761b9d 100644 --- a/selfdrive/modeld/thneed/serialize.cc +++ b/selfdrive/modeld/thneed/serialize.cc @@ -63,14 +63,14 @@ void Thneed::load(const char *filename) { map g_programs; for (const auto &[name, source] : jdat["programs"].object_items()) { - if (record & THNEED_DEBUG) printf("building %s with size %zu\n", name.c_str(), source.string_value().size()); + 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["binaries"].array_items()) { string name = obj["name"].string_value(); size_t length = obj["length"].int_value(); - if (record & THNEED_DEBUG) printf("binary %s with size %zu\n", name.c_str(), length); + 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; } diff --git a/selfdrive/modeld/thneed/thneed.cc b/selfdrive/modeld/thneed/thneed.cc index aa1caa5cd..e2dc9b72f 100644 --- a/selfdrive/modeld/thneed/thneed.cc +++ b/selfdrive/modeld/thneed/thneed.cc @@ -55,12 +55,12 @@ int ioctl(int filedes, unsigned long request, void *argp) { if (thneed != NULL) { if (request == IOCTL_KGSL_GPU_COMMAND) { struct kgsl_gpu_command *cmd = (struct kgsl_gpu_command *)argp; - if (thneed->record & THNEED_RECORD) { + 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->record & THNEED_DEBUG) { + 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, @@ -70,7 +70,7 @@ int ioctl(int filedes, unsigned long request, void *argp) { 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->record & THNEED_VERBOSE_DEBUG) { + 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); @@ -78,21 +78,21 @@ int ioctl(int filedes, unsigned long request, void *argp) { printf("\n"); } - if (thneed->record & THNEED_RECORD) { + 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->record & THNEED_DEBUG) { + 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->record & THNEED_DEBUG) { + 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->record & THNEED_VERBOSE_DEBUG) { + 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; @@ -105,7 +105,7 @@ int ioctl(int filedes, unsigned long request, void *argp) { } else if (request == IOCTL_KGSL_GPUOBJ_ALLOC || request == IOCTL_KGSL_GPUOBJ_FREE) { // this happens } else { - if (thneed->record & THNEED_DEBUG) { + if (thneed->debug >= 1) { printf("other ioctl %lx\n", request); } } @@ -197,9 +197,9 @@ void CachedCommand::exec() { cache.timestamp = ++thneed->timestamp; int ret = ioctl(thneed->fd, IOCTL_KGSL_GPU_COMMAND, &cache); - if (thneed->record & THNEED_DEBUG) printf("CachedCommand::exec got %d\n", ret); + if (thneed->debug >= 1) printf("CachedCommand::exec got %d\n", ret); - if (thneed->record & THNEED_VERBOSE_DEBUG) { + if (thneed->debug >= 2) { for (auto &it : kq) { it->debug_print(false); } @@ -220,15 +220,11 @@ Thneed::Thneed(bool do_clinit) { assert(g_fd != -1); fd = g_fd; ram = make_unique(0x80000, fd); - record = THNEED_RECORD; + record = true; timestamp = -1; g_thneed = this; char *thneed_debug_env = getenv("THNEED_DEBUG"); - if (thneed_debug_env != NULL) { - int thneed_debug_level = atoi(thneed_debug_env); - record |= (thneed_debug_level >= 1) ? THNEED_DEBUG : 0; - record |= (thneed_debug_level >= 2) ? THNEED_VERBOSE_DEBUG : 0; - } + debug = (thneed_debug_env != NULL) ? atoi(thneed_debug_env) : 0; } void Thneed::stop() { @@ -267,7 +263,7 @@ void Thneed::find_inputs_outputs() { void Thneed::copy_inputs(float **finputs) { //cl_int ret; for (int idx = 0; idx < inputs.size(); ++idx) { - if (record & THNEED_DEBUG) printf("copying %lu -- %p -> %p\n", input_sizes[idx], finputs[idx], inputs[idx]); + if (debug >= 1) printf("copying %lu -- %p -> %p\n", input_sizes[idx], finputs[idx], inputs[idx]); if (finputs[idx] != NULL) memcpy(inputs[idx], finputs[idx], input_sizes[idx]); } } @@ -276,7 +272,7 @@ void Thneed::copy_output(float *foutput) { if (output != NULL) { size_t sz; clGetMemObjectInfo(output, CL_MEM_SIZE, sizeof(sz), &sz, NULL); - if (record & THNEED_DEBUG) printf("copying %lu for output %p -> %p\n", sz, output, foutput); + if (debug >= 1) printf("copying %lu for output %p -> %p\n", sz, output, foutput); 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"); @@ -293,12 +289,12 @@ void Thneed::wait() { int wret = ioctl(fd, IOCTL_KGSL_DEVICE_WAITTIMESTAMP_CTXTID, &wait); uint64_t te = nanos_since_boot(); - if (record & THNEED_DEBUG) printf("wait %d after %lu us\n", wret, (te-tb)/1000); + 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 (record & THNEED_DEBUG) tb = nanos_since_boot(); + if (debug >= 1) tb = nanos_since_boot(); // ****** copy inputs copy_inputs(finputs); @@ -325,7 +321,7 @@ void Thneed::execute(float **finputs, float *foutput, bool slow) { int i = 0; for (auto &it : cmds) { ++i; - if (record & THNEED_DEBUG) printf("run %2d @ %7lu us: ", i, (nanos_since_boot()-tb)/1000); + if (debug >= 1) printf("run %2d @ %7lu us: ", i, (nanos_since_boot()-tb)/1000); it->exec(); if ((i == cmds.size()) || slow) wait(); } @@ -341,7 +337,7 @@ void Thneed::execute(float **finputs, float *foutput, bool slow) { ret = ioctl(fd, IOCTL_KGSL_SETPROPERTY, &prop); assert(ret == 0); - if (record & THNEED_DEBUG) { + if (debug >= 1) { te = nanos_since_boot(); printf("model exec in %lu us\n", (te-tb)/1000); } @@ -359,7 +355,7 @@ void Thneed::clinit() { cl_int Thneed::clexec() { printf("Thneed::clexec: running %lu queued kernels\n", kq.size()); for (auto &k : kq) { - if (record & THNEED_RECORD) ckq.push_back(k); + if (record) ckq.push_back(k); cl_int ret = k->exec(); assert(ret == CL_SUCCESS); } @@ -397,7 +393,7 @@ cl_int thneed_clEnqueueNDRangeKernel(cl_command_queue command_queue, assert(event_wait_list == NULL); cl_int ret = 0; - if (thneed != NULL && thneed->record & THNEED_RECORD) { + if (thneed != NULL && thneed->record) { if (thneed->context == NULL) { thneed->command_queue = command_queue; clGetKernelInfo(kernel, CL_KERNEL_CONTEXT, sizeof(thneed->context), &thneed->context, NULL); @@ -419,7 +415,7 @@ cl_int thneed_clEnqueueNDRangeKernel(cl_command_queue command_queue, cl_int thneed_clFinish(cl_command_queue command_queue) { Thneed *thneed = g_thneed; - if (thneed != NULL && thneed->record & THNEED_RECORD) { + if (thneed != NULL && thneed->record) { #ifdef RUN_OPTIMIZER thneed->optimize(); #endif @@ -526,8 +522,8 @@ cl_int CLQueuedKernel::exec() { } } - if (thneed->record & THNEED_DEBUG) { - debug_print(thneed->record & THNEED_VERBOSE_DEBUG); + if (thneed->debug >= 1) { + debug_print(thneed->debug >= 2); } return clEnqueueNDRangeKernel(thneed->command_queue, diff --git a/selfdrive/modeld/thneed/thneed.h b/selfdrive/modeld/thneed/thneed.h index 1197e4c5e..b09d32b0e 100644 --- a/selfdrive/modeld/thneed/thneed.h +++ b/selfdrive/modeld/thneed/thneed.h @@ -14,10 +14,6 @@ #include "selfdrive/modeld/thneed/include/msm_kgsl.h" -#define THNEED_RECORD 1 -#define THNEED_DEBUG 2 -#define THNEED_VERBOSE_DEBUG 4 - using namespace std; namespace json11 { @@ -110,7 +106,8 @@ class Thneed { int context_id; // protected? - int record; + bool record; + int debug; int timestamp; unique_ptr ram; vector > cmds;