Birçok CUDA çekirdek işlemi, bellek bant genişliğine bağlıdır ve yeni donanımlardaki floops/depolama oranının artması, daha fazla bellek bant genişliğine bağlı çekirdeklerin ortaya çıkmasına sebep olur. Bu nedenle, kodunuzdaki bant genişliği darboğazlarını hafifletmek için adımlar atmak oldukça önemlidir. Bu yazıda, CUDA C++ dilinde vektör yüklemeleri ve depolamaları kullanarak bant genişliği kullanımını artırmayı ve çalıştırılan talimat sayısını azaltmayı nasıl gerçekleştirebileceğimi göstereceğim.
Örnek Bellek Kopyalama Çekirdeği
Öncelikle basit bir bellek kopyalama çekirdeğine bakalım.
__global__ void device_copy_scalar_kernel(int* d_in, int* d_out, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
for (int i = idx; i < N; i += blockDim.x * gridDim.x) {
d_out[i] = d_in[i];
}
}
void device_copy_scalar(int* d_in, int* d_out, int N)
{
int threads = 256;
int blocks = min((N + threads-1) / threads, MAX_BLOCKS);
device_copy_scalar_kernel<<<blocks, threads>>>(d_in, d_out, N);
}
Bu kodda, daha önceki CUDA Pro İpucu gönderilerinde anlatılan grid-stride döngüleri kullanılarak kopyalama işlemi yapılmaktadır. Şekil 1, bu çekirdeğin GB/s cinsinden bant genişliğini kopyalama boyutuna göre göstermektedir.
Bu çekirdeğin montajını cuobjdump aracıyla inceleyebiliriz.
%> cuobjdump -sass executable
Scalar kopyalama çekirdeğinin SASS’ında, aşağıdaki talimatlar bulunmaktadır:
...
LDG.E R3, desc[UR6][R2.64] ;
...
STG.E desc[UR6][R4.64], R3 ;
...
Burada LDG.E ve STG.E talimatları, global bellekten sırasıyla 32 bit veriyi yükleyip depolar.
Vektör Yüklemelerle Performansı Artırmak
Bu işlemin performansını artırmak için, LDG.E.{64,128} ve STG.E.{64,128} gibi vektörleştirilmiş yükleme ve depolama talimatlarını kullanabiliriz. Vektörleştirilen yüklemeler, toplam talimat sayısını azaltır, gecikmeyi azaltır ve bant genişliği kullanımını artırır. Vektörleştirilmiş yüklemeleri kullanmanın en kolay yolu, CUDA C++ standart başlıklarında tanımlanan vektör veri türleri olan int2, int4, float2 veya float4 gibi türleri kullanmaktır. Bu türler, birkaç değeri tek bir veri biriminde paketler.
Bu türleri C++’ta reinterpret_cast<int2*>(d_in) ile kullanarak, int işaretçisini int2 işaretçisine dönüşüm yapabiliriz; bu işaretçi, iki int değerini tek bir birim olarak işler. C99’da ise aynı şeyi cast operatörünü kullanarak yapabilirsiniz: (int2*(d_in)).
Bunların işaretçilerini derefanslamak, derleyicinin vektörleştirilmiş talimatlar üretmesini sağlar.
int2* int2Ptr = reinterpret_cast<int2*>(d_in);
int2 data = int2Ptr[0]; // İlk iki int değeri bir int2 olarak yüklenir
Ancak, önemli bir nokta var: Bu talimatlar, hizalanmış verilere ihtiyaç duyar. Cihaz üzerinde tahsis edilen bellek, otomatik olarak veri türünün boyutuna katlanır şekilde hizalanır. Ancak, işaretçinizi kaydırmanız durumunda, kaydırmanın da hizalı olması gerekir. Örneğin, reinterpret_cast<int2*>(d_in+1) geçersizdir çünkü d_in+1, sizeof(int2) katlarına hizalanmamıştır. Ancak, hizalı bir offset kullanarak dizilere güvenle kaydırabilirsiniz; örneğin, reinterpret_cast<int2*>(d_in+2).
Yapılar kullanarak da vektörleştirilmiş yüklemeler oluşturabilirsiniz; ancak yapının boyutunun iki katına (power of two) düşmesi gerekmektedir.
struct Foo {int a, int b, double c}; // 16 byte boyutunda
Foo *x, *y;
…
x[i]=y[i];
Çift olmayan boyutlar, verilerin uyumlu olması açısından gereksiz paddingle sonuçlanabilir.
Vektör Yüklemeleri İçeren Bellek Kopyalama Çekirdeği
Artık vektörleştirilmiş talimatlar üretmeyi öğrendiğimize göre, bellek kopyalama çekirdeğini vektör yüklemeleri kullanarak değiştirebiliriz.
__global__ void device_copy_vector2_kernel(int* d_in, int* d_out, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
for (int i = idx; i < N/2; i += blockDim.x * gridDim.x) {
reinterpret_cast<int2*>(d_out)[i] = reinterpret_cast<int2*>(d_in)[i];
}
// Yalnızca bir iplikte, son elementi işler (varsa)
if (idx==N/2 && N%2==1)
d_out[N-1] = d_in[N-1];
}
void device_copy_vector2(int* d_in, int* d_out, int n) {
threads = 256;
blocks = min((N/2 + threads-1) / threads, MAX_BLOCKS);
device_copy_vector2_kernel<<<blocks, threads>>>(d_in, d_out, N);
}
Bu çekirdek yalnızca birkaç değişiklik içermektedir. İlk olarak, döngü şimdi N/2 kez çalışıyor, çünkü her tekrar iki öğeyi işliyor. İkincisi, kopyalamada yukarıda açıklanan dönüştürme tekniği kullanılmaktadır. Üçüncüsü, N’nin 2’ye tam bölünmemesi durumunda ortaya çıkabilecek kalan öğeleri işlemek için bir durum kontrolü eklenmiştir. Son olarak, önceki scalar çekirdeğimizden yarı kadar daha az iplik başlatıyoruz.
SASS’ı incelediğimizde şu değişiklikleri görüyoruz:
...
LDG.E.64 R2, desc[UR4][R2.64] ;
...
STG.E.64 desc[UR4][R4.64], R2 ;
...
Burada derleyici, LDG.E.64 ve STG.E.64 üretti. Geri kalan talimatlar aynı kalmıştır. Ancak, döngü yalnızca N/2 kez çalıştığı için, çalıştırılan talimat sayısında da yarı yarıya bir azalma olacaktır. Bu, talimat sınırı veya gecikme sınırlı bir çekirdek için oldukça önemli bir iyileştirmedir.
Dört Vektörlü Kopyalama Çekirdeği
Ayrıca, kopyalama çekirdeğinin bir vektör4 versiyonunu da yazabiliriz.
__global__ void device_copy_vector4_kernel(int* d_in, int* d_out, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
for(int i = idx; i < N/4; i += blockDim.x * gridDim.x) {
reinterpret_cast<int4*>(d_out)[i] = reinterpret_cast<int4*>(d_in)[i];
}
// Yalnızca bir iplikte, final öğelerini işler (varsa)
int remainder = N%4;
if (idx==N/4 && remainder!=0) {
while(remainder) {
int idx = N - remainder--;
d_out[idx] = d_in[idx];
}
}
}
void device_copy_vector4(int* d_in, int* d_out, int N) {
int threads = 256;
int blocks = min((N/4 + threads-1) / threads, MAX_BLOCKS);
device_copy_vector4_kernel<<<blocks, threads>>>(d_in, d_out, N);
}
Karşılık gelen SASS değişiklikleri şu şekildedir:
...
LDG.E.128 R4, desc[UR4][R4.64] ;
...
STG.E.128 desc[UR4][R8.64], R4 ;
...
Burada LDG.E.128 ve STG.E.128 üretilmiştir. Bu kod sürümü, talimat sayısını dört kat azaltmıştır. Üç çekirdek için genel performansı Şekil 2’de görebilirsiniz.
Hemen hemen her durumda, vektörleştirilmiş yüklemeler, scalar yüklemelerden daha tercih edilir. Ancak, vektörleştirilmiş yüklemelerin register baskısını artırdığı ve genel paralelliği azalttığı dikkat edilmelidir. Eğer çekirdeğiniz zaten register ile sınırlıysa veya çok düşük bir paralelliğe sahipse, scalar yüklemeleri tercih etmelisiniz. Ayrıca, daha önce belirtildiği gibi, işaretçiniz hizalanmamışsa veya veri türü boyutu iki katı (power of two) değilse, vektörleştirilmiş yüklemeleri kullanamazsınız.
Vektörleştirilmiş yüklemeler, mümkün olduğunda kullanmanız gereken temel bir CUDA optimizasyonudur; çünkü bant genişliğini artırır, talimat sayısını azaltır ve gecikmeyi düşürür. Bu yazıda, mevcut çekirdeklerinize vektörleştirilmiş yüklemeleri göreceğiniz gibi, oldukça az değişiklikle nasıl ekleyebileceğinizi gösterdim.
Bu blog yazısının bir versiyonu 4 Aralık 2013 tarihinde yayınlandı. Güncel GPU’larda görülen davranışları yansıtacak şekilde güncellenmiştir.