> For the complete documentation index, see [llms.txt](https://chenxiaowei.gitbook.io/heterogeneous-computing-with-opencl2-0/llms.txt). Markdown versions of documentation pages are available by appending `.md` to page URLs; this page is available as [Markdown](https://chenxiaowei.gitbook.io/heterogeneous-computing-with-opencl2-0/13.0-chinese/13.3-chinese.md). # 13.3 Haskell中使用OpenCL Haskell是一种纯函数式语言，其属于类\[标准ML(SML)]\( 通过多方面对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. \[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,