NvCloth 1.1.1 Release. (22392725)v1.1.1

29 files changed, 251 insertions, 56 deletions

diff --git a/NvCloth/CmakeGenerateProjects.bat b/NvCloth/CmakeGenerateProjects.bat
index a5891c9..fe50f54 100644
--- a/NvCloth/CmakeGenerateProjects.bat
+++ b/NvCloth/CmakeGenerateProjects.bat
@@ -30,38 +30,38 @@ REM Generate projects here
 rmdir /s /q compiler\vc11win32-cmake\
 mkdir compiler\vc11win32-cmake\
 pushd compiler\vc11win32-cmake\
-cmake ..\cmake\windows -G "Visual Studio 11 2012" -AWin32 -DTARGET_BUILD_PLATFORM=Windows -DPX_GENERATE_GPU_PROJECTS=1 -DCUDA_TOOLKIT_ROOT_DIR="%CUDA_PATH_%" -DSTATIC_WINCRT=1 -DPX_OUTPUT_DLL_DIR=%PX_OUTPUT_ROOT%\Bin\vc11win32-cmake -DPX_OUTPUT_LIB_DIR=%PX_OUTPUT_ROOT%\Lib\vc11win32-cmake -DPX_OUTPUT_EXE_DIR=%PX_OUTPUT_ROOT%\Bin\vc11win32-cmake
+cmake ..\cmake\windows -G "Visual Studio 11 2012" -AWin32 -DTARGET_BUILD_PLATFORM=windows -DPX_GENERATE_GPU_PROJECTS=1 -DCUDA_TOOLKIT_ROOT_DIR="%CUDA_PATH_%" -DSTATIC_WINCRT=1 -DPX_OUTPUT_DLL_DIR=%PX_OUTPUT_ROOT%\Bin\vc11win32-cmake -DPX_OUTPUT_LIB_DIR=%PX_OUTPUT_ROOT%\Lib\vc11win32-cmake -DPX_OUTPUT_EXE_DIR=%PX_OUTPUT_ROOT%\Bin\vc11win32-cmake
 popd
 
 rmdir /s /q compiler\vc11win64-cmake\
 mkdir compiler\vc11win64-cmake\
 pushd compiler\vc11win64-cmake\
-cmake ..\cmake\windows -G "Visual Studio 11 2012" -Ax64 -DTARGET_BUILD_PLATFORM=Windows -DPX_GENERATE_GPU_PROJECTS=1 -DCUDA_TOOLKIT_ROOT_DIR="%CUDA_PATH_%" -DSTATIC_WINCRT=1 -DPX_OUTPUT_DLL_DIR=%PX_OUTPUT_ROOT%\Bin\vc11win64-cmake -DPX_OUTPUT_LIB_DIR=%PX_OUTPUT_ROOT%\Lib\vc11win64-cmake -DPX_OUTPUT_EXE_DIR=%PX_OUTPUT_ROOT%\Bin\vc11win64-cmake
+cmake ..\cmake\windows -G "Visual Studio 11 2012" -Ax64 -DTARGET_BUILD_PLATFORM=windows -DPX_GENERATE_GPU_PROJECTS=1 -DCUDA_TOOLKIT_ROOT_DIR="%CUDA_PATH_%" -DSTATIC_WINCRT=1 -DPX_OUTPUT_DLL_DIR=%PX_OUTPUT_ROOT%\Bin\vc11win64-cmake -DPX_OUTPUT_LIB_DIR=%PX_OUTPUT_ROOT%\Lib\vc11win64-cmake -DPX_OUTPUT_EXE_DIR=%PX_OUTPUT_ROOT%\Bin\vc11win64-cmake
 popd
 
 
 rmdir /s /q compiler\vc12win32-cmake\
 mkdir compiler\vc12win32-cmake\
 pushd compiler\vc12win32-cmake\
-cmake ..\cmake\windows -G "Visual Studio 12 2013" -AWin32 -DTARGET_BUILD_PLATFORM=Windows -DPX_GENERATE_GPU_PROJECTS=1 -DCUDA_TOOLKIT_ROOT_DIR="%CUDA_PATH_%" -DSTATIC_WINCRT=1 -DPX_OUTPUT_DLL_DIR=%PX_OUTPUT_ROOT%\Bin\vc12win32-cmake -DPX_OUTPUT_LIB_DIR=%PX_OUTPUT_ROOT%\Lib\vc12win32-cmake -DPX_OUTPUT_EXE_DIR=%PX_OUTPUT_ROOT%\Bin\vc12win32-cmake
+cmake ..\cmake\windows -G "Visual Studio 12 2013" -AWin32 -DTARGET_BUILD_PLATFORM=windows -DPX_GENERATE_GPU_PROJECTS=1 -DCUDA_TOOLKIT_ROOT_DIR="%CUDA_PATH_%" -DSTATIC_WINCRT=1 -DPX_OUTPUT_DLL_DIR=%PX_OUTPUT_ROOT%\Bin\vc12win32-cmake -DPX_OUTPUT_LIB_DIR=%PX_OUTPUT_ROOT%\Lib\vc12win32-cmake -DPX_OUTPUT_EXE_DIR=%PX_OUTPUT_ROOT%\Bin\vc12win32-cmake
 popd
 
 rmdir /s /q compiler\vc12win64-cmake\
 mkdir compiler\vc12win64-cmake\
 pushd compiler\vc12win64-cmake\
-cmake ..\cmake\windows -G "Visual Studio 12 2013" -Ax64 -DTARGET_BUILD_PLATFORM=Windows -DPX_GENERATE_GPU_PROJECTS=1 -DCUDA_TOOLKIT_ROOT_DIR="%CUDA_PATH_%" -DSTATIC_WINCRT=1 -DPX_OUTPUT_DLL_DIR=%PX_OUTPUT_ROOT%\Bin\vc12win64-cmake -DPX_OUTPUT_LIB_DIR=%PX_OUTPUT_ROOT%\Lib\vc12win64-cmake -DPX_OUTPUT_EXE_DIR=%PX_OUTPUT_ROOT%\Bin\vc12win64-cmake
+cmake ..\cmake\windows -G "Visual Studio 12 2013" -Ax64 -DTARGET_BUILD_PLATFORM=windows -DPX_GENERATE_GPU_PROJECTS=1 -DCUDA_TOOLKIT_ROOT_DIR="%CUDA_PATH_%" -DSTATIC_WINCRT=1 -DPX_OUTPUT_DLL_DIR=%PX_OUTPUT_ROOT%\Bin\vc12win64-cmake -DPX_OUTPUT_LIB_DIR=%PX_OUTPUT_ROOT%\Lib\vc12win64-cmake -DPX_OUTPUT_EXE_DIR=%PX_OUTPUT_ROOT%\Bin\vc12win64-cmake
 popd
 
 rmdir /s /q compiler\vc14win32-cmake\
 mkdir compiler\vc14win32-cmake\
 pushd compiler\vc14win32-cmake\
-cmake ..\cmake\windows -G "Visual Studio 14 2015" -AWin32 -DTARGET_BUILD_PLATFORM=Windows -DPX_GENERATE_GPU_PROJECTS=1 -DCUDA_TOOLKIT_ROOT_DIR="%CUDA_PATH_%" -DSTATIC_WINCRT=1 -DPX_OUTPUT_DLL_DIR=%PX_OUTPUT_ROOT%\Bin\vc14win32-cmake -DPX_OUTPUT_LIB_DIR=%PX_OUTPUT_ROOT%\Lib\vc14win32-cmake -DPX_OUTPUT_EXE_DIR=%PX_OUTPUT_ROOT%\Bin\vc14win32-cmake
+cmake ..\cmake\windows -G "Visual Studio 14 2015" -AWin32 -DTARGET_BUILD_PLATFORM=windows -DPX_GENERATE_GPU_PROJECTS=1 -DCUDA_TOOLKIT_ROOT_DIR="%CUDA_PATH_%" -DSTATIC_WINCRT=1 -DPX_OUTPUT_DLL_DIR=%PX_OUTPUT_ROOT%\Bin\vc14win32-cmake -DPX_OUTPUT_LIB_DIR=%PX_OUTPUT_ROOT%\Lib\vc14win32-cmake -DPX_OUTPUT_EXE_DIR=%PX_OUTPUT_ROOT%\Bin\vc14win32-cmake
 popd
 
 rmdir /s /q compiler\vc14win64-cmake\
 mkdir compiler\vc14win64-cmake\
 pushd compiler\vc14win64-cmake\
-cmake ..\cmake\windows -G "Visual Studio 14 2015" -Ax64 -DTARGET_BUILD_PLATFORM=Windows -DPX_GENERATE_GPU_PROJECTS=1 -DCUDA_TOOLKIT_ROOT_DIR="%CUDA_PATH_%" -DSTATIC_WINCRT=1 -DPX_OUTPUT_DLL_DIR=%PX_OUTPUT_ROOT%\Bin\vc14win64-cmake -DPX_OUTPUT_LIB_DIR=%PX_OUTPUT_ROOT%\Lib\vc14win64-cmake -DPX_OUTPUT_EXE_DIR=%PX_OUTPUT_ROOT%\Bin\vc14win64-cmake
+cmake ..\cmake\windows -G "Visual Studio 14 2015" -Ax64 -DTARGET_BUILD_PLATFORM=windows -DPX_GENERATE_GPU_PROJECTS=1 -DCUDA_TOOLKIT_ROOT_DIR="%CUDA_PATH_%" -DSTATIC_WINCRT=1 -DPX_OUTPUT_DLL_DIR=%PX_OUTPUT_ROOT%\Bin\vc14win64-cmake -DPX_OUTPUT_LIB_DIR=%PX_OUTPUT_ROOT%\Lib\vc14win64-cmake -DPX_OUTPUT_EXE_DIR=%PX_OUTPUT_ROOT%\Bin\vc14win64-cmake
 popd
 
 
diff --git a/NvCloth/GenerateProjectsLinux.sh b/NvCloth/GenerateProjectsLinux.sh
index e34e51f..56d1788 100644
--- a/NvCloth/GenerateProjectsLinux.sh
+++ b/NvCloth/GenerateProjectsLinux.sh
@@ -19,23 +19,23 @@ export PX_OUTPUT_ROOT="$PWD"
 rm -r -f compiler/linux64-debug-cmake/
 mkdir compiler/linux64-debug-cmake/
 cd compiler/linux64-debug-cmake/
-cmake ../cmake/Linux -G "Unix Makefiles" -DTARGET_BUILD_PLATFORM=Linux -DCMAKE_BUILD_TYPE=debug -DPX_GENERATE_GPU_PROJECTS=0 -DPX_OUTPUT_DLL_DIR=$PX_OUTPUT_ROOT/bin/linux64-cmake -DPX_OUTPUT_LIB_DIR=$PX_OUTPUT_ROOT/lib/linux64-cmake -DPX_OUTPUT_EXE_DIR=$PX_OUTPUT_ROOT/bin/linux64-cmake
+cmake ../cmake/Linux -G "Unix Makefiles" -DTARGET_BUILD_PLATFORM=linux -DCMAKE_BUILD_TYPE=debug -DPX_GENERATE_GPU_PROJECTS=0 -DPX_OUTPUT_DLL_DIR=$PX_OUTPUT_ROOT/bin/linux64-cmake -DPX_OUTPUT_LIB_DIR=$PX_OUTPUT_ROOT/lib/linux64-cmake -DPX_OUTPUT_EXE_DIR=$PX_OUTPUT_ROOT/bin/linux64-cmake
 cd ../../
 
 rm -r -f compiler/linux64-checked-cmake/
 mkdir compiler/linux64-checked-cmake/
 cd compiler/linux64-checked-cmake/
-cmake ../cmake/Linux -G "Unix Makefiles" -DTARGET_BUILD_PLATFORM=Linux -DCMAKE_BUILD_TYPE=checked -DPX_GENERATE_GPU_PROJECTS=0 -DPX_OUTPUT_DLL_DIR=$PX_OUTPUT_ROOT/bin/linux64-cmake -DPX_OUTPUT_LIB_DIR=$PX_OUTPUT_ROOT/lib/linux64-cmake -DPX_OUTPUT_EXE_DIR=$PX_OUTPUT_ROOT/bin/linux64-cmake
+cmake ../cmake/Linux -G "Unix Makefiles" -DTARGET_BUILD_PLATFORM=linux -DCMAKE_BUILD_TYPE=checked -DPX_GENERATE_GPU_PROJECTS=0 -DPX_OUTPUT_DLL_DIR=$PX_OUTPUT_ROOT/bin/linux64-cmake -DPX_OUTPUT_LIB_DIR=$PX_OUTPUT_ROOT/lib/linux64-cmake -DPX_OUTPUT_EXE_DIR=$PX_OUTPUT_ROOT/bin/linux64-cmake
 cd ../../
 
 rm -r -f compiler/linux64-profile-cmake/
 mkdir compiler/linux64-profile-cmake/
 cd compiler/linux64-profile-cmake/
-cmake ../cmake/Linux -G "Unix Makefiles" -DTARGET_BUILD_PLATFORM=Linux -DCMAKE_BUILD_TYPE=profile -DPX_GENERATE_GPU_PROJECTS=0 -DPX_OUTPUT_DLL_DIR=$PX_OUTPUT_ROOT/bin/linux64-cmake -DPX_OUTPUT_LIB_DIR=$PX_OUTPUT_ROOT/lib/linux64-cmake -DPX_OUTPUT_EXE_DIR=$PX_OUTPUT_ROOT/bin/linux64-cmake
+cmake ../cmake/Linux -G "Unix Makefiles" -DTARGET_BUILD_PLATFORM=linux -DCMAKE_BUILD_TYPE=profile -DPX_GENERATE_GPU_PROJECTS=0 -DPX_OUTPUT_DLL_DIR=$PX_OUTPUT_ROOT/bin/linux64-cmake -DPX_OUTPUT_LIB_DIR=$PX_OUTPUT_ROOT/lib/linux64-cmake -DPX_OUTPUT_EXE_DIR=$PX_OUTPUT_ROOT/bin/linux64-cmake
 cd ../../
 
 rm -r -f compiler/linux64-release-cmake/
 mkdir compiler/linux64-release-cmake/
 cd compiler/linux64-release-cmake/
-cmake ../cmake/Linux -G "Unix Makefiles" -DTARGET_BUILD_PLATFORM=Linux -DCMAKE_BUILD_TYPE=release -DPX_GENERATE_GPU_PROJECTS=0 -DPX_OUTPUT_DLL_DIR=$PX_OUTPUT_ROOT/bin/linux64-cmake -DPX_OUTPUT_LIB_DIR=$PX_OUTPUT_ROOT/lib/linux64-cmake -DPX_OUTPUT_EXE_DIR=$PX_OUTPUT_ROOT/bin/linux64-cmake
+cmake ../cmake/Linux -G "Unix Makefiles" -DTARGET_BUILD_PLATFORM=linux -DCMAKE_BUILD_TYPE=release -DPX_GENERATE_GPU_PROJECTS=0 -DPX_OUTPUT_DLL_DIR=$PX_OUTPUT_ROOT/bin/linux64-cmake -DPX_OUTPUT_LIB_DIR=$PX_OUTPUT_ROOT/lib/linux64-cmake -DPX_OUTPUT_EXE_DIR=$PX_OUTPUT_ROOT/bin/linux64-cmake
 cd ../../
\ No newline at end of file
diff --git a/NvCloth/GenerateProjectsOsx.sh b/NvCloth/GenerateProjectsOsx.sh
index 5f0e3c7..5c57f42 100644
--- a/NvCloth/GenerateProjectsOsx.sh
+++ b/NvCloth/GenerateProjectsOsx.sh
@@ -19,13 +19,13 @@ export PX_OUTPUT_ROOT="$PWD"
 rm -r -f compiler/osx32-cmake/
 mkdir compiler/osx32-cmake/
 cd compiler/osx32-cmake/
-cmake ../cmake/Mac -G Xcode -DTARGET_BUILD_PLATFORM=Mac -DPX_32BIT=1 -DPX_OUTPUT_LIB_DIR=$PX_OUTPUT_ROOT/Lib/osx32-cmake 
+cmake ../cmake/Mac -G Xcode -DTARGET_BUILD_PLATFORM=mac -DPX_32BIT=1 -DPX_OUTPUT_LIB_DIR=$PX_OUTPUT_ROOT/Lib/osx32-cmake 
 cd ../../
 
 rm -r -f compiler/osx64-cmake/
 mkdir compiler/osx64-cmake/
 cd compiler/osx64-cmake/
-cmake ../cmake/Mac -G Xcode -DTARGET_BUILD_PLATFORM=Mac -DPX_OUTPUT_LIB_DIR=$PX_OUTPUT_ROOT/Lib/osx64-cmake 
+cmake ../cmake/Mac -G Xcode -DTARGET_BUILD_PLATFORM=mac -DPX_OUTPUT_LIB_DIR=$PX_OUTPUT_ROOT/Lib/osx64-cmake 
 cd ../../
 
 
diff --git a/NvCloth/ReleaseNotes.txt b/NvCloth/ReleaseNotes.txt
index 5751a25..7dfa9c0 100644
--- a/NvCloth/ReleaseNotes.txt
+++ b/NvCloth/ReleaseNotes.txt
@@ -1,9 +1,9 @@
 NvCloth Release notes.
 
-[1.1.0]
+[1.1.1]
 Supported platforms:
-* Windows (CPU, CUDA, DX11, RenderCloth)
-* Mac (beta, tested on Sierra, RenderCloth)
+* Windows (CPU, CUDA, DX11)
+* Mac (beta, tested on Sierra)
 * Linux (beta, tested on Ubuntu 16.04.1 LTS x64)
 * PlayStation 4 (beta) with PS4 SDK 4.5
 * Xbox one (CPU, DX11, beta)
@@ -17,6 +17,15 @@ Supported compilers (via cmake 3.7)
 * Xbox one: Visual studio 14 (2015)
 
 Fixed:
+Air drag/lift did react to flow speed only linearly, model was changed to be more physically correct.
+Fixed issue where scaling units would result in different air drag/lift behavior (by adding fluid density parameter).
+Fixed "unused typedef `__t100`" warning in xcode.
+
+Known issues:
+DirectX 11 solver has issues with ccd sphere/capsule collision.
+
+[1.1.0]
+Fixed:
 DirectX 11 solver now supports triangle collision shapes and air drag/lift.
 Fixed bug where changing friction on cloth already in simulation would cause assert or crash.
 Optimization flag now working on PS4 release build.
diff --git a/NvCloth/compiler/cmake/Linux/CMakeLists.txt b/NvCloth/compiler/cmake/Linux/CMakeLists.txt
index 4d8d85a..9425dde 100644
--- a/NvCloth/compiler/cmake/Linux/CMakeLists.txt
+++ b/NvCloth/compiler/cmake/Linux/CMakeLists.txt
@@ -4,10 +4,10 @@ include(../common/CMakeLists.txt)
 
 
 IF(NOT DEFINED TARGET_BUILD_PLATFORM) # Not defined, default to Linux
-	SET(TARGET_BUILD_PLATFORM "Linux")
+	SET(TARGET_BUILD_PLATFORM "linux")
 ENDIF()
 
-SET(PLATFORM_LIST Linux)
+SET(PLATFORM_LIST linux)
 
 IF (NOT ${TARGET_BUILD_PLATFORM} IN_LIST PLATFORM_LIST)
 	MESSAGE(FATAL_ERROR "Invalid platform:" ${TARGET_BUILD_PLATFORM})
@@ -47,7 +47,7 @@ FIND_PACKAGE(PxShared REQUIRED)
 #  such as - we don't want to be installing any built artifacts on this server, etc. So for now we hack it.
 
 # Add PxShared as a dependency so that we can use project references
-ADD_SUBDIRECTORY(${PXSHARED_ROOT_DIR}/src/compiler/cmake/Linux "${CMAKE_CURRENT_BINARY_DIR}/pxshared_bin")
+ADD_SUBDIRECTORY(${PXSHARED_ROOT_DIR}/src/compiler/cmake/linux "${CMAKE_CURRENT_BINARY_DIR}/pxshared_bin")
 
 # Include all of the projects
 INCLUDE(NvCloth.cmake)
diff --git a/NvCloth/compiler/cmake/Mac/CMakeLists.txt b/NvCloth/compiler/cmake/Mac/CMakeLists.txt
index 05a42ff..4e1c124 100644
--- a/NvCloth/compiler/cmake/Mac/CMakeLists.txt
+++ b/NvCloth/compiler/cmake/Mac/CMakeLists.txt
@@ -4,10 +4,10 @@ include(../common/CMakeLists.txt)
 
 
 IF(NOT DEFINED TARGET_BUILD_PLATFORM) # Not defined, default to Mac
-	SET(TARGET_BUILD_PLATFORM "Mac")
+	SET(TARGET_BUILD_PLATFORM "mac")
 ENDIF()
 
-SET(PLATFORM_LIST Mac)
+SET(PLATFORM_LIST mac)
 
 IF (NOT ${TARGET_BUILD_PLATFORM} IN_LIST PLATFORM_LIST)
 	MESSAGE(FATAL_ERROR "Invalid platform:" ${TARGET_BUILD_PLATFORM})
diff --git a/NvCloth/compiler/cmake/windows/CMakeLists.txt b/NvCloth/compiler/cmake/windows/CMakeLists.txt
index 79b0236..81d8289 100644
--- a/NvCloth/compiler/cmake/windows/CMakeLists.txt
+++ b/NvCloth/compiler/cmake/windows/CMakeLists.txt
@@ -4,10 +4,10 @@ include(../common/CMakeLists.txt)
 
 
 IF(NOT DEFINED TARGET_BUILD_PLATFORM) # Not defined, default to Windows
-	SET(TARGET_BUILD_PLATFORM "Windows")
+	SET(TARGET_BUILD_PLATFORM "windows")
 ENDIF()
 
-SET(PLATFORM_LIST Windows)
+SET(PLATFORM_LIST windows)
 
 IF (NOT ${TARGET_BUILD_PLATFORM} IN_LIST PLATFORM_LIST)
 	MESSAGE(FATAL_ERROR "Invalid platform:" ${TARGET_BUILD_PLATFORM})
diff --git a/NvCloth/extensions/src/ClothFabricCooker.cpp b/NvCloth/extensions/src/ClothFabricCooker.cpp
index f0e4dea..2d800e9 100644
--- a/NvCloth/extensions/src/ClothFabricCooker.cpp
+++ b/NvCloth/extensions/src/ClothFabricCooker.cpp
@@ -97,7 +97,9 @@ nv::cloth::Range<const T> CreateRange(typename nv::cloth::Vector<T>::Type const&
 template<typename T, typename U>
 nv::cloth::Range<const T> CreateRangeF(typename nv::cloth::Vector<U>::Type const& vector, int offset = 0)
 {
+#ifndef _LIBCPP_HAS_NO_STATIC_ASSERT
 	static_assert(sizeof(T) == sizeof(U), "Type T and U need to be of the same size");
+#endif
 	const T* begin = reinterpret_cast<const T*>(vector.begin()+offset);
 	const T* end = reinterpret_cast<const T*>(vector.end());
 
diff --git a/NvCloth/include/NvCloth/Cloth.h b/NvCloth/include/NvCloth/Cloth.h
index 3d783e9..ac957f7 100644
--- a/NvCloth/include/NvCloth/Cloth.h
+++ b/NvCloth/include/NvCloth/Cloth.h
@@ -376,6 +376,10 @@ class Cloth : public UserAllocated
 	virtual void setLiftCoefficient(float) = 0;
 	///Returns value set with setLiftCoefficient().
 	virtual float getLiftCoefficient() const = 0;
+	/** /brief Sets the fluid density used for air drag/lift calculations. */
+	virtual void setFluidDensity(float) = 0;
+	///Returns value set with setFluidDensity().
+	virtual float getFluidDensity() const = 0;
 
 	/* self collision */
 
diff --git a/NvCloth/samples/CmakeGenerateProjects.bat b/NvCloth/samples/CmakeGenerateProjects.bat
index d63cde8..8c6d88c 100644
--- a/NvCloth/samples/CmakeGenerateProjects.bat
+++ b/NvCloth/samples/CmakeGenerateProjects.bat
@@ -14,6 +14,9 @@ SET PATH=%PATH%;"%CMAKE_PATH_F%"
 
 REM Make sure the various variables that we need are set
 
+call "../scripts/locate_cuda.bat" CUDA_PATH_
+echo CUDA_PATH_ = %CUDA_PATH_%
+
 IF EXIST %~dp0..\Externals\CMakeModules (
 	set GW_DEPS_ROOT=%~dp0..\
 )
@@ -30,13 +33,13 @@ REM Generate projects here
 rmdir /s /q compiler\vc14win32-cmake\
 mkdir compiler\vc14win32-cmake\
 pushd compiler\vc14win32-cmake\
-cmake ..\.. -G "Visual Studio 14 2015" -AWin32 -DTARGET_BUILD_PLATFORM=Windows -DSTATIC_WINCRT=0 -DBL_DLL_OUTPUT_DIR=%OUTPUT_ROOT%\bin\vc14win32-cmake -DBL_LIB_OUTPUT_DIR=%OUTPUT_ROOT%\lib\vc14win32-cmake -DBL_EXE_OUTPUT_DIR=%OUTPUT_ROOT%\bin\vc14win32-cmake
+cmake ..\.. -G "Visual Studio 14 2015" -AWin32 -DCUDA_TOOLKIT_ROOT_DIR="%CUDA_PATH_%" -DTARGET_BUILD_PLATFORM=windows -DSTATIC_WINCRT=0 -DBL_DLL_OUTPUT_DIR=%OUTPUT_ROOT%\bin\vc14win32-cmake -DBL_LIB_OUTPUT_DIR=%OUTPUT_ROOT%\lib\vc14win32-cmake -DBL_EXE_OUTPUT_DIR=%OUTPUT_ROOT%\bin\vc14win32-cmake
 popd
 
 rmdir /s /q compiler\vc14win64-cmake\
 mkdir compiler\vc14win64-cmake\
 pushd compiler\vc14win64-cmake\
-cmake ..\.. -G "Visual Studio 14 2015" -Ax64 -DTARGET_BUILD_PLATFORM=Windows -DSTATIC_WINCRT=0 -DBL_DLL_OUTPUT_DIR=%OUTPUT_ROOT%\bin\vc14win64-cmake -DBL_LIB_OUTPUT_DIR=%OUTPUT_ROOT%\lib\vc14win64-cmake -DBL_EXE_OUTPUT_DIR=%OUTPUT_ROOT%\bin\vc14win64-cmake
+cmake ..\.. -G "Visual Studio 14 2015" -Ax64 -DCUDA_TOOLKIT_ROOT_DIR="%CUDA_PATH_%" -DTARGET_BUILD_PLATFORM=windows -DSTATIC_WINCRT=0 -DBL_DLL_OUTPUT_DIR=%OUTPUT_ROOT%\bin\vc14win64-cmake -DBL_LIB_OUTPUT_DIR=%OUTPUT_ROOT%\lib\vc14win64-cmake -DBL_EXE_OUTPUT_DIR=%OUTPUT_ROOT%\bin\vc14win64-cmake
 popd
 
 
diff --git a/NvCloth/samples/SampleBase/Main.cpp b/NvCloth/samples/SampleBase/Main.cpp
index 433ed88..f4107c3 100644
--- a/NvCloth/samples/SampleBase/Main.cpp
+++ b/NvCloth/samples/SampleBase/Main.cpp
@@ -12,7 +12,9 @@
 #include <sstream>
 
 #include <Windows.h>
-
+#include <iostream>
+#include <io.h>
+#include <fcntl.h>
 
 using namespace std;
 
@@ -22,6 +24,8 @@ int WINAPI wWinMain(HINSTANCE hInstance, HINSTANCE hPrevInstance, LPWSTR lpCmdLi
 #if defined(DEBUG) | defined(_DEBUG)
 	_CrtSetDbgFlag(_CRTDBG_ALLOC_MEM_DF | _CRTDBG_LEAK_CHECK_DF);
 	AllocConsole();
+	FILE* fp;
+	freopen_s(&fp, "CONOUT$", "w", stdout);
 #endif
 
 	SampleConfig config;
diff --git a/NvCloth/samples/SampleBase/scene/scenes/ScaledScene.cpp b/NvCloth/samples/SampleBase/scene/scenes/ScaledScene.cpp
new file mode 100644
index 0000000..f2a5b82
--- /dev/null
+++ b/NvCloth/samples/SampleBase/scene/scenes/ScaledScene.cpp
@@ -0,0 +1,98 @@
+/*
+* Copyright (c) 2008-2017, NVIDIA CORPORATION.  All rights reserved.
+*
+* NVIDIA CORPORATION and its licensors retain all intellectual property
+* and proprietary rights in and to this software, related documentation
+* and any modifications thereto.  Any use, reproduction, disclosure or
+* distribution of this software and related documentation without an express
+* license agreement from NVIDIA CORPORATION is strictly prohibited.
+*/
+
+#include "ScaledScene.h"
+#include "Scene/SceneController.h"
+#include <NvClothExt/ClothFabricCooker.h>
+#include "ClothMeshGenerator.h"
+#include <NvCloth/Fabric.h>
+#include <NvCloth/Solver.h>
+#include <NvCloth/Cloth.h>
+#include <NvCloth/Factory.h>
+#include "Renderer.h"
+#include "renderer/RenderUtils.h"
+ 
+DECLARE_SCENE_NAME(ScaledScene,"Scaled Scene")
+
+void ScaledScene::onInitialize()
+{
+	///////////////////////////////////////////////////////////////////////
+	ClothMeshData clothMesh;
+
+	physx::PxMat44 transform = PxTransform(PxVec3(-2.f, 13.f, 0.f)*0.0, PxQuat(PxPi / 6.f, PxVec3(1.f, 0.f, 0.f)));
+	clothMesh.GeneratePlaneCloth(600.f, 700.f, 49, 59, false, transform);
+	clothMesh.AttachClothPlaneByAngles(49, 59);
+
+	mClothActor = new ClothActor;
+	nv::cloth::ClothMeshDesc meshDesc = clothMesh.GetClothMeshDesc();
+	{
+		mClothActor->mClothRenderMesh = new ClothRenderMesh(meshDesc);
+		mClothActor->mClothRenderable = getSceneController()->getRenderer().createRenderable(*(static_cast<IRenderMesh*>(mClothActor->mClothRenderMesh)), *getSceneController()->getDefaultMaterial());
+		mClothActor->mClothRenderable->setColor(getRandomPastelColor());
+		mClothActor->mClothRenderable->setScale(physx::PxVec3(0.01, 0.01, 0.01));
+		mClothActor->mClothRenderable->setTransform(PxTransform(PxVec3(-2.f, 13.f, 0.f),physx::PxQuat(physx::PxIdentity)));
+	}
+
+	nv::cloth::Vector<int32_t>::Type phaseTypeInfo;
+	mFabric = NvClothCookFabricFromMesh(getSceneController()->getFactory(), meshDesc, physx::PxVec3(0.0f, -9.8f, 0.0f), &phaseTypeInfo, false);
+	trackFabric(mFabric);
+
+	// Initialize start positions and masses for the actual cloth instance
+	// (note: the particle/vertex positions do not have to match the mesh description here. Set the positions to the initial shape of this cloth instance)
+	std::vector<physx::PxVec4> particlesCopy;
+	particlesCopy.resize(clothMesh.mVertices.size());
+
+	physx::PxVec3 center = transform.transform(physx::PxVec3(0.0f, 0.0f, 0.0f));
+	for (int i = 0; i < (int)clothMesh.mVertices.size(); i++)
+	{
+		// To put attachment point closer to each other
+		if(clothMesh.mInvMasses[i] < 1e-6)
+			clothMesh.mVertices[i] = (clothMesh.mVertices[i] - center) * 0.85f + center;
+
+		particlesCopy[i] = physx::PxVec4(clothMesh.mVertices[i], clothMesh.mInvMasses[i]); // w component is 1/mass, or 0.0f for anchored/fixed particles
+	}
+
+	// Create the cloth from the initial positions/masses and the fabric
+	mClothActor->mCloth = getSceneController()->getFactory()->createCloth(nv::cloth::Range<physx::PxVec4>(&particlesCopy[0], &particlesCopy[0] + particlesCopy.size()), *mFabric);
+	particlesCopy.clear(); particlesCopy.shrink_to_fit();
+
+	mClothActor->mCloth->setGravity(physx::PxVec3(0.0f, -980.0f, 0.0f));
+
+	// Setup phase configs
+	std::vector<nv::cloth::PhaseConfig> phases(mFabric->getNumPhases());
+	for (int i = 0; i < (int)phases.size(); i++)
+	{
+		phases[i].mPhaseIndex = i;
+		phases[i].mStiffness = 1.0f;
+		phases[i].mStiffnessMultiplier = 1.0f;
+		phases[i].mCompressionLimit = 1.0f;
+		phases[i].mStretchLimit = 1.0f;
+	}
+	mClothActor->mCloth->setPhaseConfig(nv::cloth::Range<nv::cloth::PhaseConfig>(&phases.front(), &phases.back()));
+	mClothActor->mCloth->setDragCoefficient(0.1f);
+	mClothActor->mCloth->setLiftCoefficient(0.1f);
+	//mClothActor->mCloth->setWindVelocity(physx::PxVec3(50, 0.0, 50.0));
+	mClothActor->mCloth->setFluidDensity(1.0f / powf(100, 3));
+
+	mSolver = getSceneController()->getFactory()->createSolver();
+	trackSolver(mSolver);
+	trackClothActor(mClothActor);
+
+	// Add the cloth to the solver for simulation
+	addClothToSolver(mClothActor, mSolver);
+	
+	{
+		IRenderMesh* mesh = getSceneController()->getRenderer().getPrimitiveRenderMesh(PrimitiveRenderMeshType::Plane);
+		Renderable* plane = getSceneController()->getRenderer().createRenderable(*mesh, *getSceneController()->getDefaultPlaneMaterial());
+		plane->setTransform(PxTransform(PxVec3(0.f, 0.f, 0.f), PxQuat(PxPiDivTwo, PxVec3(0.f, 0.f, 1.f))));
+		plane->setScale(PxVec3(1000.f));
+		trackRenderable(plane);
+	}
+}
diff --git a/NvCloth/samples/SampleBase/scene/scenes/ScaledScene.h b/NvCloth/samples/SampleBase/scene/scenes/ScaledScene.h
new file mode 100644
index 0000000..eb47c36
--- /dev/null
+++ b/NvCloth/samples/SampleBase/scene/scenes/ScaledScene.h
@@ -0,0 +1,32 @@
+/*
+* Copyright (c) 2008-2017, NVIDIA CORPORATION.  All rights reserved.
+*
+* NVIDIA CORPORATION and its licensors retain all intellectual property
+* and proprietary rights in and to this software, related documentation
+* and any modifications thereto.  Any use, reproduction, disclosure or
+* distribution of this software and related documentation without an express
+* license agreement from NVIDIA CORPORATION is strictly prohibited.
+*/
+
+#ifndef SCALED_SCENE_H
+#define SCALED_SCENE_H
+
+#include "scene/Scene.h"
+
+class ScaledScene : public Scene
+{
+public:
+
+	ScaledScene(SceneController* sceneController):Scene(sceneController) {}
+
+	virtual void onInitialize() override;
+
+private:
+	nv::cloth::Fabric* mFabric;
+	nv::cloth::Solver* mSolver;
+	ClothActor* mClothActor;
+
+};
+
+
+#endif
\ No newline at end of file
diff --git a/NvCloth/samples/SampleBase/scene/scenes/WindScene.cpp b/NvCloth/samples/SampleBase/scene/scenes/WindScene.cpp
index 232316d..d84fb87 100644
--- a/NvCloth/samples/SampleBase/scene/scenes/WindScene.cpp
+++ b/NvCloth/samples/SampleBase/scene/scenes/WindScene.cpp
@@ -35,7 +35,7 @@ void WindScene::Animate(double dt)
 
 		for(int i = 0; i < 3; i++)
 		{
-			physx::PxVec3 wind = physx::PxVec3(dvx, vy, dvz);
+			physx::PxVec3 wind = physx::PxVec3(dvx, vy, dvz)/5.0f;
 			mClothActor[i]->mCloth->setWindVelocity(wind);
 		}
 	}
diff --git a/NvCloth/samples/compiler/cmake/SampleBase.cmake b/NvCloth/samples/compiler/cmake/SampleBase.cmake
index e3aa2cc..cee4779 100644
--- a/NvCloth/samples/compiler/cmake/SampleBase.cmake
+++ b/NvCloth/samples/compiler/cmake/SampleBase.cmake
@@ -110,6 +110,8 @@ SET(SCENES_FILES
 	${SB_SCENE_SOURCE_DIR}/scenes/CCDScene.h
 	${SB_SCENE_SOURCE_DIR}/scenes/MultiSolverScene.cpp
 	${SB_SCENE_SOURCE_DIR}/scenes/MultiSolverScene.h
+	${SB_SCENE_SOURCE_DIR}/scenes/ScaledScene.cpp
+	${SB_SCENE_SOURCE_DIR}/scenes/ScaledScene.h
 )
 
 SET(UI_FILES
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