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:
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.  During: resolving callee type: Function(<class 'numba.cuda.stubs.atomic.compare_and_swap'>)  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!
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.