memoria física en dispositivos AMD: local vs privado

Estoy escribiendo un algoritmo en OpenCL en el que necesitaría que cada unidad de trabajo recuerde una porción justa de datos, digamos algo entre long[70] y long[200] más o menos por kernel.memoria física en dispositivos AMD: local vs privado

Los dispositivos AMD recientes tienen memoria 32 KiB __local, que es (para la cantidad dada de datos por kernel) suficiente para almacenar la información de 20-58 unidades de trabajo. Sin embargo, por lo que entiendo de la arquitectura (y especialmente de this drawing), cada núcleo de sombreador también tiene una cantidad dedicada de memoria privada. Sin embargo, no encuentro su tamaño.

¿Alguien puede decirme cómo saber cuánta memoria privada tiene cada kernel?

Tengo mucha curiosidad sobre el HD7970, ya que planeo comprar algunos de estos pronto.

Editar: Problema resuelto, la respuesta es here en el apéndice D.

Fuente

2012-02-17 user1111929

No creo que la memoria privada esté dedicada por núcleo: se correlaciona con el archivo de registro, que es por recurso de unidad de cómputo. Cada elemento de trabajo obtiene registros asignados desde el archivo de registro de la unidad de cómputo, cuántos son necesarios determina el número de frentes de onda en vuelo en cualquier instante dado. – talonmies

Del famoso dibujo visto en todas partes http://www.codeproject.com/KB/showcase/Memory-Spaces/image001.jpg Llegué a la conclusión de que la memoria privada es físicamente diferente de la memoria local __, ¿no? – user1111929

Sí, son físicamente diferentes. La memoria privada se asigna al archivo de registro de la unidad informática, la memoria local para calcular la memoria compartida a nivel de unidad en la mayoría de los dispositivos AMD modernos. Algunas primeras GPU compatibles con OpenCL no tenían en la memoria compartida, y la memoria local era solo SDRAM. Tampoco es por núcleo, y cuánto usa por artículo de trabajo para grupos privados y por grupo de trabajo para efectos locales el número de frentes de onda concurrentes que se ejecutan por unidad de cómputo. – talonmies

La respuesta fue dada por el usuario talonmies en los comentarios, así que lo escribiré en una nueva respuesta aquí para cerrar la pregunta.

Estos valores se pueden encontrar en el Apéndice D de la APLICACIÓN DE AMD Guía de programación de OpenCL http://developer.amd.com/sdks/amdappsdk/assets/amd_accelerated_parallel_processing_opencl_programming_guide.pdf (existe un documento similar para nVidia). Aparentemente, un registro es de 128 bits (4x32) para dispositivos AMD y hay 16384 registros para todos los dispositivos modernos de alta gama, por lo que es un notable 256 KB por unidad de cómputo.

Fuente

2012-03-01 12:16:47 user1111929

Creo que busca la memoria __local. Eso es a lo que se refiere 32KB de almacenamiento de datos local. No creo que pueda sondear el dispositivo para obtener la cantidad de memoria privada.

Puede pasar una referencia NULL long * cl_mem para asignar la memoria. Creo que es mejor usar una cantidad estática de memoria por WI. Suponiendo que se requerirá mucho [200] para cada elemento de trabajo, usaría el siguiente código. También sería una buena idea dividir el trabajo en grupos que tienen los mismos (o similares) requisitos de memoria, para aprovechar al máximo la memoria LDS.

void __kernel(__local long* localMem, const int localMemPerItem 
     //more args... 
     ) 
{ 
    //host has 'passed' localMemPerItem*get_local_size() long values in as locamMem 
    //this work item has access to all of it, but can choose to restrict 
    //itself to only the portion it needs. 
    //work group size will be limited to CL_DEVICE_LOCAL_MEM_SIZE/(8*localMemPerItem) 
    int startIndex=localMemPerItem*get_local_id(0); 
    //use localMem[startIndex+ ... ] 
}

Fuente

2012-02-17 18:41:18 mfa

No puede sondearlo, pero ¿existe? Del famoso dibujo visto en todas partes http://www.codeproject.com/KB/showcase/Memory-Spaces/image001.jpg Supuse que hay un conjunto físicamente separado de registros privados en cada unidad de trabajo. ¿No? Esperaba que de alguna manera fuera mejor que una limitación CL_DEVICE_LOCAL_MEM_SIZE/(8 * localMemPerItem), ya que deja aproximadamente la mitad de los núcleos sin usar. Acceder a la memoria global probablemente sea demasiado lento, aunque solo esté incrementando un contador. – user1111929

Encontré algo más de información acerca de los tamaños de registro de ciprés, caimán y fermi aquí: http://www.realworldtech.com/page.cfm?ArticleID=RWT121410213827&p=11 Deberías poder modificar algunas vars privadas de tamaño decente en ese tamaño . Creo que el LDS seguirá siendo tu mejor opción. – mfa

Para responder qué tan grande es registrar el archivo en una tarjeta de la serie 79xx, desde su base en la arquitectura GCN es de 64 KB de acuerdo con la imagen en AnandTech: http://www.anandtech.com/print/5261

Para responder a su pregunta de cómo averiguar cómo mucha memoria usa cada kernel ... puedes mirar ejecutar AMD APP Profiler en tu kernel, te dice en la sección de ocupación de kernel cuánto espacio utiliza el kernel.

Fuente

2012-02-20 14:50:24 kiranputtur

¿De verdad? Eso es raro. Pensé que había encontrado la respuesta, pero es diferente. En la guía de programación de AMD OpenCL http://developer.amd.com/sdks/amdappsdk/assets/amd_accelerated_parallel_processing_opencl_programming_guide.pdf en el Apéndice D, está el tamaño total del archivo de registro, y está listado como 256 KB para todos los dispositivos modernos. ¿Cuál es correcto ahora? : S – user1111929

Creo que ambos son correctos. Tal como lo entiendo, en la arquitectura de GCN, una unidad SIMD tiene un archivo de registro de 64 kb, y hay 4 unidades SIMD por unidad de cálculo, es decir. 4 * 64kb = 256kb de archivo de registro total por unidad de cómputo. – talonmies

memoria física en dispositivos AMD: local vs privado

Respuesta

Cuestiones relacionadas