OpenCL with CLOO

CLOO 是一个对 OpenCL 的 .Net 封装，可以让 .Net/Mono 程序充分使用 OpenCL 的优势，易用、开源。

今天用 OpenCL 中的 “Hello World” 程序 - 矩阵乘法，来简单介绍一下 OpenCL。

OpenCL 是一个开放的工业标准，既然是开放的，那么，所有厂商就可以提供自己的实现，例如英特尔，英伟达等等。也正因为如此，也导致在同一台机器上可能存在多个支持不同版本的硬件。例如，英伟达的GPU到现在也才支持 OpenCL 1.2 版本，但是OpenCL都已经出到2.x版本了。

我们可以查看当前设备中，有哪些厂商提供了 OpenCL 的支持，以及运算平台是啥。例如，在 Surface Book (with Nvidia GPU) 上，我们调用以下代码看看：

//获取所有平台
var platforms = ComputePlatform.Platforms;

foreach(var platform in platforms)
{
    Console.WriteLine($"{platform.Name},{platform.Version}");

    //获取该平台下的计算设备
    var devices = platform.QueryDevices();
    foreach(var device in devices)
    {
        Console.WriteLine($"    Device:{device.Name}");
    }
}

输出：

Intel(R) OpenCL, OpenCL 2.0
    Device:Intel(R) HD Graphics 520
    Device:Intel(R) Core(TM) i7-6600U CPU @ 2.60GHz

NVIDIA CUDA, OpenCL 1.2 CUDA 8.0.0
    Device:GeForce GPU

可以看出这台机器上有两个支持 OpenCL 的平台，英特尔的CPU和英伟达的GPU，其中，INTEL CPU 里又有两个计算设备 CPU 和集成的核显 HD520。
毫无疑问，在 OpenCL 上，英伟达的GPU可以提供比 Intel CPU 高得多的性能，毕竟，GPU的流处理器数量要比CPU上那可怜的核显上的要多得多。

主机代码(host)和核心代码(kernel)

OpenCL分有主机代码和核心代码，相对于核心代码，主机代码用于对环境进行初始化，例如配置运行平台，计算设备等的。而核心代码，就是运行在指定计算设备的代码，例如GPU。这就好像DirectX上的HLSL一样，HLSL的代码只运行在GPU上。

OpenCL是一个动态编译的框架，就是说核心代码是在运行时才被编译的，程序运行时，在核心代码还没被编译前我们都可以更改核心代码。

矩阵乘法

矩阵乘法，以最普通的算法来进行计算，是时间复杂度为O(n^3)的算法，时间都花在做重复的加法和乘法运算，这种情景最适合用GPU来进行处理了。
先来看一下常规的算法：

//假设为两个规模为Rank的方形矩阵
int matrixA[Rank][Rank];
int sum[Rank][Rank];

void Mul()
{
    for (int i = 0; i < Rank; i++)
        for (int j = 0; j < Rank; j++)
            for (int k = 0; k < Rank; k++)
                sum[i][j] += matrixA[i][k] * matrixA[k][j];
}

分析：

我们要让这段代码在GPU上快速运行，当然，我们不能直接就让这段代码作为我们的OpenCL核心代码，因为这样并不能发挥GPU同步计算的优势。我们让想办法将整个流程拆分为各个可以同时进行的子运算，让GPU同步并行计算，以此降低计算耗时。

如下，是这个程序的OpenCL核心代码：

kernel void MatrixMul(
	global write_only int* result,
	global read_only int* matrix,
	int rank)
{
    int gx=get_global_id(0);
    int gy=get_global_id(1);

    int loc=gx*rank+gy;

    for(int i=0;i<rank;i++)
    {
        int leftLoc=gx*rank+i;
        int rightLoc=i*rank+gy;
        result[loc]+=matrix[leftLoc]*matrix[rightLoc];
    }
}

这就是实际上跑在GPU上的代码。前面提到，OpenCL是动态编译的框架，这段代码你可以放在任何一个文本文件中，但是有一个注意的是，只有kernel标志的函数才是程序入口。

而下面的就是主机代码：

using System;
using System.IO;
using System.Linq;
using System.Runtime.InteropServices;
using Cloo;
using System.Diagnostics;

namespace MatrixDemo
{
    class Program
    {
        static void Main(string[] args)
        {
            Matrix_mul_opencl(500);
            Matrix_mul_opencl(1000);
            Matrix_mul_opencl(1500);
        }

        static void Matrix_mul_opencl(int r)
        {
            //选取设备
            var platform = ComputePlatform.Platforms.FirstOrDefault();
            var device = platform.Devices.FirstOrDefault();

            //设置相关上下文
            var properties = new ComputeContextPropertyList(platform);
            var context = new ComputeContext(new[] { device }, properties, null, IntPtr.Zero);

            //命令队列，用于控制执行的代码
            ComputeCommandQueue commands = new ComputeCommandQueue(context, context.Devices[0],
                ComputeCommandQueueFlags.None);

            //读取opencl代码
            var code = File.ReadAllText("matrix.cl");

            //编译
            var program = new ComputeProgram(context, code);
            try
            {
                program.Build(null, null, null, IntPtr.Zero);
            }
            catch (Exception ex)
            {
                throw;
            }

            //创建核心代码，就是cl代码中以kernel标识，函数签名为MatrixMul的函数
            var kernel = program.CreateKernel("MatrixMul");

            //矩阵规模
            int rank = r;
            //储存计算结果的数组
            var result = new ComputeBuffer<int>(context, ComputeMemoryFlags.WriteOnly, rank * rank);
            var matrix = CreateMatrix(context, rank);

            //创建的核心代码函数以这种方式来传参
            kernel.SetMemoryArgument(0, result);
            kernel.SetMemoryArgument(1, matrix);
            kernel.SetValueArgument(2, rank);

            Console.WriteLine($"运行平台: {platform.Name}\n运行设备： {device.Name}\n矩阵规模: {rank}x{rank}");

            Stopwatch sw = Stopwatch.StartNew();

            //执行代码
            commands.Execute(kernel, null, new long[] { rank, rank }, null, null);

            int[] resultArray = new int[rank * rank];
            var arrHandle = GCHandle.Alloc(resultArray, GCHandleType.Pinned);

            //读取数据
            commands.Read(result, true, 0, rank * rank, arrHandle.AddrOfPinnedObject(), null);

            var elapsed = sw.Elapsed;
            Console.WriteLine($"耗时: {elapsed.TotalMilliseconds} ms\n");

            arrHandle.Free();

            kernel.Dispose();
        }

        //创建实验矩阵
        static ComputeBuffer<int> CreateMatrix(ComputeContext context, int rank)
        {
            int size = rank * rank;
            int[] datas = new int[size];
            for (int i = 0; i < size; i++)
            {
                datas[i] = i % 10;
            }
            var matrix = new ComputeBuffer<int>(context, ComputeMemoryFlags.CopyHostPointer, datas);
            return matrix;
        }
    }
}

输出：

运行平台: NVIDIA CUDA
运行设备： GeForce GTX 1060 6GB
矩阵规模: 500x500
耗时: 28.7764 ms

运行平台: NVIDIA CUDA
运行设备： GeForce GTX 1060 6GB
矩阵规模: 1000x1000
耗时: 224.5405 ms

运行平台: NVIDIA CUDA
运行设备： GeForce GTX 1060 6GB
矩阵规模: 1500x1500
耗时: 1024.3869 ms

对于矩阵规模为500x500的计算，耗时28毫秒，1500的规模则需要1024毫秒，这性能，比用常规方法跑在CPU上，耗时要18秒的情况不知道高到哪里去了。但是有个诡异的事情是，在Windows 10上，当我尝试将矩阵规模提升到2000或以上时，opencl会抛出OutOfResource异常，但是相同规模的矩阵乘法运算，用C++ AMP计算跑在相同的设备上却没有问题，恩，或许这个跟英伟达的显卡驱动或者Windows 10的内存机制有关，有空再讲。

那还有没有更快的方法呢？当然有了，我们下篇文章再谈。