高性能计算

Hybridizer:在 GPU 上运行高性能 C#

 
Figure 1. The Hybridizer Pipeline.
图 1 。杂交剂管道。

Hybridizer 是一个来自 Altimesh 的编译器,它允许您从 C 代码或. NET 程序集编程 GPUs 和其他加速器。使用修饰符号来表示并行性, Hybridizer 生成针对多核 CPUs 和 GPUs 优化的源代码或二进制文件。在这篇博文中,我们展示了 CUDA 的目标。

图 1 显示了混合器编译管道。使用 Parallel.For 等并行化模式,或者像在 CUDA 中那样显式地分配并行工作,您可以从加速器的计算马力中获益,而无需了解其内部架构的所有细节。

[EntryPoint]
public static void Run(double[] a, double[] b, int N)
{
    Parallel.For(0, N, i => { a[i] += b[i]; });
}

您可以使用 NVIDIA Nsight Visual Studio Edition 在 GPU 上调试和分析这段代码。Hybridizer 实现高级 C#功能,包括虚拟功能和泛型。

哪里可以获取 Hybridizer

Hybridizer 有 两个版本

在提供自动默认行为的同时, Hybridizer 在每个阶段都提供了完全的开发人员控制,允许您重用现有的特定于设备的代码、现有的外部库或定制的手工代码片段。

调试和分析

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

Figure 2: Debugging C# code running on the GPU with Hybridizer and NVIDIA Nsight Visual Studio Edition.
图 2 :使用 Hybridizer 和 NVIDIA Nsight VisualStudio Edition 调试运行在 GPU 上的 C 代码。

您可以在复杂的项目中集成 Hybridizer ,即使在代码不可用或混淆的库中也是如此,因为 Hybridizer 操作的是 MSIL 字节码。我们在   我们的博客帖子 中展示了这种能力,即在不修改库的情况下,用杂交子加速大型图像处理库。在 MSIL 字节码上操作还支持在. Net 虚拟机上构建的各种语言,例如 VB . Net 和 F 。

所有这些灵活性并不是以性能损失为代价的。正如我们的   benchmark 所说明的,杂交器生成的代码可以执行与手写代码一样好的性能。您可以使用性能分析器,如 NVIDIA Nsight 和 NVIDIA 可视化探查器来测量生成的二进制文件的性能,性能指标引用原始源代码(例如 C )。

一个简单的例子:曼德尔布洛特

作为第一个示例,我们演示了在 NVIDIA GeForce GTX 1080 Ti GPU ( Pascal 体系结构;计算能力 6 . 1 )上运行的 Mandelbrot 分形的渲染。

Mandelbrot C 代码

下面的代码片段显示了纯 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 属性告诉杂交器生成一个 CUDA 内核。多维数组映射到内部类型,而 Parallel2D.For 映射到 2D 执行网格。给出几行样板代码,我们在 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 所示。

Figure 3. Profiling Mandelbrot C# code in the CUDA source view.
图 3 。在 CUDA 源代码视图中分析 Mandelbrot C #代码。

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

至于性能,这个例子达到了峰值计算 FLOP / s 的 72 . 5% ,这是用 CUDA C++ 手写的相同代码的 83% 。

Figure 4: Profiler output showing the GPU utilization and execution efficiency of the Mandelbrot code on the GPU. It achieves nearly as good efficiency as handwritten CUDA C++ code.
图 4 : Profiler 输出显示了 GPU 上 Mandelbrot 代码的利用率和执行效率。它达到了与手写 CUDA C ++代码差不多的良好效率。

使用杂交器提供的扩展控制,可以从 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 ++代码执行一致,达到峰值触发器的 87% ,如图 5 所示。

Figure 5: Profiling the hand-optimized Mandelbrot C# code.
图 5 :分析手动优化的 Mandelbrot C 代码。

泛型与虚函数

Hybridizer 在设备功能上支持 泛型和虚函数调用 。现代编程语言的这些基本概念促进了代码模块化并提高了表达能力。然而, C 型的类型解析是在运行时完成的,这引入了一些性能上的惩罚。席。 NET- 泛型可以在保持灵活性的同时实现更高的性能: FixZER 将泛型映射到 C ++模板,这些模板在编译时被解决,允许函数内联和过程间优化。另一方面,虚拟函数中的方法被映射到另一个虚拟函数中。

模板实例化提示由两个属性 HybridTemplateConceptHybridRegisterTemplate 提供给混合器(在设备代码中触发实际的模板实例化),另一个使用模板映射。基准依赖于一个公开下标运算符的公共接口 IMyArray

[HybridTemplateConcept]
public interface IMyArray {

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

这些操作员必须与设备功能“混合”。为此,我们将 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];
        }
    }
}

因为我们在 ab 上调用下标运算符,所以在 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 所示。

Figure 6. A virtual function call in PTX.
图 6 。 PTX 中的虚函数调用。

这个版本的算法消耗 32 个寄存器,达到 271 GB / s 的带宽,如图 7 所示。在相同的硬件上, CUDA 工具箱中的 bandwidthTest 示例达到 352Gb / s 。

Figure 7. Low achieved bandwidth due to virtual function calls.
图 7 。由于虚拟函数调用,实现的带宽较低。

虚函数表导致更大的寄存器压力,并阻止内联。

一般呼叫

我们用泛型编写了第二个版本,要求杂交子生成模板代码。

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

Figure 8. Using generic parameters generates inline function calls rather than virtual function table lookups.
图 8 。使用泛型参数生成内联函数调用,而不是虚拟函数表查找。

泛型参数的性能要好得多,达到 339gb / s ,性能提高了 25% (图 9 ),比 bandwidthTest 提高了 96% 。

Figure 9. Generics achieve higher bandwidth due to function inlining.
图 9 。由于函数内联,泛型实现了更高的带宽。

开始使用杂交剂

Hybridizer 支持多种 C 特性,允许代码分解和表达能力。 Visual Studio 与   Nsight (调试器和探查器)的集成为您提供了一个安全高效的开发环境。 Hybridizer 即使在非常复杂、高度定制的代码上也能获得出色的 GPU 性能。

您可以从 Visual Studio 市场 下载 杂交者基本要素 。看看 github 上的 我们的 SDK

 

 

Tags