[AA-part3] CUDA architecture

Lecture notes of Architecture and algorithms

Course: https://ucilnica.fri.uni-lj.si/course/view.php?id=89
Lecturer: doc. dr. Tomaž Dobravec
Language: Slovenian
Date: 2014-05-12

Course overview (part 3):

  • hardware and programming aspects of CUDA architecture
  • solving standard problems and a comparison between serial (CPU) and parallel (GPU) implementation
  • basics of the OpenCL programming environment
  • CUDA architecture in the context of the OpenCL environment

Učili se bomo CUDA. Seminarska naloga bo napisati program v Java in CUDA, lahko pa tudi Python in OpenCL. Rok in zagovor seminarskih naj bodo 9.6. po pedagoški delavnici.

Project idea: given Java skeleton for an image processing application, implement Seam Carving using CUDA. Otherwise propose your project.

Arhitektura CUDA

Razno:

  • host – CPU
  • device – GPU
  • SM – streaming multi-processor, ima svoj cache, izvaj isto kodo
  • celo v skupkih izvajajo isto kodo
  • SP – posamezen procesor
  • kernel/ščepec – osnovna komponenta kode
  • thread/nit – izvajajo kernele, potrebno podati kako se naj izvajajo
  • 1 blok niti se vedno izvrši na istem SM
  • razporedi niti v bloke po največ 512 niti, ki med seboj lahko komunicirajo, z ostalimi pa ne morejo sodelovati (saj se lahko izvajajo hkrati)
  • programer pripravi mrezo blokov, ki se potem izvaja, na koncu veš, da se bodo vsi bloki izvršili
  • pomnilnik se ne briše med različnimi programi

Date: 2014-05-19

Recommended tool is Nsight (CUDA Eclipse-like IDE):

  • always check for returned errors
  • files *.cu

Memory model:

  • registri (32-bitni) so na vsak blok (8192), nato se razporedijo med threadi, največ 16 na en ščepec (kernel) pri polni obremenitvi, če compiler ne uspe, ga da v lokalni pomnilnik
  • lokalen pomnilnik se uporablja znotraj ščepca, počasen del DRAMa, skupaj 8kB
  • shared memory uporabljajo vse niti istega ščepca, če podatek več kot enkrat potrebuješ, se ga splača prekopirati sem, to je ključ hitrega programa, hiter
  • global (device) memory, lahko uporablja za komunikacijo med ščepci, počasen DRAM, skupaj par GB
  • host (CPU) memory

Kernel grid:

  • dokler ne konča izvajanja enega kernela, ne gre na drugega
  • lightweight context switching med bloki niti, zato lahko več blokov hkrati naloži
  • bistvo veliko dobro obteženih blokov
  • velikost bloka >=32, saj sicer izvajanje v snopih (warp) manj učinkovito

Java native interface

Date: 2014-05-26

V Javi ni možno dostopati do posebnih sistemskih klicev kot je CUDA, je pa možno z JNI poklicati C knjižnico, ki nato kliče naprej.

// -Djava.library.path=...
static {
    System.loadLibrary("JNIFirst");
}
private native static int sestej(int a, int b);
cd src
javah -jni JNI
# (generates JNI.h from JNI.java)
# (prepare JNI.c)
gcc -I$JNI_INCLUDE -c JNI.c -o JNI.o
gcc -dynamiclib -o libJNIFirst.jnilib JNI.o

Concept: Shift

How to shift all elements of a vector for 1 element to the left using CUDA.

Note: Avoid simultaneous read and write and synchronize execution using __syncthreads() (must be outside a conditional clause).

__global__ void shiftLeft(int *a) {
    __shared__ int mem[N];  // for all blocks
    int idx = threadIdx.x;
    mem[idx] = a[idx];
    __syncthreads();  // must be outside conditional clauses
    if(idx < N - 1)
        a[idx] = mem[idx + 1];
}

Concept: Reduction

How to sum all elements of a vector using CUDA. Parallel operation is called reduction.

__global__ void reduce(int *a) {
    int i = threadIdx.x;
    for(int stride=1; stride < N; stride *= 2) {
        if(i % stride == 0)
            a[2*i] += a[2 * i + stride];
        __syncthreads();
    }
}

Note: Using shared memory instead of global would speed it up.

Note: In each iteration we create holes in the vector that still use resources in a warp. By compacting and leaving holes out we can improve this.

__global__ void reduce(int *a) {
    int i = threadIdx.x;
    extern __shared__ int mem[];
    mem[i] = a[2 * i] + a[2 * i + 1];
    for(int stride=N/4; stride >= 1; stride /= 2) {
        __syncthreads();
        if(i < stride)
            mem[i] += mem[i + stride];
    }
    __syncthreads();
    a[0] = mem[0];
}

Concept: Histogram

How to count number of occurences of characters in a string (i.e. histogram).

Note: Problem is simultaneous read and write for a character. Solution is a special CUDA atomic operation.

Note: Neighbouring threads should access neighbouring memory locations, because DRAM works faster when accessed in blocks.

__global__ void histogram(unsigned char *str, int len, unsigned int *hist) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = blockDim.x * gridDim.x;
    while(i < len) {
        //hist[niz[i] - 'a']++;  // wrong
        atomicAdd(&hist[niz[i] - 'a'], 1);
        i += stride;
    }
}

Oddaja projekta (izvorna koda, kratko poročilo) in prezentacija 16.6.2014 ob 14:00.