统一CUDA Python 生态系统
Python 在科学、工程、资料分析和深度学习应用生态系统中扮演关键角色。长期以来，NVIDIA 皆致力于协助Python 生态系统利用GPU 的加速大规模平行效能，提供标准化函数库、工具和应用程式。如今，我们已经改善了Python 程式码的可移植性和相容性，进一步朝简化开发人员体验迈进。
我们的目标是以单一标准低阶介面集合，协助统一Python CUDA 生态系统，提供全面地覆盖和从Python 存取CUDA 主机的API。我们希望能提供生态系统基础，让不同的加速函数库彼此互通。最重要的是，Python 开发人员可以更轻松地使用NVIDIA GPU。
CUDA Python：漫长且曲折的道路
截至目前为止，想要透过Python 存取CUDA 和NVIDIA GPU 仅能使用第三方软体，例如Numba、CuPy、Scikit-CUDA、RAPIDS、PyCUDA、PyTorch 或TensorFlow。他们都在CUDA API 与Python 之间编写各自的互通层。
NVIDIA 发布的CUDA Python，可以让这些平台供应商专注于各自的附加价值产品与服务。NVIDIA 同时希望能降低其他Python 开发人员使用NVIDIA GPU 的门槛。CUDA Python 初版包含用于CUDA 驱动程式和执行阶段API 的Cython 与Python 包装函式。
我们可能会在未来版本中，提供用于CUDA 函式库如cuBLAS、cuFFT、cuDNN、nvJPEG 等的Pythonic 物件模型和包装函式。未来版本可能会与GitHub 上的开放原始码一并提供，或透过PIP 和Conda 封装。
CUDA Python 工作流程
由于Python 是一种解译语言，必须先设法将装置程式码编译成PTX，然后撷取将要在应用程式中呼叫的函数。相比之下理解CUDA Python 不是最重要的，但是需要了解Parallel Thread Execution（PTX）是一种低阶虚拟机器和指令集架构（instruction set architecture，ISA）。以字串形式建构装置程式码，并使用CUDA C++ 执行阶段编译函式库NVRTC进行编译。使用NVIDIA驱动程式API，在GPU 上手动建立CUDA 脉络及所有的必要资源，然后启动已编译CUDA C++ 程式码，并从GPU 撷取结果。现在，您已经大致了解，接着将进入平行程式设计的常用范例：SAXPY。
首先，从CUDA Python 套件汇入驱动程式API和NVRTC模组。在此范例中，将资料从主机复制到设备。需要NumPy在主机上储存资料。
import cuda_driver as cuda # Subject to change before releaseimport nvrtc # Subject to change before releaseimport numpy as np错误检查是程式码开发的基本最佳做法，且已提供了程式码范例。为求精简，省略了范例中的错误检查。在未来版本中，可能会使用Python 物件模型自动引发例外。
def ASSERT_DRV(err): if isinstance(err, cuda.CUresult): if err != cuda.CUresult.CUDA_SUCCESS: raise RuntimeError("Cuda Error: {}".format(err)) elif isinstance(err, nvrtc.nvrtcResult): if err != nvrtc.nvrtcResult.NVRTC_SUCCESS: raise RuntimeError("Nvrtc Error: {}".format(err)) else: raise RuntimeError("Unknown error type: {}".format(err))常见之做法是在转译单位的顶部附近编写CUDA 核心，所以接下来将编写此部分。使用三引号包住整个核心，以形成字串。之后使用NVRTC 编译字串。这是CUDA Python 中唯一需要理解CUDA C++ 的部分。若需要更多资讯，请参阅An Even Easier Introduction to CUDA (https://developer.nvidia.com/blog/even-easier-introduction-cuda/)。
saxpy = """\extern "C" __global__void saxpy(float a, float *x, float *y, float *out, size_t n){ size_t tid = blockIdx.x * blockDim.x + threadIdx.x; if (tid < n) { out[tid] = a * x[tid] + y[tid]; }}"""继续将核心编译成PTX。请记住，这是使用NVRTC，在执行阶段执行。NVRTC 有三个基本步骤：
- 从字串建立程式。
- 编译程式。
- 从已编译程式中撷取PTX。
在以下程式码范例中，针对运算能力75或Turing 架构进行编译，并启用FMAD。如果编译失败，请使用 nvrtcGetProgramLog 撷取编译记录，以取得其他资讯。
# Create programerr, prog = nvrtc.nvrtcCreateProgram(str.encode(saxpy), b"saxpy.cu", 0, [], [])
# Compile programopts = [b"--fmad=false", b"--gpu-architecture=compute_75"]err, = nvrtc.nvrtcCompileProgram(prog, 2, opts)
# Get PTX from compilationerr, ptxSize = nvrtc.nvrtcGetPTXSize(prog)ptx = b" " * ptxSizeerr, = nvrtc.nvrtcGetPTX(prog, ptx)在使用PTX 或在GPU 上执行任何工作之前，必须先建立CUDA context。CUDA context类似于设备的主机处理序。在以下程式码范例中，将驱动程式API 初始化，以存取NVIDIA 驱动程式和GPU。其次，将运算设备0 的控点传递至cuCtxCreate，以指定该GPU 建立context。在建立context之后，可以继续使用NVRTC 编译CUDA 核心。
# Initialize CUDA Driver APIerr, = cuda.cuInit(0)
# Retrieve handle for device 0err, cuDevice = cuda.cuDeviceGet(0)
# Create contexterr, context = cuda. cuCtxCreate(0, cuDevice)在设备0 上建立CUDA context之后，将先前产生的PTX 载入至模块。模块类似于设备的动态载入函数库。在载入至模块之后，使用cuModuleGetFunction 撷取特定核心。多个核心常驻于PTX 中不是罕见的情形。
# Load PTX as module data and retrieve functionptx = np.char.array(ptx)err, module = cuda.cuModuleLoadData(ptx.ctypes.get_data())err, kernel = cuda.cuModuleGetFunction(module, b"saxpy")之后，准备所有资料及传输至GPU。为了提高应用程式效能，可以在装置上输入资料，以省略资料传输。为了能完整理解，此范例将示范如何将资料输入与输出设备。
NUM_THREADS = 512 # Threads per blockNUM_BLOCKS = 32768 # Blocks per grid
a = np.array([2.0], dtype=np.float32)n = np.array(NUM_THREADS * NUM_BLOCKS, dtype=np.uint32)bufferSize = n * a.itemsize
hX = np.random.rand(n).astype(dtype=np.float32)hY = np.random.rand(n).astype(dtype=np.float32)hOut = np.zeros(n).astype(dtype=np.float32)为SAXPY 转换设备建立输入资料a、x、y 之后必须分配资源，以使用cuMemAlloc 储存资料。想要在运算与资料移动之间允许更多重叠时，请使用非同步函式cuMemcpyHtoDAsync。它会在命令执行后，立即将控制权交还给CPU。
Python 没有自然的指标概念，但是cuMemcpyHtoDAsync 需要void*。因此，XX.ctypes.get_data 会撷取与XX 有关的指标值。
err, dXclass = cuda.cuMemAlloc(bufferSize)err, dYclass = cuda.cuMemAlloc(bufferSize)err, dOutclass = cuda.cuMemAlloc(bufferSize)
err, stream = cuda.cuStreamCreate(0)
err, = cuda.cuMemcpyHtoDAsync( dXclass, hX.ctypes.get_data(), bufferSize, stream)err, = cuda.cuMemcpyHtoDAsync( dYclass, hY.ctypes.get_data(), bufferSize, stream)在完成资料准备和资源分配之后，即可启动核心。想要将装置上的资料位置传递至核心执行设备时，必须撷取装置指标。在以下程式码范例中，int(dXclass) 会重试dXclass 的指标值，即CUdeviceptr，并使用np.array 分配记忆体大小，以储存该值。
如同cuMemcpyHtoDAsync，cuLaunchKernel 在引数清单中需要void**。在先前的程式码范例中，建立void** 的方式是取得各个引数的void* 值，并将其放入各自的连续记忆体中。
# The following code example is not intuitive# Subject to change in a future releasedX = np.array([int(dXclass)], dtype=np.uint64)dY = np.array([int(dYclass)], dtype=np.uint64)dOut = np.array([int(dOutclass)], dtype=np.uint64)
args = [a, dX, dY, dOut, n]args = np.array([arg.ctypes.get_data() for arg in args], dtype=np.uint64)现在可以启动核心：
err, = cuda.cuLaunchKernel( kernel, NUM_BLOCKS, # grid x dim 1, # grid y dim 1, # grid z dim NUM_THREADS, # block x dim 1, # block y dim 1, # block z dim 0, # dynamic shared memory stream, # stream args.ctypes.get_data(), # kernel arguments 0, # extra (ignore))
err, = cuda.cuMemcpyDtoHAsync( hOut.ctypes.get_data(), dOutclass, bufferSize, stream)err, = cuda.cuStreamSynchronize(stream)cuLaunchKernel 函式取得已编译的模块核心和执行配置参数。在与资料传输相同的资料流中启动装置程式码。可以确保仅会在完成资料传输后，执行核心运算，因为资料流中的所有API 呼叫及核心启动都已经序列化。在将资料传回主机的呼叫之后，使用cuStreamSynchronize 暂停CPU 执行，直至完成指定资料流中的所有运算。
# Assert values are same after running kernelhZ = a * hX + hYif not np.allclose(hOut, hZ): raise ValueError("Error outside tolerance for host-device vectors")执行资料验证以确保正确性，并透过记忆体清理完成程式码。
err, = cuda.cuStreamDestroy(stream)err, = cuda.cuMemFree(dXclass)err, = cuda.cuMemFree(dYclass)err, = cuda.cuMemFree(dOutclass)err, = cuda.cuModuleUnload(module)err, = cuda.cuCtxDestroy(context)效能
效能是在应用程式中，以GPU 为目标的主要驱动力。因此，相较于C++ 版本，上述程式码如何呢？如表1 所示，结果几乎相同。NVIDIA NSight Systems是使用于撷取核心效能，以及CUDA Events是使用于应用程式效能。
使用以下命令剖析应用程式：
nsys profile -s none -t cuda --stats=true <executable>
CUDA Python 与用于CUDA 应用程式的互动式核心剖析工具NVIDIA Nsight Compute也相容。它让您可以详细了解核心效能。这在尝试将效能最大化时很实用（图1）。
图1：CUDA Python 范例的Nsight Compute CLI 输出萤幕撷取画面。
CUDA Python 入门
CUDA Python 即将推出，并随附API 的详细说明、安装注意事项、新功能和范例。