Nicolás Wolovick, $Date: 2012-06-20 13:08:26 -0300 (Wed, 20 Jun 2012) $, $Revision: 3623 $
· Grilla bidimensional M de tamaño total N = L×L.
· Cada celda contiene un spin, M(i,j) ∈ (0..q].
· Los valores de q van desde 2 hasta algunos cientos.
· Para q=2, es el Modelo de Ising.
La energía H está dada por el Hamilitoniano:
Esto es:
· Para cada celda i,j acumular los primeros vecinos.
· Se acumula el delta de Kronecker δ(i,j) = i==j
.
· Esto sobre todas las casillas.
Dinámica probabilística, usando el algoritmo Metrópolis.
Un Paso de Monte Carlo (MCS) está dado por el mapeo a todas las casillas i,j de:
Notar que:
Ejemplo evolución de la energía: q=24, L=200. Notar como el sistema pasa de totalmente desordenado (máxima energía) a ordenado (mínima energía) de manera muy rápida: transición de fase.
Para q≤4 esta discontinuidad no existe. Da la pauta que es un modelo muy rico a pesar de su simplicidad.
L=2048, q=9.
Total:
Por cada celda:
fp32
sorteado uniformemente en [0,1).Para el ejemplo anterior totaliza Doce Teras de números aleatorios.
>>> 1650000*2048*2048*2.0 / (1<<40)
12.5885009765625
Necesitamos un RNG:
Procesar las celdas puede tener diversos órdendes.
Optamos por checkerboard.
Estado: 64 bits.
Siguiente elemento de secuencia: x{n+1} = (xn × a + cn ) mod b.
Donde:
· a es el multiplicador,
· b la base, fijada a (1<<32)
para operar bitwise
· cn el acarreo del módulo anterior.
Cada multiplicador bueno genera un secuencia independiente de RNG, pero los multiplicadores a tienen que satisfacer.
goodmult(a) ≡ prime(a × 2^32 − 1) ∧ prime((a × 2^32 − 2)/2)
Obtener estos goodmult es computacionalmente intensivo: computarlos y almacenarlos.
¿Cuántos RNG independientes aka goodmult(a) se puede tener?
Tensión entre:
· Paralelismo.
· Memoria ocupada por el estado de los RNG y calidad de las secuencias.
Para esta implementación fijamos.
Para trabajar tamaños mayores 512×512 ...
Un hilo computa múltiples casillas, una por frame.
Adherimos a la práctica "Compute multiple outputs per thread": Volkov y Baxter.
Primeros vecinos divide en dos subconjuntos independientes: blancas y negras.
Dentro de cada subconjunto no hay dependencia de datos.
Se pueden procesar en paralelo.
Conocido como Red-Black Gauss-Seidel.
updateCUDA
1 unsigned int i;
2 // move thread RNG state to registers. Thanks to Carlos Bederián for pointing this out.
3 unsigned long long rng_state = d_x[tid];
4 const unsigned int rng_const = d_a[tid];
5 for (unsigned int iFrame=0; iFrame<(L/FRAME); iFrame++) {
6 i = iOriginal + (FRAME/2)*iFrame;
7 unsigned int j;
8 for (unsigned int jFrame=0; jFrame<(L/FRAME); jFrame++) {
9 j = jOriginal + FRAME*jFrame;
10
11 spin_old = write[i*L + j];
12
13 // computing h_before
14 spin_neigh_n = read[i*L + j]; spin_neigh_e = read[i*L + (j+1)%L];
15 spin_neigh_w = read[i*L + (j-1+L)%L];
16 spin_neigh_s = read[((i+(2*(color^(j%2))-1)+L/2)%(L/2))*L + j];
17 h_before = - (spin_old==spin_neigh_n) - (spin_old==spin_neigh_e)
18 - (spin_old==spin_neigh_w) - (spin_old==spin_neigh_s);
19 // new spin
20 spin_new = (spin_old + (byte)(1 + rand_MWC_co(&rng_state, &rng_const)*(Q-1))) % Q;
21 // h after taking new spin
22 h_after = - (spin_new==spin_neigh_n) - (spin_new==spin_neigh_e)
23 - (spin_new==spin_neigh_w) - (spin_new==spin_neigh_s);
24
25 delta_E = h_after - h_before;
26 float p = rand_MWC_co(&rng_state, &rng_const);
27 if (delta_E<=0 || p<=expf(-delta_E/temp)) {
28 write[i*L + j] = spin_new;
29 }
30 }
31 }
32 d_x[tid] = rng_state; // store RNG state into global again
updateCUDA
, highlightsexpf(-delta_E/temp)
se puede tabular para cada temperatura, delta_E
tiene 9 valores posibles.int change = delta_E<=0 || p<=expf(-delta_E/temp);
write[i*L + j] = (change)*spin_new + (1-change)*spin_old;
expf(-delta_E/temp)
(recuerdo que la ganancia era interesante). gcc-4.5
e icc-11
.1 #pragma omp parallel shared(read,write,table_expf_temp)
2 {
3 unsigned int tid = omp_get_thread_num();
4 unsigned long long x_l = x[tid]; // load RNG state into local storage
5 unsigned int a_l = a[tid];
6 #pragma omp for
7 for (unsigned int i=0; i<L/2; i++) {
8 ...
Notar que usamos la misma idea de mover el estado a registros o memoria per-thread. Sino, problemas gruesos de false sharing.
El scaling es perfecto (a partir de cierto tamaño).
Notar que hay barras de error, pero son tan chicas que se ocultan en el punto.
nvcc-3.2 -O3
(circa 2010).gcc-4.4.5 -O3 -ffast-math -march=native -funroll-loops
..cu
corriendo en una placa Fermi, no tiene diferencias si lo compilamos para CC 1.3 y CC 2.0.-Xptxas -dlcm=cg -Xptxas -dscm=cg
: deshabilitar la cache L1. (ver manual de PTX). 1 2011-10-13 08:55 bc
2 * [r2907] potts3.cu: Dont use predication when calculating neighbor
3 indices
4 2011-10-13 05:28 bc
5 * [r2906] potts3.cu: Unpacked red/black split
6 2011-10-13 00:04 bc
7 * [r2905] potts3.cu: Pasar el color del update a template argument
8 2011-06-22 07:01 bc
9 * [r2792] potts3.cu: Rectangular tiles
10 Remove asserts from timing commands
11 2011-05-02 12:14 nicolasw
12 * [r2607] potts3.cu: Un *1 de más que encontró Javier Uranga
C++
para generar dos versiones de updateCUDA
, una para blancas y otra para negras.Potts_orig
Corrida actual sobre Tesla C2070.
Probamos algunas cosas, como por ejemplo compilar con nvcc-4.0
(Open64 backend) vs. nvcc-4.2
(LLVM backend).
CUDA-4.0
Q L spinflip stddev
9 512 0.422714 0.00286214
9 1024 0.294745 0.00165411
9 2048 0.259109 0.00131451
9 4096 0.244878 0.00108468
9 8192 0.236913 0.000977784
9 16384 0.233761 0.00104113
9 32768 0.238528 0.00123815
CUDA-4.2
9 512 0.423447 0.00246027
9 1024 0.282109 0.000991892
9 2048 0.25659 0.000240049
9 4096 0.243577 0.000213649
9 8192 0.236021 0.000164541
9 16384 0.234637 0.000110928
9 32768 0.239124 0.000582012
· Se comparta casi igual en cuanto a tiempos. Notar que si mejora la desviación estándar en nvcc-4.2
.
· Los tiempos son mayores que los reportados para GTX 480 (0.147ns). Las Tesla tienen menor frecuencia para los SM y para la memoria.
Potts_int
Intentamos cambiar todos los unsigned int
a int
. En scan_warp
producía efectos benéficos.
CUDA-4.0
Q L spinflip stddev
9 512 0.4268 0.0025893
9 1024 0.327349 0.00150298
9 2048 0.290854 0.00140329
9 4096 0.27214 0.00111708
9 8192 0.264136 0.00108918
9 16384 0.260818 0.00106291
9 32768 0.261503 0.000845782
CUDA-4.2
9 512 0.442921 0.00349175
9 1024 0.305273 0.00077226
9 2048 0.297691 0.00022777
9 4096 0.288527 0.000144492
9 8192 0.283241 8.79968e-05
9 16384 0.28131 8.16863e-05
9 32768 0.282043 0.000379256
· Solo empeora.
· Acá si se nota algo de diferencia entre los compiladores.
Potts_charly
Versión "de la noche después" de Carlos Bederián.
CUDA-4.0
Q L spinflip stddev
9 512 0.320011 0.00632055
9 1024 0.255473 0.00358806
9 2048 0.224165 0.00214458
9 4096 0.214321 0.00229368
9 8192 0.210495 0.002451
9 16384 0.208602 0.00248176
9 32768 0.208539 0.00221192
CUDA-4.2
9 512 0.327076 0.00926237
9 1024 0.226146 0.00335537
9 2048 0.221503 0.000969065
9 4096 0.213315 0.000593172
9 8192 0.210459 0.000496381
9 16384 0.209569 0.000602072
9 32768 0.21212 0.000555135
· 15% de ganancia, el mejor código hasta ahora.
· Nota que sigue siendo consistente la menor stddev de nvcc-4.2
.
· No hay diferencia entre los compiladores.
Potts_charly
-2Algunos cambios menores.
CUDA-4.2 (int)
9 512 0.352479 0.00793455
9 1024 0.243674 0.00314107
9 2048 0.233309 0.0018368
9 4096 0.221555 0.000589989
9 8192 0.217872 0.00119905
9 16384 0.215092 0.000878846
9 32768 0.216102 0.000658448
CUDA-4.2 (-Xptxas -dlcm=cg -Xptxas -dscm=cg)
9 512 0.342793 0.0105217
9 1024 0.22178 0.00296488
9 2048 0.215611 0.000488704
9 4096 0.209739 0.000243782
9 8192 0.206776 0.000155276
9 16384 0.20567 0.000127431
9 32768 0.205315 0.000125488
· Nuevamente cambiar a int
no mejora.
· 2% de mejora si deshabilitamos la cache L1.
Potts-noncompact
Finalmente comprobamos que una versión sin compactar produce una menor utilización del BW de memoria.
unsigned int
9 512 0.44261 0.00905423
9 1024 0.334046 0.0046096
9 2048 0.291685 0.00199911
9 4096 0.272963 0.00220491
9 8192 0.265586 0.00233615
9 16384 0.264296 0.00232231
9 32768 0.265268 0.0021042
int
9 512 0.441877 0.00915216
9 1024 0.334246 0.00464585
9 2048 0.29183 0.00292973
9 4096 0.27297 0.00219631
9 8192 0.265794 0.00242677
9 16384 0.26442 0.00271465
9 32768 0.265333 0.00218667
· El código es nuevo y no fue verificado aun.
· Pierde un 10% de performance.
(si, esto es una invitación a un trabajo final)
int
es como poco. 1/4 del BW de memoria real.uint
.Table of Contents | t |
---|---|
Exposé | ESC |
Full screen slides | e |
Presenter View | p |
Source Files | s |
Slide Numbers | n |
Toggle screen blanking | b |
Show/hide slide context | c |
Notes | 2 |
Help | h |