Tiny GPU: Verilog ile uygulanmış minimal GPU

 (github.com/adam-maj)
2 puan yazan GN⁺ 2024-04-27 | 1 yorum | WhatsApp'ta paylaş
  • tiny-gpu, GPU’nun donanım düzeyinde nasıl çalıştığını temelden öğrenmek için Verilog tabanlı minimal bir GPU uygulamasıdır; grafiklere özel donanımdan çok GPGPU ve ML hızlandırıcılarının ortak ilkelerine odaklanır
  • Uygulama, belgelenmiş 15’ten az Verilog dosyası, mimari ve ISA belgeleri, matris toplama/çarpma kernel’ları, kernel simülasyonu ve yürütme izi desteğinden oluşur
  • GPU aynı anda tek bir kernel çalıştırır; program belleği ve veri belleğini yükleyip thread_count değerini ayarladıktan sonra start sinyalini yükselterek kernel’ı başlatır
  • Basitleştirme amacıyla her çekirdek aynı anda bir block işler; her thread’in ALU, LSU, PC ve register file’ı vardır, ancak tüm thread’lerin her komuttan sonra aynı PC’ye yakınsadığı varsayılır
  • Modern GPU’lardaki çok katmanlı önbellek, paylaşımlı bellek, bellek birleştirme, pipelining, warp scheduling, branch divergence ve barrier gibi özelliklerin çoğu dışarıda bırakılarak öğrenme amaçlı yapı önceliklendirilmiştir

tiny-gpu’nun çözmeye çalıştığı sorun

  • CPU’lar için mimariden kontrol sinyallerine kadar öğrenilebilecek çok sayıda kaynak vardır; ancak modern GPU’ların düşük seviyeli teknik ayrıntıları, rekabetçi pazar nedeniyle büyük ölçüde özel mülkiyet olarak kalmıştır
  • GPU programlama konusunda çok kaynak olsa da, GPU’nun donanım düzeyinde nasıl çalıştığını öğrenmeye yönelik kaynak neredeyse yoktur
  • Açık kaynak GPU uygulamaları olan Miaow ve VeriGPU, işlevsel tamamlılığı ve çalışmayı hedeflediği için yapıları karmaşıktır
  • tiny-gpu, üretim sınıfı ekran kartlarının birçok karmaşıklığını ortadan kaldırır ve modern donanım hızlandırıcılarında ortak olan temel öğelere odaklanır
    • GPU mimarisinin önemli bileşenleri
    • SIMD programlama modelinin donanımda nasıl uygulandığı
    • GPU’nun sınırlı bellek bant genişliğiyle nasıl başa çıktığı

Genel mimari

  • tiny-gpu aynı anda yalnızca bir kernel çalıştıracak şekilde tasarlanmıştır
  • Kernel yürütme prosedürü şöyledir
    • Global program belleğine kernel kodunu yükleme
    • Gerekli verileri veri belleğine yükleme
    • Aygıt kontrol register’ına çalıştırılacak thread sayısını yazma
    • start sinyalini high yaparak kernel yürütmeyi başlatma
  • GPU şu birimlerden oluşur
    • Aygıt kontrol register’ı
    • Dispatcher
    • Değişken sayıda compute core
    • Veri belleği ve program belleği için bellek kontrolcüsü
    • Önbellek

Kernel yürütme ve thread dağıtımı

  • Aygıt kontrol register’ı, kernel yürütme metadata’sını saklar; tiny-gpu’da yalnızca çalıştırılacak toplam thread sayısı olan thread_count saklanır
  • Dispatcher, kernel başladığında thread’leri birden fazla compute core’a dağıtır
    • Paralel çalıştırılabilecek thread gruplarını block olarak oluşturur
    • İşlemek üzere uygun çekirdeklere block gönderir
    • Tüm block’ların işlenmesi bittiğinde kernel yürütmenin tamamlandığını bildirir
  • Basitleştirilmiş çekirdek aynı anda bir block işler
  • Her thread’in kendine ait ALU, LSU, PC ve register file’ı vardır
  • Thread komut yürütmesini bu kaynaklar üzerinde yönetmek, GPU’nun zor problemlerinden biridir

Bellek yapısı ve kontrolcü

  • GPU, harici global bellekle arayüz kuracak şekilde tasarlanmıştır; basitleştirme için veri belleği ve program belleği ayrılmıştır
  • Veri belleği özellikleri
    • 8 bit adresleme
    • Toplam 256 satır
    • 8 bit veri
    • Her satır 256’dan küçük bir değer saklar
  • Program belleği özellikleri
    • 8 bit adresleme
    • Toplam 256 satır
    • 16 bit veri
    • ISA’ya göre her komut 16 bittir
  • Bellek kontrolcüsü, çekirdeklerden gelen bellek isteklerini izler, istekleri gerçek harici bellek bant genişliğine göre sınırlar ve yanıtları doğru kaynağa iletir
  • Her bellek kontrolcüsü, global bellek bant genişliğine bağlı olarak sabit sayıda kanala sahiptir
  • Önbellek üzerinde çalışılan bir özelliktir; harici bellekten getirilen verileri aygıt içindeki SRAM’de saklayarak sonraki isteklerde daha hızlı getirmeyi ve bellek bant genişliğinin yeni veriler için kullanılmasını sağlar

Çekirdek iç yapısı

  • Her çekirdekte thread yürütmesini yöneten tek bir scheduler bulunur
  • tiny-gpu scheduler’ı bir block’un komutlarını sonuna kadar çalıştırdıktan sonra yeni bir block alır ve tüm thread’lerin komutlarını senkronize sırayla yürütür
  • Daha gelişmiş scheduler’larda pipelining ve warp scheduling ile kaynak kullanım oranı artırılabilir
  • Scheduler’ın başlıca kısıtı, global bellekten veri yükleme ve veri saklama sırasında oluşan gecikmedir
    • Çoğu komut senkron olarak çalıştırılabilir
    • LDR ve STR gibi load-store işlemleri asenkron olduğundan, komut yürütme uzun bekleme süreleri etrafında düzenlenmelidir
  • Fetcher, mevcut program counter’daki komutu program belleğinden asenkron olarak getirir
  • Decoder, getirilen komutu thread yürütmesi için kontrol sinyallerine çözer
  • Her thread’in register file’ı, hesaplama sırasında kullanılan verileri tutar ve SIMD desenini mümkün kılar
    • Salt okunur register’larda %blockIdx, %blockDim, %threadIdx bulunur
    • Kernel, yerel thread ID’sine göre farklı verilerle çalıştırılabilir
  • Her thread’in ALU’su ADD, SUB, MUL, DIV aritmetik komutlarını işler
  • CMP, iki register farkının sonucunun negatif, 0 veya pozitif olup olmadığını çıkarır ve sonucu PC biriminin NZP register’ına kaydeder
  • Her thread’in LSU’su global veri belleğine erişir ve LDR/STR ile asenkron bellek bekleme sürelerini yönetir
  • Her thread’in PC’si bir sonraki yürütülecek komutu belirler
    • Varsayılan olarak her komutta 1 artar
    • BRnzp, önceki CMP tarafından ayarlanan NZP register koşulu sağlanıyorsa belirli bir program belleği satırına dallanır
    • Döngüler ve koşullu ifadeler bu şekilde uygulanır
  • tiny-gpu, basitleştirme amacıyla tüm thread’lerin her komuttan sonra aynı PC’ye yakınsadığını varsayar
  • Gerçek GPU’larda tekil thread’ler farklı PC’lere dallanabilir; bu durumda birlikte işlenen thread grubu birden çok yürütme akışına ayrılır ve branch divergence oluşur

ISA

  • tiny-gpu, matris toplama ve matris çarpma gibi proof-of-concept niteliğindeki basit kernel’ları çalıştırmak için 11 komutluk bir ISA uygular
  • Desteklenen komutlar
    • BRnzp: NZP koşulu sağlanırsa başka bir program belleği satırına atlar
    • CMP: iki register değerini karşılaştırır ve sonucu NZP register’ına kaydeder
    • ADD, SUB, MUL, DIV: tensor matematiği için temel aritmetik işlemler
    • LDR: global bellekten veri yükler
    • STR: global belleğe veri yazar
    • CONST: register’a sabit değer yükler
    • RET: mevcut thread yürütmesini sonlandırma sinyali
  • Her register 4 bit ile belirtilir; toplam 16 register vardır
    • R0 ile R12 arasındaki 13 register, okunup yazılabilen serbest register’lardır
    • Son 3’ü SIMD için gereken %blockIdx, %blockDim, %threadIdx değerlerini sağlayan salt okunur özel register’lardır

Yürütme akışı

  • Her çekirdek komut yürütürken şu denetim akışı adımlarını izler
    • FETCH: Mevcut PC’nin sıradaki komutunu getirir
    • DECODE: Komutu kontrol sinyallerine çözer
    • REQUEST: LDR veya STR gerekiyorsa global bellekten veri ister
    • WAIT: Gerekirse global bellek yanıtını bekler
    • EXECUTE: Veri üzerinde hesaplamayı yürütür
    • UPDATE: Register file’ı ve NZP register’ını günceller
  • Bu denetim akışı basitlik ve anlaşılabilirlik için yapılandırılmıştır
  • Gerçek uygulamalarda bazı adımlar sıkıştırılarak işlem süresi optimize edilebilir veya pipelining ile çekirdek kaynaklarında birden fazla komutun yürütmesi koordine edilebilir
  • Her thread, kendine ait register file’daki veriler üzerinde aynı yürütme yolunu izleyerek hesaplama yapar
  • CPU diyagramlarına benzese de, %blockIdx, %blockDim, %threadIdx değerlerinin salt okunur register’larda bulunarak SIMD işlevselliğini mümkün kılmasıyla ayrılır

Örnek kernel’lar

  • ISA’nın proof-of-concept’i için matris toplama ve matris çarpma kernel’ları yazılmıştır
  • Depodaki test dosyaları, bu kernel’ları GPU üzerinde tamamen simüle edebilir ve veri belleği durumunu ve tam yürütme izini oluşturabilir

  • Matris toplama

    • matadd.asm, iki adet 1 x 8 matrisi toplar
    • 8 elemanın her biri için toplama ayrı bir thread’de gerçekleştirilir
    • SIMD programlamayı göstermek için %blockIdx, %blockDim, %threadIdx register’larını kullanır
    • Asenkron bellek yönetimini içermek için LDR ve STR komutlarını kullanır

  • Matris çarpma

    • matmul.asm, iki adet 2 x 2 matrisi çarpar
    • İlgili satır ve sütunun iç çarpımını eleman bazında hesaplar
    • Thread içi dallanmayı göstermek için CMP ve BRnzp kullanır
    • Tüm dallar yeniden yakınsadığı için mevcut tiny-gpu uygulamasında çalışır

Simülasyon

  • Kernel simülasyonunu çalıştırmak için iverilog ve cocotb gerekir
  • Hazırlık adımları
    • brew install icarus-verilog ve pip3 install cocotb ile Verilog derleyicisini ve cocotb’yi kurun
    • sv2v’nin en son sürümünü indirip arşivden çıkarın ve binary’yi $PATH’e ekleyin
    • Depo kökünde mkdir build çalıştırın
  • Kernel simülasyonları make test_matadd ve make test_matmul ile çalıştırılır
  • Çalıştırma sonuçları test/logs içindeki log dosyalarına yazılır
    • Başlangıç veri belleği durumu
    • Kernel’ın tam yürütme izi
    • Nihai veri belleği durumu
  • Her log dosyasının başında giriş matrisleri, sonundaki nihai veri belleğinde ise sonuç matrisi görünür
  • Yürütme izi, her cycle’da tüm çekirdeklerdeki tüm thread’lerin yürütme durumunu içerir
    • Mevcut komut
    • PC
    • Register değerleri
    • Durum bilgisi

Bilerek dışarıda bırakılan gelişmiş GPU özellikleri

  • tiny-gpu, basitleştirme amacıyla modern GPU’lardaki performans ve işlev geliştirmelerinin çoğunu dışarıda bırakır

  • Çok katmanlı önbellek ve paylaşımlı bellek

    • Modern GPU’lar global bellek erişimini azaltmak için birden fazla önbellek katmanı kullanır
    • tiny-gpu, istek yapan kaynak ile bellek kontrolcüsü arasında son verileri saklayan tek bir önbellek katmanı uygular
    • Çok katmanlı önbellekler, sık kullanılan verileri kullanım yerine daha yakın önbelleğe alarak yükleme süresini azaltır
    • GPU’lar, aynı block’taki thread’lerin paylaşılan sonuçları birbirine aktarabilmesi için paylaşımlı bellek de kullanabilir

  • Bellek birleştirme

    • Paralel çalışan birden fazla thread, matrisin bitişik elemanları gibi ardışık adreslere sıkça erişir
    • Bellek birleştirme, kuyrukta biriken bellek isteklerini analiz ederek bitişik istekleri tek bir işlemde birleştirir
    • Amaç, adresleme için harcanan zamanı azaltmak ve istekleri birlikte işlemektir

  • Pipelining

    • tiny-gpu’nun çekirdeği, bir thread grubunun bir komutunun yürütmesi bitmeden sonraki komuta başlamaz
    • Modern GPU’lar, bağımlılığı olan komutlarda sıralı yürütmeyi garanti ederken birden fazla ardışık komut yürütmesini stream eder
    • Asenkron bellek isteklerini bekleme gibi durumlarda çekirdek kaynaklarının boşta kalmaması için kaynak kullanım oranını artırır

  • Warp scheduling

    • Block, birlikte çalıştırılabilir thread partileri olan warp’lara bölünür
    • Bir warp beklemedeyken başka bir warp’ın komutlarını çalıştırarak tek bir çekirdekte birden fazla warp aynı anda işlenir
    • Pipelining’e benzer, ancak farklı thread’lerin komutlarını ele alır

  • Branch divergence

    • tiny-gpu, tek bir partideki tüm thread’lerin her komuttan sonra aynı PC’de olduğunu varsayar
    • Gerçekte, veriye bağlı olarak tekil thread’ler farklı satırlara dallanabilir
    • Farklı PC’ye sahip thread’ler ayrı yürütme akışlarına bölünür ve yeniden yakınsadıkları noktanın da yönetilmesi gerekir

  • Senkronizasyon ve barrier

    • Modern GPU’lar, aynı block’taki thread gruplarının belirli bir noktaya hepsi ulaşana kadar beklemesi için barrier koyabilir
    • Thread’lerin paylaşılan veri alışverişi yapması gerektiğinde, veri işlemenin tamamlandığını garanti etmek için kullanışlıdır

Sonraki çalışmalar

  • Gelecekteki iyileştirme maddeleri şunlardır
    • Basit bir komut önbelleği ekleme
    • Tiny Tapeout 7’de GPU’yu kullanabilecek bir adapter oluşturma
    • Temel branch divergence ekleme
    • Temel bellek birleştirme ekleme
    • Temel pipelining ekleme
    • Cycle süresini iyileştirmek için denetim akışı ve register kullanımını optimize etme
    • Grafik işlevlerini göstermek için temel grafik kernel’ı yazma veya basit grafik donanımı ekleme
  • Depoyu iyileştirmek isteyen kullanıcılar PR ile katkıda bulunabilir

İlgili okumalar

1 yorum

 
GN⁺ 2024-04-27
Hacker News yorumları
  • GPU pazarı son derece rekabetçi olduğundan, modern mimarilerin düşük seviyeli teknik ayrıntılarının çoğu gizli kalıyor.
    İstisna olarak Intel, çok sayıda GPU teknik dokümanı yayımlıyor: https://kiwitree.net/~lina/intel-gfx-docs/prm/
    i810/815 kılavuzları da internette bulunabiliyor; 855/910/915/945'in eksik olduğu 965 öncesi dönemdeki tuhaf boşluk dışında dokümantasyon oldukça istikrarlı sayılır.
  • Gerçekten harika bir proje; böyle donanım projelerinin açık şekilde yürütüldüğünü görmek güzel.
    Ancak bence bu daha çok bir SIMD yardımcı işlemcisine yakın.
    GPU diyebilmek için en azından bir tür görüntü çıkışı olması gerektiğini düşünüyorum.
    Son dönemde Nvidia ve benzerlerinin yalnızca sunucuya yönelik grafik mimarisi türevlerini de GPU diye satmasıyla terimin epey gevşediğini biliyorum; ancak GPU tasarımında grafik kısmı hâlâ karmaşıklığın önemli bir bölümünü oluşturuyor.
    • Grafik işliyorsa, çıkış olmasa bile GPU sayılabileceğini düşünüyorum.
      Çıkış vermeyen bir GPU da hâlâ işe yarar.
      İş yerimde orta seviye Quadro bulunan yaklaşık 75 iş istasyonu var; kartlarda yalnızca mini-DisplayPort var, şirket ise sadece HDMI kablosu aldığı için hepsi entegre grafiğe bağlı.
      Yine de o kartlar yazılımı hızlandırıyor ve grafik işliyor; sadece ekrana görüntü vermiyorlar.
  • Güzel. Açık çekirdek GPU çalışmalarını güçlü biçimde destekliyorum.
    Başka bir örnek de var: https://github.com/jbush001/NyuziProcessor
    • Bu tür açık çekirdek işlemcilerden biri için asgari bir CUDA uygulaması olsa iyi olurdu.
      TSMC veya başka bir foundry'de böyle bir işlemciyi ekonomik olarak üretmek için ne kadar hacim gerekir?
  • Gerçekten mükemmel bir proje.
    FPGA denemek istiyorum ama açıkçası nereden başlayacağımı kestirmek bile zor; tüm alan da epey göz korkutucu geliyor.
    Nihai hedefim LLM'ler için bir hızlandırıcı kart yapmak; tamamen keyfi belirlenmiş bir hedef olsa da bu projeyle çok kesişen yönleri olacak gibi, muhtemelen yalnızca daha büyük modelleri yüklemek için bellek offloading kısmı farklı olur.
    • Zihnindeki çerçeveyi değiştirmen gerekiyor.
      FPGA'ya giriş birden fazla alt tekniğe bölünmeli ve beklentiler de ayarlanmalı.
      Bir yazılım mühendisinden en başta ilkelerden yola çıkarak tüm bir bilgisayar yapmasını, komut kümesi mimarisi yazmasını, makine dilini anlamasını, bunu assembly'ye dönüştürmesini ve Python koduyla uygulama geliştirmek için bir programlama dili bile geliştirmesini beklemeyiz.
      En üstten başlayıp yığının aşağısına inmek doğru yaklaşım.
      Karmaşıklığı soyutlayıp hazır IP'lerle sistem kurmaya odaklanırsan FPGA tasarımı oldukça kolaydır.
      Genelde MATLAB gibi şeyler önerilir; çünkü reference design'ı olan bir DevKit üzerinde HDL Coder ile ilk uygulamanı oluşturabilirsin.
      Aksi halde dijital hesaplama mimarisi, Verilog, timing, transceiver/giriş-çıkış, pin planlama, Quartus/Vivado, simülasyon/doğrulama, gömülü sistemler vb. öğrenmek gibi devasa bir yük ortaya çıkar.
      Özetle sistem düzeyi tasarımdan başlayıp tak-çalıştır IP'leri getirerek en üst seviyede bağlamayı öğrenmeli, ardından o modülü hazır bir referans tasarıma yerleştirmelisin.
      Sonra katmanları yavaş yavaş soyarak alttaki karmaşıklığı görünür hâle getirebilirsin.
    • Ben de aynı durumdayım ve planım şöyle:
      1. Harris, Harris'in Digital Design and Computer Architecture. (2022). Elsevier kitabını okumak: https://doi.org/10.1016/c2019-0-00213-0
      2. Yazarın RVFpga kursunu izleyerek FPGA üzerinde gerçek bir RISC-V CPU yapmak: https://www.youtube.com/watch?v=ePv3xD3ZmnY
    • Şu yolu öneririm:
      1. Eğitim amaçlı depo https://github.com/yuri-panchul/basics-graphics-music'i klonlamak. Sıfırdan Verilog öğrenenler için basit alıştırmalar derlemesi; GPU geliştirme için Imagination'da çalışmış Yuri Panchul tarafından yazılmış.
      2. Desteklenen onlarca FPGA kartından birini ve tuş, LED gibi aksesuarları edinmek.
      3. Yosys ve ilgili araçları kurmak.
      4. lab01 DeMorgan'dan başlayıp depodaki alıştırmaların olabildiğince çoğunu yapmak.
        Harris&Harris'i okurken alıştırmaları paralel yürütebilirsin.
        Alıştırmaları ve kitabı bitirdiğinde kendi projenize başlama zamanı gelmiş olur.
        Bu arada HackerMojo'da haftalık buluşmalar da yapılıyor; Valley'de olmasan bile Zoom üzerinden katılabilirsin.
    • Hangi aşamada olduğunu bilmiyorum ama dijital mantığı ve CPU/GPU mimarisini daha iyi anlamamda şu kaynaklar yardımcı oldu:
      1. https://learn.saylor.org/course/CS301
      2. https://www.coursera.org/learn/comparch
  1. https://hdlbits.01xz.net/wiki/Main_Page
  • LLM'leri hızlandırmak istiyorsanız önce mimariyi bilmeniz gerekir
    Oradan başlayabilirsiniz
    Donanım aslında kolay kısım da sayılır; üretim tarafında ise zor kısımdır
  • Buradaki sıralı always bloğunda non-blocking atama ve blocking atama işleçlerini karıştırarak kullanmanın bir nedeni var mı?
    • Bu yerel değişken gibi görünüyor
    • Simülasyon ve sentez sonuçlarının birebir örtüşmesine fazla takılmıyorsanız böyle yapmak sorun değil
  • Uzun zaman önce VHDL ile buna benzer bir şey yapmıştım
    Birçok açık kaynak HDL projesinin toplandığı opencores adlı bir site vardı
    Günümüzde HPC düzeyinde, büyük ölçekli dağıtık HDL simülatörleri arasında iyi bir seçenek var mı merak ediyorum
    RTL düzeyi simülasyonda modern GPU'lardan yararlanmak mantıklı görünüyor
  • ALU, DIV komutunu donanım düzeyinde doğrudan mı uyguluyor?
    Modern CUDA çekirdekleri gibi yerlerde bölme işleminin gerçek bir komut olarak bulunması normal mi, yoksa genelde yazılımla mı emüle edilir?
    Gerçek donanım bölme devresi çok fazla alan kapladığı için GPU ALU'sunda olmasını beklemezdim
    Verilog'da DIV: begin alu_out_reg <= rs / rt; end şeklinde tek satır yazmak çok kolay, ama o tek satır silikonda çok yer tüketir
    Sadece Verilog'u simüle ediyorsanız bu gerçeği göremeyebilirsiniz
    • Bu sadece birinin Verilog öğrenme projesi
      Proje simülasyonda duruyor; gerçek donanıma dönüştürmek için çok daha fazla iş gerekir
  • Yine grafik işlevi olmayan bir “GPU”
    Kişisel olarak böyle şeylerin başka bir adla anılması gerektiğini düşünüyorum
    • İlk soru zaten CPU ile GPU'nun neden ayrıldığı
      İkisi arasındaki mesafe daralıyor ve her iki taraf da birbirinin özelliklerini ekliyor, ancak hâlâ ciddi farklar var
      Bence bunun Amdahl yasasıyla ilgisi var [0]
      Bu anlamda CPU'ya gecikme için optimize edilmiş işlemci, GPU'ya ise aktarım hacmi için optimize edilmiş işlemci denebilir
      Daha somut olarak [1] CPU'yu uzun ve derin veri bağımlılığı işlemcisi, GPU'yu ise geniş ve düz veri bağımlılığı işlemcisi olarak da adlandırabiliriz
      [0]: https://en.wikipedia.org/wiki/Amdahl%27s_law
      [1]: https://en.wikipedia.org/wiki/Data_dependency
    • TPU, yani tensör işleme birimi denebilir
      Tensör yalnızca n boyutlu bir dizidir
      Üzerine yazılım veya firmware koyup GPU gibi davranmasını sağlayabilirsiniz
    • Bir 'ekran bağdaştırıcısı' yapma projesine başlamayı düşünüyordum, ama daha başlamadan UEFI'nin GOP sürücüsü ile ekran bağdaştırıcısı arasındaki iletişim protokolünü çözemediğim için takıldım
      EDK2 kaynaklarından parçaları birleştirmeye çalıştım, ancak ne kadarının QEMU'ya özgü olduğu belirsiz
    • MPU, yani matris işleme birimi demek yeterli
    • Yerleşmekte olan terimin AIA, yani AI hızlandırıcı olduğunu düşünüyorum
  • tiny-gpunun tüm iş parçacıklarının her komuttan sonra aynı program sayacında “yakınsadığını” varsayması fazla safça bir basitleştirme
    Gerçek GPU'larda tek tek iş parçacıkları farklı PC'lere dallanabilir ve başlangıçta birlikte işlenen iş parçacığı gruplarının ayrı yürütmelere bölündüğü dal sapması ortaya çıkar
    Silikondan GPU yapmadan önce GPU programlamayı denemek iyi olurdu
    Üstelik SIMD demek de pek doğru gelmiyor
    Bu kişi, daha önce başkalarının devrelerini birbirine bağlayıp LED'i yanıp söndürerek CPU yaptığını söyleyen kişi
    • İlki, her yürütmede __syncthreads() çağırmakla aynı şey sayılmaz mı?