GPU çekirdekleri yazarken bellek yönetimi, dikkate almanız gereken en önemli performans özelliklerinden biridir. Bu yazıda, global belleği ve onun performansını etkileyen önemli yönleri ele alacağız.
Global Bellek Nedir?
CUDA cihazında birden fazla bellek türü bulunmaktadır ve her birinin farklı kapsamı, ömrü, ve önbellek davranışları vardır. Global bellek (diğer adıyla cihaz belleği), CUDA cihazlarındaki ana bellek alanıdır. Bu bellek, cihaz DRAM’inde yer alır ve CPU sistemlerindeki RAM’e benzer şekilde çalışır. “Global” terimi, bu bellek alanının kapsamını ifade eder ve hem ana bilgisayar (host) hem de bir çekirdek ızgarasındaki tüm thread’ler tarafından erişilebilir ve değiştirilebilir.
Global bellek, global kapsamda __device__ tanımlama spesifikatorü ile statik olarak ya da cudaMalloc() veya cudaMallocManaged() gibi CUDA çalışma zamanı API’leri ile dinamik olarak tahsis edilebilir. Bellek, ana bilgisayardan cihaza cudaMemcpy() ile aktarılabilir ve cudaFree() ile serbest bırakılabilir. Bu tahsisatlar, serbest bırakılana kadar kalıcıdır.
Ayrıca, Unified Memory kullanılarak da global bellek tahsis veya serbest bırakılabilir. Global bellek tahsis-atma, serbest bırakma ve cihaza veri taşıma işlemleri karmaşık konulardır ve gelecekteki bir yazıda ele alınacaktır. Biz bu yazıda, CUDA çekirdeklerinde global bellek kullanmanın performans etkilerine odaklanacağız.
Global Bellek Koalescing
Global belleğe erişim performansı üzerine konuşmaya başlamadan önce CUDA yürütme modelini daha iyi anlamamız gerekiyor. Thread’lerin bloklar halinde gruplandığını ve bu blokların GPU üzerindeki çok işlemcilerle eşleştirildiğini tartıştık. Bu execute işlemlerinde, thread’ler daha küçük gruplara, yani warp‘lara (32 thread) bölünür. GPU üzerindeki çok işlemciler, her bir warp için SIMT (Single Instruction Multiple Threads) tarzında talimatlar yürütmektedir.
CUDA’da global belleğe erişim sağlarken dikkate almanız gereken kritik bir nokta, aynı warp içindeki farklı thread’lerin eriştiği bellek konumlarının ilişkinin nasıl olduğudur. Bu bellek erişimlerinin düzeni, bellek erişim verimliliğini ve genel uygulama performansını doğrudan etkiler.
Global bellek, 32 baytlık bellek işlemleri ile erişilir. Bir CUDA thread’i global bellekten veri talep ettiğinde, aynı warp içindeki tüm thread’lerden gelen bellek erişimleri, minimum sayıda bellek işlemi haline koalesce edilir. Gerekli bellek işlemi sayısı, her bir thread’in eriştiği kelimenin boyutuna ve thread’ler arasındaki bellek adreslerinin dağılımına bağlıdır.
Aşağıdaki kod, bir warp içindeki ardışık thread’lerin ardışık 4 baytlık veri öğelerine eriştiği bir durumu göstermektedir ve böylelikle optimal bir bellek erişim düzeni yaratılmıştır. Warp tarafından gönderilen tüm yüklemeler, bellekten dört adet 32 baytlık sektörde karşılanabilir. Bu durum, bellek bant genişliğinin en verimli şekilde kullanılmasını sağlar.
__global__ void coalesced_access(float* input, float* output, int n) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid
Diğer taraftan, thread’ler büyük adımlarla bellek erişimi sağlarsa, her bir bellek işlemi gereğinden fazla veri alır. Her bir thread’in talep ettiği 4 baytlık her öğe için global bellekten tam bir 32 baytlık segment alınır ve bu segmentin çoğunluk bölümü kullanılmaz.
Strided Erişim ve Bellek Bant Genişliği
Global belleğe erişim sırasında stride (adım) terimi, bir warp’ın thread’lerinin eriştiği ardışık bellek konumları arasındaki mesafeyi ifade eder. Yukarıda belirtildiği gibi, bu faktör, GPU’nun bellek bant genişliği üzerindeki etkinliği üzerinde önemli bir etkiye sahiptir.
Strided erişimlerin GPU performansını nasıl etkilediğini anlamak için, farklı adım uzunluklarına sahip çekirdeklerin bant genişliği ölçümlerini gösterebiliriz. Bu ölçümler, aksesuar bellek erişimlerinin nasıl gerçekleştirdiğini daha iyi anlamamız için faydalıdır.
Özetle, GPU belleğinin verimli kullanımı, en iyi performansı elde etmek için önemli kriterlerden biridir. Optimal global bellek performansı, koalesce edilmiş bellek erişimlerine dayanır. Global belleğe yapılan strided erişimleri minimize edin ve her zaman Nsight Compute ile GPU çekirdeklerinizi profil alarak bellek erişimlerinin koalesce edildiğinden emin olun. Bu yaklaşım, GPU kodunuzdan en yüksek performansı elde etmenize yardımcı olacaktır.
Sonuç
Bu yazı, Mark Harris’in 2013’te yayınlanan bir gönderisinin güncellemesidir.
