13.3 Haskell中使用OpenCL

Haskell是一种纯函数式语言，其属于类标准ML(SML))模型函数语言的一种。与其他已经提过的语言不同，Haskell(或SML)编程是通过函数式进行描述，应用会通过表达式的参数对表达式做出对应的判断。通常，编程的顺序不同会导致不同的结果。这会使外部声明的值没有进行初始化。这就能看出Haskell类语言的主要优势和劣势。因为Haskell在编程时的劣势很突出，并且其复杂的类型系统，通常会让一些有过C、C++、Java经验的开发者在第一次使用时，有些难以驾驭的感觉。不过，这些问题会在并行程序中解决，例如这个例子，其表达式计算出来的结果是相互独立的，所以其函数式定义是线程安全的。因此，Haskell开发者社区中，逐渐涌现出很多有意思的并行程序。对Haskell感兴趣的读者可以读去一下Hutto写的这本使用Haskell编程的书籍[1]，以及Meijer在微软频道9中的相关教学视频[2]。

通过多方面对Haskell类型系统的论证，不过对于内嵌DSL的设计来说其Haskell类型系统是一个很不错的平台，其能提供相应的抽象化模型，这样的模型能够自动的为GPU编译源码。Accelerate[3]和Obsidian[4]是两个非常不错的例子。不过，本书着重与使用底层OpenCL进行编程，所以我们依旧只会关注，如何让Haskell的编程者通过OpenCL使用GPU。Haskell使用OpenCL，除了带来性能收益，还能获得：

  • OpenCL能够分担目前CPU多线程库的工作负载
  • 高级Haskell语言能够减少OpenCL主机端代码的复杂度，可以创建更加高效的开发环境

目前，有很多Haskell程序已经使用包装好的OpenCL API；不过，我们需要使用FFI的方式对OpenCL进行绑定。另外，我们需要更加简单的使用OpenCL，并且还要能发挥OpenCL强大的计算能力。为了达到这一目标，我们会推荐HOpenCL[5]，其为一个开源项目库，提供了对底层OpenCL的包装，并且提供高级的Haskell接口来访问OpenCL API，其消除了很多与OpenCL平台之间的交互，并且比其他的Haskell实现，额外的提供了更强的静态保证。本节剩下的内容，我们将着重与高级的API；不过，对低级API感兴趣的读者，可以去查看HOpenCL的手册和开发者文档。这里需要注意的是，HOpenCL只支持OpenCL 1.2的API。OpenCL 2.0中所添加的新特性，还未加入HOpenCL中。

我们依旧会使用向量相加作为例子。内核代码不会改变，并且直接嵌入到一个字符串中，而需要改变的则是Haskell主机端代码。

13.3.1 模块结构

HOpenCL将实现的一小部分模块放入Langauage.OpenCL结构体中。

  • Language.OpenCL.Host.Constants:定义了OpenCL核心API所使用的基本类型
  • Language.OpenCL.Host.Core:定义了底层OpenCL核心API
  • Language.OpenCL.GLInterop:定义了OpenGL交互的API
  • Language.OpenCL.Host:定义了高级OpenCL API

后面的几节中将介绍高级API的相关内容，这里我们将会提到主要API的使用，以及相关的注意事项。对底层实现感兴趣的读者，可以阅读HOpenCL的手册和开发者文档[5]。

13.3.2 环境

很多OpenCL函数需要一个上下文对象，其用来定义OpenCL的执行环境；或是需要一个命令队列对象，提交到队列中的任务将在指定的OpenCL环境中执行。很多OpenCL代码中，有些参数函数作为“噪音”——只是技术上需要，其不会对代码进行很大的修改。为了获得这些信息，HOpenCL提供了两个类，Contextual和Queued。这两个类型的实例可以传入相关的OpenCL API中，执行相应的任务。

通常，使用HOpenCL的应用会使用嵌入计算的方式，这就需要在其他计算式中进行计算——例如，将Queued计算嵌入Contextual计算中，然后尝试将其联系在一起。这里提供的with函数，就是用来完成这件事的：

  1. with :: Wraps t m n => t -> m u -> n u

13.3.3 引用计数

所有OpenCL对象的声明周期，都不会定义在一个C代码范围内，C API提供对应的操作，手动的减少引用计数(比如：clRetainContext/clReleaseContext)。HOpenCL使用LifeSpan来完成这一概念的定义，并提供相应的retain和release操作：

  1. retain :: (LifeSpan t, MonadIO m) => t -> m ()
  2. release :: (LifeSpan t, MonadIO m) => t -> m ()

using函数处理构造和释放新(引用计数的)对象。其能够自动对OpenCL对象的声明周期进行管理：

  1. using :: (Lifespan t m, CatchIO m) => m t -> (t -> m u) -> m u

为了更加简单的使用OpenCL上下文(Context)和命令队列(CommandQueue)，其会自动的在HOpenCL中进行引用计数，withNew操作将with和using的功能融合在一起：

  1. withNew :: (Wraps t m n, Lifespan t, CatchIO n) => n t -> m u -> n u

13.3.4 平台和设备

与platforms相关的API函数，可以用来在给定系统中，查找可用平台：

  1. platforms :: MonadIO m => m [Platform]

与C API不同，这里无需对platform查找函数调用两次；HOpenCL将会自动获取全部平台信息。这里唯一麻烦的地方在于，平台信息的返回值为monad m，其是一个MonadIO类的实例，platforms的结果包含在这个实例当中。OpenCL在执行操作时有一定的约束，对于monad对象只能进行输入或输出操作。所有HOpenCL中的OpenCL操作都适用于此限制，所以通过API的方式获取平台信息是不安全的操作，因此需要顺序执行一些操作。

检查完平台信息之后，可以对(?)操作符进行重载，用来决定我们使用对应平台上的哪种实现。例如，下面的代码就代表我们选择第一个平台作为实现平台，且答应出相应供应商：

  1. (p:_) <- platforms
  2. putStrLn .("Platform is by: "++) =<< p ? PlatformVendor

通常，任意OpenCL对象信息都需要通过clGetXXXInfo获取，这里的XXX代表着对应OpenCL类型，这里可以这样实现：

  1. (?) :: MonadIO m => tv -> qt u -> m u

为了平台需要，我们将(?)操作符的类型改一下：

  1. (?) :: MonadIO m => Platform -> PlatformInfo u -> m u

简单的对OpenCL C++包装API的实现(clGetXXXInfo)，可以通过(?)操作符进行相关信息的返回(需要额外的层提供明确的静态类型)。例如，例子中的PlatformVendor，其返回值的类型就是Haskell中的String类型。

devices函数返回与一个平台相关的一系列设备。其将平台对象和设备类型作为参数传入。设备类型只能传入GPU、CPU或ALL。和platforms一样，可以通过(?)操作符对设备信息进行检索：

  1. deviceOfType :: MonadIO m => Platform -> [DeviceType] -> m [Device]

13.3.5 运行环境

如之前所述，主机需要内核执行在另外一个设备上。为了达到这个目的，上下文对象需要在主机端进行配置，并且需要传入命令和数据到设备端。

上下文

context函数可以根据平台和一组设备对象创建出一个上下文对象：

  1. context :: MonadIO m => Platform -> [Device] -> m Context

如果需要严格控制上下文的生命周期——例如，进行图像交互——然后，通过使用contextFromProperties函数将属性传入上下文：

  1. contextFromProperties :: MonadIO m => ContextProperties -> [Device] -> m Context

上下文属性也可以传noProperties(其定义了一组空属性值)，pushContextProperty(其可以添加一个已创建上下文的属性值)。noProperties和pushContextProperty作为Language.OpenCL.Host.Core结构中的一部分:

  1. noProperties :: ContextProperties
  2. pushContextProperty :: ContextProperty t u => t u -> u -> ContextProperties -> ContextProperties

命令队列

要向设备提交命令，就需要创建命令队列。queue函数可以通过当前Contextual创建一个命令队列：

  1. queuue :: Contextual m => Device -> m CommandQueue

命令队列创建后，引用计数开始，并且会向指定Contextual类实例中的设备进行命令的提交。queue函数的实现通常会合并withNew函数，通过嵌入当前上下文创建命令队列：

  1. withNew (queue gpu) $
  2.  __computation dependent on newly created command queue

内存对象

buffer函数将会分配一个OpenCL内存对象，并假设其使用的默认标识。函数bufferWithFlags会通过用户指定的内存标识(MemFlag定义在Language.OpenCL.Host.Constatns中)分配一个内存对象：

  1. buffer :: (Storable t, Contextural m) => Int -> m (Buffer t)
  2. bufferWithFlags :: (Storable t, Contextual m) => Int -> [MemFlag] -> m (Buffer t)

内存对象要和相关的上下文对象相关联，使用using函数可以进行相应的关联操作。

数据从主机传到设备端使用writeTo函数，数据中设备端写回主机端使用readFrom：

  1. readFrom :: (Readable cl hs, Storable t, Queued m) => cl t -> Int -> Int -> m (hs t)
  2. writeTo :: (Writable cl hs, Storable t, Queued m) => cl t -> Int -> hs t -> m Event

创建OpenCL程序对象

OpenCL程序在运行时可以通过两个函数进行编译，programFromSource和buildProgram。先通过源码创建一个OpenCL程序对象，然后对程序对象进行编译。

  1. programFromSource :: Contextual m => String -> m Program
  2. buildProgram :: MonadIO m => Program -> [Device] -> String -> m()

OpenCL内核

内核通过函数kenrel创建：

  1. kernel :: MonadIO m => Program -> String -> m Kernel

参数需要逐个通过函数fixArgument传入。不过，通常参数会在内核在调用前在进行参数传递，并且HOpenCL提供内核invoke函数：

  1. fixArgument :: (KernelArgument a, MonadIO m) => Kernel -> Int -> a -> m()
  2. invoke :: KernelInvocation r => kernel -> r

HOpenCL还提供了另外一种内核调用方式，其可以将内核认为是闭合的，通过setArgs函数对内核的参数进行设置(这种方式在多线程的上下文中十分有用)：

  1. setArgs :: Kernel -> [Co.kernel -> Int -> IO ()] -> Invocation

通过一次调用invoke函数，并不能能够完全将一个内核入队；因此，invoke函数需要和overRange函数一起使用，其会将执行域和结果作为一个事件进行入队：

  1. overRange :: Queued m => Invocation -> ([Int], [Int], [Int]) -> m Event

向量相加的实现源码

下面就是使用HOpenCL实现的向量相加源码：

  1. module VecAdd where
  3. import Language.OpenCL.Host
  4. import Language.OpenCL.Host.FFI
  6. import Control.Monad.Trans (lift0)
  8. source = 
  9. "__kernel void vec add                                 \n" ++
  10. "  __global int *C, __global int *A, __global int *B){ \n" ++
  11. "  int tid = get_global_id(0);                         \n" ++
  12. " C[tid] = A[tid] + B[tid];                            \n" ++
  13. "}"
  15. elements = 2048 :: Int
  17. main = do (p:_) <- platforms
  18.           [gpu] <- devicesOfType p [GPU]
  19.           withNew (context p [gpu]) $
  20.               using (programFromSource source) $ \p ->
  21.               using (buffer elements) $ \inBufA ->
  22.               using (buffer elements) $ \inBufB ->
  23.               using (buffer elements) $ \outBuf ->
  24.                     do { buildProgram p [gpu] ""
  25.                        ; using (kenrel p "vecadd") $ \vecadd ->
  26.                          withNew (queue gpu) $
  27.                            do writeTo inBufA 0 [0.. elements - 1]
  28.                               writeTo inBufB 0 [0.. elements - 1]
  29.                               invoke vecadd outBuf inBufA inBufB
  30.                                     'overRange' ([0], [elements], [1])
  31.                               (x::[Int]) <- readFrom outBuf 0 elements
  32.                               liftIO (if and $ zipWith (\a b -> a == b+b))
  33.                                                     x [0.. elements - 1]
  34.                                   then print "Output is correct"
  35.                                   else print "Output is incorrect")
  36.                     }

[1] G. Hutton. Programming in Haskell, Cambridge University Press, Cambridge, 2007.

[2] E. Meijer, Functional Programming Fundamentals. Channel 9 Lectures, 2009. http://channel9.msdn.com/Series/C9-Lectures-Erik-Meijer-Functional-Programming-Fundamentals/Lecture-Series-Erik-Meijer-Functional-Programming-Fundamentals-Chapter-1

[3] M.M. Chakravarty, G. Keller, S. Lee, T.L. McDonell, V.Grover, Accelerating Haskell array codes with multicore GPUs, in: Processdings on the Sixth Workshop on Declarative Aspects of Multicore Programming, ACM DAMP’11, New York, NY, 2011, pp.3-14

[4] J. Svensson, M. Sheeran, K. Claessen, Obsidian: a domain specific embedded language for parallel programming of graphics processors. in: S.-B. Scholz, O. Chitil(Eds), Implementation and Application of Functional Languages, Lecture Notes in Computer Science, vol.5836, Springer, Berlin/Heidelberg, 2011, pp.156-173

[5] B.R Gaster, J. Garrett Morris, HOpenCL, 2012, https://github.com/bgaster/hopencl.git