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.