CUDA deaktiviert den L1-Cache nur für eine Variable

Gibt es eine Möglichkeit auf CUDA 2.0-Geräten, den L1-Cache nur für eine bestimmte Variable zu deaktivieren? Ich weiß, dass man L1-Cache zur Kompilierzeit deaktivieren kann, indem man für alle Speicheroperationen das Flag -Xptxas -dlcm=cg zu nvcc hinzufügt. Ich möchte jedoch den Cache nur für Speicherlesevorgänge bei einer bestimmten globalen Variablen deaktivieren, sodass der gesamte Rest des Speichers gelesen wird, um den L1-Cache zu durchlaufen.CUDA deaktiviert den L1-Cache nur für eine Variable

Basierend auf einer Suche, die ich im Web gemacht habe, ist eine mögliche Lösung durch PTX-Assembly-Code.

Quelle

2012-09-23 zeus2

Wie bereits erwähnt Sie Inline-PTX verwenden können, hier ein Beispiel:

__device__ __inline__ double ld_gbl_cg(const double *addr) { 
    double return_value; 
    asm("ld.global.cg.f64 %0, [%1];" : "=d"(return_value) : "l"(addr)); 
    return return_value; 
}

Sie leicht das für .F32 durch Vertauschen .f64 variieren kann (float) oder .s32 (int) usw., die Einschränkung von return_value "= d" für "= f" (float) oder "= r" (int) etc. Beachten Sie, dass die letzte Einschränkung vor (addr) - "l" - 64-Bit-Adressierung bedeutet, wenn Sie 32-Bit-Adressierung verwenden, sollte es "r" sein.

Quelle

2012-09-23 20:39:48 Reguj

Danke! Das funktioniert großartig! – zeus2

@Reguj, wird dies nicht von NVIDIA-Kopfzeilen bereitgestellt? – einpoklum

[this] (https://nvlabs.github.io/cub/classcub_1_cache_modified_input_iterator.html#details) kann von Interesse sein –

Inline PTX kann zum Laden und Speichern der Variablen verwendet werden. Die Befehle ld.cg und st.cg cachen nur Daten in L2. Die Cache-Operatoren sind in Abschnitt 8.7.8.1 Cache-Operatoren des Dokuments PTX ISA 2.3 beschrieben. Die Anweisungen oder das Interesse sind ld und st. Inline-PTX ist in Using Inline PTX Assembly in CUDA beschrieben.

Quelle

2012-09-23 16:56:22

Wenn Sie die Variable als volatile deklarieren, wird sie nur im L2-Cache auf Fermi-GPUs zwischengespeichert. Beachten Sie, dass einige Compileroptimierungen, z. B. das Entfernen wiederholter Ladevorgänge, bei flüchtigen Variablen nicht ausgeführt werden, da der Compiler davon ausgeht, dass sie möglicherweise von einem anderen Thread geschrieben werden.

Quelle

2012-09-23 17:59:02 Heatsink

Ich glaube nicht, dass das Programmiermodell irgendeine Darstellung über die Cache-Fähigkeit von flüchtigen Variablen macht. – ArchaeaSoftware

@Archaea Die Fermi-Architektur macht das Zwischenspeichern flüchtiger Daten aufgrund der fehlenden Cachekohärenz unmöglich. Nachdem ich in der CUDA-Dokumentation in der Vergangenheit auf Fehler gestoßen bin, halte ich CUDAs Speichermodell-Dokumentation nicht für zuverlässig. – Heatsink

Ich versuchte die Lösung mit volatile variable Deklaration und es hat nicht funktioniert. Es scheint, dass die Variable erneut zwischengespeichert wird. – zeus2

CUDA deaktiviert den L1-Cache nur für eine Variable

Antwort

Verwandte Themen