Debug CUDA kernel printing pointers

Numba Community Support
1

Dear all,

I am trying to understand how to debug when using @cuda.jit, however I came up with some doubts, more specifically how my arrays mutate within the several threads. So to test, I desire to understand the values of my array in a determined absolute position loop, when I debug I and print(A) and for my surprise I get a pointer.
So my questions are:

  1. How can I print my array values starting through the pointer?

(Please notice, to debug I have to execute with the SET “NUMBA_ENABLE_CUDASIM=1”)
Here is the code:

from numba import cuda

import numpy as np

import os
@cuda.jit

def vec_add(A, B, out):

    x = cuda.threadIdx.x

    bx = cuda.blockIdx.x

    bdx = cuda.blockDim.x

    i = bx * bdx + x

    if i==10:

        from pdb import set_trace; set_trace()

        print(A)#print(A) in (Pdb) returns a pointer <numba.cuda.simulator.cudadrv.devicearray.FakeWithinKernelCUDAArray object at 0x00000288E604F040>

        print(A[i])#Prints 0.0, It is reading from the mutated A array.

    for k in range(len(A)):

        A[k]=0

    out[i] = A[i] + B[i]

a=np.ones(1000)

b=np.arange(1000)

c=np.arange(1000)

threadsperblock = 32

blockspergrid = (a.size + (threadsperblock - 1)) // threadsperblock

vec_add[blockspergrid, threadsperblock](a,b,c)

print(c)

Many thanks!

2

EDIT 2: Trying to answer your question. Are you looking to see values of A? if so, you can try A._item. (Not sure if i have correctly answered your question. )

EDIT1: I am using windows and i ran “SET NUMBA_ENABLE_CUDASIM=1” before running the program and it worked.

Hi, just wondering how exactly did you set NUMBA_ENABLE_CUDASIM to 1?

I tried the following

import os
from pdb import set_trace

os.environ['NUMBA_ENABLE_CUDASIM'] = "1"

@cuda.jit
def atomic_compare_and_swap(res, old, ary, fill_val):
    gid = cuda.grid(1)
    if gid < res.size:
        set_trace()
        out = cuda.atomic.compare_and_swap(res[gid:], fill_val, ary[gid])
        old[gid] = out

But, I keep getting the following error

raise TypingError(msg, loc=inst.loc)
numba.core.errors.TypingError: Failed in nopython mode pipeline (step: nopython frontend)
Untyped global name ‘set_trace’: cannot determine Numba type of <class ‘function’>

File “test_atomic_cas.py”, line 12:
def atomic_compare_and_swap(res, old, ary, fill_val):

if gid < res.size: set_trace() ^
3

If you want to do

os.environ['NUMBA_ENABLE_CUDASIM'] = "1"

Then you need to do it before you import Numba, because Numba reads the environment variable and initializes the real target or the simulator accordingly on import.

1 Like