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.