working on osx

.github/workflows/test.yml

@@ -66,6 +66,8 @@ jobs:
         run: pip install .
       - name: Test gpuctypes import
         run: python3 -c "import gpuctypes.${{ matrix.task }} as ${{ matrix.task }}"
+      - name: Run device specific test
+        run: python3 test/test_${{ matrix.task }}.py -v
       - uses: actions/upload-artifact@v3
         with:
           name: Generated ${{ matrix.task }} Header

gpuctypes/opencl.py

@@ -117,7 +117,10 @@ class Union(ctypes.Union, AsDictMixin):
 
 
 _libraries = {}
-_libraries['libOpenCL.so'] = ctypes.CDLL('/usr/lib/x86_64-linux-gnu/libOpenCL.so')
+try:
+    _libraries['libOpenCL.so'] = ctypes.CDLL('/usr/lib/x86_64-linux-gnu/libOpenCL.so')
+except OSError:
+    _libraries['libOpenCL.so'] = ctypes.CDLL('/System/Library/Frameworks/OpenCL.framework/Versions/A/OpenCL')
 c_int128 = ctypes.c_ubyte*16
 c_uint128 = c_int128
 void = None
@@ -613,13 +616,13 @@ def char_pointer_cast(string, encoding='utf-8'):
 # CL_VERSION_MINOR_MASK = ((1<<(10)) # macro
 # CL_VERSION_PATCH_MASK = ((1<<(12)) # macro
 # def CL_VERSION_MAJOR(version):  # macro
-#    return ((version)>>((10)+(12)))  
+#    return ((version)>>((10)+(12)))
 # def CL_VERSION_MINOR(version):  # macro
-#    return (((version)>>(12))&((1<<(10)))  
+#    return (((version)>>(12))&((1<<(10)))
 # def CL_VERSION_PATCH(version):  # macro
-#    return ((version)&((1<<(12)))  
+#    return ((version)&((1<<(12)))
 # def CL_MAKE_VERSION(major, minor, patch):  # macro
-#    return ((((major)&((1<<(10)))<<((10)+(12)))|(((minor)&((1<<(10)))<<(12))|((patch)&((1<<(12))))  
+#    return ((((major)&((1<<(10)))<<((10)+(12)))|(((minor)&((1<<(10)))<<(12))|((patch)&((1<<(12))))
 class struct__cl_platform_id(Structure):
     pass
 

test/ctypes_helpers.py

@@ -0,0 +1,8 @@
+from typing import List
+import ctypes
+
+c_char_p_p = ctypes.POINTER(ctypes.POINTER(ctypes.c_char))
+def to_char_p_p(options: List[str]) -> c_char_p_p:
+  c_options = (ctypes.POINTER(ctypes.c_char) * len(options))()
+  c_options[:] = [ctypes.create_string_buffer(o.encode("utf-8")) for o in options]
+  return c_options

test/test_cuda.py

@@ -0,0 +1,7 @@
+import unittest
+
+class TestCUDA(unittest.TestCase):
+  pass
+
+if __name__ == '__main__':
+  unittest.main()

test/test_hip.py

@@ -1,18 +1,19 @@
-from typing import List
 import unittest
+import os
 import ctypes
 import gpuctypes.hip as hip
+from ctypes_helpers import to_char_p_p
 
-def check(status):
-  #{ctypes.string_at(hip.hipGetErrorString(hip.hipGetLastError()))}
-  if status != 0: raise RuntimeError(f"HIP Error {status}")
+def expectedFailureIf(condition):
+  def wrapper(func):
+    if condition: return unittest.expectedFailure(func)
+    else: return func
+  return wrapper
+
+CI = os.getenv("CI", "") != ""
 
-# put in ctypes_helpers.py in tinygrad
-c_char_p_p = ctypes.POINTER(ctypes.POINTER(ctypes.c_char))
-def to_char_p_p(options: List[str]) -> c_char_p_p:
-  c_options = (ctypes.POINTER(ctypes.c_char) * len(options))()
-  c_options[:] = [ctypes.create_string_buffer(o.encode("utf-8")) for o in options]
-  return c_options
+def check(status):
+  if status != 0: raise RuntimeError(f"HIP Error {status}, {ctypes.string_at(hip.hipGetErrorString(status)).decode()}")
 
 def get_hip_bytes(arg, get_sz, get_str) -> bytes:
   sz = ctypes.c_size_t()
@@ -22,28 +23,38 @@ def get_hip_bytes(arg, get_sz, get_str) -> bytes:
   return ctypes.string_at(mstr, size=sz.value)
 
 class TestHIP(unittest.TestCase):
+  def test_has_methods(self):
+    assert hip.hipMalloc is not None
+    assert hip.hiprtcCompileProgram is not None
+    assert hip.hipGetDeviceProperties is not None
+
+  @expectedFailureIf(CI)
   def test_malloc(self):
     ptr = ctypes.c_void_p()
     check(hip.hipMalloc(ctypes.byref(ptr), 100))
     assert ptr.value != 0
     check(hip.hipFree(ptr))
 
+  @expectedFailureIf(CI)
   def test_get_device_properties(self) -> hip.hipDeviceProp_t:
     device_properties = hip.hipDeviceProp_t()
     check(hip.hipGetDeviceProperties(device_properties, 0))
     print(device_properties.gcnArchName)
     return device_properties
 
+  @expectedFailureIf(CI)
   def test_compile_fail(self):
     prg = "void test() { {"
     prog = hip.hiprtcProgram()
     check(hip.hiprtcCreateProgram(ctypes.pointer(prog), prg.encode(), "<null>".encode(), 0, None, None))
     status = hip.hiprtcCompileProgram(prog, 0, None)
     assert status != 0
+    print(f"compile fail returned {status}")
     log = get_hip_bytes(prog, hip.hiprtcGetProgramLogSize, hip.hiprtcGetProgramLog).decode()
     assert len(log) > 10
     print(log)
 
+  @expectedFailureIf(CI)
   def test_compile(self):
     prg = "void test() { }"
     prog = hip.hiprtcProgram()

test/test_opencl.py

@@ -0,0 +1,51 @@
+import ctypes
+import unittest
+from ctypes_helpers import to_char_p_p
+import gpuctypes.opencl as cl
+
+def check(status):
+  if status != 0: raise RuntimeError(f"OpenCL Error {status}")
+
+class TestOpenCL(unittest.TestCase):
+  @classmethod
+  def setUpClass(cls):
+    num_platforms = ctypes.c_uint32()
+    platform_array = (cl.cl_platform_id * 1)()
+    check(cl.clGetPlatformIDs(1, platform_array, ctypes.byref(num_platforms)))
+    assert num_platforms.value > 0, "didn't get platform"
+
+    device_array = (cl.cl_device_id * 1)()
+    num_devices = ctypes.c_uint32()
+    check(cl.clGetDeviceIDs(platform_array[0], cl.CL_DEVICE_TYPE_DEFAULT, 1, device_array, ctypes.byref(num_devices)))
+    assert num_devices.value > 0, "didn't get device"
+
+    status = ctypes.c_int32()
+    cls.context = cl.clCreateContext(None, 1, device_array, ctypes.cast(None, cl.clCreateContext.argtypes[3]), None, ctypes.byref(status))
+    check(status.value)
+    assert cls.context is not None
+
+    cls.queue = cl.clCreateCommandQueue(cls.context, device_array[0], 0, ctypes.byref(status))
+    check(status.value)
+    assert cls.queue is not None
+
+  def test_malloc(self):
+    status = ctypes.c_int32()
+    buf = cl.clCreateBuffer(self.context, cl.CL_MEM_READ_ONLY, 4 * 5, None, ctypes.byref(status))
+    assert buf is not None
+
+  def test_create_program(self):
+    prog = """
+    __kernel void vector_add(__global const int *A, __global const int *B, __global int *C) {
+      int i = get_global_id(0);
+      C[i] = A[i] + B[i];
+    }
+    """
+    num_programs = 1
+    sizes = (ctypes.c_size_t * num_programs)()
+    sizes[0] = len(prog)
+    status = ctypes.c_int32()
+    program = cl.clCreateProgramWithSource(self.context, num_programs, to_char_p_p([prog]), sizes, ctypes.byref(status))
+    assert program is not None
+
+if __name__ == '__main__':
+  unittest.main()