在我的笔记本电脑上,我有两张显卡 - 英特尔Iris和Nvidia GeForce GT 750M.我正在尝试使用简单的向量添加OpenCL.我知道,Nvidia卡更快,可以更好地完成工作.原则上,我可以if在将NVIDIA在VENDOR属性中查找的代码中放置一个语句.但我想要有一些优雅的东西.以编程方式选择更好(更快)GPU的最佳方法是什么OpenCL C/C++？
我开发了一个实时光线跟踪器(不仅仅是一个光线连接器),它以编程方式选择了两个GPU和一个CPU,并实时渲染和平衡所有三个上的负载.我就是这样做的.
比方说,有三个设备,d1,d2,和d3.分配给每个装置的重量:w1,w2,和w3.调用要渲染的像素数n.假设一个名为的自由参数alpha.
alpha = 0.5.n1=w1*n上的像素d1,接下来 n2=w2*n的像素d2,而最后n3=w3*n像素上d3并记录时间来呈现每个deivce t1,t2和t3.vsum = n1/t1 + n2/t2 + n3/t3.w_i = alpha*w_i + (1-alpha)*n_i/t_i/vsum.值的关键alpha是允许平滑过渡.而不是根据它在某些旧重量中混合的时间重新分配所有重量.没有使用alpha我有不稳定性.alpha可以调整该值.在实践中,它可能设置在1%左右但不是0%.
我们来选一个例子吧.
我有一个GTX 590,这是一个带有两个欠频GTX580的双GPU卡.我还有一台Sandy Bridge 2600K处理器.GPU比CPU快得多.我们假设它们的速度提高了大约10倍.我们还说有900像素.
使用GPU1渲染前300个像素,使用GPU2渲染下300个像素,使用CPU1渲染最后300个像素并10 s, 10 s, and 100 s分别记录时间.因此,整个图像的一个GPU需要30秒,而单独的CPU需要300秒.GPUS合二为一15 s.
计算vsum = 30 + 30 + 3 = 63.再次重新计算重量:
w1,w2 = 0.5*(1/3) + 0.5*300/10/63 = 0.4和w3 = 0.5*(1/3) + 0.5*300/100/63 = 0.2.
渲染下一帧:GPU1为360像素,GPU2为360像素,CPU1为180像素,时间变得更加平衡11 s, 11 s, and 55 s.
在多个帧之后,该(1-alpha)术语占主导地位,直到最终权重都基于该术语.在这种情况下,权重分别变为47%(427像素),47%,6%(46像素),并且时间14 s, 14 s, 14 s分别变为.在这种情况下,CPU仅将GPU的使用时间提高一秒.
我在这个计算中假设了一个均匀的负荷.在实际光线跟踪器中,负载随扫描线和像素而变化,但算法在确定权重时保持不变.
在实践中,一旦发现重量,它们不会发生太大变化,除非场景的负荷发生显着变化,例如,如果场景的一个区域具有高折射和反射,而其余区域是漫反射的,但即使在这种情况下我也限制树深度,所以这没有戏剧性的效果.
通过循环将此方法扩展到多个设备很容易.我曾在四台设备上测试了我的光线跟踪器.两个12核Xeon CPU和两个GPU.在这种情况下,CPU具有更大的影响力,但GPU仍占主导地位.
如果有人想知道.我为每个设备创建了一个上下文,并在一个单独的线程中使用每个上下文(使用pthreads).对于三个设备,我使用了三个线程.
实际上,您可以使用它在不同供应商的同一设备上运行.例如,我在2600K上同时使用了AMD和Intel CPU驱动程序(每个驱动程序产生大约一半的帧),以查看哪个供应商更好.当我第一次这样做(2012年)时,如果我没记错的话,AMD在英特尔CPU上讽刺地击败了英特尔.
如果有人对我如何提出权重公式感兴趣,我会使用物理学的一个想法(我的背景是物理学而不是编程).
速度(v)=距离/时间.在这种情况下,distance(d)是要处理的像素数.那么总距离是
d = v1*t1 + v2*t2 + v3*t3
我们希望他们每个人都在同一时间完成
d = (v1 + v2 + v3)*t
然后得到重量定义
v_i*t = w_i*d
这使
w_i = v_i*t/d
而(t/d)中的replacement(d = (v1 + v2 + v3)*t)给出:
w_i = v_i /(v1 + v2 + v3)
很容易看出这可以推广到任意数量的设备 k
w_i = v_i/(v1 + v2 + ...v_k)
所以vsum我的算法代表"速度之和".最后,因为v_i随着时间的推移它是n_i/t_i最终给出的像素
w_i = n_i/t_i/(n1/t1 + n2/t2 + ...n_k/t_k)
这是我计算权重的公式中的第二项.
如果它只是一个向量加法并且您的应用程序驻留在主机端，那么 cpu 将获胜。或者甚至更好，集成CPU会快得多。整体性能取决于算法、opencl 缓冲区类型（use_host_ptr、read_write 等）以及计算与数据比率。即使您不复制而是固定阵列并访问，cpu 的延迟也会小于 pci-e 延迟。
如果您要使用 opengl + opencl 互操作，那么您需要知道您的计算设备是否与渲染输出设备相同。（如果你的屏幕从 igpu 获取数据，那么它是 iris，如果不是，那么它是 nvidia）
如果您只需要在 C++ 数组（主机端）上执行一些操作并以最快的方式获得结果，那么我建议您“负载平衡”。
使用 Iris pro 和两个 gt750m（其中一个超频 10%）的 Core i7-5775C 上的 4k 元素矢量相加示例
首先，为所有设备提供相同数量的 ndrange 范围。在每个计算阶段结束时，检查计时。
CPU iGPU dGPU-1 dGPU-2 oc
Intel Intel Nvidia Nvidia
1024 1024 1024 1024
34 ms 5ms 10ms 9ms
然后计算计算带宽的加权（取决于最后的 ndrange 范围）但宽松（不精确但接近）近似值并相应地更改 ndrange 范围：
Intel Intel Nvidia Nvidia
512 1536 1024 1024
16 ms 8ms 10ms 9ms
然后继续计算，直到真正稳定为止。
Intel Intel Nvidia Nvidia
256 1792 1024 1024
9ms 10ms 10ms 9ms
或者直到您可以启用更细的颗粒。
Intel Intel Nvidia Nvidia
320 1728 1024 1024
10ms 10ms 10ms 9ms
Intel Intel Nvidia Nvidia
320 1728 960 1088
10ms 10ms 10ms 10ms
^ ^
| |
| PCI-E bandwidth not more than 16 GB/s per device
closer to RAM, better bandwidth (20-40 GB/s) and less kernel overhead
您可以获取最后 10 个结果的平均值（或 PID），以消除误导平衡的尖峰，而不是仅获取最新的迭代来进行平衡。此外，缓冲区副本可能比计算花费更多时间，如果将其纳入平衡中，则可以关闭不必要/无益的设备。
如果您创建一个库，那么您就不必为您的每个新项目尝试基准测试。当您加速矩阵乘法、流体运动、sql 表连接和财务近似时，它们将在设备之间自动平衡。
对于平衡的解决方案：
如果您可以将线性系统求解为 n 个未知数（每个设备的负载）和 n 个方程（所有设备的基准结果），则可以一步找到目标负载。如果选择迭代，则需要更多步骤直到收敛。后者并不比编写基准更难。前者对我来说更难，但随着时间的推移，它应该会更有效率。
尽管仅向量相加内核不是现实世界的场景，但以下是我的系统的真实基准：
__kernel void bench(__global float * a, __global float *b, __global float *c)
{
int i=get_global_id(0);
c[i]=a[i]+b[i];
}
2560 768 768
AMD FX(tm)-8150 Eight-Core Processor Oland Pitcairn
这是经过几次迭代之后的结果（即使有额外的缓冲区副本，不使用任何主机指针，fx 也会更快）。甚至 oland GPU 也赶上了皮特凯恩，因为它们的 PCI-E 带宽是相同的。
现在有一些三角函数：
__kernel void bench(__global float * a, __global float *b, __global float *c)
{
int i=get_global_id(0);
c[i]=sin(a[i])+cos(b[i])+sin(cos((float)i));
}
1792 1024 1280
测试 gddr3-128 位与 gddr5-256 位（超频）和缓存。
__kernel void bench(__global float * a, __global float *b, __global float *c)
{
int i=get_global_id(0);
c[i]=a[i]+b[i]-a[i]-b[i]+a[i]+b[i]-b[i]-a[i]+b[i];
for(int j=0;j<12000;j++)
c[i]+=a[i]+b[i]-a[i]-b[i]+a[i]+b[i]-b[i]-a[i]+b[i];
}
256 256 3584
高计算数据比：
__kernel void bench(__global float * a, __global float *b, __global float *c)
{
int i=get_global_id(0);
c[i]=0.0f; float c0=c[i];float a0=a[i];float b0=b[i];
for(int j=0;j<12000;j++)
c0+=sin(a0)+cos(b0*a0)+cos(sin(b0)*19.95f);
c[i]=c0;
}
256 2048 1792
现在，Oland GPU 再次值得一提，即使只有 320 个核心，它还是获胜了。因为 4k 元素很容易缠绕所有 320 个核心超过 10 次，但皮特凯恩 GPU（1280 个核心）没有完全充满折叠阵列（波前），这导致执行单元占用率较低 ---> 无法隐藏延迟。我认为低负载的低端设备更好。也许当 directx-12 推出一些负载平衡器时我可以使用它，并且这个 Oland 可以计算游戏爆炸中 5000 - 10000 个粒子的物理，而皮特凯恩可以计算烟雾密度。