AMD GPU 和 Intel CPU 上的 CUDA 编程模型
本文建立在之前的高性能查询：GPU vs PLINQ vs LINQ 的基础上，并将其移植以支持 OpenCL 设备，同时增加了基准测试，方便您轻松比较性能。
引言
CUDAfy.NET 是一个功能强大的开源库，用于在 Microsoft .NET 中为 NVIDIA 图形处理单元 (GPU) 进行非图形编程。最近增加了对 AMD GPU 和 x86 CPU 的支持。这使得 CUDA 编程模型能够应用于 NVIDIA 以外的设备。您可以在以下 CodeProject 文章中更深入地了解 CUDAfy：
为了不过多重复之前文章的内容，让我们简要地问一下，为什么我们要考虑将 GPU 用于非图形任务？原因是，对于正确的任务，GPU 拥有巨大的计算能力，这些能力在游戏、视频处理和 CAD 套件之外相对未被充分利用。最初，在地震学、金融、物理学、生物学等特定市场，我们看到 GPU 的使用越来越多，以加速计算密集型任务。然而，由于 GPU 的编程方法与传统编程不同，收益必须足够大，才能使付出努力是值得的。NVIDIA 的 CUDA 平台极大地简化了开发。认为它仍然太复杂是不对的，因为大规模并行编程确实需要一种不同的看待事物的方式。如何将一个算法分割成 1000 多个核心，必须也应该不同。不要被 Intel 对其 Xeon Phi 的宣传所欺骗。仅仅因为每个核心大致上是 x86，并不意味着您可以直接运行您的应用程序并获得改进。另外，也要警惕过快地加入 OpenCL 阵营。是的，理论上它允许针对 AMD GPU、Intel CPU、NVIDIA GPU 甚至 FPGA，但任何比较过 CUDA 运行时与 OpenCL 的人都会很快发现 OpenCL 有多么复杂和冗长。
CUDA 编程模型仍然是处理大规模并行架构的最佳方式。CUDAfy 等工具通过以清晰简单的方式暴露编程接口，降低了 .NET 开发者的门槛，甚至比 CUDA 本身更清晰。CUDA 仍然存在一个问题——它只支持 NVIDIA GPU。直到现在。CUDAfy.NET 已更新，允许在 AMD GPU 和 x86 CPU 上使用 CUDA 模型。使用 Bittware 的一些 FPGA 板也应该成为可能。这是如何实现的？在底层，CUDAfy 可以利用 OpenCL。这需要安装免费的 NVIDIA、AMD 和/或 Intel OpenCL SDK。您这些具有 CUDA 经验的人无需学习太多新知识即可针对所有这些其他处理器。
本文建立在之前的 高性能查询：GPU vs PLINQ vs LINQ 的基础上，并将其移植以支持 OpenCL 设备，同时增加了基准测试，方便您轻松比较性能。准备好进行 LINQ vs PLINQ vs CUDA vs AMD GPU vs Intel CPU 的巅峰对决了吗？
背景
在本文中，我们将考虑一个 GPS 轨迹点数据库。这是一个 GPS 记录器记录的所有路线的列表。这种功能通常用于公司的卡车和汽车，以监控车辆的去向——税务机关越来越多地要求提供这些信息，以证明私人和工作里程。在我们的示例中，我们希望查询 GPS 点数据库，以返回在给定开始和结束日期内，距离多个目标点一定半径内的所有点。通过计算地球表面两点之间距离的蛮力方法是一个严峻的挑战，并且应该很好地展示使用 LINQ、PLINQ CUDA 和 OpenCL 进行此操作的优缺点。
必备组件
我们将使用 CUDAfy.NET 来编程 GPU。CUDAfy 库包含在下载文件中，但如果您想使用 CUDA 或 OpenCL，您需要下载并安装
- 对于 NVIDIA GPU，您可以使用 CUDA 或 OpenCL。 在此处获取 CUDA SDK。
- 对于 AMD GPU，您可以使用 OpenCL。 在此处获取 AMD OpenCL SDK。
- 要在 Intel CPU 上运行 CUDAfy，请 下载 Intel OpenCL SDK。
具有讽刺意味的是，让 CUDA 工作是最棘手的选项，因为它还需要 Visual Studio。没有 Visual Studio，您可以通过使用 CUDA SDK 中的 OpenCL 来定位 NVIDIA GPU。有关如何开始和运行的更多信息，请参阅我之前的文章。
请参阅本文关于 GPU 上的 base64 编码，或这篇关于 CUDAfy 和 CUDA 的基本介绍。
Using the Code
此应用程序会生成用户定义的点数的随机 GPS 路线，然后允许用户输入搜索条件，包括
- 目标数量
- 日期和时间范围
- 目标半径（米）
- 然后可以使用 LINQ、PLINQ 或任何已发现的、支持 CUDA 和 OpenCL 的设备来执行相同的查询。
代码的基础是 GPS 轨迹点和 GPS 轨迹查询。这些被定义为 TrackPoint 和 TrackQuery。如果主机和 GPU 代码可以在可能的情况下共享，那将非常方便。
在 GPU 上运行的代码称为设备代码，需要对其进行标记，以便 CUDAfy 在转换为 CUDA C 代码时能够识别它。为此，请添加 Cudafy 属性。
[Cudafy]
[StructLayout(LayoutKind.Sequential)]
public struct TrackPoint
{
public TrackPoint(float lon, float lat, long timeStamp)
{
Longitude = lon;
Latitude = lat;
TimeStamp = timeStamp;
}
...
}
在 Cudafy 一个 struct 时，默认情况下所有成员都会被翻译。要忽略某个成员，请使用 CudafyIgnore 属性。在此示例中，我们有一个属性 Time。这是一个 DateTime 类型。GPU 不支持属性和 DateTime，因此我们需要将其标记为被忽略。
[CudafyIgnore]
public DateTime Time
{
get { return new DateTime(TimeStamp); }
set { TimeStamp = value.Ticks; }
}
现在与 CUDA 相比，OpenCL 相当原始和基础。构造函数和方法不能放在 struct 中。因此，我们需要更改上一篇文章中使用的代码。这不过是复制粘贴和少量编辑，但这是 OpenCL 的一个限制，CUDAfy 目前无法为您规避。这也回答了为什么继续支持 CUDA 的问题，因为 NVIDIA GPU 也支持 OpenCL——因为 CUDA 可以做得更多。
现在，敏锐的读者会注意到的另一个变化是，这里我们使用的是 float 而不是 double。这是因为与 CUDA 不同，如果设备不支持 double，代码将无法编译。它不会静默回退到单精度。由于 double 仅在高端 AMD GPU 上受支持，我们将暂时不使用它。
用于测试 GPS 点是否在目标点半径米之内的该方法定义为
public double DistanceTo(TrackPoint A, TrackPoint B)
{
float dDistance = Single.MinValue;
float dLat1InRad = A.Latitude * CONSTS.PI2;
float dLong1InRad = A.Longitude * CONSTS.PI2;
float dLat2InRad = B.Latitude * CONSTS.PI2;
float dLong2InRad = B.Longitude * CONSTS.PI2;
float dLongitude = dLong2InRad - dLong1InRad;
float dLatitude = dLat2InRad - dLat1InRad;
// Intermediate result a.
float a = GMath.Pow(Math.Sin(dLatitude / 2.0F), 2.0F) +
GMath.Cos(dLat1InRad) * GMath.Cos(dLat2InRad) *
GMath.Pow(GMath.Sin(dLongitude / 2.0F), 2.0F);
// Intermediate result c (great circle distance in Radians).
float c = 2.0F * GMath.Atan2(GMath.Sqrt(a), GMath.Sqrt(1.0F - a));
// Distance
dDistance = CONSTS.EARTHRADIUS * c;
return dDistance * 1000.0F;
}
地球表面上任意两点之间的最短距离由大圆距离给出。如果您查看机上杂志，您会看到标准 2D 地图显示您从欧洲飞往加利福尼亚的航班会经过格陵兰岛。这比出发点或到达点都远得多，因为它遵循大圆航线。计算距离需要大量的浮点数学处理。
在 CPU 上运行查询的代码如下所示。我们将目标数组以及其他查询参数传递进去。然后，使用 _track 数组的 LINQ 扩展的 Where 方法，我们返回所有 GetNearestTargetIndex 方法返回小于 255 的点。如果值为 255，则表示该点不在任何目标的半径内，否则返回离该点最近的目标的索引。标准 LINQ 将在单个处理器核心上执行。要利用更多核心，只需在 Where 方法之前插入 AsParallel 即可。
private IEnumerable<TrackPoint> GetPoints(TrackPoint[] targets,
float radius, DateTime startTime, DateTime endTime, bool parallel)
{
long targetStartTime = startTime.Ticks;
long targetEndTime = endTime.Ticks;
// Running the query in parallel is as easy as adding AsParallel().
// We're not concerned with the ordering, so it's ideal.
if(parallel)
return _track.AsParallel().Where(tp => GetNearestTargetIndex(tp,
targets, radius, targetStartTime, targetEndTime) < 255);
else
return _track.Where(tp => GetNearestTargetIndex(tp, targets,
radius, targetStartTime, targetEndTime) < 255);
}
private byte GetNearestTargetIndex(TrackPoint tp, TrackPoint[] targets,
float radius, long targetStartTime, long targetEndTime)
{
float minDist = Single.MaxValue;
byte index = 255;
// If we're not within the time range then no need to look further.
if (tp.TimeStamp >= targetStartTime && tp.TimeStamp <= targetEndTime)
{
int noTargets = targets.Length;
// Go through the targets and get the index of the closest one.
for (int i = 0; i < noTargets; i++)
{
TrackPoint target = targets[i];
float d = TrackQuery.DistanceTo(tp, target);
if (d <= radius && d < minDist)
{
minDist = d;
index = (byte)i;
}
}
}
return index;
}
现在让我们看一下将与 GPU 通信的代码。这在 TrackQuery 类中实现。构造函数负责 Cudafy 或从先前序列化的模块加载。对于 CUDA，这是一个生成 CUDA C 代码、编译该代码并将输出 (PTX) 存储到模块的过程。可以序列化此模块，以节省未来运行的时间或在没有 NVIDIA CUDA SDK 安装的用户系统上运行。对于 OpenCL，大致相同——我们生成 OpenCL 代码并将此代码存储到模块中，但我们不编译。编译在加载模块时进行。在这里，我们希望 Cudafy 两个类型：TrackQuery 和 TrackPoint。
public class TrackQuery
{
public TrackQuery(GPGPU gpu)
{
var props = gpu.GetDeviceProperties();
var name = props.PlatformName + props.Name;
// Does an already serialized valid cudafy module xml file exist. If so
// no need to re-generate it.
var mod = CudafyModule.TryDeserialize(name);
// The gpu can be either a CudaGPU or an OpenCLDevice.
// Realize that an NVIDIA GPU can be both!
// And an Intel CPU can show up as both an AMD and an
// Intel OpenCL device if both OpenCL SDKs are
// installed.
CudafyTranslator.Language = (gpu is CudaGPU) ? eLanguage.Cuda : eLanguage.OpenCL;
if (mod == null || !mod.TryVerifyChecksums())
{
// Convert the .NET code within these two types
// into either CUDA C or OpenCL C. If CUDA C then
// also compile into PTX.
mod = CudafyTranslator.Cudafy(typeof(TrackPoint), typeof(TrackQuery));
// Store to file for re-use
mod.Serialize(name);
}
try
{
// Load the module on to the device. If OpenCL then compile the source.
gpu.LoadModule(mod);
}
catch (Exception ex)
{
System.Diagnostics.Debug.WriteLine(mod.CompilerOutput);
throw;
}
_gpu = gpu;
}
为了有效地利用 GPU 进行通用编程，关键在于最大限度地减少需要在主机和 GPU 设备之间传输的数据量。在这里，我们首先加载轨迹。然后我们可以运行多个查询，而无需重新加载轨迹。由于默认情况下我们将处理 10,000,000 个 GPS 点，代表 160,000,000 字节（每个 GPS 点有两个 System.Single 和一个 System.Int64 字段：经度、纬度和时间戳），这一点很重要。轨迹以点数组的形式上传。为了清晰起见，位于 GPU 上的变量后缀为 _dev。作为一种优化，如果 GPU 上已分配的内存大于要上传的新轨迹，我们不会释放并重新分配，而只会使用所需量（_currentLength）。变量 _indexes 是一个长度等于轨迹长度的字节数组。对于每个对应的点，我们将填写 255（如果点不在任何目标的半径和日期范围内），或者填写目标索引（如果点在范围内）。
public void LoadTrack(TrackPoint[] trackPoints)
{
_currentLength = trackPoints.Length;
// If the current number of points on the GPU are less than the number of points
// we want to load then we resize the allocated memory on the GPU. We simply free
// the existing memory and allocate new memory. We need arrays on the GPU to hold
// the track points and the selected indexes. We make an array on the host to hold
// the returned indexes.
if (_trackPoints.Length < trackPoints.Length)
{
if (_trackPoints_dev != null)
{
_gpu.Free(_trackPoints_dev);
_gpu.Free(_indexes_dev);
}
_trackPoints_dev = _gpu.Allocate(trackPoints);
_indexes_dev = _gpu.Allocate<byte>(trackPoints.Length);
_indexes = new byte[trackPoints.Length];
}
_trackPoints = trackPoints;
// Copy the GPS points to the GPU.
_gpu.CopyToDevice(trackPoints, 0, _trackPoints_dev, 0, trackPoints.Length);
}
我们将从应用程序调用的方法是 SelectPoints。GPU 有多种内存形式。CPU 仅系统也有，因为 CPU 有自己的缓存，但我们从不显式使用它。然而，在 GPU 上，建议使用不同类型的内存，因为这会极大地影响性能。OpenCL 和 CUDA 在这方面非常相似，因此为 CUDAfy 添加 OpenCL 支持相对简单。最大且最慢的内存是全局内存。我们就是在那里存储的轨迹。现代 GPU 通常至少有 1GB 的这种内存。另一种内存是常量内存。它只能由主机写入。对于 GPU 来说，它是只读的，并且非常快。我们将使用它来存储我们的目标。请记住，常量内存的量非常有限（考虑 KB 而不是 MB），并且不能被释放。OpenCL 处理常量内存的方式不同，但 CUDAfy 会为您屏蔽这一点，并且可以使用标准的 CUDA 实践。据说设备函数会被启动。GPU 设备的 Launch 方法接受以下参数：
- 网格中的块数
- 每个块中的线程数（最多 1024 个）
- 要启动的函数名
- 设备上的轨迹点数组
- 轨迹点的数量
- 目标半径（米）
- 开始时间（滴答）
- 结束时间（滴答）
- 结果的字节数组
- 目标数量
启动本质上会并行启动 blocksPerGrid*ciTHREADSPERBLOCK 个线程。这个数字可能相当高，但如果 GPU 函数花费的时间超过 0.5 秒左右，那么建议将过程分成多个调用。NVIDIA 驱动程序不喜欢 GPU 被阻塞太久，并且启动会在某个点超时。OpenCL 指定块和线程的方式与 CUDA 不同，但同样，这对 CUDAfy 开发者来说是透明的。完成后，我们将 _indexes_dev 数组复制回主机，然后搜索小于 255 的值，返回相应的 TrackPoint。我们可以选择对索引做些什么，以便确定哪个目标是最接近的。这可以使用 TrackPointResult struct 来完成，该结构包含在源代码中作为示例。
public IEnumerable<TrackPoint> SelectPoints(TrackPoint[] targets,
float radius, DateTime startTime,
DateTime endTime, bool naive = false)
{
int blocksPerGrid;
// Validate the parameters and calculate how
// many blocks of threads we will need on the GPU.
// Each block of threads will execute ciTHREADSPERBLOCK threads.
Initialize(radius, startTime, endTime, out blocksPerGrid);
// Copy the targets to constant memory.
_gpu.CopyToConstantMemory(targets, 0, _targets, 0, targets.Length);
// Launch blocksPerGrid*ciTHREADSPERBLOCK threads in parallel.
// Each thread will test one GPS point against all targets.
_gpu.Launch(blocksPerGrid, ciTHREADSPERBLOCK, "SelectPointsKernel",
_trackPoints_dev, _currentLength, radius,
startTime.Ticks, endTime.Ticks, _indexes_dev, targets.Length);
// Copy the indexes array back from the GPU to the host
// and search for all points that have an index < 255.
// These correspond to GPS points lying within the search criteria.
_gpu.CopyFromDevice(_indexes_dev, 0, _indexes, 0, _currentLength);
for (int i = 0; i < _currentLength; i++)
{
byte index = _indexes[i];
if (index < 255)
yield return _trackPoints[i];
}
}
每个线程执行以下设备代码，该代码对一个点执行测试。GThread 标识线程的唯一 ID，以便线程知道要访问哪个点。OpenCL 使用一种相当不同的方法来访问 ID，这种方法可能比 CUDA 更简洁。使用 CUDAfy，您可以两全其美，因为 OpenCL 或 CUDA 的方法都可以用于 OpenCL 和 CUDA 目标。在这个例子中，我们将利用另一种形式的 GPU 内存：共享内存。这种内存是在一个块内的所有线程之间共享的。OpenCL 和 CUDA 在这里也很相似。当对全局内存的使用效率低下时，例如，以非顺序方式读取少量数据，性能会下降。在这种情况下，最好使用共享内存作为中间存储。尽管如此，新的 Fermi 架构似乎能很好地处理这段代码，并且更简单的实现 SelectPointsKernelNaive 的运行速度与下面的共享内存实现一样快。试试看。
[Cudafy]
public static void SelectPointsKernel(GThread thread, TrackPoint[] track,
int noPoints, float radius, long startTime, long endTime,
byte[] indexes, int noTargets)
{
// Here we use another form of GPU memory called shared memory.
// This is shared between threads of a single block.
// The size of the shared memory must be constant
// and here we set it to the number of threads per block.
// This can be more efficient than the naive implementation below,
// however on the latest fermi architecture there is little
// or no difference.
byte[] cache = thread.AllocateShared<byte>("cache", ciTHREADSPERBLOCK);
// Get the unique index of the thread: size of the
// block * block index + thread index.
int tid = thread.blockDim.x * thread.blockIdx.x + thread.threadIdx.x;
// Check we are not beyond the end of the points.
if (tid < noPoints)
{
TrackPoint tp = track[tid];
float minDist = Single.MaxValue;
byte index = 255;
if (tp.TimeStamp >= startTime && tp.TimeStamp <= endTime)
{
for (int i = 0; i < noTargets; i++)
{
// get the target from constant memory
TrackPoint target = _targets[i];
// Calculate distance to target and if less
// than radius and current nearest target
// set minDist and index.
float d = DistanceTo(tp, target);
if (d <= radius && d < minDist)
{
minDist = d;
index = (byte)i;
}
}
}
// set the index.
cache[thread.threadIdx.x] = index;
// Synchronize all threads in block
thread.SyncThreads();
// Write the results into global memory array
indexes[tid] = (byte)cache[thread.threadIdx.x];
}
}
基准测试
时间单位是毫秒，越低越好。结果非常令人印象深刻。我们看到，使用 PLINQ 比 LINQ 带来了非常显著的优势，但对于非常苛刻的应用，GPU 的优势尤为突出，在此示例中性能提升高达 30 倍。即使轨迹需要重新上传到 GPU（笔记本电脑上的成本为 150 毫秒），优势仍然相当可观。这非常重要，意味着在更常规的业务应用程序中使用 GPU 应该被考虑。这里的代码可以轻松地适应其他用途。当集成到（Web）服务器应用程序中时，负载可以大大减轻，从而在设备和能源方面带来可观的成本节省。
在这种情况下，同一 NVIDIA 设备上的 CUDA 和 OpenCL 之间的差异很小。但在某些算法中，差异更为显著。有趣的是，在第一代 Intel Core i5 CPU 上运行的 OpenCL 比 PLinq 的性能大约好两倍。对于 socket 1155 Ivy Bridge，我们可以期待更好的结果，因为 Intel OpenCL SDK 可以利用集成 GPU。随着我们转向 Haswell，Intel OpenCL 对于低功耗、小尺寸解决方案来说可能非常有前景。
上述基准测试是在一台强大的 Intel i7 980X 机器上运行的。我省略了 LINQ 测试，因为它会使其他时间几乎难以辨认。板载有一个 NVIDIA GTX660Ti 和一个 AMD 6570。显然，这里的性能比笔记本电脑要快一些。最有趣的结果实际上是，在 Intel i7 CPU 上运行时，AMD OpenCL SDK（黄色）远胜于 Intel OpenCL SDK（灰色）！当然，与 GPU 相比，两者都仍然处于劣势。即使是相对便宜的 AMD 6570 也表现不错。
OpenCL 能够针对 GPU、CPU 甚至 FPGA 的灵活性很有吸引力，但即便如此，如果您知道自己将使用 NVIDIA 设备，仍然有许多令人信服的理由坚持使用 CUDA。CUDA 仍然是一种更愉快的开发体验，其语言和工具更加完善，并且有大量与 CUDA 可互操作的库可用，这足以成为决定性因素。将 FFT、BLAS 和随机数生成视为 CUDA 产品仅仅是开始。许多这些库已通过 CUDAfy 访问，但仅限于 CUDA 目标。
许可证
本项目使用了 Code Project 文章 如何显示启动画面同时进行应用程序初始化 中提供的启动画面代码。
Cudafy.NET SDK 包含两个大型示例项目，其中包括光线追踪、涟漪效果和分形等。许多示例都完全支持 CUDA 和 OpenCL。它是一个双许可软件库。LGPL 版本适用于专有或开源应用程序的开发，前提是您能够遵守 GNU LGPL 版本 2.1 中包含的条款和条件。请访问 Cudafy 网站 了解更多信息。如果使用 LGPL 版本，我们恳请您考虑向 Harmony through Education 捐款。这个小型慈善机构正在帮助发展中国家的残疾儿童。请访问 Cudafy 网站的慈善页面 阅读更多信息。
历史
- 2013 年 4 月 4 日：首次发布
- 2013 年 9 月 17 日：更新至 CUDAfy V1.26