From 902657ceb030f4f41f06048c5dc9b220dc1b83e8 Mon Sep 17 00:00:00 2001 From: Ashwin Srinath Date: Wed, 19 Jan 2022 15:34:34 -0500 Subject: [PATCH 1/5] Replace the use of RMM's CUDA Python bindings with those from CUDA-Python --- python/raft/common/cuda.pxd | 24 +++++----------------- python/raft/common/cuda.pyx | 38 ++++++++++++++++++++--------------- python/raft/common/handle.pxd | 3 +-- python/raft/common/handle.pyx | 9 +++++---- 4 files changed, 33 insertions(+), 41 deletions(-) diff --git a/python/raft/common/cuda.pxd b/python/raft/common/cuda.pxd index e407213f44..0459cb96af 100644 --- a/python/raft/common/cuda.pxd +++ b/python/raft/common/cuda.pxd @@ -1,5 +1,5 @@ # -# Copyright (c) 2019, NVIDIA CORPORATION. +# Copyright (c) 2019-2022, NVIDIA CORPORATION. # # Licensed under the Apache License, Version 2.0 (the "License"); # you may not use this file except in compliance with the License. @@ -14,23 +14,9 @@ # limitations under the License. # -# cython: profile=False -# distutils: language = c++ -# cython: embedsignature = True -# cython: language_level = 3 +from cuda.ccudart cimport cudaStream_t +cdef class Stream: + cdef cudaStream_t s -# Populate this with more typedef's (eg: events) as and when needed -cdef extern from * nogil: - ctypedef void* _Stream "cudaStream_t" - ctypedef int _Error "cudaError_t" - - -# Populate this with more runtime api method declarations as and when needed -cdef extern from "cuda_runtime_api.h" nogil: - _Error cudaStreamCreate(_Stream* s) - _Error cudaStreamDestroy(_Stream s) - _Error cudaStreamSynchronize(_Stream s) - _Error cudaGetLastError() - const char* cudaGetErrorString(_Error e) - const char* cudaGetErrorName(_Error e) + cdef cudaStream_t getStream(self) diff --git a/python/raft/common/cuda.pyx b/python/raft/common/cuda.pyx index 0b97eeba67..d00fd88af0 100644 --- a/python/raft/common/cuda.pyx +++ b/python/raft/common/cuda.pyx @@ -1,5 +1,5 @@ # -# Copyright (c) 2020-2021, NVIDIA CORPORATION. +# Copyright (c) 2020-2022, NVIDIA CORPORATION. # # Licensed under the Apache License, Version 2.0 (the "License"); # you may not use this file except in compliance with the License. @@ -19,10 +19,21 @@ # cython: embedsignature = True # cython: language_level = 3 +from cuda.ccudart cimport ( + cudaStream_t, + cudaError_t, + cudaSuccess, + cudaStreamCreate, + cudaStreamDestroy, + cudaStreamSynchronize, + cudaGetLastError, + cudaGetErrorString, + cudaGetErrorName +) class CudaRuntimeError(RuntimeError): def __init__(self, extraMsg=None): - cdef _Error e = cudaGetLastError() + cdef cudaError_t e = cudaGetLastError() cdef bytes errMsg = cudaGetErrorString(e) cdef bytes errName = cudaGetErrorName(e) msg = "Error! %s reason='%s'" % (errName.decode(), errMsg.decode()) @@ -47,26 +58,22 @@ cdef class Stream: """ # NOTE: - # If we store _Stream directly, this always leads to the following error: - # "Cannot convert Python object to '_Stream'" + # If we store cudaStream_t directly, this always leads to the following error: + # "Cannot convert Python object to 'cudaStream_t'" # I was unable to find a good solution to this in reasonable time. Also, # since cudaStream_t is a pointer anyways, storing it as an integer should # be just fine (although, that certainly is ugly and hacky!). - cdef size_t s def __cinit__(self): - if self.s != 0: - return - cdef _Stream stream - cdef _Error e = cudaStreamCreate(&stream) - if e != 0: + cdef cudaStream_t stream + cdef cudaError_t e = cudaStreamCreate(&stream) + if e != cudaSuccess: raise CudaRuntimeError("Stream create") - self.s = stream + self.s = stream def __dealloc__(self): self.sync() - cdef _Stream stream = <_Stream>self.s - cdef _Error e = cudaStreamDestroy(stream) + cdef cudaError_t e = cudaStreamDestroy(self.s) if e != 0: raise CudaRuntimeError("Stream destroy") @@ -76,10 +83,9 @@ cdef class Stream: could raise exception due to issues with previous asynchronous launches """ - cdef _Stream stream = <_Stream>self.s - cdef _Error e = cudaStreamSynchronize(stream) + cdef cudaError_t e = cudaStreamSynchronize(self.s) if e != 0: raise CudaRuntimeError("Stream sync") - def getStream(self): + cdef cudaStream_t getStream(self): return self.s diff --git a/python/raft/common/handle.pxd b/python/raft/common/handle.pxd index d2ae0a401d..8415b7e3d7 100644 --- a/python/raft/common/handle.pxd +++ b/python/raft/common/handle.pxd @@ -1,5 +1,5 @@ # -# Copyright (c) 2020, NVIDIA CORPORATION. +# Copyright (c) 2020-2022, NVIDIA CORPORATION. # # Licensed under the Apache License, Version 2.0 (the "License"); # you may not use this file except in compliance with the License. @@ -21,7 +21,6 @@ from libcpp.memory cimport shared_ptr -from .cuda cimport _Stream from rmm._lib.cuda_stream_view cimport cuda_stream_view from rmm._lib.cuda_stream_pool cimport cuda_stream_pool from libcpp.memory cimport shared_ptr diff --git a/python/raft/common/handle.pyx b/python/raft/common/handle.pyx index 1accf9e679..c4faabb610 100644 --- a/python/raft/common/handle.pyx +++ b/python/raft/common/handle.pyx @@ -1,5 +1,5 @@ # -# Copyright (c) 2020, NVIDIA CORPORATION. +# Copyright (c) 2020-2022, NVIDIA CORPORATION. # # Licensed under the Apache License, Version 2.0 (the "License"); # you may not use this file except in compliance with the License. @@ -24,9 +24,10 @@ from libcpp.memory cimport shared_ptr from rmm._lib.cuda_stream_view cimport cuda_stream_per_thread from rmm._lib.cuda_stream_view cimport cuda_stream_view -from .cuda cimport _Stream, _Error, cudaStreamSynchronize +from .cuda cimport Stream from .cuda import CudaRuntimeError + cdef class Handle: """ Handle is a lightweight python wrapper around the corresponding C++ class @@ -51,7 +52,7 @@ cdef class Handle: del handle # optional! """ - def __cinit__(self, stream=None, n_streams=0): + def __cinit__(self, stream: Stream=None, n_streams=0): self.n_streams = n_streams if n_streams > 0: self.stream_pool.reset(new cuda_stream_pool(n_streams)) @@ -64,7 +65,7 @@ cdef class Handle: self.stream_pool)) else: # this constructor constructs a handle on user stream - c_stream = cuda_stream_view(<_Stream> stream.getStream()) + c_stream = cuda_stream_view(stream.getStream()) self.c_obj.reset(new handle_t(c_stream, self.stream_pool)) From d7f00745fcc34bd15819998683e7d703450abd9d Mon Sep 17 00:00:00 2001 From: Ashwin Srinath Date: Wed, 19 Jan 2022 15:44:07 -0500 Subject: [PATCH 2/5] Add cuda-python to dev env --- conda/environments/raft_dev_cuda11.5.yml | 1 + 1 file changed, 1 insertion(+) diff --git a/conda/environments/raft_dev_cuda11.5.yml b/conda/environments/raft_dev_cuda11.5.yml index 152f3a8db5..c6d9f3fbf5 100644 --- a/conda/environments/raft_dev_cuda11.5.yml +++ b/conda/environments/raft_dev_cuda11.5.yml @@ -6,6 +6,7 @@ channels: - conda-forge dependencies: - cudatoolkit=11.5 +- cuda-python >=11.5,<12.0 - clang=11.1.0 - clang-tools=11.1.0 - rapids-build-env=22.02.* From c7666ad76d8968a7d4c7213554b79dbfdc561ed8 Mon Sep 17 00:00:00 2001 From: Ashwin Srinath Date: Wed, 19 Jan 2022 15:51:16 -0500 Subject: [PATCH 3/5] Remove older comment about cudaStream_t attribute --- python/raft/common/cuda.pyx | 8 -------- 1 file changed, 8 deletions(-) diff --git a/python/raft/common/cuda.pyx b/python/raft/common/cuda.pyx index d00fd88af0..8d2f6d8c69 100644 --- a/python/raft/common/cuda.pyx +++ b/python/raft/common/cuda.pyx @@ -56,14 +56,6 @@ cdef class Stream: stream.sync() del stream # optional! """ - - # NOTE: - # If we store cudaStream_t directly, this always leads to the following error: - # "Cannot convert Python object to 'cudaStream_t'" - # I was unable to find a good solution to this in reasonable time. Also, - # since cudaStream_t is a pointer anyways, storing it as an integer should - # be just fine (although, that certainly is ugly and hacky!). - def __cinit__(self): cdef cudaStream_t stream cdef cudaError_t e = cudaStreamCreate(&stream) From 8446e2f022b22235f15d524c53e73011d4486b06 Mon Sep 17 00:00:00 2001 From: Ashwin Srinath Date: Wed, 19 Jan 2022 16:31:58 -0500 Subject: [PATCH 4/5] Style --- python/raft/common/cuda.pyx | 3 ++- python/raft/common/handle.pyx | 2 +- 2 files changed, 3 insertions(+), 2 deletions(-) diff --git a/python/raft/common/cuda.pyx b/python/raft/common/cuda.pyx index 8d2f6d8c69..b5be977bcc 100644 --- a/python/raft/common/cuda.pyx +++ b/python/raft/common/cuda.pyx @@ -19,7 +19,7 @@ # cython: embedsignature = True # cython: language_level = 3 -from cuda.ccudart cimport ( +from cuda.ccudart cimport( cudaStream_t, cudaError_t, cudaSuccess, @@ -31,6 +31,7 @@ from cuda.ccudart cimport ( cudaGetErrorName ) + class CudaRuntimeError(RuntimeError): def __init__(self, extraMsg=None): cdef cudaError_t e = cudaGetLastError() diff --git a/python/raft/common/handle.pyx b/python/raft/common/handle.pyx index c4faabb610..661c5b5f23 100644 --- a/python/raft/common/handle.pyx +++ b/python/raft/common/handle.pyx @@ -52,7 +52,7 @@ cdef class Handle: del handle # optional! """ - def __cinit__(self, stream: Stream=None, n_streams=0): + def __cinit__(self, stream: Stream = None, n_streams=0): self.n_streams = n_streams if n_streams > 0: self.stream_pool.reset(new cuda_stream_pool(n_streams)) From 4de51ed3b0e37c6ca868d68399d3e0301447d838 Mon Sep 17 00:00:00 2001 From: Ashwin Srinath Date: Wed, 19 Jan 2022 16:33:01 -0500 Subject: [PATCH 5/5] 0 -> cudaSuccess --- python/raft/common/cuda.pyx | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/python/raft/common/cuda.pyx b/python/raft/common/cuda.pyx index b5be977bcc..c3c90936aa 100644 --- a/python/raft/common/cuda.pyx +++ b/python/raft/common/cuda.pyx @@ -67,7 +67,7 @@ cdef class Stream: def __dealloc__(self): self.sync() cdef cudaError_t e = cudaStreamDestroy(self.s) - if e != 0: + if e != cudaSuccess: raise CudaRuntimeError("Stream destroy") def sync(self): @@ -77,7 +77,7 @@ cdef class Stream: launches """ cdef cudaError_t e = cudaStreamSynchronize(self.s) - if e != 0: + if e != cudaSuccess: raise CudaRuntimeError("Stream sync") cdef cudaStream_t getStream(self):