{"action":"created","issue":{"url":"https://api.github.com/repos/LLNL/RAJA/issues/91","repository_url":"https://api.github.com/repos/LLNL/RAJA","labels_url":"https://api.github.com/repos/LLNL/RAJA/issues/91/labels{/name}","comments_url":"https://api.github.com/repos/LLNL/RAJA/issues/91/comments","events_url":"https://api.github.com/repos/LLNL/RAJA/issues/91/events","html_url":"https://github.com/LLNL/RAJA/pull/91","id":164607578,"number":91,"title":"[WIP] CUDA Reductions split into 3 phases to support boxsegment","user":{"login":"jonesholger","id":12015127,"avatar_url":"https://avatars.githubusercontent.com/u/12015127?v=3","gravatar_id":"","url":"https://api.github.com/users/jonesholger","html_url":"https://github.com/jonesholger","followers_url":"https://api.github.com/users/jonesholger/followers","following_url":"https://api.github.com/users/jonesholger/following{/other_user}","gists_url":"https://api.github.com/users/jonesholger/gists{/gist_id}","starred_url":"https://api.github.com/users/jonesholger/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/jonesholger/subscriptions","organizations_url":"https://api.github.com/users/jonesholger/orgs","repos_url":"https://api.github.com/users/jonesholger/repos","events_url":"https://api.github.com/users/jonesholger/events{/privacy}","received_events_url":"https://api.github.com/users/jonesholger/received_events","type":"User","site_admin":false},"labels":[],"state":"open","locked":false,"assignee":null,"assignees":[],"milestone":null,"comments":5,"created_at":"2016-07-08T20:13:35Z","updated_at":"2016-08-18T22:40:47Z","closed_at":null,"pull_request":{"url":"https://api.github.com/repos/LLNL/RAJA/pulls/91","html_url":"https://github.com/LLNL/RAJA/pull/91","diff_url":"https://github.com/LLNL/RAJA/pull/91.diff","patch_url":"https://github.com/LLNL/RAJA/pull/91.patch"},"body":"Since boxsegment has an arbitrary number of \"participating\" threads it is necessary to split the processing into phases where we can control the number of threads used in each phase. The code in this PR has equivalent performance to the unsplit case (verified by a 90 second timing test)\r\n\r\nInitialization:  all threads\r\nLoad data: some threads\r\nFolding: all threads\r\n\r\nThe technique borrows from @trws  PR #70 -- thanks Tom \r\n\r\nInitialization utilizes the device side copy constructor, and to make sure this is activate at the right time we needed to modify the kernels \r\n\r\nThe reduction operators just load shared memory\r\n\r\nThe reduction dtors do the actual folding\r\n\r\n```\r\ntemplate <typename LOOP_BODY>                                                                                                                                                                                                    \r\n__global__ void forall_cuda_kernel(LOOP_BODY loop_body,                                                                                                                                                                          \r\n                                   Index_type begin,                                                                                                                                                                             \r\n                                   Index_type len) {                                                                                                                                                                             \r\n  auto body = loop_body;         <--- do this before any extent logic to setup copy constructors                                                                                                                                                                                                \r\n  Index_type ii = blockDim.x * blockIdx.x + threadIdx.x;                                                                                                                                                                         \r\n  if (ii < len) {                                                                                                                                                                                                                \r\n    body(begin + ii);                                                                                                                                                                                                            \r\n  }                                                                                                                                                                                                                              \r\n}       \r\n```\r\n\r\nAnother modification is the use of \"dynamic\" shared memory, so addressing equivalent shared memory offsets can be done across the three device routines (phases)\r\n\r\nI'm thinking we can perform a dynamic shared memory allocation by keeping track of reduction ID's in scope with their type sizes and BLOCK_SIZE. Right now it's one chunk just to demonstrate feasibility. \r\n\r\nOne last thing was to modify the operator call so that it returns a reference; avoiding an extra copy. \r\n\r\n"},"comment":{"url":"https://api.github.com/repos/LLNL/RAJA/issues/comments/240878797","html_url":"https://github.com/LLNL/RAJA/pull/91#issuecomment-240878797","issue_url":"https://api.github.com/repos/LLNL/RAJA/issues/91","id":240878797,"user":{"login":"jonesholger","id":12015127,"avatar_url":"https://avatars.githubusercontent.com/u/12015127?v=3","gravatar_id":"","url":"https://api.github.com/users/jonesholger","html_url":"https://github.com/jonesholger","followers_url":"https://api.github.com/users/jonesholger/followers","following_url":"https://api.github.com/users/jonesholger/following{/other_user}","gists_url":"https://api.github.com/users/jonesholger/gists{/gist_id}","starred_url":"https://api.github.com/users/jonesholger/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/jonesholger/subscriptions","organizations_url":"https://api.github.com/users/jonesholger/orgs","repos_url":"https://api.github.com/users/jonesholger/repos","events_url":"https://api.github.com/users/jonesholger/events{/privacy}","received_events_url":"https://api.github.com/users/jonesholger/received_events","type":"User","site_admin":false},"created_at":"2016-08-18T22:40:47Z","updated_at":"2016-08-18T22:40:47Z","body":"@trws I like your approach especially wrt simplifying copy ctor, and moving \"all\" of the work into dtor, and removing the need for dynamic shared memory - very nice.  I plan to explore this further when I get back from vacation.  I didn't think the class variable m_reduced_val would be fast to accumulate into in the operator += on a thread basis. Also, thinking about other use cases where you could call += multiple times for the same index (thread) in a forall, and finish with dtor."}}
{"actor":{"display_login":"jonesholger"}}