CUDA 性能：分支和共享内存

Question

我想问两个关于性能的问题。我一直无法创建简单的代码来说明。

问题 1：非发散分支有多贵？在我的代码中，它似乎甚至超过了 4 个非 fma FLOPS 的等价物。请注意，我说的是 BRA PTX 代码，其中已经计算了谓词

问题 2：我已经阅读了很多关于共享内存性能的文章，并且一些像 a Dr Dobbs article 这样的文章甚至声称它可以和寄存器一样快（只要访问良好）。在我的代码中，块内的所有线程都访问相同的共享变量。我相信在这种情况下，共享内存是以广播模式访问的，不是吗？它应该以这种方式达到寄存器的性能吗？是否有什么特殊的事情需要考虑才能让它发挥作用？

编辑：我已经能够构建一些简单的代码，为我的查询提供更多洞察力

这里是

#include <math.h>
#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <float.h>
#include "cuComplex.h"
#include "time.h"
#include "cuda_runtime.h"
#include <iostream>
using namespace std;

__global__ void test()
{
__shared__ int t[1024];
   int v=t[0];
    bool b=(v==-1);
    bool c=(v==-2);
    int myValue=0;
    for (int i=0;i<800;i++)
    {
#if 1
            v=i;
#else
            v=t[i];
#endif

#if 0
            if (b) {
                    printf("abs");
            }
#endif
            if (c)
            {
                    printf ("IT HAPPENED");
                    v=8;
            }
            myValue+=v;

    }
    if (myValue==1000)
            printf ("IT HAPPENED");



}
int main(int argc, char *argv[])
{
    cudaEvent_t event_start,event_stop;
    float timestamp;
float4  *data;
    // Initialise
    cudaDeviceReset();
    cudaSetDevice(0);
dim3 threadsPerBlock;
dim3 blocks;
 threadsPerBlock.x=32;
 threadsPerBlock.y=32;
 threadsPerBlock.z=1;
 blocks.x=1;
 blocks.y=1000;
 blocks.z=1;
 cudaEventCreate(&event_start);
 cudaEventCreate(&event_stop);
cudaEventRecord(event_start, 0);
test<<<blocks,threadsPerBlock,0>>>();
    cudaEventRecord(event_stop, 0);
    cudaEventSynchronize(event_stop);
    cudaEventElapsedTime(&timestamp, event_start, event_stop);
    printf("Calculated in %f", timestamp);
}

我在 GTX680 上运行此代码。

现在结果如下..

如果按原样运行则需要 5.44 毫秒

如果我将第一个 #if 条件更改为 0（这将启用从共享内存中读取），它将需要 6.02 毫秒。虽然不多，但对我来说仍然不够

如果我启用第二个 #if 条件（插入一个永远不会评估为真的分支），它会在 9.647040 毫秒内运行。性能下降非常大。是什么原因，有什么办法？

我还稍微更改了代码以进一步检查共享内存

代替

__shared__ int t[1024]

我做到了

__shared__ int2 t[1024] 

无论我在哪里访问 t[]，我都只是访问 t[].x。性能进一步下降到 10 毫秒..（另外 400 微秒）为什么会发生这种情况？

问候 丹尼尔

Answer 1

您确定您的内核是受计算限制还是内存限制？如果您的内核受计算限制，您的第一个问题将最相关，而如果您的内核受内存限制，则第二个问题最相关。如果您假设一个结果，那么您可能会得到令人困惑或难以重现的结果，而实际上是另一个结果。

(1) 我认为分支机构的成本尚未公布。您可能需要通过实验来确定您的架构。 CUDA 编程指南确实说没有“分支预测和推测执行”。

(2) 没错，当您从 warp 中的所有线程访问共享内存中的单个 32 位值时，该值是广播的。但我的猜测是，只要您不引起任何银行冲突，从所有线程访问单个值的成本与访问任何值组合的成本相同。因此，您最终会遇到从共享内存中进行单次提取的延迟。我不认为延迟的周期数已经公布。它足够短，通常很容易隐藏。

Answer 2

• 您需要记住，编译器是高度优化的。所以如果你注释掉分支，你也消除了条件的评估，不管你是否把它留在源代码中。因此，对于您的示例，四个指令的差异似乎非常合理：

  1. 加载-1,
  2. 比较v （并将结果存储在b），
  3. 测试b,
  4. 分支，

  虽然我没有编译您的示例并查看代码（这是您应该做的 - 在您的二进制文件上运行 cuobjdump -sass 并查看机器代码的实际差异。

• 使用 int2 中唯一的 .x 组件会更改共享内存中的布局，以便您从无银行冲突访问变为 2 路银行冲突，这会导致您的例子。 IIRC 共享内存访问的延迟大约为 30 个周期，这通常很容易被其他线程隐藏（正如 Roger 已经提到的）。