AI 技术讲座精选：技术前沿——CUDAnative.jl 支持 GPU 原生编程
AI 技术讲座精选：技术前沿——CUDAnative.jl 支持 GPU 原生编程
【AI100 导读】首款拥有 GPU 原生编程功能的 Julia 编程语言公测版终于发布了！本文介绍了如何编写像 GPU 一样的并行加速程序。
经过两年缓慢但却稳定的发展，我们最终发布了首款拥有 GPU 原生编程功能的 Julia 编程语言的公测版。虽然仍然存在某些方面的限制，但是现在运用 Julia 编写 CUDA 核心程序已经得以实现。相应地，使用 Julia 高级语言特性编写高性能的 GPU 代码也成为可能。
本篇文章中演示的编程支持是由低级构块组成的，而这些构块与 CUDA C 语言处于相同的抽象层次。如果你知道（或者想了解）如何编写像 GPU 这样的并行加速程序，那么你一定会对本篇感兴趣的，不过你需要处理棘手的运行特性和通信语义。
你可以通过安装 CUDAnative.jl 轻松地把 GPU 支持添加到已安装的 Julia 编程语言中。这个程序包建立在 Julia 编译程序测试接口上，而且特意创建 LLVM.jl 和 CUDAdrv.jl 这两个用来编译和执行代码的程序包。这些功能非常新并且尚未测试过，所以为了将来能正式推出 Julia1.0，我们需要你们的帮助和反馈来优化并完善其接口。
如何开始
CUDAnative.jl 提高了版本和平台的兼容性，而且它是 Julia 编译程序以及底层 LLVM 架构的紧密集成。关于这个公测版，在 Linux 或者 macOS 操作系统中只支持由源代码构建的 Julia 0.6 版。幸运的是，从源代码处安装的 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 上执行该程序。请一定要查看示例，或者继续阅读更多的内容介绍。
矢量加法
通过一个小型经典的 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 硬件的相互作用成为可能。例如，它会提供 CuArra 数组类型，能够负责内存管理，并且能够与 Julia 的垃圾回收器进行集成，利用 GPU 事件执行 @elapsed 等等。这就意味着，为了与 CUDA 驱动进行交互作用，我们需要形成一个强大的基础，但是不需要使用最新版本的 Julia。在 CUDArt.jl 运行环境下，我们可以使用版本稍微高一点的 Julia，而不是建立在 CUDA 运行的驱动 API 上，但是目前仍未与 CUDArt.jl 进行集成。
同时，CUDArt.jl 负责与 GPU 原生编程有关的所有事情。其中，最重要的部分就是生成 GPU 代码的部分，大概由以下三部分组成：
- 与 Julia 交互：重新目的化编译器产生 GPU 与 LLVM 兼容的指令寄存器（CPU 库没有接收到指令，简化的异常处理，……）。
- 与 LLVM 交互（使用 LLVM.jl）：优化指令寄存器，然后编译成 PTX。
- 与 CUDA 交互（使用 CUDAdrv.jl）：把 PTX 编译成 SASS，然后把它上传到 GPU 中。
所有的步骤都隐藏在 @cuda 指令后面，初次使用该指令时，它便会生成代码编译出我们的核心程序。每一次后续调用都会重新使用这些代码，对其进行转换并上传参数1，最后运行核心程序。你可以利用运行时的反射来内省这些代码：
# 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]CUDAnative.jl 的另一个重要功能就是内在函数：特殊功能和宏指令造成的功能难题或许是无法用一般功能表述的。例如，{thread,block,grid}{Idx,Dim} 功能支持访问每个工作层的尺寸和索引。@cuStaticSharedMem和 @cuDynamicSharedMem 宏指令可以创建局部分享内存，同时 @cuprintf 可以从核心程序内部显示格式化字符串。不仅如此，还有许多数学功能可以供我们使用；这些功能可以替代标准库中近似的功能。
有什么遗漏吗？
正如我所说的，我们并没有讲解 Julia 语言的所有特性。例如，它现在无法从运行的 Julia C 库（aka. 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 的驱动API，但是大多数的 CUDAnative.jl 会效仿或者复刻 CUDA C 语言。但是，我们没有记录哪些部分的 API 是被覆盖的，或者说，我们没有记录抽象的表现过程，所以你需要参考 CUDAnative 和 CUDAdrv 库中的示例和测试结果。
另一个示例：平行归约
这是一个更加复杂的示例，我们来看一下 Kepler-generation GPUs 的平行归约。
这是一个典型并且良好优化过的 GPU 的实现过程，在每个执行级别利用快速通讯原语。例如，曲线里的线程以 SIMD-like 为核心一同执行，可以通过相互的域分享数据。在块级别中，线程都归集到同一个核心处，但是没有必要一起执行，这就意味着他们需要通过局部储存器的核心进行交流。在其他高级别中，只有 GPU 的 DRAM 储存器是一个可使用的通讯媒介。
Julia 版本的这种算法看起来与 CUDA 的原始算法非常相似：这和预期的一样，因为 CUDAnative.jl 是 CUDA C 语言的副本。新版本更加通用，擅长缩减运算符和数字类型。就像我们过去经常用的 Julia 常规代码， @cuda 宏指令会根据参数类型进行即时编译并且派送到正确的分工区域。
那它的运行情况怎么样呢？结果非常好！下面是 CUDAnative.jl 和 CUDA C 运行情况对比表，我们利用 BenchmarkTools.jl 来监测其执行的时间。因为参数传递缺失，所以曲线呈现缓慢持续上扬态势（注意对数刻度），最后形成了如下的线状图表。
我们还打算比较 CUDA 工具包中的工具。例如，利用 NVIDIA 的可视化探查器，描述出 Julia 核心程序。或者利用 cuda-memcheck 探测界外访问3。
$ 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试一下！
如果你曾从事 GPCs 或者 CUDA 的研发，或者你拥有的程序包能够从 GPU 的加速中获益，那么请你看一下或者试一下 CUDAnative.jl! 在推出 Julia 1.0 之前，我们要完成其优化开发工作，并且需要完成其基本的框架。所以，如果有任何问题请务必向我们反馈！我们需要你的支持！
我想帮忙
甚至更好！贡献的方式有很多种，例如，通过观察单个程序包的事件追踪器来完善支持工作。
- CUDAnative.jl https://github.com/JuliaGPU/CUDAnative.jl/issues
- CUDAdrv.jl https://github.com/JuliaGPU/CUDAdrv.jl/issues
- LLVM.jl https://github.com/maleadt/LLVM.jl/issues
在这些程序包中，任何一个程序包在任何时候都需要更好的 API 予以覆盖。并且它们也需要相关的记录，以覆盖和解释那些已经执行过了的程序包。
本文作者 Tim Besard 根特大学的研究员，致力于使用来自高级编程语言的 GPU 的编译技术，是 Avid Linux用户，积极从事开源工作。
本文由 AI100 编译，转载需得到本公众号同意。
编译：AI100
原文链接：http://julialang.org/blog/2017/03/cudanative