MATLAB ile GPU Hesaplama

“MATLAB ile GPU Hesaplama” başlıklı bu sayfayı sürekli güncelleyerek MATLAB ile CUDA destekli grafik işlemcilerle GPU Hesaplama yapmak isteyen kişilere giriş seviyesinde bilgi vermeyi amaçlamaktayım.

gpuDeviceCount komutu ile sistemimizde CUDA destekli kaç tane grafik kartı olduğunu sorguluyoruz. Genellikle 1 tanedir, bazı sistemlerde birden çok kart takılarak, bu kartları kullanarak hesaplama yapılırken hangi karta yönelik işlem yapılacağı seçilmesi gerekir. Eğer birden fazla ise hangi kartın default olduğunu da belirtmek gerekir. Veya her işlemden önce hangi karta yönelik işlem yapılacağı ilgili kartın numarası kullanılarak belirtilebilir.

gpuDevice komutu sistemimizdeki CUDA destekli grafik kartın programlama esnasında bizim için gerekli özelliklerini veren başlangıç komutudur.

Komutun çıktısı benim grafik kartım için aşağıdaki şekilde olmuştur:

Burada;

Name: ‘GeForce GT 820M’ kartımızın ismini vermektedir.
Index: 1 1 tane kartımız olduğundan 1 numaralı indekse sahip kart olduğunu belirtiyor. Eğer 2 tane olsa idi 2.kartı seçmek için gpuDevice(2) şeklinde bir komut kullanacaktık.
ComputeCapability: ‘2.1’ CUDA ile programlama yapacak olanlar için en önemli parametrelerden birisidir. Zira bu hesaplama kapasitesi parametresi sürekli geliştiriliyor ve buna bağlı olarak bazı özelliklerde değişiyor. Örneğin dinamik paralellik ComputeCapability: ‘3.5’ ve sonrası sürümler için geçerlidir. 25 Şubat 2017 itibariyle en yüksek ComputeCapability: ‘6.1’ şeklinde bildirilmiştir. Güncel durumu aşağıdaki bağlantıdan kontrol edebilirsiniz:
https://developer.nvidia.com/cuda-gpus
SupportsDouble: 1 bilgisi kartın double hassasiyetli sayıları desteklediğini belirtmektedir.
DriverVersion: 8 sürücü versiyonunu göstermektedir.
ToolkitVersion: 7.5000 sistemde yüklü olan CUDA versiyonunu göstermektedir.
MaxThreadsPerBlock: 1024 bir block‘taki maksimum çalışabilecek thread sayısını belirtmektedir.
MaxShmemPerBlock: 49152 bir bloğun kullanabileceği shared memory boyutunu belirtir. 49152 byte / 1024 = 48KB.
MaxThreadBlockSize: [1024 1024 64] her bir boyuttaki block’ta bulunabilecek maksimum thread sayısı. 3 boyutlu block’lar tanımlanabileceğinden bu limitlere dikkat edilmesi gerekmektedir.
MaxGridSize: [65535 65535 65535] her boyuttaki maksimum grid boyutunu belirtmektedir.
SIMDWidth: 32 eşzamanlı işlem gören thread sayısını bildirmektedir.
TotalMemory: 2.1475e+09 grafik kartın toplam belleğini byte cinsinden vermektedir.
AvailableMemory: 1.7375e+09 grafik kartın kullanılabilecek bellek miktarını byte cinsinden vermektedir. Çalışırken zaman zaman kontrol edilmeli, oluşturulan değişkenlerin manuel olarak silinmesi gerekiyorsa silinmelidir. Bazı durumlarda hata çıktısı üretmektedir. Bu gibi durumlarda kod ile GPU resetlenmeli, bu fayda vermezse sistem kapatılıp açılmalıdır 🙂
MultiprocessorCount: 2 fiziksel olarak grafik kart üzerinde bulunan çokişlemcili yapıya işaret etmektedir. Buradaki işlemciler vektör işlemcidir.

Veriyi işleme şekli olarak SIMD (tek işlem çok veri) olarak tasarlanan vektör işlemcileri, en kısa tanımıyla, bir işlemcinin (processor), vektör şeklindeki ardışık veriler üzerinde bir adımda işlem yapabildiği işlemcilerdir.

ClockRateKHz: 1250000 GPU’nun saat hızıdır.

Clock Speed veya Clock Rate, işlemcinin saniye başına ürettiği sorgu veya titreşim sayısını belirtmek amacıyla kullanılan bir kavramdır. Clock speed oranı daha yüksek olan bir işlemci daha kısa süre içerisinde daha çok işlem gerçekleştirebilme kapasitesine sahiptir. Yani, clock speed değeri yüksek olan bir işlemci clock speed değeri düşük olan işlemciye oranla hem daha çok performans sunar hem de daha hızlı işlem gerçekleştirebilir.

ComputeMode: ‘Default’ grafik kartının hangi hesaplama modunda olduğunu belirtir.

‘Default‘ – grafik kart kısıtlı değildir ve aynı anda birden çok uygulama tarafından kullanılabilir. MATLAB cihazı diğer MATLAB oturumları(sessions) veya işçiler(workers) gibi diğer uygulamalarla paylaşabilir.

‘Exclusive thread‘ veya ‘Exclusive process‘ – grafik kart, aynı anda sadece bir uygulama tarafından kullanılabilir. Cihaz MATLAB’da seçilse de, diğer MATLAB oturumları(sessions) veya çalışanlar(workers) dahil diğer uygulamalar tarafından kullanılamaz.

‘Prohibited‘ Cihaz kullanılamaz.

GPUOverlapsTransfers: 1 cihazın örtüşen transferleri destekleyip desteklemediğini gösterir.

KernelExecutionTimeout: 1 cihazın uzun süren threadleri durdurup durduramayacağını gösterir. Eğer 1 ise işletim sistemi CUDA çekirdeğinin çalıştırması için izin verilen süreye bir üst sınır koyar ve bundan sonra CUDA sürücüsü çekirdeği zaman aşımına uğratır ve bir hata döndürür.

CanMapHostMemory: 1 aygıtın bilgisayar belleğini CUDA adres alanına haritalama yapmasını destekleyip desteklemediğini gösterir.

DeviceSupported: 1 cihazın desteklendiğini belirtir, bazı grafik kartları CUDA’yı desteklememektedir.

DeviceSelected: 1 cihazın seçili cihaz olduğunu belirtir.

Bu özellikleri kontrol etmek için aşağıdaki şekilde bir atama yapmak yeterlidir.

aygit=gpuDevice;
kullanilabilirbellek=aygit.AvailableMemory;

reset(gpudev) komutu ile GPU resetlenir ve bellek boşaltılmış olur.

g = gpuDevice(1);
reset(g);

wait(gpudev) komutu GPU’daki işlemler bitene kadar bekleme yapılmasını sağlar. Böylelikle GPU’da işlenip kullanmamız gereken veriler var ise bu işlemin bitip verilen tam olarak hazır olduğundan emin olmamız için wait komutunu kullanmalıyız. Aksi halde CPU’dan GPU’ya paslanan işin bitip bitmediği ile ilgili bir geri dönüt yoktur. Bunun kullanılmadığı nokta ise gather komutunun kullanılmasıdır. gather komutu ile GPU’daki veriyi RAM’e çektiğimiz için zaten işlemin bittiğini biliriz ve ekstradan wait komutunu kullanmamıza gerek kalmaz.

gpuArray komutu bana göre MATLAB ile CUDA programlayı en cazipleştiren yaklaşımdır. Zira tek bir komut ile GPU belleğine istediğiniz veriyi gönderebiliyorsunuz ki bunu C ile kodlamaya kalktığınız ise önce cudaMalloc ile yer açmalı, cudaMemcpy ile veriyi RAM’den GPU belleğine kopyalamalı, sonra ilgili veri işlendikten sonra GPU bellekten RAM’e taşımalı, GPU’da işiniz bittiğinde ise cudaFree ile ilgili bellek alanını boşaltmanız gerekmektedir.

Giriş verisi gpuArray olarak tanımlandığı için çıktısı da gpuArray olarak verilir. Kullanacağımız formata(single, double, int8 vb.) kendimiz yeniden çevirmemiz gerekmektedir. gather komutu ile GPU belleğindeki veriyi RAM’e taşıyabiliriz.

classUnderlying komutu ile verilerin sınıfını öğrenebiliriz.

10×10’luk tek hassasiyetli rastgele oluşturulmuş sayıların GPU belleğine gönderilip, kareleri alınarak geri RAM’e alınmasını sağlayan kod ve ekran çıktıları aşağıdadır:

X = rand(10,'single');
G = gpuArray(X);
classUnderlying(G)
G2 = G .* G;         
whos G2   
G1 = gather(G2);
whos G1

Oluşturma ve gönderme işlemini tek satırda yapmak için:
G = gpuArray(ones(100,’uint32′)); kullanılabilir.

Daha kullanışlı ve daha efektif olarak bazı komutların direk GPU içerisinde çalışması sağlanmıştır. Bu komutlara ihtiyaca göre eklemeler yapılmaktadır. Bu yaklaşım ile CPU-GPU arasında veri taşıma sırasında kaybedilen zaman kazanılmış olur. methods(‘gpuArray’) komutu ile güncel olarak kullanılabilecek komut listesi elde edilebilir.

Örneğin:
II = eye(1024,’int32′,’gpuArray’);
size(II) komutu ile 1024×1024’lik int32 tipinde veri içeren birim matris oluşturmaya yaramaktadır.

Kendi çalışmalarımda da kullandığım ve genel manada kullanılabilecek önemli bir durum ise rastgele sayı üretme fonksiyonlarıdır. GPU’da direk rastgele sayı üretmeye yarayan komutlar GPU’da işlem yaparken çok faydalı olacaktır.

Matlab ortamında GPU’daki varsayılan rastgele sayı üretme yaklaşımı ile CPU’daki farklıdır. Eğer CPU ve GPU’da aynı rastgele sayı üretme yaklaşımlarının kullanılmasını istiyorsak CPU’daki üretim şeklini değiştirmemiz gerekmektedir.

Aşağıdaki kod ile CPU ve GPU’daki rastgele sayı üretme teknikleri eşitlenir ve üretilen sayıların eşit olup olmadığı kontrol edilerek, eşit olduğu görülebilir.

seed=0; n=4;

cpu_stream = RandStream('CombRecursive','Seed',seed,'NormalTransform','Inversion');
RandStream.setGlobalStream(cpu_stream);

gpu_stream = parallel.gpu.RandStream('CombRecursive','Seed',seed);
parallel.gpu.RandStream.setGlobalStream(gpu_stream);

r = rand(n);             %  CPU
R = rand(n,'gpuArray');  %  GPU
OK = isequal(r,R)

Şu an için Matlab ile GPU hesaplama yaparken rastgele sayı üretmek için aşağıdaki 3 yaklaşım kullanılabilir.

parallel.gpu.RandStream(‘combRecursive’)
parallel.gpu.RandStream(‘Philox4x32-10’)
parallel.gpu.RandStream(‘Threefry4x64-20’)

combined multiplicative recursive generator (MRG32K3A) varsayılan ve yaygın olan kullanılan bir rastgele sayı üretme tekniğidir.

Normal dağıtılan rastgele sayılar üretmek içinde:

sc = RandStream('CombRecursive','NormalTransform','Inversion','Seed',1);
RandStream.setGlobalStream(sc)

sg = parallel.gpu.RandStream('CombRecursive','NormalTransform','Inversion','Seed',1);
parallel.gpu.RandStream.setGlobalStream(sg);

Rc = randn(1,4)
Rg = randn(1,4,'gpuArray')

komutları kullanılır. Yukarıdaki kodun çıktısı:

Rc = -0.0108 -0.7577 -0.8159 0.4742
Rg = -0.0108 -0.7577 -0.8159 0.4742

şeklindedir.

parallel.gpu.rng komutu ile aktif olan rastgele sayı üretme metodu görülebilir.

gpuArray olarak tanımlanmış verilerin karakteristiğini öğrenmek için aşağıdaki komutlar kullanılabilir.

G = rand(10,20,10,'gpuArray');
s=size(G); % Boyutların sınırlarını verir
c=classUnderlying(G); % Verinin tipini verir
e=existsOnGPU(G); % GPU'da olup olmadığını verir
i=isreal(G); % Verinin real türünde olup olmadığını verir
l=length(G); % En büyük dizinin boyutunu verir
n=ndims(G); % Kaç tane dizi(kaç boyutlu) olduğunu verir

Çıktısı:

CUDAKernel objesi oluşturarak Matlab ortamında GPU ile hesaplama yaptırılabilir. Bunun için C kodları ile yazılmış bir .cu dosyasını .ptx dosyasına compile etmek ve bunu sisteme tanıtmak yeterlidir. Aşağıdaki bağlantıda

PTX dosyası nasıl oluşturulur? anlatılmaktadır.

.cu ve .ptx dosyamız hazır ise:

k = parallel.gpu.CUDAKernel(‘myfun.ptx’,’myfun.cu’); komutu ile objemizi oluşturuyoruz.

Dikkat: CUDAKernel objeleri kaydedilemez veya yüklenemez.

Kernel kodu yazarken bazı kısıtlamalar söz konusudur. Bunlar;
-Kernel __global__ void mykernel(inputs …) ile başlar.
-Tüm girdiler sayı veya işaretçi ve const etiketli olabilir.
-Kernel herhangi bir değer döndürmez, sadece girişlerindeki verileri işler.
-Kernel herhangi bir bellek biçimi tahsis edemez, bu nedenle kernel çalıştırılmadan önce tüm çıktılar önceden ayrılmış olmalıdır. Bu nedenle, kerneli çalıştırmadan önce tüm çıktıların boyutları bilinmelidir.
-Prensip olarak, kernele gönderilen tüm pointerlar çıktı verileri içerebileceğinden const olarak tanımlanmazlar. {Çıktı verisi almak için genel yaklaşım budur.} Kernelin üreteceği çıktı da bir girdi olarak verilmektedir.

Kaç kernel’de hesaplama yapılacağını belirtmek için k.ThreadBlockSize = [10,1,1]; kodu kullanıldığı zaman ilgili kod 1 block içindeki aynı zamanda koşturulacak olan thread sayısının 10 olduğunu belirtir. Grid-Block-Thread sayıları CUDA mimarisi içinde belirlenmesi önemli ve zor olarak parametrelerdendir. GPU’nun fiziksel sınırlarını aşmadan, ihtiyaca en uygun yapı kurulmalıdır. Genellikle verilerin boyutlarıyla bağlantılı sayılar seçilir.

Sphere fonksiyonu için GPU’da hesaplama yaptıracak olursak:
Sphere.cu dosyamız:

__global__ void Obj(double * objt, double * nt,int N,int D)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    double total = 0.00;
    for (int j = 0; j < D; j++){
    total=(nt[i+(j*N)]*nt[i+(j*N)]);
    objt[i]=objt[i]+total;
}    
}	

şeklinde oluşturulabilir.

10×4’lük 1’den başlayıp 40’a kadar sayıların doldurduğu giriş verimizi göndererek 4 boyut için f(x)=x^2 değerini hesaplayıp toplam olarak döndüren kod aşağıdadır.

N=10;
D=4;
for i=1:N
    for j=1:D
        nt(i,j)=j+((i-1)*D);
    end
end
objt=zeros(1,N,'gpuArray');
nt=gpuArray(nt);
k = parallel.gpu.CUDAKernel('Sphere.ptx','Sphere.cu','Obj');
k.ThreadBlockSize = N;
result = feval(k,objt,nt,N,D);
objt=gather(result);

Giriş verisi:

Yapılan işlem her bir satırdaki sayıları sırasıyla kendisiyle çarpıp toplayıp fonksiyon sonucu olarak döndürmektir. Örneğin ilk satır için:

1*1+2*2+3*3+4*4=30’dur.

k.ThreadBlockSize = N; komutu ile N adet thread çalıştırmış oluyoruz. Daha önceden indeks yapısını kurduğumuz .cu dosyası sayesinde gerekli hesaplamalar yapılıp, sonuç bize geri döndürülmüştür. Örneğin N=10 olduğundan ilgili programın çıktısı:
30 174 446 846 1374 2030 2814 3726 4766 5934
şeklinde olmuştur. Aynı kodda k.ThreadBlockSize = N; kısmını k.ThreadBlockSize = 5; yaptığımızda;
30 174 446 846 1374 0 0 0 0 0 şeklinde;
k.ThreadBlockSize = 20; yaptığımızda:
30 174 446 846 1374 2030 2814 3726 4766 5934
şeklinde sonuç vermektedir. Buradan çıkaracağımız sonuca göre thread-block-grid sayısı hesaplamak istediğimiz verilerle örtüşmelidir. Fazla olduğu takdirde bir sorun olmasa da eksik olduğu takdirde yanlış hesaplamalar yapacaktır.

Giriş ve çıkış verilerinin boyutları ile thread-block-grid sayılarına dikkat etmek gerekmektedir.

Yukarıdaki CUDAKernel objesinin özelliklerini inceleyecek olursak:

MaxNumLHSArguments: 2 adet çıktı değişkeni olduğunu, NumRHSArguments: 4 adet girdiği değişkeni olduğunu belirtir.

ArgumentTypes: {‘inout double vector’ ‘inout double vector’ ‘in int32 scalar’ ‘in int32 scalar’} satırı ise giriş ve çıkış verilerinin hangilerinin giriş, hangilerinin çıkış olduğunu ve tiplerini bildirir.

SharedMemorySize: 0 Her bir iş parçacığı bloğunun kullanabileceği dinamik paylaşılan bellek miktarı (bayt cinsinden). Her bir iş parçacığı bloğunda kullanılabilir bir paylaşılan bellek bölgesi bulunur. Bu bölgenin boyutu, mevcut kartlarda ~ 16 kB ile sınırlıdır ve çok işlemcili kayıtlarla paylaşılır. Tüm belleğin olduğu gibi, kernel başlatılmadan önce bunun tahsis edilmesi gerekir. Bu paylaşılan bellek bölgesinin boyutunun, iş parçacığı bloğunun boyutuna bağlı olması da yaygındır. Bu değeri kernel üzerinde ayarlamak, bir bloğun her bir iş parçacığının bu mevcut paylaşılan bellek bölgesine erişmesini sağlar.

setConstantMemory komutu ile her thread’in ulaşabileceği bir alan olan constant memory alanına değişkenler tanımlanabilir. Örneğin;

__constant__ int32_t N1;
__constant__ int N2; // Assume ‘int’ is 32 bits
__constant__ double CONST_DATA[256];

şeklinde sabit değişkenler;

KERN = parallel.gpu.CUDAKernel(ptxFile,cudaFile);

setConstantMemory(KERN,'N1',int32(10));
setConstantMemory(KERN,'N2',int32(10));
setConstantMemory(KERN,'CONST_DATA',1:10);

şeklinde tanımlanabilir.

GPU Hesaplama’nın sorunları

GPU hesaplama hızlanma sağlarken bazı maliyetleri beraberinde getirmektedir.

1.Bellek Erişimi: GPU kart, PCI Express veriyolu ile anakarta bağlı olduğundan CPU ile iletişimi bu yol üzerinden olmaktadır. O yüzden CPU, RAM’e GPU’ya göre daha çabuk erişmektedir. Veriler işlenmek için GPU’ya gönderilip yeniden alınması gerekmektedir. Bu iletişim yükü doğurmaktadır.

2.CUDA kernel kodlaması C veya Fortran dilinde yapıldığından bu beceriye sahip olunması gerekmektedir.

3.Kodun kullanılacak GPU’ya göre optimize edilmesi gerekmektedir.

devam edebilir…