Merge branch 'master' into mempool_impl

.gitmodules

@@ -1,7 +1,7 @@
-[submodule "cupy/core/include/cupy/cub"]
-	path = cupy/core/include/cupy/cub
+[submodule "cupy/_core/include/cupy/cub"]
+	path = cupy/_core/include/cupy/cub
 	url = https://github.com/NVlabs/cub.git
 	branch = 1.8.0
-[submodule "cupy/core/include/cupy/jitify"]
-	path = cupy/core/include/cupy/jitify
+[submodule "cupy/_core/include/cupy/jitify"]
+	path = cupy/_core/include/cupy/jitify
 	url = https://github.com/NVIDIA/jitify.git

.pfnci/script.sh

@@ -68,7 +68,7 @@ main() {
       python3.7 -m pip install cython numpy
 
       wget -qO - http://repo.radeon.com/rocm/rocm.gpg.key | apt-key add -
-      echo 'deb [arch=amd64] http://repo.radeon.com/rocm/apt/debian/ xenial main' | tee /etc/apt/sources.list.d/rocm.list
+      echo 'deb [arch=amd64] http://repo.radeon.com/rocm/apt/4.0.1/ xenial main' | tee /etc/apt/sources.list.d/rocm.list
 
       # Uninstall CUDA to ensure it's a clean ROCm environment
       # https://docs.nvidia.com/cuda/cuda-installation-guide-linux/index.html#removing-cuda-tk-and-driver
@@ -148,11 +148,13 @@ prepare_docker() {
   run gcloud auth configure-docker
 }
 
+KNOWN_BASE_BRANCHES="master v8 v9"
+
 # is_known_base_branch returns 0 only if the given branch name is a known
 # base development branch.
 is_known_base_branch() {
   local branch="${1##refs/heads/}"
-  for BASE_BRANCH in master v7 v8; do
+  for BASE_BRANCH in ${KNOWN_BASE_BRANCHES}; do
     if [ "${branch}" = "${BASE_BRANCH}" ]; then
       return 0
     fi
@@ -162,7 +164,7 @@ is_known_base_branch() {
 
 # get_base_branch returns the base development branch for the current HEAD.
 get_base_branch() {
-  for BASE_BRANCH in master v7 v8; do
+  for BASE_BRANCH in ${KNOWN_BASE_BRANCHES}; do
     git merge-base --is-ancestor "origin/${BASE_BRANCH}" HEAD && echo "${BASE_BRANCH}" && return 0
   done
   echo "Base branch of HEAD is not valid." >&2

.pfnci/windows/_flexci.ps1

@@ -43,7 +43,7 @@ function ActivateCUDA($version) {
 }
 
 function IsPullRequestTest() {
-    return ${Env:FLEXCI_BRANCH}.StartsWith("refs/pull/")
+    return ${Env:FLEXCI_BRANCH} -ne $null -and ${Env:FLEXCI_BRANCH}.StartsWith("refs/pull/")
 }
 
 function PrioritizeFlexCIDaemon() {

README.md

@@ -33,8 +33,8 @@ Choose the right package for your platform.
 | CUDA 10.2 | `pip install cupy-cuda102`     |
 | CUDA 11.0 | `pip install cupy-cuda110`     |
 | CUDA 11.1 | `pip install cupy-cuda111`     |
-| CUDA 11.2 | `pip install cupy-cuda112` (see [#4704](https://github.com/cupy/cupy/issues/4704) for Linux instructions) |
-| ROCm 4.0  | `pip install cupy-rocm-4-0` (experimental; see [docs](https://docs.cupy.dev/en/latest/install_rocm.html) for details) |
+| CUDA 11.2 | `pip install cupy-cuda112`     |
+| ROCm 4.0  | `pip install cupy-rocm-4-0` (experimental; see [docs](https://docs.cupy.dev/en/latest/install.html#using-cupy-on-amd-gpu-experimental) for details) |
 
 See the [Installation Guide](https://docs.cupy.dev/en/stable/install.html) if you are using Conda/Anaconda or to build from source.
 

cupy/__init__.py

@@ -13,10 +13,10 @@
 
 
 try:
-    from cupy import core  # NOQA
+    from cupy import _core  # NOQA
 except ImportError as e:
-    # core is a c-extension module.
-    # When a user cannot import core, it represents that CuPy is not correctly
+    # _core is a c-extension module.
+    # When a user cannot import _core, it represents that CuPy is not correctly
     # built.
     _exc_info = _sys.exc_info()
     _msg = ('''\
@@ -49,7 +49,7 @@ def is_available():
 __version__ = _version.__version__
 
 
-import cupy.core.fusion  # NOQA
+import cupy._core.fusion  # NOQA
 from cupy import fft  # NOQA
 from cupy import linalg  # NOQA
 from cupy import polynomial  # NOQA
@@ -60,8 +60,8 @@ def is_available():
 
 
 # import class and function
-from cupy.core import ndarray  # NOQA
-from cupy.core import ufunc  # NOQA
+from cupy._core import ndarray  # NOQA
+from cupy._core import ufunc  # NOQA
 
 
 # =============================================================================
@@ -715,7 +715,7 @@ def isscalar(element):
 # -----------------------------------------------------------------------------
 # Undocumented functions
 # -----------------------------------------------------------------------------
-from cupy.core import size  # NOQA
+from cupy._core import size  # NOQA
 
 
 def ndim(a):
@@ -742,16 +742,16 @@ def ndim(a):
 from cupy._util import clear_memo  # NOQA
 from cupy._util import memoize  # NOQA
 
-from cupy.core import ElementwiseKernel  # NOQA
-from cupy.core import RawKernel  # NOQA
-from cupy.core import RawModule  # NOQA
-from cupy.core._reduction import ReductionKernel  # NOQA
+from cupy._core import ElementwiseKernel  # NOQA
+from cupy._core import RawKernel  # NOQA
+from cupy._core import RawModule  # NOQA
+from cupy._core._reduction import ReductionKernel  # NOQA
 
 # -----------------------------------------------------------------------------
 # DLPack
 # -----------------------------------------------------------------------------
 
-from cupy.core import fromDlpack  # NOQA
+from cupy._core import fromDlpack  # NOQA
 
 
 def asnumpy(a, stream=None, order='C'):
@@ -806,13 +806,13 @@ def get_array_module(*args):
     """
     for arg in args:
         if isinstance(arg, (ndarray, _cupyx.scipy.sparse.spmatrix,
-                            cupy.core.fusion._FusionVarArray,
-                            cupy.core.new_fusion._ArrayProxy)):
+                            cupy._core.fusion._FusionVarArray,
+                            cupy._core.new_fusion._ArrayProxy)):
             return _cupy
     return _numpy
 
 
-fuse = cupy.core.fusion.fuse
+fuse = cupy._core.fusion.fuse
 
 disable_experimental_feature_warning = False
 

cupy/_binary/elementwise.py

@@ -1,22 +1,22 @@
-from cupy import core
+from cupy import _core
 
 
-bitwise_and = core.bitwise_and
+bitwise_and = _core.bitwise_and
 
 
-bitwise_or = core.bitwise_or
+bitwise_or = _core.bitwise_or
 
 
-bitwise_xor = core.bitwise_xor
+bitwise_xor = _core.bitwise_xor
 
 
-bitwise_not = core.invert
+bitwise_not = _core.invert
 
 
-invert = core.invert
+invert = _core.invert
 
 
-left_shift = core.left_shift
+left_shift = _core.left_shift
 
 
-right_shift = core.right_shift
+right_shift = _core.right_shift

cupy/_binary/packing.py

@@ -1,8 +1,8 @@
 import cupy
-from cupy import core
+from cupy import _core
 
 
-_packbits_kernel = core.ElementwiseKernel(
+_packbits_kernel = _core.ElementwiseKernel(
     'raw T myarray, raw int32 myarray_size', 'uint8 packed',
     '''for (int j = 0; j < 8; ++j) {
         int k = i * 8 + j;
@@ -42,7 +42,7 @@ def packbits(myarray):
     return _packbits_kernel(myarray, myarray.size, packed)
 
 
-_unpackbits_kernel = core.ElementwiseKernel(
+_unpackbits_kernel = _core.ElementwiseKernel(
     'raw uint8 myarray', 'T unpacked',
     'unpacked = (myarray[i / 8] >> (7 - i % 8)) & 1;',
     'unpackbits_kernel'

cupy/_core/__init__.py

@@ -0,0 +1,74 @@
+from cupy._core import core  # NOQA
+from cupy._core import internal  # NOQA
+
+
+# internal APIs for testing and developement
+from cupy._core._accelerator import set_reduction_accelerators  # NOQA
+from cupy._core._accelerator import set_routine_accelerators  # NOQA
+from cupy._core._accelerator import get_reduction_accelerators  # NOQA
+from cupy._core._accelerator import get_routine_accelerators  # NOQA
+
+
+# import class and function
+from cupy._core._kernel import create_ufunc  # NOQA
+from cupy._core._kernel import ElementwiseKernel  # NOQA
+from cupy._core._kernel import ufunc  # NOQA
+from cupy._core._reduction import create_reduction_func  # NOQA
+from cupy._core._reduction import ReductionKernel  # NOQA
+from cupy._core._routines_binary import bitwise_and  # NOQA
+from cupy._core._routines_binary import bitwise_or  # NOQA
+from cupy._core._routines_binary import bitwise_xor  # NOQA
+from cupy._core._routines_binary import invert  # NOQA
+from cupy._core._routines_binary import left_shift  # NOQA
+from cupy._core._routines_binary import right_shift  # NOQA
+from cupy._core._routines_linalg import _mat_ptrs  # NOQA
+from cupy._core._routines_linalg import dot  # NOQA
+from cupy._core._routines_linalg import get_compute_type  # NOQA
+from cupy._core._routines_linalg import _matmul  # NOQA
+from cupy._core._routines_linalg import set_compute_type  # NOQA
+from cupy._core._routines_linalg import tensordot_core  # NOQA
+from cupy._core._routines_logic import create_comparison  # NOQA
+from cupy._core._routines_logic import equal  # NOQA
+from cupy._core._routines_logic import greater  # NOQA
+from cupy._core._routines_logic import greater_equal  # NOQA
+from cupy._core._routines_logic import less  # NOQA
+from cupy._core._routines_logic import less_equal  # NOQA
+from cupy._core._routines_logic import not_equal  # NOQA
+from cupy._core._routines_manipulation import array_split  # NOQA
+from cupy._core._routines_manipulation import broadcast  # NOQA
+from cupy._core._routines_manipulation import broadcast_to  # NOQA
+from cupy._core._routines_manipulation import concatenate_method  # NOQA
+from cupy._core._routines_manipulation import moveaxis  # NOQA
+from cupy._core._routines_manipulation import rollaxis  # NOQA
+from cupy._core._routines_manipulation import size  # NOQA'
+from cupy._core._routines_math import absolute  # NOQA
+from cupy._core._routines_math import add  # NOQA
+from cupy._core._routines_math import angle  # NOQA
+from cupy._core._routines_math import conjugate  # NOQA
+from cupy._core._routines_math import divide  # NOQA
+from cupy._core._routines_math import floor_divide  # NOQA
+from cupy._core._routines_math import imag  # NOQA
+from cupy._core._routines_math import multiply  # NOQA
+from cupy._core._routines_math import negative  # NOQA
+from cupy._core._routines_math import power  # NOQA
+from cupy._core._routines_math import real  # NOQA
+from cupy._core._routines_math import remainder  # NOQA
+from cupy._core._routines_math import sqrt  # NOQA
+from cupy._core._routines_math import subtract  # NOQA
+from cupy._core._routines_math import true_divide  # NOQA
+from cupy._core._routines_statistics import nanmax  # NOQA
+from cupy._core._routines_statistics import nanmin  # NOQA
+from cupy._core.core import _internal_ascontiguousarray  # NOQA
+from cupy._core.core import _internal_asfortranarray  # NOQA
+from cupy._core.core import array  # NOQA
+from cupy._core.core import ascontiguousarray  # NOQA
+from cupy._core.core import asfortranarray  # NOQA
+from cupy._core.core import divmod  # NOQA
+from cupy._core.core import elementwise_copy  # NOQA
+from cupy._core.core import elementwise_copy_where  # NOQA
+from cupy._core.core import ndarray  # NOQA
+from cupy._core.dlpack import fromDlpack  # NOQA
+from cupy._core.internal import complete_slice  # NOQA
+from cupy._core.internal import get_size  # NOQA
+from cupy._core.raw import RawKernel  # NOQA
+from cupy._core.raw import RawModule  # NOQA

cupy/core/_carray.pyx → cupy/_core/_carray.pyx

@@ -3,7 +3,7 @@ import os
 from libcpp cimport vector
 
 from cupy.cuda cimport function
-from cupy.core cimport internal
+from cupy._core cimport internal
 
 
 cdef class CArray(function.CPointer):

cupy/core/_cub_reduction.pxd → cupy/_core/_cub_reduction.pxd

@@ -1,6 +1,6 @@
-from cupy.core._carray cimport shape_t
-from cupy.core._kernel cimport _TypeMap
-from cupy.core.core cimport ndarray
+from cupy._core._carray cimport shape_t
+from cupy._core._kernel cimport _TypeMap
+from cupy._core.core cimport ndarray
 
 
 cdef bint _try_to_call_cub_reduction(

cupy/core/_cub_reduction.pyx → cupy/_core/_cub_reduction.pyx

@@ -1,12 +1,12 @@
-from cupy.core._carray cimport shape_t
-from cupy.core cimport _kernel
-from cupy.core cimport _optimize_config
-from cupy.core cimport _reduction
-from cupy.core cimport _scalar
-from cupy.core.core cimport compile_with_cache
-from cupy.core.core cimport ndarray
-from cupy.core.core cimport _internal_ascontiguousarray
-from cupy.core cimport internal
+from cupy._core._carray cimport shape_t
+from cupy._core cimport _kernel
+from cupy._core cimport _optimize_config
+from cupy._core cimport _reduction
+from cupy._core cimport _scalar
+from cupy._core.core cimport compile_with_cache
+from cupy._core.core cimport ndarray
+from cupy._core.core cimport _internal_ascontiguousarray
+from cupy._core cimport internal
 from cupy.cuda cimport cub
 from cupy.cuda cimport function
 from cupy.cuda cimport memory
@@ -16,7 +16,7 @@ import math
 import string
 import sys
 from cupy import _environment
-from cupy.core._kernel import _get_param_info
+from cupy._core._kernel import _get_param_info
 from cupy.cuda import driver
 from cupy import _util
 
@@ -356,7 +356,7 @@ cpdef inline tuple _can_use_cub_block_reduction(
     return (axis_permutes_cub, contiguous_size, full_reduction)
 
 
-# similar to cupy.core._kernel._get_kernel_params()
+# similar to cupy._core._kernel._get_kernel_params()
 cdef str _get_cub_kernel_params(tuple params, tuple arginfos):
     cdef _kernel.ParameterInfo p
     cdef _kernel._ArgInfo arginfo

cupy/core/_fusion_interface.py → cupy/_core/_fusion_interface.py

@@ -1,10 +1,10 @@
 import numpy
 
-from cupy.core._dtype import get_dtype
+from cupy._core._dtype import get_dtype
 import cupy
-from cupy.core import _fusion_thread_local
-from cupy.core import core
-from cupy.core._scalar import get_typename
+from cupy._core import _fusion_thread_local
+from cupy._core import core
+from cupy._core._scalar import get_typename
 
 
 _thread_local = _fusion_thread_local.thread_local
@@ -37,7 +37,7 @@ class _VariableProxy:
     """
 
     def __init__(self, content):
-        assert isinstance(content, cupy.core._fusion_variable._TraceVariable)
+        assert isinstance(content, cupy._core._fusion_variable._TraceVariable)
         self.content = content
 
     def __neg__(self):

cupy/core/_fusion_kernel.pyx → cupy/_core/_fusion_kernel.pyx

@@ -4,24 +4,24 @@ import string
 from libcpp cimport vector
 import numpy
 
-from cupy.core cimport _carray
-from cupy.core.core cimport _ndarray_init
-from cupy.core.core cimport compile_with_cache
-from cupy.core.core cimport ndarray
-from cupy.core cimport internal
-from cupy.core cimport _routines_manipulation as _manipulation
+from cupy._core cimport _carray
+from cupy._core.core cimport _ndarray_init
+from cupy._core.core cimport compile_with_cache
+from cupy._core.core cimport ndarray
+from cupy._core cimport internal
+from cupy._core cimport _routines_manipulation as _manipulation
 from cupy_backends.cuda.api cimport driver
 from cupy.cuda cimport function
 from cupy_backends.cuda.api cimport runtime
-from cupy.core cimport _reduction
+from cupy._core cimport _reduction
 
-from cupy.core import _dtype
+from cupy._core import _dtype
 from cupy import _util
-from cupy.core import _fusion_op
-from cupy.core._fusion_variable import _AbstractDim
-from cupy.core._fusion_variable import _TraceVariable
-from cupy.core._fusion_variable import _TraceScalar
-from cupy.core._fusion_variable import _TraceArray
+from cupy._core import _fusion_op
+from cupy._core._fusion_variable import _AbstractDim
+from cupy._core._fusion_variable import _TraceVariable
+from cupy._core._fusion_variable import _TraceScalar
+from cupy._core._fusion_variable import _TraceArray
 from cupyx.jit import _codeblock