Uso constante de la memoria en el código CUDA

Question

No me las puedo imaginar, ¿cuál es la mejor manera de garantizar que la memoria utilizada en mi núcleo sea constante? Hay una pregunta similar en http://stackoverflow...r-pleasant-way. Estoy trabajando con GTX580 y compilando solo para capacidad 2.0. Mi kernel parece

__global__ Foo(const int *src, float *result) {...} 

ejecuto el siguiente código en el host:

cudaMalloc(src, size); 
cudaMemcpy(src, hostSrc, size, cudaMemcpyHostToDevice); 
Foo<<<...>>>(src, result); 

el camino alternativo es añadir

__constant__ src[size]; 

a .CU de archivos, eliminar src puntero desde el kernel y ejecuta

cudaMemcpyToSymbol("src", hostSrc, size, 0, cudaMemcpyHostToDevice); 
Foo<<<...>>>(result); 

¿Son estas dos formas equivalentes o la primera no garantiza el uso de la memoria constante en lugar de la memoria global? tamaño cambia dinámicamente por lo que la segunda forma no es útil en mi caso.

Answer 1

La segunda forma es garantizar que la matriz se compile en la memoria constante CUDA y se acceda correctamente a través de la memoria caché constante. Pero debes preguntarte cómo se accederá al contenido de esa matriz dentro de un bloque de hilos. Si cada subproceso tendrá acceso uniformemente a la matriz, habrá una ventaja de rendimiento al usar memoria constante, porque hay un mecanismo de difusión de la memoria caché constante (también ahorra ancho de banda de memoria global porque la memoria constante se almacena en la memoria DRAM y la memoria caché) reduce el recuento de transacciones de DRAM). Pero si el acceso es aleatorio, puede haber una serialización de acceso a la memoria local que afectará negativamente el rendimiento.

Las cosas típicas que podrían ser buenos ajustes para la memoria __constant__ serían los coeficientes del modelo, los pesos y otros valores constantes que deben establecerse en tiempo de ejecución. En las GPU de Fermi, la lista de argumentos del núcleo se almacena en la memoria constante, por ejemplo. Pero si el contenido es de acceso no uniforme y el tipo o tamaño de los miembros no es constante de llamada a llamada, entonces es preferible la memoria global normal.

También tenga en cuenta que existe un límite de 64 kb de memoria constante por contexto de GPU, por lo que no es práctico almacenar grandes cantidades de datos en la memoria constante. Si necesita una gran cantidad de almacenamiento de solo lectura con caché, podría valer la pena intentar vincular los datos a una textura y ver cómo es el rendimiento. En las tarjetas pre-Fermi, generalmente produce una ganancia de rendimiento útil, en Fermi los resultados pueden ser menos predecibles en comparación con la memoria global debido a la mejora del diseño del caché en esa arquitectura.

Answer 2

El primer método garantizará que la memoria esté constante dentro de la función Foo. Los dos no son equivalentes, el segundo lo garantiza allí después de su inicialización. Si necesita algo más dinámico que lo que necesita para usar algo similar a la primera manera.