【发布时间】：2013-05-02 01:20:02
【问题描述】：
我一直在使用 C# + Cudafy（C# -> CUDA 或 OpenCL 翻译器）开发的波形模拟器效果很好，除了运行 OpenCL CPU 版本（Intel 驱动程序，15 " MacBook Pro Retina i7 2.7GHz, GeForce 650M (Kepler, 384 cores)) 的速度大约是 GPU 版本的四倍。
（无论我使用 CL 还是 CUDA GPU 后端都会发生这种情况。OpenCL GPU 和 CUDA 版本的性能几乎相同。）
为了澄清，对于一个示例问题：
- OpenCL CPU 1200 赫兹
- OpenCL GPU 320 Hz
- CUDA GPU -~330 Hz
我不知道为什么 CPU 版本会快 比 GPU。在这种情况下，在 CPU 和 GPU 上执行（在 CL 情况下）的内核代码是相同的。我在初始化期间选择了 CPU 或 GPU 设备，但除此之外，一切都是相同的。
编辑
这是启动其中一个内核的 C# 代码。 （其他的非常相似。）
public override void UpdateEz(Source source, float Time, float ca, float cb)
{
var blockSize = new dim3(1);
var gridSize = new dim3(_gpuEz.Field.GetLength(0),_gpuEz.Field.GetLength(1));
Gpu.Launch(gridSize, blockSize)
.CudaUpdateEz(
Time
, ca
, cb
, source.Position.X
, source.Position.Y
, source.Value
, _gpuHx.Field
, _gpuHy.Field
, _gpuEz.Field
);
}
还有，Cudafy 生成的相关 CUDA 核函数如下：
extern "C" __global__ void CudaUpdateEz(float time, float ca, float cb, int sourceX, int sourceY, float sourceValue, float* hx, int hxLen0, int hxLen1, float* hy, int hyLen0, int hyLen1, float* ez, int ezLen0, int ezLen1)
{
int x = blockIdx.x;
int y = blockIdx.y;
if (x > 0 && x < ezLen0 - 1 && y > 0 && y < ezLen1 - 1)
{
ez[(x) * ezLen1 + ( y)] = ca * ez[(x) * ezLen1 + ( y)] + cb * (hy[(x) * hyLen1 + ( y)] - hy[(x - 1) * hyLen1 + ( y)]) - cb * (hx[(x) * hxLen1 + ( y)] - hx[(x) * hxLen1 + ( y - 1)]);
}
if (x == sourceX && y == sourceY)
{
ez[(x) * ezLen1 + ( y)] += sourceValue;
}
}
为了完整起见，这里是用于生成 CUDA 的 C#：
[Cudafy]
public static void CudaUpdateEz(
GThread thread
, float time
, float ca
, float cb
, int sourceX
, int sourceY
, float sourceValue
, float[,] hx
, float[,] hy
, float[,] ez
)
{
var i = thread.blockIdx.x;
var j = thread.blockIdx.y;
if (i > 0 && i < ez.GetLength(0) - 1 && j > 0 && j < ez.GetLength(1) - 1)
ez[i, j] =
ca * ez[i, j]
+
cb * (hy[i, j] - hy[i - 1, j])
-
cb * (hx[i, j] - hx[i, j - 1])
;
if (i == sourceX && j == sourceY)
ez[i, j] += sourceValue;
}
显然，这个内核中的if 很糟糕，但即使是由此产生的管道停顿也不应该导致如此极端的性能增量。
唯一让我感到震惊的另一件事是我使用了一个蹩脚的网格/块分配方案 - 即，网格是要更新的数组的大小，每个块都是一个线程。我确信这对性能有一些影响，但我看不到它导致它是 CPU 上运行的 CL 代码速度的 1/4。啊！
【问题讨论】：
-
你有一些代码示例可以分享吗？
-
@EricBainville 当然——你想要 C#、CUDA 或 CL 内核，还是什么？ （这是一个半中型应用程序。我不想将 20k 行代码粘贴到 SO 中）
-
我没有看到任何迹象表明 cuda 内核每个块使用超过 1 个线程（没有使用
threadIdx.x或threadIdx.y）。此外，启动指定每个块 1 个线程。这意味着大约 97% 的 GPU 功能未被使用。我对 cudafy 了解不多，所以我不知道你是否可以控制它，但我一点也不惊讶 cuda 代码运行速度不快。 -
发布有关性能的问题时，请发布可重现的或使用 gridDim、blockDim 和所有循环执行的次数来注释源。启动 1 个线程的块不太可能允许 CPU 实现对代码进行矢量化。在 NVIDIA GPU 上，您将以远低于 1/32 的计算效率执行，而在 AMD GPU 上，您将以低于 1/64 的计算效率执行。我建议您分析 GPU 代码。
-
@HansRudel NVIDIA GPU 有关 SIMT 架构和执行模型的更多信息，请参阅 CUDA C 编程指南中的 SIMT Architecture。在计算能力为 1.0-3.5 的设备上，WARP_SIZE 为 32 个线程。指定具有 1 个线程的块将导致硬件执行 1 个具有 1 个活动线程和 31 个非活动线程的扭曲，从而导致 3% 的效率。 AMD GPU 在称为波前的 64 个线程组中管理和执行指令。
标签： c# cuda opencl cudafy.net