技术帖:GPU上的高性能C#——Hybridizer

桂林seo / 随笔杂谈 / 时间:2017-12-25 19:59

Hybridizer是来自Altimesh的编译器,可以让人们采用C#代码或.NET程序集编程GPU和其他加速器。Hybridizer使用修饰符号来表达并行性,可以生成针对多核CPU和GPU优化的源代码或二进制文件。在这篇博文中演示了CUDA的目标。

图1 Hybridizer编译管线

图1显示了Hybridizer编译管线。使用Parallel.For之类的并行化模式,或者像在CUDA中一样明确地分配并行工作,可以从加速器的计算能力中受益,而无需了解其内部架构的所有细节。下面是一个使用Parallel.For和lambda的简单示例。

人们可以使用NVIDIA Nsight Visual Studio Edition在GPU上调试和分析这些代码。 Hybridizer实现了先进的C#功能,其中包括虚函数和泛型。

在哪里可以获得Hybridizer

Hybridizer有两个版本:

Hybridizer Software Suite:支持CUDA,AVX,AVX2,AVX512目标和输出源代码。这个源代码需要被审查,这在某些企业(如投资银行)是强制性的。Hybridizer软件套件根据客户要求获得许可。

Hybridizer Essentials:仅启用CUDA目标并仅输出二进制文件。 Hybridizer Essentials是一个免费的Visual Studio扩展程序,没有硬件限制。人们可以在GitHub上找到一组基本代码示例和教育资料。这些样本也可以用来重现其性能结果。

调试和分析

使用调试信息进行编译时,可以在Microsoft Visual Studio中调试Hybridizer C#/ .NET代码,同时在目标硬件上运行优化代码。例如,用C#编写的程序可以在Visual Studio中的C#文件中创建一个断点,可以浏览驻留在GPU上的本地变量和对象数据。

图2 使用Hybridizer和NVIDIA Nsight Visual Studio Edition调试GPU上运行的C#代码

人们可以在复杂项目中集成Hybridizer,即使在代码不可用或模糊的库中,这是因为Hybridizer在MSIL字节码上运行。在博客文章中展示了使用Hybridizer加速AForge图像处理库而没有修改库的能力。在MSIL字节码上运行也支持在.Net虚拟机之上构建的各种语言,比如VB.Net和F#。

所有这些灵活性不会以牺牲性能损失为代价。正如基准测试所示,Hybridizer产生的代码可以像手写代码一样执行。人们可以使用性能分析器(例如NVIDIA Nsight和NVIDIA Visual Profiler)来测量生成的二进制文件的性能,其性能指标指的是原始源代码(例如C#)。

举一个简单的例子:Mandelbrot

作为第一个例子,演示了在NVIDIA GeForce GTX 1080 Ti GPU(Pascal架构,计算能力6.1)上运行的Mandelbrot分形的渲染。

Mandelbrot C#代码

以下代码片断显示了plain C#。它在CPU上平稳运行,没有任何性能损失,因为大多数代码都是属性的修改,在运行时没有任何影响(例如Run方法中的EntryPoint属性)。

[EntryPoint]

public static void Run(float[,] result)

{

int size = result.GetLength(0);

Parallel2D.For(0, size, 0, size, (i, j) => {

float x = fromX + i * h;

float y = fromY + j * h;

result[i, j] = IterCount(x, y);

});

}

public static float IterCount(float cx, float cy)

{

float result = 0.0F;

float x = 0.0f, y = 0.0f, xx = 0.0f, yy = 0.0f;

while (xx + yy <= 4.0f && result < maxiter) {

xx = x * x;

yy = y * y;

float xtmp = xx - yy + cx;

y = 2.0f * x * y + cy;

x = xtmp;

result++;

}

return result;

}

EntryPoint属性告诉Hybridizer生成一个CUDA内核。多维数组映射到内部类型,而Parallel2D.For映射到2D执行网格。给定几行Boilerplate(样板)代码,可以透明地在GPU上运行这些代码。

float[,] result = new float[N,N];

HybRunner runner = HybRunner.Cuda("Mandelbrot_CUDA.dll").SetDistrib(32, 32, 16, 16, 1, 0);

dynamic wrapper = runner.Wrap(new Program());

wrapper.Run(result);

剖析

我们使用Nvidia Nsight Visual Studio Edition分析器来分析此代码。将C#代码链接到CUDA源代码视图中的PTX,如图3所示。

图3在CUDA源代码视图中分析Mandelbrot C#代码

分析器允许与CUDA C ++代码相同的调查级别。

就性能而言,这个例子达到峰值计算FLOP / s的72.5%。这是CUDA C ++人工编写的相同代码的83%。

图4 Profiler输出显示了GPU上Mandelbrot代码的GPU利用率和执行效率。它实现的效率几乎与人工编写CUDA C ++代码一样高效。 使用Hybridizer提供的扩展控件,可以从C#代码中获得更好的性能。如下面的代码所示,其语法与CUDA C ++非常相似。[EntryPoint]

public static void Run(float[] result)

{

for (int i = threadIdx.y + blockIdx.y * blockDim.y; i < N; i += blockDim.y * gridDim.y)

{

for (int j = threadIdx.x + blockIdx.x * blockDim.x; j < N; j += blockDim.x * gridDim.x)

{

float x = fromX + i * h;

float y = fromY + j * h;

result[i * N + j] = IterCount(x, y);

}

}

}

在这个案例中,生成的代码和人工编写的CUDA C ++代码的性能完全相同,达到峰值FLOP/s的87%,如图5所示。

图5分析人工优化的Mandelbrot C#代码

泛型和虚函数

Hybridizer支持设备功能中的泛型和虚函数调用。现代编程语言的这些基本概念有助于代码模块化并提高表达能力。但是,C#中的类型解析是在运行时完成的,这会导致一些性能损失。.NET的泛型可以在保持灵活性的同时实现更高的性能:Hybridizer将泛型映射到C++模板,C ++模板在编译时解析,允许函数内联和过程间优化。另一方面,虚函数调用被映射到其中实例方法被注册的虚函数表。

通过两个属性HybridTemplateConcept和HybridRegisterTemplate(在设备代码中触发实际的模板实例化)给模板实例化提示。作为一个例子,我们来看看两个版本中的一个简单的stream benchmark,一个使用虚函数调用,另一个使用模板映射。该基准依赖于一个通用的接口IMyArray暴露出下标运算符:

[HybridTemplateConcept]

public interface IMyArray {

double this[int index] { get; set; }

}

这些操作符必须与设备功能"Hybridized(杂交)"。为此,我们把Kernel属性放在实现类中。

public class MyArray : IMyArray {

double[] _data;

public MyArray(double[] data) {

_data = data;

}

[Kernel]

public double this[int index] {

get { return _data[index]; }

set { _data[index] = value; }

}

}

虚拟功能调用

在第一个版本中,使用接口编写了一个流算法,没有进一步提示编译器。public class MyAlgorithmDispatch {

IMyArray a, b;

public MyAlgorithmDispatch(IMyArray a, IMyArray b) {

this.a = a;

this.b = b;

}

[Kernel]

public void Add(int n) {

IMyArray a = this.a;

IMyArray b = this.b;

for (int k = threadIdx.x + blockDim.x * blockIdx.x;

k < n;

k += blockDim.x * gridDim.x) {

a[k] += b[k];

}

}

}

因为把a和b上的下标运算符称为接口,所以在MSIL中有一个callvirt。IL_002a: ldloc.3

IL_002b: ldloc.s 4

IL_002d: callvirt instance float64 Mandelbrot.IMyArray::get_Item(int32)

IL_0032: ldloc.1

IL_0033: ldloc.2

IL_0034: callvirt instance float64 Mandelbrot.IMyArray::get_Item(int32)

IL_0039: add

IL_003a: callvirt instance void Mandelbrot.IMyArray::set_Item(int32, float64检查生成的二进制表明Hybridizer在虚函数表中生成了一个查找,如图6所示。

图6. PTX中的虚函数调用 这个版本的算法消耗32个寄存器,并获得271GB/s的带宽,如图7所示。在同一硬件上,CUDA Toolkit中的带宽测试示例达到352GB/s。

图7由于虚函数调用而实现的低带宽

虚函数表会导致更多的注册压力,并防止内联。

通用要求

采用泛型写了第二个版本,要求Hybridizer生成模板代码。

[HybridRegisterTemplate(Specialize = typeof(MyAlgorithm))]

public class MyAlgorithm where T : IMyArray

{

T a, b;

[Kernel]

public void Add(int n)

{

T a = this.a;

T b = this.b;

for (int k = threadIdx.x + blockDim.x * blockIdx.x;

k < n;

k += blockDim.x * gridDim.x)

a[k] += b[k];

}

}

public MyAlgorithm(T a, T b)

{

this.a = a;

this.b = b;

}

}

使用RegisterTemplate属性,Hybridizer将生成相应的模板实例。然后生成内联函数的调用,如图8所示。

图8使用泛型参数生成内联函数调用,而不是虚函数表查找

其通用参数性能要好得多,达到339GB/s,性能提高25%(如图9所示),带宽测试为96%。

图9 由于函数内联,泛型实现了更高的带宽开始使用Hybridizer

Hybridizer支持各种C#特性,允许代码分解和表达。Visual Studio和Nsight(调试器和分析器)中的集成为人们提供了一个安全高效的开发环境。即使在非常复杂的高度定制的代码上,Hybridizer也可以实现出色的GPU性能。

人们可以从Visual Studio Marketplace下载Hybridizer Essentials。查看在github上的SDK。


桂林SEO半杯酒博客文章,转载请注明原文网址摘自 http://www.mna5.com/suibizatan/663.html,谢谢配合!

阅读延展

微信扫一扫,关注我们
1
3