c++ - crear - asignando memoria compartida

crear procesos en linux

Question

Estoy tratando de asignar memoria compartida usando un parámetro constante pero obteniendo un error. mi núcleo se ve así:

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

y recibo un error que dice

error: la expresión debe tener un valor constante

¡Contar es const! ¿Por qué recibo este error? ¿Y cómo puedo evitar esto?

Answer 1

CUDA admite la asignación dinámica de memoria compartida. Si declaras el núcleo así:

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

y luego pase el número de bytes requeridos como el tercer argumento del lanzamiento del kernel

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

luego, 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, necesita usar punteros para compensar dentro de esa asignación única. Además, tenga en cuenta que al usar punteros, la memoria compartida utiliza palabras de 32 bits y todas las asignaciones deben estar alineadas con palabras de 32 bits, independientemente del tipo de asignación de la memoria compartida.

Answer 2

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

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

dónde :

• Dg es de tipo dim3 y especifica la dimensión y el tamaño de la grilla ...
• Db es del tipo dim3 y especifica la dimensión y el tamaño de cada bloque ...
• Ns es del tipo size_t y especifica el número de bytes en la memoria compartida que se asigna dinámicamente por bloque para esta llamada, además de la memoria estáticamente asignada. esta memoria asignada dinámicamente 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 ...

Entonces, al usar el parámetro dinámico Ns, el usuario puede especificar el tamaño total de la memoria compartida que puede usar una función del kernel, sin importar cuántas variables compartidas haya en este kernel.

Answer 3

No puede declarar una variable compartida como esta ..

__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 arreglar la memoria compartida en un bloque (y no utilizarlo por completo) lo llevará al cambio de contexto con memoria global (alta latencia), por lo tanto, un rendimiento deficiente ...

Hay una buena solución a este problema para declarar

extern __shared__ int a[];

y asignando la memoria al llamar kernel desde la memoria como

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

pero también debería molestarse aquí porque si está usando más memoria en bloques de la que está asignando en kernel, obtendrá resultados inesperados.

Answer 4

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 inicio 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: a los punteros a la memoria compartida dinámicamente se les da la misma dirección. Uso dos matrices de memoria compartidas para ilustrar cómo configurar manualmente dos matrices en la memoria compartida.

Answer 5

const no significa "constante", significa "solo lectura".

Una expresión constante es algo cuyo valor es conocido por el compilador en tiempo de compilación.