Why are streams in NUMBA CUDA API not improving my code performance?

Viewed 63

I'm trying to speed up my code (a simple box blur filter), and in order to do so, I've tried to improve it by using streams (avoiding the std. stream). I don't have much experience with Numba CUDA API, but I had some experience in CUDA C.

As a quick side note, the code is not clean I was just testing some functionalities.

Here is how my kernel is declared:

# Box blur for RGB image with RGB pixel values in [0, 255]
@cuda.jit
def imageProcessing(img, x, blurIntensity):
    pos = cuda.grid(1)
    if pos < img.size:
        mult = 0
        div = 0
        for i in range(1, blurIntensity):
            newX = x * (i * 3)
            if pos - newX >= 0:
                mult += img[pos - newX]
                div += 1
            if pos - newX - 3 >= 0:
                mult += img[pos - newX - 3]
                div += 1
            if pos - newX + 3 >= 0:
                mult += img[pos - newX + 3]
                div += 1
            if pos + newX < img.size:
                mult += img[pos + newX]
                div += 1
            if pos + newX + 3 < img.size:
                mult += img[pos + newX + 3]
                div += 1
            if pos + newX - 3 < img.size:
                mult += img[pos + newX - 3]
                div += 1
            if pos - (i * 3) >= 0:
                mult += img[pos - (i * 3)]
                div += 1
            if pos + (i * 3) < img.size:
                mult += img[pos + (i * 3)]
                div += 1
        if mult > 0:
            img[pos] = mult / div

and then the code from where I call the kernel:

original_shape = self.img.shape
img = self.img.flatten()

seconds = time.time()
stream = cuda.stream()
n = original_shape[0] * 3
chunks = math.ceil(img.size / n)
threadsperblock = 32
blockspergrid = n // 32
with stream.auto_synchronize():
    for i in range(chunks):
        start = i * n
        end = start + n
        splittedImage = img[start:end]
        d_img = cuda.to_device(splittedImage, stream=stream)
        imageProcessing[blockspergrid, threadsperblock, stream](d_img, original_shape[0], intensitySlider.get())
        d_img.copy_to_host(img[start:end], stream=stream)
img = np.reshape(img, original_shape)
self.display_image(image=img)
self.filtered_image = img

I know that a possible improvement is to use a 2D grid, unwrap the loop inside the kernel, and so many more. What I wish to understand is why this code (theoretically asynchronous) runs slower than that piece of code (that should be synchronous over the NULL stream):

original_shape = self.img.shape
img = self.img.flatten()

seconds = time.time()
threadsperblock = 32
blockspergrid = (img.size + (threadsperblock - 1)) // threadsperblock
d_img = cuda.to_device(img)
imageProcessing[blockspergrid, threadsperblock](d_img, original_shape[0], intensitySlider.get())
d_img.copy_to_host(img)

I went through the documentation but couldn't find an answer to this, am I missing something?

I did some rough statistics and the async code runs more than 2 times slower than the synchronous one.

0 Answers
Related