early float4 stuff for binary
MODIFY
test/test_llops.py
test/test_llops.py
@@ -20,8 +20,8 @@ class TestBenchmarkCL(unittest.TestCase):
shape = (1024,1024)
buf = GPUBuffer(shape, hostbuf=np.ones(shape, dtype=np.float32))
def fxn():
- unary_op('1.01*a', buf, buf)
- its_sec = timeit(fxn, 100000, done=lambda: sync())
+ unary_op('1.01f*a', buf, buf)
+ its_sec = timeit(fxn, 100000, done=sync)
print(f"unary op (no sync) {its_sec:.2f} its/sec")
self.assertGreater(its_sec, 10000)
@@ -29,7 +29,7 @@ class TestBenchmarkCL(unittest.TestCase):
shape = (1,)
buf = GPUBuffer(shape, hostbuf=np.ones(shape, dtype=np.float32))
def fxn():
- unary_op('1.01*a', buf, buf)
+ unary_op('1.01f*a', buf, buf)
sync()
its_sec = timeit(fxn, 1000)
print(f"unary op tiny {its_sec:.2f} its/sec")
@@ -39,12 +39,22 @@ class TestBenchmarkCL(unittest.TestCase):
shape = (1024,1024)
buf = GPUBuffer(shape, hostbuf=np.ones(shape, dtype=np.float32))
def fxn():
- unary_op('1.01*a', buf, buf)
+ unary_op('1.01f*a', buf, buf)
sync()
its_sec = timeit(fxn, 1000)
print(f"unary op {its_sec:.2f} its/sec")
self.assertGreater(its_sec, 1000)
+ def test_benchmark_binary_nosync(self):
+ shape = (1024,1024)
+ buf_a = GPUBuffer(shape, hostbuf=np.ones(shape, dtype=np.float32))
+ buf_b = GPUBuffer(shape, hostbuf=np.ones(shape, dtype=np.float32))
+ def fxn():
+ binary_op('a+b', buf_a, buf_b, buf_a)
+ its_sec = timeit(fxn, 10000, done=sync)
+ print(f"binary op (no sync) {its_sec:.2f} its/sec")
+ self.assertGreater(its_sec, 1000)
+
def test_benchmark_binary(self):
shape = (1024,1024)
buf_a = GPUBuffer(shape, hostbuf=np.ones(shape, dtype=np.float32))

MODIFY
tinygrad/llops/opencl.py
tinygrad/llops/opencl.py
@@ -70,19 +70,21 @@ def get_binop_prg(code, complist):
if complist[i][j]:
idx_exprs[j] = "idx_ret%d + d%d*(%s)" % (i, i, idx_exprs[j])
- prg = """__kernel void binop(__global const float *x_g, __global const float *y_g, __global float *res_g"""+args+""") {
+ dtype = ["float", "float", "float"]
+ prg = """__kernel void binop(__global const """+dtype[0]+""" *x_g, __global const """+dtype[1]+""" *y_g, __global """+dtype[2]+""" *res_g"""+args+""") {
int gid0 = get_global_id(0);"""+compute_idx_rets+"""
- float a = x_g["""+idx_exprs[0]+"""];
- float b = y_g["""+idx_exprs[1]+"""];
+ """+dtype[0]+""" a = x_g["""+idx_exprs[0]+"""];
+ """+dtype[1]+""" b = y_g["""+idx_exprs[1]+"""];
res_g[gid0] = """+code+""";\n}"""
- return cl.Program(cl_ctx, prg).build()
+ return cl.Program(cl_ctx, prg).build(), dtype[2] == "float4"
def binary_op(code, x, y, ret):
shape_ret, dimlist, complist = binary_broadcast(x.shape, y.shape, True)
assert tuple(shape_ret) == tuple(ret.shape)
prod_list = np.array(dimlist, dtype=i32)[-1::-1].cumprod(dtype=i32)[-1::-1] # take cumprod from back to front
- prg = get_binop_prg(code, tuple(complist))
- prg.binop(cl_queue, [prod_list[0]] if len(dimlist) > 0 else [1], None, x.cl, y.cl, ret.cl, *dimlist, *(prod_list[1:]))
+ prg, is_float4 = get_binop_prg(code, tuple(complist))
+ kernel_size = [(roundup(prod_list[0])//4) if is_float4 else prod_list[0]] if len(dimlist) > 0 else [1]
+ prg.binop(cl_queue, kernel_size, None, x.cl, y.cl, ret.cl, *dimlist, *(prod_list[1:]))
return ret
def reduce_op(code, inp, ret, start="0.0"):

MODIFY
tinygrad/ops/ops_gpu.py
tinygrad/ops/ops_gpu.py
@@ -18,7 +18,7 @@ class UnaryOp(Function):
class ReLU(UnaryOp):
fop = 'max(a, (float)0.)'
- bop = 'b * (a >= 0)'
+ bop = 'b * max(sign(a), (float)0.)'
class Log(UnaryOp):
fop = 'log(a)'
@@ -52,7 +52,7 @@ class Max(Function):
def backward(ctx, grad_output):
input, ret = ctx.saved_tensors
- ret2 = binary_op("1.0*(a==b)", input, ret, Buffer(input.shape))
+ ret2 = binary_op("1.0f*(a==b)", input, ret, Buffer(input.shape))
div = reduce_op("out += a", ret2, Buffer(grad_output.shape), start="1e-10")
binary_op("a/b", ret2, div, ret2)
return binary_op('a*b', ret2, grad_output, ret2)

early float4 stuff for binary