Section author: Vedran Miletić, Kristijan Lenković

Python modul PyCUDA: zbrajanje vektora¶

CUDA C/C++ kod zrna spremamo u zbroj.cu

uočimo da se indeks dohvaća iz varijable threadIdx.x

__global__ void zbroj_vektora (float *dest, float *a, float *b)
{
  const int i = threadIdx.x;
  dest[i] = a[i] + b[i];
}

dohvaćanje zrna u Python

mod = SourceModule(open("zbroj.cu").read())

zbroj_vektora = mod.get_function("zbroj_vektora")

inicijalizacija dva vektora čije su sve vrijednosti jednake 1 i vektora čje su sve vrijednosti jednake 0 i jednake je duljine kao ova dva
```
a = np.ones(400, dtype=np.float32)
b = np.ones(400, dtype=np.float32)

result_gpu = np.zeros_like(a)
```

Note

Pripazite na usklađenost Pythonovih numpy tipova podataka sa C-ovim tipovima podataka koji se koriste u kernelu (ovdje je to numpy.float32 s float). Potpun popis može se naći među često postavljanim pitanjima u vezi PyCUDA-e.

izvođenje koda na GPU-u (službena dokumentacija)
- prvo se navode argumenti kernela istim redom kao u definiciji kernela u C-u
- block i grid su stvari kojima se bavimo kasnije, ali uočite da stavljamo istu vrijednost kao što je veličina vektora
```
zbroj_vektora(drv.Out(result_gpu), drv.In(a), drv.In(b), block=(400,1,1), grid=(1,1))
```
- drv.In(), drv.Out() i drv.InOut() pretvaraju numpy polja u polja s kojima GPU može manipulirati (službena dokumentacija objašnjava razliku između te tri funkcije i kada je potrebno koju od njih koristiti)
izvođenje zbrajanja na CPU-u (jednostavnost sintakse osigurava preopterećeni operator zbrajanja od strane numpy modula)
```
result_cpu = a + b
```

ispis rezultata i provjera jednakosti po pojedinim poljima

print("CPU rezultat\n", result_cpu)
print("GPU rezultat\n", result_gpu)
print("CPU i GPU daju isti rezultat?\n", result_cpu == result_gpu)

Zadatak

Prilagodite kod primjera tako da računa zbroj oblika \(2a + b\).
Prilagodite kod primjera tako da računa zbroj oblika \(2a + b + 3c\).

Nit, blok i rešetka¶

Zadatak

Prilagodite kod prethodnog primjera tako da vektori imaju 500 elemenata (umjesto 400).

Zrna se pokreću na rešetci. Rešetka se sastoji od blokova, koji se sastoje od niti.

Postoji definirana varijabla gridDim i ima komponente x, y; opisuje veličinu rešetke.
Postoji definirana varijabla blockIdx i ima komponente x, y; daje indeks bloka u rešetci.
Postoji definirana varijabla blockDim i ima komponente x, y, z; opisuje veličinu bloka.
postoji definirana varijabla threadIdx i ima komponente x, y, z; daje indeks niti u bloku.

Indeks i ovisi o indeksu bloka (blockIdx), dimenziji bloka (blockDim) i indeksu niti (threadIdx).

__global__ void vector_sum (float *dest, float *a, float *b)
{
  const int i = blockIdx.x * blockDim.x + threadIdx.x;
  dest[i] = a[i] + b[i];
}

Svaki od dva bloka sada ima 250 niti i one pokrivaju čitav vektor od 500 komponenata.

zbroj_vektora(drv.Out(result_gpu), drv.In(a), drv.In(b), block=(250,1,1), grid=(2,1))

Zadatak

Provjerite možete li se s dva bloka izvesti zbrajanje vektora od 3000 elemenata. (Napomena: čim je zadano ovako, vjerojatno ne možete.)
Izvedite zbrajanje vektora od 3000 elemenata u 3 bloka.

Note

GPU-i zasnovani na Tesla arhitekturi podržavaju maksimalno 512 niti po bloku, dok GPU-i zasnovani na arhitekturama Fermi, Keppler i Maxwell podržavaju maksimalno 1024 niti po bloku.

Sposobnosti GPU uređaja možemo saznati pomoću idućeg koda

def device_meminfo():
    (free, total) = drv.mem_get_info()
    print("=== Global memory occupancy for device 0 ===")
    print("Free: %d" % free)
    print("Total: %d" % total)
    print("Percentage free: %f%%" % (float(free) * 100 / float(total)))

    for devicenum in range(drv.Device.count()):
        device = drv.Device(devicenum)
        attrs = device.get_attributes()

        print("\n=== Attributes for device %d ===" % devicenum)
        for (key, value) in attrs.iteritems():
            print("%s: %s" % (str(key), str(value)))

Atomične operacije¶

Atomične aritmetičke operacije su:

atomicAdd()
atomicSub()
atomicExch()
atomicMin()
atomicMax()
atomicInc()
atomicDec()
atomicCAS()

Atomične bitovne operacije su:

atomicAnd()
atomicOr()
atomicXor()

Todo

Ovdje nedostaje objašnjenje i zadatak.