0

パフォーマンスについて 2 つの質問をしたいと思います。説明するための簡単なコードを作成できませんでした。

質問 1: 非発散分岐のコストはどれくらいですか? 私のコードでは、fma 以外の 4 つの FLOPS に相当する値を超えることさえあるようです。述語がすでに計算されているBRA PTXコードについて話していることに注意してください

質問 2: 私は共有メモリのパフォーマンスについて多くのことを読んできました。Dr 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.44msかかる

最初の #if 条件を 0 に変更すると (共有メモリからの読み取りが有効になります)、6.02 ミリ秒かかります。それ以上ではありませんが、それでも十分ではありません。

2 番目の #if 条件 (true と評価されない分岐を挿入する) を有効にすると、9.647040ms で実行されます。パフォーマンスの低下は非常に大きいです。原因と対処法は？

また、コードを少し変更して、共有メモリでさらにチェックを行いました

それ以外の

__shared__ int t[1024]

やった

__shared__ int2 t[1024]

t[] にアクセスする場所はどこでも、t[].x にアクセスするだけです。パフォーマンスがさらに 10 ミリ秒に低下しました..(さらに 400 マイクロ秒) なぜこれが発生するのでしょうか?

よろしくダニエル

4

2 に答える 2

1

カーネルが計算バウンドかメモリバウンドかを判断しましたか? 最初の質問は、カーネルがコンピューティング バウンドの場合に最も関連性が高く、2 番目の質問は、カーネルがメモリ バウンドの場合に最も関連性があります。一方を仮定すると、混乱したり再現が困難な結果が得られる可能性がありますが、他方はそうです。

(1) 支店の費用は公表されていないと思います。あなたのアーキテクチャのために実験的にそれを決定することになるかもしれません。CUDA プログラミング ガイドには、「分岐予測も投機的実行もありません」と書かれています。

(2) ワープ内のすべてのスレッドから共有メモリ内の 1 つの 32 ビット値にアクセスすると、その値はブロードキャストされます。しかし、すべてのスレッドから単一の値にアクセスすることは、バンクの競合が発生しない限り、値の任意の組み合わせにアクセスすることと同じコストになると思います。そのため、共有メモリからの単一のフェッチのレイテンシが発生します。レイテンシのサイクル数は公開されていないと思います。通常は簡単に隠れるほど短いです。

于 2012-11-05T23:24:33.310 に答える
0

  • コンパイラーが高度に最適化されていることを覚えておく必要があります。したがって、ブランチをコメントアウトすると、ソースコードに残すかどうかに関係なく、条件付きの評価も削除されます。したがって、4つの命令の違いは、あなたの例にとって非常に妥当であるように思われます。

    1. 負荷-1、
    2. それと比較vして（そして結果をに保存してb）、
    3. テストb、
    4. ブランチ、

    私はあなたの例をコンパイルしてコードを見ていませんが（これはあなたがすべきことです-cuobjdump -sassあなたのバイナリで実行し、マシンコードの実際の違いを見てください。

  • .x変更のコンポーネントのみを使用するint2と、共有メモリのレイアウトが変更され、バンク競合のないアクセスから双方向のバンク競合に移行します。これにより、例ではさらに速度が低下します。IIRCは、共有メモリアクセスのレイテンシーが30サイクルのオーダーであり、通常、他のスレッドによって簡単に隠されます（Rogerがすでに述べたように）。

于 2012-11-08T22:00:47.740 に答える