Edge Detection using PyOpenCL

Question

I was having an issue while writing a python script to detect the edge using PyOpenCl. I am new to OpenCL and while trying I encountered a problem which after many tweaks, I am unable to solve. below is the python code:

edge.py:

import numpy as np
import pyopencl as cl
from PIL import Image
from time import time


def getKernel(krnl):
    kernel = open(krnl).read()
    return kernel


def findedges(p,d,image):

    data = np.asarray(image).astype(np.uint8)

    platform = cl.get_platforms()[p]
    device = platform.get_devices()[d]
    cntx = cl.Context([device])
    queue = cl.CommandQueue(cntx)

    mf = cl.mem_flags
    im = cl.Buffer(cntx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=data)
    out = cl.Buffer(cntx,mf.WRITE_ONLY,data.nbytes)

    prgm = cl.Program(cntx,getKernel('edge.c')%(data.shape[1],data.shape[0])).build()

    prgm.detectedge(queue,data.shape,None,im,out)

    result = np.empty_like(data)

    cl.enqueue_copy(queue,result,out)
    result = result.astype(np.uint8)
    print(result)

    img = Image.fromarray(result)
    #img.show()
    img.save('coinsedge.png')


if __name__ == '__main__':

    image = Image.open('coins.png')
    #(1,0) is my platform 1, device 0 = "AMD gpu"  
    #(0,0) for intel processor 
    findedges(1,0,image)

and my kernel file : edge.c

__kernel void detectedge(__global int *im,__global int *out){
      int j = get_global_id(1);
      int i = get_global_id(0);
      int width = %d;
      int rown = %d;
      int value;


              value = -im[(i)*width + j] -  0* im[(i)*width + j+1] + im[(i)*width + j+2]
                      -2*im[(i+1)*width + j] +  0*im[(i+1)*width + j+1] + 2*im[(i+1)*width + j+2]
                      -im[(i+2)*width + j] -  0*im[(i+2)*width + j+1] + im[(i+2)*width + j+2];

              value = (value < 0   ? 0   : value);
              value = (value > 255 ? 255 : value);
              out[i*width + j] = value;

  }

Now there are no run-time warnings/errors but the output is what I was not expecting. Here is the input and its output: This is my input image:

And this is my output:

Answer 1

I think the issue is with the size of the input and output buffers you are passing to the kernel. The image is one uint8 per pixel and you sized your buffer according to data.nbytes. The kernel is reading and writing 4 byte pixels.

I only had to change this line and it works:

data = np.asarray(image).astype(np.int32)