c++ - mundo - ¿Cómo y cuándo debería usar el puntero lanzado con la API de cuda?
hola mundo en cuda (3)
Aquí hay una explicación sobre el puntero inclinado y el relleno en cuda.
Memoria lineal vs memoria acolchada
Primero, comencemos con la razón de la existencia de memoria no lineal. Al asignar memoria con cudaMalloc, el resultado es como una asignación con malloc, tenemos un trozo de memoria contigua del tamaño especificado y podemos poner todo lo que queramos en él. Si queremos asignar un vector de 10000 flotantes, simplemente hacemos:
float* myVector;
cudaMalloc(&myVector,10000*sizeof(float));
y luego acceda al elemento ith de myVector por indexación clásica:
float element = myVector[i];
y si queremos acceder al siguiente elemento, solo hacemos:
float next_element = myvector[i+1];
Funciona muy bien porque acceder a un elemento justo al lado del primero es (por razones que no conozco y no deseo ser por el momento) barato.
Las cosas se vuelven un poco diferentes cuando usamos nuestra memoria como una matriz 2D. Digamos que nuestro vector flotante 10000 es, de hecho, una matriz de 100x100. Podemos asignarlo usando la misma función cudaMalloc, y si queremos leer la i-ésima fila, hacemos:
float* myArray;
cudaMalloc(&myArray,10000*sizeof(float));
int row[100]; // number of columns
for (int j=0; j<100; ++j)
row[j] = myArray[i*100+j];
Alineación de palabras
Por lo tanto, debemos leer la memoria de myArray + 100 * i en myArray + 101 * i-1. La cantidad de operaciones de acceso a la memoria dependerá de la cantidad de palabras de memoria que tome esta fila. El número de bytes en una palabra de memoria depende de la implementación. Para minimizar el número de accesos a la memoria al leer una sola fila, debemos asegurarnos de que comenzamos la fila al comienzo de una palabra, por lo tanto, debemos rellenar la memoria de cada fila hasta el comienzo de una nueva.
Conflictos bancarios
Otra razón para las matrices de relleno es el mecanismo bancario en cuda, con respecto al acceso a la memoria compartida. Cuando la matriz se encuentra en la memoria compartida, se divide en varios bancos de memoria. Dos hilos cuda pueden acceder a él simultáneamente, siempre que no accedan a la memoria que pertenece al mismo banco de memoria. Dado que generalmente deseamos tratar cada fila en paralelo, podemos asegurarnos de que podamos acceder a ella de forma simulada rellenando cada fila al comienzo de un nuevo banco.
Ahora, en lugar de asignar la matriz 2D con cudaMalloc, usaremos cudaMallocPitched:
size_t pitch;
float* myArray;
cudaMallocPitch(&myArray,&pitch,100*sizeof(float),100);//width in bytes by height
Tenga en cuenta que el tono aquí es el valor de retorno de la función: cudaMallocPitch verifica qué debería ser en su sistema y devuelve el valor apropiado. Lo que hace cudaMallocPitch es lo siguiente:
- Asigna la primera fila.
- Compruebe si el número de bytes asignados lo hace alineado correctamente ( es decir , es un múltiplo de 128).
- De lo contrario, asigne más bytes para alcanzar el siguiente múltiplo de 128. El tono es entonces el número de bytes asignados para una sola fila, incluidos los bytes adicionales (bytes de relleno).
- Reiterate para cada fila.
Al final, normalmente hemos asignado más memoria de la necesaria porque cada fila ahora tiene el tamaño del tono, y no el tamaño de w * sizeof (flotante).
Pero ahora, cuando queremos acceder al siguiente elemento en una columna, debemos hacer:
float next_column_element = myArray[(j+1)*pitch+i];
El desplazamiento en bytes entre dos columnas sucesivas no se puede deducir del tamaño de nuestra matriz, por eso queremos mantener el tono devuelto por cudaMallocPitch. Y dado que el tono es un múltiplo del tamaño de relleno (por lo general, el mayor tamaño de palabra y tamaño de banco), funciona de maravilla. Hurra.
Copiando datos a / desde memoria inclinada
Ahora que sabemos cómo crear y acceder a un elemento individual en una matriz creada por cudaMallocPitch, es posible que deseemos copiar una parte completa de ella a otra memoria, lineal o no.
Digamos que queremos copiar una nuestra matriz en una matriz de 100x100 asignada en nuestro host con malloc:
float* host_memory = (float*)malloc(100*100*sizeof(float));
Si usamos cudaMemcpy, copiaremos toda la memoria asignada con cudaMallocPitch, incluidos los bytes acolchados entre cada fila. Lo que debemos hacer para evitar la memoria de relleno es copiar cada fila una a una. Podemos hacerlo manualmente:
for (size_t i=0;i<100;++i) {
cudaMemcpy(host_memory[i*100],myArray[pitch*i],
100*sizeof(float),cudaMemcpyDeviceToHost);
}
O podemos decirle a la API de cuda que solo queremos la memoria útil de la memoria que asignamos con bytes de relleno para su conveniencia, así que si pudiera manejar su propio desorden automáticamente, sería realmente muy agradable, gracias. Y aquí ingresa cudaMemcpy2D:
cudaMemcpy2D(host_memory,100*sizeof(float)/*destination pitch*/,myArray,pitch,
100*sizeof(float)/*width*/,100/*heigth*/,cudaMemcpyDeviceToHost);
Ahora la copia se hará automáticamente. Copiará la cantidad de bytes especificados en ancho (aquí: 100xsizeof (float)), heigth time (en este caso: 100), omitiendo los bytes de tono cada vez que salte a la siguiente fila. Tenga en cuenta que aún debemos proporcionar el tono para la memoria de destino porque también puede rellenarse. Aquí no es así, entonces el tono es igual al tono de una matriz no acolchada: es el tamaño de una fila. Tenga en cuenta también que el parámetro width en la función memcpy se expresa en bytes, pero el parámetro heigth se expresa en número de elementos. Esto se debe a la forma en que se realiza la copia, de alguna manera como escribí la copia manual anterior: el ancho es el tamaño de cada copia a lo largo de una fila (elementos que están contiguos en la memoria) y la altura es el número de veces que esta operación debe ser logrado. (Estas inconsistencias en unidades, como físico, me molestan mucho).
Tratar con matrices 3D
Las matrices 3D no son diferentes a las matrices 2D en realidad, no se incluye relleno adicional. Una matriz 3D es simplemente una matriz clásica en 2D de filas acolchadas. Es por eso que cuando se asigna una matriz 3D, solo se obtiene un paso que es la diferencia en el recuento de bytes entre los puntos sucesivos a lo largo de una fila. Si desea acceder a puntos sucesivos a lo largo de la dimensión de profundidad, puede multiplicar el tono de forma segura por el número de columnas, lo que le proporciona el punto de corte.
La API api para acceder a la memoria 3D es ligeramente distinta a la de la memoria 2D, pero la idea es la misma:
- Al usar cudaMalloc3D, recibirá un valor de tono que debe guardar cuidadosamente para el acceso posterior a la memoria.
- Al copiar un fragmento de memoria 3D, no puede usar cudaMemcpy a menos que esté copiando una sola fila. Debe usar cualquier otro tipo de copia proporcionada por la utilidad cuda que tenga en cuenta el tono.
- Cuando copie sus datos en / desde la memoria lineal, debe proporcionar un tono a su puntero aunque sea irrelevante: este tono es del tamaño de una fila, expresado en bytes.
- Los parámetros de tamaño se expresan en bytes para el tamaño de fila y en el número de elementos para la columna y la dimensión de profundidad.
Comprendo bastante bien cómo asignar y copiar memoria lineal con cudaMalloc() y cudaMemcpy() . Sin embargo, cuando quiero usar las funciones CUDA para asignar y copiar matrices 2D o 3D, a menudo me confunden los diversos argumentos, especialmente con respecto a los punteros pitchados que siempre están presentes cuando se trata de matrices 2D / 3D. La documentación es buena para proporcionar algunos ejemplos sobre cómo usarlos, pero supone que estoy familiarizado con la noción de relleno y tono, que no soy.
Normalmente termino retocando los diversos ejemplos que encuentro en la documentación o en otro lugar en la web, pero la depuración ciega que sigue es bastante dolorosa, así que mi pregunta es:
¿Qué es un lanzamiento? ¿Como lo uso? ¿Cómo puedo asignar y copiar matrices 2D y 3D en CUDA?
En la respuesta de Ernest_Galbrun
float next_column_element = myArray[(j+1)*pitch+i];
necesitan ser
float next_column_element = *((float*)((char*)myArray + (j+1) * pitch) + i);
como en documentation
Y como señaló @RobertCrovella,
float next_column_element = myArray[(j+1)*pitch/sizeof(float)+i];
tampoco es un camino correcto.
En la respuesta de Hefesto
Si usamos cudaMemcpy, copiaremos toda la memoria asignada con cudaMallocPitch, incluidos los bytes acolchados entre cada fila. Lo que debemos hacer para evitar la memoria de relleno es copiar cada fila una a una. Podemos hacerlo manualmente:
for (size_t i=0;i<100;++i) { cudaMemcpy(host_memory[i*100],myArray[pitch*100], 100*sizeof(float),cudaMemcpyDeviceToHost); }
Aquí, la dirección de "memoria de origen" debe ser myArray[i*pitch] lugar de myArray[pitch*100] .