He pasado por muchos foros y nvidia manual, pero no pude entender lo que es __threadfence() y el uso de él?CUDA __threadfence()
Gracias.
He pasado por muchos foros y nvidia manual, pero no pude entender lo que es __threadfence() y el uso de él?CUDA __threadfence()
Gracias.
Normalmente, no hay garantía de que si un bloque escribe algo en la memoria global, el otro bloque "lo verá". Tampoco se garantiza el orden de las escrituras en la memoria global, con la excepción del bloque que lo emitió.
Hay dos excepciones:
Imagínese, que un bloque produce algunos datos, y luego usa operación atómica de marque una bandera que los datos están allí. Pero es posible que el otro bloque vea la bandera, pero leerá datos incorrectos o incompletos.
__threadfence
función detiene el hilo actual hasta que sus escrituras a la memoria global se garanticen para ser visibles por todos los otros hilos en la red. Por lo tanto, si haces algo como:
__threadfence()
se garantiza que si el otro bloque ve la bandera, también verá los datos.
Más información: Guía de programación Cuda, capítulos B.2.4 y B.5
¿Qué tal __syncthreads(), hace que la garantía de que los accesos a memoria por cualquier tema en el bloque será visible a todos los hilos en el bloque? –
'__syncthreads()' es más fuerte que '__threadfence_block()'. Después de '__syncthreads()' tiene la garantía de que todas las escrituras de memoria compartida/global antes de la barrera son visibles por todos los hilos después de la barrera. Sin embargo, '__syncthreads()' tiene un efecto solo en un bloque y no se dan garantías entre subprocesos de diferentes bloques. – CygnusX1
L1 de otros SMs será coherente con los datos almacenados? ¿O todavía necesita especificar cargas de alcance global (L1-no-cacheable)? – maxbc