Considere el siguiente código:¿Hay alguna forma de establecer el valor predeterminado para la matriz de memoria compartida?
__global__ void kernel(int *something) {
extern __shared__ int shared_array[];
// Some operations on shared_array here.
}
¿Es posible establecer toda shared_array a algún valor - por ejemplo, 0 - sin direccionar explícitamente cada celda en algún hilo?
No. La memoria compartida no está inicializada. Usted tiene que inicializar de alguna manera a sí mismo, de una manera u otra ...
De CUDA C Guía de programación 3.2, sección B.2.4.2, párrafo 2:
__shared__variables no pueden tener una inicialización como parte de su declaración.
Esto también descarta constructores no triviales por defecto para variables compartidas.
Puede inicializar arrays compartidos de manera eficiente en paralelo como esto
// if SHARED_SIZE == blockDim.x, eliminate this loop
for (int i = threadIdx.x; i < SHARED_SIZE; i += blockDim.x)
shared_array[i] = INITIAL_VALUE;
__syncthreads();
Sí, se puede. Puede especificar que el primer hilo en el bloque establece que, mientras que los otros de NO por ejemplo .:
extern __shared__ unsigned int local_bin[]; // Size specified in kernel call
if (threadIdx.x == 0) // Wipe on first thread - include " && threadIdx.y == 0" and " && threadIdx.z == 0" if threadblock has 2 or 3 dimensions instead of 1.
{
// For-loop to set all local_bin array indexes to specified value here - note you cannot use cudaMemset as it translates to a kernel call itself
}
// Do stuff unrelated to local_bin here
__syncthreads(); // To make sure the memset above has completed before other threads start writing values to local_bin.
// Do stuff to local_bin here
Lo ideal es hacer tanto trabajo como sea posible antes de que los syncthreads llaman, ya que esto permite todos los demás hilos para hacer su trabajo antes de que se complete el memset; obviamente, esto solo importa si el trabajo tiene el potencial de tener tiempos de finalización del hilo bastante diferentes, por ejemplo, si hay una bifurcación condicional. Tenga en cuenta que para el thread 0 "setting" for-loop, necesita haber pasado el tamaño de la matriz local_bin como un parámetro al kernel para que sepa el tamaño de la matriz que está iterando.
Gracias. Usé algo similar a esto en mi implementación final. – rayryeng
esto pierde los beneficios de la paralelización, siempre trate de usar threadIdx y BlockIdx tanto como sea posible – ejectamenta
Creo que tiene razón en general, pero depende de si el valor de inicialización se puede calcular a partir de esos o si tiene un número numérico más complejo/etc patrón de asignación. – metamorphosis
eso es sólo el caso si usted tiene un bloque 1D, por supuesto. Diciendo simplemente que los novatos no caigan en trampas obvias. También me pregunto cuánto aumentará float4, que es otro truco que todavía se aplica a los dispositivos más nuevos y cuánto beneficio proporciona junto con este tipo de init de fusión de memoria. Sidenote, si está cargando dentro de un núcleo 2D o 3D, es importante saber que están divididos en urdimbres, como una matriz de [z] [y] [x] dentro de la memoria. Deje que [x] los hilos se escriban más cerca unos de otros y los diferentes en su [z] más lejos. –
Así que lo he intentado y sí, usando reinterpret_cast para copiar en trozos de 16 bytes como en https://devblogs.nvidia.com/parallelforall/cuda-pro-tip-increase-performance-with-vectorized-memory-access/aún obtiene un poco mejores resultados. Además, es importante establecer __restrict__ en los datos que desea copiar: fácil + 15% de aumento de rendimiento, casi tanto como los fragmentos alineados de 16 bytes. –
Tenga en cuenta que la fusión de la memoria puede significar una diferencia de velocidad de 16x fácilmente. –