 __constant__ int2 portionD;
 __constant__ int flagD;
 __constant__ int numRigidBodiesD;
+__constant__ int numFlexBodiesD;
 __constant__ int startRigidMarkersD;
 __constant__ int startFlexMarkersD;
 __constant__ int numRigid_SphMarkersD;

						float(tStep)/stepEnd,
						delT);

		ApplyBoundary(posRadD, rhoPresMuD, numAllMarkers, posRigidD, velMassRigidD, numRigidBodies);
		//************



	CUT_CHECK_ERROR("Kernel execution failed: UpdateKernelBoundary");
}
//--------------------------------------------------------------------------------------------------------------------------------
void ApplyBoundarySPH_Markers(
		thrust::device_vector<real3> & posRadD,
		thrust::device_vector<real4> & rhoPresMuD,
		int numAllMarkers) {
	uint nBlock_NumSpheres, nThreads_SphMarkers;
	computeGridSize(numAllMarkers, 256, nBlock_NumSpheres, nThreads_SphMarkers);
	ApplyPeriodicBoundaryXKernel<<<nBlock_NumSpheres, nThreads_SphMarkers>>>(R3CAST(posRadD), R4CAST(rhoPresMuD));
	cudaThreadSynchronize();
	CUT_CHECK_ERROR("Kernel execution failed: ApplyPeriodicBoundaryXKernel");
	// these are useful anyway for out of bound particles
	ApplyPeriodicBoundaryYKernel<<<nBlock_NumSpheres, nThreads_SphMarkers>>>(R3CAST(posRadD), R4CAST(rhoPresMuD));
	cudaThreadSynchronize();
	CUT_CHECK_ERROR("Kernel execution failed: ApplyPeriodicBoundaryXKernel");
	ApplyPeriodicBoundaryZKernel<<<nBlock_NumSpheres, nThreads_SphMarkers>>>(R3CAST(posRadD), R4CAST(rhoPresMuD));
	cudaThreadSynchronize();
	CUT_CHECK_ERROR("Kernel execution failed: ApplyPeriodicBoundaryXKernel");
}
//--------------------------------------------------------------------------------------------------------------------------------
void ApplyBoundaryRigid(
		thrust::device_vector<real3> & posRigidD,
		int numRigidBodies) {
	uint nBlock_NumRigids, nThreads_RigidBodies;
	computeGridSize(numRigidBodies, 128, nBlock_NumRigids, nThreads_RigidBodies);
	cudaMemcpyToSymbolAsync(numRigidBodiesD, &numRigidBodies, sizeof(numRigidBodies)); //can be defined outside of the kernel, and only once
	ApplyPeriodicBoundaryXKernel_RigidBodies<<<nBlock_NumRigids, nThreads_RigidBodies>>>(R3CAST(posRigidD));
	cudaThreadSynchronize();
	CUT_CHECK_ERROR("Kernel execution failed: UpdateKernelRigid");
	ApplyPeriodicBoundaryYKernel_RigidBodies<<<nBlock_NumRigids, nThreads_RigidBodies>>>(R3CAST(posRigidD));
	cudaThreadSynchronize();
	CUT_CHECK_ERROR("Kernel execution failed: UpdateKernelRigid");
	ApplyPeriodicBoundaryZKernel_RigidBodies<<<nBlock_NumRigids, nThreads_RigidBodies>>>(R3CAST(posRigidD));
	cudaThreadSynchronize();
	CUT_CHECK_ERROR("Kernel execution failed: UpdateKernelRigid");
}

//--------------------------------------------------------------------------------------------------------------------------------
void ApplyBoundaryFlex(
		thrust::device_vector<real3> & ANCF_NodesD,
		thrust::device_vector<int2> & ANCF_ReferenceArrayNodesOnBeamsD,
		int numFlex) {
	uint nBlock_NumFlex, nThreads_Flex;
	computeGridSize(numFlex, 128, nBlock_NumFlex, nThreads_Flex);
	cudaMemcpyToSymbolAsync(numFlexBodiesD, &numFlex, sizeof(numFlex)); //can be defined outside of the kernel, and only once
	ApplyPeriodicBoundaryKernel_FlexBodies<<<nBlock_NumFlex, nThreads_Flex>>>(R3CAST(ANCF_NodesD), I2CAST(ANCF_ReferenceArrayNodesOnBeamsD)); // x,y,z all implemented in a single kernel
	cudaThreadSynchronize();
}
//--------------------------------------------------------------------------------------------------------------------------------
void ApplyBoundary(
		thrust::device_vector<real3> & posRadD,
		thrust::device_vector<real4> & rhoPresMuD,
		int numAllMarkers,
		thrust::device_vector<real3> & posRigidD,
		int numRigidBodies,
		thrust::device_vector<real3> & ANCF_NodesD,
		thrust::device_vector<int2> & ANCF_ReferenceArrayNodesOnBeamsD,
		int numFlex) {

	ApplyBoundarySPH_Markers(posRadD, rhoPresMuD, numAllMarkers);
	ApplyBoundaryRigid(posRigidD, numRigidBodies);

}
//--------------------------------------------------------------------------------------------------------------------------------
void FindPassesFromTheEnd(
