     uint numBlocks, numThreads;
 
     computeGridSize((int)numObjectsH->numAllMarkers, 256, numBlocks, numThreads);
-    ApplyPeriodicBoundaryXKernel<<<numBlocks, numThreads>>>
-        (mR4CAST(sphMarkersD->posRadD), mR4CAST(sphMarkersD->rhoPresMuD));
-    cudaDeviceSynchronize();
-    cudaCheckError();
-
-    ApplyPeriodicBoundaryYKernel<<<numBlocks, numThreads>>>
-        (mR4CAST(sphMarkersD->posRadD), mR4CAST(sphMarkersD->rhoPresMuD));
-    cudaDeviceSynchronize();
-    cudaCheckError();
-
-    ApplyPeriodicBoundaryZKernel<<<numBlocks, numThreads>>>
-        (mR4CAST(sphMarkersD->posRadD), mR4CAST(sphMarkersD->rhoPresMuD));
+    ApplyPeriodicBoundaryXKernel<<<numBlocks, numThreads>>>(
+        mR4CAST(sphMarkersD->posRadD), mR4CAST(sphMarkersD->rhoPresMuD),
+        U1CAST(fsiSystem->fsiGeneralData->activityIdentifierD));
+    cudaDeviceSynchronize();
+    cudaCheckError();
+
+    ApplyPeriodicBoundaryYKernel<<<numBlocks, numThreads>>>(
+        mR4CAST(sphMarkersD->posRadD), mR4CAST(sphMarkersD->rhoPresMuD),
+        U1CAST(fsiSystem->fsiGeneralData->activityIdentifierD));
+    cudaDeviceSynchronize();
+    cudaCheckError();
+
+    ApplyPeriodicBoundaryZKernel<<<numBlocks, numThreads>>>(
+        mR4CAST(sphMarkersD->posRadD), mR4CAST(sphMarkersD->rhoPresMuD),
+        U1CAST(fsiSystem->fsiGeneralData->activityIdentifierD));
     cudaDeviceSynchronize();
     cudaCheckError();
 


// -----------------------------------------------------------------------------
// Kernel to apply periodic BC along z
__global__ void ApplyPeriodicBoundaryZKernel(Real4* posRadD, Real4* rhoPresMuD) {
    uint index = blockIdx.x * blockDim.x + threadIdx.x;
    if (index >= numObjectsD.numAllMarkers) {
        return;
    }
    Real4 rhoPresMu = rhoPresMuD[index];
    if (fabs(rhoPresMu.w) < .1) {
        return;


    // these are useful anyway for out of bound particles
    ApplyPeriodicBoundaryYKernel<<<numBlocks, numThreads>>>
        (mR4CAST(sphMarkersD->posRadD),mR4CAST(sphMarkersD->rhoPresMuD),
         U1CAST(fsiSystem->fsiGeneralData->activityIdentifierD));
    cudaDeviceSynchronize();
    cudaCheckError();

    ApplyPeriodicBoundaryZKernel<<<numBlocks, numThreads>>>
        (mR4CAST(sphMarkersD->posRadD), mR4CAST(sphMarkersD->rhoPresMuD),
         U1CAST(fsiSystem->fsiGeneralData->activityIdentifierD));
    cudaDeviceSynchronize();
    cudaCheckError();
}
