0

我有两个几乎相同的 OpenCL 内核，我想计算它们在 GFLOPS 中的性能。内核 #1 是：

__kernel void Test41(__global float *data, __global float *rands, int index, int rand_max){

    float16 temp;
    int gid = get_global_id(0);

    temp = data[gid];
    temp = (float) rands[1] * temp;
    temp = (float) rands[2] * temp;
    temp = (float) rands[3] * temp;
    temp = (float) rands[4] * temp;
    .
    .
    .
    temp = (float) rands[497] * temp;
    temp = (float) rands[498] * temp;
    temp = (float) rands[499] * temp;
    data[gid] = temp.s0;

}

第二个内核是：

__kernel void Test42(__global float *data, __global float *rands, int index, int rand_max){

    float16 temp[500];
    int gid = get_global_id(0);

    temp[0] = data[gid];
    temp[1] = (float) rands[1] * temp[0];
    temp[2] = (float) rands[2] * temp[1];
    temp[3] = (float) rands[3] * temp[2];
    temp[4] = (float) rands[4] * temp[3];
    .
    .
    .
    temp[497] = (float) rands[497] * temp[496];
    temp[498] = (float) rands[498] * temp[497];
    temp[499] = (float) rands[499] * temp[498];
    data[gid] = temp[index].s0;

}

正如您在代码中看到的，我使用的流大小为 16。每个内核都有 500 行操作，其中每个内核只执行一个浮点操作。我还总共部署了大约 1048576 个内核，因此我将有大约 1048576 个工作项并行执行。

为了计算我所做的翻牌：

flops = #numWorkItems(1048576) * (500) * StreamSize(16) / timeTaken;

不幸的是，对于第一个内核，我得到了大约 1.4 TFLOP，但对于第二个内核，我得到了 38 GFLOP。我无法解释这个巨大的差距。使用 temp 向量而不是单个 temp 似乎很重要。似乎真正的应用程序大多像第二个内核。第一个内核对于真正的应用程序来说太简单了。

谁能帮助我了解这里到底发生了什么以及第二个内核性能如何达到第一个？一般来说，如果我要对我的设备进行基准测试，我是否应该期望看到接近理论值的性能？

PS 我知道我需要将 rands 复制到 __local 内存中，但现在让我们跳过它。

4

2 回答 2

0

正如@pmdj 在评论中所建议的那样，第二个内核的主要问题是寄存器压力：您使用了大量的硬件寄存器，这减少了同时执行的工作组的数量。一般来说，大型私有数组在 OpenCL/CUDA 中不是一个好主意。在这种情况下，编译器几乎无法优化性能。您可以为数组使用本地内存，但是您需要添加适当的同步来访问它。

于 2017-05-09T08:31:55.500 回答
-1

有两个可能的问题：

  • 您将float16临时缓冲区声明为__private{这是 OpenCL 中的默认值}，并且很可能它将分配在global具有相当高的访问延迟的内存空间中。您可能会尝试将其声明为__local float16适合您的设备本地内存。
  • 添加temp缓冲区给编译器带来了一些问题......原始代码在某些 GPU 架构（例如 Intel）上很容易矢量化，并且您通过添加store+来添加人工依赖项load

我实际上会为此提交一份关于编译器的问题报告。编译器应该很容易找出依赖关系，进行优化甚至摆脱temp缓冲区。

于 2017-05-06T23:21:34.483 回答