diff options
| author | Marijn Tamis <[email protected]> | 2019-04-01 14:21:09 +0200 |
|---|---|---|
| committer | Marijn Tamis <[email protected]> | 2019-04-01 14:21:09 +0200 |
| commit | d243404d4ba88bcf53f7310cc8980b4efe38c19f (patch) | |
| tree | dcc8ce2904e9f813e03f71f825c4d3c9ec565d91 /NvCloth/src/cuda/CuSolverKernel.cu | |
| parent | Add new SetSpheres and SetPlanes api's to bring them in line with setTriangles. (diff) | |
| download | nvcloth-1.1.6.tar.xz nvcloth-1.1.6.zip | |
1.1.6 Release.1.1.6
Diffstat (limited to 'NvCloth/src/cuda/CuSolverKernel.cu')
| -rw-r--r-- | NvCloth/src/cuda/CuSolverKernel.cu | 3 |
1 files changed, 2 insertions, 1 deletions
diff --git a/NvCloth/src/cuda/CuSolverKernel.cu b/NvCloth/src/cuda/CuSolverKernel.cu index 1ad5896..86f038d 100644 --- a/NvCloth/src/cuda/CuSolverKernel.cu +++ b/NvCloth/src/cuda/CuSolverKernel.cu @@ -851,7 +851,7 @@ __device__ void applyImpulse(SharedParticleData::ParticleReferenceType pos, cons 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 + // 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)); @@ -1394,6 +1394,7 @@ __launch_bounds__(512, 1) for (int32_t i = threadIdx.x; i < configDataSize; i += blockDim.x) gSharedUnsigned[i] = reinterpret_cast<const uint32_t*>(gClothData.mPhaseConfigs)[i]; + Pointer<Shared, uint32_t> scratchPtr = Pointer<Shared, uint32_t>( gSharedUnsigned + configDataSize + 4 * gFrameData.mNumSharedPositions * gClothData.mNumParticles); |