Unter welchen Umständen sollten Sie das Schlüsselwort volatile mit dem gemeinsamen Speicher eines CUDA-Kerns verwenden? Ich verstehe, dass volatile dem Compiler sagt, niemals Werte zwischenzuspeichern, aber meine Frage betrifft das Verhalten mit einem Shared Array:
Brauche ich products , um in diesem Fall volatil zu sein? Auf jeden Array-Eintrag wird nur von einem einzigen Thread zugegriffen, außer am Ende, wo alles von Thread 0 gelesen wird. Ist es möglich, dass der Compiler das gesamte Array zwischenspeichern kann, und ich brauche es also volatile , oder wird es nur Cache-Elemente?
Danke!
Wenn Sie ein gemeinsam genutztes Array nicht als volatile deklarieren, kann der Compiler die Speicherorte im Shared Memory optimieren, indem er sie in Registern (deren Geltungsbereich für einen einzelnen Thread spezifisch ist) für jeden beliebigen Thread anordnet Wählen. Dies gilt unabhängig davon, ob Sie auf dieses bestimmte freigegebene Element von nur einem Thread zugreifen oder nicht. Wenn Sie Shared Memory als Kommunikationsmedium zwischen Threads eines Blocks verwenden, empfiehlt es sich daher, volatile zu deklarieren.
Offensichtlich, wenn jeder Thread nur auf seine eigenen Elemente im Shared Memory zugreift, und niemals auf solche, die mit einem anderen Thread verbunden sind, spielt dies keine Rolle, und die Compiler-Optimierung wird nichts kaputt machen.
In Ihrem Fall, wo Sie einen Codeabschnitt haben, in dem jeder Thread auf seine eigenen Elemente des gemeinsamen Speichers zugreift und der einzige Thread-Inter-Zugriff an einem gut bekannten Ort stattfindet, könnten Sie ein Memory-Fence-Funktion , um den Compiler zu veranlassen, alle Werte zu löschen, die vorübergehend gespeichert sind Register, zurück zum Shared Array. Sie könnten also denken, dass __threadfence_block() nützlich sein könnte, aber in Ihrem Fall __syncthreads() ist bereits eine Memory-Fencing-Funktion eingebaut . Daher reicht Ihr Aufruf __syncthreads() aus, um die Thread-Synchronisierung zu erzwingen und alle im Register zwischengespeicherten Werte im gemeinsam genutzten Speicher in den gemeinsamen Speicher zu übertragen.
Wenn diese Reduzierung am Ende Ihres Codes von Leistungsproblemen betroffen ist, können Sie eine Parallelreduzierungsmethode verwenden, um die Geschwindigkeit zu erhöhen.
Tags und Links compiler-construction volatile cuda gpu gpgpu