00 > Bir Block, Bir Thread

Bu derste olabilecek en basit CUDA programını çalıştıracaksın. Kernel tek bir block ve tek bir thread kullanır, yani hiç paralellik yok. Böylece işler karmaşıklaşmadan ilk çıktını görürsün. Sonraki her ders bu programı biraz değiştiriyor.

[!NOTE] Bu sayfadaki tüm çıktılar, Ubuntu 24 üzerinde CUDA 13.0 ile çalışan bir NVIDIA L40S'ten alındı.

GPU ve CPU

CPU'da bir fonksiyon tek bir çekirdekte bir kez çalışır. GPU'da ise bir kernel aynı anda birçok kez çalışır. Kernel, GPU'da çalışan bir fonksiyondur ve onun çalışan her kopyasına thread denir. Kaç thread çalışacağını iki sayı belirler. Biri block sayısı, diğeri block başına thread sayısı.

Örneğin 3'er thread'li 2 block, 2 x 3 = 6 thread başlatır ve bu 6 thread'in hepsi aynı kernel kodunu çalıştırır.

__global__ Ne Demek

__global__ void printIDs() { ... }

__global__, bir fonksiyonu GPU kernel'ı olarak işaretler. Derleyici onu CPU için değil, GPU için derler. Onu CPU çağırır ama GPU'da çalışır.

[!NOTE] İki niteleyici daha var. __device__ GPU'da çalışır ve yalnızca GPU kodundan çağrılabilir. __host__ normal bir CPU fonksiyonudur ve yalnızca CPU'dan çağrılabilir.

Başlatma Ayarı <<<blocks, threads>>>

printIDs<<<1, 1>>>();
//          ^  ^
//  blocks -+  +- threads per block

<<<...>>> sözdizimine çalıştırma ayarı denir ve fonksiyon adı ile argüman listesi arasına yazılır. İlk sayı block sayısı, ikinci sayı block başına thread sayısıdır. <<<1, 1>>>, tek thread'li tek bir block demektir ve toplam thread sayısı 1 x 1 = 1.

Thread, Block, Grid

Her kernel başlatması üç seviye oluşturur.

• Thread en küçük birimdir. Bir thread, kernel'ın bir kopyasını çalıştırır.
• Block, tek bir SM üzerinde çalışan bir grup thread'dir. SM, GPU'nun içindeki birçok küçük işlemciden biridir. Aynı block'taki thread'ler bellek paylaşabilir.
• Grid, bir kernel başlatmasındaki tüm block'lardır. Bir başlatma, bir grid.

[!NOTE] Bir GPU'da birçok SM vardır. Bu derslerde kullanılan L40S'te 142 tane var. Bir block asla iki SM'ye bölünmez, ama farklı block'lar aynı anda farklı SM'lerde çalışabilir. Ders 02 bundan yararlanıyor.

blockIdx.x ve threadIdx.x

printf("Block ID: %d  Thread ID: %d", blockIdx.x, threadIdx.x);

blockIdx.x, bu thread'in içinde bulunduğu block'un indeksidir. threadIdx.x, bu thread'in kendi block'u içindeki indeksidir. İkisi de 0'dan başlar. İkisinin de .x, .y ve .z parçaları vardır, çünkü grid'ler ve block'lar 1D, 2D ya da 3D olabilir. 1D işlerde yalnızca .x kullanırsın. <<<1, 1>>> ile ikisi de her zaman 0'dır.

Header Dosyaları

• cuda_runtime.h, CUDA runtime API'sidir. cudaDeviceSynchronize() ve hata kontrol fonksiyonları burada tanımlanır.
• stdio.h standart C'dir ve printf için gereklidir.
• device_launch_parameters.h, MSVC ya da bazı IDE'leri kullanırken blockIdx, threadIdx, blockDim ve gridDim'i editöre tanıtır. nvcc'nin buna ihtiyacı yoktur, ama zararı da yoktur.

cudaDeviceSynchronize()

Kernel başlatmaları asenkrondur. CPU kernel'ı başlatır ve hemen bir sonraki satıra geçer. cudaDeviceSynchronize() olmadan main() biter ve program, GPU daha hiçbir şey yazdıramadan kapanır. Bu fonksiyon, CPU'nun bütün GPU işleri bitene kadar beklemesini sağlar.

[!WARNING] cudaDeviceSynchronize()'ı unutursan program yine hatasız derlenir ve çalışır. Sadece hiçbir şey yazdırmaz.

Kod

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>

__global__ void printIDs()
{
    printf("\nBlock ID: %d  ===  Thread ID: %d", blockIdx.x, threadIdx.x);
}

int main()
{
    printIDs<<<1, 1>>>();
    cudaDeviceSynchronize();
    return 0;
}

• Üç #include satırı, yukarıda anlatılan header'ları yükler.
• printIDs kernel'dır. Her thread kendi block ID'sini ve thread ID'sini yazdırır. Metnin başındaki \n, her çıktının yeni bir satırda başlamasını sağlar.
• printIDs<<<1, 1>>>();, kernel'ı tek thread'li tek bir block ile başlatır.
• cudaDeviceSynchronize(); GPU'yu bekler, böylece yazdırılanlar program bitmeden ekrana gelir.

Kod Gezintisi

Programı, boş bir dosyadan başlayarak yazacağın sırayla adım adım geç.

Derle ve Çalıştır

İlk komut kodu derleyip bir programa dönüştürür. İkinci komut onu çalıştırır.

nvcc -o first_kernel first_kernel.cu
./first_kernel

• nvcc, CUDA derleyicisidir. Dosyanın hem CPU hem GPU kısmını derler.
• -o first_kernel, programa first_kernel adını verir. Bu olmazsa adı a.out olur.
• first_kernel.cu, yukarıdaki kodu içeren kaynak dosyadır. CUDA kaynak dosyaları .cu ile biter.
• ./first_kernel programı çalıştırır. ./, shell'e programı içinde bulunduğun klasörde aramasını söyler.

Çıktı

Block ID: 0  ===  Thread ID: 0

• Tek satır var, çünkü tek thread var ve her thread bir kez yazdırır.
• İki ID de 0, çünkü tek block ve tek thread 0 indeksini alır.
• Çıktı her çalıştırmada aynıdır, çünkü tek thread'in yarışacağı başka bir thread yoktur.

Kendin Yaz

GPU'dan seni selamlayan, tek thread'li bir kernel'ı sıfırdan yaz.

1. Aşağıdaki iskeletle bir hello.cu dosyası oluştur.
2. blockIdx.x ve threadIdx.x kullanarak Hello from block 0, thread 0 yazdıran hello kernel'ını yaz.
3. Onu tek thread'li tek bir block ile başlat ve CPU'nun onu beklemesini sağla.

#include "cuda_runtime.h"
#include <stdio.h>

// TODO: write the kernel hello() that prints its block ID and thread ID

int main()
{
    // TODO: launch hello with 1 block of 1 thread
    // TODO: wait for the GPU to finish
    return 0;
}

#include "cuda_runtime.h"
#include <stdio.h>

__global__ void hello()
{
    printf("Hello from block %d, thread %d\n", blockIdx.x, threadIdx.x);
}

int main()
{
    hello<<<1, 1>>>();
    cudaDeviceSynchronize();
    return 0;
}

Sözlük

• CUDA (Compute Unified Device Architecture): NVIDIA'nın, kendi kodunu GPU'da çalıştırmanı sağlayan platformu.
• GPU (Graphics Processing Unit, grafik işlem birimi): kernel'ları çalıştıran, binlerce küçük çekirdekli işlemci.
• CPU (Central Processing Unit, merkezi işlem birimi): ana işlemci. main()'i çalıştırır ve kernel'ları başlatır.
• kernel: GPU'da çalışan bir fonksiyon. Onu bir kez yazarsın, GPU onu aynı anda birçok thread'de çalıştırır.
• thread: en küçük çalışma birimi. Bir thread, kernel'ın çalışan bir kopyasıdır ve kendi ID'si vardır.
• block: tek bir SM üzerinde çalışan bir grup thread. Bu thread'ler paylaşımlı bellek üzerinden veri paylaşabilir.
• grid: tek bir kernel çağrısıyla başlatılan tüm block'lar.
• SM (Streaming Multiprocessor, akış çoklu işlemcisi): GPU'nun içindeki işlemcilerden biri. Block'lar SM'lerde çalışır.
• __global__: derleyiciye bu fonksiyonun bir GPU kernel'ı olduğunu söyler. CPU onu çağırır, GPU çalıştırır.
• blockIdx.x: o anki thread'in içinde bulunduğu block'un indeksi. 0'dan başlar.
• threadIdx.x: o anki thread'in kendi block'u içindeki indeksi. 0'dan başlar.
• cudaDeviceSynchronize(): CPU'nun, GPU tüm işini bitirene kadar beklemesini sağlar.
• asenkron: CPU beklemez. GPU'ya bir komut gönderir ve hemen yoluna devam eder.
• API (Application Programming Interface, uygulama programlama arayüzü): bir kütüphanenin sunduğu fonksiyonlar kümesi, burada CUDA runtime fonksiyonları.
• 3D (three-dimensional, üç boyutlu): x, y ve z boyutları olan bir şekil. 1D yalnızca x, 2D ise x ve y kullanır.
• derleyici (compiler): kaynak kodunu makinenin çalıştırabileceği bir programa dönüştüren araç. CUDA'da bu nvcc'dir.
• çalıştırma ayarı (execution configuration): kernel başlatmasındaki, kaç thread çalışacağını belirleyen <<<blocks, threads>>> kısmı.
• MSVC (Microsoft Visual C++): Microsoft'un Windows'taki C ve C++ derleyicisi.
• IDE (Integrated Development Environment, tümleşik geliştirme ortamı): Visual Studio gibi, derleme ve hata ayıklama araçları olan bir editör.