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 代码的部分,大概由以下三部分组成:

  1. 与 Julia 交互:重新目的化编译器产生 GPU 与 LLVM 兼容的指令寄存器(CPU 库没有接收到指令,简化的异常处理,……)。
  2. 与 LLVM 交互(使用 LLVM.jl):优化指令寄存器,然后编译成 PTX。
  3. 与 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


原文发布于微信公众号 - AI科技大本营(rgznai100)

原文发表时间:2017-03-18

本文参与腾讯云自媒体分享计划,欢迎正在阅读的你也加入,一起分享。

发表于

我来说两句

0 条评论
登录 后参与评论

相关文章

来自专栏开源FPGA

基于Xilinx FPGA的视频图像采集系统

         本站点博客将逐步迁移至http://ninghechuanblogs.cn/   本篇要分享的是基于Xilinx FPGA的视频图像采集系统,...

55690
来自专栏非著名程序员

超级实用:小而精,优而美的开源库,开年第一篇有意思的分享

昨天是元旦,新年的第一天,估计细心的读者可能发现了,我昨天发起了一个抽奖送书活动,而且发布的时间正好是1月1号11点11分。这可不是一个巧合,这是我故意为之,是...

22260
来自专栏数据的力量

google的搜索技巧

17280
来自专栏逍遥剑客的游戏开发

VR中物理的网络同步

之前做VR游戏时也是尝试了几种物理的同步方案, 最近看到Oculus Blog上也分享了一些, 经验, 做个笔记.。

34360
来自专栏互联网高可用架构

谷歌大数据的三驾马车资料下载

33140
来自专栏微信公众号:Java团长

对高并发流量控制的一点思考

滑动窗口的意思是说把固定时间片,进行划分,并且随着时间的流逝,进行移动,这样就巧妙的避开了计数器的临界点问题。也就是说这些固定数量的可以移动的格子,将会进行计...

26330
来自专栏二进制文集

思维导图学 Maven

14920
来自专栏Linyb极客之路

对高并发流量控制的一点思考

在实际项目中,曾经遭遇过线上5W+QPS的峰值,也在压测状态下经历过10W+QPS的大流量请求,本篇博客的话题主要就是自己对高并发流量控制的一点思考。

11730
来自专栏极客猴

学爬虫之道

Django 已经算是入门,所以自己把学习目标转到爬虫。自己接下来会利用三个月的时间来专攻 Python 爬虫。这几天,我使用“主题阅读方法”阅读 Python...

11020
来自专栏SDNLAB

软件定义光网络故障恢复与资源分配

前言 传统IP分组交换网使用域内路由协议(Interior Gateway Protocol,IGP)和域间路由协议(Border Gateway Protoc...

35190

扫码关注云+社区

领取腾讯云代金券