optimization - mejorar - Técnicas para reducir la latencia de transferencia de datos de CPU a GPU
como bajar ms (3)
He estado buscando formas de reducir la latencia causada por la transferencia de datos desde la CPU y la GPU. Cuando comencé a usar CUDA noté que la transferencia de datos entre la CPU y la GPU tomó algunos segundos, pero realmente no me importó porque esto no es realmente una preocupación para los pequeños programas que he estado escribiendo. De hecho, la latencia probablemente no sea un gran problema para la gran mayoría de los programas que utilizan GPU, incluidos los videojuegos, porque aún son mucho más rápidos que si se hubieran ejecutado en la CPU.
Sin embargo, soy un poco entusiasta de HPC y comencé a preocuparme por la dirección de mis estudios cuando vi la enorme discrepancia entre los FLOPS máximos teóricos de Tianhe-I y el rendimiento real medido de LINPACK. Esto ha levantado mi preocupación sobre si estoy tomando el camino correcto de la carrera.
El uso de la memoria fija (bloqueo de página) a través del uso de la función cudaHostAlloc () es un método para reducir la latencia (bastante efectivo), pero ¿hay otras técnicas que no conozco? Y para ser claros, estoy hablando de optimizar el código, no el hardware en sí (eso es NVIDIA y los trabajos de AMD).
Solo como una pregunta secundaria, soy consciente de que Dell y HP venden los servidores de Tesla. Tengo curiosidad sobre qué tan bien una GPU aprovecha una aplicación de base de datos, donde necesitaría una lectura constante del disco duro (HDD o SSD), una operación que solo la CPU puede realizar,
Puede usar cudaMemcpyAsync() para superponer el trabajo que está haciendo en la CPU con la transferencia de memoria. Esto no reducirá la latencia de una transferencia de datos, pero podría mejorar el rendimiento general de un algoritmo. Hay algo de información al respecto en la guía CUDA C Best Practices .
Hay algunas formas de abordar la sobrecarga de comunicación de la CPU-GPU: espero que a eso se refiera con latencia y no con la latencia de la transferencia. Tenga en cuenta que, deliberadamente, utilicé el término dirección en lugar de reducir, ya que no es necesario que reduzca la latencia si puede ocultarla. También tenga en cuenta que estoy mucho más familiarizado con CUDA, por lo que a continuación solo me refiero a CUDA, pero algunas características también están disponibles en OpenCL.
Como mencionaste , la memoria bloqueada en la página tiene el mismo propósito de aumentar. Además, se puede mapear la memoria de host bloqueada en la página a la GPU, mecanismo que permite el acceso directo de los datos asignados desde el kernel GPU sin la necesidad de una transferencia de datos adicional. Este mecanismo se denomina transferencia de copia cero y es útil si los datos se leen / escriben solo una vez acompañados de una cantidad sustancial de computación y para las GPU sin memoria separada (móvil). Sin embargo, si el núcleo que accede a los datos copiados al cero no está fuertemente vinculado al cálculo y, por lo tanto, la latencia del acceso a los datos no se puede ocultar, la memoria bloqueada pero no la memoria mapeada será más eficiente. Además, si los datos no encajan en la memoria de la GPU, la copia cero seguirá funcionando.
Tenga en cuenta que una cantidad excesiva de memoria de página bloqueada puede causar una desaceleración grave en el lado de la CPU.
Al abordar el problema desde un ángulo diferente, como mencionó tkerwin, la transferencia asíncrona (mientras el hilo de la CPU habla con la GPU) es la clave para ocultar la latencia de transferencia de la CPU-GPU superponiendo el cálculo en la CPU con la transferencia. Esto se puede lograr con cudaMemcpyAsync() así como también con copia cero con ejecución asincrónica del kernel.
Uno puede llevar esto aún más lejos al usar múltiples flujos para superponer la transferencia con la ejecución del kernel. Tenga en cuenta que la programación de secuencias puede necesitar atención especial para una buena superposición; Las tarjetas Tesla y Quadro tienen un motor dual DMA que permite la transferencia de datos simultánea desde y hacia la GPU. Además, con CUDA 4.0 se hizo más fácil usar una GPU de múltiples subprocesos de CPU, por lo que en un código de CPU de subprocesos múltiples cada subproceso puede enviar sus propios datos a la GPU y lanzar kernels más fácilmente.
Finalmente, GMAC implementa un modelo de memoria compartida asimétrica para CUDA. Una de sus características más interesantes son los modelos de coherencia que proporciona, en particular la actualización lenta y continua que permite la transferencia de solo los datos modificados en la CPU de forma bloqueada.
Para más detalles, ver el siguiente artículo: Gelado et al. - Un modelo de memoria compartida asimétrica distribuida para sistemas paralelos heterogéneos .
Si la latencia es un problema, podría valer la pena analizar los intercambios que puede realizar con la arquitectura de fusión de AMD. La latencia que obtiene se minimiza drásticamente y, en algunos casos, puede ser más rápida que las transferencias de CPU desde la RAM. Sin embargo, si obtiene un golpe de rendimiento con el uso de la GPU no discreta adelgazada.