CUDA Hafıza Organizasyonu

1. Yerel Bellek (Local Memory)

Program çalışırken kullanıldığından kullanıcıya açık bir bellek yapısı değildir. Buradaki veriler, sahip olunan her bir thread’in çalışması bitene kadar saklanır. Her local bellek yapısı, ait olduğu thread’e özeldir. Diğer thread’ler tarafından ulaşılamazlar. Kendi thread’leri tarafından üzerine veri yazılabilir ve üzerinden veri okunabilir.

</a>

2. Paylaşımlı Bellek (Shared Memory)

Grafik İşlem Birimi’nin sahip olduğu tüm çekirdek kümelerinde bulunmaktadır. Yalnızca aynı block’ta bulunan thread’ler bu belleğe veri yazıp okuyabilirler. Aynı block içerisindeki thread’ler çakışmadığı sürece çok hızlı çalışmaktadır.

</a>

3. Genel Bellek (Global Memory)

Çalışan bütün thread’ler tarafından erişilebilen , üzerinden veri okunabilen ve üzerine veri yazılabilen bir bellek türüdür. Bu belleğin bant genişliği düşüktür.

</a>

Global bellek, Host ile Device arasındaki temel haberleşme birimidir. Uzun erişme süreleri vardır. Alttaki şekilde global bellek, paylaşımlı bellek ve her bir thread’e ait olan kaydedici hafıza alanları görülmektedir. Ok yönleri , ilgili bellek gözlerinden veri okuyup-yazma izinlerini belirtmektedir.

</a>

4. Sabit Bellek (Constant Memory)

Read-only(üzerinden sadece okuma yapılabilen)bir bellek modelidir. Kernel çalıştırıldığında değişmeyecek veriler için kullanılması avantaj sağlar. 64 KB’lık büyüklüğe sahiptir. NVIDIA’nın değişmez bellek modelini sunmasındaki amaç bellek bant genişliğinin global memory’ye göre daha az olmasındandır. Eğer geliştirilen uygulamada veriler sadece okunacaksa ve yazılmayacaksa bu bellek modeli kullanılmalıdır. Değişmez bellekten dinamik olarak yer alınmasına izin verilmemektedir. Değişmez belleğe yalnızca MİB fonksiyonları tarafından veri yazılabilir.

Değişmez bellek üzerine veri kopyalamak için cudaMemcpyToSymbol() fonksiyonu kullanılır. cudaMemcpy() fonksiyonundan farkı, cudaMemcpy() fonksiyonun global bellek üzerine kopyalama yapması ve cudaMemcpyToSymbol() fonksiyonunun değişmez bellek üzerine veri kopyalamasıdır.

Constant Memory Kullanılması Tavsiye Edilen Durumlar

• Giriş verimizin “execution” süresince değişmeyeceğini bildiğimiz durumlarda
• Tüm thread’lerimiz memory alanının aynı adres alanından veriye erişecekse. Örneğin; bir thread block’undaki tüm threadler aynı hafıza adres alanını işaret ediyorsa

Constant Memory Kullanılması Tavsiye Edilmeyen Durumlar

• Giriş verimizin “execution” süresince değişeceğini biliyorsak
• Tüm thread’lerimiz memory alanının aynı adres alanından veriye erişmeyecekse
• Verimiz Read-only değilse

Constant Memory Performans Değerlendirmesi 64 KB’lık constant memory’den okumanın , global memory’den okumaya göre 3 avantajı vardır:

1) Constant memory’den 1 okuma ile yakınlardaki(nearby) thread’lere broadcast yapılabilmesi ile 15 adet okumadan tasarruf edilir.

2) Constant memory “cached” bir hafıza alanıdır, bu yüzden aynı adresten okumalar herhangi bir ek bellek trafiğine maruz kalmayacaktır.

3) Ön belleğe sahip olduğu için global bellekten daha hızlıdır.

5. Doku Bellek (Texture Memory)

Doku belleği de değişmez bellek gibi read-only yapıya sahiptir. Bu yapı performansı artırır ve bellek trafiğini azaltır.

NVIDIA, texture birimini klasik OpenGL ve DirectX rendering iş hattı için tasarlamıştır. Texture Memory bunu kullanışlı kılmak için bazı özelliklere sahiptir.

Constant memory gibi, texture memory de chip üzerinde önbellekte(cached)dir. Bu yapı bazı durumlarda DRAM’de daha az talep trafiği ile daha etkin düzeyde bantgenişliği olanağı sunar. Özellikle, doku önbellekleri grafik uygulamaları için tasarlanmıştır.

Referanslar

1. NVIDIA,(2012) “NVIDIA CUDA C Programming Guide Version 4.2”
2. BRADLEY, T. (2010) ,“Advanced CUDA Optimization:1-Introduction”, NVIDIA Corporotion
3. “Introduction to CUDA 5.0” http://www.3dgep.com/introduction-to-cuda-5-0/
4. “Constant Memory in CUDA” , http://cuda-programming.blogspot.com.tr/2013/01/what-is-constant-memory-in-cuda.html
5. SANDERS, J. ,KANDROT, E. (2010) “CUDA by Example:An Introduction to General-Purpose GPU Programming”, NVIDIA Corporotion