[HPC#] Hybridizer：GPU 上的高性能 C#

Hybridizer 有两个版本：

• Hybridizer Software Suite（Hybridizer 软件套件）：启用 CUDA、AVX、AVX2、AVX512 目标并输出源代码。 该源代码可以进行审查，这在某些业务（例如投资银行）中是强制性的。 Hybridizer 软件套件可根据要求按客户许可。

• Hybridizer Essentials：仅启用 CUDA 目标并仅输出二进制文件。 Hybridizer Essentials 是一款免费的 Visual Studio 扩展，没有硬件限制。 您可以在 GitHub 上找到一组基本代码示例和教育材料。 这些样本还可以作为重现我们的性能结果的一种方式。

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

调试和分析

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

图 2：使用 Hybridizer 和 NVIDIA Nsight Visual Studio 版本调试在 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# 代码

  以下代码片段显示了纯 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 执行网格。 给定几行样板代码，我们在 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%。 这是 83% 的相同代码，用 CUDA 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 所示。

图 5：分析手动优化的 Mandelbrot C# 代码。

泛型和虚函数

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

  模板实例化提示通过两个属性 HybridTemplateConcept 和 HybridRegisterTemplate （触发设备代码中的实际模板实例化）提供给 Hybridizer。 作为示例，让我们看一下两个版本的简单流基准测试，一个使用虚拟函数调用，另一个使用模板映射。 该基准测试依赖于公开下标运算符的通用接口 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];
        }
    }
}

由于我们在被视为接口的 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 个寄存器并实现 271 GB/s 的带宽，如图 7 所示。 在相同的硬件上，CUDA 工具包中的带宽测试示例达到 352 GB/s。

虚函数表会导致更多的寄存器压力，并阻止内联。

泛型调用

我们用泛型编写了第二个版本，要求 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. 使用泛型参数生成内联函数调用，而不是虚拟函数表查找。

通用参数的性能要好得多，达到了 339 GB/s，性能提升了 25%（图 9），带宽测试达到了 96%。

开始使用Hybridizer

  Hybridizer 支持多种 C# 功能，允许代码分解和表达能力。 Visual Studio 和 Nsight（调试器和分析器）内的集成为您提供了一个安全且高效的开发环境。 即使在非常复杂、高度定制的代码上，Hybridizer 也能实现出色的 GPU 性能。

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