diff options
| author | Marijn Tamis <[email protected]> | 2018-05-03 18:22:48 +0200 |
|---|---|---|
| committer | Marijn Tamis <[email protected]> | 2018-05-03 18:22:48 +0200 |
| commit | ca32c59a58d37c1822e185a2d5f3d0d3e8943593 (patch) | |
| tree | b06b9eec03f34344ef8fc31aa147b2714d3962ee /NvCloth/src/cuda | |
| parent | Forced rename of platform folders in cmake dir. Git didn't pick this up before. (diff) | |
| download | nvcloth-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.h | 2 | ||||
| -rw-r--r-- | NvCloth/src/cuda/CuCollision.h | 7 | ||||
| -rw-r--r-- | NvCloth/src/cuda/CuFabric.cpp | 3 | ||||
| -rw-r--r-- | NvCloth/src/cuda/CuFactory.cpp | 7 | ||||
| -rw-r--r-- | NvCloth/src/cuda/CuSolverKernelBlob.h | 0 |
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 |