Cuda atomics cambiar la bandera

Tengo un código de serie que hace algo como esto

if( ! variable ) { do some initialization here variable = true; } 

Entiendo que esto funciona perfectamente bien en serie y solo se ejecutará una vez. ¿Qué operación atómica sería la correcta aquí en CUDA?

Me parece que lo que quieres es una “sección crítica” en tu código. Una sección crítica permite que un hilo ejecute una secuencia de instrucciones mientras evita que otro hilo o bloque de hilos ejecute esas instrucciones.

Una sección crítica se puede usar para controlar el acceso a un área de memoria, por ejemplo, para permitir el acceso sin conflictos a esa área mediante un solo hilo.

Las atómicas por sí mismas solo pueden usarse para una operación muy limitada, básicamente única, en una sola variable. Pero los átomos pueden usarse para construir una sección crítica.

Debe usar el siguiente código en su núcleo para controlar el acceso de subprocesos a una sección crítica:

 __syncthreads(); if (threadIdx.x == 0) acquire_semaphore(&sem); __syncthreads(); //begin critical section // ... your critical section code goes here //end critical section __syncthreads(); if (threadIdx.x == 0) release_semaphore(&sem); __syncthreads(); 

Antes del kernel define estas funciones auxiliares y la variable del dispositivo:

 __device__ volatile int sem = 0; __device__ void acquire_semaphore(volatile int *lock){ while (atomicCAS((int *)lock, 0, 1) != 0); } __device__ void release_semaphore(volatile int *lock){ *lock = 0; __threadfence(); } 

He probado y utilizado con éxito el código anterior. Tenga en cuenta que básicamente arbitra entre los bloques de hilos usando el hilo 0 en cada bloque de hilos como un solicitante. Debería condicionar aún más (ej if (threadIdx.x < ...) ) su código de sección crítica si solo quiere un hilo en el bloque de hilos ganador para ejecutar el código de sección crítica.

Tener múltiples hilos dentro de un warp arbitrate para un semáforo presenta complejidades adicionales, por lo que no recomiendo ese enfoque. En su lugar, haga que cada bloque de hilos arbitre como lo he mostrado aquí, y luego controle su comportamiento dentro del bloque de hilos ganador usando los métodos comunes de comunicación / sincronización de __syncthreads() por ejemplo, __syncthreads() , memoria compartida, etc.)

Tenga en cuenta que esta metodología será costosa para el rendimiento. Solo debe usar secciones críticas cuando no pueda encontrar la manera de paralelizar su algoritmo.

Finalmente, una palabra de advertencia. Como en cualquier architecture paralela roscada, el uso inadecuado de secciones críticas puede conducir a un punto muerto. En particular, hacer suposiciones sobre el orden de ejecución de bloques de hilos y / o urdimbres dentro de un hilo de rosca es un enfoque defectuoso.