1. 概述

本文档主要用于指导用户在曦云® 系列 GPU 上安装部署和使用 mcPy。

1.1. mcPy 介绍

mcPy 是一个兼容 NumPy 并使用曦云系列 GPU 进行计算加速的 Python 工具库。 mcPy 的目标是给 Python 用户提供 GPU 的计算加速能力，用户只需要执行简单的代码替换就可以将原有的 NumPy 应用程序代码迁移到曦云系列 GPU 上运行。

mcPy 支持 ndarray 数据对象并提供相关的 API 接口，这些 API 接口兼容与之对应的NumPy API，不同之处在于 mcPy 底层由曦云系列 GPU 提供计算加速。 mcPy 提供的 API 接口如下：

  • N 维数组（ndarray）： cupy.ndarray

    • 数据类型（dtypes）：boolean（bool）、integer（int8, int16, int32, int64, uint8, uint16, uint32, uint64）、float（float16, float32, float64）和complex（complex64）

    • 支持与 numpy.ndarray 相同的编程方式，包括基础索引、高级索引和广播

  • mcPy例程

    • 模块级函数： cupy.*

    • 线性代数函数： cupy.linalg.*

    • 快速傅里叶变换： cupy.fft.*

    • 随机数生成器： cupy.random.*

mcPy 的 API 构建在 MXMACA SDK 之上，包括 mcBLAS、mcFFT、mcSPARSE、mcSOLVER、mcRAND、mcThrust、mcCUB，并由这些底层 SDK 提供最佳的性能。

mcPy 支持对 ndarray 对象应用自定义的核函数（kernel function），包括：

  • 内核模版（Kernel Template）：在单个核函数中快速定义 element-wise操作和 reduction 操作

  • 原始内核（Raw Kernel）：快速导入现有的MXMACA C++代码

  • 即时编译（JIT Compile）：从 Python 代码中编译生成内核

  • 内核融合（Kernel Fusion）：将多个 mcPy 操作融合成一个内核

mcPy 实现了用于数据交换和互操作的标准API，包括 DLPack、__array_ufunc__（NEP 13）、__array_function__（NEP 18）和数组 API 标准（Array API Standard NEP 47）。 基于这些标准的 API 和协议，mcPy 可以很容易地和 NumPy、PyTorch、TensorFlow、MPI4Py 以及其他支持相同 API 和协议的 Python 库进行数据互操作。

1.2. 软件包信息

mcPy 提供了 mxc500-mcpy-<VERSION>-linux-x86_64.tar.xz 和 mxc500-mcpy-<VERSION>-linux-aarch64.tar.xz，可以使用命令 tar -xJf 进行解压，解压后的二进制软件包信息说明参见表 1.1。

表 1.1 mcPy发布的软件包

软件包类型

文件名示例

说明

Python wheel包

numbax-x.y.z.n+bu.v.w.m-cp38-cp38-linux_x86_64.whl

适用于x86_64和Python 3.8的numbax Python二进制扩展包

Python wheel包

numbax-x.y.z.n+bu.v.w.m-cp310-cp310-linux_x86_64.whl

适用于x86_64和Python 3.10的numbax Python二进制扩展包

Python wheel包

mcpy-x.y.z.n+bu.v.w.m-cp38-cp38-linux_x86_64.whl

适用于x86_64和Python 3.8的mcPy Python二进制扩展包

Python wheel包

mcpy-x.y.z.n+bu.v.w.m-cp310-cp310-linux_x86_64.whl

适用于x86_64和Python 3.10的mcPy Python二进制扩展包

Python wheel包

numbax-x.y.z.n+bu.v.w.m-cp38-cp38-linux_aarch64.whl

适用于aarch64和Python 3.8的numbax Python二进制扩展包

Python wheel包

numbax-x.y.z.n+bu.v.w.m-cp310-cp310-linux_aarch64.whl

适用于aarch64和Python 3.10的numbax Python二进制扩展包

Python wheel包

mcpy-x.y.z.n+bu.v.w.m-cp38-cp38-linux_aarch64.whl

适用于aarch64和Python 3.8的mcPy Python二进制扩展包

Python wheel包

mcpy-x.y.z.n+bu.v.w.m-cp310-cp310-linux_aarch64.whl

适用于aarch64和Python 3.10的mcPy Python二进制扩展包

备注

  • x.y.z.n 表示对应包的软件发布版本，bu.v.w.m 表示基于版本号为 u.v.w.m 的 MXMACA 进行编译的构建号。

  • mcPy 运行依赖 numbax。

  • 当前支持 x86_64 和 aarch64 平台的 Python 3.8 和 3.10，应选取适配具体 Python 版本的软件包。

2. 安装部署

2.1. 依赖关系

  • Host支持x86_64和aarch64，且glibc≥2.27

  • Host 操作系统

    • x86_64平台支持Ubuntu 18.04 LTS、Ubuntu 20.04 LTS、Ubuntu 22.04 LTS、Centos 8

    • aarch64平台支持Kylin Linux advanced server V10

  • Device硬件仅支持曦云系列 GPU

  • MXMACA SDK，配套版本以《MXMAP发布说明》为准

  • 支持 Python 3.8 或 3.10

2.1.1. Python 依赖

mcPy API接口兼容 NumPy 1.24 版本的API，mcPy安装时依赖如下 Python 库：

  • NumPy≥1.21，且≤1.27

  • FastRLock≥0.5

2.2. mcPy 安装与卸载

本节描述的安装与卸载方法以x86_64平台为例。

2.2.1. wheel包

安装

  • Python 3.8

    $ pip install numbax-x.y.z.n+bu.v.w.m-cp38-cp38-linux_x86_64.whl
$ pip install mcpy-x.y.z.n+bu.v.w.m-cp38-cp38-linux_x86_64.whl

  • Python 3.10

    $ pip install numbax-x.y.z.n+bu.v.w.m-cp310-cp310-linux_x86_64.whl
$ pip install mcpy-x.y.z.n+bu.v.w.m-cp310-cp310-linux_x86_64.whl

卸载

$ pip uninstall mcpy
$ pip uninstall numbax

升级

  • Python 3.8

    $ pip install -U numbax-x.y.z.n+bu.v.w.m-cp38-cp38-linux_x86_64.whl
$ pip install -U mcpy-x.y.z.n+bu.v.w.m-cp38-cp38-linux_x86_64.whl

  • Python 3.10

    $ pip install -U numbax-x.y.z.n+bu.v.w.m-cp310-cp310-linux_x86_64.whl
$ pip install -U mcpy-x.y.z.n+bu.v.w.m-cp310-cp310-linux_x86_64.whl

重新安装

  • Python 3.8

    $ pip uninstall numbax
$ pip uninstall mcpy
$ pip install numbax-x.y.z.n+bu.v.w.m-cp38-cp38-linux_x86_64.whl --no-cache-dir
$ pip install mcpy-x.y.z.n+bu.v.w.m-cp38-cp38-linux_x86_64.whl --no-cache-dir

  • Python 3.10

    $ pip uninstall numbax
$ pip uninstall mcpy
$ pip install numbax-x.y.z.n+bu.v.w.m-cp310-cp310-linux_x86_64.whl --no-cache-dir
$ pip install mcpy-x.y.z.n+bu.v.w.m-cp310-cp310-linux_x86_64.whl --no-cache-dir

2.3. 环境变量设置

mcPy 正常运行依赖如下环境变量：

  • MACA_PATH：指向 MXMACA 包的安装路径

  • MXMACA 安装和运行所需要的其他环境变量

2.4. 常见问题

关于 glibc 版本配套：mcPy运行时的 glibc 版本需 ≥2.27，否则可能存在符号缺失的问题。可通过 ldd --version 命令查询当前环境中的 glibc 版本。

3. mcPy 用户指南

3.1. mcPy 基础

3.1.1. mcPy N维数组基础

如下代码中 cp 是 cupy 的缩写， np 是 numpy 的缩写：

>>> import numpy as np
>>> import cupy as cp

cupy.ndarray 是 mcPy 的核心类，它是 numpy.ndarray 的替代库。

>>> x_gpu = cp.array([1, 2, 3])

上述代码中 x_gpu 是 cupy.ndarray 的一个实例，如代码所示，mcPy 的代码语法和 NumPy 的语法相同。mcPy和 NumPy 的主要区别在于 mcPy array 实例的数据保存在 GPU 之上。

绝大部分的 array API 和 NumPy API 相同，以欧几里得范数（Euclidean norm），即 L2 范数为例说明， numpy.linalg.norm() 在 CPU 上计算。

>>> x_cpu = np.array([1, 2, 3])
>>> l2_cpu = np.linalg.norm(x_cpu)

使用 mcPy，可以用类似的方式在 GPU 上完成相同的计算：

>>> x_gpu = cp.array([1, 2, 3])
>>> l2_gpu = cp.linalg.norm(x_gpu)

mcPy 在 cupy.ndarray 对象上实现了许多功能函数，这些功能函数是 NumPy API 的一个子集。具备 NumPy 的相关知识将有利于快速掌握 mcPy，用户可阅读 NumPy使用文档 熟悉了解 NumPy。

3.1.2. 当前设备（Current Device）

mcPy 中有“当前设备”的概念，当前设备是指保存 ndarray 数据以及对 ndarray 数据执行计算的 GPU 设备。默认情况下当前设备的 ID 为 0。如下代码中将在GPU 0 上创建一个 array 对象 x_on_gpu0。

>>> x_on_gpu0 = cp.array([1, 2, 3, 4, 5])

如果要切换到其他 GPU 设备上，可以使用 Device 上下文管理器：

>>> with cp.cuda.Device(1):
...    x_on_gpu1 = cp.array([1, 2, 3, 4, 5])
>>> x_on_gpu0 = cp.array([1, 2, 3, 4, 5])

除了多 GPU 特性和设备到设备拷贝特性之外的所有 mcPy 操作都将在当前设备上执行。

一般情况下，mcPy 函数假设 array 保存于当前设备上。如果传入 API 的 array 对象不是位于当前设备上，API的执行结果将取决于具体的硬件配置，不推荐这样编码。

如果 array 所位于的设备不是当前设备，mcPy 将尝试在这两个设备之间建立 P2P 内存访问通道并将 array 从另外一个设备拷贝到当前设备上。如果 P2P 功能无法正常执行，上述尝试将导致失败，并抛出 ValueError 异常。

cupy.ndarray.device 属性标识了 array 的数据保存在哪个设备上。

>>> with cp.cuda.Device(1):
...     x = cp.array([1, 2, 3, 4, 5])
>>> x.device
<GPU Device 1>

3.1.3. 当前流（Current Stream）

与“当前设备”相关联的另一个概念是“当前流”。当前流可以避免在每个 API 接口参数中显式传递流对象，这样有利于保持和 NumPy 相同的接口参数。mcPy 中，将涉及到 GPU 的操作（例如数据传输和内核加载）加入当前流的任务队列中，同一个流的任务队列中的多个操作将串行执行，但是从 host 的角度看，这些操作和 host 的操作是并行执行的。

mcPy 中默认的当前流是 0 号流，0号流也被称为默认流（default stream），它在整个 GPU 设备内唯一。可以通过 cupy.cuda.Stream 接口修改当前流。可以通过 cupy.cuda.get_current_stream 接口查询当前流。

mcPy 的当前流是在基于每个线程、每个设备进行管理的，这意味着在不同的 Python 线程或不同的设备上，当前流（如果不是默认流）可能不同。

3.1.4. 数据传输（Data Transfer）

3.1.4.1. 将数据移动到设备

cupy.asarray() 可以用来传输一个 numpy.ndarray、一个 list 或者任何可以传递给 numpy.array() 的对象到当前设备：

>>> x_cpu = np.array([1, 2, 3])
>>> x_gpu = cp.asarray(x_cpu) # move the data to the current device.

cupy.asarray() 接受 cupy.ndarray 对象作为入参，所以可以使用这个接口在不同设备之间移动 array 数据。

>>> with cp.cuda.Device(0):
...     x_gpu_0 = cp.ndarray([1, 2, 3]) # create an array in GPU 0
>>> with cp.cuda.Device(1):
...     x_gpu_1 = cp.asarray(x_gpu_0) # move the array to GPU 1

备注

cupy.asarray() 尽量不做数据拷贝，如果传入的 array 对象位于当前设备，该接口将会返回输入的 array 对象本身。

如果需要在这种情况下实现数据拷贝，可以使用 cupy.array() 并传入 copy=True 参数。实际上 cupy.asarray() 等同于 cupy.array(arr, dtype, copy=False)。

3.1.4.2. 将数据从设备移动到主机

可通过 cupy.asnumpy() 将设备上的 array 对象移动到 host 上：

>>> x_gpu = cp.array([1, 2, 3]) # create an array in the current device
>>> x_cpu = cp.asnumpy(x_gpu) # move the array to the host.

或者使用 cupy.ndarray.get()：

>>> x_cpu = x_gpu.get()

3.1.5. 编写CPU/GPU无关代码

mcPy 和 NumPy 之间的兼容性使得开发者可以编写 CPU/GPU 无关代码。为了实现上述目标，mcPy 实现了 cupy.get_array_module() 函数，如果该函数的任何一个入参位于 GPU，该函数将返回 mcpy 模块，否则返回 numpy 模块。以下代码演示了如何编写计算 loglp 的 CPU/GPU 无关代码：

>>> # Stable implementation of log(1 + exp(x))
>>> def softplus(x):
...     xp = cp.get_array_module(x) # 'xp' is a standard usage in the community
...     print("Using:", xp.__name__)
...     return xp.maximum(0, x) + xp.log1p(xp.exp(-abs(x)))

如果需要同时操作 CPU 和 GPU 的 array 对象，可能需要将不同的 array 对象显式移动到相同的设备上（ CPU 或者 GPU ），mcPy提供 cupy.asnumpy() 和 cupy.asarray()，可以实现上述功能。代码示例如下：

>>> x_cpu = np.array([1, 2, 3])
>>> y_cpu = np.array([4, 5, 6])
>>> x_cpu + y_cpu
array([5, 7, 9])
>>> x_gpu = cp.asarray(x_cpu)
>>> x_gpu + y_cpu
Traceback (most recent call last):
...
TypeError: Unsupported type <class 'numpy.ndarray'>
>>> cp.asnumpy(x_gpu) + y_cpu
array([5, 7, 9])
>>> cp.asnumpy(x_gpu) + cp.asnumpy(y_cpu)
array([5, 7, 9])
>>> x_gpu + cp.asarray(y_cpu)
array([5, 7, 9])
>>> cp.asarray(x_gpu) + cp.asarray(y_cpu)
array([5, 7, 9])

cupy.asnumpy() 方法返回 NumPy array（ host 上的 array ），对应的 cupy.asarray() 方法返回 mcPy array（当前设备上的 array）。两个方法都接受不限制来源的输入，即：这两个方法接受任何可以转换为 array 对象的输入，而不限制输入对象位于 host 还是device。

3.2. 用户自定义内核

mcPy 支持自定义三种不同的内核：elementwise 内核、reduction 内核和原始内核。本章节描述如何定义并调用上述各种内核。

3.2.1. Elementwise 内核基本概念

可通过 ElementwiseKernel 类定义 elementwise 内核，使用该类实例的 __call__ 方法调用该实例定义的 GPU 内核。

Elementwise 内核的定义由 4 个部分构成：输入参数列表、输出参数列表、一段主体循环代码和一个内核名称。以下代码演示了如何计算 \(f(x,y)=(x-y)^2\) 的内核：

>>> squared_diff = cp.ElementwiseKernel(
...    'float32 x, float32 y',
...    'float32 z',
...    'z = (x - y) * (x - y)',
...    'squared_diff')

参数列表由一个或者多个参数定义（逗号分隔）构成，每个参数定义由一个 类型标识 和一个 参数名称 构成。NumPy 数据类型的名称可以用来做参数定义中的类型标识。

备注

n、i 以及下划线 _ 开头的变量名预留为内部专用。

可以在标量和向量上调用 elementwise 内核，并且支持广播：

>>> x = cp.arange(10, dtype=np.float32).reshape(2, 5)
>>> y = cp.arange(5, dtype=np.float32)
>>> squared_diff(x, y)
array([[ 0., 0., 0., 0., 0.],
    [25., 25., 25., 25., 25.]], dtype=float32)
>>> squared_diff(x, 5)
array([[25., 16., 9., 4., 1.],
    [ 0., 1., 4., 9., 16.]], dtype=float32)

可以显式指定输出参数（输入参数的旁边）：

>>> z = cp.empty((2, 5), dtype=np.float32)
>>> squared_diff(x, y, z)
array([[ 0., 0., 0., 0., 0.],
       [25., 25., 25., 25., 25.]], dtype=float32)

3.2.2. 通用类型的内核

如果类型标识是单个字符，该单个字符为类型占位符（type placeholder），通过这种方式可以定义通用类型的内核。上述 squared_diff 内核可以修改为通用类型的内核：

>>> squared_diff_generic = cp.ElementwiseKernel(
...     'T x, T y',
...     'T z',
...     'z = (x - y) * (x - y)',
...     'squared_diff_generic')

内核定义中相同字符的类型占位符表示相同的类型。这些占位符的实际类型由实际的参数类型决定。 ElementwiseKernel 类首先检查输出参数，然后检查输入参数，以确定实际类型。如果在内核调用中没有给出输出参数，那么只使用输入参数来确定类型。

类型占位符可以在循环体代码中使用：

>>> squared_diff_generic = cp.ElementwiseKernel(
...     'T x, T y',
...     'T z',
...     '''
...         T diff = x - y;
...         z = diff * diff;
...     ''',
...     'squared_diff_generic')

内核定义中可以使用多种类型占位符。例如，上面的 squared_diff_generic 内核可以使用多种类型占位符而进一步泛化：

>>> squared_diff_super_generic = cp.ElementwiseKernel(
...     'X x, Y y',
...     'Z z',
...     'z = (x - y) * (x - y)',
...     'squared_diff_super_generic')

请注意，这个内核需要显式指定输出参数的类型，因为 Z 类型不能从输入参数中自动确定。

3.2.3. 原始参数标识符

ElementwiseKernel 类自动通过广播进行索引，这对于定义大多数 elementwise 计算都很有用。但有些情况下，用户需要为一些参数编写一个带有手动索引的内核，此时可以通过在类型说明符之前添加 raw 关键字来告诉 ElementwiseKernel 类使用手动索引。

可以使用特殊变量 i 和方法 _ind.size() 进行手动索引。 i 表示循环中的索引。 _ind.size() 表示要应用 elementwise 操作的元素总数， _ind.size() 表示广播操作后的大小。

例如，将两个向量中的一个向量反转，再把两个向量相加的内核函数可以编写为：

>>> add_reverse = cp.ElementwiseKernel(
...     'T x, raw T y', 'T z',
...     'z = x + y[_ind.size() - i - 1]',
...     'add_reverse')

以上示例是刻意为之的，其实无需定义新的核函数，可以通过 z = x + y[::-1] 实现。一个被 raw 修饰的参数可以作为一个 array 。索引运算符 y[_ind.size() - i - 1] 涉及对 y 的索引计算，因此 y 可以是任意形状和步长的。

广播中不涉及被 raw 修饰的参数。如果要将所有参数标记为 raw 参数，则必须在调用时指定 size 参数，该参数定义 _ind.size() 的值。

3.2.4. Reduction 内核

可以通过 ReductionKernel 类定义 reduction 内核。可以通过定义内核代码的四个部分来使用它：

  1. 标识值（Identity value）：reduction操作接口的初始值

  2. 映射表达式（Mapping expression）：在每个元素归约之前应用该表达式进行转换计算

  3. 归约表达式（Reduction expression）：进行归约操作的算子表达式，内置的特殊变量 a 和 b 用于表达式的操作数

  4. 映射后表达式（Post mapping expression）：用于对归约操作的结果进行数据转换，内置特殊变量 a 作为其输入。该表达式中的输出参数应与 ReductionKernel 类定义中的输出参数名称一致

ReductionKernel 类自动插入高效灵活的 reduction 实现所需的其他代码片段。

例如，在指定轴上计算 L2 范数的内核可以编写为：

>>> l2norm_kernel = cp.ReductionKernel(
...     'T x', # input params
...     'T y', # output params
...     'x * x', # map
...     'a + b', # reduce
...     'y = sqrt(a)', # post-reduction map
...     '0', # identity value
...     'l2norm' # kernel name
...     )
>>> x = cp.arange(10, dtype=np.float32).reshape(2, 5)
>>> l2norm_kernel(x, axis=1)
array([ 5.477226 , 15.9687195], dtype=float32)

备注

如果要在 reduction 内核中使用 raw 说明符，则需要遵循一些约束和限制：要求执行归约的轴必须位于 shape 的头部。即，如果想使用 raw 说明符修饰一个或多个参数，那么归约的 axis 参数必须是 0 或从 0 开始的连续递增整数序列，如 (0, 1)、 (0, 1, 2) 等。

3.2.5. 原始内核

可以通过 RawKernel 类定义原始内核。通过使用原始内核，可以从原始的 C kernel 代码定义内核。

RawKernel 对象允许使用 mcLaunchKernel 接口调用内核。即：可以控制网格大小，块大小，共享内存大小和流。

>>> add_kernel = cp.RawKernel(r'''
... extern "C" __global__
... void my_add(const float* x1, const float* x2, float* y) {
...     int tid = blockDim.x * blockIdx.x + threadIdx.x;
...     y[tid] = x1[tid] + x2[tid];
... }
... ''', 'my_add')
>>> x1 = cp.arange(25, dtype=cp.float32).reshape(5, 5)
>>> x2 = cp.arange(25, dtype=cp.float32).reshape(5, 5)
>>> y = cp.zeros((5, 5), dtype=cp.float32)
>>> add_kernel((5,), (5,), (x1, x2, y)) # grid, block and arguments
>>> y
array([[ 0., 2., 4., 6., 8.],
       [10., 12., 14., 16., 18.],
       [20., 22., 24., 26., 28.],
       [30., 32., 34., 36., 38.],
       [40., 42., 44., 46., 48.]], dtype=float32)

还可以创建在复数数组上运行的原始内核：

>>> complex_kernel = cp.RawKernel(r'''
... #include <cupy/complex.cuh>
... extern "C" __global__
... void my_func(const complex<float>* x1, const complex<float>* x2,
...              complex<float>* y, float a) {
...     int tid = blockDim.x * blockIdx.x + threadIdx.x;
...     y[tid] = x1[tid] + a * x2[tid];
... }
... ''', 'my_func')
>>> x1 = cupy.arange(25, dtype=cupy.complex64).reshape(5, 5)
>>> x2 = 1j*cupy.arange(25, dtype=cupy.complex64).reshape(5, 5)
>>> y = cupy.zeros((5, 5), dtype=cupy.complex64)
>>> complex_kernel((5,), (5,), (x1, x2, y, cupy.float32(2.0))) # grid, block and arguments
>>> y
array([[ 0. +0.j, 1. +2.j, 2. +4.j, 3. +6.j, 4. +8.j],
       [ 5.+10.j, 6.+12.j, 7.+14.j, 8.+16.j, 9.+18.j],
       [10.+20.j, 11.+22.j, 12.+24.j, 13.+26.j, 14.+28.j],
       [15.+30.j, 16.+32.j, 17.+34.j, 18.+36.j, 19.+38.j],
       [20.+40.j, 21.+42.j, 22.+44.j, 23.+46.j, 24.+48.j]],
       dtype=complex64)

虽然鼓励对复数使用 complex<T> 类型（如上所示，可通过引入 <cupy/complex.cuh> 来获得），但对于已经使用 mcComplex.h 中函数编写的 MXMACA 代码，无需自己进行转换：只需在创建 RawKernel 实例时设置 translate_cuComplex=True 选项即可。

MXMACA 内核属性可以通过访问属性字典或直接访问 RawKernel 对象的属性来检索；后者也可以用于设置某些属性：

>>> add_kernel = cp.RawKernel(r'''
... extern "C" __global__
... void my_add(const float* x1, const float* x2, float* y) {
...     int tid = blockDim.x * blockIdx.x + threadIdx.x;
...     y[tid] = x1[tid] + x2[tid];
... }
... ''', 'my_add')
>>> add_kernel.attributes
{'max_threads_per_block': 1024, 'shared_size_bytes': 0, 'const_size_bytes': 0, 'local_size_bytes': 0, 'num_regs': 10, 'ptx_version': 70, 'binary_version': 70, 'cache_mode_ca': 0, 'max_dynamic_shared_size_bytes': 49152, 'preferred_shared_memory_carveout': -1}
>>> add_kernel.max_dynamic_shared_size_bytes
49152
>>> add_kernel.max_dynamic_shared_size_bytes = 50000 # set a new value for the attribute
>>> add_kernel.max_dynamic_shared_size_bytes
50000

RawKernel支持动态并行。只需要为 RawKernel 的 options 参数提供链接标志（如 -dc ）。

如果内核依赖于 C++ 标准库标头，如 <type_traits>，那么很可能会遇到编译错误。在这种情况下，请尝试在创建RawKernel实例时通过设置 jitify=True 来启用 mcPy 的 Jitify 支持。 它提供了基本的 C++ 标准支持来纠正常见错误。

备注

  • 原始内核没有返回值。需要将输入数组和输出数组作为参数进行传递。

  • 在内核代码中使用 printf() 时，可能需要同步流以查看输出。如果使用默认流，则可以使用 cupy.cuda.Stream.null.synchronize()。

  • 在上面的所有示例中，在 extern "C" 块中声明内核，表明使用了 C 链接。这是为了确保内核名称不会被破坏，以便可以按名称检索它们。

3.2.6. 内核参数

Python 基本类型和 NumPy 标量将通过值传递给内核。内核中的 Array 参数（以指针形式传参），必须保证指针指向 mcPy ndarray 类型数组。 mcPy 不会对传递给内核的参数（包括参数的类型和数量）进行验证。

特别要注意的是，在传递 mcPy ndarray 时，它的 dtype 应该与 C 源代码的函数定义中声明的参数的类型匹配（除非用户有意转换数组）。

例如， cupy.float32 和 cupy.uint64 数组必须分别传递给类型为 float* 和无符号 long long* 的参数。 mcPy 不直接支持 float3* 等非基元类型的数组，但不禁止在内核中将 float* 或 void* 强制转换为 float3*。

Python基本数据类型 int， float， complex， bool 分别与 long long， double， cuDoubleComplex， bool 一一对应。

Size 都为 1 的 NumPy 标量和 NumPy 向量以值传递的方式传入内核。 这意味着可以通过值传递任何基本 NumPy 类型，如 numpy.int8 或 numpy.float64，前提是内核参数的大小匹配。 匹配 mcPy/NumPy 数据类型和 C 类型，可参见表 3.1。

表 3.1 mcPy/NumPy数据类型和C类型

mcPy/NumPy类型

C类型

字节数 (bytes)

bool

bool

1

int8

char, signed char

1

int16

short, signed short

2

int32

int, signed int

4

int64

long long, signed long long

8

uint8

unsigned char

1

uint16

unsigned short

2

uint32

unsigned int

4

uint64

unsigned long long

8

float16

half

2

float32

float

4

float64

double

8

complex64

float2, cuFloatComplex, complex<float>

8

MXMACA 实现机制保证 host 和 device 上基本类型的大小始终匹配。然而 size_t、 ptrdiff_t、 intptr_t、 uintptr_t、 long、 signed long 和 unsigned long 的项目大小取决于平台。 要传递任何内置的向量类型（如 float3 或者用户自定义的向量结构）作为内核参数（前提是它与设备端内核参数类型匹配），请参见 3.2.7 用户自定义数据类型。

3.2.7. 用户自定义数据类型

通过自定义 NumPy 数据类型，可以使用自定义类型（复合类型，例如 struct ，或包含 struct 的 struct） 作为内核参数。在执行此操作时，需要匹配 host 和 device 结构的内存布局。MXMACA 实现机制保证 host 和 device 上基本类型的大小始终匹配。然而，它可能要求 device 内存对齐。这意味着，对于复合类型，结构成员偏移量可能与预期的不同。

当内核参数按值传递时，MXMACA驱动程序将从 NumPy 对象数据指针的开头开始精确复制 sizeof(param_type) 个字节，其中 param_type 是内核中的参数类型。必须通过定义相应的 NumPy dtypes 来匹配 param_type 的内存布局，例如：大小、对齐和结构填充/打包。

对于内置数据类型如 int2 和 double4，以及其他具有命名成员的压缩结构，可以直接定义这样的 NumPy dtypes ：

>>> import numpy as np
>>> names = ['x', 'y', 'z']
>>> types = [np.float32]*3
>>> float3 = np.dtype({'names': names, 'formats': types})
>>> arg = np.random.rand(3).astype(np.float32).view(float3)
>>> print(arg)
[(0.9940819, 0.62873816, 0.8953669)]
>>> arg['x'] = 42.0
>>> print(arg)
[(42., 0.62873816, 0.8953669)]

这里的 arg 可以直接用作内核参数。当不需要命名字段时，用户可能更喜欢使用此语法来定义压缩结构，如向量或矩阵：

>>> import numpy as np
>>> float5x5 = np.dtype({'names': ['dummy'], 'formats': [(np.float32,(5,5))]})
>>> arg = np.random.rand(25).astype(np.float32).view(float5x5)
>>> print(arg.itemsize)
100

这里 arg 表示一个 100 字节的标量（即大小为 1 的 NumPy 数组），可以通过值传递给任何内核。内核参数在一个专用的 4kB 内存组中按值传递，该内存组有自己的带广播的缓存。因此，内核参数总大小的上限为4kB。需要注意的是，此专用内存组未与设备 __constant__ 内存空间共享。

目前，mcPy 没有提供任何辅助例程来创建用户定义的复合类型。然而，这样的复合类型可以使用 NumPy dtype 偏移量和 itemsize 功能递归构建。

备注

不能使用 arg[N] 语法类型直接将静态数组作为内核参数传递，其中 N 是编译时常数。 __global__void kernel(float arg[5]) 的函数定义被编译器视为 __global__valid kernel(float*arg)。如果要按值将五个浮点值传递给内核，则需要定义一个自定义结构 struct float5 { float val[5]; }; 并将内核签名修改为 __global__void kernel(float5 arg)。

3.2.8. 原始modules

对于处理大型原始 MXMACA 源码或加载现有的曦云系列二进制文件， RawModule 类可能更方便。 RawModule 类可以通过 MXMACA 源代码或曦云系列二进制文件的路径进行初始化，接受 RawKernel 中的大多数参数。 然后可以通过调用 get_function() 方法来检索所需的内核，该方法返回一个 RawKernel 实例，该实例可以以 3.2.5 原始内核 所述的方式被调用。

>>> loaded_from_source = r'''
... extern "C"{
...
... __global__ void test_sum(const float* x1, const float* x2, float* y, \
...                          unsigned int N)
... {
...     unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;
...     if (tid < N)
...     {
...         y[tid] = x1[tid] + x2[tid];
...     }
... }
...
... __global__ void test_multiply(const float* x1, const float* x2, float* y, \
...                               unsigned int N)
... {
...     unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;
...     if (tid < N)
...     {
...         y[tid] = x1[tid] * x2[tid];
...     }
... }
...
... }'''

>>> module = cp.RawModule(code=loaded_from_source)
>>> ker_sum = module.get_function('test_sum')
>>> ker_times = module.get_function('test_multiply')
>>> N = 10
>>> x1 = cp.arange(N**2, dtype=cp.float32).reshape(N, N)
>>> x2 = cp.ones((N, N), dtype=cp.float32)
>>> y = cp.zeros((N, N), dtype=cp.float32)
>>> ker_sum((N,), (N,), (x1, x2, y, N**2)) # y = x1 + x2
>>> assert cp.allclose(y, x1 + x2)
>>> ker_times((N,), (N,), (x1, x2, y, N**2)) # y = x1 * x2
>>> assert cp.allclose(y, x1 * x2)

上面关于在 RawKernel 中使用复数的说明也适用于 RawModule。

对于需要访问全局符号（如常量内存）的内核，可以使用 get_global() 方法。

为了支持 C++ 模板内核， RawModule 还提供了一个“名称表达式”(name_expressions) 参数。通过名称表达式参数可以指定一个模板特化（template specializations）列表，以便 mcPy 可以按类型生成和检索相应的内核：

>>> code = r'''
... template<typename T>
... __global__ void fx3(T* arr, int N) {
...     unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x;
...     if (tid < N) {
...         arr[tid] = arr[tid] * 3;
...     }
... }
... '''
>>>
>>> name_exp = ['fx3<float>', 'fx3<double>']
>>> mod = cp.RawModule(code=code, options=('-std=c++11',),
...     name_expressions=name_exp)
>>> ker_float = mod.get_function(name_exp[0]) # compilation happens here
>>> N=10
>>> a = cp.arange(N, dtype=cp.float32)
>>> ker_float((1,), (N,), (a, N))
>>> a
array([ 0., 3., 6., 9., 12., 15., 18., 21., 24., 27.], dtype=float32)
>>> ker_double = mod.get_function(name_exp[1])
>>> a = cp.arange(N, dtype=cp.float64)
>>> ker_double((1,), (N,), (a, N))
>>> a
array([ 0., 3., 6., 9., 12., 15., 18., 21., 24., 27.])

用于初始化 RawModule 实例和检索内核的“名称表达式”中使用的内核名称（例如上述示例中的 fx3）是原始（未混淆的）内核名称，并且“名称表达式”中需要明确指定了所有模板参数。

3.2.9. 内核融合

cupy.fuse() 是一个融合函数的装饰器。这个装饰器可以用于定义 elementwise 或 reduction 内核，比 ElementwiseKernel 或 ReductionKernel 更容易使用。

通过使用这个装饰器，可以按照如下代码的方式定义 squared_diff 内核：

>>> @cp.fuse()
... def squared_diff(x, y):
...     return (x - y) * (x - y)

上面的内核可以在标量、NumPy 数组或 mcPy 数组上调用，就像原始函数一样。

>>> x_cp = cp.arange(10)
>>> y_cp = cp.arange(10)[::-1]
>>> squared_diff(x_cp, y_cp)
array([81, 49, 25, 9, 1, 1, 9, 25, 49, 81])
>>> x_np = np.arange(10)
>>> y_np = np.arange(10)[::-1]
>>> squared_diff(x_np, y_np)
array([81, 49, 25, 9, 1, 1, 9, 25, 49, 81])

在第一次调用函数时，融合函数基于参数的抽象信息（例如，它们的 dtypes 和 ndim ）分析原始函数，并创建和缓存实际的内核。从具有相同输入类型的第二次函数调用开始，融合函数调用先前缓存的内核，因此强烈建议重用相同的修饰函数，而不是修饰多次定义的局部函数。

cupy.fuse() 还支持简单 reduction 内核。

>>> @cp.fuse()
... def sum_of_products(x, y):
...     return cp.sum(x * y, axis = -1)

可以使用 kernel_name 关键字参数指定内核名称：

>>> @cp.fuse(kernel_name='squared_diff')
... def squared_diff(x, y):
...     return (x - y) * (x - y)

备注

目前， cupy.fuse() 只能融合简单的 elementwise 操作和 reduction 操作。不支持其他大多数例程，例如 cupy.matmul()、 cupy.reshape()。

3.2.10. JIT 编译内核

cupyx.jit.rawkernel 装饰器可以从 Python 函数创建原始内核。

在本节中，用装饰器包装的 Python 函数被称为 目标函数 。

目标函数由基本标量运算组成，用户必须在目标函数的代码中手动实现并行加速。mcPy 的数组运算，因为会自动进行并行运算，而不被支持，例如 add()、 sum()。 如果需要基于此类数组函数自定义内核，参见 3.2.9 内核融合。

3.2.10.1. 基本使用方式

下面是一个简短的例子，说明如何编写一个 cupyx.jit.rawkernel，使用 grid-stride 循环将值从 x 复制到 y：

>>> from cupyx import jit
>>>
>>> @jit.rawkernel()
... def elementwise_copy(x, y, size):
...     tid = jit.blockIdx.x * jit.blockDim.x + jit.threadIdx.x
...     ntid = jit.gridDim.x * jit.blockDim.x
...     for i in range(tid, size, ntid):
...         y[i] = x[i]

>>> size = cupy.uint32(2 ** 22)
>>> x = cupy.random.normal(size=(size,), dtype=cupy.float32)
>>> y = cupy.empty((size,), dtype=cupy.float32)

>>> elementwise_copy((128,), (1024,), (x, y, size)) # RawKernel style
>>> assert (x == y).all()

如上所示，启动内核的调用参数列表中，前两项分别是网格大小和块大小。 grid （形如(128,)）是网格的大小， block （形如(1024,)） 是每个线程块的尺寸。 在具有预定网格大小和块大小的 GPU 上启动 MXMACA 内核，需要对 MXMACA 编程模型有基本了解，详细信息可参见《曦云® 系列通用计算 GPU 快速上手指南》。

编译将推迟到第一次函数调用。mcPy 的 JIT 编译器在调用时推断参数的类型，并缓存编译后的内核以加快后续任意调用。

3.2.10.2. 基本设计原则

mcPy 的 JIT 编译器通过 Python AST 生成 device 内核代码。为了避免性能下降，mcPy中没有使用 Python 字节码来生成内核代码。在将目标函数转换为字节码时，目标函数的 for 循环和其他控制语句会完全转换为跳转指令，所以由 Python 字节码生成的 device 内核源代码不会被 device 编译器有效地优化。

3.2.10.3. 类型规则

局部变量的类型是在函数的第一次赋值时推断出来的。第一次赋值必须在函数体内的顶层实现，即，它不能在 if / else 代码块或 for 循环中。

3.2.10.4. 使用限制

JIT 在 Python 的交互式解释器（REPL）中不起作用，因为编译器需要获取目标函数的源代码。

3.3. 快速傅里叶变换

mcPy（ cupy.fft ）涵盖了 NumPy 中提供的完整快速傅立叶变换（FFT）功能。除了那些可以按原样使用的高级 API 之外，mcPy 还提供了额外的特性：

  • 提供 mcFFT 库为 GPU 提供的额外例程

  • 更好地控制 FFT 例程的性能和行为

其中一些功能是实验性的（可能会更改、弃用或删除）。

3.3.1. 用户管理的 FFT 计划

出于性能原因，用户可能希望自己创建、重用和管理 FFT 计划。mcPy 为此提供了一个高级的实验性 API get_fft_plan()。用户指定要执行的 FFT 变换，并将根据输入生成计划。

import cupy as cp
from cupyx.scipy.fft import get_fft_plan

a = cp.random.random((4, 64, 64)).astype(cp.complex64)
plan = get_fft_plan(a, axes=(1, 2), value_type='C2C') # for batched, C2C, 2D transform

返回的计划可以显式用作 cupyx.scipy.fft API 的参数：

import cupyx.scipy.fft

# the rest of the arguments must match those used when generating the plan
out = cupyx.scipy.fft.fft2(a, axes=(1, 2), plan=plan)

或作为 cupy.fft API的上下文管理器：

with plan:
    # the arguments must match those used when generating the plan
    out = cp.fft.fft2(a, axes=(1, 2))

3.3.2. FFT 计划缓存

然而，在某些情况下，用户可能不想自己管理 FFT 计划。此外，计划也可以在 mcPy 的例程中进行内部重用，用户管理的计划不适用于此。因此，mcPy 提供了一个内置的计划缓存，默认情况下是启用的。计划缓存以每个设备、每个线程为基础进行，可以通过 get_plan_cache() API检索。

>>> import cupy as cp
>>>
>>> cache = cp.fft.config.get_plan_cache()
>>> cache.show_info()
------------------- cuFFT plan cache (device 0) -------------------
cache enabled? True
current / max size : 0 / 16 (counts)
current / max memsize: 0 / (unlimited) (bytes)
hits / misses: 0 / 0 (counts)

cached plans (most recently used first):

>>> # perform a transform, which would generate a plan and cache it
>>> a = cp.random.random((4, 64, 64))
>>> out = cp.fft.fftn(a, axes=(1, 2))
>>> cache.show_info() # hit = 0
------------------- cuFFT plan cache (device 0) -------------------
cache enabled? True
current / max size : 1 / 16 (counts)
current / max memsize: 262144 / (unlimited) (bytes)
hits / misses: 0 / 1 (counts)

cached plans (most recently used first):
key: ((64, 64), (64, 64), 1, 4096, (64, 64), 1, 4096, 105, 4, 'C', 2, None), plan type: PlanNd, memory usage: 262144

>>> # perform the same transform again, the plan is looked up from cache and reused
>>> out = cp.fft.fftn(a, axes=(1, 2))
>>> cache.show_info() # hit = 1
------------------- cuFFT plan cache (device 0) -------------------
cache enabled? True
current / max size : 1 / 16 (counts)
current / max memsize: 262144 / (unlimited) (bytes)
hits / misses: 1 / 1 (counts)

cached plans (most recently used first):
key: ((64, 64), (64, 64), 1, 4096, (64, 64), 1, 4096, 105, 4, 'C', 2, None), plan type: PlanNd, memory usage: 262144
>>> # clear the cache
>>> cache.clear()
>>> cp.fft.config.show_plan_cache_info() # = cache.show_info(), for all devices
=============== cuFFT plan cache info (all devices) ===============
------------------- cuFFT plan cache (device 0) -------------------
cache enabled? True
current / max size : 0 / 16 (counts)
current / max memsize: 0 / (unlimited) (bytes)
hits / misses: 0 / 0 (counts)

cached plans (most recently used first):

返回的 PlanCache 对象有其他方法进行更精细的控制，例如设置缓存大小（按计数或按内存使用情况）。如果大小设置为0，则缓存将被禁用。

如上所示，每个 FFT 计划都分配了一个相关的工作区域。如果发生内存不足错误，可能需要检查、清除或限制计划缓存。

备注

不缓存 get_fft_plan() 返回的计划。

3.3.3. FFT 回调

mcFFT 提供 FFT 回调，用于将预处理和/或后处理内核与 FFT 例程合并，从而减少对全局存储器的访问。此功能得到了 mcPy 的实验性支持。用户需要提供自定义的加载和/或将内核存储为字符串，并通过 set_cufft_callbacks() 设置上下文管理器。请注意，加载（存储）内核指针必须命名为 d_loadCallbackPtr （ d_storeCallbackPtr ）。

import cupy as cp

# a load callback that overwrites the input array to 1
code = r'''
__device__ cufftComplex CB_ConvertInputC(
    void *dataIn,
    size_t offset,
    void *callerInfo,
    void *sharedPtr)
{
    cufftComplex x;
    x.x = 1.;
    x.y = 0.;
    return x;
}
__device__ cufftCallbackLoadC d_loadCallbackPtr = CB_ConvertInputC;
'''

a = cp.random.random((64, 128, 128)).astype(cp.complex64)

# this fftn call uses callback
with cp.fft.config.set_cufft_callbacks(cb_load=code):
    b = cp.fft.fftn(a, axes=(1,2))

# this does not use
c = cp.fft.fftn(cp.ones(shape=a.shape, dtype=cp.complex64), axes=(1,2))

# result agrees
assert cp.allclose(b, c)

# "static" plans are also cached, but are distinct from their no-callback counterparts
cp.fft.config.get_plan_cache().show_info()

备注

在内部，此功能需要为每对不同的加载和存储内核重新编译一个 Python 模块。因此，第一次调用将非常缓慢，如果回调可以在后续计算中重复使用，则此成本将摊销。编译后的模块缓存在磁盘上，默认位置为 $HOME/.cupy/callback_cache，这个路径可以通过环境变量 CUPY_CACHE_DIR 覆盖。

3.3.4. 多 GPU FFT

mcPy 目前为多 GPU FFT 提供了两种 实验性 支持。

备注

使用多个 GPU 来执行 FFT 并不能保证更高的性能。如果变换适合 1 个 GPU，则应避免使用多个。

第一种支持是使用高级的 fft() 和 ifft() API，这需要输入数组驻留在其中一个参与的 GPU 上。多 GPU 计算是在底层进行的，计算结束后，计算结果位于进行计算的 device 上。目前仅支持 1D 的复数到复数（C2C）变换；不支持复数到实数（C2R）或实数到复数（R2C）转换，如 rfft()。可以对转换进行批处理（batch size >1），也可以不进行批处理。

import cupy as cp

cp.fft.config.use_multi_gpus = True
cp.fft.config.set_cufft_gpus([0, 1]) # use GPU 0 & 1
shape = (64, 64) # batch size = 64
dtype = cp.complex64
a = cp.random.random(shape).astype(dtype) # reside on GPU 0

b = cp.fft.fft(a) # computed on GPU 0 & 1, reside on GPU 0

如果需要执行 2D/3D 变换（例如： fftn()）而不是 1D 变换（例如： fft()）， 2D/3D 变换也可以在多 GPU 上执行，但在这种情况下，在底层沿变换轴以循环方式完成计算（和 NumPy 的算法完全一致），这可能会导致次优性能。

第二种用法是使用低级的私有 mcPy API。需要构造一个 Plan1d 对象，并像使用 mcFFT 在C/C++中编程一样使用它。使用这种方法，输入 array 可以作为 numpy.ndarray 驻留在 host 上，因此其大小可以比单个 GPU 所能容纳的大得多，这是运行多GPU FFT的主要原因之一。

import numpy as np
import cupy as cp

# no need to touch cp.fft.config, as we are using low-level API

shape = (64, 64)
dtype = np.complex64
a = np.random.random(shape).astype(dtype) # reside on CPU

if len(shape) == 1:
    batch = 1
    nx = shape[0]
elif len(shape) == 2:
    batch = shape[0]
    nx = shape[1]

# compute via cuFFT
cufft_type = cp.cuda.cufft.CUFFT_C2C # single-precision c2c
plan = cp.cuda.cufft.Plan1d(nx, cufft_type, batch, devices=[0,1])
out_cp = np.empty_like(a) # output on CPU
plan.fft(a, out_cp, cufft.CUFFT_FORWARD)

out_np = numpy.fft.fft(a) # use NumPy's fft
# np.fft.fft alway returns np.complex128
if dtype is numpy.complex64:
    out_np = out_np.astype(dtype)

# check result
assert np.allclose(out_cp, out_np, rtol=1e-4, atol=1e-7)

备注

如果通过高级 API 自动生成，则会缓存多 GPU 计划，但如果通过低级 API 手动生成，则不会缓存。

3.3.5. 半精度 FFT

mcFFT 提供了 mcfftXtMakePlanMany 和 mcfftXtExec 例程，以支持广泛的 FFT 需求，包括 64 位索引和半精度 FFT。mcPy 通过新的（尽管是私有的） XtPlanNd API 为该功能提供了实验支持。对于半精度 FFT，在支持的硬件上，其速度可能是单精度 FFT 的两倍。不过，NumPy 还没有为半精度复数（即 numpy.complex32 ）提供必要的基础设施，因此该功能的实现步骤比一般情况更复杂。

import cupy as cp
import numpy as np

shape = (1024, 256, 256) # input array shape
idtype = odtype = edtype = 'E' # = numpy.complex32 in the future

# store the input/output arrays as fp16 arrays twice as long, as complex32 is not yet available
a = cp.random.random((shape[0], shape[1], 2*shape[2])).astype(cp.float16)
out = cp.empty_like(a)

# FFT with mcFFT
plan = cp.cuda.cufft.XtPlanNd(shape[1:],
                              shape[1:], 1, shape[1]*shape[2], idtype,
                              shape[1:], 1, shape[1]*shape[2], odtype,
                              shape[0], edtype,
                              order='C', last_axis=-1, last_size=None)

plan.fft(a, out, cp.cuda.cufft.CUFFT_FORWARD)

# FFT with NumPy
a_np = cp.asnumpy(a).astype(np.float32) # upcast
a_np = a_np.view(np.complex64)
out_np = np.fft.fftn(a_np, axes=(-2,-1))
out_np = np.ascontiguousarray(out_np).astype(np.complex64) # downcast
out_np = out_np.view(np.float32)
out_np = out_np.astype(np.float16)

# don't worry about accruacy for now, as we probably lost a lot during casting
print('ok' if cp.mean(cp.abs(out - cp.asarray(out_np))) < 0.1 else 'not ok')

3.4. 内存管理

默认情况下，mcPy 使用 内存池 进行内存分配。内存池通过减轻内存分配和 CPU/GPU 同步的开销，显著提高了性能。

mcPy 中有两个不同的内存池：

  • 设备内存池（GPU设备内存），用于 GPU 内存分配。

  • 固页内存池（non-swappable CPU内存），用于 CPU 到 GPU 的数据传输。

备注

监视内存使用情况时（例如，对 GPU 内存使用 mx-smi 或对 CPU 内存使用 ps ），即使在 array 实例超出范围后，内存也不会被释放。这是一种预期的行为，因为默认内存池“缓存”分配的内存块。

为了更方便地使用固页内存，在 cupyx 中还提供了一些高级API，包括 cupyx.empty_pined()、 cupyx.empty_like_pinned()、 cupyx.zeros_pined() 和 cupyx.zeros_like_pinned()。它们返回由固页内存支持的 NumPy array。如果正在使用 mcPy 的固页内存池，则从该池中分配固定内存。

备注

mcPy 提供了 FFT 计划高速缓存，如果使用 FFT 和相关功能，则该高速缓存会占用一部分设备内存。可以通过收缩或禁用缓存来释放占用的设备内存。

3.4.1. 内存池操作

内存池实例提供有关内存分配的统计信息。要访问默认内存池实例，请使用 cupy.get_default_memory_pool() 和 cupy.get_default_pinned_memory_pool()。还可以释放内存池中所有未使用的内存块。代码示例如下：

import cupy
import numpy

mempool = cupy.get_default_memory_pool()
pinned_mempool = cupy.get_default_pinned_memory_pool()

# Create an array on CPU.
# NumPy allocates 400 bytes in CPU.
a_cpu = numpy.ndarray(100, dtype=numpy.float32)
print(a_cpu.nbytes) # 400

# You can access statistics of these memory pools.
print(mempool.used_bytes()) # 0
print(mempool.total_bytes()) # 0
print(pinned_mempool.n_free_blocks()) # 0

# Transfer the array from CPU to GPU.
# This allocates 400 bytes from the device memory pool, and another 400
# bytes from the pinned memory pool. The allocated pinned memory will be
# released just after the transfer is complete. Note that the actual
# allocation size may be rounded to larger value than the requested size
# for performance.
a = cupy.array(a_cpu)
print(a.nbytes) # 400
print(mempool.used_bytes()) # 512
print(mempool.total_bytes()) # 512
print(pinned_mempool.n_free_blocks()) # 1

# When the array goes out of scope, the allocated device memory is released
# and kept in the pool for future reuse.
a = None # (or `del a`)
print(mempool.used_bytes()) # 0
print(mempool.total_bytes()) # 512
print(pinned_mempool.n_free_blocks()) # 1

# You can clear the memory pool by calling `free_all_blocks`.
mempool.free_all_blocks()
pinned_mempool.free_all_blocks()
print(mempool.used_bytes()) # 0
print(mempool.total_bytes()) # 0
print(pinned_mempool.n_free_blocks()) # 0

3.4.2. 限制显存使用量

可以使用 CUPY_GPU_MEMORY_LIMIT 环境变量来硬限制可以分配的 GPU 内存量。

# Set the hard-limit to 1 GiB:
#   $ export CUPY_GPU_MEMORY_LIMIT="1073741824"
# You can also specify the limit in fraction of the total amount of memory
# on the GPU. If you have a GPU with 2 GiB memory, the following is
# equivalent to the above configuration.
#   $ export CUPY_GPU_MEMORY_LIMIT="50%"

import cupy
print(cupy.get_default_memory_pool().get_limit()) # 1073741824

也可以使用 cupy.cuda.MemoryPool.set_limit() 设置限制（或覆盖通过环境变量指定的值）。 这样，可以为每个 GPU 设置不同的容量限制。

import cupy

mempool = cupy.get_default_memory_pool()

with cupy.cuda.Device(0):
    mempool.set_limit(size=1024**3) # 1 GiB

with cupy.cuda.Device(1):
    mempool.set_limit(size=2*1024**3) # 2 GiB

备注

MXMACA 在内存池之外分配一些 GPU 内存（如 MXMACA 上下文、库句柄等）。根据使用情况，此类内存可能需要一到几百 MB。这将不计入限额。

3.4.3. 变更内存池

可以使用自己的内存分配器而不是默认的内存池，方法是将内存分配函数传递给 cupy.cuda.set_allocator()/ cupy.cuda.set_pined_memory_allocattor()。内存分配器函数接受 1 个参数（内存请求大小以字节为单位），并返回 cupy.cuda.MemoryPointer/ cupy.cuda.PinnedMemoryPointer。

mcPy 提供了两个这样的分配器，用于在 GPU 上使用托管内存（managed memory）和流序内存（stream ordered memory）。要启用由托管内存支持的内存池，可以构造一个新的 MemoryPool 实例，其分配器设置为 malloc_managed()，如下所示：

import cupy

# Use managed memory
cupy.cuda.set_allocator(cupy.cuda.MemoryPool(cupy.cuda.malloc_managed).malloc)

如果不创建 MemoryPool 实例，而是直接将 malloc_managed() 传递给 set_allocator()，则当内存释放时，它将立即释放回系统。

mcPy 为流序内存分配器提供了一个实验接口。与 mcPy 的内存池类似，流序内存分配器也以流顺序的方式从内存池异步分配/释放内存。关键区别在于，它是 MXMACA 驱动程序中实现的内置功能，因此同一进程中的其他 MXMACA 应用程序可以轻松地从同一池中分配内存。

要启用管理流序内存的内存池，可以构造一个新的 MemoryAsyncPool 实例：

import cupy

# Use asynchronous stream ordered memory
cupy.cuda.set_allocator(cupy.cuda.MemoryAsyncPool().malloc)

# Create a custom stream
s = cupy.cuda.Stream()

# This would allocate memory asynchronously on stream s
with s:
    a = cupy.empty((100,), dtype=cupy.float64)

在这种情况下，不使用 MemoryPool 类。 MemoryAsyncPool 采用与 MemoryPool 不同的输入参数来指示要使用哪个池。

如果不创建 MemoryAsyncPool 实例，而是直接将 malloc_async() 传递给 set_allocator()，则将使用设备的当前内存池。

当使用流序内存时，需要开发者自己使用流和事件 API 来维护正确的流语义。在解除分配时，内存会在分配的流（第一次尝试）或任何当前 mcPy 流（第二次尝试）上异步释放。允许在释放分配给它的所有内存之前销毁分配内存的流。

此外，应用程序/库内部使用 mcMalloc （MXMACA的默认同步分配器）可能会与流序内存分配器发生意外的相互作用。具体来说，释放到内存池的内存可能不会立即对 mcMalloc 可见，从而导致潜在的内存不足错误。 在这种情况下，可以调用 free_all_blocks()，也可以手动执行（事件/流/设备）同步，然后重试。

目前 MemoryAsyncPool 接口是实验性的。特别是，虽然其 API 与 MemoryPool 的 API 基本相同，但由于 MXMACA 的限制，该池的一些方法需要足够新的驱动程序（当然，还需要支持的硬件、MXMACA版本和平台）。

可以通过以下代码禁用默认内存池。请确保在执行任何其他 mcPy 操作之前执行此操作。

import cupy

# Disable memory pool for device memory (GPU)
cupy.cuda.set_allocator(None)

# Disable memory pool for pinned memory (CPU).
cupy.cuda.set_pinned_memory_allocator(None)

3.5. 高性能最佳实践

本章节介绍了一些提高 mcPy 性能的技巧和建议。

3.5.1. 基准测试

在尝试优化代码之前，首先确定性能瓶颈是非常重要的。为了帮助设置基线基准，mcPy提供了一个有用的实用程序 cupyx.profiler.benchmark()，用于统计 Python 函数在 CPU 和 GPU 上的运行时间：

>>> from cupyx.profiler import benchmark
>>>
>>> def my_func(a):
...     return cp.sqrt(cp.sum(a**2, axis=-1))
...
>>> a = cp.random.random((256, 1024))
>>> print(benchmark(my_func, (a,), n_repeat=20))
my_func : CPU: 44.407 us +/- 2.428 (min: 42.516 / max: 53.098) us GPU-0: 181.565 us +/- 1.853 (min: 180.288 / max: 188.608) us

由于 GPU 与 CPU 是异步运行，因此 GPU 编程中的一个常见缺陷是错误地使用 CPU 计时实用程序（如 Python 标准库中的 time.perf_counter() 或 IPython 中的 %timeit ）来测量所用时间，而这些程序对 GPU 的运行时间一无所知。 cupyx.profiler.benchmark() 通过在要测量的函数前后在当前流上设置 MXMACA 事件并在结束事件上进行同步来解决这一问题。下面简要介绍 cupyx.profiler.benchmark() 中的内部操作：

>>> import time
>>> start_gpu = cp.cuda.Event()
>>> end_gpu = cp.cuda.Event()
>>>
>>> start_gpu.record()
>>> start_cpu = time.perf_counter()
>>> out = my_func(a)
>>> end_cpu = time.perf_counter()
>>> end_gpu.record()
>>> end_gpu.synchronize()
>>> t_gpu = cp.cuda.get_elapsed_time(start_gpu, end_gpu)
>>> t_cpu = end_cpu - start_cpu

此外， cupyx.profiler.benchmark() 执行一些预热运行，以减少时间波动并排除首次调用的开销。

3.5.2. 一次性开销

在对 mcPy 代码进行基准测试时要注意这些开销。

3.5.2.1. 上下文初始化

在进程中第一次调用 mcPy 函数可能需要几秒钟的时间。这是因为 MXMACA 驱动程序会在 MXMACA 应用程序中的第一个 API 调用期间创建 GPU 上下文。

3.5.2.2. 内核编译

mcPy 使用动态内核生成机制。当需要调用内核时，它编译针对给定参数的维度和数据类型优化的内核代码，将其发送到 GPU 设备，并执行内核。

mcPy 在进程中缓存发送到 GPU 设备的内核代码，这减少了后续调用时的内核编译时间。

编译后的代码也缓存在目录 ${HOME}/.cupy/kernel_cache 中（可以通过设置 CUPY_CACHE_DIR 环境变量重新指定缓存目录）。这使得在整个过程中可重用编译后的内核二进制文件。

3.5.3. 使用 mcCUB 对归约和其他例程加速

对于归约操作，如 sum()、 prod()、 amin()、 amax()、 argmin()、 argmax()，以及基于它们构建的更多例程，mcPy附带了自己的实现，这样就可以开箱即用了。然而，有其他一些专门的库来进一步加速这些程序，如mcCUB。

为了在适用的情况下支持更高性能的后端，mcPy引入了一个环境变量 CUPY_ACCELERATORS，使用户可以指定所需的后端（以及尝试调用后端的顺序）。例如，考虑在 256 三维数组上求和：

>>> from cupyx.profiler import benchmark
>>> a = cp.random.random((256, 256, 256), dtype=cp.float32)
>>> print(benchmark(a.sum, (), n_repeat=100))
sum : CPU: 12.101 us +/- 0.694 (min: 11.081 / max: 17.649) us GPU-0:10174.898 us +/-180.551 (min:10084.576 / max:10595.936) us

可以看到（在这个 GPU 上）运行大约需要 10 毫秒。然而，如果使用 CUPY_ACCELERATORS=cub Python 启动 Python 会话，将免费获得约 100 倍的加速（仅约 0.1 毫秒）：

>>> print(benchmark(a.sum, (), n_repeat=100))
sum : CPU: 20.569 us +/- 5.418 (min: 13.400 / max: 28.439) us GPU-0: 114.740 us +/- 4.130 (min: 108.832 / max: 122.752) us

mcCUB 是与 mcPy 一起提供的后端。它还加速了其他例程，如 inclusive 扫描（例如 cumsum() ）、直方图、稀疏矩阵向量乘法和归约内核。在 mcCUB 后端不适用的情况下，将返回到 mcPy 的默认实现。

虽然一般情况下 mcCUB 加速的归约操作执行速度更快，但根据数据布局可能会有例外。特别是，mcCUB reductions 仅支持在连续轴上进行 reduce。在任何情况下，建议执行一些基准测试，以确定 mcCUB 是否提供更好的性能。

备注

默认情况下，mcPy 使用 mcCUB。要关闭它，需要显式指定环境变量 CUPY_ACCELERATORS=""。

3.6. 互操作性

mcPy 可以与其他库一起使用。

3.6.1. NumPy

cupy.ndarray 实现了 __array_ufunc__ 接口。有关详细信息，请参见 NumPy相关文档 。这使得 NumPy ufunc 可以直接在 mcPy 数组上进行操作， __array_ufunc__ 功能需要 NumPy ≥1.13。

import cupy
import numpy

arr = cupy.random.randn(1, 2, 3, 4).astype(cupy.float32)
result = numpy.sum(arr)
print(type(result)) # => <class 'cupy._core.core.ndarray'>

cupy.ndarray 也实现了 __array_function__ 接口。这使得使用 NumPy 的代码可以直接在 mcPy 数组上进行操作。 __array_function__ 功能需要 NumPy ≥1.16 ；从 NumPy 1.17 开始，默认情况下会启用 __array_function__。

3.6.2. mpi4py

mpi4py 是消息传递接口（MPI）库的 Python 包装器。

MPI 是用于高性能进程间通信的最广泛使用的标准。最近，包括 MPICH、Open MPI 和 MVAPICH 在内的多家 MPI 供应商已将其支持扩展到 MPI-3.1 标准之外，以实现“GPU感知”，即，将 GPU 设备指针直接传递给 MPI 调用以避免 host 和 device 之间的显式数据移动。

有了在 mcPy 中实现的 DLPack 数据交换协议（参见 3.6.6 DLPack），mpi4py现在为将 mcPy 数组传递给 MPI 调用提供了（实验性）支持，前提是 mpi4py 是针对 GPU 感知的 MPI 实现构建的。以下是从 mpi4py 教程中借来的一个简单示例代码：

# To run this script with N MPI processes, do
# mpiexec -n N python this_script.py

import cupy
from mpi4py import MPI

comm = MPI.COMM_WORLD
size = comm.Get_size()

# Allreduce
sendbuf = cupy.arange(10, dtype='i')
recvbuf = cupy.empty_like(sendbuf)
comm.Allreduce(sendbuf, recvbuf)
assert cupy.allclose(recvbuf, sendbuf*size)

此新功能是从 mpi4py 3.1.0 开始添加的。有关更多信息，请访问 mpi4py相关文档 。

3.6.3. PyTorch

PyTorch 是一种机器学习框架，提供高性能、可微分的张量运算。

PyTorch 可以实现 mcPy 和 PyTorch 之间的零拷贝数据交换。唯一需要注意的是，PyTorch默认情况下会创建 CPU 张量，这些张量没有驻留在 GPU 显存中，用户需要确保张量在交换之前已经在 GPU 上。

>>> import cupy as cp
>>> import torch
>>>
>>> # convert a torch tensor to an mcPy array
>>> a = torch.rand((4, 4), device='maca')
>>> b = cp.asarray(a)
>>> b *= b
>>> b
array([[0.8215962 , 0.82399917, 0.65607935, 0.30354425],
[0.422695 , 0.8367199 , 0.00208597, 0.18545236],
[0.00226746, 0.46201342, 0.6833052 , 0.47549972],
[0.5208748 , 0.6059282 , 0.1909013 , 0.5148635 ]], dtype=float32)
>>> a
tensor([[0.8216, 0.8240, 0.6561, 0.3035],
[0.4227, 0.8367, 0.0021, 0.1855],
[0.0023, 0.4620, 0.6833, 0.4755],
[0.5209, 0.6059, 0.1909, 0.5149]], device='gpu:0')
>>> # check the underlying memory pointer is the same
>>>
>>> # convert an mcPy array to a torch tensor
>>> a = cp.arange(10)
>>> b = torch.as_tensor(a, device='maca')
>>> b += 3
>>> b
tensor([ 3, 4, 5, 6, 7, 8, 9, 10, 11, 12], device='gpu:0')
>>> a
array([ 3, 4, 5, 6, 7, 8, 9, 10, 11, 12])

PyTorch 还支持通过 DLPack 进行零拷贝数据交换，更多详情，请参见 3.6.6 DLPack。

import cupy
import torch

from torch.utils.dlpack import to_dlpack
from torch.utils.dlpack import from_dlpack

# Create a PyTorch tensor.
tx1 = torch.randn(1, 2, 3, 4).cuda()

# Convert it into a DLPack tensor.
dx = to_dlpack(tx1)

# Convert it into an mcPy array.
cx = cupy.from_dlpack(dx)

# Convert it back to a PyTorch tensor.
tx2 = from_dlpack(cx.toDlpack())

3.6.4. 在 PyTorch 中使用用户自定义内核

有了 DLPack 协议，使用 mcPy 用户定义内核在 PyTorch 中实现功能变得非常简单。以下是 PyTorch autograd 函数的示例，该函数使用 cupy.RawKernel 计算对数的前向和后向传递。

import cupy
import torch


cupy_custom_kernel_fwd = cupy.RawKernel(
    r"""
extern "C" __global__
void cupy_custom_kernel_fwd(const float* x, float* y, int size) {
    int tid = blockDim.x * blockIdx.x + threadIdx.x;
    if (tid < size)
        y[tid] = log(x[tid]);
}
""",
    "cupy_custom_kernel_fwd",
)


cupy_custom_kernel_bwd = cupy.RawKernel(
    r"""
extern "C" __global__
void cupy_custom_kernel_bwd(const float* x, float* gy, float* gx, int size) {
    int tid = blockDim.x * blockIdx.x + threadIdx.x;
    if (tid < size)
        gx[tid] = gy[tid] / x[tid];
}
""",
    "cupy_custom_kernel_bwd",
)


class CuPyLog(torch.autograd.Function):
    @staticmethod
    def forward(ctx, x):
        ctx.input = x
        # Enforce contiguous arrays to simplify RawKernel indexing.
        cupy_x = cupy.ascontiguousarray(cupy.from_dlpack(x.detach()))
        cupy_y = cupy.empty(cupy_x.shape, dtype=cupy_x.dtype)
        x_size = cupy_x.size
        bs = 128
        cupy_custom_kernel_fwd(
            (bs,), ((x_size + bs - 1) // bs,), (cupy_x, cupy_y, x_size)
        )
        # the ownership of the device memory backing cupy_y is implicitly
        # transferred to torch_y, so this operation is safe even after
        # going out of scope of this function.
        torch_y = torch.from_dlpack(cupy_y)
        return torch_y

    @staticmethod
    def backward(ctx, grad_y):
        # Enforce contiguous arrays to simplify RawKernel indexing.
        cupy_input = cupy.from_dlpack(ctx.input.detach()).ravel()
        cupy_grad_y = cupy.from_dlpack(grad_y.detach()).ravel()
        cupy_grad_x = cupy.zeros(cupy_grad_y.shape, dtype=cupy_grad_y.dtype)
        gy_size = cupy_grad_y.size
        bs = 128
        cupy_custom_kernel_bwd(
            (bs,),
            ((gy_size + bs - 1) // bs,),
            (cupy_input, cupy_grad_y, cupy_grad_x, gy_size),
        )
        # the ownership of the device memory backing cupy_grad_x is implicitly
        # transferred to torch_y, so this operation is safe even after
        # going out of scope of this function.
        torch_grad_x = torch.from_dlpack(cupy_grad_x)
        return torch_grad_x

备注

仅在 PyTorch 1.10+ 中添加的（新的）dlpack 数据交换协议中支持直接向 cupy.from_dlpack() 提供 torch.Tensor。 对于早期版本，需要使用 torc.utils.dlpack.to_dlpack() 打包 Tensor，示例如上所示。

3.6.5. 内存管理RMMX

RMMX（类 RAPID 内存管理器的扩展库）提供高度可配置的内存分配器。

RMMX 提供了一个接口，允许 mcPy 从 RMMX 内存池而不是从 mcPy 自己的池中分配内存。它可以简单地设置为：

import cupy
import rmmx

cupy.cuda.set_allocator(rmmx.rmm_cupy_allocator)

有时，可能需要一个性能更高的分配器。RMMX提供了一个切换分配器的选项：

import cupy
import rmmx

rmm.reinitialize(pool_allocator=True) # can also set init pool size etc here
cupy.cuda.set_allocator(rmm.rmm_cupy_allocator)

3.6.6. DLPack

DLPack 是一种张量结构规范，用于在框架之间共享张量。

mcPy 支持从 DLPack 数据结构（ cupy.from_DLPack() 和 cupy.ndarray.toDpack() ）导入和导出。

代码示例如下：

import cupy

# Create an mcPy array.
cx1 = cupy.random.randn(1, 2, 3, 4).astype(cupy.float32)

# Convert it into a DLPack tensor.
dx = cx1.toDlpack()

# Convert it back to an mcPy array.
cx2 = cupy.from_dlpack(dx)

TensorFlow 也支持 DLpack，因此 mcPy 和 TensorFlow 之间可以通过 DLpack 进行零拷贝数据交换：

>>> import tensorflow as tf
>>> import cupy as cp
>>>
>>> # convert a TF tensor to an mcPy array
>>> with tf.device('/GPU:0'):
...     a = tf.random.uniform((10,))
...
>>> a
<tf.Tensor: shape=(10,), dtype=float32, numpy=
array([0.9672388 , 0.57568085, 0.53163004, 0.6536236 , 0.20479882,
       0.84908986, 0.5852566 , 0.30355775, 0.1733712 , 0.9177849 ],
       dtype=float32)>
>>> a.device
'/job:localhost/replica:0/task:0/device:GPU:0'
>>> cap = tf.experimental.dlpack.to_dlpack(a)
>>> b = cp.from_dlpack(cap)
>>> b *= 3
>>> b
array([1.4949363 , 0.60699713, 1.3276931 , 1.5781245 , 1.1914308 ,
       2.3180873 , 1.9560868 , 1.3932796 , 1.9299742 , 2.5352407 ],
       dtype=float32)
>>> a
<tf.Tensor: shape=(10,), dtype=float32, numpy=
array([1.4949363 , 0.60699713, 1.3276931 , 1.5781245 , 1.1914308 ,
       2.3180873 , 1.9560868 , 1.3932796 , 1.9299742 , 2.5352407 ],
       dtype=float32)>
>>>
>>> # convert an mcPy array to a TF tensor
>>> a = cp.arange(10)
>>> cap = a.toDlpack()
>>> b = tf.experimental.dlpack.from_dlpack(cap)
>>> b.device
'/job:localhost/replica:0/task:0/device:GPU:0'
>>> b
<tf.Tensor: shape=(10,), dtype=int64, numpy=array([0, 1, 2, 3, 4, 5, 6, 7, 8, 9])>
>>> a
array([0, 1, 2, 3, 4, 5, 6, 7, 8, 9])

请注意，在 TensorFlow 中，所有张量都是不可变的，因此在后一种情况下， b 中的任何变化都不能反映在 mcPy 数组 a 中。

从 DLPack v0.5 开始，为了确保正确性，上述方法（隐式）要求用户确保此类转换（导入和导出 mcPy array）必须发生在同一个 GPU stream 上。如果有疑问，可以通过调用 cupy.cuda.get_current_tream() 来获取当前使用的流。

3.6.6.1. DLPack 数据交换协议

为了避免用户管理的流和 DLPack 张量对象，DLPack 数据交换协议提供了一种将责任从用户转移到库的机制。任何兼容的对象（如 cupy.ndarray ） 都必须实现一对方法 __dlpack__ 和 __dlpack_device__。函数 cupy.from_dlpack() 接受这样的对象，并返回一个 cupy.ndarray，该对象可以在 mcPy 的当前流上安全访问。同样， cupy.ndarray 可以通过任何兼容库的 from_dlpack() 函数导出。

备注

mcPy 使用 CUPY_DLPACK_EXPORT_VERSION 来控制如何处理由 MXMACA 托管内存支持的张量。

3.6.7. 显存指针

3.6.7.1. Import

mcPy 提供 UnownedMemory API，允许与其他库中分配的 GPU 设备内存进行互操作。

# Create a memory chunk from raw pointer and its size.
mem = cupy.cuda.UnownedMemory(140359025819648, 1024, owner=None)

# Wrap it as a MemoryPointer.
memptr = cupy.cuda.MemoryPointer(mem, offset=0)

# Create an ndarray view backed by the memory pointer.
arr = cupy.ndarray((16, 16), dtype= cupy.float32, memptr=memptr)
assert arr.nbytes <= arr.data.mem.size

需要指定正确的 shape、数据类型、stride 和 order，以便匹配所创建的 ndarray 视图。

UnownedMemory API 不管理内存分配的生命周期。必须确保指针在被 mcPy 使用时是活跃的。如果指针生存期由 Python 对象管理，则可以将其传递给 UnownedMemory 的 owner 参数，以保留对该对象的引用。

3.6.7.2. Export

可以将在 mcPy 中分配的内存指针传递给其他库。

arr = cupy.arange(10)
print(arr.data.ptr, arr.nbytes) # => (140359025819648, 80)

当 ndarray （ 上例中的 arr ）被析构时，将释放 mcPy 分配的内存。当指针被其他库使用时，必须使 ndarray 实例处于活跃状态。

3.6.8. MXMACA 流指针

3.6.8.1. Import

mcPy 提供了 ExternalStream API，允许与其他库中创建的 GPU 流进行互操作。

import torch

# Create a stream on PyTorch.
s = torch.cuda.Stream()

# Switch the current stream in PyTorch.
with torch.cuda.stream(s):
    # Switch the current stream in CuPy, using the pointer of the stream created in PyTorch.
    with cupy.cuda.ExternalStream(s.cuda_stream):
        # This block runs on the same CUDA stream.
        torch.arange(10, device='cuda')
        cupy.arange(10)

ExternalStream API 不管理流的生命周期。必须确保流指针在被 mcPy 使用时是活跃的。

还需要确保 ExternalStream 对象在创建流的设备上使用。如果在创建 ExternalStream 时传递 device_id 参数，mcPy 可以验证这一点。

3.6.8.2. Export

可以将在 mcPy 中创建的流传递给其他库。

s = cupy.cuda.Stream()
print(s.ptr, s.device_id) # => (93997451352336, 0)

当流被销毁时，将销毁 MXMACA 流。当指针被其他库使用时，必须使 Stream 实例处于活跃状态。

3.7. mcPy 和 NumPy 的差异

mcPy 的接口兼容 NumPy 的接口，但是也有一些不同之处。

3.7.1. float 到 integer 类型转换

浮点到整数的某些强制转换行为在 C++ 规范中没有定义，例如从负浮点转换到无符号整数以及从无穷大转换到整数。NumPy 的行为取决于 CPU 体系结构。这是英特尔 CPU 上的结果：

>>> np.array([-1], dtype=np.float32).astype(np.uint32)
array([4294967295], dtype=uint32)
>>> cupy.array([-1], dtype=np.float32).astype(np.uint32)
array([0], dtype=uint32)

>>> np.array([float('inf')], dtype=np.float32).astype(np.int32)
array([-2147483648], dtype=int32)
>>> cupy.array([float('inf')], dtype=np.float32).astype(np.int32)
array([2147483647], dtype=int32)

3.7.2. Random 方法支持 dtype 参数

NumPy 的随机值生成器不支持 dtype 参数，而是始终返回 float64 值。由于 mcPy 使用的 mcRAND 同时支持 float32 和float64，所以mcPy支持 dtype 参数。

>>> np.random.randn(dtype=np.float32)
Traceback (most recent call last):
  File "<stdin>", line 1, in <module>
TypeError: randn() got an unexpected keyword argument 'dtype'
>>> cupy.random.randn(dtype=np.float32)
array(0.10689262300729752, dtype=float32)

3.7.3. 越界索引

默认情况下，使用整数数组索引时，mcPy 处理越界索引的方式与 NumPy 不同。NumPy 通过引发一个错误来处理它们，但 mcPy 会从另一端重新开始索引。

>>> x = np.array([0, 1, 2])
>>> x[[1, 3]] = 10
Traceback (most recent call last):
  File "<stdin>", line 1, in <module>
IndexError: index 3 is out of bounds for axis 1 with size 3
>>> x = cupy.array([0, 1, 2])
>>> x[[1, 3]] = 10
>>> x
array([10, 10, 2])

3.7.4. 多次引用同一位置

当整数数组多次引用同一位置时，mcPy的 __setitem__ 的行为与 NumPy 不同。在这种情况下，实际存储的值是未定义的。这是一个 mcPy 的例子：

>>> a = cupy.zeros((2,))
>>> i = cupy.arange(10000) % 2
>>> v = cupy.arange(10000).astype(np.float32)
>>> a[i] = v
>>> a
array([ 9150., 9151.])

NumPy 存储引用重复位置的元素中最后一个元素对应的值。

>>> a_cpu = np.zeros((2,))
>>> i_cpu = np.arange(10000) % 2
>>> v_cpu = np.arange(10000).astype(np.float32)
>>> a_cpu[i_cpu] = v_cpu
>>> a_cpu
array([9998., 9999.])

3.7.5. 零维数组

3.7.5.1. 归约方法

NumPy 的归约函数（例如 numpy.sum() ）返回标量值（例如 numpy.float32 ）。然而，mcPy 的对应函数返回零维 cupy.ndarray。这是因为 mcPy 标量值（如 cupy.float32 ）是 NumPy 标量的别名，并在 CPU 内存中分配。如果返回这些类型，则需要在 GPU 和 CPU 之间进行同步。如果要使用标量值，请显式强制转换返回的数组。

>>> type(np.sum(np.arange(3))) == np.int64
True
>>> type(cupy.sum(cupy.arange(3))) == cupy.ndarray
True

3.7.5.2. 类型提升

在具有两个或多个操作数的函数中，mcPy 自动提升 cupy.ndarray 的 dtype，返回值的 dtype 由输入的 dtype 决定。这与 NumPy 在运算数包含零维数组时的提升原则不同。零维 numpy.ndarray 如果出现在 NumPy 函数的运算数中，则被视为标量值。此标量的数值决定输出值的 dtype。

>>> (np.array(3, dtype=np.int32) * np.array([1., 2.], dtype=np.float32)).dtype
dtype('float32')
>>> (np.array(300000, dtype=np.int32) * np.array([1., 2.], dtype=np.float32)).dtype
dtype('float64')
>>> (cupy.array(3, dtype=np.int32) * cupy.array([1., 2.], dtype=np.float32)).dtype
dtype('float64')

3.7.6. 数据类型

mcPy 数组的数据类型不能是非数字型，比如字符串或者对象。

3.7.7. 仅接受 mcPy 数组或标量的通用函数

与 NumPy 不同，mcPy 中的通用函数仅接受 mcPy 数组或标量。它们不接受其他对象（例如，lists 或 numpy.ndarray ）。

>>> np.power([np.arange(5)], 2)
array([[ 0, 1, 4, 9, 16]])

>>> cupy.power([cupy.arange(5)], 2)
Traceback (most recent call last):
  File "<stdin>", line 1, in <module>
TypeError: Unsupported type <class 'list'>

3.7.8. Random 数组种子被散列到一个数字

与 NumPy 一样，mcPy 的 RandomState 对象接受数字或完整的 NumPy 数组作为种子。

>>> seed = np.array([1, 2, 3, 4, 5])
>>> rs = cupy.random.RandomState(seed=seed)

但与 NumPy 不同的是，数组种子将被散列到一个数字，因此可能不会向底层随机数生成器传递太多的随机性。

3.7.9. NaN（not-a-number）处理

默认情况下，mcPy 的归约函数（例如， cupy.sum() ）处理复数形式的 NaN 与 NumPy 的对应函数不同：

>>> a = [0.5 + 3.7j, complex(0.7, np.nan), complex(np.nan, -3.9), complex(np.nan, np.nan)]
>>>
>>> a_np = np.asarray(a)
>>> print(a_np.max(), a_np.min())
(0.7+nanj) (0.7+nanj)
>>>
>>> a_cp = cp.asarray(a_np)
>>> print(a_cp.max(), a_cp.min())
(nan-3.9j) (nan-3.9j)

原因是 reduction 是以跨步的方式执行的，因此它不能确保正确的比较顺序，并且不能遵循 NumPy 规则始终传播第一个遇到的 NaN。请注意，当启用 mcCUB 时（这是 mcPy 的默认设置）此差异不适用。

3.7.10. 连续性和步长

为了提供最佳性能，无法保证生成的 ndarray 的连续性（contiguity）与 NumPy 输出的连续性相匹配。

>>> a = np.array([[1, 2], [3, 4]], order='F')
>>> print((a + a).flags.f_contiguous)
True
>>> a = cp.array([[1, 2], [3, 4]], order='F')
>>> print((a + a).flags.f_contiguous)
False

4. mcPy 编程接口

mcPy 的 API 接口兼容 CuPy 接口，详细信息可参见 CuPy API参考文档 。

5. 附录

术语/缩略语

全称

说明

FFT

Fast Fourier Transform

快速傅里叶变换

JIT Compile

Just-in-time Compile

即时编译

MPI

Message Passing Interface

消息传递接口

P2P

peer-to-peer

PCIe P2P是 PCIe 的一种特性，使两个 PCIe 设备之间可以直接传输数据

REPL

Read-Eval-Print Loop

交互式解释器