Repository: cupy/cupy
Pull request: Add appropriate pointer types in Cython to the contribute guide
Followup of #1952. Fixes #1913 further. 

This PR is encouraged in https://github.com/cupy/cupy/issues/1913#issuecomment-529298038. The main change is that most of the types in `cupy.cuda.driver` (such as `Function` and `Stream`) are pointers, and thus should be represented by `intptr_t`, see https://github.com/cupy/cupy/issues/1913#issuecomment-529262513. Once it's fixed, the changes should propagate outward to other modules.

***Challenges***: 
1. ~~In `cupy.cuda.memory` it seems streams are used in the memory mappings. I'm trying to figure out which `size_t` is meant to present a stream pointer there...~~
2. It's likely this PR would still not be enough to clear up everything. I expect an incremental approach is more appropriate.

If Jenkins can be kicked off to test the current PR, it'd be great!

Thanks.

Problem statement:
Use correct integer types in Cython
Currently incorrect integer types are used in Cython implementation.
It's less portable and perhaps more importantly, confusing for developers.

For example:
* Pointers should be `intptr_t`.
* Memory sizes should be `size_t`.
* Memory offsets should be `ptrdiff_t`.

Hints from issue discussion:
Hi @ninoshi, when I was working on PRs that would touch the driver wrappers, I noticed that in `driver.pxd`/`driver.pyx` there are many types that should have been changed from `size_t` to `intptr_t` but were left untouched in #1952, which made me thought that using `size_t` was the deliberately chosen convention for drivers, but looking at this issue I doubt it's the case. 

Specifically, I am referring to the driver functions involving the following types:
https://github.com/cupy/cupy/blob/c670b6f6f1282665c9985180c10f498d8591f909/cupy/cuda/driver.pxd#L12-L18
all of which are just pointers but were represented using `size_t`. Is this an oversight? If so I can make a PR to fix them.
I didn't deliberately leave them. I think I simply overlooked, so please feel free to fix them!

Commit message: improved language in contribution guide

Co-authored-by: Sean Farley <sean@farley.io>

Patch to review:
diff --git a/cupy/core/raw.pyx b/cupy/core/raw.pyx
index 1a94562fd90..7d0ea226838 100644
--- a/cupy/core/raw.pyx
+++ b/cupy/core/raw.pyx
@@ -256,6 +256,6 @@ cdef class RawModule:
             name (str): Name of the texture reference.
 
         Returns:
-            size_t: A ``CUtexref`` handle, to be passed to :class:`~cupy.cuda.texture.TextureReference`.
+            intptr_t: A ``CUtexref`` handle, to be passed to :class:`~cupy.cuda.texture.TextureReference`.
         '''  # noqa
         return self.module.get_texref(name)
diff --git a/cupy/cuda/cufft.pxd b/cupy/cuda/cufft.pxd
index b30c903d848..9168d2eaaa0 100644
--- a/cupy/cuda/cufft.pxd
+++ b/cupy/cuda/cufft.pxd
@@ -2,7 +2,7 @@ cdef extern from *:
     ctypedef float Float 'cufftReal'
     ctypedef double Double 'cufftDoubleReal'
     ctypedef int Result 'cufftResult_t'
-    ctypedef size_t Handle 'cufftHandle'
+    ctypedef int Handle 'cufftHandle'
     ctypedef int Type 'cufftType_t'
 
 
diff --git a/cupy/cuda/cufft.pyx b/cupy/cuda/cufft.pyx
index bcdacdc7854..08ed8ec53ce 100644
--- a/cupy/cuda/cufft.pyx
+++ b/cupy/cuda/cufft.pyx
@@ -1,4 +1,5 @@
 cimport cython  # NOQA
+from libc.stdint cimport intptr_t
 import numpy
 import threading
 
@@ -145,17 +146,17 @@ class Plan1d(object):
             result = cufftSetStream(plan, <driver.Stream>stream)
         check_result(result)
         if self.fft_type == CUFFT_C2C:
-            execC2C(self.plan, a.data, out.data, direction)
+            execC2C(plan, a.data.ptr, out.data.ptr, direction)
         elif self.fft_type == CUFFT_R2C:
-            execR2C(self.plan, a.data, out.data)
+            execR2C(plan, a.data.ptr, out.data.ptr)
         elif self.fft_type == CUFFT_C2R:
-            execC2R(self.plan, a.data, out.data)
+            execC2R(plan, a.data.ptr, out.data.ptr)
         elif self.fft_type == CUFFT_Z2Z:
-            execZ2Z(self.plan, a.data, out.data, direction)
+            execZ2Z(plan, a.data.ptr, out.data.ptr, direction)
         elif self.fft_type == CUFFT_D2Z:
-            execD2Z(self.plan, a.data, out.data)
+            execD2Z(plan, a.data.ptr, out.data.ptr)
         else:
-            execZ2D(self.plan, a.data, out.data)
+            execZ2D(plan, a.data.ptr, out.data.ptr)
 
     def _output_dtype_and_shape(self, a):
         shape = list(a.shape)
@@ -282,9 +283,9 @@ class PlanNd(object):
             result = cufftSetStream(plan, <driver.Stream>stream)
         check_result(result)
         if self.fft_type == CUFFT_C2C:
-            execC2C(self.plan, a.data, out.data, direction)
+            execC2C(plan, a.data.ptr, out.data.ptr, direction)
         elif self.fft_type == CUFFT_Z2Z:
-            execZ2Z(self.plan, a.data, out.data, direction)
+            execZ2Z(plan, a.data.ptr, out.data.ptr, direction)
         else:
             raise NotImplementedError('only C2C and Z2Z implemented')
 
@@ -309,39 +310,39 @@ class PlanNd(object):
             raise ValueError('output shape mismatch')
 
 
-cpdef execC2C(size_t plan, size_t idata, size_t odata, int direction):
+cpdef execC2C(Handle plan, intptr_t idata, intptr_t odata, int direction):
     with nogil:
         result = cufftExecC2C(plan, <Complex*>idata, <Complex*>odata,
                               direction)
     check_result(result)
 
 
-cpdef execR2C(size_t plan, size_t idata, size_t odata):
+cpdef execR2C(Handle plan, intptr_t idata, intptr_t odata):
     with nogil:
         result = cufftExecR2C(plan, <Float*>idata, <Complex*>odata)
     check_result(result)
 
 
-cpdef execC2R(size_t plan, size_t idata, size_t odata):
+cpdef execC2R(Handle plan, intptr_t idata, intptr_t odata):
     with nogil:
         result = cufftExecC2R(plan, <Complex*>idata, <Float*>odata)
     check_result(result)
 
 
-cpdef execZ2Z(size_t plan, size_t idata, size_t odata, int direction):
+cpdef execZ2Z(Handle plan, intptr_t idata, intptr_t odata, int direction):
     with nogil:
         result = cufftExecZ2Z(plan, <DoubleComplex*>idata,
                               <DoubleComplex*>odata, direction)
     check_result(result)
 
 
-cpdef execD2Z(size_t plan, size_t idata, size_t odata):
+cpdef execD2Z(Handle plan, intptr_t idata, intptr_t odata):
     with nogil:
         result = cufftExecD2Z(plan, <Double*>idata, <DoubleComplex*>odata)
     check_result(result)
 
 
-cpdef execZ2D(size_t plan, size_t idata, size_t odata):
+cpdef execZ2D(Handle plan, intptr_t idata, intptr_t odata):
     with nogil:
         result = cufftExecZ2D(plan, <DoubleComplex*>idata, <Double*>odata)
     check_result(result)
diff --git a/cupy/cuda/driver.pxd b/cupy/cuda/driver.pxd
index e05891d5f58..f8f336404cd 100644
--- a/cupy/cuda/driver.pxd
+++ b/cupy/cuda/driver.pxd
@@ -95,30 +95,30 @@ cpdef devicePrimaryCtxRelease(Device dev)
 # Context management
 ###############################################################################
 
-cpdef size_t ctxGetCurrent() except? 0
-cpdef ctxSetCurrent(size_t ctx)
-cpdef size_t ctxCreate(Device dev) except? 0
-cpdef ctxDestroy(size_t ctx)
+cpdef intptr_t ctxGetCurrent() except? 0
+cpdef ctxSetCurrent(intptr_t ctx)
+cpdef intptr_t ctxCreate(Device dev) except? 0
+cpdef ctxDestroy(intptr_t ctx)
 
 ###############################################################################
 # Module load and kernel execution
 ###############################################################################
 
-cpdef size_t linkCreate() except? 0
-cpdef linkAddData(size_t state, int input_type, bytes data, unicode name)
-cpdef bytes linkComplete(size_t state)
-cpdef linkDestroy(size_t state)
-cpdef size_t moduleLoad(str filename) except? 0
-cpdef size_t moduleLoadData(bytes image) except? 0
-cpdef moduleUnload(size_t module)
-cpdef size_t moduleGetFunction(size_t module, str funcname) except? 0
-cpdef size_t moduleGetGlobal(size_t module, str varname) except? 0
-cpdef size_t moduleGetTexRef(size_t module, str texrefname) except? 0
+cpdef intptr_t linkCreate() except? 0
+cpdef linkAddData(intptr_t state, int input_type, bytes data, unicode name)
+cpdef bytes linkComplete(intptr_t state)
+cpdef linkDestroy(intptr_t state)
+cpdef intptr_t moduleLoad(str filename) except? 0
+cpdef intptr_t moduleLoadData(bytes image) except? 0
+cpdef moduleUnload(intptr_t module)
+cpdef intptr_t moduleGetFunction(intptr_t module, str funcname) except? 0
+cpdef intptr_t moduleGetGlobal(intptr_t module, str varname) except? 0
+cpdef intptr_t moduleGetTexRef(intptr_t module, str texrefname) except? 0
 cpdef launchKernel(
     intptr_t f, unsigned int grid_dim_x, unsigned int grid_dim_y,
     unsigned int grid_dim_z, unsigned int block_dim_x,
     unsigned int block_dim_y, unsigned int block_dim_z,
-    unsigned int shared_mem_bytes, size_t stream, intptr_t kernel_params,
+    unsigned int shared_mem_bytes, intptr_t stream, intptr_t kernel_params,
     intptr_t extra)
 
 ###############################################################################
@@ -132,16 +132,17 @@ cpdef funcSetAttribute(intptr_t func, int attribute, int value)
 # Texture reference
 ###############################################################################
 
-cpdef size_t texRefSetAddress(size_t texref, size_t dptr, size_t nbytes)
-cpdef texRefSetAddress2D(size_t texref, size_t desc, size_t dptr, size_t Pitch)
-cpdef texRefSetAddressMode(size_t texref, int dim, int am)
-cpdef texRefSetArray(size_t texref, size_t array)
-cpdef texRefSetBorderColor(size_t texref, pBorderColor)
-cpdef texRefSetFilterMode(size_t texref, int fm)
-cpdef texRefSetFlags(size_t texref, unsigned int Flags)
-cpdef texRefSetFormat(size_t texref, int fmt, int NumPackedComponents)
-cpdef texRefSetMaxAnisotropy(size_t texref, unsigned int maxAniso)
-cpdef paramSetTexRef(size_t func, size_t texref)
+cpdef size_t texRefSetAddress(intptr_t texref, intptr_t dptr, size_t nbytes)
+cpdef texRefSetAddress2D(intptr_t texref, intptr_t desc, intptr_t dptr,
+                         size_t Pitch)
+cpdef texRefSetAddressMode(intptr_t texref, int dim, int am)
+cpdef texRefSetArray(intptr_t texref, intptr_t array)
+cpdef texRefSetBorderColor(intptr_t texref, pBorderColor)
+cpdef texRefSetFilterMode(intptr_t texref, int fm)
+cpdef texRefSetFlags(intptr_t texref, unsigned int Flags)
+cpdef texRefSetFormat(intptr_t texref, int fmt, int NumPackedComponents)
+cpdef texRefSetMaxAnisotropy(intptr_t texref, unsigned int maxAniso)
+cpdef paramSetTexRef(intptr_t func, intptr_t texref)
 
 ###############################################################################
 # Occupancy
diff --git a/cupy/cuda/driver.pyx b/cupy/cuda/driver.pyx
index e9918a2c709..40249f5ebd1 100644
--- a/cupy/cuda/driver.pyx
+++ b/cupy/cuda/driver.pyx
@@ -145,27 +145,27 @@ cpdef devicePrimaryCtxRelease(Device dev):
 # Context management
 ###############################################################################
 
-cpdef size_t ctxGetCurrent() except? 0:
+cpdef intptr_t ctxGetCurrent() except? 0:
     cdef Context ctx
     with nogil:
         status = cuCtxGetCurrent(&ctx)
     check_status(status)
-    return <size_t>ctx
+    return <intptr_t>ctx
 
-cpdef ctxSetCurrent(size_t ctx):
+cpdef ctxSetCurrent(intptr_t ctx):
     with nogil:
         status = cuCtxSetCurrent(<Context>ctx)
     check_status(status)
 
-cpdef size_t ctxCreate(Device dev) except? 0:
+cpdef intptr_t ctxCreate(Device dev) except? 0:
     cdef Context ctx
     cdef unsigned int flags = 0
     with nogil:
         status = cuCtxCreate(&ctx, flags, dev)
     check_status(status)
-    return <size_t>ctx
+    return <intptr_t>ctx
 
-cpdef ctxDestroy(size_t ctx):
+cpdef ctxDestroy(intptr_t ctx):
     with nogil:
         status = cuCtxDestroy(<Context>ctx)
     check_status(status)
@@ -175,15 +175,15 @@ cpdef ctxDestroy(size_t ctx):
 # Module load and kernel execution
 ###############################################################################
 
-cpdef size_t linkCreate() except? 0:
+cpdef intptr_t linkCreate() except? 0:
     cpdef LinkState state
     with nogil:
         status = cuLinkCreate(0, <CUjit_option*>0, <void**>0, &state)
     check_status(status)
-    return <size_t>state
+    return <intptr_t>state
 
 
-cpdef linkAddData(size_t state, int input_type, bytes data, unicode name):
+cpdef linkAddData(intptr_t state, int input_type, bytes data, unicode name):
     cdef const char* data_ptr = data
     cdef size_t data_size = len(data) + 1
     cdef bytes b_name = name.encode()
@@ -195,7 +195,7 @@ cpdef linkAddData(size_t state, int input_type, bytes data, unicode name):
     check_status(status)
 
 
-cpdef bytes linkComplete(size_t state):
+cpdef bytes linkComplete(intptr_t state):
     cdef void* cubinOut
     cdef size_t sizeOut
     with nogil:
@@ -204,48 +204,48 @@ cpdef bytes linkComplete(size_t state):
     return bytes((<char*>cubinOut)[:sizeOut])
 
 
-cpdef linkDestroy(size_t state):
+cpdef linkDestroy(intptr_t state):
     with nogil:
         status = cuLinkDestroy(<LinkState>state)
     check_status(status)
 
 
-cpdef size_t moduleLoad(str filename) except? 0:
+cpdef intptr_t moduleLoad(str filename) except? 0:
     cdef Module module
     cdef bytes b_filename = filename.encode()
     cdef char* b_filename_ptr = b_filename
     with nogil:
         status = cuModuleLoad(&module, b_filename_ptr)
     check_status(status)
-    return <size_t>module
+    return <intptr_t>module
 
 
-cpdef size_t moduleLoadData(bytes image) except? 0:
+cpdef intptr_t moduleLoadData(bytes image) except? 0:
     cdef Module module
     cdef char* image_ptr = image
     with nogil:
         status = cuModuleLoadData(&module, image_ptr)
     check_status(status)
-    return <size_t>module
+    return <intptr_t>module
 
 
-cpdef moduleUnload(size_t module):
+cpdef moduleUnload(intptr_t module):
     with nogil:
         status = cuModuleUnload(<Module>module)
     check_status(status)
 
 
-cpdef size_t moduleGetFunction(size_t module, str funcname) except? 0:
+cpdef intptr_t moduleGetFunction(intptr_t module, str funcname) except? 0:
     cdef Function func
     cdef bytes b_funcname = funcname.encode()
     cdef char* b_funcname_ptr = b_funcname
     with nogil:
         status = cuModuleGetFunction(&func, <Module>module, b_funcname_ptr)
     check_status(status)
-    return <size_t>func
+    return <intptr_t>func
 
 
-cpdef size_t moduleGetGlobal(size_t module, str varname) except? 0:
+cpdef intptr_t moduleGetGlobal(intptr_t module, str varname) except? 0:
     cdef Deviceptr var
     cdef size_t size
     cdef bytes b_varname = varname.encode()
@@ -253,24 +253,24 @@ cpdef size_t moduleGetGlobal(size_t module, str varname) except? 0:
     with nogil:
         status = cuModuleGetGlobal(&var, &size, <Module>module, b_varname_ptr)
     check_status(status)
-    return <size_t>var
+    return <intptr_t>var
 
 
-cpdef size_t moduleGetTexRef(size_t module, str texrefname) except? 0:
+cpdef intptr_t moduleGetTexRef(intptr_t module, str texrefname) except? 0:
     cdef TexRef texref
     cdef bytes b_refname = texrefname.encode()
     cdef char* b_refname_ptr = b_refname
     with nogil:
         status = cuModuleGetTexRef(&texref, <Module>module, b_refname_ptr)
     check_status(status)
-    return <size_t>texref
+    return <intptr_t>texref
 
 
 cpdef launchKernel(
         intptr_t f, unsigned int grid_dim_x, unsigned int grid_dim_y,
         unsigned int grid_dim_z, unsigned int block_dim_x,
         unsigned int block_dim_y, unsigned int block_dim_z,
-        unsigned int shared_mem_bytes, size_t stream, intptr_t kernel_params,
+        unsigned int shared_mem_bytes, intptr_t stream, intptr_t kernel_params,
         intptr_t extra):
     with nogil:
         status = cuLaunchKernel(
@@ -311,7 +311,7 @@ cpdef funcSetAttribute(intptr_t f, int attribute, int value):
 # Texture reference
 ###############################################################################
 
-cpdef size_t texRefSetAddress(size_t texref, size_t dptr, size_t nbytes):
+cpdef size_t texRefSetAddress(intptr_t texref, intptr_t dptr, size_t nbytes):
     cdef size_t ByteOffset
     with nogil:
         status = cuTexRefSetAddress(&ByteOffset, <TexRef>texref,
@@ -320,7 +320,7 @@ cpdef size_t texRefSetAddress(size_t texref, size_t dptr, size_t nbytes):
     return ByteOffset
 
 
-cpdef texRefSetAddress2D(size_t texref, size_t desc, size_t dptr,
+cpdef texRefSetAddress2D(intptr_t texref, intptr_t desc, intptr_t dptr,
                          size_t Pitch):
     with nogil:
         status = cuTexRefSetAddress2D(<TexRef>texref, <const Array_desc*>desc,
@@ -328,20 +328,20 @@ cpdef texRefSetAddress2D(size_t texref, size_t desc, size_t dptr,
     check_status(status)
 
 
-cpdef texRefSetAddressMode(size_t texref, int dim, int am):
+cpdef texRefSetAddressMode(intptr_t texref, int dim, int am):
     with nogil:
         status = cuTexRefSetAddressMode(<TexRef>texref, dim, <Address_mode>am)
     check_status(status)
 
 
-cpdef texRefSetArray(size_t texref, size_t array):
+cpdef texRefSetArray(intptr_t texref, intptr_t array):
     with nogil:
         status = cuTexRefSetArray(<TexRef>texref, <Array>array,
                                   CU_TRSA_OVERRIDE_FORMAT)
     check_status(status)
 
 
-cpdef texRefSetBorderColor(size_t texref, pBorderColor):
+cpdef texRefSetBorderColor(intptr_t texref, pBorderColor):
     cdef vector.vector[float] colors
     for i in range(4):
         colors.push_back(pBorderColor[i])
@@ -350,32 +350,32 @@ cpdef texRefSetBorderColor(size_t texref, pBorderColor):
     check_status(status)
 
 
-cpdef texRefSetFilterMode(size_t texref, int fm):
+cpdef texRefSetFilterMode(intptr_t texref, int fm):
     with nogil:
         status = cuTexRefSetFilterMode(<TexRef>texref, <Filter_mode>fm)
     check_status(status)
 
 
-cpdef texRefSetFlags(size_t texref, unsigned int Flags):
+cpdef texRefSetFlags(intptr_t texref, unsigned int Flags):
     with nogil:
         status = cuTexRefSetFlags(<TexRef>texref, Flags)
     check_status(status)
 
 
-cpdef texRefSetFormat(size_t texref, int fmt, int NumPackedComponents):
+cpdef texRefSetFormat(intptr_t texref, int fmt, int NumPackedComponents):
     with nogil:
         status = cuTexRefSetFormat(<TexRef>texref, <Array_format>fmt,
                                    NumPackedComponents)
     check_status(status)
 
 
-cpdef texRefSetMaxAnisotropy(size_t texref, unsigned int maxAniso):
+cpdef texRefSetMaxAnisotropy(intptr_t texref, unsigned int maxAniso):
     with nogil:
         status = cuTexRefSetMaxAnisotropy(<TexRef>texref, maxAniso)
     check_status(status)
 
 
-cpdef paramSetTexRef(size_t func, size_t texref):
+cpdef paramSetTexRef(intptr_t func, intptr_t texref):
     with nogil:
         status = cuParamSetTexRef(<Function>func, CU_PARAM_TR_DEFAULT,
                                   <TexRef>texref)
diff --git a/cupy/cuda/memory.pyx b/cupy/cuda/memory.pyx
index 49278829680..3463a6eb8e5 100644
--- a/cupy/cuda/memory.pyx
+++ b/cupy/cuda/memory.pyx
@@ -204,7 +204,7 @@ cdef class _Chunk:
         mem (~cupy.cuda.Memory): The device memory buffer.
         offset (int): An offset bytes from the head of the buffer.
         size (int): Chunk size in bytes.
-        stream_ptr (size_t): Raw stream handle of cupy.cuda.Stream
+        stream_ptr (intptr_t): Raw stream handle of cupy.cuda.Stream
 
     Attributes:
         mem (Memory): The device memory buffer.
@@ -220,7 +220,7 @@ cdef class _Chunk:
         readonly BaseMemory mem
         readonly ptrdiff_t offset
         readonly size_t size
-        readonly size_t stream_ptr
+        readonly intptr_t stream_ptr
         public _Chunk prev
         public _Chunk next
 
@@ -230,7 +230,7 @@ cdef class _Chunk:
         self._init(mem, offset, size, stream_ptr)
 
     cdef _init(self, BaseMemory mem, ptrdiff_t offset,
-               size_t size, Py_ssize_t stream_ptr):
+               size_t size, intptr_t stream_ptr):
         assert mem.ptr != 0 or offset == 0
         self.mem = mem
         self.offset = offset
@@ -627,7 +627,8 @@ cdef class PooledMemory(BaseMemory):
 cdef int _index_compaction_threshold = 512
 
 
-cdef _compact_index(SingleDeviceMemoryPool pool, size_t stream_ptr, bint free):
+cdef _compact_index(SingleDeviceMemoryPool pool, intptr_t stream_ptr,
+                    bint free):
     # need self._free_lock
     cdef list arena, new_arena
     cdef set free_list, keep_list
@@ -665,7 +666,7 @@ cdef _compact_index(SingleDeviceMemoryPool pool, size_t stream_ptr, bint free):
 
 
 cdef object _get_chunk(SingleDeviceMemoryPool pool, size_t size,
-                       size_t stream_ptr):
+                       intptr_t stream_ptr):
     # need self._free_lock
     cdef set free_list
     cdef size_t i, index, length
@@ -856,13 +857,13 @@ cdef class SingleDeviceMemoryPool:
     cdef:
         object _allocator
 
-        # Map from memory pointer of the chunk (size_t) to the corresponding
+        # Map from memory pointer of the chunk (intptr_t) to the corresponding
         # Chunk object. All chunks currently allocated to the application from
         # this pool are stored.
         # `_in_use_lock` must be acquired to access it.
         dict _in_use
 
-        # Map from stream pointer (int) to its arena (list) for the stream.
+        # Map from stream pointer (intptr_t) to its arena (list) for the stream
         # `_free_lock` must be acquired to access it.
         dict _free
 
@@ -883,8 +884,8 @@ cdef class SingleDeviceMemoryPool:
 
         # Map from stream pointer to its arena index.
         # `_free_lock` must be acquired to access it.
-        map.map[size_t, vector.vector[size_t]] _index
-        map.map[size_t, vector.vector[int8_t]] _flag
+        map.map[intptr_t, vector.vector[size_t]] _index
+        map.map[intptr_t, vector.vector[int8_t]] _flag
 
     def __init__(self, allocator=_malloc):
         self._in_use = {}
@@ -898,7 +899,7 @@ cdef class SingleDeviceMemoryPool:
 
         self.set_limit(**(self._parse_limit_string()))
 
-    cpdef list _arena(self, size_t stream_ptr):
+    cpdef list _arena(self, intptr_t stream_ptr):
         """Returns appropriate arena (list of bins) of a given stream.
 
         All free chunks in the stream belong to one of the bin in the arena.
@@ -910,7 +911,7 @@ cdef class SingleDeviceMemoryPool:
             self._free[stream_ptr] = ret = []
         return ret
 
-    cdef inline vector.vector[size_t]* _arena_index(self, size_t stream_ptr):
+    cdef inline vector.vector[size_t]* _arena_index(self, intptr_t stream_ptr):
         """Returns appropriate arena sparse index of a given stream.
 
         Each element of the returned vector is an index value of the arena
@@ -923,7 +924,7 @@ cdef class SingleDeviceMemoryPool:
         """
         return &self._index[stream_ptr]
 
-    cdef vector.vector[int8_t]* _arena_flag(self, size_t stream_ptr):
+    cdef vector.vector[int8_t]* _arena_flag(self, intptr_t stream_ptr):
         """Returns appropriate arena used flag list of a given stream.
 
         Caller is responsible to acquire `_free_lock`.
@@ -1041,7 +1042,7 @@ cdef class SingleDeviceMemoryPool:
 
     cpdef free_all_blocks(self, stream=None):
         """Free all **non-split** chunks"""
-        cdef size_t stream_ptr
+        cdef intptr_t stream_ptr
 
         with LockAndNoGc(self._free_lock):
             # free blocks in all arenas
diff --git a/cupy/cuda/nccl.pyx b/cupy/cuda/nccl.pyx
index 14a99bdd3b4..58666250bf8 100644
--- a/cupy/cuda/nccl.pyx
+++ b/cupy/cuda/nccl.pyx
@@ -5,6 +5,7 @@ Wrapper for NCCL: Optimized primiteive for collective multi-GPU communication
 """
 cimport cython  # NOQA
 
+from libc.stdint cimport intptr_t
 from libcpp cimport vector
 
 from cupy.cuda cimport driver
@@ -366,8 +367,8 @@ cdef class NcclCommunicator:
         check_status(status)
         return ranks
 
-    def allReduce(self, size_t sendbuf, size_t recvbuf,
-                  size_t count, int datatype, int op, size_t stream):
+    def allReduce(self, intptr_t sendbuf, intptr_t recvbuf,
+                  size_t count, int datatype, int op, intptr_t stream):
         with nogil:
             status = _ncclAllReduce(<void*>sendbuf, <void*>recvbuf,
                                     count, <ncclDataType_t>datatype,
@@ -375,8 +376,8 @@ cdef class NcclCommunicator:
                                     <driver.Stream>stream)
         check_status(status)
 
-    def reduce(self, size_t sendbuf, size_t recvbuf,
-               size_t count, int datatype, int op, int root, size_t stream):
+    def reduce(self, intptr_t sendbuf, intptr_t recvbuf,
+               size_t count, int datatype, int op, int root, intptr_t stream):
         with nogil:
             status = _ncclReduce(<void*>sendbuf, <void*>recvbuf,
                                  count, <ncclDataType_t>datatype,
@@ -384,8 +385,8 @@ cdef class NcclCommunicator:
                                  <driver.Stream>stream)
         check_status(status)
 
-    def broadcast(self, size_t sendbuff, size_t recvbuff, int count,
-                  int datatype, int root, size_t stream):
+    def broadcast(self, intptr_t sendbuff, intptr_t recvbuff, int count,
+                  int datatype, int root, intptr_t stream):
         if NCCL_VERSION_CODE < 2200:
             # ncclBroadcast is not available in NCCL 2.1 or older.
             if self.rank_id() == root and sendbuff != recvbuff:
@@ -400,16 +401,16 @@ cdef class NcclCommunicator:
                                     self._comm, <driver.Stream>stream)
         check_status(status)
 
-    def bcast(self, size_t buff, int count, int datatype,
-              int root, size_t stream):
+    def bcast(self, intptr_t buff, int count, int datatype,
+              int root, intptr_t stream):
         with nogil:
             status = _ncclBcast(<void*>buff, count,
                                 <ncclDataType_t>datatype, root,
                                 self._comm, <driver.Stream>stream)
         check_status(status)
 
-    def reduceScatter(self, size_t sendbuf, size_t recvbuf,
-                      size_t recvcount, int datatype, int op, size_t stream):
+    def reduceScatter(self, intptr_t sendbuf, intptr_t recvbuf,
+                      size_t recvcount, int datatype, int op, intptr_t stream):
         with nogil:
             status = _ncclReduceScatter(<void*>sendbuf, <void*>recvbuf,
                                         recvcount, <ncclDataType_t>datatype,
@@ -417,8 +418,8 @@ cdef class NcclCommunicator:
                                         <driver.Stream>stream)
         check_status(status)
 
-    def allGather(self, size_t sendbuf, size_t recvbuf, size_t count,
-                  int datatype, size_t stream):
+    def allGather(self, intptr_t sendbuf, intptr_t recvbuf, size_t count,
+                  int datatype, intptr_t stream):
         with nogil:
             status = _ncclAllGather(<void*>sendbuf, <void*>recvbuf,
                                     count, <ncclDataType_t>datatype,
diff --git a/cupy/cuda/nvrtc.pxd b/cupy/cuda/nvrtc.pxd
index baaba6bfc70..31c7c2e4124 100644
--- a/cupy/cuda/nvrtc.pxd
+++ b/cupy/cuda/nvrtc.pxd
@@ -1,3 +1,5 @@
+from libc.stdint cimport intptr_t
+
 
 ###############################################################################
 # Types
@@ -17,9 +19,9 @@ cpdef tuple getVersion()
 # Program
 ###############################################################################
 
-cpdef size_t createProgram(unicode src, unicode name, headers,
-                           include_names) except? 0
-cpdef destroyProgram(size_t prog)
-cpdef compileProgram(size_t prog, options)
-cpdef unicode getPTX(size_t prog)
-cpdef unicode getProgramLog(size_t prog)
+cpdef intptr_t createProgram(unicode src, unicode name, headers,
+                             include_names) except? 0
+cpdef destroyProgram(intptr_t prog)
+cpdef compileProgram(intptr_t prog, options)
+cpdef unicode getPTX(intptr_t prog)
+cpdef unicode getProgramLog(intptr_t prog)
diff --git a/cupy/cuda/nvrtc.pyx b/cupy/cuda/nvrtc.pyx
index 5b4c180585d..2c40782a3c0 100644
--- a/cupy/cuda/nvrtc.pyx
+++ b/cupy/cuda/nvrtc.pyx
@@ -68,8 +68,8 @@ cpdef tuple getVersion():
 # Program
 ###############################################################################
 
-cpdef size_t createProgram(unicode src, unicode name, headers,
-                           include_names) except? 0:
+cpdef intptr_t createProgram(unicode src, unicode name, headers,
+                             include_names) except? 0:
     cdef Program prog
     cdef bytes b_src = src.encode()
     cdef const char* src_ptr = b_src
@@ -88,17 +88,17 @@ cpdef size_t createProgram(unicode src, unicode name, headers,
             &prog, src_ptr, name_ptr, num_headers, &(header_vec[0]),
             &(include_name_vec[0]))
     check_status(status)
-    return <size_t>prog
+    return <intptr_t>prog
 
 
-cpdef destroyProgram(size_t prog):
+cpdef destroyProgram(intptr_t prog):
     cdef Program p = <Program>prog
     with nogil:
         status = nvrtcDestroyProgram(&p)
     check_status(status)
 
 
-cpdef compileProgram(size_t prog, options):
+cpdef compileProgram(intptr_t prog, options):
     cdef int option_num = len(options)
     cdef vector.vector[const char*] option_vec
     cdef option_list = [opt.encode() for opt in options]
@@ -111,7 +111,7 @@ cpdef compileProgram(size_t prog, options):
     check_status(status)
 
 
-cpdef unicode getPTX(size_t prog):
+cpdef unicode getPTX(intptr_t prog):
     cdef size_t ptxSizeRet
     cdef bytes ptx
     cdef char* ptx_ptr
@@ -130,7 +130,7 @@ cpdef unicode getPTX(size_t prog):
     return ptx.decode('UTF-8')
 
 
-cpdef unicode getProgramLog(size_t prog):
+cpdef unicode getProgramLog(intptr_t prog):
     cdef size_t logSizeRet
     cdef bytes log
     cdef char* log_ptr
diff --git a/cupy/cuda/runtime.pxd b/cupy/cuda/runtime.pxd
index 49e12a3b5e4..22c89f053f4 100644
--- a/cupy/cuda/runtime.pxd
+++ b/cupy/cuda/runtime.pxd
@@ -8,8 +8,8 @@ from libc.stdint cimport intptr_t, uintmax_t
 cdef class PointerAttributes:
     cdef:
         public int device
-        public size_t devicePointer
-        public size_t hostPointer
+        public intptr_t devicePointer
+        public intptr_t hostPointer
         public int isManaged
         public int memoryType
 
@@ -328,35 +328,35 @@ cpdef freeArray(intptr_t ptr)
 cpdef memGetInfo()
 cpdef memcpy(intptr_t dst, intptr_t src, size_t size, int kind)
 cpdef memcpyAsync(intptr_t dst, intptr_t src, size_t size, int kind,
-                  size_t stream)
+                  intptr_t stream)
 cpdef memcpyPeer(intptr_t dst, int dstDevice, intptr_t src, int srcDevice,
                  size_t size)
 cpdef memcpyPeerAsync(intptr_t dst, int dstDevice,
                       intptr_t src, int srcDevice,
-                      size_t size, size_t stream)
+                      size_t size, intptr_t stream)
 cpdef memcpy2D(intptr_t dst, size_t dpitch, intptr_t src, size_t spitch,
                size_t width, size_t height, MemoryKind kind)
 cpdef memcpy2DAsync(intptr_t dst, size_t dpitch, intptr_t src, size_t spitch,
                     size_t width, size_t height, MemoryKind kind,
-                    size_t stream)
+                    intptr_t stream)
 cpdef memcpy2DFromArray(intptr_t dst, size_t dpitch, intptr_t src,
                         size_t wOffset, size_t hOffset, size_t width,
                         size_t height, int kind)
 cpdef memcpy2DFromArrayAsync(intptr_t dst, size_t dpitch, intptr_t src,
                              size_t wOffset, size_t hOffset, size_t width,
-                             size_t height, int kind, size_t stream)
+                             size_t height, int kind, intptr_t stream)
 cpdef memcpy2DToArray(intptr_t dst, size_t wOffset, size_t hOffset,
                       intptr_t src, size_t spitch, size_t width, size_t height,
                       int kind)
 cpdef memcpy2DToArrayAsync(intptr_t dst, size_t wOffset, size_t hOffset,
                            intptr_t src, size_t spitch, size_t width,
-                           size_t height, int kind, size_t stream)
+                           size_t height, int kind, intptr_t stream)
 cpdef memcpy3D(intptr_t Memcpy3DParmsPtr)
-cpdef memcpy3DAsync(intptr_t Memcpy3DParmsPtr, size_t stream)
+cpdef memcpy3DAsync(intptr_t Memcpy3DParmsPtr, intptr_t stream)
 cpdef memset(intptr_t ptr, int value, size_t size)
-cpdef memsetAsync(intptr_t ptr, int value, size_t size, size_t stream)
+cpdef memsetAsync(intptr_t ptr, int value, size_t size, intptr_t stream)
 cpdef memPrefetchAsync(intptr_t devPtr, size_t count, int dstDevice,
-                       size_t stream)
+                       intptr_t stream)
 cpdef memAdvise(intptr_t devPtr, size_t count, int advice, int device)
 cpdef PointerAttributes pointerGetAttributes(intptr_t ptr)
 
@@ -365,21 +365,21 @@ cpdef PointerAttributes pointerGetAttributes(intptr_t ptr)
 # Stream and Event
 ###############################################################################
 
-cpdef size_t streamCreate() except? 0
-cpdef size_t streamCreateWithFlags(unsigned int flags) except? 0
-cpdef streamDestroy(size_t stream)
-cpdef streamSynchronize(size_t stream)
-cpdef streamAddCallback(size_t stream, callback, intptr_t arg,
+cpdef intptr_t streamCreate() except? 0
+cpdef intptr_t streamCreateWithFlags(unsigned int flags) except? 0
+cpdef streamDestroy(intptr_t stream)
+cpdef streamSynchronize(intptr_t stream)
+cpdef streamAddCallback(intptr_t stream, callback, intptr_t arg,
                         unsigned int flags=*)
-cpdef streamQuery(size_t stream)
-cpdef streamWaitEvent(size_t stream, size_t event, unsigned int flags=*)
-cpdef size_t eventCreate() except? 0
-cpdef size_t eventCreateWithFlags(unsigned int flags) except? 0
-cpdef eventDestroy(size_t event)
-cpdef float eventElapsedTime(size_t start, size_t end) except? 0
-cpdef eventQuery(size_t event)
-cpdef eventRecord(size_t event, size_t stream)
-cpdef eventSynchronize(size_t event)
+cpdef streamQuery(intptr_t stream)
+cpdef streamWaitEvent(intptr_t stream, intptr_t event, unsigned int flags=*)
+cpdef intptr_t eventCreate() except? 0
+cpdef intptr_t eventCreateWithFlags(unsigned int flags) except? 0
+cpdef eventDestroy(intptr_t event)
+cpdef float eventElapsedTime(intptr_t start, intptr_t end) except? 0
+cpdef eventQuery(intptr_t event)
+cpdef eventRecord(intptr_t event, intptr_t stream)
+cpdef eventSynchronize(intptr_t event)
 
 
 ##############################################################################
diff --git a/cupy/cuda/runtime.pyx b/cupy/cuda/runtime.pyx
index eb75e84d7e7..f3811e802ea 100644
--- a/cupy/cuda/runtime.pyx
+++ b/cupy/cuda/runtime.pyx
@@ -350,7 +350,7 @@ cpdef memcpy(intptr_t dst, intptr_t src, size_t size, int kind):
 
 
 cpdef memcpyAsync(intptr_t dst, intptr_t src, size_t size, int kind,
-                  size_t stream):
+                  intptr_t stream):
     with nogil:
         status = cudaMemcpyAsync(
             <void*>dst, <void*>src, size, <MemoryKind>kind,
@@ -367,7 +367,7 @@ cpdef memcpyPeer(intptr_t dst, int dstDevice, intptr_t src, int srcDevice,
 
 
 cpdef memcpyPeerAsync(intptr_t dst, int dstDevice, intptr_t src, int srcDevice,
-                      size_t size, size_t stream):
+                      size_t size, intptr_t stream):
     with nogil:
         status = cudaMemcpyPeerAsync(<void*>dst, dstDevice, <void*>src,
                                      srcDevice, size, <driver.Stream> stream)
@@ -382,7 +382,7 @@ cpdef memcpy2D(intptr_t dst, size_t dpitch, intptr_t src, size_t spitch,
 
 cpdef memcpy2DAsync(intptr_t dst, size_t dpitch, intptr_t src, size_t spitch,
                     size_t width, size_t height, MemoryKind kind,
-                    size_t stream):
+                    intptr_t stream):
     with nogil:
         status = cudaMemcpy2DAsync(<void*>dst, dpitch, <void*>src, spitch,
                                    width, height, kind, <driver.Stream>stream)
@@ -400,7 +400,7 @@ cpdef memcpy2DFromArray(intptr_t dst, size_t dpitch, intptr_t src,
 
 cpdef memcpy2DFromArrayAsync(intptr_t dst, size_t dpitch, intptr_t src,
                              size_t wOffset, size_t hOffset, size_t width,
-                             size_t height, int kind, size_t stream):
+                             size_t height, int kind, intptr_t stream):
     with nogil:
         status = cudaMemcpy2DFromArrayAsync(<void*>dst, dpitch, <Array>src,
                                             wOffset, hOffset, width, height,
@@ -420,7 +420,7 @@ cpdef memcpy2DToArray(intptr_t dst, size_t wOffset, size_t hOffset,
 
 cpdef memcpy2DToArrayAsync(intptr_t dst, size_t wOffset, size_t hOffset,
                            intptr_t src, size_t spitch, size_t width,
-                           size_t height, int kind, size_t stream):
+                           size_t height, int kind, intptr_t stream):
     with nogil:
         status = cudaMemcpy2DToArrayAsync(<Array>dst, wOffset, hOffset,
                                           <void*>src, spitch, width, height,
@@ -435,7 +435,7 @@ cpdef memcpy3D(intptr_t Memcpy3DParmsPtr):
     check_status(status)
 
 
-cpdef memcpy3DAsync(intptr_t Memcpy3DParmsPtr, size_t stream):
+cpdef memcpy3DAsync(intptr_t Memcpy3DParmsPtr, intptr_t stream):
     with nogil:
         status = cudaMemcpy3DAsync(<Memcpy3DParms*>Memcpy3DParmsPtr,
                                    <driver.Stream> stream)
@@ -448,14 +448,14 @@ cpdef memset(intptr_t ptr, int value, size_t size):
     check_status(status)
 
 
-cpdef memsetAsync(intptr_t ptr, int value, size_t size, size_t stream):
+cpdef memsetAsync(intptr_t ptr, int value, size_t size, intptr_t stream):
     with nogil:
         status = cudaMemsetAsync(<void*>ptr, value, size,
                                  <driver.Stream> stream)
     check_status(status)
 
 cpdef memPrefetchAsync(intptr_t devPtr, size_t count, int dstDevice,
-                       size_t stream):
+                       intptr_t stream):
     with nogil:
         status = cudaMemPrefetchAsync(<void*>devPtr, count, dstDevice,
                                       <driver.Stream> stream)
@@ -483,26 +483,26 @@ cpdef PointerAttributes pointerGetAttributes(intptr_t ptr):
 # Stream and Event
 ###############################################################################
 
-cpdef size_t streamCreate() except? 0:
+cpdef intptr_t streamCreate() except? 0:
     cdef driver.Stream stream
     status = cudaStreamCreate(&stream)
     check_status(status)
-    return <size_t>stream
+    return <intptr_t>stream
 
 
-cpdef size_t streamCreateWithFlags(unsigned int flags) except? 0:
+cpdef intptr_t streamCreateWithFlags(unsigned int flags) except? 0:
     cdef driver.Stream stream
     status = cudaStreamCreateWithFlags(&stream, flags)
     check_status(status)
-    return <size_t>stream
+    return <intptr_t>stream
 
 
-cpdef streamDestroy(size_t stream):
+cpdef streamDestroy(intptr_t stream):
     status = cudaStreamDestroy(<driver.Stream>stream)
     check_status(status)
 
 
-cpdef streamSynchronize(size_t stream):
+cpdef streamSynchronize(intptr_t stream):
     with nogil:
         status = cudaStreamSynchronize(<driver.Stream>stream)
     check_status(status)
@@ -512,11 +512,11 @@ cdef _streamCallbackFunc(driver.Stream hStream, int status,
                          void* func_arg) with gil:
     obj = <object>func_arg
     func, arg = obj
-    func(<size_t>hStream, status, arg)
+    func(<intptr_t>hStream, status, arg)
     cpython.Py_DECREF(obj)
 
 
-cpdef streamAddCallback(size_t stream, callback, intptr_t arg,
+cpdef streamAddCallback(intptr_t stream, callback, intptr_t arg,
                         unsigned int flags=0):
     func_arg = (callback, arg)
     cpython.Py_INCREF(func_arg)
@@ -527,52 +527,52 @@ cpdef streamAddCallback(size_t stream, callback, intptr_t arg,
     check_status(status)
 
 
-cpdef streamQuery(size_t stream):
+cpdef streamQuery(intptr_t stream):
     return cudaStreamQuery(<driver.Stream>stream)
 
 
-cpdef streamWaitEvent(size_t stream, size_t event, unsigned int flags=0):
+cpdef streamWaitEvent(intptr_t stream, intptr_t event, unsigned int flags=0):
     with nogil:
         status = cudaStreamWaitEvent(<driver.Stream>stream,
                                      <driver.Event>event, flags)
     check_status(status)
 
 
-cpdef size_t eventCreate() except? 0:
+cpdef intptr_t eventCreate() except? 0:
     cdef driver.Event event
     status = cudaEventCreate(&event)
     check_status(status)
-    return <size_t>event
+    return <intptr_t>event
 
-cpdef size_t eventCreateWithFlags(unsigned int flags) except? 0:
+cpdef intptr_t eventCreateWithFlags(unsigned int flags) except? 0:
     cdef driver.Event event
     status = cudaEventCreateWithFlags(&event, flags)
     check_status(status)
-    return <size_t>event
+    return <intptr_t>event
 
 
-cpdef eventDestroy(size_t event):
+cpdef eventDestroy(intptr_t event):
     status = cudaEventDestroy(<driver.Event>event)
     check_status(status)
 
 
-cpdef float eventElapsedTime(size_t start, size_t end) except? 0:
+cpdef float eventElapsedTime(intptr_t start, intptr_t end) except? 0:
     cdef float ms
     status = cudaEventElapsedTime(&ms, <driver.Event>start, <driver.Event>end)
     check_status(status)
     return ms
 
 
-cpdef eventQuery(size_t event):
+cpdef eventQuery(intptr_t event):
     return cudaEventQuery(<driver.Event>event)
 
 
-cpdef eventRecord(size_t event, size_t stream):
+cpdef eventRecord(intptr_t event, intptr_t stream):
     status = cudaEventRecord(<driver.Event>event, <driver.Stream>stream)
     check_status(status)
 
 
-cpdef eventSynchronize(size_t event):
+cpdef eventSynchronize(intptr_t event):
     with nogil:
         status = cudaEventSynchronize(<driver.Event>event)
     check_status(status)
@@ -590,9 +590,9 @@ cdef _ensure_context():
 
     See discussion on https://github.com/cupy/cupy/issues/72 for details.
     """
-    cdef size_t status
-    status = <size_t>cpython.PyThread_get_key_value(_context_initialized)
-    if status == 0:
+    cdef void* status = NULL
+    status = cpython.PyThread_get_key_value(_context_initialized)
+    if status == NULL:
         # Call Runtime API to establish context on this host thread.
         memGetInfo()
         cpython.PyThread_set_key_value(_context_initialized, <void *>1)
diff --git a/cupy/cuda/stream.pxd b/cupy/cuda/stream.pxd
index bdc23dc9fe3..1e396b631d4 100644
--- a/cupy/cuda/stream.pxd
+++ b/cupy/cuda/stream.pxd
@@ -1,2 +1,5 @@
-cdef size_t get_current_stream_ptr()
+from libc.stdint cimport intptr_t
+
+
+cdef intptr_t get_current_stream_ptr()
 cpdef get_current_stream()
diff --git a/cupy/cuda/stream.pyx b/cupy/cuda/stream.pyx
index f64db0a7b06..921514be602 100644
--- a/cupy/cuda/stream.pyx
+++ b/cupy/cuda/stream.pyx
@@ -21,11 +21,11 @@ cdef class _StreamThreadLocal:
         return <_StreamThreadLocal>tls
 
     cdef set_current_stream(self, stream):
-        self.current_stream = <void*><size_t>stream.ptr
+        self.current_stream = <void*><intptr_t>stream.ptr
         self.current_stream_ref = weakref.ref(stream)
 
     cdef set_current_stream_ref(self, stream_ref):
-        self.current_stream = <void*><size_t>stream_ref().ptr
+        self.current_stream = <void*><intptr_t>stream_ref().ptr
         self.current_stream_ref = stream_ref
 
     cdef get_current_stream(self):
@@ -45,14 +45,14 @@ cdef class _StreamThreadLocal:
         return self.current_stream
 
 
-cdef size_t get_current_stream_ptr():
+cdef intptr_t get_current_stream_ptr():
     """C API to get current CUDA stream pointer.
 
     Returns:
-        size_t: The current CUDA stream pointer.
+        intptr_t: The current CUDA stream pointer.
     """
     tls = _StreamThreadLocal.get()
-    return <size_t>tls.get_current_stream_ptr()
+    return <intptr_t>tls.get_current_stream_ptr()
 
 
 cpdef get_current_stream():
@@ -81,7 +81,7 @@ class Event(object):
             processes.
 
     Attributes:
-        ~Event.ptr (size_t): Raw stream handle. It can be passed to
+        ~Event.ptr (intptr_t): Raw stream handle. It can be passed to
             the CUDA Runtime API via ctypes.
 
     """
@@ -161,7 +161,7 @@ class Stream(object):
             the NULL stream.
 
     Attributes:
-        ~Stream.ptr (size_t): Raw stream handle. It can be passed to
+        ~Stream.ptr (intptr_t): Raw stream handle. It can be passed to
             the CUDA Runtime API via ctypes.
 
     """
@@ -177,10 +177,11 @@ class Stream(object):
             self.ptr = runtime.streamCreate()
 
     def __del__(self):
+        cdef void* current_ptr
         if self.ptr:
             tls = _StreamThreadLocal.get()
             current_ptr = tls.get_current_stream_ptr()
-            if <size_t>self.ptr == <size_t>current_ptr:
+            if <void*>self.ptr == current_ptr:
                 tls.set_current_stream(self.null)
             runtime.streamDestroy(self.ptr)
         # Note that we can not release memory pool of the stream held in CPU
diff --git a/cupy/cuda/texture.pxd b/cupy/cuda/texture.pxd
index 01d3b83f6fb..3c8b00a5536 100644
--- a/cupy/cuda/texture.pxd
+++ b/cupy/cuda/texture.pxd
@@ -42,7 +42,7 @@ cdef class TextureObject:
 
 cdef class TextureReference:
     cdef:
-        readonly size_t texref
+        readonly intptr_t texref
         readonly ResourceDescriptor ResDesc
         readonly TextureDescriptor TexDesc
 
diff --git a/cupy/cuda/texture.pyx b/cupy/cuda/texture.pyx
index 12a68a834b9..fa4381443c6 100644
--- a/cupy/cuda/texture.pyx
+++ b/cupy/cuda/texture.pyx
@@ -509,8 +509,8 @@ cdef class TextureReference:
     driver API is used under the hood).
 
     Args:
-        texref (size_t): a handle to the texture reference declared in the CUDA
-            source code. This can be obtained by calling
+        texref (intptr_t): a handle to the texture reference declared in the
+            CUDA source code. This can be obtained by calling
             :meth:`~cupy.RawModule.get_texref`.
         ResDesc (ResourceDescriptor): an intance of the resource descriptor.
         TexDesc (TextureDescriptor): an instance of the texture descriptor.
@@ -533,7 +533,7 @@ cdef class TextureReference:
     '''  # noqa
     # Basically, this class translates from the Runtime API's descriptors
     # to the driver API calls.
-    def __init__(self, size_t texref, ResourceDescriptor ResDesc,
+    def __init__(self, intptr_t texref, ResourceDescriptor ResDesc,
                  TextureDescriptor TexDesc):
         cdef ResourceDesc* ResDescPtr = <ResourceDesc*>(ResDesc.ptr)
         cdef TextureDesc* TexDescPtr = <TextureDesc*>(TexDesc.ptr)
@@ -546,14 +546,15 @@ cdef class TextureReference:
         self.TexDesc = TexDesc
 
         if ResDescPtr.resType == runtime.cudaResourceTypeArray:
-            driver.texRefSetArray(texref, <size_t>(ResDescPtr.res.array.array))
+            driver.texRefSetArray(texref,
+                                  <intptr_t>(ResDescPtr.res.array.array))
             ChDesc = ResDesc.cuArr.desc
             arr_format, num_channels = \
                 self._get_format(ChDesc.get_channel_format())
             driver.texRefSetFormat(self.texref, arr_format, num_channels)
         elif ResDescPtr.resType == runtime.cudaResourceTypeLinear:
             driver.texRefSetAddress(texref,
-                                    <size_t>(ResDescPtr.res.linear.devPtr),
+                                    <intptr_t>(ResDescPtr.res.linear.devPtr),
                                     ResDescPtr.res.linear.sizeInBytes)
             ChDesc = ResDesc.chDesc
             arr_format, num_channels = \
@@ -567,9 +568,10 @@ cdef class TextureReference:
             arr_desc.NumChannels = num_channels
             arr_desc.Height = ResDescPtr.res.pitch2D.height
             arr_desc.Width = ResDescPtr.res.pitch2D.width
-            driver.texRefSetAddress2D(texref, <size_t>(&arr_desc),
-                                      <size_t>(ResDescPtr.res.pitch2D.devPtr),
-                                      ResDescPtr.res.pitch2D.pitchInBytes)
+            driver.texRefSetAddress2D(
+                texref, <intptr_t>(&arr_desc),
+                <intptr_t>(ResDescPtr.res.pitch2D.devPtr),
+                ResDescPtr.res.pitch2D.pitchInBytes)
             # don't call driver.texRefSetFormat() here!
         else:  # TODO(leofang): mipmap
             raise ValueError("mpimap array is not supported yet.")
diff --git a/docs/source/contribution.rst b/docs/source/contribution.rst
index e26278620f9..f02723465ab 100644
--- a/docs/source/contribution.rst
+++ b/docs/source/contribution.rst
@@ -261,6 +261,15 @@ Please note the followings when writing the document.
   original one, users should explicitly describe only what is implemented in
   the document.
 
+For changes that modify or add new Cython files, please make sure the pointer types follow these guidelines (`GH-1913 <https://github.com/cupy/cupy/issues/1913>`_).
+
+* Pointers should be ``intptr_t``.
+* Memory sizes should be ``size_t``.
+* Memory offsets should be ``ptrdiff_t``.
+
+.. note::
+
+     We are incrementally enforcing the above rules, so some existing code may not follow the above guidelines, but please ensure all new contributions do.
 
 .. _testing-guide:
 


Code review comments:
diff_hunk: @@ -261,6 +261,15 @@ Please note the followings when writing the document.
   original one, users should explicitly describe only what is implemented in
   the document.
 
+For changes that modify or add new Cython files, please make sure the pointer types follow these guidelines (`GH-1913 <https://github.com/cupy/cupy/issues/1913>`_).
+
+* Pointers should be ``intptr_t``.
path: docs/source/contribution.rst
text: @user1:
I think we could mention `void*` as well.
How about something like:
"Pointers should be `void*` if it's only used from within Cython, or `intptr_t` if it's exposed to Python code.

(Please feel free to rephrase)

@author:
@user1 I am not sure about the usage of `void*`. It seems in CuPy's Cython codes they're all `ctypedef`'d? Anyway, how about this?
```suggestion
* Pointers should be ``void*`` if only used within Cython, or ``intptr_t`` if exposed to the Python space.
```

@author:
Applied the above patch, please take a look. Thanks.

@user1:
Thanks, sounds better!
Not all pointers are `ctypedef`s, including raw data pointers.