asignación de memoria compartida

CUDA admite la asignación dinámica de memoria compartida. Si declara el núcleo de esta manera:

__global__ void Kernel(const int count)
{
    extern __shared__ int a[];
}

Y luego pasar el número de bytes necesarios como el tercer argumento del lanzamiento del núcleo

Kernel<<< gridDim, blockDim, a_size >>>(count)

Entonces a se puede dimensionar en tiempo de ejecución. Tenga en cuenta que el tiempo de ejecución solo admite una única asignación declarada dinámicamente por bloque. Si necesita más, también necesita usar punteros para compensaciones dentro de esa asignación única. También tenga en cuenta al usar punteros que la memoria compartida usa 32 bits palabras, y todas las asignaciones deben estar alineadas con palabras de 32 bits, independientemente del tipo de asignación de memoria compartida.

Opción uno: declarar memoria compartida con valor constante (no es lo mismo que const)

__global__ void Kernel(int count_a, int count_b)
{
    __shared__ int a[100];
    __shared__ int b[4];
}

Opción dos: declarar la memoria compartida dinámicamente en la configuración de lanzamiento del kernel:

__global__ void Kernel(int count_a, int count_b)
{
    extern __shared__ int *shared;
    int *a = &shared[0]; //a is manually set at the beginning of shared
    int *b = &shared[count_a]; //b is manually set at the end of a
}

sharedMemory = count_a*size(int) + size_b*size(int);
Kernel <<<numBlocks, threadsPerBlock, sharedMemory>>> (count_a, count_b);

nota: Los punteros a memoria dinámicamente compartida son todos dados la misma dirección. Utilizo dos matrices de memoria compartida para ilustrar cómo configurar manualmente dos matrices en memoria compartida.

De la " Guía de Programación CUDA C": La configuración de ejecución se especifica insertando una expresión de la forma:

<<<Dg, Db, Ns, S>>>

Donde:

• Dg es de tipo dim3 y especifica las dimensiones y el tamaño de la cuadrícula ...
• Db es de tipo dim3 y especifica la dimensión y el tamaño de cada bloque ...
• Ns es de tipo size_t y especifica el número de bytes en la memoria compartida que se dinámicamente asignado por bloque para esta llamada además de la memoria estáticamente asignada. esta memoria dinámicamente asignada es utilizada por cualquiera de las variables declaradas como una matriz externa como se menciona en _ _ shared _ _ ; Ns es un argumento opcional que por defecto es 0;
• S es de tipo cudaStream_t y especifica la secuencia asociada ...

Por lo tanto, mediante el uso del parámetro dinámico Ns, el usuario puede especificar el tamaño total de memroy compartido que una función del núcleo puede usar, no importa cuántas variables compartidas haya en este núcleo.

No se puede declarar variable compartida de esta manera..

__shared__ int a[count];

Aunque si está lo suficientemente seguro sobre el tamaño máximo de la matriz a, entonces puede declarar directamente como

__shared__ int a[100];

Pero en este caso debería preocuparse por cuántos bloques hay en su programa , ya que fijar la memoria compartida a un bloque ( y no ser utilizado completamente), le llevará a cambiar de contexto con memoria global( alta latencia) , por lo tanto, un rendimiento pobre...

Hay una buena solución a este problema declarar

extern __shared__ int a[];

Y asignar la memoria mientras se llama al núcleo desde la memoria como

Kernel<<< gridDim, blockDim, a_size >>>(count)

Pero también debería molestarse aquí porque si está utilizando más memoria en bloques de la que está asignando en el núcleo , obtendrá resultados inesperados.