cuda - 随机访问时共享内存中的预期存储库冲突数

让 A 成为共享内存中正确对齐的 32 位整数数组。

如果单个 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?

请假设特斯拉或费米架构。我不想详述 Kepler 的 32 位和 64 位 bank 配置。另外，为简单起见，让我们假设所有随机数都是不同的(因此没有广播机制)。

我的直觉暗示了一个介于 4 和 6 之间的数字，但我想对它进行一些数学评估。

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

最佳答案

我假设 fermi 32-bank 共享内存，其中每 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(共享内存中的 bank 数)、WARPSIZE(机器的扭曲大小)和 NSAMPLE(为评估模型而生成的随机共享内存访问数)来调整运行。

关于cuda - 随机访问时共享内存中的预期存储库冲突数，我们在Stack Overflow上找到一个类似的问题：https://stackoverflow.com/questions/12823597/