两个几乎相同的OpenCL内核之间的性能差距

Question

我有两个几乎相同的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 TFLOPs的，但对于第二个内核，我得到38个GFLOPS。我无法解释这个巨大的差距。使用一个温度矢量而不是单个温度似乎是一个巨大的交易。也似乎真正的应用程序大多像第二个内核。第一个内核对于真正的应用来说太简单了。

任何人都可以帮助我理解这里究竟发生了什么，以及第二个内核性能如何能够达到第一个？一般来说，如果我要对我的设备进行基准测试，我希望看到性能接近理论值吗？

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

Answer 1

有两个可能的问题：

• 你宣布float16临时缓冲区为__private {这是OpenCL的默认}，并最有可能将在global内存空间与相当高的访问延迟进行分配。如果它适合您的设备本地内存，您可能会尝试将其声明为__local float16。
• 添加temp缓冲创造了编译器的一些问题...原代码是在某些GPU架构易于量化的（英特尔为例），并通过添加store + load

其实我提交添加人工的依赖为编译器发布报告。编译器应该很容易找出依赖关系，进行优化，甚至摆脱您的temp缓冲区。

Answer 2

由于@pmdj在评论中提出，第二个内核的主要问题是注册压力：您正在使用大量硬件寄存器，这会减少同时执行的工作组数量。 一般来说，大型私有数组在OpenCL/CUDA中并不是一个好主意。 在这种情况下，编译器可以做的很少，以优化性能。 您可以使用本地内存作为阵列，但是您需要添加适当的同步来访问它。