Skip to content

CUDA NEDİR?

Jump to bottom
Nezihe Sözen edited this page Mar 24, 2019 · 3 revisions

1. GİRİŞ

Grafik İşlem Birimi(GPU)’nin temel görevi bilgisayarda oluşturulan görüntülerin ekrana verilmesini sağlamaktır.Bu yüzden ilk GPU’lar sadece bu görevi yerine getirmekteydi. Zaman içerisinde Merkezi İşlem Birimi(CPU)’nin karşılaşılan büyük hesaplama problemlerinde yetersiz kaması üzerine GPU’nun donanımsal paralelliğinden yararlanma fikri ortaya çıkmıştır. GPU'ların programlanabilir bir arayüze sahip olmasını ve yüksek seviyeli dillerle programlanabilmesini sağlamak için GPGPU modeli oluşturulmuştur.

CUDA(Compute Unified Device Architecture), NVIDIA firmasının 2006 yılında GPU'nun donanımsal hesaplama gücünden faydalanmak amacıyla sunduğu paralel hesaplama mimarisidir. Linux, Windows ve Mac Osx platformları üzerinde çalışabilmektedir.FORTRAN, C/C++ ve Python gibi dilleri destekleyen bir API’dir.Rakiplerine göre avantajları paylaşımlı bellek kullanımı, GPU'dan daha hızlı veri okuma ve bit düzeyinde işlem yapılabilmesine olanak sağlaması olarak sayılabilir.

GPU'nun CPU'dan farkı, SIMD(Single Instruction Multiple Data-Tek Komut Çoklu Veri) mimarisine sahip olmasıdır.Alttaki şekilde CPU(soldaki) ve GPU(sağdaki)’ların donanımsal olarak karşılaştırılması gösterilmiştir. CPU hesaplamalarını seri bir şekilde gerçekleştirirken, GPU hesaplamalarını yapısal olarak paralel olması sebebiyle paralel bir şekilde gerçekleştirmektedir.

2. İŞ AKIŞI

CUDA mimarisinde geliştirilen uygulamalar sadece GPU üzerinde çalışmazlar. Öncelikle CPU tarafından kontrol edilen ana bellek üzerinden grafik kartı üzerindeki belleğe kopyalanması gereklidir. GPU belleğindeki veri CUDA iş parçacıkları tarafından yürütülerek paralel olarak hesaplanması tamamlanır ve ardından tekrar ana belleğe gönderilerek işlem sonlandırılır. Bu akış modeli alttaki şekilde görülmektedir.

CPU üzerinde çalışan kod parçaları ile GPU üzerinde çalışan kod parçaları farklıdır. CUDA mimarisi ile uygulama geliştirillirken bu yapıya uygun bir şekilde kod yazılır. Alttaki şekilde CUDA ile geliştirilmiş bir programın genel çalışma mantığı ifade edilmektedir. Burada “Host” CPU olarak ve “Device” ise GPU olarak düşünülmelidir. Seri kodlar CPU üzerinde yürütülür. “Kernel” olarak adlandırılan paralel kod parçaları ise GPU üzerinde yürütülür. Host üzerinde çalıştırılan komut tektir ve aynı anda tek bir thread(iş parçacığı) işlenir. Device üzerinde çalıştırılan komutlar ise birden fazla thread’e bölünerek işlenir.

ÖLÇEKLENEBİLİR PROGRAMLAMA

Gelişen GPU teknolojisi ile birlikte donanımsal farklılıklar da ortaya çıkmaktadır. CUDA mimarisi, NVIDIA’nın Geforce 8800 ve sonrasında çıkan modellerinde desteklenmektedir. Piyasaya çıkan her yeni ürün öncekilerden farklı çekirdek sayılarına sahiptir. Ancak CUDA ile yazılmış bir uygulama , GPU çekirdekleri ile ilgili bir işlem yapmak zorunda değildir. Bu özelliğe GPU'nun ölçeklenebilirlik özelliği denilmektedir. Alttaki şekilde çift çekirdekli bir GPU ile 4 çekirdekli bir GPU'nun programa nasıl ayak uydurdukları görülmektedir.

CUDA PROGRAMLAMA MİMARİSİ

1. Kernel

CUDA’da geliştirilen bir kodun GPU tarafında çalışacak olan kısmına “kernel” adı verilmektedir. GPU, veri kümesinin her bir elemanı için birer kernel kopyası oluşturulur. Bu kernel kopyalarına “thread” adı verilmektedir. Kernel kodu Host tarafından çağrılır ve Device üzerinde yürütülür. Kernel kodu “global” ile nitelendirilir. Kernel, thread dizilimleri(thread arrays) tarafından çalıştırılır.

2. Thread

Thread, CUDA mimarisindeki en küçük programlama parçacığıdır. Block’lar içerisinde 1B’lu, 2B’lu veya 3B’lu olabilirler. Kendi aralarında eş zamanlı olarak aynı kod parçasını çalıştırmaktadırlar. Thread’ler Block’lar içerisinde dizilerek gruplanırlar. Farklı Block’lardaki Thread’ler ortak çalışmazlar.

Her bir thread’in kendine ait Program Sayacı(PC), kaydedicisi(register) ve durum kaydedicisi(State Register) vardır. Her thread’in Block içerisinde kendine ait bir ID(kimlik,indis)’si vardır. Bu indisler; “threadIdx.x”, “threadIdx.y” ve “threadIdx.z” şeklinde ifade edilmektedir.

3. Block

Block yapısı, paralel olarak çalışan thread’lerden oluşmaktadır. Grid içerisinde tektirler. Grid içerisinde 1B’lu, 2B’lu veya 3B’lu olabilirler. Grid’ler içerisinde dizilerek gruplanırlar. Her bir Block’un , Grid içerisinde kendine ait bir indisi vardır. Bu indisler “blockIdx.x”, “blockIdx.y” ve “blockIdx.z” şeklindedir. İçerisinde barındırdığı thread satır ve sütun sayısına göre boyutlanırlar ve bu boyutlar “blockDim.x”, “blockDim.y” , “blockDim.z” terimleri ile ifade edilmektedir. “blockDim” terimi “threadsPerBlock” terimi ile denktir ve her bir block’taki thread anlamına gelmektedir.

4. Grid

Grid, Block’ların bir araya gelerek oluşturdukları yapılardır. Her bir Kernel çağrısı bir Grid oluşturur. Yani paralel olarak çalıştırılan kod parçası, kernel çağrısı ile Device üzerinde çalıştırılırken her bir kopyası için bir Grid oluşturulmaktadır. Grid’ler 1B’lu veya 2B’lu olabilirler.Hesaplama kapasitesi 2.0 ve yukarısı sürümlerde gridler 3B’lu da olmaktadırlar. Bu boyutlar “gridDim.x” , “gridDim.y” ve “gridDim.z” şeklinde ifade edilmektedir. “gridDim” terimi ile “blocksPerGrid” terimi denktirler. Alttaki şekilde thread hiyerşisi görülmektedir. Her bir Grid Block’lardan, her bir Block ise Thread’lerden oluşmaktadır.

Aşağıdaki 2x2’lik Block ve 4x3’lük Thread yapısı göz önüne alınırsa:

  • gridDim.x = 3 (block’ların sütun sayısı kadar)

  • gridDim.y = 2 (block’ların satır sayısı kadar)

  • blockDim.x= 4 (thread’lerin sütun sayısı kadar)

  • blockDim.y= 3 (thread’lerin satır sayısı kadar) verileri elde edilir.

CUDA HAFIZA ORGANİZASYONU

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.

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.

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.

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.

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.
  1. Constant memory “cached” bir hafıza alanıdır, bu yüzden aynı adresten okumalar herhangi bir ek bellek trafiğine maruz kalmayacaktır.
  1. Ö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.

BELLEK ERİŞİM FONKSİYON NİTELEYİCİLERİ

1. global

Kernel kodu , global ile nitelendirilir. main() içerisinde kernel çağrısı yapılırken kernel fonksiyon adı yazılır ve “<<< >>>” açılı parantezleri yazılır. Bu yazım şekli CUDA’ya özgüdür. Açılı parantezler arasına yazılan ifadeler sırasıyla blocksPerGrid(her bir grid’teki block sayısı) ve threadsPerBlock(her bir block’taki thread sayısı) bilgilerini götermektedir.Kısacası bir kernel çağrısı:

Kernel_fonksiyon_adı<<<blocksPerGrid,threadsPerBlock>>>(fonk_alacağı_parametreler); şeklindedir.

global ile nitelenmiş fonksiyonunun dönüş tipi void olmalıdır. Bu fonksiyon MİB tarafından çağrılabilir, GİB tarafından çağrılamaz. Başka bir global fonksiyonu tarafından da çağrılamaz. GİB üzerinde işlem yürütmektedir.

“global” ile nitelenmiş fonksiyonlar asenkron olarak çalışırlar. Yani aygıt(GPU) üzerinde çalışması bitmeden döner.

2. device

GPU üzerinde çalışırlar ve CPU tarafından çağrılamazlar. Yalnızca GPU tarafından çağrılabilir.

3. host

Nitelediği fonksiyonun sadece CPU tarafından çağrılabildiğini ve sadece CPU üzerinde çalıştırılabildiğini ifade etmektedir. Varsayılan fonksiyon niteleyicisi “host” olarak belirlenmiştir.Yani hiçbir niteleyici yazılmasa bile direkt olarak “host” niteleyicisi olduğu varsayılmaktadır. Bazı durumlarda “device” niteleyicisi ile birlikte kullanımları da vardır. “global” niteleyicisi ile birlikte kullanımı sözkonus değildir.

DEĞİŞKEN NİTELEYİCİLERİ

1. device

GİB üzerinde konumlandırılmış bir değişkeni niteler.Bir uygulamada yaşam süresi bellidir. Global bellek üzerinde yer almaktadır. Tüm thread’ler tarafından yürütülebilir.

2. constant

Seçime bağlı olarak “device” niteleyicisi ile birlikte kullanılabilir. Değişmez bellek üzerinde yer alan, belli bir yaşam süresi olan uygulamalarda yer alan ve tüm thread’ler tarafından erişilebilen değişkenleri tanımlamaya olanak sağlar.

3. shared

Seçime bağlı olarak “device” niteleyicisi ile birlikte kullanılabilir. Paylaşımlı bellek üzerinde bir thread block’unda yer alan, belli bir yaşam süresi olan block içerisinde bulunan ve sadece aynı block üzerinde yer alan thread’lerin erişebildiği değişkenleri nitelemektedir.

BELLEK YÖNETİMİ

Bir CUDA uygulamasında CPU ve GPU ayrı bellek alanlarını kullanırlar. Programın çalıştırılması sırasında bu bellek alışverişinin sağlanması için fonksiyonlar geliştirilmiştir. Alttaki şekilde “Host(CPU)” ile “Device (GPU)” arasındaki veri aktarımı görülmektedir.

1. Hafıza Ayırma Fonksiyonları

cudaMalloc (void ** pointer, size_t nbytes);

C dilindeki Malloc ile eşdeğerdir, ancak hafıza GİB üzerinde ayrılır .

cudaMemset (void * pointer, int value, size_t count);

İstenilen değer, istenildiği kadar bellek alanına yerleştirilebilir.

cudaFree (void* pointer);

C dilindeki Free ile eşdeğerdir, ancak burada GİB üzerinden ayrılan hafıza serbest bırakılır.

2. Veri Kopyalama Fonksiyonları

cudaMemcpy( void *dst, void *src, size_t nbytes, enum cudaMemcpyKind_direction);

Burada “dst” ile kopyalanacak verinin hedefi, “src” ile kopyalanacak verinin kaynağı, “nbytes” ile kopyalanacak verinin boyutu ve “cudaMemcpyKind_direction” ile kopyalama yönü belirtilmektedir. “enum” veri tipi ile ifade edilen değişkenler daha önceden belirlenmiş kopyalama yönleridir. Buna göre veri alışverişinin yönü CPU'dan GPU'ya, GPU'dan CPU'ya ya da GPU’nun kendi içerisinde olmaktadır. CPU’dan CPU’ya bir yön bulunmamaktadır çünkü bu durumda uygulama heterojen değil klasik bir CPU uygulaması (seri hesaplama) olmaktadır.

  • cudaMemcpyHostToDevice (CPU->GPU)
  • cudaMemcpyDeviceToHost (GPU->CPU)
  • cudaMemcpyDeviceToDevice(GPU->GPU)

Kaynaklar

  1. “CUDA Teknolojisi” http://www.hwa.com.tr/nvidia/wpcontent/uploads/2013/08/cuda_teknolojisi_3.jpg
  2. NVIDIA,(2012) “NVIDIA CUDA C Programming Guide Version 4.2”
  3. AKÇAY, M., ŞEN, B., ORAK, İ.M., ÇELİK,A. (2011), “Paralel Hesaplama ve CUDA”, 6. Uluslar arası İleri Teknolojiler Sempozyumu (İATS’11)
  4. KIRK, D. ve HWU, W.W. (2010) “Programming Massively Parallel Processor”, Morgan Kaufmann Publishers
  5. BRADLEY, T. (2010) ,“Advanced CUDA Optimization:1-Introduction”, NVIDIA Corporotion
  6. “Introduction to CUDA 5.0” http://www.3dgep.com/introduction-to-cuda-5-0/
  7. “Constant Memory in CUDA” , http://cuda-programming.blogspot.com.tr/2013/01/what-is-constant-memory-in-cuda.html
  8. SANDERS, J. ,KANDROT, E. (2010) “CUDA by Example:An Introduction to General-Purpose GPU Programming”, NVIDIA Corporotion
Clone this wiki locally