spacecruft
/
nopenpilot
1
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

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

Co-authored-by: Comma Device <device@comma.ai>
pull/23939/merge
George Hotz 2022-03-10 16:33:35 -08:00 committed by GitHub
parent 2c0d68120f
commit 5c5a56c5e6
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
5 changed files with 29 additions and 36 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
selfdrive/modeld/runners/thneedmodel.cc
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
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
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
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
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;