title: Have trouble using numba atomic operation functions (cuda.atomic.compare_and_swap)
body: <p>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:</p>
  
  <p><a href="https://i.stack.imgur.com/9WZEP.png" rel="nofollow noreferrer">enter image description here</a></p>
  
  <pre><code>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 &gt;= 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()
  </code></pre>
  
  <p>Then I got this error:</p>
  
  <pre><code>
  TypingError: Invalid use of Function(&lt;class 'numba.cuda.stubs.atomic.compare_and_swap'&gt;) 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(&lt;class 'numba.cuda.stubs.atomic.compare_and_swap'&gt;)
  [2] During: typing of call at /home/qinyu/test.py (20)
  </code></pre>
  
  <p>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!</p>