CUDA

CUDA Nedir
CUDA grafik motoruna gidecek veriyi hazırlamak için kullanılan bir hesaplama motoru gibi düşünülebilir. GPU'lar SIMD (Single Instruction Multiple Data) modeline göre tasarlanmışlardır. Veriyi küçük parçalara bölerek paralel çalıştırmaya uygundurlar. OpenCL'deki computing unit ile Cuda  core farklı şeyler.

Program Yapısı
Kod device (GPU'da çalışır), host (CPU'da çalışır) ve global(CPU'da çalışır) ayrılıyor. Aşağıda host'un seri, device kodunun ise paralel çaşıltığı görülebilir.

Her device bloklara ayrılır. Her blok içinde ise thread'ler çalışır. Yani bloklar thread'leri gruplarlar.
Aşağıdaki şeklde bir bloğun içi görülebilir.

CUDA ile Dynamic Parallelism
GPU içindeki bir thread'in kendi kendine başka bir thread yaratabilmesi anlamına geliyor. Örnek:

Açıklaması ise şöyle:
Dynamic Parallelism on Kepler GPU dynamically spawns new threads by adapting to the data without going back to the CPU, greatly simplifying GPU programming and accelerating a broader set of popular algorithms.

Sürücü
Linux'ta Nvidia CUDA sürücüsünün sürümü şöyle öğrenilir.

$cat /proc/driver/nvidia/version

NVRM version: NVIDIA UNIX x86_64 Kernel Module 304.125 Mon Dec 1 19:58:28 PST 2014
GCC version:  gcc version 4.8.2 (Ubuntu 4.8.2-19ubuntu1) 

Benim sistemimde Kernel Module 340.96  ve sene olarak 2015 yazıyor.

nvdia-smi komutu
Şuna benzer bir çıktı verir.

+------------------------------------------------------+                       
| NVIDIA-SMI 352.79     Driver Version: 352.79         |                       
|-------------------------------+----------------------+----------------------+
| GPU  Name        Persistence-M| Bus-Id        Disp.A | Volatile Uncorr. ECC |
| Fan  Temp  Perf  Pwr:Usage/Cap|         Memory-Usage | GPU-Util  Compute M. |
|===============================+======================+======================|
|   0  GeForce GTX 780 Ti  Off  | 0000:01:00.0     N/A |                  N/A |
| 30%   43C    P2    N/A /  N/A |    877MiB /  3071MiB |     N/A      Default |
+-------------------------------+----------------------+----------------------+
|   1  GeForce GTX 780 Ti  Off  | 0000:03:00.0     N/A |                  N/A |
| 29%   25C    P8    N/A /  N/A |     11MiB /  3071MiB |     N/A      Default |
+-------------------------------+----------------------+----------------------+

deviceQuery komutu
Şuna benzer bir çıktı verir.


Derleyici
CUDA derleyicisi nvcc. CUDA kodları .cu uzantısı ile biter. Örneğin hello.cu.

Basit bir kodu

#include <cstdlib>
#include <cstdio>
#include <cuda.h>

using namespace std;

__global__ void mykernel(void) {
}

int main(void) {
mykernel<<<1,1>>>();
printf("CPU Hello World!\n");
return 0;
} 

Şöyle derleriz.

$ nvcc hello.cu
$ a.out
Hello World!

--compiler-options seçeneği
Şöyle yaparız.

nvcc --ptxas-options=-v --compiler-options '-fPIC' --shared
lots of object files
lots of -L/wherever -lwhatever options
-o libOutput.so

-o seçeneği
Oluşmasını istediğimiz dosya ismidir.

$ nvcc -o t1056 t1056.cu

-version
Şu çıktıyı alırız.

nvcc: NVIDIA (R) Cuda compiler driver
 Copyright (c) 2005-2015 NVIDIA Corporation
 Built on Tue_Aug_11_14:27:32_CDT_2015
 Cuda compilation tools, release 7.5, V7.5.17

cuda-memcheck
Varsa bellek hatasını gösterir.

$ nvcc -o t1074 t1074.cu
$ cuda-memcheck ./t1074
========= CUDA-MEMCHECK
...
========= ERROR SUMMARY: 0 errors

Kodlama Mantığı
Mantık şöyle.
1. Önce veri host bellekten device belleğe kopyalanır.
2. Hesaplama yapılır.
3. Veri bu sefer device bellekten, host belleğe kopyalanır.

metodlar
Metodları kullanbilmek için şu dosya include edilir.

#include<cuda.h>

__global__
GPU'da koşacak kodu belirtir. Örnekte her işlemci array'e 7 yazar. Bu örnekte blockIdx.x yukarıda bahsedilen blok numarasıdır. threadIdx.x ise blok içindeki thread numarasıdır.

__global__ void kernel(int *array)
{
  int index = blockIdx.x * blockDim.x + threadIdx.x;

  array[index] = 7;
}

Kod şöyle çağırılır.

int main(void)
{
  ...
  kernel<<<grid_size,block_size>>>(device_array);
  ...
  
} 

blockIdx, blockDim, threadIdx
cuda metodunu çağırırken <<<...>>> şeklinde kodlanır. Şöyle yaparız.

kernel<<<1, 267>>>(...);

İlk parametre kaç tane blok olacağını belirtir. İkinci parametre blok içinde kaç tane thread olacağını belirtir.

Sabit sayılar yerine dim3 değişkenleri de kullanılabilir.

dim3 dimBlocks(numBlocks);
dim3 dimThreads(numThreadsPerBlock);
kernel<<<dimBlocks, dimThreads>>>(...);

x,y,z şöyle hesaplanır. x ve y en fazla 1024 olabilir.  z ise en fazla 64 olabilir.

__global__ void kernel(...)
{
  unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
  unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
  unsigned int z = blockIdx.z * blockDim.z + threadIdx.z;
  ...
}

__shared__
Tam nasıl kullanıldığını anlamadım ama şöyle tanımlanıyor.

extern __shared__ float myArray[];

cudaDeviceReset metodu
Uygulama başlarken şöyle yaparız.

int main(void) {

  cudaDeviceReset();
  cudaDeviceSynchronize();
  ...
}

cudaSetDeviceFlags metodu
Şöyle yaparız.

cudaSetDeviceFlags(cudaDeviceMapHost);

Cuda Bellek Yönetimi
İki farklı yöntem var.
1. Zero Copy
2. Device Copy

Zero Copy yöntemini anlamadım ancak sanırım daha hızlı çünkü device ve host bellekleri arasında kopyalama yapılmıyor.

Device Copy yönteminde device ve host bellekleri arasında kopyalama yapılıyor.

cudaHostAlloc metodu
Şöyle yaparız.

float *psrcMat;
cudaHostAlloc((void **)&psrcMat, 10*sizeof(float), cudaHostAllocMapped);

Bu bellek alanına şöyle yapılabilir.

float *d_psrcMat;
//Map device to host pointers
cudaHostGetDevicePointer((void **)&d_psrcMat, (void *)psrcMat, 0);

cudaFreeHost
cudaHostAllıc ile host üzerinde yaratılan bellek alanının free edilmesi gerekir. Şöyle yaparız.

cudaFreeHost(psrcMat);

cudaMalloc metodu
C'deki malloc'tan farklı olarak hata kodu döner. Şöyle kullanılabilir.

void *fixed_cudaMalloc(size_t len)
{
  void *p;
  if (cudaMalloc(&p, len) == success_code) 
    return p;
  return 0;
}

cduGetLastError ile beraber şöyle kullanılabilir.

// Number of bytes in the matrix.
int bytes = 9 *sizeof(float);
// Pointers to the device arrays
float *matrixd=NULL; 

// Allocate memory on the device to store matrix
cudaMalloc ((void**) &matrixd, bytes);
cudaError_t status = cudaGetLastError ();  //Check the error
if (status != cudaSuccess) {                     
  cudaFree (matrixd);                     //Free call for memory
  ...
}

cudaFree metodu
cudaMalloc ile device üzerinde yaratılan bellek alanının free edilmesi gerekir.

int main(void)
{
  int num_bytes = 256;

  // pointers to host & device arrays
  int *device_array = 0;

  // cudaMalloc a device array
  cudaMalloc((void**)&device_array, num_bytes);

  ....

  // deallocate memory
  cudaFree(device_array);
} 

cudaGetErrorString metodu
Eğer istenirse hata kotu string'e çevrilebilir. Şöyle yaparız.

cudaError_t code = ...;
if (code != cudaSuccess) 
{
  fprintf(stderr,"Error: %s\n", cudaGetErrorString(code));
  ...
}

cudaMemcpy metodu

Host bellekten device belleğe ya de device bellekten host bellğe kopyalamak için kullanılır. Şöyle yaparız.

bool* deviceFlag;
cudaMalloc(&deviceFlag, sizeof(bool));
cudaMemcpy(deviceFlag, &hostFlag, sizeof(bool), cudaMemcpyHostToDevice);

Elimizde bir struct olsun ve bu struct'ı ile dolu bir de vector olsun.

struct Atomic_Orbital{
      L  ag;
      Cont cont;
};
vector<Atomic_Orbital> v;

v.resize(100);
for(int i=0; i<100; i++)
{
  ...
}

Bu vector'ü host bellekten device belleğe şöyle kopyalarız

Atomic_Orbital *devPtr;
size_t size = B.size() * sizeof(Atomic_Orbital);
cudaMalloc(&devPtr, size);
cudaMemcpy(devPtr, B.data(), size, cudaMemcpyHostToDevice);

cudaMallocPitch ve cudaMemcpy2D metodları
Sanırım bir şekilde beraber kullanılıyorlar.

cudaMallocPitch(...);
cudaMemcpy2D(...);

cudaDeviceSyncronize metodu
kernel<<<...,...>>> çağrısından sonra yapılır.

kernel<<<...,...>>>(...);
cudaDeviceSynchronize();

cudaThreadSyncronize metodu
kernel<<<...>>> çağrısından sonra yapılır.

kernel<<<...,...>>>(...);
cudaThreadSynchronize();

cudaEvent_t
Event şöyle tanımlanır

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

Daha sonra şöyle yaparız.

cudaEventRecord(start,0);
...//Call something
cudaEventRecord(stop,0);
cudaEventSynchronize(stop);
float milliseconds = 0;
cudaEventElapsedTime(&milliseconds, start, stop);

cudaStream_t
Şöyle yaparız.

cudaStream_t stream1;   
cudaStreamCreate(&stream1);

thrust

thrust yazısına taşıdım.