{"action":"created","issue":{"url":"https://api.github.com/repos/ComputationalRadiationPhysics/picongpu/issues/889","labels_url":"https://api.github.com/repos/ComputationalRadiationPhysics/picongpu/issues/889/labels{/name}","comments_url":"https://api.github.com/repos/ComputationalRadiationPhysics/picongpu/issues/889/comments","events_url":"https://api.github.com/repos/ComputationalRadiationPhysics/picongpu/issues/889/events","html_url":"https://github.com/ComputationalRadiationPhysics/picongpu/issues/889","id":81328071,"number":889,"title":"DirectionalSplitting not twitstet blockDim","user":{"login":"psychocoderHPC","id":4037125,"avatar_url":"https://avatars.githubusercontent.com/u/4037125?v=3","gravatar_id":"","url":"https://api.github.com/users/psychocoderHPC","html_url":"https://github.com/psychocoderHPC","followers_url":"https://api.github.com/users/psychocoderHPC/followers","following_url":"https://api.github.com/users/psychocoderHPC/following{/other_user}","gists_url":"https://api.github.com/users/psychocoderHPC/gists{/gist_id}","starred_url":"https://api.github.com/users/psychocoderHPC/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/psychocoderHPC/subscriptions","organizations_url":"https://api.github.com/users/psychocoderHPC/orgs","repos_url":"https://api.github.com/users/psychocoderHPC/repos","events_url":"https://api.github.com/users/psychocoderHPC/events{/privacy}","received_events_url":"https://api.github.com/users/psychocoderHPC/received_events","type":"User","site_admin":false},"labels":[{"url":"https://api.github.com/repos/ComputationalRadiationPhysics/picongpu/labels/question","name":"question","color":"cc317c"}],"state":"open","locked":false,"assignee":{"login":"Heikman","id":5159590,"avatar_url":"https://avatars.githubusercontent.com/u/5159590?v=3","gravatar_id":"","url":"https://api.github.com/users/Heikman","html_url":"https://github.com/Heikman","followers_url":"https://api.github.com/users/Heikman/followers","following_url":"https://api.github.com/users/Heikman/following{/other_user}","gists_url":"https://api.github.com/users/Heikman/gists{/gist_id}","starred_url":"https://api.github.com/users/Heikman/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/Heikman/subscriptions","organizations_url":"https://api.github.com/users/Heikman/orgs","repos_url":"https://api.github.com/users/Heikman/repos","events_url":"https://api.github.com/users/Heikman/events{/privacy}","received_events_url":"https://api.github.com/users/Heikman/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":89,"closed_issues":715,"state":"open","created_at":"2013-08-06T16:38:57Z","updated_at":"2015-05-27T07:44:17Z","due_on":null,"closed_at":null},"comments":1,"created_at":"2015-05-27T07:39:28Z","updated_at":"2015-05-27T07:47:04Z","closed_at":null,"body":"The current implementation of `Directional Splitting` twist the axes and run three kernel to calculate new field. The used algorithm `ForEachBlock` had one compile time vector with the Blocksize (equal to SuperCellSize).\r\nhttps://github.com/ComputationalRadiationPhysics/picongpu/blob/dev/src/picongpu/include/fields/MaxwellSolver/DirSplitting/DirSplitting.hpp#L78\r\n\r\nMy question is: Is this right that the `BlockDim` is not twistet? For me it looks like we get memory access in the GUARD were we only calculate the CORE+BORDER domain. Might be also out of memory access in some cases.\r\n\r\nIMO and my understanding of the implementation we must twist the `BlockDim` in the same way we does it with the `gridSize`. \r\nThe reason is that our default SuperCell is 8x8x4 cells. We twist the fields thus every direction (x/y/z) is in one run at in kernel direction x. The kernel traverse over x and the shared memory size is create out of `BlockDim` by extent the x direction with two cells. In the default SuperCell case 10x8x4. If in one of the three runs field direction `z` is mapped to `x` and the simulation size is `24x24x12` `12` is mapped to `x` and traversed with a block 8x8x4 threads.\r\n\r\n@Heikman this might be very important   for @ax3l "},"comment":{"url":"https://api.github.com/repos/ComputationalRadiationPhysics/picongpu/issues/comments/105804940","html_url":"https://github.com/ComputationalRadiationPhysics/picongpu/issues/889#issuecomment-105804940","issue_url":"https://api.github.com/repos/ComputationalRadiationPhysics/picongpu/issues/889","id":105804940,"user":{"login":"psychocoderHPC","id":4037125,"avatar_url":"https://avatars.githubusercontent.com/u/4037125?v=3","gravatar_id":"","url":"https://api.github.com/users/psychocoderHPC","html_url":"https://github.com/psychocoderHPC","followers_url":"https://api.github.com/users/psychocoderHPC/followers","following_url":"https://api.github.com/users/psychocoderHPC/following{/other_user}","gists_url":"https://api.github.com/users/psychocoderHPC/gists{/gist_id}","starred_url":"https://api.github.com/users/psychocoderHPC/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/psychocoderHPC/subscriptions","organizations_url":"https://api.github.com/users/psychocoderHPC/orgs","repos_url":"https://api.github.com/users/psychocoderHPC/repos","events_url":"https://api.github.com/users/psychocoderHPC/events{/privacy}","received_events_url":"https://api.github.com/users/psychocoderHPC/received_events","type":"User","site_admin":false},"created_at":"2015-05-27T07:47:03Z","updated_at":"2015-05-27T07:47:03Z","body":"I change this to a bug `cuda-memcheck` confirmed that we get out of memory access.\r\n\r\nDirectional Splitting without particles and laser plane wave:\r\n```\r\ncuda-memcheck --report-api-errors no  ./picongpu -d 1 1 1 -g 24 64 12 -s 1 --periodic 1 0 1\r\n```\r\n```BASH\r\n========= CUDA-MEMCHECK\r\n========= Invalid __global__ read of size 4\r\n=========     at 0x000006f0 in _ZN5PMacc9algorithm6kernel6detail18kernelForeachBlockINS2_13Sph\r\nericMapperILi3ENS_4math2CT6VectorIN4mpl_10integral_cIiLi8EEESA_NS9_IiLi4EEEEENS8_5void_EEENS_6\r\ncursor6CursorINSF_14MarkerAccessorINSG_INSF_17TwistAxesAccessorINSF_12BufferCursorINS5_6Vector\r\nIfLi3ENS5_16StandartAccessorENS5_17StandartNavigatorENS5_6detail17Vector_componentsEEELi3EEENS\r\n6_3IntILi2ELi0ELi1EEEEENSF_2CT18TwistAxesNavigatorISS_Li3EEESQ_EEEENSF_15CursorNavigatorESX_EE\r\nS10_N8picongpu18DirSplittingKernelISC_EEEEvT_T0_T1_T2_\r\n=========     by thread (7,3,2) in block (0,0,13)\r\n=========     Address 0x13045afabc is out of bounds\r\n=========     Saved host backtrace up to driver entry point at kernel launch time\r\n=========     Host Frame:/opt/pkg/devel/cuda/6.5/lib64/libcuda.so.1 (cuLaunchKernel + 0x2cd) [\r\n0x150cad]\r\n=========     Host Frame:/opt/pkg/devel/cuda/6.5/lib64/libcudart.so.6.5 [0xf5d8]\r\n=========     Host Frame:/opt/pkg/devel/cuda/6.5/lib64/libcudart.so.6.5 (cudaLaunch + 0x143) [\r\n0x36833]\r\n=========     Host Frame:./picongpu (_ZN5PMacc9algorithm6kernel12ForeachBlockINS_4math2CT6Vect\r\norIN4mpl_10integral_cIiLi8EEES8_NS7_IiLi4EEEEESA_EclINS_4zone11SphericZoneILi3EEENS_6cursor6Cu\r\nrsorINSG_14MarkerAccessorINSH_INSG_17TwistAxesAccessorINSG_12BufferCursorINS3_6VectorIfLi3ENS3\r\n_16StandartAccessorENS3_17StandartNavigatorENS3_6detail17Vector_componentsEEELi3EEENS4_3IntILi\r\n2ELi0ELi1EEEEENSG_2CT18TwistAxesNavigatorIST_Li3EEESR_EEEENSG_15CursorNavigatorESY_EES11_N8pic\r\nongpu18DirSplittingKernelISA_EEEEvRKT_T0_T1_RKT2_ + 0x2cc) [0xc47ec]\r\n=========     Host Frame:./picongpu (_ZNK8picongpu12dirSplitting12DirSplitting20update_beforeC\r\nurrentEj + 0xc09) [0xcb029]\r\n=========     Host Frame:./picongpu (_ZN8picongpu12MySimulation10runOneStepEj + 0xff) [0xcb73f\r\n]\r\n=========     Host Frame:./picongpu (_ZN5PMacc16SimulationHelperILj3EE15startSimulationEv + 0x\r\n1c0) [0xb0f90]\r\n=========     Host Frame:./picongpu (_ZN8picongpu17SimulationStarterINS_21InitialiserControlle\r\nrENS_16PluginControllerENS_12MySimulationEE5startEv + 0x90) [0xe3950]\r\n=========     Host Frame:./picongpu (main + 0x62) [0x777e2]\r\n=========     Host Frame:/lib/x86_64-linux-gnu/libc.so.6 (__libc_start_main + 0xf5) [0x21ec5]\r\n=========     Host Frame:./picongpu [0x77baf]\r\n=========\r\n========= Invalid __global__ read of size 4\r\n```"}}