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) 上,我们调用以下代码看看:

1
2
3
4
5
6
7
8
9
10
11
12
13
14
//获取所有平台
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}");
}
}

输出:

1
2
3
4
5
6
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来进行处理了。
先来看一下常规的算法:

1
2
3
4
5
6
7
8
9
10
11
//假设为两个规模为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核心代码:

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
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标志的函数才是程序入口。

而下面的就是主机代码:

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
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;
}
}
}

输出:

1
2
3
4
5
6
7
8
9
10
11
12
13
14
运行平台: 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的内存机制有关,有空再讲。

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

评论