__Threadfence impliziert die Wirkung von __syncThreads?

Question

Ich implementiere parallele Reduktion in CUDA.

Der Kernel hat eine __syncthreads warten auf alle Threads zu vervollständigen 2 Lesevorgänge aus dem gemeinsamen Speicher, die dann die Summe in den gemeinsamen Speicher zurückschreiben würde.

Sollte ich eine __threadfence_block verwenden, um sicherzustellen, dass Schreibvorgänge im gemeinsam genutzten Speicher für alle Threads für die nächste Iteration sichtbar sind, oder verwenden Sie __syncthreads wie in NVIDIA's example angegeben?

Answer 1

__syncthreads() impliziert eine Memory Fence-Funktion. Dies ist in der documentation behandelt:

wartet, bis alle Fäden im Faden-Block erreicht haben diesen Punkt und alle globalen und gemeinsam genutzten Speicherzugriffe, die durch diesen Fäden vor __syncthreads() sind für alle Threads im Block gemacht .

Also in diesem Fall wäre es nicht notwendig sein __threadfence_block() zusätzlich zu verwenden, um __syncthreads()

können Sie keine threadfence Funktion für die Ausführung Barriere in der üblichen allgemeinen parallel Reduktion ersetzen. Die Ausführungsbarriere (__syncthreads()) wird zusätzlich zur Memory Fencing-Funktion benötigt. Im allgemeinen Fall ist es im Allgemeinen notwendig zu warten, bis alle Threads eine bestimmte Reduktionsrunde ausgeführt haben, bevor mit der nächsten Runde fortgefahren wird; __threadfence_block() selbst wird keine Warps zwingen zu warten, während andere Warps eine bestimmte Runde der Reduktion ausführen.

Daher ist in der Regel __syncthreads() erforderlich, und vorausgesetzt, Sie haben es ordnungsgemäß verwendet, ist die __threadfence_block() in der Regel nicht erforderlich.

__syncthreads() bedeutet __threadfence_block().

__threadfence_block() bedeutet nicht __syncthreads()