{"action":"created","issue":{"url":"https://api.github.com/repos/ComputationalRadiationPhysics/picongpu/issues/858","labels_url":"https://api.github.com/repos/ComputationalRadiationPhysics/picongpu/issues/858/labels{/name}","comments_url":"https://api.github.com/repos/ComputationalRadiationPhysics/picongpu/issues/858/comments","events_url":"https://api.github.com/repos/ComputationalRadiationPhysics/picongpu/issues/858/events","html_url":"https://github.com/ComputationalRadiationPhysics/picongpu/issues/858","id":73998363,"number":858,"title":"AdiosWriter performance (follow up of mailing list post)","user":{"login":"anshumang","id":2358929,"avatar_url":"https://avatars.githubusercontent.com/u/2358929?v=3","gravatar_id":"","url":"https://api.github.com/users/anshumang","html_url":"https://github.com/anshumang","followers_url":"https://api.github.com/users/anshumang/followers","following_url":"https://api.github.com/users/anshumang/following{/other_user}","gists_url":"https://api.github.com/users/anshumang/gists{/gist_id}","starred_url":"https://api.github.com/users/anshumang/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/anshumang/subscriptions","organizations_url":"https://api.github.com/users/anshumang/orgs","repos_url":"https://api.github.com/users/anshumang/repos","events_url":"https://api.github.com/users/anshumang/events{/privacy}","received_events_url":"https://api.github.com/users/anshumang/received_events","type":"User","site_admin":false},"labels":[{"url":"https://api.github.com/repos/ComputationalRadiationPhysics/picongpu/labels/question","name":"question","color":"cc317c"}],"state":"closed","locked":false,"assignee":{"login":"ax3l","id":1353258,"avatar_url":"https://avatars.githubusercontent.com/u/1353258?v=3","gravatar_id":"","url":"https://api.github.com/users/ax3l","html_url":"https://github.com/ax3l","followers_url":"https://api.github.com/users/ax3l/followers","following_url":"https://api.github.com/users/ax3l/following{/other_user}","gists_url":"https://api.github.com/users/ax3l/gists{/gist_id}","starred_url":"https://api.github.com/users/ax3l/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/ax3l/subscriptions","organizations_url":"https://api.github.com/users/ax3l/orgs","repos_url":"https://api.github.com/users/ax3l/repos","events_url":"https://api.github.com/users/ax3l/events{/privacy}","received_events_url":"https://api.github.com/users/ax3l/received_events","type":"User","site_admin":false},"milestone":{"url":"https://api.github.com/repos/ComputationalRadiationPhysics/picongpu/milestones/2","html_url":"https://github.com/ComputationalRadiationPhysics/picongpu/milestones/Open%20Beta","labels_url":"https://api.github.com/repos/ComputationalRadiationPhysics/picongpu/milestones/2/labels","id":395839,"number":2,"title":"Open Beta","description":"Tasks to complete until the release of our open beta (version 0.1)\r\n\r\nThe [beta](https://en.wikipedia.org/wiki/Software_release_life_cycle#Beta) release shall be useful for *physicist* that used a Particle-in-Cell code *before*. Experience with *compiling libraries from source* and a *good Linux knowledge* is required for the initial installation.\r\n","creator":{"login":"ax3l","id":1353258,"avatar_url":"https://avatars.githubusercontent.com/u/1353258?v=3","gravatar_id":"","url":"https://api.github.com/users/ax3l","html_url":"https://github.com/ax3l","followers_url":"https://api.github.com/users/ax3l/followers","following_url":"https://api.github.com/users/ax3l/following{/other_user}","gists_url":"https://api.github.com/users/ax3l/gists{/gist_id}","starred_url":"https://api.github.com/users/ax3l/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/ax3l/subscriptions","organizations_url":"https://api.github.com/users/ax3l/orgs","repos_url":"https://api.github.com/users/ax3l/repos","events_url":"https://api.github.com/users/ax3l/events{/privacy}","received_events_url":"https://api.github.com/users/ax3l/received_events","type":"User","site_admin":false},"open_issues":88,"closed_issues":686,"state":"open","created_at":"2013-08-06T16:38:57Z","updated_at":"2015-05-07T14:07:37Z","due_on":null,"closed_at":null},"comments":5,"created_at":"2015-05-07T14:02:04Z","updated_at":"2015-05-07T20:21:21Z","closed_at":"2015-05-07T14:07:37Z","body":"@ax3l @psychocoderHPC \r\nI was wondering if any of you had noticed my post on the picongpu-users mailing list. If not, here's the post (copied) -\r\n\r\nI ran some measurements on the ADIOSWriter plugin and wanted to check if there are some reference numbers to validate against. I could only run it on a M2090 so numbers might not agree but still wanted to get a ballpark comparison.\r\n\r\nExperiment description:\r\n* -g 128 128 128\r\n* -d 1 1 1\r\n* Single node\r\n* KelvinHelmholtzInstability\r\n\r\nPerformance data:\r\n* Avg simulation timestep : 2.1sec\r\n* ADIOSWriter : 338sec\r\n    * Field : 1.8sec\r\n    * Species1 : 165sec\r\n        * kernel 'copySpecies' : 163.6sec\r\n    * Species2 : 165.1sec\r\n        * kernel 'copySpecies' : 163.6sec\r\n\r\nQuestions:\r\n* Why is the destination buffer (deviceFrame) of 'copySpecies' alloc'd on host pinned memory and not on device memory?\r\n* Does the 163sec of execution time of 'copySpecies' for the chosen simulation size look reasonable even for an M2090?\r\n*  If the source buffer (speciesTmp->getDeviceParticlesBox()) is copied to host memory and a CPU version of 'copySpecies' is run instead, would it be same semantically?\r\n-----> To do the above, I measured the following - \r\n           speciesTmp->synchronize();\r\n           cudaDeviceSynchronize();\r\n        It comes to -\r\n             Species1 : 369 ms\r\n             Species2 : 443 ms"},"comment":{"url":"https://api.github.com/repos/ComputationalRadiationPhysics/picongpu/issues/comments/100005812","html_url":"https://github.com/ComputationalRadiationPhysics/picongpu/issues/858#issuecomment-100005812","issue_url":"https://api.github.com/repos/ComputationalRadiationPhysics/picongpu/issues/858","id":100005812,"user":{"login":"anshumang","id":2358929,"avatar_url":"https://avatars.githubusercontent.com/u/2358929?v=3","gravatar_id":"","url":"https://api.github.com/users/anshumang","html_url":"https://github.com/anshumang","followers_url":"https://api.github.com/users/anshumang/followers","following_url":"https://api.github.com/users/anshumang/following{/other_user}","gists_url":"https://api.github.com/users/anshumang/gists{/gist_id}","starred_url":"https://api.github.com/users/anshumang/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/anshumang/subscriptions","organizations_url":"https://api.github.com/users/anshumang/orgs","repos_url":"https://api.github.com/users/anshumang/repos","events_url":"https://api.github.com/users/anshumang/events{/privacy}","received_events_url":"https://api.github.com/users/anshumang/received_events","type":"User","site_admin":false},"created_at":"2015-05-07T20:21:21Z","updated_at":"2015-05-07T20:21:21Z","body":"Couldn't find a reply button on the mailing list viewer....so replying here with only the context ...\r\n\r\nAH : the \"copySpecies\" kernel basically *is* a deep-copy.\r\nAG : but to run it, the simulation has to be blocked, no?...have you considered, \r\n(1) deep copying all the data using the DMA engines from GPU memory to host memory \r\n(2) then run the strided access kernel on the CPU to populate a contiguous target buffer for adios_write\r\n?\r\n\r\n>   - what version of PIConGPU are you running\r\n>     (beta-rc6, release-0.1.0, latest dev?)\r\nMaster, last commit : f0051b7 Hotfix: Race Condition in MW Solver & Pusher #604 on Jan 5, 2015\r\n>   - what system/compilers/third-party libraries are you using?\r\nOS : RHEL 6.6\r\ng++ : 4.4.7\r\nnvcc : 6.5\r\nadios : 1.8.0\r\n\r\n> \r\n> In the current development (dev) version of PIConGPU, we do not use\r\n> double-buffering for particles any more. For that, we nowadays use\r\n> *mapped* memory for access from the host side, e.g., for dumps.\r\nDoes this mean the host side can access parts of the species data before the copySpecies kernel completes?\r\n \r\n> \r\n> How many time steps did you run?\r\n10\r\n> How often did you dump (adios.period)?\r\n1\r\n> How many particles did you use in total?\r\n~100 macro particles/gpu, 1 GPU, 1 node (-d 1 1 1)\r\n> \r\n> Also, the M2090 (\"Fermi\") generation is \"pretty old\" meaning it might be\r\n> possible that the performance of the Fermi implementation for mapped\r\n> memory is not comparable any more to modern hardware (\"Kepler\" is out\r\n> since Nov/2012). Please use modern hardware for benchmarks, we won't\r\n> optimize for that generation (even if we still support general operation\r\n> on it).\r\n> \r\n> We did not benchmark `copySpecies` in detail yet, but since we only use\r\n> it on Kepler the \"163.6sec\" sound a bit unrealistic (or way off). Can\r\n> you compare it to your Kepler cards in your node, pls?\r\n> \r\nWill get back with this shortly.\r\n> \r\n> The\r\n>   speciesTmp->synchronize();\r\n> is without implementation - it would usually synchronize data from\r\n> device to host but is not implemented for particles (due to their nested\r\n> memory structure, see above).\r\n> \r\n> So basically what you measured is the time of a\r\n>   cudaDeviceSynchronize();\r\n> and *all the kernels that were still running* before that, which can be\r\n> anything.\r\n> \r\n> Also: before you measure kernels, always run a\r\n>   cudaDeviceSynchronize();\r\n> else you get the unpredictable load of asynchronous kernels that you\r\n> don't intended to measure ;)\r\nHmmm...interesting...I do have a leading and trailing cudaDeviceSynchronize() like the following -\r\n            cudaDeviceSynchronize();\r\n            speciesTmp->synchronize();\r\n            cudaDeviceSynchronize();\r\nI am surprised that the duration for the above is ~400 ms even though the actual cudaMemcpy does not occur. Any ideas why this is happening?"}}