Escrituras concurrentes en la misma ubicación de memoria global

Question

Tengo varios bloques, cada uno con algunos enteros en una matriz de memoria compartida de tamaño 512. ¿Cómo puedo verificar si la matriz en cada bloque contiene un cero como elemento?

Lo que estoy haciendo es crear una matriz que reside en la memoria global. El tamaño de esta matriz depende del número de bloques y se inicializa en 0. De ahí que cada bloque escriba en a[blockid] = 1 si la matriz de memoria compartida contiene un cero.

Mi problema es cuando tengo varios hilos en un solo bloque escribiendo al mismo tiempo. Es decir, si la matriz en la memoria compartida contiene más de un cero, varios hilos escribirán a[blockid] = 1. ¿Esto generaría algún problema?

En otras palabras, ¿Sería un problema si 2 hilos escriben exactamente el mismo valor en el mismo elemento de matriz en la memoria global?

Answer 1

En el modelo de ejecución CUDA, no hay garantías de que todos los escritura simultánea de hilos en el mismo bloque en la misma ubicación de memoria global tendrá éxito. Al menos una escritura funcionará, pero el modelo de programación no garantiza cuántas transacciones de escritura se producirán, ni en qué orden ocurrirán si se ejecuta más de una transacción.

Si esto es un problema, entonces un mejor enfoque (desde el punto de vista de la corrección) sería tener solo un hilo de cada bloque para la escritura global. Puede usar un indicador de memoria compartida configurado atómicamente o una operación de reducción para determinar si el valor debe establecerse. Lo que elija puede depender de cuántos ceros haya. Cuantos más ceros haya, más atractiva será la reducción. CUDA incluye warp level __any() y __all() operadores que pueden integrarse en una reducción booleana muy eficiente en unas pocas líneas de código.

Answer 2

Sí, será un problema llamado Race Condition.
Debe tener en cuenta synchronizing acceso a los datos globales a través de process Semaphores

Answer 3

Si bien no es un mutex o semáforo, CUDA contiene un primative de sincronización que puede utilizar para serializar el acceso a un segmento de código dado o ubicación de memoria. Mediante la función __syncthreads(), puede crear una barrera para que cualquier hilo dado bloquee en el punto de la llamada al comando hasta que todos los hilos en un bloque determinado hayan ejecutado el comando __syncthreads(). De esta forma, esperemos que serialice el acceso a la ubicación de su memoria y evite una situación en la que dos hilos necesiten escribir en la misma ubicación de memoria al mismo tiempo. La única advertencia es que todos los subprocesos tienen que ejecutar en algún momento __syncthreads(), de lo contrario terminará con una situación de bloqueo. Por lo tanto, no coloque la llamada dentro de alguna instrucción if condicional donde algunos subprocesos nunca ejecuten el comando. Si aborda su problema de esta manera, deberá haber alguna disposición para los hilos que inicialmente no llaman al __syncthreads() para llamar a la función más tarde a fin de evitar un interbloqueo.

Answer 4

Para un programa CUDA, si varios hilos en una urdimbre de escritura en el mismo lugar, entonces la ubicación se ser actualizados pero es ocurren indefinidoscuántas veces la ubicación se actualiza (es decir, el número de escrituras reales en serie) y es undefinedcuyo thread escribirá último (es decir, qué thread ganará la carrera).

Para dispositivos de capacidad de cálculo 2.x, si hay varios hilos en un warp que escriben en la misma dirección, entonces solo un hilo ejecutará la escritura, cuyo hilo no está definido.

Desde el CUDA C Programming Guide sección F.4.2:

Si una instrucción no atómica ejecutado por un túnel escribe en la misma ubicación en la memoria global para más de uno de los hilos de la urdimbre, sólo un hilo realiza una escritura y qué hilo lo hace no está definido.

Consulte también la sección 4.1 de la guía para obtener más información.

En otras palabras, si todos los hilos que escriben en una ubicación determinada escriben el mismo valor, entonces es seguro.