2d Stencil Operation For cuda

0 views
Skip to first unread message

Ebad Ali

unread,
Oct 5, 2016, 12:08:39 PM10/5/16
to Numba Public Discussion - Public

Hi.
I am trying to run 2d Stencil Operation For cuda.
Trying to create LBP Textures From 256*256 image.


Any one addressing similar problem ?
 
 

Siu Kwan Lam

unread,
Oct 6, 2016, 9:26:57 AM10/6/16
to Numba Public Discussion - Public
Stencil operation using numba.cuda.jit is possible.  If you have a serial version, a naive numba-cuda impl would be to use the loop body as the cuda kernel, which is launched on a 2d thread-grid:

serial version:

for x in range(xsize):
    for y in range(ysize):
        stencil_code(x, y)

cuda version:

@cuda.jit
def kernel(...):
    x, y = cuda.grid(2)
    if x < xsize and y < ysize:
        stencil_code(x, y)


Then, you can convert it into a block algorithm and take advantage of cuda shared memory for better performance.  
(note: we don't have support for cuda texture.)



--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.
To unsubscribe from this group and stop receiving emails from it, send an email to numba-users...@continuum.io.
To post to this group, send email to numba...@continuum.io.
To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/2b9f1ad4-e5ea-489f-9a4f-5caee196e9c1%40continuum.io.
For more options, visit https://groups.google.com/a/continuum.io/d/optout.
--
Siu Kwan Lam
Software Engineer
Continuum Analytics

Ebad Ali

unread,
Oct 6, 2016, 9:52:53 AM10/6/16
to numba...@continuum.io
Hi,

Thank you for the Response.
Although i have implemented it.

For LBP Texture Some thing like:

@jit([void(int32[:,:], int32[:])], target='cuda')
def unsharp_masking(arry, hist):

    # We have 10*10 threads per block
    A = cuda.shared.array(shape=(32,32), dtype=int32)
    # H = cuda.shared.array(BIN_COUNT, dtype=int32)

    x,y = cuda.grid(2)

    ty = cuda.threadIdx.x
    tx = cuda.threadIdx.y

    A[ty,tx] = arry[x,y]


    cuda.syncthreads()

    threadCountX = A.shape[0] - 1
    threadCountY = A.shape[1] - 1
    # If within x range and y range then calculate the LBP discriptor along
    # with histogram value to specific bin

    # Other wise Ignore the Value
    if (ty > 0 and  (threadCountX-ty) > 0 ) and (tx > 0 and (threadCountY-tx) > 0):
    #     # You can do the Processing here. ^_^
        code = 0
        #  We need to make sure that each value is accessable to each thread
        #  TODO: make them atomic
        center = A[ty, tx]
        # Lets try averaging,
        code += A[ty-1][tx-1]*-1
        code += A[ty][tx-1]*-2
        code += A[ty+1][tx-1]*-1
        code += A[ty+1][tx]*-2
        code += A[ty+1][tx+1]*-1
        code += A[ty][tx+1]*-2
        code += A[ty-1][tx+1]*-1
        code += A[ty-1][tx-1]*-2

        code = code / 16

        code = ( code - center)

        A[ty,tx] = code

        # cuda.atomic.add(A, (ty,tx),code)
        cuda.syncthreads()

        val  = A[ty,tx]
        cuda.atomic.add(arry, (x,y),val)

Thanks again.

--------------
EBAD ALI

On Thu, Oct 6, 2016 at 6:26 PM, Siu Kwan Lam <s...@continuum.io> wrote:
Stencil operation using numba.cuda.jit is possible.  If you have a serial version, a naive numba-cuda impl would be to use the loop body as the cuda kernel, which is launched on a 2d thread-grid:

serial version:

for x in range(xsize):
    for y in range(ysize):
        stencil_code(x, y)

cuda version:

@cuda.jit
def kernel(...):
    x, y = cuda.grid(2)
    if x < xsize and y < ysize:
        stencil_code(x, y)


Then, you can convert it into a block algorithm and take advantage of cuda shared memory for better performance.  
(note: we don't have support for cuda texture.)


On Wed, Oct 5, 2016 at 11:08 AM Ebad Ali <ebad...@gmail.com> wrote:

Hi.
I am trying to run 2d Stencil Operation For cuda.
Trying to create LBP Textures From 256*256 image.


Any one addressing similar problem ?
 
 

--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.
To unsubscribe from this group and stop receiving emails from it, send an email to numba-users+unsubscribe@continuum.io.
--
Siu Kwan Lam
Software Engineer
Continuum Analytics

--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.
To unsubscribe from this group and stop receiving emails from it, send an email to numba-users+unsubscribe@continuum.io.

To post to this group, send email to numba...@continuum.io.
To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/CAHjo7%3DW%2BDb_CZVWoy7AqRVb8dGeOL5KUxcMieMcD7ZPiacnLsA%40mail.gmail.com.

Siu Kwan Lam

unread,
Oct 6, 2016, 10:38:41 AM10/6/16
to numba...@continuum.io
Sorry, I am confused.  Do you have any specific problem?  It seems like you already have something working.

To unsubscribe from this group and stop receiving emails from it, send an email to numba-users...@continuum.io.
--
Siu Kwan Lam
Software Engineer
Continuum Analytics

--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.
To unsubscribe from this group and stop receiving emails from it, send an email to numba-users...@continuum.io.

To post to this group, send email to numba...@continuum.io.
To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/CAHjo7%3DW%2BDb_CZVWoy7AqRVb8dGeOL5KUxcMieMcD7ZPiacnLsA%40mail.gmail.com.

--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.
To unsubscribe from this group and stop receiving emails from it, send an email to numba-users...@continuum.io.

To post to this group, send email to numba...@continuum.io.
To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/CAKzNkZcomrbZ3yCQbC_goJKW7S0FMkQQJtbjqeZH_j-FJYkP-Q%40mail.gmail.com.

For more options, visit https://groups.google.com/a/continuum.io/d/optout.

Ebad Ali

unread,
Oct 6, 2016, 10:42:00 AM10/6/16
to numba...@continuum.io

Sorry, mind my english.

I have just shared the solution of 2d stencil code. May be you guys could add similar example in git repo.

Sorry about the confusion.


To unsubscribe from this group and stop receiving emails from it, send an email to numba-users+unsubscribe@continuum.io.
--
Siu Kwan Lam
Software Engineer
Continuum Analytics

--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.
To unsubscribe from this group and stop receiving emails from it, send an email to numba-users+unsubscribe@continuum.io.

To post to this group, send email to numba...@continuum.io.
To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/CAHjo7%3DW%2BDb_CZVWoy7AqRVb8dGeOL5KUxcMieMcD7ZPiacnLsA%40mail.gmail.com.

--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.
To unsubscribe from this group and stop receiving emails from it, send an email to numba-users+unsubscribe@continuum.io.

To post to this group, send email to numba...@continuum.io.
To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/CAKzNkZcomrbZ3yCQbC_goJKW7S0FMkQQJtbjqeZH_j-FJYkP-Q%40mail.gmail.com.
For more options, visit https://groups.google.com/a/continuum.io/d/optout.
--
Siu Kwan Lam
Software Engineer
Continuum Analytics

--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.
To unsubscribe from this group and stop receiving emails from it, send an email to numba-users+unsubscribe@continuum.io.

To post to this group, send email to numba...@continuum.io.
To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/CAHjo7%3DU2Oxn0DztCvz2SMcknxA4bthQttcQNex%2BnkgbjtCqfWw%40mail.gmail.com.

Siu Kwan Lam

unread,
Oct 6, 2016, 10:49:27 AM10/6/16
to numba...@continuum.io
Oh, Thanks! =)

To unsubscribe from this group and stop receiving emails from it, send an email to numba-users...@continuum.io.
--
Siu Kwan Lam
Software Engineer
Continuum Analytics

--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.
To unsubscribe from this group and stop receiving emails from it, send an email to numba-users...@continuum.io.

To post to this group, send email to numba...@continuum.io.
To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/CAHjo7%3DW%2BDb_CZVWoy7AqRVb8dGeOL5KUxcMieMcD7ZPiacnLsA%40mail.gmail.com.

--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.
To unsubscribe from this group and stop receiving emails from it, send an email to numba-users...@continuum.io.

To post to this group, send email to numba...@continuum.io.
To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/CAKzNkZcomrbZ3yCQbC_goJKW7S0FMkQQJtbjqeZH_j-FJYkP-Q%40mail.gmail.com.
For more options, visit https://groups.google.com/a/continuum.io/d/optout.
--
Siu Kwan Lam
Software Engineer
Continuum Analytics

--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.
To unsubscribe from this group and stop receiving emails from it, send an email to numba-users...@continuum.io.

To post to this group, send email to numba...@continuum.io.
To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/CAHjo7%3DU2Oxn0DztCvz2SMcknxA4bthQttcQNex%2BnkgbjtCqfWw%40mail.gmail.com.

For more options, visit https://groups.google.com/a/continuum.io/d/optout.

--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.
To unsubscribe from this group and stop receiving emails from it, send an email to numba-users...@continuum.io.

To post to this group, send email to numba...@continuum.io.
To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/CAKzNkZfqgmtb8-QtV7O1dBMNQgHUTvQUtzZemmcy6VCwmF_CEg%40mail.gmail.com.

For more options, visit https://groups.google.com/a/continuum.io/d/optout.
Reply all
Reply to author
Forward
0 new messages