{"action":"created","issue":{"url":"https://api.github.com/repos/torch/cutorch/issues/128","labels_url":"https://api.github.com/repos/torch/cutorch/issues/128/labels{/name}","comments_url":"https://api.github.com/repos/torch/cutorch/issues/128/comments","events_url":"https://api.github.com/repos/torch/cutorch/issues/128/events","html_url":"https://github.com/torch/cutorch/pull/128","id":66783440,"number":128,"title":"GPU multi-stream support in cutorch","user":{"login":"soumith","id":1310570,"avatar_url":"https://avatars.githubusercontent.com/u/1310570?v=3","gravatar_id":"","url":"https://api.github.com/users/soumith","html_url":"https://github.com/soumith","followers_url":"https://api.github.com/users/soumith/followers","following_url":"https://api.github.com/users/soumith/following{/other_user}","gists_url":"https://api.github.com/users/soumith/gists{/gist_id}","starred_url":"https://api.github.com/users/soumith/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/soumith/subscriptions","organizations_url":"https://api.github.com/users/soumith/orgs","repos_url":"https://api.github.com/users/soumith/repos","events_url":"https://api.github.com/users/soumith/events{/privacy}","received_events_url":"https://api.github.com/users/soumith/received_events","type":"User","site_admin":false},"labels":[],"state":"closed","locked":false,"assignee":null,"milestone":null,"comments":14,"created_at":"2015-04-07T05:50:08Z","updated_at":"2015-04-10T16:52:04Z","closed_at":"2015-04-07T16:20:01Z","pull_request":{"url":"https://api.github.com/repos/torch/cutorch/pulls/128","html_url":"https://github.com/torch/cutorch/pull/128","diff_url":"https://github.com/torch/cutorch/pull/128.diff","patch_url":"https://github.com/torch/cutorch/pull/128.patch"},"body":"This implements a version of the CUDA stream/event API in cutorch.\r\n\r\nOriginal Author: @wickedfoo\r\n\r\nBy default streams are not allocated, but reserveStreams can be used to create a number of streams that are available on every device. This API does not expose explicit stream or event creation/destruction due to complexity, but it gives all the functionality that one would need via allowing one to construct single- or multi-device one-way and two-way barriers at the lua level\r\n\r\n###Functions:\r\n- `cutorch.reserveStreams(n)`: creates n user streams for use on every device.\r\n- `n = cutorch.getNumStreams()`: returns the number of user streams available on every device. By `default`, this is `0`, meaning only the default stream (stream 0) is available.\r\n- `cutorch.setStream(n)`: specifies that the current stream active for the current device (or any other device) is `n`. This is preserved across device switches. `1-N` are user streams, `0` is the default stream.\r\n- `n = cutorch.getStream()`: returns the current stream active. By default, returns `0`.\r\n- `cutorch.setDefaultStream()`: an alias for `cutorch.setStream(0)`\r\n- `cutorch.streamWaitFor(streamWaiting, {streamsToWaitOn...})`: A `1-to-N-way` barrier. `streamWaiting` will wait for the list of streams specified to finish executing all kernels/events/barriers. Does not block any of the streamsToWaitOn. Current device only.\r\n- `cutorch.streamWaitForMultiDevice(deviceWaiting, streamWaiting, {[device]={streamsToWaitOn...}...})`: (`deviceWaiting`, `streamWaiting`) will wait on the list of (`device`, `streams`...) pairs; handles single or multiple device. `cutorch.streamWaitForMultiDevice, a, b, {[a]={streams...}})` is equivalent to `cutorch.setDevice(a); cutorch.streamWaitFor(b, {streams...})`.\r\n- `cutorch.streamBarrier({streams...})`: an `N-to-N-way` barrier between all the streams; all streams will wait for the completion of all other streams on the current device only. More efficient than creating the same N-to-N-way dependency via `streamWaitFor`.\r\n- `cutorch.streamBarrierMultiDevice({[device]={streamsToWaitOn...}...})`: As with streamBarrier but allows barriers between streams on arbitrary devices. Creates a cross-device N-to-N-way barrier between all (device, stream) values listed.\r\n- `cutorch.streamSynchronize(stream)`: equivalent to `cudaStreamSynchronize(stream)` for the current device. Blocks the CPU until stream completes its queued kernels/events.\r\n\r\n###Tests:\r\n- Tests added to cutorch for streamWaitFor(), streamWaitForMultiDevice(), streamBarrier() and streamBarrierMultiDevice().\r\n- Small fixes for deviceReset\r\n\r\n###Additional:\r\n- All cutorch has been converted to be adhering to this new streams API with the exception of thrust calls. That is also coming soon.\r\n- cunn calls are already converted as well, will sync that soon.\r\n- Multi-device checks have been added to ALL cutorch Tensor functions. If two tensors are on different devices, barring the copy operation, all other operations will assertively fail. cunn checks are also done and will be synced soon.\r\n\r\nAfter cunn fixes go in as well, containers like nn.Concat and nn.Parallel can be made truly parallel-ly executable.\r\nFixes for the cudnn package are also ready and will be synced soon.\r\n\r\ncc: @dominikgrewe @clementfarabet "},"comment":{"url":"https://api.github.com/repos/torch/cutorch/issues/comments/91619469","html_url":"https://github.com/torch/cutorch/pull/128#issuecomment-91619469","issue_url":"https://api.github.com/repos/torch/cutorch/issues/128","id":91619469,"user":{"login":"dominikgrewe","id":1619348,"avatar_url":"https://avatars.githubusercontent.com/u/1619348?v=3","gravatar_id":"","url":"https://api.github.com/users/dominikgrewe","html_url":"https://github.com/dominikgrewe","followers_url":"https://api.github.com/users/dominikgrewe/followers","following_url":"https://api.github.com/users/dominikgrewe/following{/other_user}","gists_url":"https://api.github.com/users/dominikgrewe/gists{/gist_id}","starred_url":"https://api.github.com/users/dominikgrewe/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/dominikgrewe/subscriptions","organizations_url":"https://api.github.com/users/dominikgrewe/orgs","repos_url":"https://api.github.com/users/dominikgrewe/repos","events_url":"https://api.github.com/users/dominikgrewe/events{/privacy}","received_events_url":"https://api.github.com/users/dominikgrewe/received_events","type":"User","site_admin":false},"created_at":"2015-04-10T16:52:03Z","updated_at":"2015-04-10T16:52:03Z","body":"Agreed, we should always enable p2p access if available. But we should also make sure that our code works if we don't have p2p access. I'll file a bug for the non-contiguous copy."}}