debug is env var, and simpler faster pad2d
MODIFY
tinygrad/opsgpu.py
tinygrad/opsgpu.py
@@ -300,25 +300,23 @@ class Pad2D(Function):
prg = clbuild(ctx.cl_ctx, """
__kernel void pad2d(
__global const float *input, __global float *output,
- int cin, int py, int px, int oy, int ox, int iy, int ix
+ int py, int px, int oy, int ox, int iy, int ix
)
{
- int B = get_global_id(0);
- int C = get_global_id(1);
- int Y = get_global_id(2);
+ int BC = get_global_id(0);
+ int Y = get_global_id(1);
+ int X = get_global_id(2);
- int iptr = B*cin*iy*ix + C*iy*ix + Y*ix;
- int optr = B*cin*oy*ox + C*oy*ox + (Y+py)*ox + px;
+ int iptr = BC*iy*ix + Y*ix + X;
+ int optr = BC*oy*ox + (Y+py)*ox + px + X;
- for (int x = 0; x < ix; x++) {
- output[optr+x] = input[iptr+x];
- }
+ output[optr] = input[iptr];
}
""")
ctx.save_for_backward(padding)
- prg.pad2d(ctx.cl_queue, [bs, cin, iy], None,
+ prg.pad2d(ctx.cl_queue, [bs*cin, iy, ix], None,
x, ret,
- np.int32(cin), np.int32(padding[2]), np.int32(padding[0]),
+ np.int32(padding[2]), np.int32(padding[0]),
np.int32(oy), np.int32(ox), np.int32(iy), np.int32(ix)
)
return ret

MODIFY
tinygrad/tensor.py
tinygrad/tensor.py
from inspect import signature
import numpy as np
import time
+import os
try:
import pyopencl as cl
GPU = True
except ImportError:
# no GPU support
GPU = False
-DEBUG = False
+DEBUG = os.getenv("DEBUG", None) is not None
cl_ctx, cl_queue = None, None
def require_init_gpu():

debug is env var, and simpler faster pad2d