don't recopy backing
MODIFY
accel/opencl/ops_opencl.py
accel/opencl/ops_opencl.py
@@ -155,6 +155,7 @@ class OpenCLBuffer(GPUBuffer):
if self._backing is not None:
self._buf = CLBuffer(4*roundup(prod(self._backing.shape)))
CL.enqueue_copy(self._buf.cl, self._backing, is_blocking=False)
+ self._backing = None
elif self.st.contiguous:
self._buf = CLBuffer(4*roundup(prod(self.shape)))
@@ -162,7 +163,7 @@ class OpenCLBuffer(GPUBuffer):
self._buf = CLBuffer(4*roundup(prod(self._image.shape)*4))
if self._backing is not None:
CL.enqueue_copy(self._buf.cl, self._backing, is_blocking=False)
- #self._backing = None
+ self._backing = None
#print(f"converting {self.shape} back to buffer, image shape is {self._image.shape}")
CLProgram("from_image", f"""
__kernel void from_image(

MODIFY
accel/opencl/preprocessing.py
accel/opencl/preprocessing.py
@@ -56,6 +56,9 @@ def preprocessing_op(x,w,C):
C = C._replace(out_shape = (C.bs*C.oy, C.ox*C.cout//4, 4))
#x = contiguous(ctx, x, x.shapetracker) if not x.shapetracker.contiguous else x
#w = contiguous(ctx, w, w.shapetracker) if not w.shapetracker.contiguous else w
+
+ # okay but unneeded
+ #x = x.contiguous_op()
w = w.contiguous_op()
# early realize on the weights

MODIFY
tinygrad/llops/ops_gpu.py
tinygrad/llops/ops_gpu.py
@@ -125,7 +125,7 @@ class GPUBuffer:
def unary_op(x, op:UnaryOps): return type(x)(x.shape)._processing_op([("A", x)], GPUBuffer.code_for_op[op])
def binary_op(x, op:BinaryOps, y:GPUBuffer): return type(x)(x.shape)._processing_op([("A", x), ("B", y)], GPUBuffer.code_for_op[op])
- def contiguous_op(x): return x if x.st.contiguous else x.unary_op(UnaryOps.NOOP)
+ def contiguous_op(x): return type(x)(x.shape, x, x._backing) if x.st.contiguous else x.unary_op(UnaryOps.NOOP)
def movement_op(x, op:MovementOps, arg) -> GPUBuffer: return type(x)(ShapeTracker(x.st).movement_op(op, arg), x)
def reduce_op(x, op:ReduceOps, new_shape:Tuple[int, ...]): return type(x)(new_shape)._processing_op([("A", x)], code="acc", earlycode=GPUBuffer.code_for_op[op], earlybufs=set("A"), op=op)

MODIFY
tinygrad/ops.py
tinygrad/ops.py
@@ -125,6 +125,8 @@ def _realize_loadops(self:LazyBuffer) -> Tuple[DeviceBuffer, List[DeviceBuffer],
elif self.op.op == LoadOps.CONTIGUOUS:
real_src = self.op.src[0].realize(self.device)
return real_src.contiguous_op(), [real_src], LoadOps
+ else:
+ assert NotImplementedError(f"unknown LoadOp {self.op.op}")
def _realize_movementops(self:LazyBuffer) -> Tuple[DeviceBuffer, List[DeviceBuffer], OpType]:
real_src = self.op.src[0].realize(self.device)

don't recopy backing