【AI100 导读】首款拥有 GPU 原生编程功能的 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 代码的部分,大概由以下三部分组成:
所有的步骤都隐藏在 @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 之前,我们要完成其优化开发工作,并且需要完成其基本的框架。所以,如果有任何问题请务必向我们反馈!我们需要你的支持!
我想帮忙
甚至更好!贡献的方式有很多种,例如,通过观察单个程序包的事件追踪器来完善支持工作。
在这些程序包中,任何一个程序包在任何时候都需要更好的 API 予以覆盖。并且它们也需要相关的记录,以覆盖和解释那些已经执行过了的程序包。
本文作者 Tim Besard 根特大学的研究员,致力于使用来自高级编程语言的 GPU 的编译技术,是 Avid Linux用户,积极从事开源工作。
本文由 AI100 编译,转载需得到本公众号同意。
编译:AI100
原文链接:http://julialang.org/blog/2017/03/cudanative