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)
  

• 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)
  

Nit, blok i rešetka

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))

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()

Author: Vedran Miletić, Kristijan Lenković