align with latest design

Author: leofang

cuda_core/cuda/core/__init__.py

@@ -2,9 +2,9 @@
 #
 # SPDX-License-Identifier: LicenseRef-NVIDIA-SOFTWARE-LICENSE
 
-from cuda.core._compiler import Compiler
 from cuda.core._device import Device
 from cuda.core._event import EventOptions
 from cuda.core._launcher import LaunchConfig, launch
+from cuda.core._program import Program
 from cuda.core._stream import Stream, StreamOptions
 from cuda.core._version import __version__

cuda_core/cuda/core/_device.py

@@ -8,7 +8,7 @@
 
 from cuda import cuda, cudart
 from cuda.core._utils import handle_return, ComputeCapability, CUDAError, \
-                           precondition
+                             precondition
 from cuda.core._context import Context, ContextOptions
 from cuda.core._memory import _DefaultAsyncMempool, Buffer, MemoryResource
 from cuda.core._stream import default_stream, Stream, StreamOptions
@@ -50,7 +50,7 @@ def __new__(cls, device_id=None):
     def _check_context_initialized(self, *args, **kwargs):
         if not self._has_inited:
             raise CUDAError("the device is not yet initialized, "
-                            "perhaps you forgot to call .use() first?")
+                            "perhaps you forgot to call .set_current() first?")
 
     @property
     def device_id(self) -> int:
@@ -120,14 +120,14 @@ def __int__(self):
     def __repr__(self):
         return f"<Device {self._id} ({self.name})>"
 
-    def use(self, ctx: Context=None) -> Union[Context, None]:
+    def set_current(self, ctx: Context=None) -> Union[Context, None]:
         """
         Entry point of this object. Users always start a code by
         calling this method, e.g.
         
         >>> from cuda.core import Device
         >>> dev0 = Device(0)
-        >>> dev0.use()
+        >>> dev0.set_current()
         >>> # ... do work on device 0 ...
         
         The optional ctx argument is for advanced users to bind a

cuda_core/cuda/core/_event.py

@@ -13,7 +13,7 @@
 
 @dataclass
 class EventOptions:
-    disable_timing: Optional[bool] = False
+    enable_timing: Optional[bool] = False
     busy_waited_sync: Optional[bool] = False
     support_ipc: Optional[bool] = False
 
@@ -37,8 +37,9 @@ def _init(options: Optional[EventOptions]=None):
 
         options = check_or_create_options(EventOptions, options, "Event options")
         flags = 0x0
-        self._timing_disabled = self._busy_waited = False
-        if options.disable_timing:
+        self._timing_disabled = False
+        self._busy_waited = False
+        if not options.enable_timing:
             flags |= cuda.CUevent_flags.CU_EVENT_DISABLE_TIMING
             self._timing_disabled = True
         if options.busy_waited_sync:
@@ -91,4 +92,4 @@ def is_done(self) -> bool:
 
     @property
     def handle(self) -> int:
-        return self._handle
+        return int(self._handle)

cuda_core/cuda/core/_memory.py

@@ -75,7 +75,7 @@ def device_id(self) -> int:
             return self._mr.device_id
         raise NotImplementedError
 
-    def copy_to(self, dst: Buffer=None, stream=None) -> Buffer:
+    def copy_to(self, dst: Buffer=None, *, stream) -> Buffer:
         # Copy from this buffer to the dst buffer asynchronously on the
         # given stream. The dst buffer is returned. If the dst is not provided,
         # allocate one from self.memory_resource. Raise an exception if the
@@ -92,7 +92,7 @@ def copy_to(self, dst: Buffer=None, stream=None) -> Buffer:
             cuda.cuMemcpyAsync(dst._ptr, self._ptr, self._size, stream._handle))
         return dst
 
-    def copy_from(self, src: Buffer, stream=None):
+    def copy_from(self, src: Buffer, *, stream):
         # Copy from the src buffer to this buffer asynchronously on the
         # given stream. Raise an exception if the stream is not provided. 
         if stream is None:

cuda_core/cuda/core/_memoryview.pyx

@@ -31,6 +31,17 @@ cdef class GPUMemoryView:
     readonly: bool = None
     obj: Any = None
 
+    def __init__(self, obj=None, stream_ptr=None):
+        if obj is not None:
+            # populate self's attributes
+            if check_has_dlpack(obj):
+                view_as_dlpack(obj, stream_ptr, self)
+            else:
+                view_as_cai(obj, stream_ptr, self)
+        else:
+            # default construct
+            pass
+
     def __repr__(self):
         return (f"GPUMemoryView(ptr={self.ptr},\n"
               + f"              shape={self.shape},\n"
@@ -57,22 +68,27 @@ cdef str get_simple_repr(obj):
     return obj_repr
 
 
+cdef bint check_has_dlpack(obj) except*:
+    cdef bint has_dlpack
+    if hasattr(obj, "__dlpack__") and hasattr(obj, "__dlpack_device__"):
+        has_dlpack = True
+    elif hasattr(obj, "__cuda_array_interface__"):
+        has_dlpack = False
+    else:
+        raise RuntimeError(
+            "the input object does not support any data exchange protocol")
+    return has_dlpack
+
+
 cdef class _GPUMemoryViewProxy:
 
     cdef:
         object obj
         bint has_dlpack
 
     def __init__(self, obj):
-        if hasattr(obj, "__dlpack__") and hasattr(obj, "__dlpack_device__"):
-            has_dlpack = True
-        elif hasattr(obj, "__cuda_array_interface__"):
-            has_dlpack = False
-        else:
-            raise RuntimeError(
-                "the input object does not support any data exchange protocol")
         self.obj = obj
-        self.has_dlpack = has_dlpack
+        self.has_dlpack = check_has_dlpack(obj)
 
     cpdef GPUMemoryView view(self, stream_ptr=None):
         if self.has_dlpack:
@@ -81,7 +97,7 @@ cdef class _GPUMemoryViewProxy:
             return view_as_cai(self.obj, stream_ptr)
 
 
-cdef GPUMemoryView view_as_dlpack(obj, stream_ptr):
+cdef GPUMemoryView view_as_dlpack(obj, stream_ptr, view=None):
     cdef int dldevice, device_id, i
     cdef bint device_accessible, versioned, is_readonly
     dldevice, device_id = obj.__dlpack_device__()
@@ -144,7 +160,7 @@ cdef GPUMemoryView view_as_dlpack(obj, stream_ptr):
         dl_tensor = &dlm_tensor.dl_tensor
         is_readonly = False
 
-    cdef GPUMemoryView buf = GPUMemoryView()
+    cdef GPUMemoryView buf = GPUMemoryView() if view is None else view
     buf.ptr = <intptr_t>(dl_tensor.data)
     buf.shape = tuple(int(dl_tensor.shape[i]) for i in range(dl_tensor.ndim))
     if dl_tensor.strides:
@@ -226,7 +242,7 @@ cdef object dtype_dlpack_to_numpy(DLDataType* dtype):
     return numpy.dtype(np_dtype)
 
 
-cdef GPUMemoryView view_as_cai(obj, stream_ptr):
+cdef GPUMemoryView view_as_cai(obj, stream_ptr, view=None):
     cdef dict cai_data = obj.__cuda_array_interface__
     if cai_data["version"] < 3:
         raise BufferError("only CUDA Array Interface v3 or above is supported")
@@ -235,7 +251,7 @@ cdef GPUMemoryView view_as_cai(obj, stream_ptr):
     if stream_ptr is None:
         raise BufferError("stream=None is ambiguous with view()")
 
-    cdef GPUMemoryView buf = GPUMemoryView()
+    cdef GPUMemoryView buf = GPUMemoryView() if view is None else view
     buf.obj = obj
     buf.ptr, buf.readonly = cai_data["data"]
     buf.shape = cai_data["shape"]

cuda_core/cuda/core/_module.py

@@ -30,14 +30,14 @@ def __init__(self):
     @staticmethod
     def _from_obj(obj, mod):
         assert isinstance(obj, (cuda.CUkernel, cuda.CUfunction))
-        assert isinstance(mod, Module)
+        assert isinstance(mod, ObjectCode)
         ker = Kernel.__new__(Kernel)
         ker._handle = obj
         ker._module = mod
         return ker
 
 
-class Module:
+class ObjectCode:
 
     __slots__ = ("_handle", "_code_type", "_module", "_loader", "_sym_map")
     _supported_code_type = ("cubin", "ptx", "fatbin")

cuda_core/cuda/core/_compiler.py cuda_core/cuda/core/_program.pycuda_core/cuda/core/_compiler.py renamed to cuda_core/cuda/core/_program.py

@@ -4,10 +4,10 @@
 
 from cuda import nvrtc
 from cuda.core._utils import handle_return
-from cuda.core._module import Module
+from cuda.core._module import ObjectCode
 
 
-class Compiler:
+class Program:
 
     __slots__ = ("_handle", "_backend", )
     _supported_code_type = ("c++", )
@@ -26,6 +26,8 @@ def __init__(self, code, code_type):
             self._handle = handle_return(
                 nvrtc.nvrtcCreateProgram(code.encode(), b"", 0, [], []))
             self._backend = "nvrtc"
+        else:
+            raise NotImplementedError
 
     def __del__(self):
         self.close()
@@ -72,7 +74,7 @@ def compile(self, target_type, options=(), name_expressions=(), logs=None):
 
             # TODO: handle jit_options for ptx?
 
-            return Module(data, target_type, symbol_mapping=symbol_mapping)
+            return ObjectCode(data, target_type, symbol_mapping=symbol_mapping)
 
     @property
     def backend(self):