aboutsummaryrefslogtreecommitdiff
path: root/NvCloth/src/cuda
diff options
context:
space:
mode:
authorMarijn Tamis <[email protected]>2018-05-03 18:22:48 +0200
committerMarijn Tamis <[email protected]>2018-05-03 18:22:48 +0200
commitca32c59a58d37c1822e185a2d5f3d0d3e8943593 (patch)
treeb06b9eec03f34344ef8fc31aa147b2714d3962ee /NvCloth/src/cuda
parentForced rename of platform folders in cmake dir. Git didn't pick this up before. (diff)
downloadnvcloth-ca32c59a58d37c1822e185a2d5f3d0d3e8943593.tar.xz
nvcloth-ca32c59a58d37c1822e185a2d5f3d0d3e8943593.zip
NvCloth 1.1.4 Release. (24070740)
Diffstat (limited to 'NvCloth/src/cuda')
-rw-r--r--NvCloth/src/cuda/CuCloth.h2
-rw-r--r--NvCloth/src/cuda/CuCollision.h7
-rw-r--r--NvCloth/src/cuda/CuFabric.cpp3
-rw-r--r--NvCloth/src/cuda/CuFactory.cpp7
-rw-r--r--NvCloth/src/cuda/CuSolverKernelBlob.h0
5 files changed, 12 insertions, 7 deletions
diff --git a/NvCloth/src/cuda/CuCloth.h b/NvCloth/src/cuda/CuCloth.h
index 8dc3082..f7d8268 100644
--- a/NvCloth/src/cuda/CuCloth.h
+++ b/NvCloth/src/cuda/CuCloth.h
@@ -131,7 +131,7 @@ class CuCloth : protected CuContextLock, public ClothImpl<CuCloth>
uint32_t getSharedMemorySize() const; // without particle data
// expects transformed configs, doesn't call notifyChanged()
- void setPhaseConfigInternal(Range<const PhaseConfig>);
+ void setPhaseConfigInternal(Range<const PhaseConfig> configs);
Range<physx::PxVec4> push(CuConstraints&);
void clear(CuConstraints&);
diff --git a/NvCloth/src/cuda/CuCollision.h b/NvCloth/src/cuda/CuCollision.h
index aeb2bda..f9b69f7 100644
--- a/NvCloth/src/cuda/CuCollision.h
+++ b/NvCloth/src/cuda/CuCollision.h
@@ -1160,9 +1160,10 @@ __device__ void apply(ParticleDataT& particleData, const int4& indices, const fl
posX.y += delta.y * weights.x; posY.y += delta.y * weights.y; posZ.y += delta.y * weights.z;
posX.z += delta.z * weights.x; posY.z += delta.z * weights.y; posZ.z += delta.z * weights.z;
- particleData(indices.x) = posX;
- particleData(indices.y) = posY;
- particleData(indices.z) = posZ;
+ //if(particle is not static) store new position
+ if(particleData(indices.x,3)!= 0) particleData(indices.x) = posX;
+ if(particleData(indices.y,3)!= 0) particleData(indices.y) = posY;
+ if(particleData(indices.z,3)!= 0) particleData(indices.z) = posZ;
}
#else
template <typename PointerT>
diff --git a/NvCloth/src/cuda/CuFabric.cpp b/NvCloth/src/cuda/CuFabric.cpp
index 957f912..6794fa5 100644
--- a/NvCloth/src/cuda/CuFabric.cpp
+++ b/NvCloth/src/cuda/CuFabric.cpp
@@ -165,7 +165,8 @@ uint32_t cloth::CuFabric::getNumStiffnessValues() const
uint32_t cloth::CuFabric::getNumSets() const
{
- return uint32_t(mSets.size() - 1);
+ return uint32_t(physx::PxMax(0, (int)mSets.size() - 1));
+
}
uint32_t cloth::CuFabric::getNumIndices() const
diff --git a/NvCloth/src/cuda/CuFactory.cpp b/NvCloth/src/cuda/CuFactory.cpp
index 397262c..df52dcd 100644
--- a/NvCloth/src/cuda/CuFactory.cpp
+++ b/NvCloth/src/cuda/CuFactory.cpp
@@ -97,6 +97,7 @@ void cloth::checkSuccessImpl(CUresult err, const char* file, const int line)
ADD_CASE(CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES);
ADD_CASE(CUDA_ERROR_LAUNCH_TIMEOUT);
ADD_CASE(CUDA_ERROR_LAUNCH_INCOMPATIBLE_TEXTURING);
+ ADD_CASE(CUDA_ERROR_ILLEGAL_ADDRESS);
default:
ADD_CASE(CUDA_ERROR_UNKNOWN);
#undef ADD_CASE
@@ -115,8 +116,9 @@ uint32_t getMaxThreadsPerBlock(CUcontext context)
CUdevice device;
checkSuccess(cuCtxGetDevice(&device));
- int major = 0;
+ int major = 0, minor = 0;
checkSuccess(cuDeviceGetAttribute(&major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, device));
+ checkSuccess(cuDeviceGetAttribute(&minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, device));
checkSuccess(cuCtxPopCurrent(nullptr));
@@ -149,7 +151,7 @@ cloth::Fabric* cloth::CuFactory::createFabric(uint32_t numParticles, Range<const
Range<const float> tetherLengths, Range<const uint32_t> triangles)
{
return NV_CLOTH_NEW(CuFabric)(*this, numParticles, phaseIndices, sets, restvalues, stiffnessValues, indices, anchors, tetherLengths, triangles,
- getNextFabricId());
+ getNextFabricId());
}
cloth::Cloth* cloth::CuFactory::createCloth(Range<const PxVec4> particles, Fabric& fabric)
@@ -210,6 +212,7 @@ void cloth::CuFactory::extractFabricData(const Fabric& fabric, Range<uint32_t> p
if (!sets.empty())
{
+ // need to skip copying the first element
NV_CLOTH_ASSERT(sets.size() == cuFabric.mSets.size() - 1);
const uint32_t* deviceSets = cuFabric.mSets.begin().get();
copyToHost(deviceSets + 1, deviceSets + cuFabric.mSets.size(), sets.begin());
diff --git a/NvCloth/src/cuda/CuSolverKernelBlob.h b/NvCloth/src/cuda/CuSolverKernelBlob.h
new file mode 100644
index 0000000..e69de29
--- /dev/null
+++ b/NvCloth/src/cuda/CuSolverKernelBlob.h