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

Question

假设在共享内存中有一个32位整数的正确对齐数组A。

如果单个warp尝试随机获取A的元素，那么预期的bank冲突次数是多少？

换句话说：

__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中抽象出来，并作为一个数学问题来呈现。我将其搜索为生日悖论的扩展，但那里有一些非常可怕的公式，也没有找到最终的公式。希望有更简单的方法...

- CygnusX1

---

好问题，点赞！（我们需要在cuda标签上更多地点赞好问题！） - harrism

Answer 1

在数学中，这被认为是一个“球放入箱子”的问题——32个球随机落入32个箱子。您可以列举可能的模式并计算它们的概率以确定分布情况。然而，天真的方法行不通，因为模式数量巨大：(63!)/(32!)(31!)几乎是一百万亿。

但是，如果您递归地构建解决方案并使用条件概率，就有可能解决这个问题。

请查看名为“Multinomial/Dirichlet和Multivariate Hypergeometric频率的最大值、最小值和范围的确切分布”的论文，作者是Charles J. Corrado。

在下面的例子中，我们从最左边的桶开始，计算每个可能掉落到该桶中的球的概率。然后我们向右移动一个位置，并根据已经使用的球和桶的数量来确定每个可能在该桶中的球数的条件概率。

对于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

- A. Webb

---

我忘了提到，但应该插入一个很棒的博客，通过harrism的帖子来到这里。 - A. Webb

非常棒的答案，谢谢A。顺便说一下，提到的博客是John D Cook的，恰巧最近有一篇与这个问题相关的文章。 - harrism

Answer 2

我将尝试一个数学答案，虽然我还没有完全正确。

基本上，您想知道，在对齐的__shared__数组中，给定warp内随机32位字索引，“最大数量的地址映射到单个bank的期望值是多少？”

如果我认为这个问题类似于哈希，那么它涉及到预期的最大项数，这些项将哈希到单个位置，并且此文档显示了将n个项哈希到n个桶中的上限为O(log n / log log 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;
}

使用程序中找到的百分比作为概率，期望最大值是产品之和sum(i * probability(i))，对于i从1到32。我计算出期望值为3.529（与ahmad的数据相匹配）。它不是超出预期太远，但2.788应该是一个上限。由于上限以big-O符号给出，我猜测有一个常数因子被省略了。但这是我目前所能做到的。

未解决问题：这个常数因子足够解释吗？是否可能计算出n = 32的常数因子？如果能协调这些，并/或者找到一个32个银行和32个并行线程的预期最大银行冲突度的闭合形式解，那将会很有趣。

这是一个非常有用的主题，因为它可以帮助建立模型和预测共享内存寻址时的性能。

- harrism

Answer 3

我假设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（机器的warp大小）和NSAMPLE（生成以评估模型的随机共享内存访问次数）来调整运行。

- lashgar

---

哈！在数学失败的地方，实验起作用了 :) 我仍然想看到一些解决方案的数学公式，但如果我找不到任何东西，我就会接受这种方法。 - CygnusX1

rand()是一个相当糟糕的随机数生成器，在Windows上只有16位，而在Linux上则为32位... - harrism

@harrism：我已经在Linux上报告了这个数字。我认为rand()足够好，因为我们只需要rand()%32。 - lashgar

这个分析对我来说似乎非常明智和有用。作为一个小观察，我认为在 Fermi 的情况下，如果 2（或更多）个线程访问数组中的完全相同位置，那不是银行冲突，而是通过广播机制解决的。但从统计学的角度来看，我认为这不会对结果产生太大或根本没有影响。而且数组越大，发生这种情况的可能性就越小。 - Robert Crovella

似乎很难得到一个数学公式，我只能接受这个答案。但如果有人有更好的想法，请分享！ - CygnusX1