# 13.3 Haskell中使用OpenCL

Haskell是一种纯函数式语言，其属于类\[标准ML(SML)]\([https://en.wikipedia.org/wiki/ML\_(programming\_language))模型函数语言的一种。与其他已经提过的语言不同，Haskell(或SML)编程是通过函数式进行描述，应用会通过表达式的参数对表达式做出对应的判断。通常，编程的顺序不同会导致不同的结果。这会使外部声明的值没有进行初始化。这就能看出Haskell类语言的主要优势和劣势。因为Haskell在编程时的劣势很突出，并且其复杂的类型系统，通常会让一些有过C、C++、Java经验的开发者在第一次使用时，有些难以驾驭的感觉。不过，这些问题会在并行程序中解决，例如这个例子，其表达式计算出来的结果是相互独立的，所以其函数式定义是线程安全的。因此，Haskell开发者社区中，逐渐涌现出很多有意思的并行程序。对Haskell感兴趣的读者可以读去一下Hutto写的这本使用Haskell编程的书籍\[1\]，以及Meijer在微软频道9中的相关教学视频\[2\]。](https://en.wikipedia.org/wiki/ML_\(programming_language\)\)%E6%A8%A1%E5%9E%8B%E5%87%BD%E6%95%B0%E8%AF%AD%E8%A8%80%E7%9A%84%E4%B8%80%E7%A7%8D%E3%80%82%E4%B8%8E%E5%85%B6%E4%BB%96%E5%B7%B2%E7%BB%8F%E6%8F%90%E8%BF%87%E7%9A%84%E8%AF%AD%E8%A8%80%E4%B8%8D%E5%90%8C%EF%BC%8CHaskell\(%E6%88%96SML\)%E7%BC%96%E7%A8%8B%E6%98%AF%E9%80%9A%E8%BF%87%E5%87%BD%E6%95%B0%E5%BC%8F%E8%BF%9B%E8%A1%8C%E6%8F%8F%E8%BF%B0%EF%BC%8C%E5%BA%94%E7%94%A8%E4%BC%9A%E9%80%9A%E8%BF%87%E8%A1%A8%E8%BE%BE%E5%BC%8F%E7%9A%84%E5%8F%82%E6%95%B0%E5%AF%B9%E8%A1%A8%E8%BE%BE%E5%BC%8F%E5%81%9A%E5%87%BA%E5%AF%B9%E5%BA%94%E7%9A%84%E5%88%A4%E6%96%AD%E3%80%82%E9%80%9A%E5%B8%B8%EF%BC%8C%E7%BC%96%E7%A8%8B%E7%9A%84%E9%A1%BA%E5%BA%8F%E4%B8%8D%E5%90%8C%E4%BC%9A%E5%AF%BC%E8%87%B4%E4%B8%8D%E5%90%8C%E7%9A%84%E7%BB%93%E6%9E%9C%E3%80%82%E8%BF%99%E4%BC%9A%E4%BD%BF%E5%A4%96%E9%83%A8%E5%A3%B0%E6%98%8E%E7%9A%84%E5%80%BC%E6%B2%A1%E6%9C%89%E8%BF%9B%E8%A1%8C%E5%88%9D%E5%A7%8B%E5%8C%96%E3%80%82%E8%BF%99%E5%B0%B1%E8%83%BD%E7%9C%8B%E5%87%BAHaskell%E7%B1%BB%E8%AF%AD%E8%A8%80%E7%9A%84%E4%B8%BB%E8%A6%81%E4%BC%98%E5%8A%BF%E5%92%8C%E5%8A%A3%E5%8A%BF%E3%80%82%E5%9B%A0%E4%B8%BAHaskell%E5%9C%A8%E7%BC%96%E7%A8%8B%E6%97%B6%E7%9A%84%E5%8A%A3%E5%8A%BF%E5%BE%88%E7%AA%81%E5%87%BA%EF%BC%8C%E5%B9%B6%E4%B8%94%E5%85%B6%E5%A4%8D%E6%9D%82%E7%9A%84%E7%B1%BB%E5%9E%8B%E7%B3%BB%E7%BB%9F%EF%BC%8C%E9%80%9A%E5%B8%B8%E4%BC%9A%E8%AE%A9%E4%B8%80%E4%BA%9B%E6%9C%89%E8%BF%87C%E3%80%81C++%E3%80%81Java%E7%BB%8F%E9%AA%8C%E7%9A%84%E5%BC%80%E5%8F%91%E8%80%85%E5%9C%A8%E7%AC%AC%E4%B8%80%E6%AC%A1%E4%BD%BF%E7%94%A8%E6%97%B6%EF%BC%8C%E6%9C%89%E4%BA%9B%E9%9A%BE%E4%BB%A5%E9%A9%BE%E9%A9%AD%E7%9A%84%E6%84%9F%E8%A7%89%E3%80%82%E4%B8%8D%E8%BF%87%EF%BC%8C%E8%BF%99%E4%BA%9B%E9%97%AE%E9%A2%98%E4%BC%9A%E5%9C%A8%E5%B9%B6%E8%A1%8C%E7%A8%8B%E5%BA%8F%E4%B8%AD%E8%A7%A3%E5%86%B3%EF%BC%8C%E4%BE%8B%E5%A6%82%E8%BF%99%E4%B8%AA%E4%BE%8B%E5%AD%90%EF%BC%8C%E5%85%B6%E8%A1%A8%E8%BE%BE%E5%BC%8F%E8%AE%A1%E7%AE%97%E5%87%BA%E6%9D%A5%E7%9A%84%E7%BB%93%E6%9E%9C%E6%98%AF%E7%9B%B8%E4%BA%92%E7%8B%AC%E7%AB%8B%E7%9A%84%EF%BC%8C%E6%89%80%E4%BB%A5%E5%85%B6%E5%87%BD%E6%95%B0%E5%BC%8F%E5%AE%9A%E4%B9%89%E6%98%AF%E7%BA%BF%E7%A8%8B%E5%AE%89%E5%85%A8%E7%9A%84%E3%80%82%E5%9B%A0%E6%AD%A4%EF%BC%8CHaskell%E5%BC%80%E5%8F%91%E8%80%85%E7%A4%BE%E5%8C%BA%E4%B8%AD%EF%BC%8C%E9%80%90%E6%B8%90%E6%B6%8C%E7%8E%B0%E5%87%BA%E5%BE%88%E5%A4%9A%E6%9C%89%E6%84%8F%E6%80%9D%E7%9A%84%E5%B9%B6%E8%A1%8C%E7%A8%8B%E5%BA%8F%E3%80%82%E5%AF%B9Haskell%E6%84%9F%E5%85%B4%E8%B6%A3%E7%9A%84%E8%AF%BB%E8%80%85%E5%8F%AF%E4%BB%A5%E8%AF%BB%E5%8E%BB%E4%B8%80%E4%B8%8BHutto%E5%86%99%E7%9A%84%E8%BF%99%E6%9C%AC%E4%BD%BF%E7%94%A8Haskell%E7%BC%96%E7%A8%8B%E7%9A%84%E4%B9%A6%E7%B1%8D\[1]%EF%BC%8C%E4%BB%A5%E5%8F%8AMeijer%E5%9C%A8%E5%BE%AE%E8%BD%AF%E9%A2%91%E9%81%939%E4%B8%AD%E7%9A%84%E7%9B%B8%E5%85%B3%E6%95%99%E5%AD%A6%E8%A7%86%E9%A2%91\[2]%E3%80%82)

通过多方面对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函数，就是用来完成这件事的：

```haskell
with :: Wraps t m n => t -> m u -> n u
```

## 13.3.3 引用计数

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

```haskell
retain :: (LifeSpan t, MonadIO m) => t -> m ()
release :: (LifeSpan t, MonadIO m) => t -> m ()
```

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

```haskell
using :: (Lifespan t m, CatchIO m) => m t -> (t -> m u) -> m u
```

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

```haskell
withNew :: (Wraps t m n, Lifespan t, CatchIO n) => n t -> m u -> n u
```

## 13.3.4 平台和设备

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

```haskell
platforms :: MonadIO m => m [Platform]
```

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

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

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

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

```haskell
(?) :: MonadIO m => tv -> qt u -> m u
```

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

```haskell
(?) :: MonadIO m => Platform -> PlatformInfo u -> m u
```

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

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

```haskell
deviceOfType :: MonadIO m => Platform -> [DeviceType] -> m [Device]
```

## 13.3.5 运行环境

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

**上下文**

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

```haskell
context :: MonadIO m => Platform -> [Device] -> m Context
```

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

```haskell
contextFromProperties :: MonadIO m => ContextProperties -> [Device] -> m Context
```

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

```haskell
noProperties :: ContextProperties
pushContextProperty :: ContextProperty t u => t u -> u -> ContextProperties -> ContextProperties
```

**命令队列**

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

```haskell
queuue :: Contextual m => Device -> m CommandQueue
```

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

```haskell
withNew (queue gpu) $
 __computation dependent on newly created command queue
```

**内存对象**

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

```haskell
buffer :: (Storable t, Contextural m) => Int -> m (Buffer t)
bufferWithFlags :: (Storable t, Contextual m) => Int -> [MemFlag] -> m (Buffer t)
```

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

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

```haskell
readFrom :: (Readable cl hs, Storable t, Queued m) => cl t -> Int -> Int -> m (hs t)
writeTo :: (Writable cl hs, Storable t, Queued m) => cl t -> Int -> hs t -> m Event
```

**创建OpenCL程序对象**

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

```haskell
programFromSource :: Contextual m => String -> m Program
buildProgram :: MonadIO m => Program -> [Device] -> String -> m()
```

**OpenCL内核**

内核通过函数kenrel创建：

```haskell
kernel :: MonadIO m => Program -> String -> m Kernel
```

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

```haskell
fixArgument :: (KernelArgument a, MonadIO m) => Kernel -> Int -> a -> m()
invoke :: KernelInvocation r => kernel -> r
```

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

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

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

```haskell
overRange :: Queued m => Invocation -> ([Int], [Int], [Int]) -> m Event
```

**向量相加的实现源码**

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

```haskell
module VecAdd where

import Language.OpenCL.Host
import Language.OpenCL.Host.FFI

import Control.Monad.Trans (lift0)

source = 
"__kernel void vec add                                 \n" ++
"  __global int *C, __global int *A, __global int *B){ \n" ++
"  int tid = get_global_id(0);                         \n" ++
" C[tid] = A[tid] + B[tid];                            \n" ++
"}"

elements = 2048 :: Int

main = do (p:_) <- platforms
          [gpu] <- devicesOfType p [GPU]
          withNew (context p [gpu]) $
              using (programFromSource source) $ \p ->
              using (buffer elements) $ \inBufA ->
              using (buffer elements) $ \inBufB ->
              using (buffer elements) $ \outBuf ->
                    do { buildProgram p [gpu] ""
                       ; using (kenrel p "vecadd") $ \vecadd ->
                         withNew (queue gpu) $
                           do writeTo inBufA 0 [0.. elements - 1]
                              writeTo inBufB 0 [0.. elements - 1]
                              invoke vecadd outBuf inBufA inBufB
                                    'overRange' ([0], [elements], [1])
                              (x::[Int]) <- readFrom outBuf 0 elements
                              liftIO (if and $ zipWith (\a b -> a == b+b))
                                                    x [0.. elements - 1]
                                  then print "Output is correct"
                                  else print "Output is incorrect")
                    }
```

\[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>


---

# Agent Instructions: Querying This Documentation

If you need additional information that is not directly available in this page, you can query the documentation dynamically by asking a question.

Perform an HTTP GET request on the current page URL with the `ask` query parameter:

```
GET https://chenxiaowei.gitbook.io/heterogeneous-computing-with-opencl2-0/13.0-chinese/13.3-chinese.md?ask=<question>
```

The question should be specific, self-contained, and written in natural language.
The response will contain a direct answer to the question and relevant excerpts and sources from the documentation.

Use this mechanism when the answer is not explicitly present in the current page, you need clarification or additional context, or you want to retrieve related documentation sections.
