1

我正在执行如下所示的基准测试

CHECK( context = clCreateContext(props, 1, &device, NULL, NULL, &_err); );
CHECK( queue = clCreateCommandQueue(context, device, 0, &_err); );
#define SYNC() clFinish(queue)
#define LAUNCH(glob, loc, kernel) OCL(clEnqueueNDRangeKernel(queue, kernel, 2,\
                                                             NULL, glob, loc,\
                                                             0, NULL, NULL))

/* Build program, set arguments over here */


START;
for (int i = 0; i < iter; i++) {
    LAUNCH(global, local, plus_kernel);
}
SYNC();
STOP;
printf("Time taken (plus) : %lf\n", uSec / iter);

START;
for (int i = 0; i < iter; i++) {
    LAUNCH(global, local, minus_kernel);
}
SYNC();
STOP;
printf("Time taken (minus): %lf\n", uSec / iter);

START;
for (int i = 0; i < iter; i++) {
    LAUNCH(global, local, plus_kernel);
    LAUNCH(global, local, minus_kernel);
}
SYNC();
STOP;
printf("Time taken (both) : %lf\n", uSec / iter);

结果看起来很奇怪：

Time taken (plus) : 31.450000
Time taken (minus): 28.120000
Time taken (both) : 2256.380000

START, 并且STOP只是启动和停止计时器的宏。这是相关的宏。

我不确定为什么排队是内核减慢它们的速度（并且仅在 AMD GPU 上）！

编辑我正在使用 Radeon 7970

编辑两个内核都在独立内存上运行。这里还有系统信息。

操作系统：Ubuntu 11.10

fglrx信息：

display: :0  screen: 0
OpenGL vendor string: Advanced Micro Devices, Inc.
OpenGL renderer string: AMD Radeon HD 7900 Series 
OpenGL version string: 4.2.11762 Compatibility Profile Context
4

1 回答 1

1

我认为答案与在较新的 GPU 上缓存数据有关（特别是 Radeon 7970，它使用了下一代图形计算 (GCN)架构。

这种架构的优点之一是它的缓存功能（此时有点接近 CPU 缓存）。如果您执行这样的调用：

PLUS
PLUS 
PLUS
....

然后是驻留在 GPU 内部缓存中的内存。另一方面，如果您拨打这样的电话：

PLUS
MINUS
PLUS 
MINUS
...

如果两个内核具有与之关联的不同内存对象，则数据会从每个 CU 上的硬件设备中踢出，导致需要从非常缓慢的全局内存中引入它们。

两种简单的方法来测试是否是这种情况：

  1. 仅运行具有不同迭代次数的 Plus。随着迭代次数的增加，平均时间将减少，因为第一次运行（将数据带入）的成本已摊销。此外，您应该注意到第一次之后的所有调用应该是相对相等的。

  2. 使 Plus 和 Minus 内核在相同的内存对象上运行。如果减速的原因是因为内存对象的缓存，那么总体运行时间应该是 PLUS 和 MINUS 的单个运行时间的平均值（可能取决于实验 1）。

让我知道你是否知道这是否真的如此！

于 2012-09-20T22:11:50.290 回答