Julia GPU：Julia语言如何让程序员轻松使用GPU

加入Not a Monad Tutorial电报小组或频道，讨论编程、计算机科学和论文。那里见！

如果您正在寻找优秀的工程师，请给我发一封电子邮件到[email protected]，或者您也可以通过Twitter联系我，电子邮件地址是@Federation iccarrone。

我们生活在一个每天都在创建越来越多的数据的时代，也有新的技术和复杂的算法试图从这些数据中提取出最大的价值。因此，CPU能力正在接近其计算能力的瓶颈。GPU计算很早以前就开启了高性能和并行计算的新范式，但直到最近才开始大规模用于数据科学。在这次采访中，JuliaGPU项目的主要贡献者之一Tim Besard不仅从性能角度，而且从用户角度，深入探讨了有关GPU计算的一些细节，以及使Julia成为适合这类任务的语言的特性。

请给我们介绍一下你自己。你的背景是什么？你现在的职位是什么？

我一直对系统编程感兴趣，在获得CS学位后，我有机会开始在比利时根特大学攻读博士学位，就在2012年左右Julia第一次获释的时候。该语言似乎很有趣，因为我想获得一些LLVM的经验，所以我决定将一些图像处理研究代码从MATLAB和C++移植到Julia。我们的目标是与C++版本的性能相匹配，但它的一些内核是在CUDAC…中实现的。所以很明显，Julia需要一个GPU后端！

当然，这说起来容易做起来难，我的博士学位主要是关于实现后端和(重新)构造现有的Julia编译器以促进这些额外的后端。现在我在Julia Computing，在那里我仍然在做所有与GPU相关的事情。

JuliaGPU是我们用来对Julia中的GPU相关资源进行分组的名称：有一个托管大多数软件包的GitHub组织，有一个为新用户指路的网站，我们有JuliaGPU项目的CI基础设施，有一个松弛频道和话语类别，等等。

所有这些都是为了让各种用户更容易使用GPU。目前的技术往往设置了很大的进入壁垒：CUDA安装相当棘手，许多用户不熟悉C和C++，等等。通过我们作为JuliaGPU组织的一部分开发的软件，我们的目标是使GPU的使用变得容易，而不会阻碍优化或使用硬件必须提供的低级功能的能力。

GPU计算是指使用GPU(一种最初设计用于图形处理的设备)执行通用计算。由于CPU性能不再像过去那样稳步提高，这一点变得更加重要。取而代之的是，越来越多地使用诸如GPU或FPGA之类的专用设备来提高某些计算的性能。在使用GPU的情况下，该架构非常适合执行高度并行的应用程序。机器学习网络就是这种并行应用的一个很好的例子，它们的普及是GPU变得如此重要的原因之一。

Julia的主要优点是该语言是为编译而设计的。尽管语法是高级的，但是生成的机器码是紧凑的，并且具有很好的性能特征(有关更多详细信息，请参阅本文)。这对于GPU执行是至关重要的，在GPU执行中，我们需要运行本机二进制文件，并且不能像其他语言的语义通常所要求的那样容易(或有效地)解释代码。

因为我们能够直接为GPU编译Julia，所以我们几乎可以使用该语言的所有功能来构建强大的抽象。例如，您可以定义自己的类型，在GPU数组中使用这些类型，将其与现有的抽象(如惰性转置包装)组合在一起，在GPU上访问这些类型，同时受益于自动边界检查(如果需要)，等等。

从Python程序员的角度来看，CUDA.jl与PyCUDA相比如何？它们的功能是否相同？

PyCUDA为程序员提供了对CUDAAPI的访问，并提供了更易于使用的高级Python函数。Jl提供相同的功能，但在Julia中提供。PyCUDA主页上的“hello world”与Julia几乎一模一样：

使用CUDA函数Multiply_Them(DEST，a，b)i=threadIdx().x DEST[i]=a[i]*b[i]return end a=CuArray(randn(Float32,400))b=CuArray(randn(Float32,400))DEST=相似(A)@Cuda线程=400 Multiply_Them(DEST，a，b)println(DEST-a.*b)。

有一个非常大的区别："；Multiply_Them"；这里是用Julia编写的函数，而PyCUDA使用的是用CUDA C编写的内核。原因很简单：Python不容易编译。当然，像Numba这样的项目证明了这样做的可能性很大，但最终这些都是独立的编译器，它们试图尽可能地与参考Python编译器相匹配。使用CUDA.jl，我们可以与该引用编译器集成，因此更容易保证语义一致，并在语言更改时遵循(有关更多详细信息，请参阅本文)。

不用谢。CUDA.jl面向不同类型的(GPU)程序员。如果您对编写自己的内核很有信心，那么您可以这样做，同时使用CUDA GPU必须提供的所有低级功能。但是，如果您是GPU编程的新手，您可以使用使用CUDA.jl中现有内核的高级数组操作。例如，上面的逐元素乘法可以写成：

是否需要知道如何在CUDA.jl中编码才能充分利用Julia中的GPU计算？

对于大多数用户来说并非如此。Julia有一种强大的泛型数组操作语言("；map&34；，"；Reduce"；，"；Broadcast"；，"；Columate"；等)，可以应用于所有类型的数组，包括GPU数组。这意味着您经常可以通过CUDA.jl重用为CPU开发的代码库(本文展示了一些强大的示例)。这样做通常只需要最少的更改：更改数组类型，确保使用数组操作而不是for循环，等等。

您可能需要超越这种编程风格，例如，因为您的应用程序没有清晰地映射到数组操作，使用特定的GPU功能等。在这种情况下，一些关于CUDA和GPU编程模型的基本知识就足以用CUDA.jl编写内核。

与CUDAC相比，在CUDA.jl中编写内核的体验如何？这些知识彼此之间的传递性如何？

它非常相似，这是经过设计的：我们试图使CUDA.jl中的内核抽象接近它们的CUDAC对应物，以便现有的GPU程序员熟悉编程环境。当然，通过使用高级源语言可以提高生活质量。例如，您可以像在CUDAC中那样静态和动态地分配共享内存，但是我们使用的不是原始指针，而是一个可以轻松索引的N维数组对象。NVIDIA开发人员博客中的一个示例：

__global__void staticReverse(int*d，int n){__Shared__int s[64]；int t=threadIdx.x；int tr=n-t-1；s[t]=d[t]；__syncthread()；d[t]=s[tr]；}。

此内核的CUDA.jl等效项看起来非常熟悉，但是使用数组对象而不是原始指针：

函数静态Reverse(D)s=@cuStaticSharedMem(Int，64)t=threadIdx().x tr=长度(D)-t+1 s[t]=d[t]sync_thread()d[t]=s[tr]返回端

使用数组对象有很多优点，例如，多维被极大地简化，我们可以只做#34；d[i，j]"；。但它也更安全，因为这些访问是经过边界检查的：

Julia>；a=cuArray(1：64)64元素cuArray{Int64，1}：1 2 3⋮62 63 64 julia>；@cuda thread=65 staticReverse(A)错误：内核执行期间引发异常。堆栈跟踪：[1]在abstractarray.jl：541处抛出_rangserror。

当然，边界检查不是免费的，一旦我们确定我们的代码是正确的，我们就可以向内核添加@inbound；注释，并获得我们期望的高性能代码：

Julia>；@device_code_ptx@cuda线程=64静态反向(A).可见.entry静态反向(.param.align 8.b8 d[16]){.reg.b32%r<；2>；；.reg.b64%rd<；15>；；.Shared.align 32.b8 s[512]；mov.b64%rd1，d；ld.param.u64%rd2，[%rd1]；ld.param.u64%rd3，[%rd1+8]；Mov.u32%r1，%tid.x；cvt.u64.u32%rd4，%r1；mul.wide.u32%rd5，%r1，8；add.s64%rd6，%rd5，-8；add.s64%rd7，%rd3，%rd6；ld.global.u64%rd8，[%rd7+8]；mov.u64%rd9，s；add.s64%rd10，%rd9，%rd6；st.shared.u64[%rd10+8]，%rd8；bar.sync 0；S64%rd11，%rd2，%rd4；shl.b64%rd12，%rd11，3；add.s64%rd13，%rd9，%rd12；ld.shared.u64%rd14，[%rd13+-8]；st.global.u64[%rd7+8]，%rd14；ret；}julia>；a 64元素立方体数组{Int64，1}：64 63 62⋮3 2 1。

像"；@DEVICE_CODE_PTX"；这样的工具使有经验的开发人员可以轻松地检查生成的代码，并确保编译器执行他想要的操作。

为什么拥有编译器会对CUDA.jl这样的库产生如此大的影响？(将其集成到Julia编译器的过程如何？)。

因为我们可以使用编译器，所以我们可以依赖高阶函数和其他基于用户提供的参数进行专门化的泛型抽象。这大大简化了我们的库，但也为用户提供了非常强大的工具。例如，我们仔细实现了一个使用共享内存、WARP内部函数等来执行高性能缩减的`mapduce‘函数。不过，该实现是泛型的，并且将根据函数的参数自动重新专门化(即使在运行时也是如此)：

有了这个由经验丰富的GPU程序员实现的强大的“地图缩减”抽象，其他开发人员就可以在没有这种经验的情况下创建派生的抽象。例如，让我们实现一个`count`函数，该函数计算一个谓词有多少项为真：

Count(谓词，数组)=MapReduce(谓词，+，数组)Julia>；a=CUDA.rand(Int8，4)4元素CuArray{Int8，1}：51 3 70 100 Julia>；count(iseven，a)2。

尽管没有为`Int8`类型或`iseven`谓词专门实现我们的`mapduce`实现，但Julia编译器会自动专门化该实现，从而导致内核针对此特定调用进行优化。

在为JuliaGPU开发软件包，特别是用Julia等高级编程语言编写CUDA.jl这样的低级软件包时，最大的挑战是什么？

最初的大部分工作都集中在开发工具上，这些工具使得用Julia编写低级代码成为可能。例如，我们开发了允许我们访问LLVMAPI的LLVM.jl包。最近，我们的重点已转向推广此功能，以便其他GPU后端，如AMDGPU.jl或oneAPI.jl可以从CUDA.jl的开发中受益。例如，供应商中立的数组操作现在在GPUArrays.jl中实现，而共享编译器功能现在在GPUCompiler.jl中实现。这应该可以在多个GPU后端上工作，即使大多数后端只由一个开发人员维护。

关于JuliaGPU博客中宣布的关于多设备编程的最新版本，这个新功能解决了哪些困难？这与需要大量计算资源的行业相关吗？

在工业或大型研究实验室中，MPI通常用于跨多个节点或GPU分配工作。Julia的MPI.jl支持该用例，并在必要时与CUDA.jl集成。CUDA 1.3新增的多设备功能还使得在单个进程中使用多个GPU成为可能。它很好地映射了Julia的基于任务的并发性，并使得在单个节点内分发工作变得容易：

没有任何具体的路线图，但即将到来的一个主要功能是适当支持降低精度的输入，如16位浮点。我们已经支持CUBLAS或CUDNN支持的Float16数组，但是Julia的下一个版本将使编写对这些值进行操作的内核成为可能。

除此之外，特性是原封不动的：-)一定要订阅JuliaGPU博客，在那里我们会为Julia的GPU后端的每个主要版本发布一篇简短的帖子。

一个快乐的Erlang，Rust/ML和Lisp福音派Strikeforce的成员。网络协议的RFC狂热分子。大数据和机器学习。

关于编程语言、操作系统、网络协议、人工智能和机器学习的写作、评论和采访。

一个快乐的Erlang，Rust/ML和Lisp福音派Strikeforce的成员。网络协议的RFC狂热分子。大数据和机器学习。

关于编程语言、操作系统、网络协议、人工智能和机器学习的写作、评论和采访。

媒体是一个开放的平台，1.7亿读者来这里发现有洞察力和动态的思维。在这里，专家和未被发现的声音都会深入到任何话题的核心，并将新想法浮出水面。了解更多。

关注对你来说重要的作家、出版物和话题，你会在你的主页和收件箱里看到它们。探索。

如果你有故事要讲，有知识要分享，或者有观点要提供-欢迎回家。在任何话题上发表你的想法都很容易，也很自由。在介质上写入