refactor to Program class (#4476)

* refactor to Program class

* switch to Program

* fix tests

* smaller diff

* self.p

* more tests

* fix metal test

* tests

* fix openpilot

* move that to linearizer

* p.launchdims
This commit is contained in:
George Hotz
2024-05-09 17:29:07 -07:00
committed by GitHub
parent 5de4a46f10
commit c9e84ed0da
17 changed files with 94 additions and 75 deletions

View File

@@ -10,7 +10,7 @@ from tinygrad.dtype import dtypes
# *** first, we implement the atan2 op at the lowest level ***
# `atan2_gpu` for GPUBuffers and `atan2_cpu` for CPUBuffers
from tinygrad.lazy import Buffer, create_lazybuffer
from tinygrad.device import CompiledRunner, Device
from tinygrad.device import CompiledRunner, Device, Program
from tinygrad.shape.shapetracker import ShapeTracker
# we don't always have GPU support, so the type signature is the abstract CompiledBuffer instead of GPUBuffer
@@ -21,7 +21,7 @@ def atan2_gpu(ret:Buffer, a:Buffer, b:Buffer):
int idx = get_global_id(0);
c[idx] = atan2(a[idx], b[idx]);
}"""
CompiledRunner("atan2_gpu", src, ret.device, global_size=[ret.size]).exec([ret, a, b])
CompiledRunner(Program("atan2_gpu", src, ret.device, global_size=[ret.size,1,1])).exec([ret, a, b])
def atan2_cpu(ret:Buffer, a:Buffer, b:Buffer): ret.copyin(np.require(np.arctan2(a._buf, b._buf), requirements='C').data)