PyCUDAラプラシアンフィルタ完成版を使って速度比較

前回このブログで作成したPyCUDA laplacian filterはできそこないだったので、今回新たにPyCUDA用のラプラシアンフィルタのコードを書き直した。

スポンサーリンク

Introducing the discrete Laplacian

import skimage.data
import skimage.color
from skimage.filters import laplace
import numpy as np
import matplotlib.pyplot as plt
%matplotlib inline
image = skimage.data.astronaut()
image = skimage.color.rgb2gray(image)
print(image.shape)
print(image.dtype)
(512, 512)
float64

Scikit-Image implementation

def laplace_skimage(image):
    """Applies Laplace operator to 2D image using skimage implementation. 
    Then tresholds the result and returns boolean image."""
    laplacian = laplace(image)
    thresh = np.abs(laplacian) > 0.05
    return thresh
edges = laplace_skimage(image)
edges.shape
(512, 512)
def compare(left, right):
    """Compares two images, left and right."""
    fig, ax = plt.subplots(1, 2, figsize=(20, 10))
    ax[0].imshow(left, cmap='gray')
    ax[1].imshow(right, cmap='gray')
compare(left=image, right=edges)

NumPy implementation

def laplace_numpy(image):
    """Applies Laplace operator to 2D image using our own NumPy implementation. 
    Then tresholds the result and returns boolean image."""
    laplacian = image[:-2, 1:-1] + image[2:, 1:-1] + image[1:-1, :-2] + image[1:-1, 2:] - 4*image[1:-1, 1:-1]
    thresh = np.abs(laplacian) > 0.05
    return thresh
laplace_numpy(image).shape
(510, 510)
compare(image, laplace_numpy(image))
np.allclose(laplace_skimage(image)[1:-1, 1:-1], laplace_numpy(image))
True

Cython

%load_ext cython
%%cython
import numpy as np
cimport numpy as cnp
import cython

@cython.boundscheck(False) # turn off bounds-checking for entire function
@cython.wraparound(False)  # turn off negative index wrapping for entire function
def laplace_cython(cnp.ndarray[double, ndim=2] image):
    """Applies Laplace operator to 2D image, then tresholds the result and returns boolean image.
    Cython implementation."""
    cdef int h = image.shape[0]
    cdef int w = image.shape[1]
    cdef cnp.ndarray[double, ndim=2] laplacian = np.empty((w-2, h-2), dtype=np.double)
    cdef int i, j
    for i in range(1, h-1):
        for j in range(1, w-1):
            laplacian[i-1, j-1] = image[i-1, j] + image[i+1, j] + image[i, j-1] + image[i, j+1] - 4*image[i, j]
    thresh = np.abs(laplacian) > 0.05
    return thresh
np.allclose(laplace_numpy(image), laplace_cython(image))
True

Pythran

%load_ext pythran.magic
%%pythran
#pythran export laplace_pythran_highlevel(float[][])
import numpy as np
def laplace_pythran_highlevel(image):
    """Laplace operator in NumPy for 2D images. Pythran accelerated."""
    laplacian = image[:-2, 1:-1] + image[2:, 1:-1] + image[1:-1, :-2] + image[1:-1, 2:] - 4*image[1:-1, 1:-1]
    thresh = np.abs(laplacian) > 0.05
    return thresh
np.allclose(laplace_numpy(image), laplace_pythran_highlevel(image))
True

Numba

from numba import jit

@jit(nopython=True, fastmath = True, parallel = True, nogil = True)
def laplace_numba(image):
    """Laplace operator for 2D images. Numba accelerated."""
    h = image.shape[0]
    w = image.shape[1]
    laplacian = np.empty((h - 2, w - 2))
    for i in range(1, h - 1):
        for j in range(1, w - 1):
            laplacian[i-1, j-1] = np.abs(image[i-1, j] + image[i+1, j] + image[i, j-1] + image[i, j+1] - 4*image[i, j]) > 0.05
    return laplacian
np.allclose(laplace_numpy(image), laplace_numba(image))
True

numba guvectorize

from numba import guvectorize

@guvectorize('void(float64[:, :], float64[:, :])', "(m, n)->(m, n)", \
           nopython=True, fastmath = True)
def laplace_numba_guvectorize(image, laplacian):
    """Laplace operator in NumPy for 2D images. Numba accelerated."""
    h = image.shape[0]
    w = image.shape[1]
    for i in range(1, h - 1):
        for j in range(1, w - 1):
            laplacian[i-1, j-1] = np.abs(4 * image[i, j] - image[i - 1, j] - \
               image[i + 1, j] - image[i, j + 1] - image[i, j - 1]) > 0.05
laplacian = np.empty_like(image)
laplace_numba_guvectorize(image, laplacian);
np.allclose(laplace_numpy(image), laplacian[:-2, :-2])
True

PyCUDA

import pycuda.gpuarray as gpuarray
import pycuda.driver as cuda
import pycuda.autoinit
from pycuda.compiler import SourceModule

image = image.astype(np.float32)
(h,w)=image.shape
print (h,w)
print (image.dtype)

mod_copy_texture=SourceModule(
"""
texture<float,2>tex;
__global__ void  copy_texture_kernel(float *C,float * data)
 {
  int i = threadIdx.x+(blockIdx.x*(blockDim.x));
  int j = threadIdx.y+(blockIdx.y*(blockDim.y));
  int h=C[0];
  int w=C[1];
  while(i<w-1 && j<h-1)
  {
  data[(i-1)+w*(j-1)] = abs(4*tex2D(tex,j,i)-tex2D(tex,j-1,i)-tex2D(tex,j+1,i)\
  -tex2D(tex,j,i+1)-tex2D(tex,j,i-1)) > 0.05;
  __syncthreads();
  j += blockDim.y * gridDim.y;
  i += blockDim.x * gridDim.x;
  }
}
""")
copy_texture_func = mod_copy_texture.get_function("copy_texture_kernel")
texref = mod_copy_texture.get_texref("tex")
cuda.matrix_to_texref(image , texref , order = "F")
#gpu_output = np.empty_like(image)
a_gpu = gpuarray.to_gpu(np.array([image.shape],dtype = np.float32))
c_gpu = gpuarray.to_gpu(np.empty_like(image))
copy_texture_func(a_gpu,c_gpu,\
          block=(32,32,1), grid=(h//32,w//32,1), texrefs=[texref])
512 512
float32
%timeit copy_texture_func(a_gpu,c_gpu,\
          block=(16,16,1), grid=(h//16,w//16,1), texrefs=[texref])
28.6 µs ± 195 ns per loop (mean ± std. dev. of 7 runs, 10000 loops each)
compare(c_gpu.get(), laplace_numpy(image))
np.allclose(laplace_numpy(image), c_gpu.get()[:-2, :-2])
True

Wrap-up and plots

timings = {}
for func in [laplace_skimage, laplace_numpy, laplace_cython, laplace_pythran_highlevel, laplace_numba]:
    t = %timeit -o func(image)
    timings[func.__name__] = t
2.25 ms ± 16.3 µs per loop (mean ± std. dev. of 7 runs, 100 loops each)
1.08 ms ± 17.5 µs per loop (mean ± std. dev. of 7 runs, 1000 loops each)
516 µs ± 6.31 µs per loop (mean ± std. dev. of 7 runs, 1000 loops each)
321 µs ± 604 ns per loop (mean ± std. dev. of 7 runs, 1000 loops each)
154 µs ± 262 ns per loop (mean ± std. dev. of 7 runs, 10000 loops each)
t = %timeit -o laplace_numba_guvectorize(image, laplacian);
timings['laplace_numba_guvectorize'] = t
363 µs ± 2.24 µs per loop (mean ± std. dev. of 7 runs, 1000 loops each)
t = %timeit -o copy_texture_func(a_gpu,c_gpu,\
          block=(16,16,1), grid=(h//16,w//16,1), texrefs=[texref]);
timings['laplace_pycuda'] = t
26.8 µs ± 254 ns per loop (mean ± std. dev. of 7 runs, 10000 loops each)
import pandas as pd

pd.Series({key: timings[key].average * \
 1e6 for key in timings}).to_frame(name='timings (μs)').\
sort_values(by='timings (μs)')
timings (μs)
laplace_pycuda 26.836729
laplace_numba 154.318407
laplace_pythran_highlevel 321.117471
laplace_numba_guvectorize 363.161995
laplace_cython 515.931875
laplace_numpy 1083.795276
laplace_skimage 2254.967624
fig, ax = plt.subplots(figsize=(20, 12))
plt.rcParams["font.size"] = "25"
pd.Series({key: timings[key].average * \
  1e6 for key in timings}).to_frame(name='timings (μs)').\
sort_values(by='timings (μs)').plot(kind='barh', ax=ax)
<matplotlib.axes._subplots.AxesSubplot at 0x7f2565c0a6a0>

前回のlaplace_pycudaの635.745086μ秒から、今回の26.836729μ秒は23.7倍も高速化されている。さらに前回はnp.allcloseがFalseだったが、今回はTrueになっている。


参考サイトhttps://github.com/

参考サイトOptimizing your code with NumPy, Cython, pythran and numba