I try to draw water with this algorithm.
However, I have two problems.
The first problem is with CUDA Graphics Interop. Here is the code:
__global__ void copyParticlesPos(float3* particlesD3D, float3* particlesPhysX, int stride)
{
int thread = blockIdx.x * blockDim.x + threadIdx.x;
char* pointerToParticles = (char*)particlesPhysX;
pointerToParticles += thread*stride;
#pragma unroll
for(int i = 0; i < 16; i++)
{
particlesD3D[thread + i] = *((float3*)pointerToParticles);
pointerToParticles += stride;
}
}
void PhysXSimulator::Update(cudaGraphicsResource* particles)
{
mScene->simulate(simulationTimeStep);
mScene->fetchResults(true);
float3 *mappedRes;
size_t mappedSize;
auto physxdata = mFluid->lockParticleFluidReadData(physx::PxDataAccessFlag::eDEVICE);
cudaError Err = cudaGraphicsMapResources(1, &particles);
if(Err != cudaSuccess)
{
WCHAR aaa[100];
swprintf(aaa, L"Error code: %d", Err);
//MessageBox(NULL, aaa, L"Error", MB_OK); //At second and next launches shows error code 4 or 33
}
cudaGraphicsResourceGetMappedPointer(reinterpret_cast<void**>(&mappedRes), &mappedSize, particles);
int numParticlesPerBlock = 16;
dim3 threadsPerBlock(4096 / numParticlesPerBlock);
dim3 numBlocks(mappedSize / 4096);
copyParticlesPos<<<numBlocks, threadsPerBlock>>>(mappedRes, (float3*)physxdata->positionBuffer.ptr(), physxdata->positionBuffer.stride());
cudaDeviceSynchronize();
cudaGraphicsUnmapResources(1, &particles);
physxdata->unlock();
}
When Update() computes at first time, everything is ok. But at the second and every next time, the function lockParticleFluidReadData() returns cudaErrorLaunchFailure in debug build and cudaErrorInvalidResourceHandle in release build. It looks like resource is not being unmapped properly.
The second problem is strange fluid behavior. Here is the fluid parameters:
mFluid->setRestOffset(0.45f);
mFluid->setContactOffset(0.9f);
mFluid->setMaxMotionDistance(1.0f);
mFluid->setRestitution(0.3f);
mFluid->setDynamicFriction(0.001f);
mFluid->setStaticFriction(0.0f);
mFluid->setDamping(0.0f);
mFluid->setExternalAcceleration(physx::PxVec3(0.0f, -9.8f, 0.0f));
mFluid->setParticleBaseFlag(physx::PxParticleBaseFlag::eENABLED, true);
mFluid->setParticleBaseFlag(physx::PxParticleBaseFlag::eGPU, true);
mFluid->setStiffness(45.0f);
mFluid->setViscosity(150.0f);
The fluid does not leak and pool is not being filled by water, the particles just flock to walls. Here is view in PVD(4096 particles).