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
igrid
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()
idrv.InOut()
pretvarajunumpy
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 komponentex
,y
; opisuje veličinu rešetke.Postoji definirana varijabla
blockIdx
i ima komponentex
,y
; daje indeks bloka u rešetci.Postoji definirana varijabla
blockDim
i ima komponentex
,y
,z
; opisuje veličinu bloka.postoji definirana varijabla
threadIdx
i ima komponentex
,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.