cuda segfault nvvm with opt+debug

[stuartarchibald]

Reporting a bug

• I have tried using the latest released version of Numba (most recent is
  visible in the change log (https://github.com/numba/numba/blob/master/CHANGE_LOG).
• I have included a self contained code sample to reproduce the problem.
  i.e. it's possible to run as 'python bug.py'.

This segfaults:

from numba import cuda

@cuda.jit(debug=True)
def f(b, n):
    result = 0
    for i in range(n):
        result += b

f[1, 1](1, 2)

Under cuda-gdb the backtrace points to the nvvmCompilerProgram call in libnvvm.so.

Setting opt=False or debug=False avoids the segfault.

CC @gmarkall

[danielverschueren]

Thanks for opening this ticket @stuartarchibald.

I apologise if this is not the right place to post this, but I thought it might be related and helpful to the issue:

I am experiencing some more strange issues with while/for loops in numba: matrix indexing in a while loop also creates segfaults. I have attached a minimal example here that produces the error:

from numba import cuda
import numpy as np

# 3 kernels, 2 fail on segfault
# fmat       indexes a matrix b[i,j] with 2 variables i,j; fails on segfault
# flinTwo   indexes an array b[i+j] with 2 variables i,j; passes
# flinthree indexes an array b[i+m+j] with 3 variables i,j,m; fails on segfault

@cuda.jit(debug=True)
def fmat(b, n):
    # this kernel fails... segfault
    result = 0
    j = 1
    i = 0
    while (j < b.shape[0]): # don't overflow
        while (i < n) and (i < b.shape[1]): # don't overflow
            result = b[j,i]
            i += 1
        j += 1
        i = 0

@cuda.jit(debug=True)
def flinTwo(b, m, n):
    # this kernel is happy, no faults!
    result = 0
    j = 0
    i = 0
    while (j < 2):
        while (i < n):
            if (j+i+m) < len(b): # don't overflow
                result = b[j+i]
            i += 1
        j += 1
        i = 0

@cuda.jit(debug=True)
def flinThree(b, m, n):
    # this kernel is happy, no faults!
    result = 0
    j = 0
    i = 0
    while (j < 2):
        while (i < n):
            if (j+i+m) < len(b): # don't overflow
                result = b[j+i+m]
            i += 1
        j += 1
        i = 0


a = np.asarray([[0,1,2,3,4,5],[0,1,2,3,4,5]])
d = a.flatten()
a_d = cuda.to_device(a)
d_d = cuda.to_device(d)

flinTwo[1,1](d_d, 0, 5)
print('flinTwo completes, no failures')

# this segfaults 
flinThree[1, 1](d_d, 0, 5)

# this segfaults too
fmat[1,1](a_d, 5)

I assume this issue is related to the issue on which the ticket was opened, so I post it here. Any fix or workaround would be appreciated!

[gmarkall]

Turning on debug without turning off optimization is not supported by NVVM - noted in the description of the -g option in https://docs.nvidia.com/cuda/libnvvm-api/group__compilation.html#group__compilation_1g76ac1e23f5d0e2240e78be0e63450346.

To compile with debug, opt=False also needs adding:

from numba import cuda

@cuda.jit(debug=True, opt=False)
def f(b, n):
    result = 0
    for i in range(n):
        result += b

f[1, 1](1, 2)

and

from numba import cuda
import numpy as np

# 3 kernels, 2 fail on segfault
# fmat       indexes a matrix b[i,j] with 2 variables i,j; fails on segfault
# flinTwo   indexes an array b[i+j] with 2 variables i,j; passes
# flinthree indexes an array b[i+m+j] with 3 variables i,j,m; fails on segfault

@cuda.jit(debug=True, opt=False)
def fmat(b, n):
    # this kernel fails... segfault
    result = 0
    j = 1
    i = 0
    while (j < b.shape[0]): # don't overflow
        while (i < n) and (i < b.shape[1]): # don't overflow
            result = b[j,i]
            i += 1
        j += 1
        i = 0

@cuda.jit(debug=True, opt=False)
def flinTwo(b, m, n):
    # this kernel is happy, no faults!
    result = 0
    j = 0
    i = 0
    while (j < 2):
        while (i < n):
            if (j+i+m) < len(b): # don't overflow
                result = b[j+i]
            i += 1
        j += 1
        i = 0

@cuda.jit(debug=True, opt=False)
def flinThree(b, m, n):
    # this kernel is happy, no faults!
    result = 0
    j = 0
    i = 0
    while (j < 2):
        while (i < n):
            if (j+i+m) < len(b): # don't overflow
                result = b[j+i+m]
            i += 1
        j += 1
        i = 0


a = np.asarray([[0,1,2,3,4,5],[0,1,2,3,4,5]])
d = a.flatten()
a_d = cuda.to_device(a)
d_d = cuda.to_device(d)

flinTwo[1,1](d_d, 0, 5)
print('flinTwo completes, no failures')

# this segfaults 
flinThree[1, 1](d_d, 0, 5)

# this segfaults too
fmat[1,1](a_d, 5)

both execute for me with no segfault or other issue.

For now, (prior to other fixes / reorganization of the debug kwarg to disentangle exception checking and debuginfo generation), we should emit a warning to inform the user they ought to also use opt=False.

[gmarkall]

PR #7222 adds the relevant warning.

[sgbaird]

I think this might be the issue I'm experiencing in #7326

[gmarkall]

I think this might be the issue I'm experiencing in #7326

I don't see any use of debug=True there - why do you think this is the same issue? (have I missed the use of debug?)

[sgbaird]

It was with the environment variable (unless I'm mistaken about it's usage):

os.environ["NUMBA_CUDA_DEBUGINFO"] = "1"

[gmarkall]

Ah! You're right!