diff options
| author | Marijn Tamis <[email protected]> | 2017-07-31 13:52:20 +0200 |
|---|---|---|
| committer | Marijn Tamis <[email protected]> | 2017-07-31 13:52:20 +0200 |
| commit | 223aff8b3f73bb786dce5c67b83ff55208d43969 (patch) | |
| tree | 2e3e2760cb49afbf8d9379e23e27d175bbba27aa /NvCloth/src/cuda | |
| parent | Remove unused copy of PxShared. (diff) | |
| download | nvcloth-1.1.2.tar.xz nvcloth-1.1.2.zip | |
NvCloth 1.1.2 Release. (22576033)v1.1.2
Diffstat (limited to 'NvCloth/src/cuda')
| -rw-r--r-- | NvCloth/src/cuda/CuSolverKernel.cu | 13 |
1 files changed, 10 insertions, 3 deletions
diff --git a/NvCloth/src/cuda/CuSolverKernel.cu b/NvCloth/src/cuda/CuSolverKernel.cu index edb66dc..1ad5896 100644 --- a/NvCloth/src/cuda/CuSolverKernel.cu +++ b/NvCloth/src/cuda/CuSolverKernel.cu @@ -849,10 +849,17 @@ __device__ float3 cross3(const float3& u, const float3& v) __device__ void applyImpulse(SharedParticleData::ParticleReferenceType pos, const float3& impulse) { float scale = -pos.mReferences[3]; + +#if CONVERT_ADDRESSES //Use this instead of atomicAdd function to work around compiler issue treating the pointer as global memory instead of shared memory - asm("red.shared.add.f32 [%0], %1;" :: POINTER_CONSTRAINT(pos.mReferences[0].mPtr), "f"(impulse.x * scale)); - asm("red.shared.add.f32 [%0], %1;" :: POINTER_CONSTRAINT(pos.mReferences[1].mPtr), "f"(impulse.y * scale)); - asm("red.shared.add.f32 [%0], %1;" :: POINTER_CONSTRAINT(pos.mReferences[2].mPtr), "f"(impulse.z * scale)); + asm("red.shared.add.f32 [%0], %1;" ::POINTER_CONSTRAINT(pos.mReferences[0].mPtr), "f"(impulse.x * scale)); + asm("red.shared.add.f32 [%0], %1;" ::POINTER_CONSTRAINT(pos.mReferences[1].mPtr), "f"(impulse.y * scale)); + asm("red.shared.add.f32 [%0], %1;" ::POINTER_CONSTRAINT(pos.mReferences[2].mPtr), "f"(impulse.z * scale)); +#else + atomicAdd(pos.mReferences[0].mPtr, impulse.x * scale); + atomicAdd(pos.mReferences[1].mPtr, impulse.y * scale); + atomicAdd(pos.mReferences[2].mPtr, impulse.z * scale); +#endif } __device__ void applyImpulse(GlobalParticleData::ParticleReferenceType pos, const float3& impulse) { |