Diseño de Procesadores y Evaluación de ConfiguracionesEscuela Politécnica SuperiorUniversidad de Córdoba
David López Fernández2º Ingeniería Informática
Curso 10/11
Aplicación de técnicas de optimización y paralelización
con NVIDIA CUDA en Fortran
2
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Nvidia CUDA¿Qué es CUDA?GPU vs CPUConceptos de la arquitecturaModelo de programaciónMemoria en CUDASincronización
PGI Resources: CUDA Fortran
Objetivos
Indice
ArquitecturaCPUGPU
Aplicación de optimizaciones
Loop unswitchingStrip-miningLoop distribution
Conclusiones y resultados
3
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Nvidia CUDA → ¿Qué es CUDA?
Concepto que trata de aprovechar las características y capacidades de cómputo de la GPU, que son:
Múltiples núcleos dirigidos por un gran ancho de banda de memoria.Gran paralelismo.Optimización para cálculos en coma flotante.Bajo coste.Ingente cantidad de transistores.
Procesamiento de datos.Almacenamiento en caché.Control del flujo.
4
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Nvidia CUDA → GPU vs CPU
GPU es un procesador especializado en descargar al microprocesador del renderizado 3D de los gráficos.
Paralelismo inherente en aplicaciones gráficas, al ser sus unidades de cómputo píxeles y vértices.
GPU CPU600 – 800 Mhz 3.8 – 4 Ghz
Alta especialización Propósito general
Modelo circulante Von Neumann
Optimizadas para cálculos en coma
flotante
...
Media docena de procesadores de
vértices
...
5
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Nvidia CUDA → Conceptos de la arquitectura
6
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Arquitectura hardware y software para la computación en la GPU.
Computed Unified Device Architecture.
Nvidia CUDA → Conceptos de la arquitectura
7
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Nvidia CUDA → Conceptos de la arquitectura
Acceso general a la memoria DRAM. Mayor flexibilidad.
Caché de datos en paralelo, chip de memoria compartida de alta velocidad.
8
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
El modelo de computación vincula el término de procesador SIMD con un hilo y cada multiprocesador con un bloque.Estructuras de datos en CPU deben adaptarse a otras más simples tipo matriz o vector.Todos los procesadores SIMD ejecutan el mismo trozo de código pero con diferentes datos.
Nvidia CUDA → Modelo de programación
9
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Nvidia CUDA → Modelo de programación
GPU (device) trabaja como un coprocesador junto con la CPU (host).
Una parte de una aplicación que se ejecuta varias veces, con diferentes datos, puede ser aislada dentro de una función que se ejecuta en varios hilos diferentes en la GPU (kernel).
CPU y GPU mantienen su propia DRAMHost memory.Device memory.Se pueden copiar datos de una a otra.
10
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Nvidia CUDA → Modelo de programación
Bloque. Conjunto de hilosComparten datos. Memoria compartida de rápido acceso.Puntos de sincronización en un kernel
Hilos de un bloque se suspenden para esperar al resto
Número de hilos limitado.
Hilothread ID, número de hilo dentro del bloque. Para un bloque de dos dimensiones (Dx, Dy) el thread ID de un hilo de índice (x, y) es (x+yDx) y para un bloque de tres dimensiones (Dx, Dy, Dz), el thread ID de un hilo de índice (x, y, z) es (x + yDx + zDxDy).
11
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Nvidia CUDA → Modelo de programación
Grid o rejillabloques de hilos del mismo tamaño y dimensión, que ejecutan un único kernel.Los hilos que pertenecen a bloques distintos dentro de la misma rejilla, no pueden comunicarse, sincronizarse ni compartir memoria.
block ID, número de bloque dentro del gridPara una rejilla de 2 dimensiones (Dx, Dy), el block ID de un bloque de índices (x, y) es (x +yDy).
12
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Nvidia CUDA → Modelo de programación
13
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Nvidia CUDA → Interfaz de programación
Extendida de C.
Identificadores de funciones:__device__
Se ejecuta en GPU.Sólo se puede llamar desde GPU.
__global__ (kernel)Se ejecuta en GPU.Sólo se puede llamar desde CPU.
__host__Se ejecuta en CPU.Sólo se puede llamar desde CPU.
14
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Nvidia CUDA → Interfaz de programación
Identificadores de variables__device__
Declara una variable que reside en la GPU.Reside en memoria global.Su tiempo de vida es el de la aplicación.Es accesible desde todos los hilos dentro de una rejilla y desde la CPU a través de la runtime library de CUDA.
__constant__ Declara una variable que reside en memoria constante.Su tiempo de vida es el de la aplicación.Accesible desde todos los hilos dentro de una rejilla y desde la CPU a través de la runtime library.
15
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Nvidia CUDA → Interfaz de programación
Identificadores de variables__shared__ declara una variable que reside en la memoria compartida de un bloque de hilos.
Su tiempo de vida es el del bloque.Solamente es accesible desde todos los hilos del bloque.
16
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Nvidia CUDA → Memoria en CUDA
Cada dispositivo CUDA tiene diferentes tipos de memorias:
17
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Nvidia CUDA → Memoria en CUDA
Cada hilo puede:Leer/escribir en los registros de hilo.Leer/escribir en la memoria local del hilo.Leer/escribir en la memoria compartida del bloque.Leer/escribir en la memoria global de la rejilla.Leer datos de la memoria constante.
La CPU sólo puede leer o escribir en las memorias global y contantes por medio de la API CUDA.
Tiempo de acceso alto / Ancho de banda limitado
La memoria constante sólo permite acceso para lectura por parte de la GPU y esta provista de un acceso a sus datos mucho más rápido que la memoria global.
18
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Nvidia CUDA → Memoria en CUDA
Registros y memorias compartidas:Acceso muy rápido y paralelo. Cada hilo tiene sus propios registros. Una función de kernel usa registros para acceder rápidamente a variables privadas de cada hilo, evitando congestión en memoria global.Cooperación entre hilos de un mismo bloque es posible gracias a las memorias compartidas.
Declaración de variables Memoria Scope Lifetime
Variables (escalares) Registro Thread Kernel
Arrays Global Thread Kernel
__device__ __shared__ int sharedVar; Compartida Block Kernel
__device__ int globalVar; Global Grid Aplicación
__device__ __constant__ int constVar; Constante Grid Aplicación
19
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Nvidia en CUDA → Memoria en CUDA
Si un kernel declara una variable cuyo ámbito es un hilo y es ejecutada por un millón de hilos, se creará un millón de copias de la variable, que cada hilo inicializará y usará su propia copia de la variable.
Antes de transferir datos entre CPU y GPU hay que reservar memoria:
cudaMalloc((void**)&puntero, elems(sizeof(float));cudaFree(puntero);
20
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Nvidia en CUDA → Memoria en CUDA
cudaMemcpy(d_mem, h_mem, memSize, cudaMemcpyHostToDevice);d_mem → variable en memoria global en la GPU donde secopiarán los datos que se transfieren desde la CPU.h_mem → variable en el host o CPU cuyo contenido se copiará a lamemoria global de la GPU.MemSize → tamaño de memoria que se transfiere
cudaMemcpy(h_mem, d_mem, memSize, cudaMemcpyDeviceToHost);
cudaMemcpy(d_mem1, d_mem2, memSize, cudaMemcpyDeviceToDevice);
Copia de la GPU a la GPU
21
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Nvidia en CUDA → Sincronización de hilos
__syncthreads();Un hilo esperará en el punto de la función kernel donde se haya realizado la llamada a que todos los hilos finalicen.
cudaThreadSynchronize();Similar a la anterior pero controlando la sincronización con la CPU.
22
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
PGI Resources → CUDA Fortran
Permite la programación CUDA directamente en Fortranhttp://www.pgroup.com/resources/cudafortran.htm
ExtensionesDeclaración de variables que residen en la memoria de la GPU.Asignación dinámica de datos en la memoria de la GPU.Copiar datos desde la memoria del host a la memoria de la GPU, y al revés.Invocación de subrutinas GPU desde el host.
23
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
PGI Resources → CUDA Fortran
Deviced, shared y constant
! A instantiated in GPU memory
real,device :: A(M,N)
! B is allocatable in GPU memory
real,device,allocatable :: B(:,:)
...
! Allocate B in GPU memory
allocate(B(size(A,1),size(A,2)),stat=istat)
! Copy data from A (host) to Adev (GPU)
Adev = A
! Copy data from Adev (GPU) to A (host)
A = Adev
24
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
PGI Resources → API
Usando la API
use cudafor
real, allocatable, device :: a(:)
real :: b(10), b2(2), c(10)
...
istat = cudaMalloc( a, 10 )
istat = cudaMemcpy( a, b, 10 )
istat = cudaMemcpy( a(2), b2, 2 )
istat = cudaMemcpy( c, a, 10 )
istat = cudaFree( a )
25
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
PGI Resources → Subrutina kernel
Invocación del kernel (sólo uno a la vez por GPU)Su código se ejecutará con diferentes datos en varios hilos
call mmul_kernel<<<dimGrid,dimBlock>>>( parámetros )
dimBlock = 256
if(N/dimBlock == 0) then
dimGrid = N/dimBlock
else
dimGrid = N/dimBlock + 1
end if
call kernel<<<dimGrid, dimBlock>>>( Adev, Bdev, Cdev, N, z, x)
26
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
PGI Resources → Subrutina kernel
attributes(global) subroutine kernel( A, B, C, N, z, x )
real,device :: A(N), B(N), C(N)
integer, value :: N, z, x
integer :: i
i = (blockidx%x1)*256 + threadidx%x
if( i <= N .AND. x<7) then
A(i) = A(i) + z
B(i) = A(i) + C(i)
else if( i <= N .AND. x>=8) then
A(i) = A(i) + z
B(i) = A(i) C(i)
end if
End subroutine
27
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
PGI Resources → Uso del compilador
Uso del compilador de PGI Resources
$ pgfortran fast matmul.CUF
$ ./a.out
-Mcuda enlaza ficheros objeto. -g -Mcuda emula el programa CUDA en el Host.-fast incluye f-expensive-optimizations
O2 y O3
28
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Objetivos
Implementar en Fortran optimizaciones para aumento de localidadLoop unswitchingLoop interchangeStrip-miningLoop tilingLoop distribution
Estudiar las características de CUDA
Optimizaciones sobre esta arquitectura, en FortranPGI Resources, The Portland Group
Comparar resultados
29
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Descripción de la arquitectura → CPU
Intel Core 2 DUO T75002 CPU cores
2.20 Ghz
L164 setsAsociatividad 832 Kb
L24096 setsAsociatividad 164096 Mb
30
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Descripción de la arquitectura → GPU
Nvidia GeForce 8600M GS 256 Mb
Procesadores tipo “stream” 16
Reloj central (MHz) 600
Reloj de las unidades de sombreado (MHz)
1200
Reloj de la memoria (MHz) 700
Cantidad de memoria 512Mb
Interfaz de memoria 128-bit
Tipo de memoria GDDR3, GDDR2 Transistores 298 millones
31
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Desarrollo → Loop unswitching
N = 3000000
DO i = 2,N,1
A(i) = A(i) + z
PRINT *,' A[',i, '] = ', A(i)
IF (x<7) THEN
B(i) = A(i) * C(i)
PRINT *,'x<7: B[',i,'] = ', B(i)
ELSE
B(i) = A(i1) * B(i1)
PRINT *,'x>=7: B[',i,'] = ', B(i)
END IF
END DO
32
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Desarrollo → Loop unswitching
Sin aplicarlo en CUDAInvocación kernel (N=3000000)
dimBlock = 256
if(N/dimBlock == 0) then
dimGrid = N/dimBlock
else
dimGrid = N/dimBlock + 1
end if
call kernel<<<dimGrid, dimBlock>>>( Adev,Bdev,Cdev,N,z,x)
r = cudathreadsynchronize()
33
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Desarrollo → Loop unswitching
Sin aplicarlo en CUDASubrutina kernel
attributes(global) subroutine kernel( A, B, C, N, z, x )
real,device :: A(N), B(N), C(N)
integer, value :: N, z, x
integer :: i
i = (blockidx%x1)*256 + threadidx%x
if( i <= N .AND. x<7) then
A(i) = A(i) + z
B(i) = A(i) + C(i)
else if( i <= N .AND. x>=8) then
A(i) = A(i) + z
B(i) = A(i) C(i)
end if
End subroutine
34
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Desarrollo → Loop unswitching
Aplicación en CUDA2 funciones kernel, condición antes de la llamada
dimBlock = 256
if(N/dimBlock == 0) then
dimGrid = N/dimBlock
else
dimGrid = N/dimBlock + 1
end if
if (x<7) then
call kernel_a<<<dimGrid, dimBlock>>>( Adev,Bdev,Cdev,N,z)
else
call kernel_b<<<dimGrid, dimBlock>>>( Adev,Bdev,Cdev,N,z)
end if
r = cudathreadsynchronize()
35
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Desarrollo → Loop unswitching
Aplicación en CUDASubrutina kernel_a (x<7)
attributes(global) subroutine kernel_a( A, B, C, N, z )
real,device :: A(N), B(N), C(N)
integer, value :: N, z
integer :: i
i = (blockidx%x1)*256 + threadidx%x
if( i <= N ) then
A(i) = A(i) + z
B(i) = A(i) + C(i)
end if
End subroutine
36
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Desarrollo → Loop unswitching
Aplicación en CUDASubrutina kernel_b (x<=7)
attributes(global) subroutine kernel_b( A, B, C, N, z )
real,device :: A(N), B(N), C(N)
integer, value :: N, z
integer :: i
i = (blockidx%x1)*256 + threadidx%x
if( i <= N ) then
A(i) = A(i) + z
B(i) = A(i) C(i)
end if
End subroutine
37
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Desarrollo → Loop unswitching
Resultados
CUDA > NO LOOP UNSWITCHING
Kernel time excluding data xfer: 7.6258183E-03 seconds
Megaflops excluding data xfer: 3.9340054E+08
Total time including data xfer: 0.1741080 seconds
Megaflops including data xfer: 1.7230688E+07
CUDA > LOOP UNSWITCHING
Kernel time excluding data xfer: 7.5747967E-03 seconds
Megaflops excluding data xfer: 3.9605037E+08
Total time including data xfer: 0.1214762 seconds
Megaflops including data xfer: 2.4696210E+07
38
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Desarrollo → Strip-mining
N = 300000
DO i = 1,N
A(i) = A(i) * 3 / 32
END DO
CUDA asocia el concepto de bloque al de multiprocesador.Porción del bucle -> hilos por bloque.
39
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Desarrollo → Strip-mining
Aplicación en CUDAInvocación kernel (N=300000)
dimBlock = 64
if(N/dimBlock == 0) then
dimGrid = N/dimBlock
else
dimGrid = N/dimBlock + 1
end if
call kernel<<<dimGrid, dimBlock>>>( Adev, N)
r = cudathreadsynchronize()
40
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Desarrollo → Strip-mining
Aplicación en CUDASubrutina kernel
attributes(global) subroutine kernel( A, N )
real,device :: A(N)
integer, value :: N
integer :: i
i = (blockidx%x1)*64 + threadidx%x
if( i <= N ) A(i) = A(i) * 3 /(32)
end subroutine
41
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Desarrollo → Strip-mining
Resultados
CUDA > STRIPMINING
Kernel time excluding data xfer: 4.9114227E-05 seconds
Megaflops excluding data xfer: 6.1082301E+09
Total time including data xfer: 8.4445953E-02 seconds
Megaflops including data xfer: 3552580
HOST > NO STRIPMINING
Host time without strip-mining: 5.8197975E-04 seconds
Megaflops: 515481856
42
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Desarrollo → Loop distribution
N = 300000
DO i = 1, N
AH(i) = AH(i) + c
XH(i) = XH(i)*7 + XH(i) + AH(i)
END DO
43
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Desarrollo → Loop distribution
Sin aplicarlo en CUDAInvocación kernel (N=300000)
dimBlock = 256
if(N/dimBlock == 0) then
dimGrid = N/dimBlock
else
dimGrid = N/dimBlock + 1
end if
call kernel<<<dimGrid, dimBlock>>>( Adev, Xdev, N, c)
r = cudathreadsynchronize()
44
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Desarrollo → Loop distribution
Sin aplicarlo en CUDASubrutina kernel
attributes(global) subroutine kernel( A, X, N, c )
real,device :: A(N), X(N)
integer, value :: N, c
integer :: i
i = (blockidx%x1)*256 + threadidx%x
if( i <= N ) then
A(i) = A(i) + c
X(i) = X(i)*7 + X(i) + A(i)
end if
end subroutine
45
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Desarrollo → Loop distribution
Aplicación en CUDA2 funciones kernel, una para cada bucle (fisión)
dimBlock = 256
if(N/dimBlock == 0) then
dimGrid = N/dimBlock
else
dimGrid = N/dimBlock + 1
end if
call kernel_loop1<<<dimGrid, dimBlock>>>( Adev, N, c)
call kernel_loop2<<<dimGrid, dimBlock>>>( Adev, Xdev, N, c)
r = cudathreadsynchronize()
46
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Desarrollo → Loop distribution
Aplicación en CUDASubrutina kernel_loop1
attributes(global) subroutine kernel_loop1( A, N, c )
real,device :: A(N)
integer, value :: N, c
integer :: i
i = (blockidx%x1)*256 + threadidx%x
if( i <= N ) then
A(i) = A(i) + c
end if
end subroutine
47
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Desarrollo → Loop distribution
Aplicación en CUDASubrutina kernel_loop2
attributes(global) subroutine kernel_loop2( A, X, N, c )
real,device :: X(N), A(N)
integer, value :: N, c
integer :: i
i = (blockidx%x1)*256 + threadidx%x
if( i <= N ) then
X(i) = X(i)*7 + X(i) + A(i)
end if
end subroutine
48
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Desarrollo → Loop distribution
Resultados
CUDA > NO LOOP DISTRIBUTION
Kernel time excluding data xfer: 8.5496902E-04 seconds
Megaflops excluding data xfer: 3.5089107E+08
Total time including data xfer: 9.2410088E-02 seconds
Megaflops including data xfer: 3246410.
CUDA > LOOP DISTRIBUTION
Kernel time excluding data xfer: 9.4509125E-04 seconds
Megaflops excluding data xfer: 3.1743072E+08
Total time including data xfer: 9.3610287E-03 seconds
Megaflops including data xfer: 3.2047866E+07
49
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Desarrollo → Otras optimizaciones
Loop unrollingYa lo realiza CUDA.
Loop tilingGeneralización multidimensional de strip-mining. Ya lo realiza CUDA, no tiene sentido.
Optimizaciones en las estructuras de datosEn GPU no hay caché como en la CPU, pero hay jerarquía de memoria (memoria constante, compartida, global, y registros)
No hay forma? (conocida) de medir fallos de caché en la GPU, porque las librerías PAPI son para CPU.
50
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Conclusiones y resultados
Transformación No aplicada en HOST Aplicada en HOST No aplicada enCUDA
Aplicada enCUDA
Fallos decaché
Tiempo de CPU Fallos decaché
Tiempo de CPU Tiempo de CPU Tiempo de CPU
Loop unswitching 681339 2.27661E-02seg.
680803 1.95629E-02seg.
7.62581E-03seg.
7.57479E-03 seg.
51
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Conclusiones y resultados
Transformación No aplicada en HOST Aplicada en HOST No aplicada enCUDA
Aplicada enCUDA
Fallos decaché
Tiempo de CPU Fallos decaché
Tiempo de CPU Tiempo de CPU Tiempo de CPU
Strip-mining No apl. No apl. 19745 5.81979E-04seg.
No es posible,CUDA lo realizainherente.
4.91142E-05 seg.
52
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Conclusiones y resultados
Transformación No aplicada en HOST Aplicada en HOST No aplicada enCUDA
Aplicada enCUDA
Fallos decaché
Tiempo de CPU Fallos decaché
Tiempo de CPU Tiempo de CPU Tiempo de CPU
Loop distribution 39092 1.45506E-03seg.
57699. 2.12907E-03seg.
8.54969E-04seg.
9.45091E-04 seg.
53
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Bibliografía
[1] Introduction to CUDA Fortran. PGI insider. Technical News from The Portland Group.
http://www.pgroup.com/lit/articles/insider/v1n3a2.htm
[2] Nvidia CUDA. Compute Unified Device Architecture. Programming Guide. V2.0. 06/07/2008.http://developer.download.nvidia.com/compute/cuda/2_0/docs/NVIDIA_CUDA_Programming_Gui de_2.0.pdf
[3] Optimizaciones para el aumento de la localidad. Tema 5. Apuntes de la asignatura Diseño de Procesadores y Evaluación deComputadores. Universidad de Córdoba.
[4] Miles Douglas. Part 2: Introduction to CUDA Fortran. GPU Programming with PGI CUDA Fortran and the PGI AcceleratorProgramming Model. The Portland Group.
http://www.developers-source.de/devsource/projects/devsource/media/PGI-CUDA-Tutorial.pdf
[5] CUDA Fortran. Programming Guide and Reference. Release 2011. The Portland Group.
http:// www.pgroup.com/doc/pgicudaforug.pdf
54
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Bibliografía
[6] PGI Server , PGI Workstation Installation Guide Release 2011. The Portland Group.
http://www.pgroup.com/doc/pgiwsinstall.pdf
[7] Fática, Massimiliano. High Performance Computing with nVidia CUDA. Cuda Toolkit.
[8] Scalable Parallel Programming with CUDA, in ACM Queue, VOL 6, No. 2 (March/April 2008), © ACM, 2008
[9] Rondan, Andrés. CUDA, Modelo de Programación. Departamento de Informática y Sistemas Universidad de Murcia.
http://dis.um.es/~domingo/apuntes/AlgProPar/0910/exposicion1/RondanAndres-CUDA.pdf
[10] David Kirk/NVIDIA, Wen- mei Hwu. Manual CUDA: Chapter 4 - Memory CUDA.
http://courses.ece.illinois.edu/ece498/al/textbook/Chapter4-CudaMemoryModel.pdf.
55
Aplicación de técnicas de optimización y paralelización con NVIDIA CUDA en Fortran
Preguntas
Top Related