Merge branch 'main' into IPC

Author: keenan-simpson

cuda_bindings/examples/0_Introduction/clock_nvrtc_test.py

@@ -58,6 +58,10 @@
 NUM_THREADS = 256
 
 
+def elems_to_bytes(nelems, dt):
+    return nelems * np.dtype(dt).itemsize
+
+
 def main():
     print("CUDA Clock sample")
 
@@ -75,31 +79,31 @@ def main():
     kernelHelper = common.KernelHelper(clock_nvrtc, devID)
     kernel_addr = kernelHelper.getFunction(b"timedReduction")
 
-    dinput = checkCudaErrors(cuda.cuMemAlloc(np.dtype(np.float32).itemsize * NUM_THREADS * 2))
-    doutput = checkCudaErrors(cuda.cuMemAlloc(np.dtype(np.float32).itemsize * NUM_BLOCKS))
-    dtimer = checkCudaErrors(cuda.cuMemAlloc(np.dtype(np.int64).itemsize * NUM_BLOCKS * 2))
-    checkCudaErrors(cuda.cuMemcpyHtoD(dinput, hinput, np.dtype(np.float32).itemsize * NUM_THREADS * 2))
+    dinput = checkCudaErrors(cuda.cuMemAlloc(hinput.nbytes))
+    doutput = checkCudaErrors(cuda.cuMemAlloc(elems_to_bytes(NUM_BLOCKS, np.float32)))
+    dtimer = checkCudaErrors(cuda.cuMemAlloc(timer.nbytes))
+    checkCudaErrors(cuda.cuMemcpyHtoD(dinput, hinput, hinput.nbytes))
+
+    args = ((dinput, doutput, dtimer), (None, None, None))
+    shared_memory_nbytes = elems_to_bytes(2 * NUM_THREADS, np.float32)
 
-    arr = ((dinput, doutput, dtimer), (None, None, None))
+    grid_dims = (NUM_BLOCKS, 1, 1)
+    block_dims = (NUM_THREADS, 1, 1)
 
     checkCudaErrors(
         cuda.cuLaunchKernel(
             kernel_addr,
-            NUM_BLOCKS,
-            1,
-            1,  # grid dim
-            NUM_THREADS,
-            1,
-            1,  # block dim
-            np.dtype(np.float32).itemsize * 2 * NUM_THREADS,
+            *grid_dims,  # grid dim
+            *block_dims,  # block dim
+            shared_memory_nbytes,
             0,  # shared mem, stream
-            arr,
+            args,
             0,
         )
     )  # arguments
 
     checkCudaErrors(cuda.cuCtxSynchronize())
-    checkCudaErrors(cuda.cuMemcpyDtoH(timer, dtimer, np.dtype(np.int64).itemsize * NUM_BLOCKS * 2))
+    checkCudaErrors(cuda.cuMemcpyDtoH(timer, dtimer, timer.nbytes))
     checkCudaErrors(cuda.cuMemFree(dinput))
     checkCudaErrors(cuda.cuMemFree(doutput))
     checkCudaErrors(cuda.cuMemFree(dtimer))

cuda_bindings/examples/0_Introduction/simpleCubemapTexture_test.py

@@ -6,7 +6,6 @@
 # this software and related documentation outside the terms of the EULA
 # is strictly prohibited.
 import ctypes
-import math
 import sys
 import time
 
@@ -103,18 +102,11 @@ def main():
     num_faces = 6
     num_layers = 1
     cubemap_size = width * width * num_faces
-    size = cubemap_size * num_layers * np.dtype(np.float32).itemsize
-    h_data = np.zeros(cubemap_size * num_layers, dtype="float32")
-
-    for i in range(cubemap_size * num_layers):
-        h_data[i] = i
+    h_data = np.arange(cubemap_size * num_layers, dtype="float32")
+    size = h_data.nbytes
 
     # This is the expected transformation of the input data (the expected output)
-    h_data_ref = np.zeros(cubemap_size * num_layers, dtype="float32")
-
-    for layer in range(num_layers):
-        for i in range(cubemap_size):
-            h_data_ref[layer * cubemap_size + i] = -h_data[layer * cubemap_size + i] + layer
+    h_data_ref = np.repeat(np.arange(num_layers, dtype=h_data.dtype), cubemap_size) - h_data
 
     # Allocate device memory for result
     d_data = checkCudaErrors(cudart.cudaMalloc(size))
@@ -130,10 +122,11 @@ def main():
             cudart.cudaArrayCubemap,
         )
     )
+    width_nbytes = h_data[:width].nbytes
     myparms = cudart.cudaMemcpy3DParms()
     myparms.srcPos = cudart.make_cudaPos(0, 0, 0)
     myparms.dstPos = cudart.make_cudaPos(0, 0, 0)
-    myparms.srcPtr = cudart.make_cudaPitchedPtr(h_data, width * np.dtype(np.float32).itemsize, width, width)
+    myparms.srcPtr = cudart.make_cudaPitchedPtr(h_data, width_nbytes, width, width)
     myparms.dstArray = cu_3darray
     myparms.extent = cudart.make_cudaExtent(width, width, num_faces)
     myparms.kind = cudart.cudaMemcpyKind.cudaMemcpyHostToDevice
@@ -211,23 +204,21 @@ def main():
     print(f"{cubemap_size / ((stop - start + 1) / 1000.0) / 1e6:.2f} Mtexlookups/sec")
 
     # Allocate mem for the result on host side
-    h_odata = np.zeros(cubemap_size * num_layers, dtype="float32")
+    h_odata = np.empty_like(h_data)
     # Copy result from device to host
     checkCudaErrors(cudart.cudaMemcpy(h_odata, d_data, size, cudart.cudaMemcpyKind.cudaMemcpyDeviceToHost))
 
-    print("Comparing kernel output to expected data")
-    MIN_EPSILON_ERROR = 5.0e-3
-    for i in range(cubemap_size * num_layers):
-        d = h_odata[i] - h_data_ref[i]
-        if math.fabs(d) > MIN_EPSILON_ERROR:
-            print("Failed")
-            sys.exit(-1)
-    print("Passed")
-
     checkCudaErrors(cudart.cudaDestroyTextureObject(tex))
     checkCudaErrors(cudart.cudaFree(d_data))
     checkCudaErrors(cudart.cudaFreeArray(cu_3darray))
 
+    print("Comparing kernel output to expected data")
+    MIN_EPSILON_ERROR = 5.0e-3
+    if np.max(np.abs(h_odata - h_data_ref)) > MIN_EPSILON_ERROR:
+        print("Failed")
+        sys.exit(-1)
+    print("Passed")
+
 
 if __name__ == "__main__":
     main()

cuda_bindings/examples/0_Introduction/vectorAddDrv_test.py

@@ -38,7 +38,7 @@
 def main():
     print("Vector Addition (Driver API)")
     N = 50000
-    size = N * np.dtype(np.float32).itemsize
+    nbytes = N * np.dtype(np.float32).itemsize
 
     # Initialize
     checkCudaErrors(cuda.cuInit(0))
@@ -57,18 +57,18 @@ def main():
     _VecAdd_kernel = kernelHelper.getFunction(b"VecAdd_kernel")
 
     # Allocate input vectors h_A and h_B in host memory
-    h_A = np.random.rand(size).astype(dtype=np.float32)
-    h_B = np.random.rand(size).astype(dtype=np.float32)
-    h_C = np.random.rand(size).astype(dtype=np.float32)
+    h_A = np.random.rand(N).astype(dtype=np.float32)
+    h_B = np.random.rand(N).astype(dtype=np.float32)
+    h_C = np.random.rand(N).astype(dtype=np.float32)
 
     # Allocate vectors in device memory
-    d_A = checkCudaErrors(cuda.cuMemAlloc(size))
-    d_B = checkCudaErrors(cuda.cuMemAlloc(size))
-    d_C = checkCudaErrors(cuda.cuMemAlloc(size))
+    d_A = checkCudaErrors(cuda.cuMemAlloc(nbytes))
+    d_B = checkCudaErrors(cuda.cuMemAlloc(nbytes))
+    d_C = checkCudaErrors(cuda.cuMemAlloc(nbytes))
 
     # Copy vectors from host memory to device memory
-    checkCudaErrors(cuda.cuMemcpyHtoD(d_A, h_A, size))
-    checkCudaErrors(cuda.cuMemcpyHtoD(d_B, h_B, size))
+    checkCudaErrors(cuda.cuMemcpyHtoD(d_A, h_A, nbytes))
+    checkCudaErrors(cuda.cuMemcpyHtoD(d_B, h_B, nbytes))
 
     if True:
         # Grid/Block configuration
@@ -98,7 +98,7 @@ def main():
 
     # Copy result from device memory to host memory
     # h_C contains the result in host memory
-    checkCudaErrors(cuda.cuMemcpyDtoH(h_C, d_C, size))
+    checkCudaErrors(cuda.cuMemcpyDtoH(h_C, d_C, nbytes))
 
     for i in range(N):
         sum_all = h_A[i] + h_B[i]

cuda_core/cuda/core/experimental/_context.py

@@ -15,13 +15,13 @@ class ContextOptions:
 class Context:
     __slots__ = ("_handle", "_id")
 
-    def __init__(self):
-        raise NotImplementedError("TODO")
+    def __new__(self, *args, **kwargs):
+        raise RuntimeError("Context objects cannot be instantiated directly. Please use Device or Stream APIs.")
 
-    @staticmethod
-    def _from_ctx(obj, dev_id):
+    @classmethod
+    def _from_ctx(cls, obj, dev_id):
         assert isinstance(obj, driver.CUcontext)
-        ctx = Context.__new__(Context)
+        ctx = super().__new__(cls)
         ctx._handle = obj
         ctx._id = dev_id
         return ctx

cuda_core/cuda/core/experimental/_device.py

@@ -22,13 +22,14 @@ class DeviceProperties:
     Attributes are read-only and provide information about the device.
     """
 
-    def __init__(self):
-        raise RuntimeError("DeviceProperties should not be instantiated directly")
+    def __new__(self, *args, **kwargs):
+        raise RuntimeError("DeviceProperties cannot be instantiated directly. Please use Device APIs.")
 
     __slots__ = ("_handle", "_cache")
 
-    def _init(handle):
-        self = DeviceProperties.__new__(DeviceProperties)
+    @classmethod
+    def _init(cls, handle):
+        self = super().__new__(cls)
         self._handle = handle
         self._cache = {}
         return self

cuda_core/cuda/core/experimental/_event.py

@@ -67,14 +67,14 @@ def close(self):
                 handle_return(driver.cuEventDestroy(self.handle))
                 self.handle = None
 
-    __slots__ = ("__weakref__", "_mnff", "_timing_disabled", "_busy_waited")
+    def __new__(self, *args, **kwargs):
+        raise RuntimeError("Event objects cannot be instantiated directly. Please use Stream APIs (record).")
 
-    def __init__(self):
-        raise NotImplementedError("directly creating an Event object can be ambiguous. Please call Stream.record().")
+    __slots__ = ("__weakref__", "_mnff", "_timing_disabled", "_busy_waited")
 
-    @staticmethod
-    def _init(options: Optional[EventOptions] = None):
-        self = Event.__new__(Event)
+    @classmethod
+    def _init(cls, options: Optional[EventOptions] = None):
+        self = super().__new__(cls)
         self._mnff = Event._MembersNeededForFinalize(self, None)
 
         options = check_or_create_options(EventOptions, options, "Event options")

cuda_core/cuda/core/experimental/_launcher.py

@@ -45,9 +45,6 @@ class LaunchConfig:
         Group of threads (Thread Block) that will execute on the same
         streaming multiprocessor (SM). Threads within a thread blocks have
         access to shared memory and can be explicitly synchronized.
-    stream : :obj:`~_stream.Stream`
-        The stream establishing the stream ordering semantic of a
-        launch.
     shmem_size : int, optional
         Dynamic shared-memory size per thread block in bytes.
         (Default to size 0)
@@ -58,7 +55,6 @@ class LaunchConfig:
     grid: Union[tuple, int] = None
     cluster: Union[tuple, int] = None
     block: Union[tuple, int] = None
-    stream: Stream = None
     shmem_size: Optional[int] = None
 
     def __post_init__(self):
@@ -72,12 +68,6 @@ def __post_init__(self):
             if Device().compute_capability < (9, 0):
                 raise CUDAError("thread block clusters are not supported on devices with compute capability < 9.0")
             self.cluster = self._cast_to_3_tuple(self.cluster)
-        # we handle "stream=None" in the launch API
-        if self.stream is not None and not isinstance(self.stream, Stream):
-            try:
-                self.stream = Stream._init(self.stream)
-            except Exception as e:
-                raise ValueError("stream must either be a Stream object or support __cuda_stream__") from e
         if self.shmem_size is None:
             self.shmem_size = 0
 
@@ -105,27 +95,35 @@ def _cast_to_3_tuple(self, cfg):
             raise ValueError
 
 
-def launch(kernel, config, *kernel_args):
+def launch(stream, config, kernel, *kernel_args):
     """Launches a :obj:`~_module.Kernel`
     object with launch-time configuration.
 
     Parameters
     ----------
-    kernel : :obj:`~_module.Kernel`
-        Kernel to launch.
+    stream : :obj:`~_stream.Stream`
+        The stream establishing the stream ordering semantic of a
+        launch.
     config : :obj:`~_launcher.LaunchConfig`
         Launch configurations inline with options provided by
         :obj:`~_launcher.LaunchConfig` dataclass.
+    kernel : :obj:`~_module.Kernel`
+        Kernel to launch.
     *kernel_args : Any
         Variable length argument list that is provided to the
         launching kernel.
 
     """
+    if stream is None:
+        raise ValueError("stream cannot be None, stream must either be a Stream object or support __cuda_stream__")
+    if not isinstance(stream, Stream):
+        try:
+            stream = Stream._init(stream)
+        except Exception as e:
+            raise ValueError("stream must either be a Stream object or support __cuda_stream__") from e
     if not isinstance(kernel, Kernel):
         raise ValueError
     config = check_or_create_options(LaunchConfig, config, "launch config")
-    if config.stream is None:
-        raise CUDAError("stream cannot be None")
 
     # TODO: can we ensure kernel_args is valid/safe to use here?
     # TODO: merge with HelperKernelParams?
@@ -141,7 +139,7 @@ def launch(kernel, config, *kernel_args):
         drv_cfg = driver.CUlaunchConfig()
         drv_cfg.gridDimX, drv_cfg.gridDimY, drv_cfg.gridDimZ = config.grid
         drv_cfg.blockDimX, drv_cfg.blockDimY, drv_cfg.blockDimZ = config.block
-        drv_cfg.hStream = config.stream.handle
+        drv_cfg.hStream = stream.handle
         drv_cfg.sharedMemBytes = config.shmem_size
         attrs = []  # TODO: support more attributes
         if config.cluster:
@@ -157,6 +155,6 @@ def launch(kernel, config, *kernel_args):
         # TODO: check if config has any unsupported attrs
         handle_return(
             driver.cuLaunchKernel(
-                int(kernel._handle), *config.grid, *config.block, config.shmem_size, config.stream._handle, args_ptr, 0
+                int(kernel._handle), *config.grid, *config.block, config.shmem_size, stream._handle, args_ptr, 0
             )
         )