7

共有メモリ内の 32 ビット整数Aの適切に配置された配列とします。

単一のワープが の要素をAランダムに取得しようとした場合、予想されるバンク競合の数は?

言い換えると：

__shared__ int A[N];          //N is some big constant integer
...
int v = A[ random(0..N-1) ];  // <-- expected number of bank conflicts here?

Tesla または Fermi アーキテクチャを想定してください。Kepler の 32 ビット バンク構成と 64 ビット バンク構成にこだわるつもりはありません。また、簡単にするために、すべての乱数が異なると仮定します (したがって、ブロードキャスト メカニズムはありません)。

私の直感では 4 から 6 の間の数字を示唆していますが、それを数学的に評価したいと思います。

この問題は CUDA から抽象化して、数学の問題として提示できると思います。誕生日のパラドックスの延長で探したのですが、本当に怖い公式があり、最終的な公式は見つかりませんでした。もっと簡単な方法があることを願っています...

4

3 に答える 3

6

数学では、これは「ボールがビンに入る」問題と考えられます。つまり、32 個のボールがランダムに 32 個のビンに落とされます。可能なパターンを列挙し、その確率を計算して分布を決定できます。ただし、パターンの数が非常に多いため、単純なアプローチは機能しません: (63!)/(32!)(31!) は「ほぼ」京です。

ただし、ソリューションを再帰的に構築し、条件付き確率を使用すると、対処することができます。

Charles J. Corrado による「多項式/ディリクレおよび多変量超幾何度数の最大値、最小値、および範囲の正確な分布」という論文を探してください。

以下では、左端のバケツから始めて、そこに落ちた可能性のあるボールの数ごとに確率を計算します。次に、1 つ右に移動し、既に使用されているボールとバケツの数を考慮して、そのバケツに入る可能性のある各数のボールの条件付き確率を決定します。

VBAコードについてはお詫びしますが、答えたいと思ったときに利用できるのはVBAだけでした:)。

Function nCr#(ByVal n#, ByVal r#)
    Static combin#()
    Static size#
    Dim i#, j#

    If n = r Then
        nCr = 1
        Exit Function
    End If

    If n > size Then
        ReDim combin(0 To n, 0 To n)
        combin(0, 0) = 1
        For i = 1 To n
            combin(i, 0) = 1
            For j = 1 To i
                combin(i, j) = combin(i - 1, j - 1) + combin(i - 1, j)
            Next
        Next
        size = n
    End If

    nCr = combin(n, r)
End Function

Function p_binom#(n#, r#, p#)
    p_binom = nCr(n, r) * p ^ r * (1 - p) ^ (n - r)
End Function

Function p_next_bucket_balls#(balls#, balls_used#, total_balls#, _
  bucket#, total_buckets#, bucket_capacity#)

    If balls > bucket_capacity Then
        p_next_bucket_balls = 0
    Else
        p_next_bucket_balls = p_binom(total_balls - balls_used, balls, 1 / (total_buckets - bucket + 1))
    End If
End Function

Function p_capped_buckets#(n#, cap#)
    Dim p_prior, p_update
    Dim bucket#, balls#, prior_balls#

    ReDim p_prior(0 To n)
    ReDim p_update(0 To n)

    p_prior(0) = 1

    For bucket = 1 To n
        For balls = 0 To n
            p_update(balls) = 0
            For prior_balls = 0 To balls
                p_update(balls) = p_update(balls) + p_prior(prior_balls) * _
                   p_next_bucket_balls(balls - prior_balls, prior_balls, n, bucket, n, cap)
            Next
        Next
        p_prior = p_update
    Next

    p_capped_buckets = p_update(n)
End Function

Function expected_max_buckets#(n#)
    Dim cap#

    For cap = 0 To n
        expected_max_buckets = expected_max_buckets + (1 - p_capped_buckets(n, cap))
    Next
End Function

Sub test32()
    Dim p_cumm#(0 To 32)
    Dim cap#

    For cap# = 0 To 32
        p_cumm(cap) = p_capped_buckets(32, cap)
    Next

    For cap = 1 To 32
        Debug.Print "    ", cap, Format(p_cumm(cap) - p_cumm(cap - 1), "0.000000")
    Next
End Sub

32 個のボールとバケツの場合、バケツ内のボールの予想最大数は約 3.532941 です。

アフマドのものと比較する出力:

           1            0.000000
           2            0.029273
           3            0.516311
           4            0.361736
           5            0.079307
           6            0.011800
           7            0.001417
           8            0.000143
           9            0.000012
           10           0.000001
           11           0.000000
           12           0.000000
           13           0.000000
           14           0.000000
           15           0.000000
           16           0.000000
           17           0.000000
           18           0.000000
           19           0.000000
           20           0.000000
           21           0.000000
           22           0.000000
           23           0.000000
           24           0.000000
           25           0.000000
           26           0.000000
           27           0.000000
           28           0.000000
           29           0.000000
           30           0.000000
           31           0.000000
           32           0.000000
于 2012-10-18T14:56:06.517 に答える
4

数学の答えを試してみますが、まだ正確ではありません。

__shared__基本的に知りたいのは、整列配列へのワープ内のランダムな 32 ビット ワード インデックス付けを考えると、「 1 つのバンクにマップされるワープ内のアドレスの最大数の期待値は?」ということです。

問題をハッシングに似ていると考えると、それは単一の場所にハッシュされるアイテムの予想される最大数に関連しており、このドキュメントはハッシング nの O(log n / log log n) の数の上限を示していますアイテムを n 個のバケットに入れます。（数学はかなり難しいです！）。

n = 32 の場合、約 2.788 になります (自然対数を使用)。それは問題ありませんが、ここでは期待される最大値を経験的に計算するために ahmad のプログラムを少し変更しました (コードを簡素化し、名前などを変更してわかりやすくし、いくつかのバグを修正しました)。

#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <algorithm>
#define NBANK 32
#define WARPSIZE 32
#define NSAMPLE 100000

int main(){  
    int i=0,j=0;

    int *bank=(int*)malloc(sizeof(int)*NBANK);
    int *randomNumber=(int*)malloc(sizeof(int)*WARPSIZE);
    int *maxCount=(int*)malloc(sizeof(int)*(NBANK+1));
    memset(maxCount, 0, sizeof(int)*(NBANK+1));

    for (int i=0; i<NSAMPLE; ++i) {
        // generate a sample warp shared memory access
        for(j=0; j<WARPSIZE; j++){
            randomNumber[j]=rand()%NBANK;
        }

        // check the bank conflict
        memset(bank, 0, sizeof(int)*NBANK);
        int max_bank_conflict=0;
        for(j=0; j<WARPSIZE; j++){
            bank[randomNumber[j]]++;       
        }

        for(j=0; j<WARPSIZE; j++) 
            max_bank_conflict = std::max<int>(max_bank_conflict, bank[j]);

        // store statistic
        maxCount[max_bank_conflict]++;
    }

    // report statistic
    printf("Max conflict degree %% (%d random samples)\n", NSAMPLE);
    float expected = 0;
    for(i=1; i<NBANK+1; i++) {
        float prob = maxCount[i]/(float)NSAMPLE;
        printf("%02d -> %6.4f\n", i, prob);
        expected += prob * i;
    }
    printf("Expected maximum bank conflict degree = %6.4f\n", expected);
    return 0;
}

プログラムで見つかったパーセンテージを確率として使用すると、期待される最大値は1 から 32 までの積の合計にsum(i * probability(i))なります。期待値は 3.529 と計算されます (ahmad のデータと一致します)。i極端に離れているわけではありませんが、2.788 が上限になるはずです。上限は Big-O 表記で指定されているため、一定の要素が省略されていると思います。しかし、それは現在私が得た限りです。

未解決の質問: その一定の要因はそれを説明するのに十分ですか? n = 32 の定数係数を計算することは可能ですか? これらを調整すること、および/または 32 バンクと 32 並列スレッドで予想される最大バンク競合度のクローズド フォーム ソリューションを見つけることは興味深いでしょう。

これは、共有メモリのアドレス指定が事実上ランダムである場合のパフォーマンスのモデル化と予測に役立つため、非常に有用なトピックです。

于 2012-10-12T00:28:13.960 に答える
3

後件部の各4バイトが後件部のバンクに格納されているfermi32バンク共有メモリを想定しています。次のコードを使用します。

#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#define NBANK 32
#define N 7823
#define WARPSIZE 32


#define NSAMPLE 10000

int main(){
    srand ( time(NULL) );

    int i=0,j=0;
    int *conflictCheck=NULL;
    int *randomNumber=NULL;
    int *statisticCheck=NULL;

    conflictCheck=(int*)malloc(sizeof(int)*NBANK);
    randomNumber=(int*)malloc(sizeof(int)*WARPSIZE);
    statisticCheck=(int*)malloc(sizeof(int)*(NBANK+1));
    while(i<NSAMPLE){
        // generate a sample warp shared memory access
        for(j=0; j<WARPSIZE; j++){
            randomNumber[j]=rand()%NBANK;
        }
        // check the bank conflict
        memset(conflictCheck, 0, sizeof(int)*NBANK);
        int max_bank_conflict=0;
        for(j=0; j<WARPSIZE; j++){
            conflictCheck[randomNumber[j]]++;
            max_bank_conflict = max_bank_conflict<conflictCheck[randomNumber[j]]? conflictCheck[randomNumber[j]]: max_bank_conflict;
        }
        // store statistic
        statisticCheck[max_bank_conflict]++;

        // next iter
        i++;
    }
    // report statistic
    printf("Over %d random shared memory access, there found following precentages of bank conflicts\n");
    for(i=0; i<NBANK+1; i++){
        //
        printf("%d -> %6.4f\n",i,statisticCheck[i]/(float)NSAMPLE);
    }
    return 0;
}

次の出力が得られました：

Over 0 random shared memory access, there found following precentages of bank conflicts
0 -> 0.0000
1 -> 0.0000
2 -> 0.0281
3 -> 0.5205
4 -> 0.3605
5 -> 0.0780
6 -> 0.0106
7 -> 0.0022
8 -> 0.0001
9 -> 0.0000
10 -> 0.0000
11 -> 0.0000
12 -> 0.0000
13 -> 0.0000
14 -> 0.0000
15 -> 0.0000
16 -> 0.0000
17 -> 0.0000
18 -> 0.0000
19 -> 0.0000
20 -> 0.0000
21 -> 0.0000
22 -> 0.0000
23 -> 0.0000
24 -> 0.0000
25 -> 0.0000
26 -> 0.0000
27 -> 0.0000
28 -> 0.0000
29 -> 0.0000
30 -> 0.0000
31 -> 0.0000
32 -> 0.0000

ランダムアクセスでは、3〜4ウェイの競合が最も可能性が高いと結論付けることができます。N異なる（配列内の要素の数）、NBANK（共有メモリ内のバンクの数）、WARPSIZE（マシンのワープサイズ）、およびNSAMPLE（モデルを評価するために生成されたランダム共有メモリアクセスの数）を使用して実行を調整できます。

于 2012-10-11T08:47:35.333 に答える