Numero previsto di conflitti bancari nella memoria condivisa all'accesso casuale

Question

Lasciare A una matrice allineata correttamente di numeri interi a 32 bit nella memoria condivisa.

Se un singolo ordito tenta di recuperare elementi di A a caso, qual è il numero previsto di conflitti bancari?

In altre parole:

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

Si prega di assumere architettura Tesla o Fermi. Non voglio soffermarmi sulle configurazioni bancarie da 32 bit a 64 bit di Kepler. Inoltre, per semplicità, supponiamo che tutti i numeri casuali siano diversi (quindi nessun meccanismo di trasmissione).

Il mio istinto suggerisce un numero compreso tra 4 e 6, ma mi piacerebbe trovare qualche valutazione matematica di esso.

---

Credo che il problema possa essere estratto dal CUDA e presentato come un problema di matematica. L'ho cercato come estensione di Birthday Paradox, ma ho trovato formule davvero spaventose e non ho trovato una formula finale. Spero ci sia un modo più semplice ...

Answer 1

Presumo fermi 32-bank di memoria condivisa in cui ogni 4 byte conseguenti sono memorizzati in banchi conseguenti. Utilizzando il codice seguente:

#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; 
} 

mi sono seguente output:

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 

possiamo arrivare a concludere che 3 a 4 modo conflitto è il più probabile con accesso casuale. È possibile ottimizzare l'esecuzione con diversi (numero di elementi nell'array), NBANK (numero di banchi nella memoria condivisa), WARPSIZE (dimensione del foglio di ordito della macchina) e NSAMPLE (numero di accessi di memoria condivisa casuale generati per valutare il modello).

Answer 2

Proverò una risposta matematica, anche se non ho ancora ragione.

Fondamentalmente si vuole sapere, dato l'indicizzazione casuale di parole a 32 bit all'interno di un ordito in un allineamento __shared__ allineato, "qual è lo expected value del numero massimo di indirizzi all'interno di una trama che si associano a un singolo banco?"

Se si considera il problema simile all'hash, quindi si riferisce al numero massimo previsto di elementi che eseguiranno l'hash in una singola posizione e this document shows an upper bound on that number of O(log n/log log n) per l'hashing n elementi in n bucket. (La matematica è piuttosto pelosa!).

Per n = 32, funziona a circa 2,778 (utilizzando registro naturale). Va bene, ma qui ho modificato un po 'il programma di ahmad per calcolare empiricamente il massimo previsto (ho anche semplificato il codice e i nomi modificati e così via per chiarezza e corretto alcuni bug).

#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; 
} 

Utilizzando le percentuali riscontrate nel programma come probabilità, il valore massimo previsto è la somma dei prodotti sum(i * probability(i)), per i da 1 a 32. I calcolare il valore dovrebbe essere 3.529 (partite di dati di Ahmad). Non è molto lontano, ma il 2.788 dovrebbe essere un limite superiore. Dal momento che il limite superiore è dato in notazione big-O, immagino che rimanga un fattore costante. Ma questo è quanto ho ottenuto finora.

Domande aperte: questo fattore costante è sufficiente per spiegarlo? È possibile calcolare il fattore costante per n = 32? Sarebbe interessante riconciliarli e/o trovare una soluzione di forma chiusa per il massimo grado di conflitto bancario previsto con 32 banchi e 32 thread paralleli.

Questo è un argomento molto utile, poiché può aiutare nella modellazione e nella previsione delle prestazioni quando l'indirizzamento della memoria condivisa è effettivamente casuale.

Answer 3

In matematica, questo è pensato come un problema "palle in bidoni" - 32 palline vengono casualmente cadute in 32 contenitori. È possibile enumerare i possibili modelli e calcolare le loro probabilità per determinare la distribuzione. Un approccio ingenuo non funzionerà anche se il numero di pattern è enorme: (63!)/(32!) (31!) È "quasi" un quintilione.

È possibile affrontare il problema se si crea la soluzione in modo ricorsivo e si utilizzano le probabilità condizionali.

Cercare un documento chiamato "La distribuzione esatta del massimo, minimo e la gamma di frequenze Multinomiali/Dirichlet e Multivariate Ipergeometriche" di Charles J. Corrado.

In seguito, partiamo dal bucket più a sinistra e calcoliamo le probabilità per ciascun numero di sfere che potrebbero esserci cadute. Poi ne spostiamo uno a destra e determiniamo le probabilità condizionali di ogni numero di palle che potrebbero essere in quel secchio dato il numero di palle e secchi già usati.

Ci scusiamo per il codice VBA, ma VBA era tutto ciò che avevo a disposizione quando ero motivato a rispondere :).

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 

per 32 palline e secchi, ottengo un numero massimo atteso di palline nei secchi di circa 3,532,941 mila.

uscita per confrontare Ahmad di:

  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