Número previsto de conflictos bancarios en la memoria compartida con acceso aleatorio

Question

Sea A una matriz alineada correctamente de enteros de 32 bits en la memoria compartida.

Si un warp único intenta obtener elementos de A al azar, ¿cuál es el número esperado de conflictos bancarios?

En otras palabras:

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

favor suponer Tesla Fermi o la arquitectura. No quiero detenerme en las configuraciones de 32 bits frente a las de 64 bits de Kepler. Además, para simplificar, supongamos que todos los números aleatorios son diferentes (por lo tanto, no hay mecanismo de difusión).

Mi intuición sugiere un número entre 4 y 6, pero me gustaría encontrar alguna evaluación matemática de la misma.

---

Creo que el problema puede ser abstraído de CUDA y presentado como un problema de matemáticas. Lo busqué como una extensión de Birthday Paradox, pero encontré fórmulas realmente aterradoras y no encontré una fórmula final. Espero que haya una manera más simple ...

Answer 1

Supongo fermi 32-bank shared memory donde cada 4 bytes consecutivos se almacenan en los bancos correspondientes. Utilizando el código siguiente:

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

me dieron resultado siguiente:

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 

podemos llegar a la conclusión de que 3 a 4 conflicto manera es la más probable con acceso aleatorio. Puede sintonizar la ejecución con diferentes N (cantidad de elementos en la matriz), NBANK (número de bancos en la memoria compartida), WARPSIZE (tamaño de la urdimbre de la máquina) y NSAMPLE (número de accesos aleatorios de memoria compartida generados para evaluar el modelo).

Answer 2

Voy a intentar una respuesta matemática, aunque todavía no la tengo del todo bien.

Es, básicamente, quieren saber, dada la indexación al azar palabra de 32 bits dentro de un túnel en una gama __shared__ alineados, "¿cuál es la expected value del número máximo de direcciones dentro de un túnel del que corresponde a una única banco?"

Si considero el problema similar al de hash, a continuación, se relaciona con el número máximo esperado de elementos que hash en un solo lugar, y this document shows an upper bound on that number of O(log n/log log n) de hash n elementos en n cubos. (¡La matemática es bastante peluda!).

Para n = 32, eso equivale a aproximadamente 2.788 (usando el registro natural). Está bien, pero aquí modifiqué un poco el programa de ahmad para calcular empíricamente el máximo esperado (también simplifiqué el código y modifiqué los nombres y tal para mayor claridad y corrigí algunos errores).

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

Usando los porcentajes que se encuentran en el programa como probabilidades, el valor máximo esperado es la suma de los productos sum(i * probability(i)), por i de 1 a 32. I calcular el valor esperado para ser 3,529 (coincide con los datos de Ahmad). No está muy lejos, pero se supone que el 2.788 es un límite superior. Dado que el límite superior se da en notación de O grande, creo que hay un factor constante omitido. Pero eso es actualmente lo más lejos que he llegado.

Preguntas abiertas: ¿Es ese factor constante suficiente para explicarlo? ¿Es posible calcular el factor constante para n = 32? Sería interesante conciliar estos, y/o encontrar una solución de forma cerrada para el grado de conflicto bancario máximo esperado con 32 bancos y 32 hilos paralelos.

Este es un tema muy útil, ya que puede ayudar a modelar y predecir el rendimiento cuando el direccionamiento de memoria compartida es efectivamente aleatorio.

Answer 3

En matemáticas, esto se considera un problema de "bolas en los contenedores": 32 bolas se depositan aleatoriamente en 32 contenedores. Puede enumerar los patrones posibles y calcular sus probabilidades para determinar la distribución. Sin embargo, un enfoque ingenuo no funcionará, ya que el número de patrones es enorme: (63!)/(32!) (31!) Es "casi" un quintillón.

Es posible abordar sin embargo si construyes la solución recursivamente y usas probabilidades condicionales.

Busque un artículo llamado "La distribución exacta de la máxima, mínima y el rango de Multinomial/Dirichlet y frecuencias hipergeométricas multivariantes" por Charles J. Corrado.

A continuación, comenzamos en el cubo más a la izquierda y calculamos las probabilidades para cada número de bolas que podrían haber caído en él. Luego nos movemos uno a la derecha y determinamos las probabilidades condicionales de cada número de bolas que podrían estar en ese cubo dado el número de bolas y cubos que ya se utilizaron.

Disculpas por el código VBA, pero VBA era todo lo que tenía disponible cuando estaba motivado para responder :).

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 

para 32 bolas y cubos, me sale un número máximo esperado de las bolas en los cubos de alrededor de 3,532941.

salida de comparar a 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