CUDA Mejores prácticas de memoria constante

gpu-constant-memory

Question

Os presento aquí un código

__constant__ int array[1024];

__global__ void kernel1(int *d_dst) {
   int tId = threadIdx.x + blockIdx.x * blockDim.x;
   d_dst[tId] = array[tId];
}

__global__ void kernel2(int *d_dst, int *d_src) {
   int tId = threadIdx.x + blockIdx.x * blockDim.x;
   d_dst[tId] = d_src[tId];
}

int main(int argc, char **argv) {
   int *d_array;
   int *d_src;
   cudaMalloc((void**)&d_array, sizeof(int) * 1024);
   cudaMalloc((void**)&d_src, sizeof(int) * 1024);

   int *test = new int[1024];
   memset(test, 0, sizeof(int) * 1024);

   for (int i = 0; i < 1024; i++) {
     test[i] = 100;
   }

   cudaMemcpyToSymbol(array, test, sizeof(int) * 1024);
   kernel1<<< 1, 1024 >>>(d_array);

   cudaMemcpy(d_src, test, sizeof(int) * 1024, cudaMemcpyHostToDevice);
   kernel2<<<1, 32 >>>(d_array, d_src),

   free(test);
   cudaFree(d_array);
   cudaFree(d_src);

   return 0;
}

Lo que simplemente muestra la memoria constante y el uso global de la memoria. En su ejecución, el "kernel2" se ejecuta aproximadamente 4 veces más rápido (en términos de tiempo) que "kernel1"

Según la guía de programación de Cuda C, entiendo que esto se debe a que los accesos a la memoria constante se están serializando. Lo que me lleva a la idea de que la memoria constante se puede utilizar mejor si una distorsión accede a un valor constante único como entero, flotante, doble, etc. pero no es beneficioso acceder a una matriz. En otros términos, puedo decir que una deformación debe acceder a una única dirección para obtener beneficios de optimización / aceleración beneficiosos gracias al acceso constante a la memoria. ¿Es esto correcto?

También quiero saber si mantengo una estructura en lugar de un tipo simple en mi memoria constante. Cualquier acceso a la estructura por un hilo con en una urdimbre; ¿También se considera como acceso de memoria única o más? Me refiero a que una estructura puede contener varios tipos simples y una matriz, por ejemplo; Al acceder a estos tipos simples, ¿estos accesos también son serializados o no?

La última pregunta sería, en caso de que tenga una matriz con valores constantes, a la que se deba acceder a través de diferentes hilos dentro de una deformación; para un acceso más rápido, debe mantenerse en la memoria global en lugar de la memoria constante. ¿Es eso correcto?

Cualquiera puede remitirme algún código de ejemplo donde se muestre un uso eficiente y constante de la memoria.

Saludos,

Answer 1

Puedo decir que una deformación debe acceder a una sola dirección para obtener beneficios de optimización / aceleración beneficiosos gracias al acceso constante a la memoria. ¿Es esto correcto?

Sí, esto es generalmente correcto y es el principal objetivo del uso de memoria constante / caché constante. La memoria caché constante puede proporcionar una cantidad de 32 bits por ciclo por SM. Por lo tanto, si cada hilo en un warp está accediendo al mismo valor:

int i = array[20];

entonces tendrá la oportunidad de obtener un buen beneficio de la memoria / caché constante. Si cada hilo en una urdimbre está accediendo a una cantidad única:

int i = array[threadIdx.x]; 

luego, los accesos se serializarán y el uso constante de datos será decepcionante, en cuanto al rendimiento.

También quiero saber si mantengo una estructura en lugar de un tipo simple en mi memoria constante. Cualquier acceso a la estructura por un hilo con en una urdimbre; ¿También se considera como acceso de memoria única o más?

Ciertamente puedes poner estructuras en memoria constante. Se aplican las mismas reglas:

int i = constant_struct_ptr->array[20]; 

tiene la oportunidad de beneficiarse, pero

int i = constant_struct_ptr->array[threadIdx.x];

no. Si accede al mismo elemento de estructura de tipo simple a través de subprocesos, es ideal para el uso constante de caché.

La última pregunta sería, en caso de que tenga una matriz con valores constantes, a la que se deba acceder a través de diferentes hilos dentro de una deformación; para un acceso más rápido, debe mantenerse en la memoria global en lugar de la memoria constante. ¿Es eso correcto?

Sí, si sabe que, en general, sus accesos interrumpirán la memoria constante de una cantidad de 32 bits por ciclo, por lo que probablemente sea mejor dejar los datos en la memoria global ordinaria.

Hay una variedad de códigos de muestra de cuda que demuestran el uso de los datos de __constant__ . Aquí hay algunos:

1. gráficos volumeRender
2. Filtración bilateral de imágenes
3. convolución de imágenes
4. financiar MonteCarloGPU

y hay otros

EDITAR: respondiendo a una pregunta en los comentarios, si tenemos una estructura como esta en memoria constante:

struct Simple { int a, int b, int c} s;

Y lo accedemos así:

int p = s.a + s.b + s.c;
          ^     ^     ^
          |     |     |
cycle:    1     2     3

Tendremos un buen uso de la memoria constante / caché. Cuando se compila el código C, bajo el capó generará accesos de código de máquina correspondientes a 1,2,3 en el diagrama anterior. Imaginemos que el acceso 1 ocurre primero. Dado que el acceso 1 se encuentra en la misma ubicación de memoria independiente del hilo en la urdimbre, durante el ciclo 1, todos los hilos recibirán el valor en sa y aprovecharán el caché para obtener el mejor beneficio posible. Igualmente para los accesos 2 y 3. Si por otro lado tuviéramos:

struct Simple { int a[32], int b[32], int c[32]} s;
...
int idx = threadIdx.x + blockDim.x * blockIdx.x;
int p = s.a[idx] + s.b[idx] + s.c[idx];

Esto no daría un buen uso de memoria constante / caché. En cambio, si esto fuera típico de nuestros accesos a s , probablemente tendríamos mejores ubicaciones de rendimiento en la memoria global ordinaria.