{"action":"created","issue":{"url":"https://api.github.com/repos/numba/numba/issues/1891","repository_url":"https://api.github.com/repos/numba/numba","labels_url":"https://api.github.com/repos/numba/numba/issues/1891/labels{/name}","comments_url":"https://api.github.com/repos/numba/numba/issues/1891/comments","events_url":"https://api.github.com/repos/numba/numba/issues/1891/events","html_url":"https://github.com/numba/numba/issues/1891","id":154146053,"number":1891,"title":"CUDA kernel errors when using local variable inside loop","user":{"login":"anthls","id":14816341,"avatar_url":"https://avatars.githubusercontent.com/u/14816341?v=3","gravatar_id":"","url":"https://api.github.com/users/anthls","html_url":"https://github.com/anthls","followers_url":"https://api.github.com/users/anthls/followers","following_url":"https://api.github.com/users/anthls/following{/other_user}","gists_url":"https://api.github.com/users/anthls/gists{/gist_id}","starred_url":"https://api.github.com/users/anthls/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/anthls/subscriptions","organizations_url":"https://api.github.com/users/anthls/orgs","repos_url":"https://api.github.com/users/anthls/repos","events_url":"https://api.github.com/users/anthls/events{/privacy}","received_events_url":"https://api.github.com/users/anthls/received_events","type":"User","site_admin":false},"labels":[],"state":"open","locked":false,"assignee":null,"milestone":null,"comments":5,"created_at":"2016-05-11T02:17:10Z","updated_at":"2016-05-24T00:46:42Z","closed_at":null,"body":"I have been attempting to implement a CUDA kernel and come across some strange behaviour. The behaviour relates to the use of local variables inside loops within the kernel function. The heart of the matter is that kernel code (see numbered code below) gives different results to those calculated using Numpy (see test case code at the bottom of the post). Specifically:\r\n\r\na) Using the kernel code as listed below (and exercising it using the test case) results in the follow:\r\nResult from CUDA kernel : -1394.3506603578408\r\nResult from Numpy : 300.53724387519816\r\n\r\nb) Uncommenting line 26 of the kernel code causes the Numpy and CUDA kernel results to match.\r\n\r\nUsing the CUDA simulator generates results that match in both a) and b).\r\n\r\nI am using Python 2.7.11, Anaconda 2.5.0 and Numba 0.25.0.\r\n\r\nI have created this example in order to demonstrate the problem I am seeing in a different kernel implementation so my apologies for the contrivance.\r\n\r\nAny help would be most appreciated!\r\n\r\nKernel code:\r\n```python\r\n1   @cuda.jit(device=True)\r\n2   def dev_sum(a, b):\r\n3       return a + b\r\n4   \r\n5   @cuda.jit\r\n6   def cu_sum(a, num_to_sums, b):\r\n7       sa = cuda.shared.array(shape=(512,), dtype=float64)\r\n8       tx = cuda.threadIdx.x\r\n9       bx = cuda.blockIdx.x\r\n10      bw = cuda.blockDim.x\r\n11      i = tx + bx * bw\r\n12  \r\n13      sa[tx] = 0\r\n14      if i < a.shape[0]:\r\n15          num_choices = 200\r\n16          num_to_sum = num_to_sums[i]\r\n17          l_a = cuda.local.array(shape=num_choices, dtype=float64)\r\n18  \r\n19          l_sum = float64(0.)\r\n20          for j in xrange(num_to_sum):\r\n21              l_a[j] = math.exp(a[i, j])\r\n22              l_sum += l_a[j]\r\n23  \r\n24          l_b = float64(0.)\r\n25          # uncommenting line 26 means results match those expected (i.e. from Numpy)\r\n26          #l_one = l_sum / l_sum\r\n27          for j in xrange(num_to_sum):\r\n28              # fails unless line 26 is uncommented\r\n29              l_b += math.log(l_a[j]) * l_sum / l_sum\r\n30  \r\n31          sa[tx] = l_b\r\n32          cuda.syncthreads()\r\n33          if tx == 0:\r\n34              # Uses the first thread of each block to perform the actual reduction\r\n35              s = sa[tx]\r\n36              for j in range(1, bw):\r\n37                  s = dev_sum(s, sa[j])\r\n38              b[bx] = s\r\n```\r\n\r\nTest case to exercise the kernel above and generate results for comparison:\r\n```python\r\nimport numpy as np\r\n\r\ndef test_cu_sum(self):\r\n    n = 1000 # total length of vector to be summed\r\n    np.random.seed(seed=1112)\r\n    a = np.random.uniform(low=-100, high=100, size=(n, 200)).astype(np.float64)\r\n    num_to_sums = np.random.randint(1, 200, n)\r\n\r\n    # add up a different number of items from each dimension of n\r\n    n_dim = 0\r\n    numpy_sum_a = np.float64(0.)\r\n    for num in num_to_sums:\r\n        for i in xrange(num):\r\n            numpy_sum_a += np.log(np.exp(a[n_dim, i]))\r\n        n_dim += 1\r\n\r\n    d_a = cuda.to_device(a)\r\n    grid_size = n\r\n    threads = 512\r\n    threads_per_block = (threads, 1)\r\n    blocks_per_grid = int(math.ceil(grid_size / float(threads)))\r\n    b = np.zeros((blocks_per_grid,), dtype=np.float64)\r\n    d_b = cuda.to_device(b)\r\n    d_num_to_sums = cuda.to_device(num_to_sums)\r\n\r\n    cu_sum[blocks_per_grid, threads_per_block](d_a, d_num_to_sums, d_b)\r\n    d_b.to_host()\r\n    sum_a = 0\r\n    for result_index in xrange(blocks_per_grid):\r\n        sum_a += b[result_index]\r\n\r\n    self.assertAlmostEqual(sum_a, numpy_sum_a)\r\n```"},"comment":{"url":"https://api.github.com/repos/numba/numba/issues/comments/221137116","html_url":"https://github.com/numba/numba/issues/1891#issuecomment-221137116","issue_url":"https://api.github.com/repos/numba/numba/issues/1891","id":221137116,"user":{"login":"anthls","id":14816341,"avatar_url":"https://avatars.githubusercontent.com/u/14816341?v=3","gravatar_id":"","url":"https://api.github.com/users/anthls","html_url":"https://github.com/anthls","followers_url":"https://api.github.com/users/anthls/followers","following_url":"https://api.github.com/users/anthls/following{/other_user}","gists_url":"https://api.github.com/users/anthls/gists{/gist_id}","starred_url":"https://api.github.com/users/anthls/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/anthls/subscriptions","organizations_url":"https://api.github.com/users/anthls/orgs","repos_url":"https://api.github.com/users/anthls/repos","events_url":"https://api.github.com/users/anthls/events{/privacy}","received_events_url":"https://api.github.com/users/anthls/received_events","type":"User","site_admin":false},"created_at":"2016-05-24T00:46:42Z","updated_at":"2016-05-24T00:46:42Z","body":"I have tried running the test case using the CUDA simulator (i.e. NUMBA_ENABLE_CUDASIM=1) and cannot replicate the difference in results. That is, results between those generated by Numpy and Numba match with the CUDA simulator enabled."}}
{"actor":{"display_login":"anthls"}}