Programowanie GPU: Kurs CUDA – Wprowadzenie

Procesory graficzne coraz odważniej wkraczają w świat programowania i przetwarzania strumieniowego. Przetwarzanie strumieniowe to technika pozwalająca na wykonywanie tego samego kodu równocześnie na różnych danych. Ma to zastosowanie przy obliczeniach numerycznych, gdzie np. 500 jednostek strumieniowych jest zaprzęgniętych do przemnożenia jednej macierzy przez drugą. Każda jednostka liczy odrębny fragment macierzy wynikowej i dzięki temu obliczenia stają się wielokrotnie szybsze. Taki typ przetwarzania ma zastosowanie w sztucznej inteligencji, przetwarzaniu sygnałów i obrazów, czy symulacjach naukowych. Jak narazie nie jest to alternatywa do zwykłych CPU, a jedynie pomoc przy niektórych obliczeniach. Programy sekwencyjne nadal są wykonywane na CPU, gdyż algorytmy predykcyjne i wyższe taktowanie dają CPU przewagę w obszarze obliczeń sekwencyjnych.

Koncepcji na obliczenia wykonywane przez GPU jest tyle ilu producentów tego sprzętu (m.in. NVIDIA, AMD, ATU). NVIDIA wprowadza architekturę zwaną CUDA (Compute Unified Device Architecture). Udostępnia przy tym SDK zawierający m. in. kompilator oparty na składni C.

Bardzo ważnym aspektem architekruty CUDA jest wprowadzenie koncepcji wątku (thread) oraz bloku wątków (thread block). Dane są przetwarzne przez zupełnie niezależne bloki. W zależności od architektury sprzętu bloków naraz wykonuących obliczenia może być więcej lub mnie. Każdy blok składa się z pewnej liczby wątków. Wątki w obrębie jednego bloku można synchronizować i mają dostęp do wspólnych danych dzielonych.

Poniżej przykładowa aplikacja w CUDA. Kod wyjaśniony w komentarzach.

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
#include "stdafx.h"
#include <stdio.h>
#include <cuda.h>
 
// Funkcja wykonywana równolegle na wielu danych
// przez jednostki strumieniowe GPU
__global__ void square_array(float *a, int N)
{
  // Obliczamy numer wątku na podstawie numeru bloku,
  // numeru wątku w bloku i ilości wątków w bloku
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  // Warunek, aby nie wychodzić poza zakres tablicy,
  // którą będziemy przetwarzać oraz właściwe obliczenie
  // czyli obliczenie kwadratu danego pola, za które
  // odpowiedzialny jest dany wątek
  if (idx<N) a[idx] = a[idx] * a[idx];
}
 
// Funkcja wywoływana na CPU
int main(void)
{
  float *a_h, *a_d;  // Wskaźniki do wykorzystywanych tablic
  const int N = 10;  // Długość tablic
  size_t size = N * sizeof(float); // Rozmiar tablic w bajtach
 
  a_h = (float *)malloc(size);        // Alokacja pamięci przez CPU
  cudaMalloc((void **) &a_d, size);   // Alokacja pamięci przez GPU
 
  // Inicjalizacja danych przez CPU
  for (int i=0; i<N; i++) a_h[i] = (float)i;
  // Skopiowanie danych do GPU
  cudaMemcpy(a_d, a_h, size, cudaMemcpyHostToDevice);
 
  // Ustalenie rozmiaru bloku
  int block_size = 4;
  // Obliczenie liczby bloków koniczecznych do przetworzenia całej tablicy
  int n_blocks = N/block_size + (N%block_size == 0 ? 0:1);
 
  // Specjalna składnia wywołania podprogramu dla GPU z podaniem
  // nazwy funkcji, liczby i rozmiaru bloku oraz parametrów funkcji
  square_array <<< n_blocks, block_size >>> (a_d, N);
  // Pobranie danych wynikowych z GPU do CPU
  cudaMemcpy(a_h, a_d, sizeof(float)*N, cudaMemcpyDeviceToHost);
 
  // Wypisanie wyników
  for (int i=0; i<N; i++) printf("%d %f\n", i, a_h[i]);
  // Zwalnianie
  free(a_h); cudaFree(a_d);
}