test opencl, commit to removing the crap conv code from GPU
MODIFY
.github/workflows/test.yml
.github/workflows/test.yml
@@ -96,10 +96,32 @@ jobs:
python-version: 3.8
- name: Install Dependencies
run: pip install -e '.[gpu,testing]'
- - name: Run Pytest
- run: LAZY=0 GPU=1 python -m pytest -s -v
- - name: Run Pytest (lazy)
- run: LAZY=1 GPU=1 python -m pytest -s -v
+ - name: Run Pytest (default)
+ run: GPU=1 python -m pytest -s -v
+
+ testgpu:
+ name: OpenCL Tests
+ runs-on: ubuntu-20.04
+
+ steps:
+ - name: Checkout Code
+ uses: actions/checkout@v2
+ - name: Update packages
+ run: |
+ wget -O- https://apt.repos.intel.com/intel-gpg-keys/GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB | gpg --dearmor | sudo tee /usr/share/keyrings/oneapi-archive-keyring.gpg > /dev/null
+ echo "deb [signed-by=/usr/share/keyrings/oneapi-archive-keyring.gpg] https://apt.repos.intel.com/oneapi all main" | sudo tee /etc/apt/sources.list.d/oneAPI.list
+ sudo apt-get update
+ - name: Install OpenCL
+ #run: sudo apt-get install -y pocl-opencl-icd
+ run: sudo apt-get install -y intel-oneapi-runtime-compilers intel-oneapi-runtime-opencl
+ - name: Set up Python 3.8
+ uses: actions/setup-python@v2
+ with:
+ python-version: 3.8
+ - name: Install Dependencies
+ run: pip install -e '.[gpu,testing]'
+ - name: Run Pytest (default)
+ run: OPENCL=1 python -m pytest -s -v
testmypy:
name: Mypy Tests

MODIFY
accel/opencl/ops_opencl.py
accel/opencl/ops_opencl.py
@@ -116,8 +116,12 @@ class OpenCLBuffer(GPUBuffer):
self._buf = None
return self._image
- seen = set()
SUPPORTS_PADDING = True
+ def processing_op(x, op:ProcessingOps, w:GPUBuffer, C:ConvArgs):
+ assert op == ProcessingOps.CONV, f"{op} isn't supported"
+ return type(x)(C.out_shape)._processing_op([("input", x.contiguous_op()), ("weight", w.contiguous_op())], "acc", C)
+
+ seen = set()
def _processing_op(ret, bufs: List[Tuple[str, OpenCLBuffer]]=[], code:str="acc", C=None, start="0.0"):
if C is None:
# TODO: handle an opencl conv without the conv part

MODIFY
tinygrad/llops/ops_gpu.py
tinygrad/llops/ops_gpu.py
@@ -111,11 +111,6 @@ class GPUBuffer:
def contiguous_op(x): return x 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)
- SUPPORTS_PADDING = True
- def processing_op(x, op:ProcessingOps, w:GPUBuffer, C:ConvArgs):
- assert op == ProcessingOps.CONV, f"{op} isn't supported"
- return type(x)(C.out_shape)._processing_op([("input", x.contiguous_op()), ("weight", w.contiguous_op())], "acc", C)
-
def reduce_op(x, op:ReduceOps, new_shape:Tuple[int, ...]):
return type(x)(new_shape)._processing_op([("A", x)], GPUBuffer.code_for_op[op], None, GPUBuffer.start_for_op[op])
@@ -126,29 +121,8 @@ class GPUBuffer:
# this takes a ret index to an inp index, indexing 0 on the reduced strides
# if it's not a reduce, this should be a NOOP
view = View(ret.shape, strides_for_shape(bufs[0][1].shape))
- if C is not None: # this is a conv
- ints = ''.join(f"int {x} = {getattr(C, x)};" for x in ["H", "W", "sy", "sx", "dx", "dy", "px", "py", "groups", "rcout", "cin"])
- params = [(f"int {x}", getattr(C, x)) for x in ["oy", "ox", "iy", "ix"]]
- global_size = [C.bs*C.cout, C.oy, C.ox] # [nGk, h, w]
- assert ret.shape == C.out_shape, "output shape is wrong (NOTE: you can't reduce and conv together)"
-
- # now input and weight can be anywhere in bufs
- bufs = [(x[0], x[1].contiguous_op()) if x[0] in ["input", "weight"] else x for x in bufs]
- ewbufs = [x for x in bufs if x[0] not in ["input", "weight"]]
- assert len(bufs) == len(ewbufs)+2, "input or weight missing"
-
- # TODO: is there a way to unify this with reduce? it looks very similar
- conv_src = """
- int B = gid/(groups*rcout); int g = (gid/rcout)%groups; int c = gid % rcout;
- int Y = get_global_id(1); int X = get_global_id(2); gid = gid*oy*ox + Y*ox + X;
- for (int ci = 0; ci < cin; ci++) { for (int y = 0; y < H; y++) { for (int x = 0; x < W; x++) {
- int idx_y = y*dy + Y*sy - py;
- int idx_x = x*dx + X*sx - px;
- int valid = (idx_y >= 0 && idx_y < iy && idx_x >= 0 && idx_x < ix);
- acc += valid * input_g[B*groups*cin*iy*ix + g*cin*iy*ix + ci*iy*ix + clamp(idx_y, 0, iy-1)*ix + clamp(idx_x, 0, ix-1)] * \
- weight_g[g*rcout*cin*H*W + c*cin*H*W + ci*H*W + y*W + x];
- } } }"""
- elif ret.shape != bufs[0][1].shape: # this is a reduce
+ assert C is None
+ if ret.shape != bufs[0][1].shape: # this is a reduce
# reverse operation of expand, this validates inputs
# generate loops with combined adjacent reduce axis
acc = 1
@@ -156,7 +130,7 @@ class GPUBuffer:
if stride == 0: loop.append((f"for (int axis_{len(loop)} = 0; axis_{len(loop)} < {shp}; axis_{len(loop)}++) {{", f"idx += {acc}; }} idx -= {shp*acc};"))
acc *= shp
- kernel_name = "conv" if C is not None else ("reduce" if len(loop) > 0 else "elementwise")
+ kernel_name = "reduce" if len(loop) > 0 else "elementwise"
views = {name:buf.contiguous_view_constant_fold(name) for name, buf in ewbufs}
buf_types = [f"__global const float *{name}_g" for name, _ in bufs if name not in views or views[name][1]]
conv_prg = CLProgram(kernel_name, f"""{chr(10).join([x[0] for x in views.values()])}

MODIFY
tinygrad/ops.py
tinygrad/ops.py
@@ -240,11 +240,11 @@ class LazyBuffer:
if op == MovementOps.PAD and arg == tuple((0,0) for _ in x.shape): return x
if op == MovementOps.FLIP and all(s == 1 or i not in arg for i,s in enumerate(x.shape)): return x
- # two reshapes in a row is one reshape
+ # two ops in a row is one op
if op == MovementOps.RESHAPE and x.realized is None and x.op.op == MovementOps.RESHAPE: return x.op.src[0].movement_op(op, arg)
-
- # two permutes in a row is one permute
if op == MovementOps.PERMUTE and x.realized is None and x.op.op == MovementOps.PERMUTE: return x.op.src[0].movement_op(op, tuple(x.op.arg[i] for i in arg))
+ if op == MovementOps.SHRINK and x.realized is None and x.op.op == MovementOps.SHRINK: return x.op.src[0].movement_op(op, arg)
+ if op == MovementOps.PAD and x.realized is None and x.op.op == MovementOps.PAD: return x.op.src[0].movement_op(op, tuple((b1+b2, e1+e2) for (b1,e1),(b2,e2) in zip(x.op.arg, arg)))
# some permutes are actually just reshapes
if op == MovementOps.PERMUTE and ShapeTracker(x.shape).movement_op(op, arg).contiguous: return x.movement_op(MovementOps.RESHAPE, tuple(x.shape[i] for i in arg))
@@ -253,6 +253,7 @@ class LazyBuffer:
# if this MovementOp is being applied to a BinaryOp, apply the MovementOp to all the BinaryOp inputs instead
def replace_with_movement_op(y:Union[LazyOp, LazyBuffer]) -> LazyBuffer:
if isinstance(y, LazyBuffer): return y.movement_op(op, arg)
+ assert isinstance(y.op, BinaryOps) or isinstance(y.op, UnaryOps)
return elementwise_op(y.op, *[replace_with_movement_op(z) for z in y.src])
return replace_with_movement_op(x.op)
@@ -272,7 +273,7 @@ class LazyBuffer:
# TODO: fixup C?
if NOCONV or not getattr(x.dbuffer, "SUPPORTS_PADDING", False): x = x.slice(((0, x.shape[0]), (0, x.shape[1]), (-C.py, x.shape[2]+C.py_), (-C.px, x.shape[3]+C.px_)))
- if NOCONV:
+ if NOCONV or not getattr(x.dbuffer, "processing_op", False):
# universal conv, just mul and reduce
# TODO: is there any way to replace strided with other movement ops?
x = x.movement_op(MovementOps.STRIDED, (

test opencl, commit to removing the crap conv code from GPU