65.9K
CodeProject 正在变化。 阅读更多。
Home

使用 Cudafy 在 .NET 中进行 GPGPU 编程

Nick Kopp
starIconstarIconstarIconstarIcon
emptyStarIcon
starIcon

4.95/5 (59投票s)

2011年 5月 27日

LGPL3

13分钟阅读

viewsIcon

398071

downloadIcon

5434

介绍如何使用 Cudafy.NET 在 GPU 上执行处理

 

引言

本文探讨了如何利用 GPU 从 .NET 进行通用处理。 

重要提示：请访问 CUDAfy.NET Codeplex 网站获取更新。

背景

图形处理单元 (GPU) 越来越多地用于执行非图形工作。世界上最快的超级计算机——天河一号——就使用了大量的 GPU。使用 GPU 的原因是它们提供了大规模并行架构。虽然顶级的 Intel 和 AMD 处理器提供六到八个核心，但 GPU 可以有数百个核心。此外，GPU 具有各种类型的内存，可以实现高效的寻址方案。根据算法，这可以带来巨大的性能提升，并且 100 倍甚至更高的速度提升并不罕见，甚至也并不复杂。但这不仅仅是为超级计算机准备的。即使是普通 PC 也可以利用 GPU。我正在使用一台相当便宜的 Acer 笔记本电脑（800 美元），它配备了 Intel i5 处理器和 NVIDIA GT540M GPU。这台小东西 hardly runs warm，并且可以让我标准的配备两块 NVIDIA GTX 460 的工作站与之媲美。这台工作站令人惊叹的是，在理想条件下，它可以达到 1 teraflops（不到 100GFLOPs 是因为 Intel i7 CPU）。如果您回顾超级计算机的历史，这意味着我拥有相当于 1996 年世界顶级计算机（ASCI Red）的性能。1996 年并不算太久远。到 2026 年，我们的桌子上能否拥有天河一号？

Super Computers 1996 and 2010

以上关键在于，应用程序能否加速取决于算法。并非所有应用程序都能受益，有时需要创造力。这是否需要时间和预算来完成，必须权衡。总之，任何降低利用 PC 或笔记本电脑中的超级计算机的门槛的事物都只会是好事。在通用 GPU (GPGPU) 领域，NVIDIA 的 CUDA 目前是最用户友好的。这是一种 C 的变体。编译需要使用 NVIDIA NVCC 编译器，然后该编译器会使用 Microsoft Visual C++ 编译器。它不是一门难学的语言，但确实会带来一些有趣的问题。应用程序往往首先成为 CUDA 应用程序。要扩展一个“正常”的应用程序以卸载到 GPU，需要一种不同的方法，通常使用 CUDA 驱动程序 API。您可以使用 NVCC 编译器编译模块，并将它们加载到您的应用程序中。

这一切都很好，但它留下了一种相当笨拙的方法。您有两个单独的代码库。如果您不喜欢 Visual Studio 和 .NET，这可能不是什么大问题。从 .NET 使用 CUDA 则另当别论。目前，NVIDIA 直接将 .NET 用户引导至 CUDA.NET，这是一个不错的、尽管有些薄弱的 CUDA API 封装。但是 GPU 代码仍然必须用 NVCC 单独编写和编译，并且 CUDA.NET 的工作似乎已经停止了。CUDA 3.2 引入的最新更改意味着许多功能已损坏（例如，CUDA 3.2 引入了 64 位指针和许多 API 的 v2 版本）。

作为一个忠实的 .NET 开发者，是时候纠正这个问题了，结果就是 Cudafy.NET。Cudafy 是用于描述将 CPU 代码移植到 CUDA GPU 代码的非官方动词。Cudafy.NET 允许您完全在 .NET 应用程序内部对 GPU 进行编程，并且最大限度地减少了杂乱、笨拙的操作。现在开始。

Cudafy Code Snippet

代码

本项目将帮助您设置并运行 Cudafy.NET。一些简单的例程将从标准的 .NET 应用程序在 GPU 上运行。如果您还没有，则需要先执行一些操作。首先，请确保您拥有一块相对较新的 NVIDIA 显卡，并且支持 CUDA。如果您没有，那也并非世界末日，因为 Cudafy 支持 GPGPU 仿真。仿真有助于调试，但根据您尝试并行运行的线程数量，可能会非常缓慢。如果您有一台普通 PC，可以花很少的钱买到一块 NVIDIA PCI Express CUDA GPU。由于应用程序会自动扩展，您的应用程序将在所有 CUDA GPU 上运行，从一些上网本上的最小型 GPU 到高端专用 Tesla 系列。 

然后，您需要访问 NVIDIA CUDA 网站并下载 CUDA 5.5 Toolkit。将其安装在默认位置。接下来，确保您拥有最新的 NVIDIA 驱动程序。这些可以通过 NVIDIA 更新或从 这里 获取。此处的项目是使用 Visual Studio 2010 和 C# 语言构建的，虽然 VB 和其他 .NET 语言应该也可以。Visual Studio Express 可以正常使用，但请记住，它只能创建 32 位应用程序。要使其与 Express 版本一起工作，您需要访问 Visual Studio Express 网站

  1. 下载并安装 Visual C++ 2010 Express（NVIDIA 编译器需要它）
  2. 下载并安装 Visual C# 2010 Express
  3. 下载并安装 CUDA 5.5 Toolkit 
  4. 确保 C++ 编译器 (cl.exe) 在搜索路径中（环境变量）
  5. 可能需要重新启动

NVCC 的这个设置实际上是整个过程中最困难的阶段，所以请坚持下去。仔细阅读您遇到的任何错误——最有可能的是它们与找不到 cl.exe 或没有 CUDA Toolkit 有关。

最后，要正确使用 Cudafy，需要对 CUDA 架构有基本的了解。这一点无法回避。本教程的目标不是提供这方面的内容，因此我将您引荐给 Jason Sanders 和 Edward Kandrot 的《CUDA by Example》。

使用代码 

可下载的代码提供了一个 VS2010 C# 4.0 控制台应用程序，其中包含 Cudafy 库。有关 Cudafy.NET SDK 的更多信息，请访问 网站。该应用程序在 GPU 上执行一些基本操作。向 Cudafy.NET.dll 添加了一个简单的引用。为了翻译成 CUDA C，它依赖于 SharpDevelop 的出色 .NET 反编译器 ILSpy 和 JB Evian 的 Mono.Cecil。库 ICSharpCode.Decompiler.dll、ICSharpCode.NRefactory.dll、IlSpy.dll 和 Mono.Cecil.dll 包含此功能。Cudafy.NET 目前依赖于 ILSpy 1.0.0.822 库的微小修改版本。

有各种命名空间需要处理。在 Progam.cs 中，我们使用以下命名空间

using Cudafy;
using Cudafy.Host;
using Cudafy.Translator;

要指定要在 GPU 上运行的函数，请应用 Cudafy 属性。您可以在 GPU 上运行的最简单的函数由 kernel 方法说明，并且还展示了一个稍微复杂且更有用的方法。

[Cudafy]
public static void kernel()
{
}
	   
[Cudafy]
public static void add(int a, int b, int[] c)
{
	c[0] = a + b;
}

可以使用 CudafyTranslator 从同一应用程序内将这些方法转换为 GPU 代码。这是一个围绕 ILSpy 派生的 CUDA 语言的包装器，它简单地将 .NET 代码转换为 CUDA C，并将此代码与反射信息一起封装到一个 CudafyModule 对象中。CudafyModule 有一个 Compile 方法，该方法包装了 NVIDIA NVCC 编译器。NVCC 编译的输出是 NVIDIA 称之为 PTX 的内容。这是一种 GPU 的中间语言，它允许多代 GPU 与相同的应用程序一起工作。这也被存储在 CudafyModule 中。CudafyModule 也可以序列化和反序列化到/从 XML。这种文件的默认扩展名是 *.cdfy。这种 XML 文件很有用，因为它们可以加快速度并避免将来运行时的不必要编译。CudafyModule 有检查校验和的方法——基本上，我们测试 .NET 代码自创建缓存 XML 文件以来是否已更改。如果没有，那么我们可以放心地假设反序列化的 CudafyModule 实例和 .NET 代码是同步的。

在此示例项目中，我们使用 CudafyTranslator 上的“智能”Cudafy() 方法，该方法会自动进行缓存和类型推断。它执行以下步骤：

CudafyModule km = CudafyModule.TryDeserialize(typeof(Program).Name);
if (km == null || !km.TryVerifyChecksums())
{
	km = CudafyTranslator.Cudafy(typeof(Program));
	km.Serialize();
}

第一行，我们尝试从名为 Program.cdfy 的 XML 文件进行反序列化。如果未显式指定扩展名，则会自动添加。如果文件不存在或因某种原因失败，则返回 null。相比之下，Deserialize 方法在失败时会抛出异常。如果未返回 null，则使用 TryVerifyChecksums() 验证校验和。如果文件是用具有不同校验和的程序集创建的，此方法将返回 false。如果这两项检查都失败，那么我们再次进行 cudafy。这次，我们显式传递要 cudafy 的类型。这里可以指定多个类型。最后，我们将其序列化以供将来使用。

现在我们有了一个有效的模块，可以继续了。要加载模块，我们首先需要获得所需 GPU 的句柄。这是这样完成的。GetDevice() 是重载的，并且可以包含一个设备 ID 来指定多 GPU 系统中的哪个 GPU，它返回一个抽象的 GPGPU 实例。eGPUType 枚举器可以是 Cuda 或 Emulator。在下一行中，加载模块非常明显。CudaGPU 和 EmulatedGPU 继承自 GPGPU。

_gpu = CudafyHost.GetDevice(eGPUType.Cuda);
_gpu.LoadModule(km);

要运行 kernel 方法，我们需要使用 GPU 实例上的 Launch 方法（Launch 是 GPU 启动的另一种说法，在 .NET 中相当于 Invoke）。此方法有许多重载版本。最直接和最干净的使用方式是利用 .NET 4.0 的动态语言运行时 (DLR) 的版本。

_gpu.Launch().kernel();

在这种情况下，Launch 不接受任何参数，这意味着在 GPU 上启动一个线程，它运行 kernel 方法，该方法也不接受任何参数。下面展示了一种替代的非动态启动方式。优点是第一次运行时更快。DLR 可能需要长达 50ms 来完成其魔法。值 1 的两个参数指的是线程数（基本上是 1 * 1），但稍后会详细介绍。

_gpu.Launch(1, 1, "kernel");

还提供了许多其他示例，包括强制性的“Hello, world”（在 GPU 上以 Unicode 编写）。更重要的是“Add vectors”代码，因为处理大型数据集是 GPGPU 的核心业务。我们的 addVector 方法定义如下：

[Cudafy]
public static void addVector(GThread thread, int[] a, int[] b, int[] c)
{
	// Get the id of the thread. addVector is called N times in parallel, so we need 
	// to know which one we are dealing with.
	int tid = thread.blockIdx.x;
	// To prevent reading beyond the end of the array we check that 
	// the id is less than Length
	if (tid < a.Length)
		c[tid] = a[tid] + b[tid];
}

参数 a 和 b 是输入向量，c 是结果向量。GThread 是一个意外的组成部分。由于 GPU 将并行启动许多 addVector 线程，因此我们需要在方法内部能够识别我们正在处理的线程。这可以通过 CUDA 内置变量实现，在 Cudafy 中可以通过 GThread 访问。

如果您在主机上有一个长度为 N 的数组 a，并且想在 GPU 上对其进行处理，那么您需要将数据传输到那里。我们使用 GPU 实例的 CopyToDevice 方法。

int[] a = new int[N];
int[] dev_a = _gpu.CopyToDevice(a);

这里有趣的是 CopyToDevice 的返回值。它看起来像一个整数数组。但是，如果您在调试器中将鼠标悬停在上面，您会发现它的长度为零，而不是 N。返回的是 GPU 上数据的指针。它只在 GPU 代码（您用 Cudafy 属性标记的方法）中有效。GPU 实例存储这些指针。将数据传输到 GPU 固然好，但我们也可能需要在 GPU 上为结果或中间数据分配内存。为此，我们使用 Allocate 方法。以下是为 GPU 分配 N 个整数的代码。

int[] dev_c = _gpu.Allocate(N);

启动 addVector 方法更复杂，需要参数来指定线程数，以及目标方法本身的参数。

_gpu.Launch(N, 1).addVector(dev_a, dev_b, dev_c);

线程分组为 Blocks。Blocks 分组为 Grid。这里我们启动 N 个 Blocks，每个 Block 包含 1 个线程。注意 addVector 包含一个 GThread 参数——无需将其作为参数传递。如前所述，GThread 是 Cudafy 中等效的 CUDA 内置变量，我们用它来标识线程 ID。下图展示了一个 Grid 包含 Blocks 的二维数组的示例，其中每个 Block 包含 Threads 的二维数组。

Grids Blocks Threads

另一个值得注意的有趣点是 GPU 实例上的 FreeAll 方法。GPU 上的内存通常比主机上的内存有限，所以要明智地使用。您需要显式释放内存，但是如果 GPU 实例超出作用域，其析构函数将清理 GPU 内存。

最后一个示例稍微复杂一些，说明了结构体和多维数组的使用。在文件 Struct.cs 中定义了 ComplexFloat。

[Cudafy]
public struct ComplexFloat
{
	public ComplexFloat(float r, float i)
	{
		Real = r;
		Imag = i;
	}
	public float Real;
	public float Imag;
	public ComplexFloat Add(ComplexFloat c)
	{
		return new ComplexFloat(Real + c.Real, Imag + c.Imag);
	}
}

将完整地翻译整个结构体。不需要在成员上放置属性。我们可以在主机和 GPU 代码中自由使用此结构体。在这种情况下，我们在主机上初始化这些元素的 3D 数组，然后传输到 GPU。在 GPU 上，我们执行以下操作：

[Cudafy]
public static void struct3D(GThread thread, ComplexFloat[,,] result)
{
	int x = thread.blockIdx.x;
	int y = thread.blockIdx.y;
	int z = 0;
	while (z < result.GetLength(2)) 
	{
		result[x, y, z] = result[x, y, z].Add(result[x, y, z]);
		z++;
	}
}

这次线程在一个二维 Grid 中启动，因此每个线程由 x 和 y 分量标识（我们启动的线程 Grid 等于数组 x 和 y 维的大小）。每个线程然后处理 z 维中的所有元素。维度的长度通过 .NET 数组的 GetLength 方法获得。计算将每个元素加到自身。

许可证

Cudafy.NET SDK 是双许可的软件库。LGPL 版本适用于开发专有或开源应用程序，前提是您可以遵守 GNU LGPL 版本 2.1 中包含的条款和条件。谢谢。

结束语

希望您已受到启发，进一步探索 GPGPU 编程的世界。大多数 PC 已经拥有了您 PC 中强大的协处理器，它可以补充 CPU 并提供潜在的巨大性能提升。事实上，这些提升有时是如此之大，以至于对于不了解的人来说，它们不再有意义。我知道有些开发者会故意夸大其词地“向下”报告改进…… 问题在于如何利用这种强大的性能。通过 NVIDIA 在 CUDA 方面的努力，这变得越来越容易。Cudafy.NET 进一步发展，使 .NET 开发者也能进入这个世界。

这个简短的文章涵盖的内容远不止这些。请购买一本关于该主题的好书，或查阅 NVIDIA 网站上的教程。背景知识将帮助您更好地理解 Cudafy，并构建更复杂的算法。另外，请访问 Cudafy 网站下载 SDK。SDK 包含两个大型示例项目，其中包含光线追踪、涟漪效果和分形等内容。

历史

  • 2011 年 5 月 27 日
    • 首次提交
  • 2011 年 6 月 16 日
    • 更新以使用 Cudafy.NET V1.2
    • 添加了一些照片
  • 2011 年 7 月 12 日
    • 更新以使用 Cudafy.NET V1.4
  • 2013 年 9 月 16 日
    • 更新以使用 Cudafy.NET V1.26 
© . All rights reserved.