技术预览：使用 CUDAnative.jl 进行原生 GPU 编程

经过两年缓慢但稳定的开发，我们很高兴地宣布 Julia 原生 GPU 编程功能的首个预览版本。你现在可以在 Julia 中编写 CUDA 内核，尽管存在一些限制，但这使得可以使用 Julia 的高级语言特性编写高性能 GPU 代码成为可能。

我们今天展示的编程支持包括低级构建块，位于与 CUDA C 相同的抽象级别。如果你了解（或想学习）如何对 GPU 等并行加速器进行编程，同时处理棘手的性能特征和通信语义，那么你应该对此感兴趣。

你可以通过安装 CUDAnative.jl，轻松地将 GPU 支持添加到你的 Julia 安装中（有关详细说明，请参见下文）。此软件包构建在 Julia 编译器的实验性接口之上，以及专为该目的而构建的 LLVM.jl 和 CUDAdrv.jl 软件包，用于编译和执行代码。所有这些功能都是全新的，并且经过了彻底的测试，因此我们需要你的帮助和反馈，以便在 Julia 1.0 之前改进和完善这些接口。

如何开始
"Hello World" 向量加法
1. 它是如何工作的？
2. 缺少什么？
另一个示例：并行归约
试试吧！
1. 我想帮忙
感谢

如何开始

CUDAnative.jl 与 Julia 编译器和底层的 LLVM 框架紧密集成，这使得版本和平台兼容性变得复杂。对于此预览版本，我们只支持从源代码构建的 Julia 0.6，在 Linux 或 macOS 上运行。幸运的是，从源代码安装 Julia 在主存储库的 README 中有详细说明。大多数情况下，它可以归结为以下命令

$ git clone https://github.com/JuliaLang/julia.git
$ cd julia
$ git checkout v0.6.0-pre.alpha  # or any later tag
$ make                           # add -jN for N parallel jobs
$ ./julia

从 Julia REPL 中，安装 CUDAnative.jl 及其依赖项只需使用包管理器即可。请注意，你需要使用 NVIDIA 二进制驱动程序，并安装 CUDA 工具包。

Pkg.add("CUDAnative")

# Optional: test the package
Pkg.test("CUDAnative")

此时，你可以开始编写内核，并使用 CUDAnative 的 @cuda 在 GPU 上执行它们！请务必查看示例，或继续阅读以获得更具文本性的介绍。

"Hello World" 向量加法

GPU 编程能力的典型小演示（将其视为GPU Hello World）是执行向量加法。以下代码段正是使用 Julia 和 CUDAnative.jl 完成此操作的。

using CUDAdrv, CUDAnative

function kernel_vadd(a, b, c)
    # from CUDAnative: (implicit) CuDeviceArray type,
    #                  and thread/block intrinsics
    i = (blockIdx().x-1) * blockDim().x + threadIdx().x
    c[i] = a[i] + b[i]

    return nothing
end

dev = CuDevice(0)
ctx = CuContext(dev)

# generate some data
len = 512
a = rand(Int, len)
b = rand(Int, len)

# allocate & upload on the GPU
d_a = CuArray(a)
d_b = CuArray(b)
d_c = similar(d_a)

# execute and fetch results
@cuda (1,len) kernel_vadd(d_a, d_b, d_c)    # from CUDAnative.jl
c = Array(d_c)

using Base.Test
@test c == a + b

destroy(ctx)

此示例的大部分内容并不依赖于 CUDAnative.jl，而是使用 CUDAdrv.jl 的功能。此软件包使通过 CUDA 驱动程序 API 的用户友好包装器与 CUDA 硬件进行交互成为可能。例如，它提供了一种数组类型 CuArray，负责内存管理，与 Julia 的垃圾收集器集成，使用 GPU 事件实现 @elapsed 等。它的目的是为所有与 CUDA 驱动程序的交互奠定坚实的基础，并且不需要 Julia 的最新版本。一个略微更高层次的替代方案在 CUDArt.jl 下可用，它建立在 CUDA 运行时 API 之上，但尚未与 CUDAnative.jl 集成。

同时，CUDAnative.jl 负责与原生 GPU 编程相关的所有事项。其中最显著的部分是生成 GPU 代码，它基本上包括三个阶段

与 Julia 交互：重新利用编译器以生成与 GPU 兼容的 LLVM IR（不调用 CPU 库，简化的异常，…）
与 LLVM 交互（使用 LLVM.jl）：优化 IR，并编译为 PTX
与 CUDA 交互（使用 CUDAdrv.jl）：将 PTX 编译为 SASS，并将其上传到 GPU

所有这些都隐藏在对 @cuda 的调用之后，该调用会在第一次使用时生成用于编译内核的代码。每次随后的调用都会重复使用该代码，转换和上传参数^[1]，最后启动内核。与我们在 CPU 上习惯的那样，你可以使用运行时反射来内省此代码

# CUDAnative.jl provides alternatives to the @code_ macros,
# looking past @cuda and converting argument types
julia> CUDAnative.@code_llvm @cuda (1,len) kernel_vadd(d_a, d_b, d_c)
define void @julia_kernel_vadd_68711 {
    [LLVM IR]
}

# ... but you can also invoke without @cuda
julia> @code_ptx kernel_vadd(d_a, d_b, d_c)
.visible .func julia_kernel_vadd_68729(...) {
    [PTX CODE]
}

# or manually specify types (this is error prone!)
julia> code_sass(kernel_vadd, (CuDeviceArray{Float32,2},CuDeviceArray{Float32,2},CuDeviceArray{Float32,2}))
code for sm_20
        Function : julia_kernel_vadd_68481
[SASS CODE]

^[1]

CUDAnative.jl 的另一个重要部分是内联函数：提供难以或不可能使用普通函数表达的功能的特殊函数和宏。例如，{thread,block,grid}{Idx,Dim} 函数提供对每个工作级别的大小和索引的访问。可以使用 @cuStaticSharedMem 和 @cuDynamicSharedMem 宏创建本地共享内存，而 @cuprintf 可用于从内核函数中显示格式化的字符串。许多数学函数也可用；这些函数应代替标准库中的类似函数。

缺少什么？

正如我已经暗示的那样，我们尚不支持 Julia 语言的所有功能。例如，目前无法从 Julia C 运行时库（即 libjulia.so）调用任何函数。这使得动态分配成为不可能，并削弱了异常等功能。因此，标准库的很大一部分无法在 GPU 上使用。我们当然会在未来努力改进这一点，但目前，编译器会在遇到不支持的语言特性时报错。

julia> nope() = println(42)
nope (generic function with 1 method)

julia> @cuda (1,1) nope()
ERROR: error compiling nope: emit_builtin_call for REPL[1]:1 requires the runtime language feature, which is disabled

另一个很大的差距是文档。CUDAnative.jl 的大部分内容模仿或复制 CUDA C，而 CUDAdrv.jl 包装 CUDA 驱动程序 API。但是，我们还没有记录这些 API 的哪些部分被覆盖，或者这些抽象如何工作，因此你需要参考 CUDAnative 和 CUDAdrv 存储库中的示例和测试。

另一个示例：并行归约

对于更复杂的示例，让我们看一下针对 Kepler 架构 GPU 的并行归约。这是典型的经过良好优化的 GPU 实现，它使用每级执行的快速通信原语。例如，同一个 warp 中的线程在一个类似于 SIMD 的核心上一起执行，并且可以通过彼此的寄存器共享数据。在块级别，线程分配在同一个核心上，但并不一定一起执行，这意味着它们需要通过核心本地内存进行通信。再向上一个级别，只有 GPU 的 DRAM 内存是可行的通信媒介。

此算法的 Julia 版本看起来与 CUDA 原版非常相似：这是预期的结果，因为 CUDAnative.jl 是 CUDA C 的对应物。但新版本更加通用，它专门针对归约运算符和值类型。与我们使用常规 Julia 代码习惯的那样，@cuda 宏会根据参数类型即时编译和分派到正确的专门化。

那么，它的性能如何？事实证明，相当不错！以下图表比较了 CUDAnative.jl 和 CUDA C 实现的性能^[2]，使用 BenchmarkTools.jl 来测量执行时间。少量固定开销（注意对数刻度）是由于参数传递的缺陷导致的，将会得到修复。

Performance comparison of parallel reduction implementations.

[2]	测量结果包括内存传输时间，这就是为什么没有包括 CPU 实现的原因（现实情况是，数据会尽可能长地保留在 GPU 上，这使得比较不公平）。

我们还旨在与 CUDA 工具包中的工具兼容。例如，你可以使用 NVIDIA Visual Profiler 分析 Julia 内核，或者使用 cuda-memcheck 检测越界访问。

$ cuda-memcheck julia examples/oob.jl
========= CUDA-MEMCHECK
========= Invalid __global__ write of size 4
=========     at 0x00000148 in examples/oob.jl:14:julia_memset_66041
=========     by thread (10,0,0) in block (0,0,0)
=========     Address 0x1020b000028 is out of bounds

完整的调试信息尚不可用，因此 cuda-gdb 及其同类工具将无法正常工作。

[1]	请参阅 README，了解有关这在当前情况下是多么昂贵的说明。

技术预览：使用 CUDAnative.jl 进行原生 GPU 编程

2017 年 3 月 14 日 | Tim Besard