{"action":"created","issue":{"url":"https://api.github.com/repos/inducer/pycuda/issues/46","labels_url":"https://api.github.com/repos/inducer/pycuda/issues/46/labels{/name}","comments_url":"https://api.github.com/repos/inducer/pycuda/issues/46/comments","events_url":"https://api.github.com/repos/inducer/pycuda/issues/46/events","html_url":"https://github.com/inducer/pycuda/issues/46","id":37888399,"number":46,"title":"pycuda._driver.LogicError: cuMemcpyHtoDAsync failed: invalid value","user":{"login":"adityaatluri","id":2188479,"avatar_url":"https://avatars.githubusercontent.com/u/2188479?v=3","gravatar_id":"","url":"https://api.github.com/users/adityaatluri","html_url":"https://github.com/adityaatluri","followers_url":"https://api.github.com/users/adityaatluri/followers","following_url":"https://api.github.com/users/adityaatluri/following{/other_user}","gists_url":"https://api.github.com/users/adityaatluri/gists{/gist_id}","starred_url":"https://api.github.com/users/adityaatluri/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/adityaatluri/subscriptions","organizations_url":"https://api.github.com/users/adityaatluri/orgs","repos_url":"https://api.github.com/users/adityaatluri/repos","events_url":"https://api.github.com/users/adityaatluri/events{/privacy}","received_events_url":"https://api.github.com/users/adityaatluri/received_events","type":"User","site_admin":false},"labels":[],"state":"closed","locked":false,"assignee":null,"milestone":null,"comments":7,"created_at":"2014-07-15T14:28:07Z","updated_at":"2015-08-03T10:34:54Z","closed_at":"2014-07-18T19:36:28Z","body":"Hi,\r\nI am trying to use async memcpy using pycuda. I am getting this error:\r\n```shell\r\nTraceback (most recent call last):\r\n  File \"streams.py\", line 29, in <module>\r\n    drv.memcpy_htod_async(a_gpu,a_pin)\r\npycuda._driver.LogicError: cuMemcpyHtoDAsync failed: invalid value\r\n```\r\nThe source code is:\r\n```python\r\nimport numpy as np\r\nimport pycuda.autoinit\r\nimport pycuda.driver as drv\r\n\r\nfrom pycuda.compiler import SourceModule\r\n\r\nmod = SourceModule(\"\"\"\r\n__global__ void add_them(long *dest, long *a, long *b){\r\nint tx = blockIdx.x;\r\ndest[tx] = a[tx] + b[tx];\r\n}\r\n\"\"\")\r\n\r\nadd_them = mod.get_function(\"add_them\")\r\n\r\na = np.random.randint(10, size = 65536)\r\nb = np.random.randint(10, size = 65536)\r\n\r\ndest = np.zeros_like(a)\r\na_gpu = drv.mem_alloc(a.nbytes)\r\nb_gpu = drv.mem_alloc(b.nbytes)\r\ndest_gpu = drv.mem_alloc(dest.nbytes)\r\n\r\na_pin = drv.register_host_memory(a)\r\nb_pin = drv.register_host_memory(b)\r\n\r\ndrv.memcpy_htod_async(a_gpu,a_pin)\r\ndrv.memcpy_htod_async(b_gpu,b_pin)\r\n\r\nadd_them(dest_gpu,a_gpu,b_gpu,block=(1,1,1),grid=(65536,1,1))\r\n\r\ndrv.memcpy_dtoh(dest,dest_gpu)\r\n```"},"comment":{"url":"https://api.github.com/repos/inducer/pycuda/issues/comments/127188152","html_url":"https://github.com/inducer/pycuda/issues/46#issuecomment-127188152","issue_url":"https://api.github.com/repos/inducer/pycuda/issues/46","id":127188152,"user":{"login":"sjperkins","id":3530212,"avatar_url":"https://avatars.githubusercontent.com/u/3530212?v=3","gravatar_id":"","url":"https://api.github.com/users/sjperkins","html_url":"https://github.com/sjperkins","followers_url":"https://api.github.com/users/sjperkins/followers","following_url":"https://api.github.com/users/sjperkins/following{/other_user}","gists_url":"https://api.github.com/users/sjperkins/gists{/gist_id}","starred_url":"https://api.github.com/users/sjperkins/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/sjperkins/subscriptions","organizations_url":"https://api.github.com/users/sjperkins/orgs","repos_url":"https://api.github.com/users/sjperkins/repos","events_url":"https://api.github.com/users/sjperkins/events{/privacy}","received_events_url":"https://api.github.com/users/sjperkins/received_events","type":"User","site_admin":false},"created_at":"2015-08-03T10:34:54Z","updated_at":"2015-08-03T10:34:54Z","body":"@inducer I've created two versions of this behaviour:\r\n\r\n- [CUDA C](https://gist.github.com/sjperkins/8cd836604ff7d2a2a3b3)\r\n- [PyCUDA](https://gist.github.com/sjperkins/d9e6db1b2d6038febb72)\r\n\r\nThe CUDA C version seems to run fine for me, but if `pycuda.driver.register_host_memory` is used, the PyCUDA versions fails with:\r\n\r\npycuda._driver.LogicError: cuMemcpyHtoDAsync failed: **invalid argument**\r\n\r\n If `pycuda.driver.pagelocked_empty` is used to create pinned memory instead, the PyCUDA version succeeds.\r\n\r\nIs it possible that the `wrap_cudadrv.cpp` is returning an invalid `HostPointer` from `register_host_memory`? It's the only I thing I can think of. It's unlikely because within the test case, I've got a `assert np.all(a == a_pin)`. \r\n\r\nI am aware of these use cases where the pinned host memory is mapped into the device's address space, and the address of this space is retrieved with `get_device_pointer()`. See this snippet from @fasiha's (failing) [gist](https://gist.github.com/fasiha/4036693) for e.g.\r\n\r\n```python\r\n# Time use of pinned host memory:\r\nx = drv.aligned_empty((N, N), dtype=np.float64, order='C')\r\nx = drv.register_host_memory(x, flags=drv.mem_host_register_flags.DEVICEMAP)\r\nx_gpu_ptr = np.intp(x.base.get_device_pointer())\r\n```\r\n\r\nBut I don't think this is the default behaviour. As illustrated by the CUDA C version, it should be possible to just pin the host memory without performing a corresponding mapped allocation on the device, and then just use the host pointer as a call to the asynchronous memory copy.\r\n\r\n\r\n"}}