Normál kernel írása

Áttérés tetszőleges kernelekre

Elrejtés

A Numba vektorizált műveletei nagyon szépen működnek amíg az független adatokon végzett számításokra van szükségünk. Ha viszont átlapoló adatokkal kell dolgoznunk, akkor vektorizálva nem tudjuk azt megoldani. (A konvolúció például már nem oldható meg vektorizált függvénnyel.)

Írhatunk viszont tetszőleges kernelt is. Ilyenkor a kernel írásakor pár dologra kell figyelnünk:

  • A függvény előtt "@cuda.jit" prefix lesz a vektorizálás helyett;
  • Blokkstruktúrát használunk aminek paramétereit a "cuda.threadIdx.x/y/z", "cuda.blockIdx.x/y/z", "cuda.blockDim.x/y/z", "cuda.gridDim.x/y/z" beépített változókon keresztül érjük el.
  • A kernel paraméterei nem lesznek vektorizálva, azokat minden szál ugyanúgy látja.
@cuda.jit
def add_kernel(x, y, out, n):
    tx = cuda.threadIdx.x
    bx = cuda.blockIdx.x

    tid = tx + bx * cuda.blockDim.x

    if tid < n:
        out[tid] = x[tid] + y[tid]

A másik oldalon a függvény hívásokkor is figyelembe kell vennünk néhány újdonságot.

  • Meg kell határoznunk az indított szál-struktúrát (blokk és rács méret),
  • Ezt a struktúrát át is kell adnunk a kernelnek a paraméterlista előtt szögletes zárójelek között: "kernel[blocks_per_grid, threads_per_block](...)"
    threads_per_block = 128
    blocks_per_grid = math.floor((N+threads_per_block-1)/threads_per_block)
    
    start = timer()
    
    add_kernel[blocks_per_grid, threads_per_block](A, B, C, N)

Teljes kód egyben

Elrejtés

A teljes kód egyben itt látható:

!!! Lentebb van még egy kiegészítés !!!

import numpy as np
from timeit import default_timer as timer
import math

from numba import vectorize
from numba import cuda

@cuda.jit
def add_kernel(x, y, out, n):
    tx = cuda.threadIdx.x
    bx = cuda.blockIdx.x

    tid = tx + bx * cuda.blockDim.x

    if tid < n:
        out[tid] = x[tid] + y[tid]


def main():
    N = 32000000
    
    A = np.ones(N, dtype=np.float32)
    B = np.ones(N, dtype=np.float32)
    C = np.zeros(N, dtype=np.float32)
    
    threads_per_block = 128
    blocks_per_grid = math.floor((N+threads_per_block-1)/threads_per_block)
    
    start = timer()
    
    add_kernel[blocks_per_grid, threads_per_block](A, B, C, N)
    
    vectoradd_time = timer() - start
    
    print("C[:5] = " + str(C[:5]))
    print("C[-5:] = " + str(C[-5:]))
    
    print("VectorAdd took {} seconds".format(vectoradd_time))

if __name__== "__main__":
    main()

Menedzselt memória

Elrejtés

Végül pedig, ahogy a vektorizált műveletek esetén is, itt is van lehetőségünk a memória kézzel történő kezelésére.

Ennek módja megegyezik a vektorizálásnál látottal.

import numpy as np
from timeit import default_timer as timer
import math

from numba import vectorize
from numba import cuda

@cuda.jit
def add_kernel(x, y, out, n):
    tx = cuda.threadIdx.x
    bx = cuda.blockIdx.x

    tid = tx + bx * cuda.blockDim.x

    if tid < n:
        out[tid] = x[tid] + y[tid]


def main():
    N = 32000000
    
    d_A = cuda.to_device(np.ones(N, dtype=np.float32))
    d_B = cuda.to_device(np.ones(N, dtype=np.float32))
    d_C = cuda.device_array(shape=(N,), dtype=np.float32)
    
    threads_per_block = 128
    blocks_per_grid = math.floor((N+threads_per_block-1)/threads_per_block)
    
    start = timer()
    
    add_kernel[blocks_per_grid, threads_per_block](d_A, d_B, d_C, N)
    
    vectoradd_time = timer() - start
    
    C = d_C.copy_to_host();
    
    print("C[:5] = " + str(C[:5]))
    print("C[-5:] = " + str(C[-5:]))
    
    print("VectorAdd took {} seconds".format(vectoradd_time))

if __name__== "__main__":
    main()

Feladatok

Elrejtés

  • Írjunk át egy tetszőleges korábbi példaprogramot python nyelvre!