CUDA koşut programlamanın tüm zorluklarına bir de ekran kartında çalışmanın getirdiği zorlukları da eklediğinden bazen gerçekten can sıkıcı sorunlara sebep olabiliyor. Bunları tanımlamak sorunları çözmekten daha uzun vakitler alabiliyor. Bu sebeple CUDA derleyicisi nvcc’nin bazı keyfi tercihlerini bilmek oldukça önemli, çünkü bunları ancak ve ancak derlediğiniz kodun ptx çıktısına baktığınız zaman farkedebiliyorsunuz ve hatalı bir gidiş yolu geri dönüşleri oldukça zorlaştırıyor.

Burada benim karşılaştığım bir sorun üzerinden gideceğim. Öncelikle sorunu anlayabilmek için CUDA mimarisine biraz giriş yapmak lazım. CUDA farklı bloklar içine bölünmüş threadleri koşut olarak işletebilmektedir. Her bir blok farklı bir multiprocessor’a atanmaktadır ve blok içinde zaman uyumlama yapılması mümkündür. Fakat bloklar arası zaman uyumlama gerçekleştirimi mevcut mimaride yoktur. Çünkü grafik işlemede böyle bir ihtiyaç olmadığı gibi performans düşürücü etkileri mevcuttur. Her blok içerisinde en fazla 512 thread, GPU üzerinde de en fazla 65535 ader blok çalıştırılabilmektedir. Fakat GPU üzerinde genel programlama yapacağınız zaman bu bloklar arası bariyerlere ihtiyaç doğmaktadır.

Konuya döndüğümüzde böyle bir bariyer gerçekleştirilmesi için kullanılabilecek ilk yaklaşım işi bitmiş threadleri, diğer threadleri beklemeleri için sonsuz döngüye sokmaktır(bkz. Spinlock). Fakat burada gerçekleştirilecek sonsuz döngü nvcc tarafından gereksiz olarak yorumlanarak daha derleme aşamasında kodunuzdan atılmaktadır.

1
2
3
4
5
6
7
8
9
10
11

__device__ int sense = 1;
 
__device__ void barrier(){
 
	int local_sense=0;
 
	// işlem yaptığınız kısım
 
	while(sense != local_sense);
 
}

Üstteki gibi bir kernel içinde yerel bir değişkenle kurulmuş sonsuz döngünüz hiçbir zaman çalışmayacaktır. nvcc bu yerel değişkenin değerinin thread içerisinde değiştirilmeyeceğini öngörecek ve kod parçasını atacak, global tanımlanmış değişkeni de tekrar okuma zahmetine girmeyecektir. Bunu engellemek için tek yapılması gereken yerelde tanımlanmış değişkenin tanımını alttaki gibi değiştirmektir.

1

int volatile local_sense=0;

Bu değişiklik yapıldığı vakit nvcc bu yerel değişken için her okumayı ve yazmayı anında gerçekleştirecek ve sonsuz döngüyü işletecektir. Tabi ki bu küçük ayrıntılar bazen o kadar büyük sonuçlar doğuruyor ki, dikkatten kaçmaması gereken öneme sahipler. CUDA belgelerinden tek edinebileceğiniz bilgi, yerel değişkenlerin volatile ile tanımlanmadıkları koşulda yazma ve okumalarının derleyici tarafından optimize edilebileceği, fakat sonsuz döngünün derleme esnasında yok sayılacağına dair bir bilgi yok.

Leave a Reply