CUDA en Fortran

  • Published on
    02-Dec-2014

  • View
    700

  • Download
    2

Embed Size (px)

Transcript

<p>Diseo de Procesadores y Evaluacin de ConfiguracionesEscuela Politcnica Superior Universidad de Crdoba</p> <p>Aplicacin de tcnicas de optimizacin y paralelizacin con NVIDIA CUDA en FortranDavid Lpez Fernndez2 Ingeniera Informtica Curso 10/11</p> <p>IndiceNvidia CUDAQu es CUDA? GPU vs CPU Conceptos de la arquitectura Modelo de programacin Memoria en CUDA Sincronizacin</p> <p>ArquitecturaCPU GPU</p> <p>Aplicacin de optimizacionesLoop unswitching Strip-mining Loop distribution</p> <p>PGI Resources: CUDA Fortran Objetivos</p> <p>Conclusiones y resultados2Aplicacin de tcnicas de optimizacin y paralelizacin con NVIDIA CUDA en Fortran</p> <p>Nvidia CUDA Qu es CUDA?Concepto que trata de aprovechar las caractersticas y capacidades de cmputo de la GPU, que son:Mltiples ncleos dirigidos por un gran ancho de banda de memoria. Gran paralelismo. Optimizacin para clculos en coma flotante. Bajo coste. Ingente cantidad de transistores.Procesamiento de datos. Almacenamiento en cach. Control del flujo.</p> <p>3Aplicacin de tcnicas de optimizacin y paralelizacin con NVIDIA CUDA en Fortran</p> <p>Nvidia CUDA GPU vs CPUGPU es un procesador especializado en descargar al microprocesador del renderizado 3D de los grficos. GPU600 800 Mhz Alta especializacin Modelo circulante Optimizadas para clculos en coma flotante Media docena de procesadores de vrtices</p> <p>CPU3.8 4 Ghz Propsito general Von Neumann ...</p> <p>...</p> <p>Paralelismo inherente en aplicaciones grficas, al ser sus unidades de cmputo pxeles y vrtices.4Aplicacin de tcnicas de optimizacin y paralelizacin con NVIDIA CUDA en Fortran</p> <p>Nvidia CUDA Conceptos de la arquitectura</p> <p>5Aplicacin de tcnicas de optimizacin y paralelizacin con NVIDIA CUDA en Fortran</p> <p>Nvidia CUDA Conceptos de la arquitecturaArquitectura hardware y software para la computacin en la GPU.</p> <p>Computed Unified Device Architecture.Aplicacin de tcnicas de optimizacin y paralelizacin con NVIDIA CUDA en Fortran</p> <p>6</p> <p>Nvidia CUDA Conceptos de la arquitecturaAcceso general a la memoria DRAM. Mayor flexibilidad. Cach de datos en paralelo, chip de memoria compartida de alta velocidad.</p> <p>7Aplicacin de tcnicas de optimizacin y paralelizacin con NVIDIA CUDA en Fortran</p> <p>Nvidia CUDA Modelo de programacinEl modelo de computacin vincula el trmino de procesador SIMD con un hilo y cada multiprocesador con un bloque. Estructuras de datos en CPU deben adaptarse a otras ms simples tipo matriz o vector. Todos los procesadores SIMD ejecutan el mismo trozo de cdigo pero con diferentes datos.8Aplicacin de tcnicas de optimizacin y paralelizacin con NVIDIA CUDA en Fortran</p> <p>Nvidia CUDA Modelo de programacinGPU (device) trabaja como un coprocesador junto con la CPU (host). Una parte de una aplicacin que se ejecuta varias veces, con diferentes datos, puede ser aislada dentro de una funcin 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.9Aplicacin de tcnicas de optimizacin y paralelizacin con NVIDIA CUDA en Fortran</p> <p>Nvidia CUDA Modelo de programacinBloque. Conjunto de hilosComparten datos. Memoria compartida de rpido acceso. Puntos de sincronizacin en un kernelHilos de un bloque se suspenden para esperar al resto</p> <p>Nmero de hilos limitado.</p> <p>Hilothread ID, nmero 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).</p> <p>10Aplicacin de tcnicas de optimizacin y paralelizacin con NVIDIA CUDA en Fortran</p> <p>Nvidia CUDA Modelo de programacinGrid o rejillabloques de hilos del mismo tamao y dimensin, que ejecutan un nico kernel. Los hilos que pertenecen a bloques distintos dentro de la misma rejilla, no pueden comunicarse, sincronizarse ni compartir memoria.</p> <p>block ID, nmero 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).</p> <p>11Aplicacin de tcnicas de optimizacin y paralelizacin con NVIDIA CUDA en Fortran</p> <p>Nvidia CUDA Modelo de programacin</p> <p>12Aplicacin de tcnicas de optimizacin y paralelizacin con NVIDIA CUDA en Fortran</p> <p>Nvidia CUDA Interfaz de programacinExtendida de C. Identificadores de funciones:__device__Se ejecuta en GPU. Slo se puede llamar desde GPU.</p> <p>__global__ (kernel)Se ejecuta en GPU. Slo se puede llamar desde CPU.</p> <p>__host__Se ejecuta en CPU. Slo se puede llamar desde CPU.</p> <p>13Aplicacin de tcnicas de optimizacin y paralelizacin con NVIDIA CUDA en Fortran</p> <p>Nvidia CUDA Interfaz de programacinIdentificadores de variables__device__Declara una variable que reside en la GPU. Reside en memoria global. Su tiempo de vida es el de la aplicacin. Es accesible desde todos los hilos dentro de una rejilla y desde la CPU a travs de la runtime library de CUDA.</p> <p>__constant__Declara una variable que reside en memoria constante. Su tiempo de vida es el de la aplicacin. Accesible desde todos los hilos dentro de una rejilla y desde la CPU a travs de la runtime library.14Aplicacin de tcnicas de optimizacin y paralelizacin con NVIDIA CUDA en Fortran</p> <p>Nvidia CUDA Interfaz de programacinIdentificadores 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.</p> <p>15Aplicacin de tcnicas de optimizacin y paralelizacin con NVIDIA CUDA en Fortran</p> <p>Nvidia CUDA Memoria en CUDACada dispositivo CUDA tiene diferentes tipos de memorias:</p> <p>16Aplicacin de tcnicas de optimizacin y paralelizacin con NVIDIA CUDA en Fortran</p> <p>Nvidia CUDA Memoria en CUDACada 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.</p> <p>La CPU slo puede leer o escribir en las memorias global y contantes por medio de la API CUDA.Tiempo de acceso alto / Ancho de banda limitado</p> <p>La memoria constante slo permite acceso para lectura por parte de la GPU y esta provista de un acceso a sus datos mucho ms rpido que la memoria global.Aplicacin de tcnicas de optimizacin y paralelizacin con NVIDIA CUDA en Fortran</p> <p>17</p> <p>Nvidia CUDA Memoria en CUDARegistros y memorias compartidas:Acceso muy rpido y paralelo. Cada hilo tiene sus propios registros. Una funcin de kernel usa registros para acceder rpidamente a variables privadas de cada hilo, evitando congestin en memoria global. Cooperacin entre hilos de un mismo bloque es posible gracias a las memorias compartidas.Declaracin de variables Variables (escalares) Arrays __device__ __shared__ int sharedVar; __device__ int globalVar; __device__ __constant__ int constVar; Memoria Registro Global Global Constante Scope Thread Thread Grid Grid Lifetime Kernel Kernel Kernel Aplicacin Aplicacin</p> <p>Compartida Block</p> <p>18Aplicacin de tcnicas de optimizacin y paralelizacin con NVIDIA CUDA en Fortran</p> <p>Nvidia en CUDA Memoria en CUDASi un kernel declara una variable cuyo mbito es un hilo y es ejecutada por un milln de hilos, se crear un milln 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**)&amp;puntero, elems(sizeof(float)); cudaFree(puntero);</p> <p>19Aplicacin de tcnicas de optimizacin y paralelizacin con NVIDIA CUDA en Fortran</p> <p>Nvidia en CUDA Memoria en CUDAcudaMemcpy(d_mem, h_mem, memSize, cudaMemcpyHostToDevice); d_mem variable en memoria global en la GPU donde se copiarn los datos que se transfieren desde la CPU. h_mem variable en el host o CPU cuyo contenido se copiar a la memoria global de la GPU. MemSize tamao de memoria que se transfiere cudaMemcpy(h_mem, d_mem, memSize, cudaMemcpyDeviceToHost); cudaMemcpy(d_mem1, d_mem2, memSize, cudaMemcpyDeviceToDevice);</p> <p>Copia de la GPU a la GPU</p> <p>20Aplicacin de tcnicas de optimizacin y paralelizacin con NVIDIA CUDA en Fortran</p> <p>Nvidia en CUDA Sincronizacin de hilos</p> <p>__syncthreads();Un hilo esperar en el punto de la funcin kernel donde se haya realizado la llamada a que todos los hilos finalicen.</p> <p>cudaThreadSynchronize();Similar a la anterior pero controlando la sincronizacin con la CPU.</p> <p>21Aplicacin de tcnicas de optimizacin y paralelizacin con NVIDIA CUDA en Fortran</p> <p>PGI Resources CUDA Fortran</p> <p>Permite la programacin CUDA directamente en Fortranhttp://www.pgroup.com/resources/cudafortran.htm</p> <p>ExtensionesDeclaracin de variables que residen en la memoria de la GPU. Asignacin dinmica de datos en la memoria de la GPU. Copiar datos desde la memoria del host a la memoria de la GPU, y al revs. Invocacin de subrutinas GPU desde el host.22Aplicacin de tcnicas de optimizacin y paralelizacin con NVIDIA CUDA en Fortran</p> <p>PGI Resources CUDA FortranDeviced, shared y constant!AinstantiatedinGPUmemory real,device::A(M,N) !BisallocatableinGPUmemory real,device,allocatable::B(:,:) ... !AllocateBinGPUmemory allocate(B(size(A,1),size(A,2)),stat=istat) !CopydatafromA(host)toAdev(GPU) Adev=A !CopydatafromAdev(GPU)toA(host) A=AdevAplicacin de tcnicas de optimizacin y paralelizacin con NVIDIA CUDA en Fortran</p> <p>23</p> <p>PGI Resources API</p> <p>Usando la APIusecudafor 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)24Aplicacin de tcnicas de optimizacin y paralelizacin con NVIDIA CUDA en Fortran</p> <p>PGI Resources Subrutina kernel</p> <p>Invocacin del kernel (slo uno a la vez por GPU)Su cdigo se ejecutar con diferentes datos en varios hiloscallmmul_kernel(parmetros) dimBlock=256 if(N/dimBlock==0)then dimGrid=N/dimBlock else dimGrid=N/dimBlock+1 endif callkernel(Adev,Bdev,Cdev,N,z,x) 25Aplicacin de tcnicas de optimizacin y paralelizacin con NVIDIA CUDA en Fortran</p> <p>PGI Resources Subrutina kernelattributes(global)subroutinekernel(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</p>