aboutsummaryrefslogtreecommitdiff
path: root/NvCloth/src
diff options
context:
space:
mode:
Diffstat (limited to 'NvCloth/src')
-rw-r--r--NvCloth/src/ClothBase.h1
-rw-r--r--NvCloth/src/ClothImpl.h26
-rw-r--r--NvCloth/src/IterationState.h4
-rw-r--r--NvCloth/src/SwClothData.cpp1
-rw-r--r--NvCloth/src/SwClothData.h1
-rw-r--r--NvCloth/src/SwSolverKernel.cpp25
-rw-r--r--NvCloth/src/cuda/CuClothData.cpp1
-rw-r--r--NvCloth/src/cuda/CuClothData.h1
-rw-r--r--NvCloth/src/cuda/CuSolverKernel.cu14
-rw-r--r--NvCloth/src/dx/DxClothData.cpp1
-rw-r--r--NvCloth/src/dx/DxClothData.h3
-rw-r--r--NvCloth/src/dx/DxSolverKernel.hlsl15
-rw-r--r--NvCloth/src/neon/SwCollisionHelpers.h4
-rw-r--r--NvCloth/src/scalar/SwCollisionHelpers.h4
14 files changed, 71 insertions, 30 deletions
diff --git a/NvCloth/src/ClothBase.h b/NvCloth/src/ClothBase.h
index 8d75a72..ec1ee40 100644
--- a/NvCloth/src/ClothBase.h
+++ b/NvCloth/src/ClothBase.h
@@ -74,6 +74,7 @@ void initialize(Cloth& cloth, const physx::PxVec4* pIt, const physx::PxVec4* pEn
cloth.mWind = physx::PxVec3(0.0f);
cloth.mDragLogCoefficient = 0.0f;
cloth.mLiftLogCoefficient = 0.0f;
+ cloth.mFluidDensity = 1.0f;
cloth.mEnableContinuousCollision = false;
cloth.mCollisionMassScale = 0.0f;
cloth.mFriction = 0.0f;
diff --git a/NvCloth/src/ClothImpl.h b/NvCloth/src/ClothImpl.h
index 0c2b362..1e8d9a1 100644
--- a/NvCloth/src/ClothImpl.h
+++ b/NvCloth/src/ClothImpl.h
@@ -157,6 +157,8 @@ class ClothImpl : public Cloth
virtual float getDragCoefficient() const;
virtual void setLiftCoefficient(float);
virtual float getLiftCoefficient() const;
+ virtual void setFluidDensity(float);
+ virtual float getFluidDensity() const;
virtual void setSelfCollisionDistance(float);
virtual float getSelfCollisionDistance() const;
@@ -216,6 +218,7 @@ public: //Fields shared between all cloth classes
physx::PxVec3 mWind;
float mDragLogCoefficient;
float mLiftLogCoefficient;
+ float mFluidDensity;
// sleeping
uint32_t mSleepTestInterval; // how often to test for movement
@@ -313,7 +316,8 @@ inline physx::PxVec3 ClothImpl<T>::getGravity() const
inline float safeLog2(float x)
{
- return x ? physx::shdfnd::log2(x) : -FLT_MAX_EXP;
+ NV_CLOTH_ASSERT(("safeLog2",x >= 0.0f));
+ return x > 0 ? physx::shdfnd::log2(x) : -FLT_MAX_EXP;
}
inline physx::PxVec3 safeLog2(const physx::PxVec3& v)
@@ -1214,11 +1218,30 @@ inline float ClothImpl<T>::getLiftCoefficient() const
}
template <typename T>
+inline void ClothImpl<T>::setFluidDensity(float fluidDensity)
+{
+ NV_CLOTH_ASSERT(fluidDensity < 0.f);
+ if (fluidDensity == mFluidDensity)
+ return;
+
+ mFluidDensity = fluidDensity;
+ getChildCloth()->notifyChanged();
+ wakeUp();
+}
+
+template <typename T>
+inline float ClothImpl<T>::getFluidDensity() const
+{
+ return mFluidDensity;
+}
+
+template <typename T>
inline uint32_t ClothImpl<T>::getNumSelfCollisionIndices() const
{
return uint32_t(getChildCloth()->mSelfCollisionIndices.size());
}
+
// Fixed 4505:local function has been removed
template <typename T>
inline void ClothImpl<T>::setRestPositions(Range<const physx::PxVec4> restPositions)
@@ -1255,6 +1278,7 @@ inline float ClothImpl<T>::getSelfCollisionDistance() const
template <typename T>
inline void ClothImpl<T>::setSelfCollisionStiffness(float stiffness)
{
+ NV_CLOTH_ASSERT(stiffness <= 1.0f);
float value = safeLog2(1 - stiffness);
if (value == getChildCloth()->mSelfCollisionLogStiffness)
return;
diff --git a/NvCloth/src/IterationState.h b/NvCloth/src/IterationState.h
index f199663..224e87e 100644
--- a/NvCloth/src/IterationState.h
+++ b/NvCloth/src/IterationState.h
@@ -137,7 +137,7 @@ struct IterationState
Simd4f mCurBias; // in local space
Simd4f mPrevBias; // in local space
- Simd4f mWind; // delta position per iteration
+ Simd4f mWind; // delta position per iteration (wind velocity * mIterDt)
Simd4f mPrevMatrix[3];
Simd4f mCurMatrix[3];
@@ -290,7 +290,7 @@ cloth::IterationState<Simd4f> cloth::IterationStateFactory::create(MyCloth const
result.mCurBias = transform(result.mRotationMatrix, curLinearInertia + bias) & maskXYZ;
result.mPrevBias = transform(result.mRotationMatrix, linearInertia - curLinearInertia) & maskXYZ;
- Simd4f wind = load(array(cloth.mWind)) * iterDt;
+ Simd4f wind = load(array(cloth.mWind)) * iterDt; // multiply with delta time here already so we don't have to do it inside the solver
result.mWind = transform(result.mRotationMatrix, translation - wind) & maskXYZ;
result.mIsTurning = mPrevAngularVelocity.magnitudeSquared() + cloth.mAngularVelocity.magnitudeSquared() > 0.0f;
diff --git a/NvCloth/src/SwClothData.cpp b/NvCloth/src/SwClothData.cpp
index eddd821..f102bde 100644
--- a/NvCloth/src/SwClothData.cpp
+++ b/NvCloth/src/SwClothData.cpp
@@ -80,6 +80,7 @@ cloth::SwClothData::SwClothData(SwCloth& cloth, const SwFabric& fabric)
mNumTriangles = uint32_t(fabric.mTriangles.size()) / 3;
mDragCoefficient = 1.0f - expf(stiffnessExponent * cloth.mDragLogCoefficient);
mLiftCoefficient = 1.0f - expf(stiffnessExponent * cloth.mLiftLogCoefficient);
+ mFluidDensity = cloth.mFluidDensity * 0.5f; //divide by 2 to so we don't have to compensate for double area from cross product in the solver
mStartMotionConstraints = cloth.mMotionConstraints.mStart.size() ? array(cloth.mMotionConstraints.mStart.front()) : 0;
mTargetMotionConstraints =
diff --git a/NvCloth/src/SwClothData.h b/NvCloth/src/SwClothData.h
index d2387b5..78a6f99 100644
--- a/NvCloth/src/SwClothData.h
+++ b/NvCloth/src/SwClothData.h
@@ -92,6 +92,7 @@ struct SwClothData
uint32_t mNumTriangles;
float mDragCoefficient;
float mLiftCoefficient;
+ float mFluidDensity;
// motion constraint data
const float* mStartMotionConstraints;
diff --git a/NvCloth/src/SwSolverKernel.cpp b/NvCloth/src/SwSolverKernel.cpp
index 52dfdaa..dec46d7 100644
--- a/NvCloth/src/SwSolverKernel.cpp
+++ b/NvCloth/src/SwSolverKernel.cpp
@@ -384,9 +384,13 @@ T4f calculateMaxDelta(const T4f* prevIt, const T4f* curIt, const T4f* curEnd)
template <bool IsTurning, typename T4f>
void applyWind(T4f* __restrict curIt, const T4f* __restrict prevIt, const uint16_t* __restrict tIt,
- const uint16_t* __restrict tEnd, T4f dragCoefficient, T4f liftCoefficient, T4f wind,
+ const uint16_t* __restrict tEnd, float itrDtf, float dragCoefficientf, float liftCoefficientf, float fluidDensityf, T4f wind,
const T4f (&rotation)[3])
{
+ const T4f dragCoefficient = simd4f(dragCoefficientf);
+ const T4f liftCoefficient = simd4f(liftCoefficientf);
+ const T4f fluidDensity = simd4f(fluidDensityf);
+ const T4f itrDt = simd4f(itrDtf);
const T4f oneThird = simd4f(1.0f / 3.0f);
for (; tIt < tEnd; tIt += 3)
@@ -410,7 +414,7 @@ void applyWind(T4f* __restrict curIt, const T4f* __restrict prevIt, const uint16
T4f previous = oneThird * (p0 + p1 + p2);
//offset of the triangle center, including wind
- T4f delta = current - previous + wind;
+ T4f delta = current - previous + wind; //wind is also already multiplied by dt in the iteration state so everything it in the same units
if (IsTurning)
{
@@ -423,23 +427,25 @@ void applyWind(T4f* __restrict curIt, const T4f* __restrict prevIt, const uint16
T4f normal = cross3(c2 - c0, c1 - c0);
T4f doubleArea = sqrt(dot3(normal, normal));
+ normal = normal / doubleArea;
T4f invSqrScale = dot3(delta, delta);
T4f isZero = invSqrScale < gSimd4fEpsilon;
T4f scale = rsqrt(invSqrScale);
+ T4f deltaLength = sqrt(invSqrScale);
//scale 'normalizes' delta, doubleArea normalized normal
- T4f cosTheta = dot3(normal, delta) * scale / doubleArea;
+ T4f cosTheta = dot3(normal, delta) * scale;
T4f sinTheta = sqrt(max(gSimd4fZero, gSimd4fOne - cosTheta * cosTheta));
// orthogonal to delta, in delta-normal plane, same length as delta
T4f liftDir = cross3(cross3(delta, normal), delta * scale);
// sin(theta) * cos(theta) = 0.5 * sin(2 * theta)
- T4f lift = liftCoefficient * cosTheta * sinTheta * liftDir;
- T4f drag = dragCoefficient * abs(cosTheta) * doubleArea * delta; //dragCoefficient should compensate for double area
+ T4f lift = liftCoefficient * cosTheta * sinTheta * liftDir * deltaLength / itrDt;
+ T4f drag = dragCoefficient * abs(cosTheta) * delta * deltaLength / itrDt;
- T4f impulse = (lift + drag) & ~isZero;
+ T4f impulse = (drag + lift) * fluidDensity * doubleArea & ~isZero; //fluidDensity compensates for double area
curIt[i0] = c0 - impulse * splat<3>(c0);
curIt[i1] = c1 - impulse * splat<3>(c1);
@@ -668,17 +674,14 @@ void cloth::SwSolverKernel<T4f>::applyWind()
const uint16_t* tIt = mClothData.mTriangles;
const uint16_t* tEnd = tIt + 3 * mClothData.mNumTriangles;
- T4f dragCoefficient = simd4f(mClothData.mDragCoefficient);
- T4f liftCoefficient = simd4f(mClothData.mLiftCoefficient);
-
if (mState.mIsTurning)
{
- ::applyWind<true>(curIt, prevIt, tIt, tEnd, dragCoefficient, liftCoefficient, mState.mWind,
+ ::applyWind<true>(curIt, prevIt, tIt, tEnd, mState.mIterDt, mClothData.mDragCoefficient, mClothData.mLiftCoefficient, mClothData.mFluidDensity, mState.mWind,
mState.mRotationMatrix);
}
else
{
- ::applyWind<false>(curIt, prevIt, tIt, tEnd, dragCoefficient, liftCoefficient, mState.mWind,
+ ::applyWind<false>(curIt, prevIt, tIt, tEnd, mState.mIterDt, mClothData.mDragCoefficient, mClothData.mLiftCoefficient, mClothData.mFluidDensity, mState.mWind,
mState.mRotationMatrix);
}
}
diff --git a/NvCloth/src/cuda/CuClothData.cpp b/NvCloth/src/cuda/CuClothData.cpp
index decfd2c..927997c 100644
--- a/NvCloth/src/cuda/CuClothData.cpp
+++ b/NvCloth/src/cuda/CuClothData.cpp
@@ -124,6 +124,7 @@ cloth::CuFrameData::CuFrameData(CuCloth& cloth, uint32_t numSharedPositions, con
stiffness = gSimd4fOne - exp2(logStiffness * stiffnessExponent);
mDragCoefficient = array(stiffness)[0];
mLiftCoefficient = array(stiffness)[1];
+ mFluidDensity = cloth.mFluidDensity * 0.5f; //divide by 2 to so we don't have to compensate for double area from cross product in the solver
for (int i = 0; i < 9; ++i)
mRotation[i] = array(state.mRotationMatrix[i / 3])[i % 3];
diff --git a/NvCloth/src/cuda/CuClothData.h b/NvCloth/src/cuda/CuClothData.h
index 0e4cda0..dd836fd 100644
--- a/NvCloth/src/cuda/CuClothData.h
+++ b/NvCloth/src/cuda/CuClothData.h
@@ -135,6 +135,7 @@ struct CuFrameData
// wind data
float mDragCoefficient;
float mLiftCoefficient;
+ float mFluidDensity;
float mRotation[9];
// motion constraint data
diff --git a/NvCloth/src/cuda/CuSolverKernel.cu b/NvCloth/src/cuda/CuSolverKernel.cu
index 3517193..edb66dc 100644
--- a/NvCloth/src/cuda/CuSolverKernel.cu
+++ b/NvCloth/src/cuda/CuSolverKernel.cu
@@ -867,6 +867,8 @@ __device__ void applyWind(CurrentT& current, PreviousT& previous)
{
const float dragCoefficient = gFrameData.mDragCoefficient;
const float liftCoefficient = gFrameData.mLiftCoefficient;
+ const float fluidDensity = gFrameData.mFluidDensity;
+ const float itrDt = gFrameData.mIterDt;
if (dragCoefficient == 0.0f && liftCoefficient == 0.0f)
return;
@@ -912,20 +914,22 @@ __device__ void applyWind(CurrentT& current, PreviousT& previous)
float3 normal = cross3(c2 - c0, c1 - c0);
- float doubleArea = sqrtf(dot3(normal, normal));
+ const float doubleArea = sqrtf(dot3(normal, normal));
+ normal = (1.0f / doubleArea) * normal;
float invSqrScale = dot3(delta, delta);
float scale = rsqrtf(invSqrScale);
+ float deltaLength = sqrtf(invSqrScale);
- float cosTheta = dot3(normal, delta) * scale / doubleArea;
+ float cosTheta = dot3(normal, delta) * scale;
float sinTheta = sqrtf(max(0.0f, 1.0f - cosTheta * cosTheta));
float3 liftDir = cross3(cross3(delta, normal), scale * delta);
- float3 lift = liftCoefficient * cosTheta * sinTheta * liftDir;
- float3 drag = dragCoefficient * abs(cosTheta) * doubleArea * delta;
+ float3 lift = liftCoefficient * cosTheta * sinTheta * ((deltaLength / itrDt) * liftDir);
+ float3 drag = dragCoefficient * abs(cosTheta) * ((deltaLength / itrDt) * delta);
- float3 impulse = invSqrScale < FLT_EPSILON ? make_float3(0.0f, 0.0f, 0.0f) : lift + drag;
+ float3 impulse = invSqrScale < FLT_EPSILON ? make_float3(0.0f, 0.0f, 0.0f) : fluidDensity * doubleArea * (lift + drag);
applyImpulse(current(i0), impulse);
applyImpulse(current(i1), impulse);
diff --git a/NvCloth/src/dx/DxClothData.cpp b/NvCloth/src/dx/DxClothData.cpp
index 075dc81..57049bf 100644
--- a/NvCloth/src/dx/DxClothData.cpp
+++ b/NvCloth/src/dx/DxClothData.cpp
@@ -112,6 +112,7 @@ cloth::DxFrameData::DxFrameData(DxCloth& cloth, uint32_t numSharedPositions, con
Simd4f stiffness = gSimd4fOne - exp2(logStiffness * stiffnessExponent);
mDragCoefficient = array(stiffness)[0];
mLiftCoefficient = array(stiffness)[1];
+ mFluidDensity = cloth.mFluidDensity * 0.5f; //divide by 2 to so we don't have to compensate for double area from cross product in the solver
for(int i = 0; i < 9; ++i)
mRotation[i] = array(state.mRotationMatrix[i / 3])[i % 3];
}
diff --git a/NvCloth/src/dx/DxClothData.h b/NvCloth/src/dx/DxClothData.h
index af02bc6..f91d37d 100644
--- a/NvCloth/src/dx/DxClothData.h
+++ b/NvCloth/src/dx/DxClothData.h
@@ -46,7 +46,7 @@ typedef unsigned int uint32_t;
typedef int int32_t;
#endif
-static const uint32_t MaxParticlesInSharedMem = 1972;
+static const uint32_t MaxParticlesInSharedMem = 1971;
struct DxPhaseConfig
@@ -167,6 +167,7 @@ struct DxFrameData
// wind data
float mDragCoefficient;
float mLiftCoefficient;
+ float mFluidDensity;
float mRotation[9];
// motion constraint data
diff --git a/NvCloth/src/dx/DxSolverKernel.hlsl b/NvCloth/src/dx/DxSolverKernel.hlsl
index 2ca42b3..4d91f4b 100644
--- a/NvCloth/src/dx/DxSolverKernel.hlsl
+++ b/NvCloth/src/dx/DxSolverKernel.hlsl
@@ -202,7 +202,8 @@ void applyWind(IParticles curParticles, IParticles prevParticles, uint32_t threa
{
const float dragCoefficient = gFrameData.mDragCoefficient;
const float liftCoefficient = gFrameData.mLiftCoefficient;
-
+ const float fluidDensity = gFrameData.mFluidDensity;
+ const float itrDt = gFrameData.mIterDt;
if(dragCoefficient == 0.0f && liftCoefficient == 0.0f)
return;
@@ -258,20 +259,22 @@ void applyWind(IParticles curParticles, IParticles prevParticles, uint32_t threa
float3 normal = cross(c2.xyz - c0.xyz, c1.xyz - c0.xyz);
- float doubleArea = sqrt(dot(normal, normal));
+ const float doubleArea = sqrt(dot(normal, normal));
+ normal = normal / doubleArea;
float invSqrScale = dot(delta, delta);
float scale = rsqrt(invSqrScale);
+ float deltaLength = sqrt(invSqrScale);
- float cosTheta = dot(normal, delta) * scale / doubleArea;
+ float cosTheta = dot(normal, delta) * scale;
float sinTheta = sqrt(max(0.0f, 1.0f - cosTheta * cosTheta));
float3 liftDir = cross(cross(delta, normal), scale * delta);
- float3 lift = liftCoefficient * cosTheta * sinTheta * liftDir;
- float3 drag = dragCoefficient * abs(cosTheta) * doubleArea * delta;
+ float3 lift = liftCoefficient * cosTheta * sinTheta * liftDir * deltaLength / itrDt;
+ float3 drag = dragCoefficient * abs(cosTheta) * delta * deltaLength / itrDt;
- float3 impulse = invSqrScale < 1.192092896e-07F ? float3(0.0f, 0.0f, 0.0f) : lift + drag;
+ float3 impulse = invSqrScale < 1.192092896e-07F ? float3(0.0f, 0.0f, 0.0f) : (lift + drag) * fluidDensity * doubleArea;
curParticles.atomicAdd(i0, -impulse * c0.w);
curParticles.atomicAdd(i1, -impulse * c1.w);
diff --git a/NvCloth/src/neon/SwCollisionHelpers.h b/NvCloth/src/neon/SwCollisionHelpers.h
index 0b9410b..acd8a2a 100644
--- a/NvCloth/src/neon/SwCollisionHelpers.h
+++ b/NvCloth/src/neon/SwCollisionHelpers.h
@@ -33,7 +33,7 @@
#include <arm_neon.h>
#endif
-namespace physx
+namespace nv
{
namespace cloth
{
@@ -84,4 +84,4 @@ Simd4i Gather<Simd4i>::operator()(const Simd4i* ptr) const
}
} // namespace cloth
-} // namespace physx
+} // namespace nv
diff --git a/NvCloth/src/scalar/SwCollisionHelpers.h b/NvCloth/src/scalar/SwCollisionHelpers.h
index 0d7321f..af21812 100644
--- a/NvCloth/src/scalar/SwCollisionHelpers.h
+++ b/NvCloth/src/scalar/SwCollisionHelpers.h
@@ -29,7 +29,7 @@
#pragma once
-namespace physx
+namespace nv
{
namespace cloth
{
@@ -89,4 +89,4 @@ Scalar4i Gather<Scalar4i>::operator()(const Scalar4i* ptr) const
}
} // namespace cloth
-} // namespace physx
+} // namespace nv
generated by cgit v1.2.3 (git 2.25.1) at 2026-03-12 05:57:10 +0000