CUDA wyłącza pamięć podręczną L1 tylko dla jednej zmiennej

Czy jest jakiś sposób na urządzeniach CUDA 2.0 wyłączających pamięć podręczną L1 tylko dla jednej konkretnej zmiennej? Wiem, że można wyłączyć pamięć podręczną L1 podczas kompilacji, dodając flagę -Xptxas -dlcm=cg do nvcc dla wszystkich operacji związanych z pamięcią. Jednak chcę wyłączyć pamięć podręczną tylko dla odczytów pamięci dla określonej zmiennej globalnej, tak aby cała reszta pamięci była czytana, aby przejść przez pamięć podręczną L1.CUDA wyłącza pamięć podręczną L1 tylko dla jednej zmiennej

Na podstawie wyszukiwania, które przeprowadziłem w Internecie, możliwe jest rozwiązanie poprzez kod zespołu PTX.

Źródło

2012-09-23 zeus2

Jak wspomniano powyżej, można użyć inline PTX, oto przykład:

__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; 
}

Można łatwo zmieniać się to poprzez zamianę .f64 dla .f32 (float) lub .s32 (int) itd, ograniczenie wartości zwracanej "= d" dla "= f" (float) lub "= r" (int) etc. Zauważ, że ostatnie ograniczenie przed (addr) - "l" - oznacza adresowanie 64-bitowe, jeśli używasz 32-bitowego adresowania, powinno to być "r".

Źródło

2012-09-23 20:39:48 Reguj

Dzięki! To działa świetnie! – zeus2

@Reguj, czy nigdzie nie są dostarczane przez nagłówki NVIDIA? – einpoklum

[this] (https://nvlabs.github.io/cub/classcub_1_1_cache_modified_input_iterator.html#details) może być interesujące –

Inline PTX może być użyty do załadowania i zapisania zmiennej. Instrukcje ld.cg i st.cg zawierają tylko dane pamięci podręcznej w L2. Operatory pamięci podręcznej są opisane w sekcji 8.7.8.1 Operatory buforowania w dokumencie PTX ISA 2.3. Instrukcje lub zainteresowania: ld i st. Inline PTX jest opisane w Using Inline PTX Assembly in CUDA.

Źródło

2012-09-23 16:56:22

Jeśli zadeklarujesz zmienną jako volatile, zostanie ona zapisana w pamięci podręcznej tylko w pamięci podręcznej L2 na procesorach graficznych Fermi. Zauważ, że niektóre optymalizacje kompilatora, takie jak usuwanie powtarzających się obciążeń, nie są wykonywane na zmiennych ulotnych, ponieważ kompilator zakłada, że mogą być zapisywane przez inny wątek.

Źródło

2012-09-23 17:59:02 Heatsink

Nie sądzę, aby model programowania reprezentował pamięć podręczną zmiennych ulotnych. – ArchaeaSoftware

@Archaea Architektura Fermi powoduje, że buforowanie lotnych danych nie jest możliwe z powodu braku spójności pamięci podręcznej. Po wystąpieniu błędów w dokumentacji CUDA w przeszłości, nie uważam dokumentacji modelu pamięci CUDA za wiarygodną. – Heatsink

Próbowałem rozwiązania z lotną zmienną dekloryzacją i to nie działało. Wydaje się, że zmienna jest ponownie buforowana. – zeus2

CUDA wyłącza pamięć podręczną L1 tylko dla jednej zmiennej

Odpowiedz

Powiązane problemy