sadmen
/
openpilot_comma
mirror of https://github.com/commaai/openpilot.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

thneed: a more sane way of doing record/debug (#23938)

Browse Source 
Co-authored-by: Comma Device <device@comma.ai>
old-commit-hash: 5c5a56c5e6
taco
George Hotz 3 years ago committed by GitHub
parent 0734322971
commit 7efe157848
5 changed files with 29 additions and 36 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Split View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 2
      selfdrive/modeld/runners/thneedmodel.cc
  2. 2
      selfdrive/modeld/thneed/compile.cc
  3. 4
      selfdrive/modeld/thneed/serialize.cc
  4. 50
      selfdrive/modeld/thneed/thneed.cc
  5. 7
      selfdrive/modeld/thneed/thneed.h

2
selfdrive/modeld/runners/thneedmodel.cc
Unescape View File

@ -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);

2
selfdrive/modeld/thneed/compile.cc
Unescape View File

@ -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();

4
selfdrive/modeld/thneed/serialize.cc
Unescape View File

@ -63,14 +63,14 @@ void Thneed::load(const char *filename) {
map<string, cl_program> 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;
}

50
selfdrive/modeld/thneed/thneed.cc
Unescape View File

@ -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<CachedCommand>(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<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->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<GPUMalloc>(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,

7
selfdrive/modeld/thneed/thneed.h
Unescape View File

@ -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<GPUMalloc> ram;
vector<unique_ptr<CachedIoctl> > cmds;

Write Preview
Loading…
Cancel
Save