Fix GPT2 with OpenCL backend. (#6821)

* Fix GPT2 with OpenCL backend.

* Add test for unaligned copies into OpenCL buffers.

test/test_ocl.py

@@ -1,5 +1,7 @@
 import unittest
 from tinygrad import Device
+from tinygrad.device import Buffer
+from tinygrad.dtype import dtypes
 from tinygrad.helpers import CI
 from tinygrad.runtime.ops_gpu import CLDevice, CLAllocator, CLCompiler, CLProgram
 
@@ -18,3 +20,12 @@ class TestCLError(unittest.TestCase):
     with self.assertRaises(RuntimeError) as err:
       CLProgram(device, name="", lib=CLCompiler(device, "test").compile("__kernel void test(__global int* a) { a[0] = 1; }"))
     assert str(err.exception) == "OpenCL Error -46: CL_INVALID_KERNEL_NAME"
+
+  def test_unaligned_copy(self):
+    data = list(range(65))
+    unaligned = memoryview(bytearray(data))[1:]
+    buffer = Buffer("GPU", 64, dtypes.uint8).allocate()
+    buffer.copyin(unaligned)
+    result = memoryview(bytearray(len(data) - 1))
+    buffer.copyout(result)
+    assert unaligned == result, "Unaligned data copied in must be equal to data copied out."

tinygrad/runtime/ops_gpu.py

@@ -74,6 +74,7 @@ class CLAllocator(LRUAllocator):
       check(cl.clEnqueueWriteImage(self.device.queue, dest[0], False, (ctypes.c_size_t * 3)(0,0,0),
                                    (ctypes.c_size_t * 3)(dest[1].image.shape[1],dest[1].image.shape[0],1), 0, 0, from_mv(src), 0, None, None))
     else:
+      if ctypes.addressof(ctypes.c_char.from_buffer(src)) % 16: src = memoryview(bytearray(src))
       check(cl.clEnqueueWriteBuffer(self.device.queue, dest[0], False, 0, len(src)*src.itemsize, from_mv(src), 0, None, None))
     self.device.pending_copyin.append(src)    # NOTE: these can't be freed until the GPU actually executes this command
   def copyout(self, dest:memoryview, src:Tuple[ctypes._CData, BufferOptions]):