Have trouble using numba atomic operation functions (cuda.atomic.compare_and_swap)

I am trying to use Numba to write cuda kernels for my code. And somehow I wanna use the atomic operation in part of my code and I wrote a test kernel to see how cuda.atomic.compare_and_swap works. On the documentation it says this:

enter image description here

from numba import cuda
import numpy as np


@cuda.jit
def atomicCAS(N,out1):
    idx = cuda.threadIdx.x + cuda.blockIdx.x * cuda.blockDim.x
    if idx >= N:
        return
    A = out1[idx:]
    cuda.atomic.compare_and_swap(A,idx,0)

N    = 1024
out1 = np.arange(N)
out1 = np.zeros(N)
dout1 = cuda.to_device(out1)
tpb  = 32
bpg  = int(np.ceil(N/tpb))
atomicCAS[bpg,tpb](N,dout1)
hout1 = dout1.copy_to_host()

Then I got this error:

TypingError: Invalid use of Function(<class 'numba.cuda.stubs.atomic.compare_and_swap'>) with argument(s) of type(s): (array(float64, 1d, A), int64, Literal[int](0))
 * parameterized
In definition 0:
    All templates rejected with literals.
In definition 1:
    All templates rejected without literals.
This error is usually caused by passing an argument of a type that is unsupported by the named function.
[1] During: resolving callee type: Function(<class 'numba.cuda.stubs.atomic.compare_and_swap'>)
[2] During: typing of call at /home/qinyu/test.py (20)

This is a pretty naive code and I think I feed in the write type of variables but I got this typingerror. It worked pretty well with the other atomic operations in Numba. This is the only one that does not work for me. Can somebody help me figure out the problem or is there another alternative ways to do this? Thanks!

Answer

The key in the error message is this:

array(float64, 1d, A), int64, Literal[int](0))

CUDA atomicCAS only supports integer types. You cannot pass a floating point type.