Impresiones

Si uno desea que el kernel imprima algo mientras se ejecuta se tiene que hacer uso de printf que es la función utilizada en C y CUDA para imprimir. Sin embargo hay recordar algunas cosas. En primer lugar, printf imprime cosas en la consola, es por eso que el lector notará al ejecutar la siguiente celda, que incluso cuando en el kernel le estamos ordenando que imprima, en el notebook no obtenemos nada. A primera vista parecería un error garrafal, sin embargo si comentáramos la primer linea (la que comienza con %%writefile), al revisar la consola desde la cuál abrimos el notebook se observaría que los resultados esperados se han impreso ahí. Ejecutemos la siguiente celda, con el comando mágico %%writefile para escribir el contenido de la celda en un archivo.


In [1]:
%%writefile ./Programas/saludar.py
import pycuda.driver as cuda
import pycuda.autoinit
from pycuda.compiler import SourceModule

mod = SourceModule("""
    #include <stdio.h>

    __global__ void saluda()
    {
      printf("Mi indice x es %d, mi indice en y es %d\\n", threadIdx.x, threadIdx.y);
    }
    """)

func = mod.get_function("saluda")
func(block=(4,4,1))


Writing ./Programas/saludar.py

Ahora ejecutemos el programa


In [2]:
!ipython ./Programas/saludar.py


Mi indice x es 0, mi indice en y es 0
Mi indice x es 1, mi indice en y es 0
Mi indice x es 2, mi indice en y es 0
Mi indice x es 3, mi indice en y es 0
Mi indice x es 0, mi indice en y es 1
Mi indice x es 1, mi indice en y es 1
Mi indice x es 2, mi indice en y es 1
Mi indice x es 3, mi indice en y es 1
Mi indice x es 0, mi indice en y es 2
Mi indice x es 1, mi indice en y es 2
Mi indice x es 2, mi indice en y es 2
Mi indice x es 3, mi indice en y es 2
Mi indice x es 0, mi indice en y es 3
Mi indice x es 1, mi indice en y es 3
Mi indice x es 2, mi indice en y es 3
Mi indice x es 3, mi indice en y es 3

Por último hagamos un kernel que imprima más datos acerca del thread, como su índice de bloque, la dimensión del bloque, y sus índices dentro del bloque.


In [3]:
%%writefile ./Programas/saludar_bloques.py
import pycuda.driver as cuda
import pycuda.autoinit
from pycuda.compiler import SourceModule

mod = SourceModule("""
    #include <stdio.h>

    __global__ void say_hi()
    {
      printf("Soy el thread numero %d en threadIdx.x:%d.threadIdx.y:%d  blockIdx.:%d blockIdx.y:%d blockDim.x:%d blockDim.y:%d\\n",(threadIdx.x+threadIdx.y*blockDim.x+(blockIdx.x*blockDim.x*blockDim.y)+(blockIdx.y*blockDim.x*blockDim.y)),threadIdx.x, threadIdx.y,blockIdx.x,blockIdx.y,blockDim.x,blockDim.y);
    }
    """)

func = mod.get_function("say_hi")
func(block=(4,4,1),grid=(2,2,1))


Writing ./Programas/saludar_bloques.py

Lo ejecutamos. Esperamos que se ejecuten uno por uno los bloques; veamos qué es lo que pasa.


In [4]:
!ipython ./Programas/saludar_bloques.py


Soy el thread numero 0 en threadIdx.x:0.threadIdx.y:0  blockIdx.:0 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 1 en threadIdx.x:1.threadIdx.y:0  blockIdx.:0 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 2 en threadIdx.x:2.threadIdx.y:0  blockIdx.:0 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 3 en threadIdx.x:3.threadIdx.y:0  blockIdx.:0 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 4 en threadIdx.x:0.threadIdx.y:1  blockIdx.:0 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 5 en threadIdx.x:1.threadIdx.y:1  blockIdx.:0 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 6 en threadIdx.x:2.threadIdx.y:1  blockIdx.:0 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 7 en threadIdx.x:3.threadIdx.y:1  blockIdx.:0 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 8 en threadIdx.x:0.threadIdx.y:2  blockIdx.:0 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 9 en threadIdx.x:1.threadIdx.y:2  blockIdx.:0 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 10 en threadIdx.x:2.threadIdx.y:2  blockIdx.:0 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 11 en threadIdx.x:3.threadIdx.y:2  blockIdx.:0 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 12 en threadIdx.x:0.threadIdx.y:3  blockIdx.:0 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 13 en threadIdx.x:1.threadIdx.y:3  blockIdx.:0 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 14 en threadIdx.x:2.threadIdx.y:3  blockIdx.:0 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 15 en threadIdx.x:3.threadIdx.y:3  blockIdx.:0 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 32 en threadIdx.x:0.threadIdx.y:0  blockIdx.:1 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 33 en threadIdx.x:1.threadIdx.y:0  blockIdx.:1 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 34 en threadIdx.x:2.threadIdx.y:0  blockIdx.:1 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 35 en threadIdx.x:3.threadIdx.y:0  blockIdx.:1 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 36 en threadIdx.x:0.threadIdx.y:1  blockIdx.:1 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 37 en threadIdx.x:1.threadIdx.y:1  blockIdx.:1 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 38 en threadIdx.x:2.threadIdx.y:1  blockIdx.:1 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 39 en threadIdx.x:3.threadIdx.y:1  blockIdx.:1 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 40 en threadIdx.x:0.threadIdx.y:2  blockIdx.:1 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 41 en threadIdx.x:1.threadIdx.y:2  blockIdx.:1 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 42 en threadIdx.x:2.threadIdx.y:2  blockIdx.:1 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 43 en threadIdx.x:3.threadIdx.y:2  blockIdx.:1 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 44 en threadIdx.x:0.threadIdx.y:3  blockIdx.:1 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 45 en threadIdx.x:1.threadIdx.y:3  blockIdx.:1 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 46 en threadIdx.x:2.threadIdx.y:3  blockIdx.:1 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 47 en threadIdx.x:3.threadIdx.y:3  blockIdx.:1 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 16 en threadIdx.x:0.threadIdx.y:0  blockIdx.:0 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 17 en threadIdx.x:1.threadIdx.y:0  blockIdx.:0 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 18 en threadIdx.x:2.threadIdx.y:0  blockIdx.:0 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 19 en threadIdx.x:3.threadIdx.y:0  blockIdx.:0 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 20 en threadIdx.x:0.threadIdx.y:1  blockIdx.:0 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 21 en threadIdx.x:1.threadIdx.y:1  blockIdx.:0 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 22 en threadIdx.x:2.threadIdx.y:1  blockIdx.:0 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 23 en threadIdx.x:3.threadIdx.y:1  blockIdx.:0 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 24 en threadIdx.x:0.threadIdx.y:2  blockIdx.:0 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 25 en threadIdx.x:1.threadIdx.y:2  blockIdx.:0 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 26 en threadIdx.x:2.threadIdx.y:2  blockIdx.:0 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 27 en threadIdx.x:3.threadIdx.y:2  blockIdx.:0 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 28 en threadIdx.x:0.threadIdx.y:3  blockIdx.:0 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 29 en threadIdx.x:1.threadIdx.y:3  blockIdx.:0 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 30 en threadIdx.x:2.threadIdx.y:3  blockIdx.:0 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 31 en threadIdx.x:3.threadIdx.y:3  blockIdx.:0 blockIdx.y:1 blockDim.x:4 blockDim.y:4
Soy el thread numero 16 en threadIdx.x:0.threadIdx.y:0  blockIdx.:1 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 17 en threadIdx.x:1.threadIdx.y:0  blockIdx.:1 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 18 en threadIdx.x:2.threadIdx.y:0  blockIdx.:1 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 19 en threadIdx.x:3.threadIdx.y:0  blockIdx.:1 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 20 en threadIdx.x:0.threadIdx.y:1  blockIdx.:1 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 21 en threadIdx.x:1.threadIdx.y:1  blockIdx.:1 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 22 en threadIdx.x:2.threadIdx.y:1  blockIdx.:1 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 23 en threadIdx.x:3.threadIdx.y:1  blockIdx.:1 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 24 en threadIdx.x:0.threadIdx.y:2  blockIdx.:1 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 25 en threadIdx.x:1.threadIdx.y:2  blockIdx.:1 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 26 en threadIdx.x:2.threadIdx.y:2  blockIdx.:1 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 27 en threadIdx.x:3.threadIdx.y:2  blockIdx.:1 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 28 en threadIdx.x:0.threadIdx.y:3  blockIdx.:1 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 29 en threadIdx.x:1.threadIdx.y:3  blockIdx.:1 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 30 en threadIdx.x:2.threadIdx.y:3  blockIdx.:1 blockIdx.y:0 blockDim.x:4 blockDim.y:4
Soy el thread numero 31 en threadIdx.x:3.threadIdx.y:3  blockIdx.:1 blockIdx.y:0 blockDim.x:4 blockDim.y:4

En efecto se ejecutaron uno a la vez.

Rendimiento

A continuación agregamos un código para comparar el rendimiento de varias formas de operar. Cada una es clarificada en el código.


In [5]:
import pycuda.driver as drv
import pycuda.tools
import pycuda.autoinit
import numpy
from pycuda.compiler import SourceModule
import pycuda.gpuarray as gpuarray
import pycuda.cumath
from pycuda.elementwise import ElementwiseKernel

blocks = 64
block_size = 128
valores = blocks * block_size

print "Usando", valores, "valores"

# Número de iteraciones para los cálculos
n_iter = 100000
print "Calculando %d iteraciones" % (n_iter)

# Crear dos timers
inicio = drv.Event()
fin = drv.Event()


# SourceModele

mod = SourceModule("""
__global__ void gpusin(float *dest, float *a, int n_iter)
{
  const int i = blockDim.x*blockIdx.x + threadIdx.x;
  for(int n = 0; n < n_iter; n++) {
    a[i] = sin(a[i]);
  }
  dest[i] = a[i];
}
""")

gpusin = mod.get_function("gpusin")

# creamos un arreglo 1s
a = numpy.ones(valores).astype(numpy.float32)
# creamos un arreglo para guardar el resultado
dest = numpy.zeros_like(a)

inicio.record() # comenzamos a tomar el tiempo
gpusin(drv.Out(dest), drv.In(a), numpy.int32(n_iter), grid=(blocks,1), block=(block_size,1,1) )
fin.record() # terminamos de tomar el tiempo
# calculamos cuánto duró la corrida
fin.synchronize()
segs = inicio.time_till(fin)*1e-3
print "Tiempo con SourceModule y primeros 3 resultados:"
print "%fs, %s" % (segs, str(dest[:3]))


# Usando ElementwiseKernel con sin en un ciclo for en C

kernel = ElementwiseKernel(
   "float *a, int n_iter",
   "for(int n = 0; n < n_iter; n++) { a[i] = sin(a[i]);}",
   "gpusin")

a = numpy.ones(valores).astype(numpy.float32)
a_gpu = gpuarray.to_gpu(a)
inicio.record() # comenzamos a tomar el tiempo
kernel(a_gpu, numpy.int(n_iter))
fin.record() # terminamos de tomar el tiempo
# calculamos cuánto duró la corrida
fin.synchronize()
segs = inicio.time_till(fin)*1e-3
print "Tiempo con Elementwise y primeros 3 resultados:"
print "%fs, %s" % (segs, str(a_gpu.get()[:3]))



# Elementwise haciendo el loop en Python

kernel = ElementwiseKernel(
   "float *a",
   "a[i] = sin(a[i]);",
   "gpusin")

a = numpy.ones(valores).astype(numpy.float32)
a_gpu = gpuarray.to_gpu(a)
inicio.record() # comenzamos a tomar el tiempo
for i in range(n_iter):
    kernel(a_gpu)
fin.record() # terminamos de tomar el tiempo
# calculamos cuánto duró la corrida
fin.synchronize()
segs = inicio.time_till(fin)*1e-3
print "Tiempo con Elementwise en Python y primeros 3 resultados:"
print "%fs, %s" % (segs, str(a_gpu.get()[:3]))


# GPUArray
# El resultado se copia a la memoria principal en cada iteración (esto es un cuello de botella)

a = numpy.ones(valores).astype(numpy.float32)
a_gpu = gpuarray.to_gpu(a)
inicio.record() # comenzamos a tomar el tiempo
for i in range(n_iter):
    a_gpu = pycuda.cumath.sin(a_gpu)
fin.record() # terminamos de tomar el tiempo
# calculamos cuánto duró la corrida
fin.synchronize()
segs = inicio.time_till(fin)*1e-3
print "Tiempo con GPUArray y primeros 3 resultados:"
print "%fs, %s" % (segs, str(a_gpu.get()[:3]))



# CPU 

a = numpy.ones(valores).astype(numpy.float32)
inicio.record() # comenzamos a tomar el tiempo
inicio.synchronize()

for i in range(n_iter):
    a = numpy.sin(a)

fin.record() # terminamos de tomar el tiempo
# calculamos cuánto duró la corrida
fin.synchronize()
segs = inicio.time_till(fin)*1e-3
print "Tiempo con CPU y primeros 3 resultados:"
print "%fs, %s" % (segs, str(a[:3]))


Usando 8192 valores
Calculando 100000 iteraciones
Tiempo con SourceModule y primeros 3 resultados:
0.102129s, [ 0.005477  0.005477  0.005477]
Tiempo con Elementwise y primeros 3 resultados:
0.232770s, [ 0.005477  0.005477  0.005477]
Tiempo con Elementwise en Python y primeros 3 resultados:
1.979794s, [ 0.005477  0.005477  0.005477]
Tiempo con GPUArray y primeros 3 resultados:
6.091540s, [ 0.005477  0.005477  0.005477]
tiempo con CPU y primeros 3 resultados:
4.378356s, [ 0.005477  0.005477  0.005477]

Como podemos ver, no siempre es más eficiente hacer las cosas en el GPU; obviamente esta vez no lo fue porque forzamos al código a ser ineficiente, sin embargo es usual que el programador introduzca cuellos de botella como el que introdujimos sin siquiera darse cuenta.

En los materiales adicionales se agregan unas lecturas que pueden resultar ilustrativas de este punto.