remove useless buffer
MODIFY
extra/thneed.py
extra/thneed.py
@@ -92,6 +92,9 @@ class Thneed:
bufs[o['id']] = buf
bufs_loaded[o['id']] = 'data' in o
+ # if it's loaded, it's saved
+ if 'data' in o:
+ self.buffers_to_save.add(buf)
# load in the programs (this isn't used)
prgs = {}
@@ -189,11 +192,10 @@ class Thneed:
row_pitch = (a.shape[0]*4*(2 if FLOAT16 else 4) + 63)//64 * 64
size = row_pitch * a.shape[1]
# this is *2 if float16 and *4 if float32
- buf = CLBuffer(size * (2 if FLOAT16 else 1))
+ buf = cl.Buffer(CL().cl_ctx, cl.mem_flags.READ_WRITE, size=size * (2 if FLOAT16 else 1))
# zero out the buffer
- zeros = np.zeros(size, dtype=np.uint8)
- CL.enqueue_copy(buf.cl, zeros, is_blocking=True)
+ CL.enqueue_copy(buf, b'\x00'*buf.size, is_blocking=True)
CLProgram("from_image_strided", """
__kernel void from_image_strided(read_only image2d_t in, __global float4 *out, int row_pitch) {
@@ -203,7 +205,7 @@ class Thneed:
l.x = get_global_id(0);
out[l.y*row_pitch + l.x] = read_imagef(in, smp, l);
}
- """, argdtypes=(None, None, np.int32))(a.shape, None, a, buf.cl, row_pitch//(4*(2 if FLOAT16 else 4)))
+ """, argdtypes=(None, None, np.int32))(a.shape, None, a, buf, row_pitch//(4*(2 if FLOAT16 else 4)))
# multiple of 32 isn't enough
jdat['objects'].append({
@@ -213,7 +215,7 @@ class Thneed:
if needs_load:
data = np.empty(size//(2 if FLOAT16 else 4), dtype=np.float32)
- CL.enqueue_copy(data, buf.cl, is_blocking=True)
+ CL.enqueue_copy(data, buf, is_blocking=True)
if FLOAT16: data = data.astype(np.float16)
weights.append(data.tobytes())
else:

MODIFY
openpilot/compile.py
openpilot/compile.py
@@ -129,8 +129,7 @@ def compile(dat, output_fn):
np.testing.assert_allclose(torch_out, thneed_out, atol=1e-4, rtol=1e-2)
# test loading/run thneed
- #_, new_np_inputs = get_random_input_tensors(input_shapes)
- new_np_inputs = np_inputs
+ _, new_np_inputs = get_random_input_tensors(input_shapes)
new_torch_out = run_onnx_torch(onnx_model, new_np_inputs).numpy()
nt = Thneed()
@@ -143,35 +142,7 @@ def compile(dat, output_fn):
nt.run()
new_thneed_out = np.empty((nt.outputs[0].size//4,), dtype=np.float32).reshape(tinygrad_out.shape)
CL.enqueue_copy(new_thneed_out, nt.outputs[0], is_blocking=True)
- try:
np.testing.assert_allclose(new_torch_out, new_thneed_out, atol=1e-4, rtol=1e-2)
- except AssertionError:
- # NOTE: this doesn't pass even if thneed passes
- # but the differences are in unused stuff
- print("THNEED ERROR")
- for i,(a,b) in enumerate(zip(t.cl_cache, nt.cl_cache)):
- for j,(c,d) in enumerate(zip(a[1][2:], b[1][2:])):
- if type(c) == cl.Buffer and c not in t.buffers_to_save:
- CL.enqueue_copy(c, b'\x00'*c.size)
- CL.enqueue_copy(d, b'\x00'*d.size)
- # rerun
- a[0].clprg(CL().cl_queue, *a[1])
- CL().cl_queue.finish()
- b[0].clprg(CL().cl_queue, *b[1])
- CL().cl_queue.finish()
- assert len(a[1]) == len(b[1])
- for j,(c,d) in enumerate(zip(a[1][2:], b[1][2:])):
- if type(c) != type(d):
- print("type mismatch", type(c), type(d))
- if type(c) == cl.Buffer:
- cc = np.empty((c.size//4,), dtype=np.float32)
- CL.enqueue_copy(cc, c, is_blocking=True)
- dd = np.empty((c.size//4,), dtype=np.float32)
- CL.enqueue_copy(dd, d, is_blocking=True)
- if not (cc == dd).all():
- print(f"mismatch in layer {i} arg {j}")
- np.testing.assert_allclose(cc, dd)
- assert False, "all layers match?"
print("thneed self-test passed!")
except ModuleNotFoundError:
pass

remove useless buffer