String Argument in CUDA Signature

Numba Support: How do I do ...?
1

The Numba documentation provides details regarding the types and signatures that are supported. For CUDA, it is common to declare a function that accepts numbers, arrays, and booleans with something like:

@cuda.jit(
    "(i8, f8, f8[:], i8[:], b1)"
)
def some_func(some_int, some_float, some_float_array, some_int_array, some_bool):

I was wondering if it is possible specify a string? Let’s say for a CUDA kernel that accepts an argument like:

# add CUDA signature here
def some_func(side="left"):
    if side == "left":
        # do something
    if side == "right":
        # do something else

I know this can be done for @njit but it wasn’t clear if this were possible in CUDA

2

Unfortunately not - trying to compile something that accepts a string:

from numba import cuda

@cuda.jit('void(int32[::1], unicode_type)')
def f(x, s):
    if s == 'left':
        r = 1
    elif s == 'right':
        r = 2
    else:
        r = 3
    x[0] = r

results in an error:

NotImplementedError: No definition for lowering <built-in method eq_impl of _dynfunc._Closure object at 0x7fe2cfe0b160>(unicode_type, unicode_type) -> bool

(this error message is a bug though, a NotImplementedError shouldn’t be thrown from lowering)

1 Like
3

Thank you @gmarkall! Would it be any different if it were a device function instead (i.e., a kernel that calls this function)?

@cuda.jit(device=True)
def f(x, s):
    if s == 'left':
        r = 1
    elif s == 'right':
        r = 2
    else:
        r = 3
    x[0] = r

Or would this also fall into the same fate? I guess there would also be no way to set a kernel variable to a string?

4

It would fall into the same fate - the error appeared due to a lack of support for comparing strings in CUDA.

For this use case, could an Enum class work instead? E.g.

import numpy as np
from numba import cuda
from enum import Enum


class Hand(Enum):
    LEFT = 1
    RIGHT = 2
    OTHER = 3


@cuda.jit()
def f(x, s):
    if s == Hand.LEFT:
        r = 7
    elif s == Hand.RIGHT:
        r = 31
    else:
        r = -1
    x[0] = r


x = cuda.device_array(1, np.int32)
f[1, 1](x, Hand.RIGHT)

print(x[0])
np.testing.assert_equal(x[0], 31)

runs and prints 31.

1 Like
5

Thank you! I will play with your suggestion