随机访问时共享内存中的预期银行冲突数

Question

让A是32位整数中的共享存储器中的正确对齐阵列.

如果单个warp尝试A随机获取元素,那么预期的银行冲突数量是多少？

换一种说法:

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

请假设特斯拉或费米建筑.我不想详述Kepler的32位与64位存储区配置.而且,为简单起见,我们假设所有随机数都不同(因此没有广播机制).

我的直觉表明4到6之间的数字,但我想找到一些数学评估.

---

我相信这个问题可以从CUDA中抽象出来并作为数学问题出现.我搜索它作为生日悖论的延伸,但我在那里发现了非常可怕的公式,并没有找到最终的公式.我希望有一种更简单的方法......

Answer 1

在数学方面,这被认为是"垃圾箱中的球"问题 - 32个球随机掉入32个垃圾箱.您可以枚举可能的模式并计算其概率以确定分布.虽然模式的数量很大,但天真的方法是行不通的:(63!)/(32!)(31!)几乎是一个五分法.

如果您以递归方式构建解决方案并使用条件概率,则可以解决这个问题.

查找Charles J. Corrado撰写的一篇名为"Multinomial/Dirichlet和多变量超几何频率的最大值,最小值和范围的精确分布"的论文.

在下文中,我们从最左边的桶开始,计算可能落入其中的每个球数的概率.然后我们向右移动一个并确定在给定已经使用的球和桶的数量的情况下可以在该桶中的每个球数的条件概率.

为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的预期最大球数.

输出与ahmad相比:

           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

Answer 2

我假设 fermi 32 组共享内存，其中每 4 个后续字节都存储在后续组中。使用以下代码：

#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（为评估模型而生成的随机共享内存访问数量）来调整运行。