登录 / 注册
IT168技术开发频道
IT168首页 > 技术开发 > 技术开发技术 > 正文

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

2017-12-25 13:31    it168网站 原创  作者: smart编译 编辑: 田晓旭

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

  Hybridizer:GPU上的高性能C

  图1 Hybridizer编译管线

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

  Hybridizer:GPU上的高性能C

  人们可以使用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上的本地变量和对象数据。

  Hybridizer:GPU上的高性能C

图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所示。

  Hybridizer:GPU上的高性能C

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

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

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

  Hybridizer:GPU上的高性能C

图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所示。

  Hybridizer:GPU上的高性能C

图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所示。

  Hybridizer:GPU上的高性能C

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

  Hybridizer:GPU上的高性能C

  图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所示。

  Hybridizer:GPU上的高性能C

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

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

  Hybridizer:GPU上的高性能C

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

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

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

标签: 编译器 , C#
  • IT168企业级IT168企业级
  • IT168文库IT168文库

扫一扫关注

行车视线文章推荐

首页 评论 返回顶部