     uint numBlocks, numThreads;
     computeGridSize(updatePortion.y - updatePortion.x, 256, numBlocks, numThreads);
     UpdateFluidD<<<numBlocks, numThreads>>>(
-        mR4CAST(sphMarkersD->posRadD), mR3CAST(sphMarkersD->velMasD), mR4CAST(sphMarkersD->rhoPresMuD),
-        mR3CAST(sphMarkersD->tauXxYyZzD), mR3CAST(sphMarkersD->tauXyXzYzD),
-        mR3CAST(fsiSystem.fsiGeneralData->vel_XSPH_D), mR4CAST(fsiSystem.fsiGeneralData->derivVelRhoD_old),
-        mR3CAST(fsiSystem.fsiGeneralData->derivTauXxYyZzD), mR3CAST(fsiSystem.fsiGeneralData->derivTauXyXzYzD),
-        mR4CAST(fsiSystem.fsiGeneralData->sr_tau_I_mu_i), U1CAST(fsiSystem.fsiGeneralData->activityIdentifierD),
-        U1CAST(fsiSystem.fsiGeneralData->freeSurfaceIdD), updatePortion, dT, isErrorD);
+        mR4CAST(sphMarkersD->posRadD), 
+        mR3CAST(sphMarkersD->velMasD), 
+        mR4CAST(sphMarkersD->rhoPresMuD), 
+        mR3CAST(sphMarkersD->tauXxYyZzD), 
+        mR3CAST(sphMarkersD->tauXyXzYzD), 
+        mR3CAST(fsiSystem.fsiGeneralData->vel_XSPH_D), 
+        mR4CAST(fsiSystem.fsiGeneralData->derivVelRhoD_old),
+        mR3CAST(fsiSystem.fsiGeneralData->derivTauXxYyZzD), 
+        mR3CAST(fsiSystem.fsiGeneralData->derivTauXyXzYzD),
+        mR4CAST(fsiSystem.fsiGeneralData->sr_tau_I_mu_i), 
+        U1CAST(fsiSystem.fsiGeneralData->activityIdentifierD),
+        U1CAST(fsiSystem.fsiGeneralData->freeSurfaceIdD), 
+        updatePortion, dT, isErrorD);
     cudaDeviceSynchronize();
     cudaCheckError();
     //------------------------
     cudaMemcpy(isErrorH, isErrorD, sizeof(bool), cudaMemcpyDeviceToHost);
-    if (*isErrorH == true) {
+    if (*isErrorH == true)
         throw std::runtime_error("Error! program crashed in UpdateFluidD!\n");
-    }
     cudaFree(isErrorD);
     free(isErrorH);
 }


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

    uint activity = activityIdentifierD[index];
    if (activity == 0)
        return;  // no need to do anything if it is not an active particle

    Real4 rhoPresMu = rhoPresMuD[index];
    if (fabs(rhoPresMu.w) < .1) {
        return;
    }  // no need to do anything if it is a boundary particle
    Real3 posRad = mR3(posRadD[index]);
    Real h = posRadD[index].w;


    *isErrorH = false;
    cudaMemcpy(isErrorD, isErrorH, sizeof(bool), cudaMemcpyHostToDevice);
    Update_Fluid_State<<<numBlocks, numThreads>>>(
        mR3CAST(fsiSystem.fsiGeneralData->vel_XSPH_D), 
        mR4CAST(sphMarkersD->posRadD), mR3CAST(sphMarkersD->velMasD), 
        mR4CAST(sphMarkersD->rhoPresMuD), updatePortion,
        numObjectsH->numAllMarkers, paramsH->dT, isErrorD);
    cudaDeviceSynchronize();
    cudaCheckError();

    cudaMemcpy(isErrorH, isErrorD, sizeof(bool), cudaMemcpyDeviceToHost);
    if (*isErrorH == true)
        throw std::runtime_error("Error! program crashed in Update_Fluid_State!\n");
    cudaFree(isErrorD);
    free(isErrorH);
}

// -----------------------------------------------------------------------------
// Apply periodic boundary conditions in x, y, and z directions
void ChFluidDynamics::ApplyBoundarySPH_Markers(std::shared_ptr<SphMarkerDataD> sphMarkersD) {
    uint numBlocks, numThreads;

    computeGridSize((int)numObjectsH->numAllMarkers, 256, numBlocks, numThreads);
    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();

