0

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

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

问题 2：我一直在阅读很多关于共享内存性能的文章，一些文章（例如Dobbs 博士的文章）甚至指出它可以与寄存器一样快（只要访问良好）。在我的代码中，块内的所有线程都访问相同的共享变量。我相信在这种情况下，共享内存是以广播模式访问的，不是吗？它应该以这种方式达到寄存器的性能吗？是否有任何特殊的事情需要考虑才能让它发挥作用？

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

这里是

#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 微秒）为什么会发生这种情况？

问候丹尼尔

4

2 回答 2

1

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

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

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

于 2012-11-05T23:24:33.310 回答
0

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

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

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

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

于 2012-11-08T22:00:47.740 回答