Initial commit:

PhysX 3.4.0 Update @ 21294896
APEX 1.4.0 Update @ 21275617

[CL 21300167]

28 files changed, 5813 insertions, 0 deletions

diff --git a/APEX_1.4/module/fieldsampler/cuda/include/common.h b/APEX_1.4/module/fieldsampler/cuda/include/common.h
new file mode 100644
index 00000000..02332daa
--- /dev/null
+++ b/APEX_1.4/module/fieldsampler/cuda/include/common.h
@@ -0,0 +1,31 @@
+/*
+ * Copyright (c) 2008-2015, 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 __COMMON_H__
+#define __COMMON_H__
+
+#define APEX_CUDA_MODULE_PREFIX FieldSampler_
+
+#include "ApexCuda.h"
+#include "InplaceTypes.h"
+
+#include "../include/FieldSamplerCommon.h"
+
+
+namespace physx
+{
+namespace apex
+{
+
+}
+} // namespace physx::apex
+
+#endif
diff --git a/APEX_1.4/module/fieldsampler/cuda/include/fieldsampler.h b/APEX_1.4/module/fieldsampler/cuda/include/fieldsampler.h
new file mode 100644
index 00000000..28e4896c
--- /dev/null
+++ b/APEX_1.4/module/fieldsampler/cuda/include/fieldsampler.h
@@ -0,0 +1,39 @@
+/*
+ * Copyright (c) 2008-2015, 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.
+ */
+
+
+APEX_CUDA_SURFACE_3D(surfRefGridAccum)
+
+
+APEX_CUDA_BOUND_KERNEL(APEX_CUDA_KERNEL_DEFAULT_CONFIG, clearKernel,
+                       ((APEX_MEM_BLOCK(float4), g_accumField))((APEX_MEM_BLOCK(float4), g_accumVelocity))
+                      )
+
+APEX_CUDA_BOUND_KERNEL(APEX_CUDA_KERNEL_DEFAULT_CONFIG, composeKernel,
+                       ((APEX_MEM_BLOCK(float4), g_accumField))((APEX_MEM_BLOCK(const float4), g_accumVelocity))((APEX_MEM_BLOCK(const float4), g_velocity))((float, timestep))
+                      )
+
+APEX_CUDA_FREE_KERNEL_3D(APEX_CUDA_KERNEL_DEFAULT_CONFIG, clearGridKernel,
+                       ((uint32_t, numX))((uint32_t, numY))((uint32_t, numZ))
+                      )
+
+APEX_CUDA_BOUND_KERNEL(APEX_CUDA_KERNEL_DEFAULT_CONFIG, applyParticlesKernel,
+                       ((float4*, g_velocity))((const float4*, g_outField))
+                      )
+
+#ifdef APEX_TEST
+
+APEX_CUDA_BOUND_KERNEL((), testParticleKernel,
+                       ((float4*, g_position))((float4*, g_velocity))
+					   ((uint32_t*, g_flag))
+					   ((const float4*, g_initPosition))((const float4*, g_initVelocity))
+                      )
+
+#endif
\ No newline at end of file
diff --git a/APEX_1.4/module/fieldsampler/cuda/include/fieldsamplerInc.h b/APEX_1.4/module/fieldsampler/cuda/include/fieldsamplerInc.h
new file mode 100644
index 00000000..23a9c3fb
--- /dev/null
+++ b/APEX_1.4/module/fieldsampler/cuda/include/fieldsamplerInc.h
@@ -0,0 +1,155 @@
+/*
+ * Copyright (c) 2008-2015, 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.
+ */
+
+
+APEX_CUDA_SURFACE_3D(surfRefFieldSamplerGridAccum)
+
+
+#ifdef FIELD_SAMPLER_SEPARATE_KERNELS
+
+APEX_CUDA_BOUND_KERNEL(FIELD_SAMPLER_POINTS_KERNEL_CONFIG, fieldSamplerPointsKernel,
+                       ((nvidia::fieldsampler::FieldSamplerKernelParams, baseParams))
+                       ((APEX_MEM_BLOCK(float4), accumField))
+                       ((APEX_MEM_BLOCK(float4), accumVelocity))
+                       ((APEX_MEM_BLOCK(const float4), positionMass))
+                       ((APEX_MEM_BLOCK(const float4), velocity))
+                       ((nvidia::fieldsampler::FieldSamplerParamsEx, paramsEx))
+                       ((InplaceHandle<nvidia::fieldsampler::FieldSamplerQueryParams>, queryParamsHandle))
+					   ((nvidia::fieldsampler::FieldSamplerKernelMode::Enum, kernelMode))
+                      )
+APEX_CUDA_FREE_KERNEL_3D(FIELD_SAMPLER_GRID_KERNEL_CONFIG, fieldSamplerGridKernel,
+                       ((nvidia::fieldsampler::FieldSamplerKernelParams, baseParams))
+                       ((nvidia::fieldsampler::FieldSamplerGridKernelParams, gridParams))
+                       ((nvidia::fieldsampler::FieldSamplerParamsEx, paramsEx))
+                       ((InplaceHandle<nvidia::fieldsampler::FieldSamplerQueryParams>, queryParamsHandle))
+					   ((nvidia::fieldsampler::FieldSamplerKernelMode::Enum, kernelMode))
+                      )
+
+#ifndef __CUDACC__
+#define LAUNCH_FIELD_SAMPLER_KERNEL( launchData ) \
+	const ApexCudaConstStorage& _storage_ = *getFieldSamplerCudaConstStorage(); \
+	InplaceHandle<nvidia::fieldsampler::FieldSamplerQueryParams> queryParamsHandle = _storage_.mappedHandle( launchData.queryParamsHandle ); \
+	uint32_t fieldSamplerCount = launchData.fieldSamplerArray->size(); \
+	switch( launchData.kernelType ) \
+	{ \
+	case nvidia::fieldsampler::FieldSamplerKernelType::POINTS: \
+		{ \
+			const nvidia::fieldsampler::FieldSamplerPointsKernelLaunchDataIntl& data = static_cast<const nvidia::fieldsampler::FieldSamplerPointsKernelLaunchDataIntl&>(launchData); \
+			const nvidia::fieldsampler::FieldSamplerPointsKernelArgs* kernelArgs = static_cast<const nvidia::fieldsampler::FieldSamplerPointsKernelArgs*>(data.kernelArgs); \
+			for (uint32_t i = 0, activeIdx = 0; i < fieldSamplerCount; ++i) \
+			{ \
+				const nvidia::fieldsampler::FieldSamplerWrapperGPU* wrapper = static_cast<const nvidia::fieldsampler::FieldSamplerWrapperGPU* >( (*data.fieldSamplerArray)[i].mFieldSamplerWrapper ); \
+				if (wrapper->isEnabled()) \
+				{ \
+					nvidia::fieldsampler::FieldSamplerParamsEx paramsEx; \
+					paramsEx.paramsHandle = _storage_.mappedHandle( wrapper->getParamsHandle() ); \
+					paramsEx.multiplier = (*data.fieldSamplerArray)[i].mMultiplier; \
+					nvidia::fieldsampler::FieldSamplerKernelMode::Enum kernelMode = (++activeIdx == data.activeFieldSamplerCount) ? data.kernelMode : nvidia::fieldsampler::FieldSamplerKernelMode::DEFAULT; \
+					ON_LAUNCH_FIELD_SAMPLER_KERNEL( wrapper->getInternalFieldSampler(), wrapper->getInternalFieldSamplerDesc() ); \
+					SCENE_CUDA_OBJ(this, fieldSamplerPointsKernel)( data.stream, data.threadCount, \
+						*static_cast<const nvidia::fieldsampler::FieldSamplerKernelParams*>(kernelArgs), \
+						nvidia::apex::createApexCudaMemRef(kernelArgs->accumField, size_t(data.memRefSize)), \
+						nvidia::apex::createApexCudaMemRef(kernelArgs->accumVelocity, size_t(data.memRefSize)), \
+						nvidia::apex::createApexCudaMemRef(kernelArgs->positionMass, size_t(data.memRefSize), ApexCudaMemFlags::IN), \
+						nvidia::apex::createApexCudaMemRef(kernelArgs->velocity, size_t(data.memRefSize), ApexCudaMemFlags::IN), \
+						paramsEx, queryParamsHandle, kernelMode ); \
+				} \
+			} \
+		} \
+		return true; \
+	case nvidia::fieldsampler::FieldSamplerKernelType::GRID: \
+		{ \
+			const nvidia::fieldsampler::FieldSamplerGridKernelLaunchDataIntl& data = static_cast<const nvidia::fieldsampler::FieldSamplerGridKernelLaunchDataIntl&>(launchData); \
+			const nvidia::fieldsampler::FieldSamplerGridKernelArgs* kernelArgs = static_cast<const nvidia::fieldsampler::FieldSamplerGridKernelArgs*>(data.kernelArgs); \
+			SCENE_CUDA_OBJ(this, surfRefFieldSamplerGridAccum).bindTo(*data.accumArray, ApexCudaMemFlags::IN_OUT); \
+			for (uint32_t i = 0, activeIdx = 0; i < fieldSamplerCount; ++i) \
+			{ \
+				const nvidia::fieldsampler::FieldSamplerWrapperGPU* wrapper = static_cast<const nvidia::fieldsampler::FieldSamplerWrapperGPU* >( (*data.fieldSamplerArray)[i].mFieldSamplerWrapper ); \
+				if (wrapper->isEnabled() && wrapper->getInternalFieldSamplerDesc().gridSupportType == FieldSamplerGridSupportTypeIntl::VELOCITY_PER_CELL) \
+				{ \
+					nvidia::fieldsampler::FieldSamplerParamsEx paramsEx; \
+					paramsEx.paramsHandle = _storage_.mappedHandle( wrapper->getParamsHandle() ); \
+					paramsEx.multiplier = (*data.fieldSamplerArray)[i].mMultiplier; \
+					nvidia::fieldsampler::FieldSamplerKernelMode::Enum kernelMode = (++activeIdx == data.activeFieldSamplerCount) ? data.kernelMode : nvidia::fieldsampler::FieldSamplerKernelMode::DEFAULT; \
+					ON_LAUNCH_FIELD_SAMPLER_KERNEL( wrapper->getInternalFieldSampler(), wrapper->getInternalFieldSamplerDesc() ); \
+					SCENE_CUDA_OBJ(this, fieldSamplerGridKernel)( data.stream, data.threadCountX, data.threadCountY, data.threadCountZ, \
+						*static_cast<const nvidia::fieldsampler::FieldSamplerKernelParams*>(kernelArgs), \
+						*static_cast<const nvidia::fieldsampler::FieldSamplerGridKernelParams*>(kernelArgs), \
+						paramsEx, queryParamsHandle, kernelMode ); \
+				} \
+			} \
+			SCENE_CUDA_OBJ(this, surfRefFieldSamplerGridAccum).unbind(); \
+		} \
+		return true; \
+	default: \
+		PX_ALWAYS_ASSERT(); \
+		return false; \
+	};
+#endif
+
+#else
+
+APEX_CUDA_BOUND_KERNEL(FIELD_SAMPLER_POINTS_KERNEL_CONFIG, fieldSamplerPointsKernel,
+                       ((nvidia::fieldsampler::FieldSamplerKernelParams, baseParams))
+                       ((APEX_MEM_BLOCK(float4), accumField))
+                       ((APEX_MEM_BLOCK(float4), accumVelocity))
+                       ((APEX_MEM_BLOCK(const float4), positionMass))
+                       ((APEX_MEM_BLOCK(const float4), velocity))
+                       ((InplaceHandle<nvidia::fieldsampler::FieldSamplerParamsExArray>, paramsExArrayHandle))
+                       ((InplaceHandle<nvidia::fieldsampler::FieldSamplerQueryParams>, queryParamsHandle))
+					   ((nvidia::fieldsampler::FieldSamplerKernelMode::Enum, kernelMode))
+                      )
+APEX_CUDA_FREE_KERNEL_3D(FIELD_SAMPLER_GRID_KERNEL_CONFIG, fieldSamplerGridKernel,
+                       ((nvidia::fieldsampler::FieldSamplerKernelParams, baseParams))
+                       ((nvidia::fieldsampler::FieldSamplerGridKernelParams, gridParams))
+                       ((InplaceHandle<nvidia::fieldsampler::FieldSamplerParamsExArray>, paramsExArrayHandle))
+                       ((InplaceHandle<nvidia::fieldsampler::FieldSamplerQueryParams>, queryParamsHandle))
+					   ((nvidia::fieldsampler::FieldSamplerKernelMode::Enum, kernelMode))
+                      )
+
+#ifndef __CUDACC__
+#define LAUNCH_FIELD_SAMPLER_KERNEL( launchData ) \
+	const ApexCudaConstStorage& _storage_ = *getFieldSamplerCudaConstStorage(); \
+	InplaceHandle<nvidia::fieldsampler::FieldSamplerParamsExArray> paramsExArrayHandle = _storage_.mappedHandle( launchData.paramsExArrayHandle ); \
+	InplaceHandle<nvidia::fieldsampler::FieldSamplerQueryParams> queryParamsHandle = _storage_.mappedHandle( launchData.queryParamsHandle ); \
+	switch( launchData.kernelType ) \
+	{ \
+	case nvidia::fieldsampler::FieldSamplerKernelType::POINTS: \
+		{ \
+			const nvidia::fieldsampler::FieldSamplerPointsKernelLaunchDataIntl& data = static_cast<const nvidia::fieldsampler::FieldSamplerPointsKernelLaunchDataIntl&>(launchData); \
+			const nvidia::fieldsampler::FieldSamplerPointsKernelArgs* kernelArgs = static_cast<const nvidia::fieldsampler::FieldSamplerPointsKernelArgs*>(data.kernelArgs); \
+			SCENE_CUDA_OBJ(this, fieldSamplerPointsKernel)( data.stream, data.threadCount, \
+				*static_cast<const nvidia::fieldsampler::FieldSamplerKernelParams*>(kernelArgs), \
+				nvidia::apex::createApexCudaMemRef(kernelArgs->accumField, size_t(data.memRefSize)), \
+				nvidia::apex::createApexCudaMemRef(kernelArgs->accumVelocity, size_t(data.memRefSize)), \
+				nvidia::apex::createApexCudaMemRef(kernelArgs->positionMass, size_t(data.memRefSize), ApexCudaMemFlags::IN), \
+				nvidia::apex::createApexCudaMemRef(kernelArgs->velocity, size_t(data.memRefSize), ApexCudaMemFlags::IN), \
+				paramsExArrayHandle, queryParamsHandle, data.kernelMode ); \
+		} \
+		return true; \
+	case nvidia::fieldsampler::FieldSamplerKernelType::GRID: \
+		{ \
+			const nvidia::fieldsampler::FieldSamplerGridKernelLaunchDataIntl& data = static_cast<const nvidia::fieldsampler::FieldSamplerGridKernelLaunchDataIntl&>(launchData); \
+			const nvidia::fieldsampler::FieldSamplerGridKernelArgs* kernelArgs = static_cast<const nvidia::fieldsampler::FieldSamplerGridKernelArgs*>(data.kernelArgs); \
+			SCENE_CUDA_OBJ(this, surfRefFieldSamplerGridAccum).bindTo(*data.accumArray, ApexCudaMemFlags::IN_OUT); \
+			SCENE_CUDA_OBJ(this, fieldSamplerGridKernel)( data.stream, data.threadCountX, data.threadCountY, data.threadCountZ, \
+				*static_cast<const nvidia::fieldsampler::FieldSamplerKernelParams*>(kernelArgs), \
+				*static_cast<const nvidia::fieldsampler::FieldSamplerGridKernelParams*>(kernelArgs), \
+				paramsExArrayHandle, queryParamsHandle, data.kernelMode ); \
+			SCENE_CUDA_OBJ(this, surfRefFieldSamplerGridAccum).unbind(); \
+		} \
+		return true; \
+	default: \
+		PX_ALWAYS_ASSERT(); \
+		return false; \
+	};
+#endif
+
+#endif //FIELD_SAMPLER_SEPARATE_KERNELS
diff --git a/APEX_1.4/module/fieldsampler/include/FieldBoundaryWrapper.h b/APEX_1.4/module/fieldsampler/include/FieldBoundaryWrapper.h
new file mode 100644
index 00000000..fbaf10f5
--- /dev/null
+++ b/APEX_1.4/module/fieldsampler/include/FieldBoundaryWrapper.h
@@ -0,0 +1,85 @@
+/*
+ * Copyright (c) 2008-2015, 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 __FIELD_BOUNDARY_WRAPPER_H__
+#define __FIELD_BOUNDARY_WRAPPER_H__
+
+#include "Apex.h"
+#include "ApexSDKHelpers.h"
+#include "ApexActor.h"
+#include "FieldBoundaryIntl.h"
+
+#if APEX_CUDA_SUPPORT
+#include "ApexCudaWrapper.h"
+#endif
+
+#include "FieldSamplerCommon.h"
+
+
+namespace nvidia
+{
+namespace fieldsampler
+{
+
+class FieldSamplerManager;
+
+
+class FieldBoundaryWrapper : public ApexResourceInterface, public ApexResource
+{
+public:
+	// ApexResourceInterface methods
+	void			release();
+	void			setListIndex(ResourceList& list, uint32_t index)
+	{
+		m_listIndex = index;
+		m_list = &list;
+	}
+	uint32_t	getListIndex() const
+	{
+		return m_listIndex;
+	}
+
+	FieldBoundaryWrapper(ResourceList& list, FieldSamplerManager* manager, FieldBoundaryIntl* fieldBoundary, const FieldBoundaryDescIntl& fieldBoundaryDesc);
+
+	FieldBoundaryIntl* getInternalFieldBoundary() const
+	{
+		return mFieldBoundary;
+	}
+	PX_INLINE const FieldBoundaryDescIntl& getInternalFieldBoundaryDesc() const
+	{
+		return mFieldBoundaryDesc;
+	}
+
+	void update();
+
+	const nvidia::Array<FieldShapeDescIntl>&	getFieldShapes() const
+	{
+		return mFieldShapes;
+	}
+	bool									getFieldShapesChanged() const
+	{
+		return mFieldShapesChanged;
+	}
+
+protected:
+	FieldSamplerManager*			mManager;
+
+	FieldBoundaryIntl*				mFieldBoundary;
+	FieldBoundaryDescIntl				mFieldBoundaryDesc;
+
+	nvidia::Array<FieldShapeDescIntl>	mFieldShapes;
+	bool							mFieldShapesChanged;
+};
+
+}
+} // end namespace nvidia::apex
+
+#endif
diff --git a/APEX_1.4/module/fieldsampler/include/FieldSamplerCommon.h b/APEX_1.4/module/fieldsampler/include/FieldSamplerCommon.h
new file mode 100644
index 00000000..8ef0b770
--- /dev/null
+++ b/APEX_1.4/module/fieldsampler/include/FieldSamplerCommon.h
@@ -0,0 +1,386 @@
+/*
+ * Copyright (c) 2008-2015, 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 __FIELD_SAMPLER_COMMON_H__
+#define __FIELD_SAMPLER_COMMON_H__
+
+#include "PxVec3.h"
+#include "PxVec4.h"
+#include <PxMat44.h>
+
+#include <FieldSamplerIntl.h>
+#include <FieldBoundaryIntl.h>
+
+#include "PsArray.h"
+
+#if APEX_CUDA_SUPPORT
+#pragma warning(push)
+#pragma warning(disable:4201)
+#pragma warning(disable:4408)
+
+#include <vector_types.h>
+
+#pragma warning(pop)
+#endif
+
+#define FIELD_SAMPLER_MULTIPLIER_VALUE 1
+#define FIELD_SAMPLER_MULTIPLIER_WEIGHT 2
+//0, FIELD_SAMPLER_MULTIPLIER_VALUE or FIELD_SAMPLER_MULTIPLIER_WEIGHT
+#define FIELD_SAMPLER_MULTIPLIER FIELD_SAMPLER_MULTIPLIER_WEIGHT
+
+namespace nvidia
+{
+namespace apex
+{
+	class ApexCudaArray;
+}
+
+namespace fieldsampler
+{
+
+#define VELOCITY_WEIGHT_THRESHOLD 0.00001f
+
+struct FieldSamplerExecuteArgs
+{
+	physx::PxVec3			position;
+	float			mass;
+	physx::PxVec3			velocity;
+
+	float			elapsedTime;
+	uint32_t			totalElapsedMS;
+};
+
+//struct FieldShapeParams
+#define INPLACE_TYPE_STRUCT_NAME FieldShapeParams
+#define INPLACE_TYPE_STRUCT_BASE FieldShapeDescIntl
+#define INPLACE_TYPE_STRUCT_FIELDS \
+	INPLACE_TYPE_FIELD(float,	fade)
+#include INPLACE_TYPE_BUILD()
+
+
+//struct FieldShapeGroupParams
+#define INPLACE_TYPE_STRUCT_NAME FieldShapeGroupParams
+#define INPLACE_TYPE_STRUCT_FIELDS \
+	INPLACE_TYPE_FIELD(InplaceArray<FieldShapeParams>,	shapeArray)
+#include INPLACE_TYPE_BUILD()
+
+
+//struct FieldSamplerParams
+#define INPLACE_TYPE_STRUCT_NAME FieldSamplerParams
+#define INPLACE_TYPE_STRUCT_FIELDS \
+	INPLACE_TYPE_FIELD(uint32_t,										executeType) \
+	INPLACE_TYPE_FIELD(InplaceHandleBase,									executeParamsHandle) \
+	INPLACE_TYPE_FIELD(InplaceEnum<FieldSamplerTypeIntl::Enum>,				type) \
+	INPLACE_TYPE_FIELD(InplaceEnum<FieldSamplerGridSupportTypeIntl::Enum>,	gridSupportType) \
+	INPLACE_TYPE_FIELD(float,										dragCoeff) \
+	INPLACE_TYPE_FIELD(FieldShapeParams,									includeShape) \
+	INPLACE_TYPE_FIELD(InplaceArray<InplaceHandle<FieldShapeGroupParams> >,	excludeShapeGroupHandleArray)
+#include INPLACE_TYPE_BUILD()
+
+
+//struct FieldSamplerQueryParams
+#define INPLACE_TYPE_STRUCT_NAME FieldSamplerQueryParams
+#include INPLACE_TYPE_BUILD()
+
+
+//struct FieldSamplerParams
+#define INPLACE_TYPE_STRUCT_NAME FieldSamplerParamsEx
+#define INPLACE_TYPE_STRUCT_FIELDS \
+	INPLACE_TYPE_FIELD(InplaceHandle<FieldSamplerParams>,	paramsHandle) \
+	INPLACE_TYPE_FIELD(float,						multiplier)
+#include INPLACE_TYPE_BUILD()
+
+typedef InplaceArray< FieldSamplerParamsEx, false > FieldSamplerParamsExArray;
+
+
+#if APEX_CUDA_SUPPORT || defined(__CUDACC__)
+
+struct FieldSamplerKernelType
+{
+	enum Enum
+	{
+		POINTS,
+		GRID
+	};
+};
+
+struct FieldSamplerKernelParams
+{
+	float	elapsedTime;
+	physx::PxVec3	cellSize;
+	uint32_t	totalElapsedMS;
+};
+
+struct FieldSamplerGridKernelParams
+{
+	uint32_t numX, numY, numZ;
+
+	PxMat44 gridToWorld;
+
+	float mass;
+};
+
+
+struct FieldSamplerKernelArgs : FieldSamplerKernelParams
+{
+};
+
+struct FieldSamplerPointsKernelArgs : FieldSamplerKernelArgs
+{
+	float4* accumField;
+	float4* accumVelocity;
+	const float4* positionMass;
+	const float4* velocity;
+};
+
+struct FieldSamplerGridKernelArgs : FieldSamplerKernelArgs, FieldSamplerGridKernelParams
+{
+};
+
+struct FieldSamplerKernelMode
+{
+	enum Enum
+	{
+		DEFAULT = 0,
+		FINISH_PRIMARY = 1,
+		FINISH_SECONDARY = 2
+	};
+};
+
+#endif
+
+class FieldSamplerWrapper;
+
+struct FieldSamplerInfo
+{
+	FieldSamplerWrapper*	mFieldSamplerWrapper;
+	float			mMultiplier;
+};
+
+#if APEX_CUDA_SUPPORT
+
+struct FieldSamplerKernelLaunchDataIntl
+{
+	CUstream                                        stream;
+	FieldSamplerKernelType::Enum                    kernelType;
+	const FieldSamplerKernelArgs*                   kernelArgs;
+	InplaceHandle<FieldSamplerQueryParams>          queryParamsHandle;
+	InplaceHandle<FieldSamplerParamsExArray>        paramsExArrayHandle;
+	const nvidia::Array<FieldSamplerInfo>*           fieldSamplerArray;
+	uint32_t                                    activeFieldSamplerCount;
+	FieldSamplerKernelMode::Enum                    kernelMode;
+};
+
+struct FieldSamplerPointsKernelLaunchDataIntl : FieldSamplerKernelLaunchDataIntl
+{
+	uint32_t                                    threadCount;
+	uint32_t                                    memRefSize;
+};
+
+struct FieldSamplerGridKernelLaunchDataIntl : FieldSamplerKernelLaunchDataIntl
+{
+	uint32_t									threadCountX;
+	uint32_t									threadCountY;
+	uint32_t									threadCountZ;
+	ApexCudaArray*								accumArray;
+};
+
+#endif
+
+APEX_CUDA_CALLABLE PX_INLINE float evalFade(float dist, float fade)
+{
+	float x = (1 - dist) / (fade + 1e-5f);
+	return physx::PxClamp<float>(x, 0, 1);
+}
+
+APEX_CUDA_CALLABLE PX_INLINE float evalFadeAntialiasing(float dist, float fade, float cellRadius)
+{
+	const float f = fade;
+	const float r = cellRadius;
+	const float x = dist - 1.0f;
+
+	float res = 0.0f;
+	//linear part
+	//if (x - r < -f)
+	{
+		const float a = PxMin(x - r, -f);
+		const float b = PxMin(x + r, -f);
+
+		res += (b - a);
+	}
+	//quadratic part
+	if (f >= 1e-5f)
+	{
+		//if (x - r < 0.0f && x + r > -f)
+		{
+			const float a = physx::PxClamp(x - r, -f, 0.0f);
+			const float b = physx::PxClamp(x + r, -f, 0.0f);
+
+			res += (a*a - b*b) / (2 * f);
+		}
+	}
+	return res / (2 * r);
+}
+
+
+APEX_CUDA_CALLABLE PX_INLINE float evalDistInShapeNONE(const FieldShapeDescIntl& /*shapeParams*/, const physx::PxVec3& /*worldPos*/)
+{
+	return 0.0f; //always inside
+}
+
+APEX_CUDA_CALLABLE PX_INLINE float evalDistInShapeSPHERE(const FieldShapeDescIntl& shapeParams, const physx::PxVec3& worldPos)
+{
+	const physx::PxVec3 shapePos = shapeParams.worldToShape.transform(worldPos);
+	const float radius = shapeParams.dimensions.x;
+	return shapePos.magnitude() / radius;
+}
+
+APEX_CUDA_CALLABLE PX_INLINE float evalDistInShapeBOX(const FieldShapeDescIntl& shapeParams, const physx::PxVec3& worldPos)
+{
+	const physx::PxVec3 shapePos = shapeParams.worldToShape.transform(worldPos);
+	const physx::PxVec3& halfSize = shapeParams.dimensions;
+	physx::PxVec3 unitPos(shapePos.x / halfSize.x, shapePos.y / halfSize.y, shapePos.z / halfSize.z);
+	return physx::PxVec3(physx::PxAbs(unitPos.x), physx::PxAbs(unitPos.y), physx::PxAbs(unitPos.z)).maxElement();
+}
+
+APEX_CUDA_CALLABLE PX_INLINE float evalDistInShapeCAPSULE(const FieldShapeDescIntl& shapeParams, const physx::PxVec3& worldPos)
+{
+	const physx::PxVec3 shapePos = shapeParams.worldToShape.transform(worldPos);
+	const float radius = shapeParams.dimensions.x;
+	const float halfHeight = shapeParams.dimensions.y * 0.5f;
+
+	physx::PxVec3 clampPos = shapePos;
+	clampPos.y -= physx::PxClamp(shapePos.y, -halfHeight, +halfHeight);
+
+	return clampPos.magnitude() / radius;
+}
+
+APEX_CUDA_CALLABLE PX_INLINE float evalDistInShape(const FieldShapeDescIntl& shapeParams, const physx::PxVec3& worldPos)
+{
+	switch (shapeParams.type)
+	{
+	case FieldShapeTypeIntl::NONE:
+		return evalDistInShapeNONE(shapeParams, worldPos);
+	case FieldShapeTypeIntl::SPHERE:
+		return evalDistInShapeSPHERE(shapeParams, worldPos);
+	case FieldShapeTypeIntl::BOX:
+		return evalDistInShapeBOX(shapeParams, worldPos);
+	case FieldShapeTypeIntl::CAPSULE:
+		return evalDistInShapeCAPSULE(shapeParams, worldPos);
+	default:
+		return 1.0f; //always outside
+	};
+}
+
+APEX_CUDA_CALLABLE PX_INLINE physx::PxVec3 scaleToShape(const FieldShapeDescIntl& shapeParams, const physx::PxVec3& worldVec)
+{
+	switch (shapeParams.type)
+	{
+	case FieldShapeTypeIntl::SPHERE:
+	case FieldShapeTypeIntl::CAPSULE:
+	{
+		const float radius = shapeParams.dimensions.x;
+		return physx::PxVec3(worldVec.x / radius, worldVec.y / radius, worldVec.z / radius);
+	}
+	case FieldShapeTypeIntl::BOX:
+	{
+		const physx::PxVec3& halfSize = shapeParams.dimensions;
+		return physx::PxVec3(worldVec.x / halfSize.x, worldVec.y / halfSize.y, worldVec.z / halfSize.z);
+	}
+	default:
+		return worldVec;
+	};
+}
+
+
+APEX_CUDA_CALLABLE PX_INLINE float evalWeightInShape(const FieldShapeParams& shapeParams, const physx::PxVec3& position)
+{
+	float dist = nvidia::fieldsampler::evalDistInShape(shapeParams, position);
+	return nvidia::fieldsampler::evalFade(dist, shapeParams.fade) * shapeParams.weight;
+}
+
+APEX_CUDA_CALLABLE PX_INLINE void accumFORCE(const FieldSamplerExecuteArgs& args,
+	const physx::PxVec3& field, float fieldW,
+	physx::PxVec4& accumAccel, physx::PxVec4& accumVelocity)
+{
+	PX_UNUSED(accumVelocity);
+
+	physx::PxVec3 newAccel = ((1 - accumAccel.w) * fieldW * args.elapsedTime / args.mass) * field;
+	accumAccel.x += newAccel.x;
+	accumAccel.y += newAccel.y;
+	accumAccel.z += newAccel.z;
+}
+
+APEX_CUDA_CALLABLE PX_INLINE void accumACCELERATION(const FieldSamplerExecuteArgs& args,
+	const physx::PxVec3& field, float fieldW,
+	physx::PxVec4& accumAccel, physx::PxVec4& accumVelocity)
+{
+	PX_UNUSED(accumVelocity);
+
+	physx::PxVec3 newAccel = ((1 - accumAccel.w) * fieldW * args.elapsedTime) * field;
+	accumAccel.x += newAccel.x;
+	accumAccel.y += newAccel.y;
+	accumAccel.z += newAccel.z;
+}
+
+APEX_CUDA_CALLABLE PX_INLINE void accumVELOCITY_DIRECT(const FieldSamplerExecuteArgs& args,
+	const physx::PxVec3& field, float fieldW,
+	physx::PxVec4& accumAccel, physx::PxVec4& accumVelocity)
+{
+	PX_UNUSED(args);
+
+	physx::PxVec3 newVelocity = ((1 - accumAccel.w) * fieldW) * field;
+	accumVelocity.x += newVelocity.x;
+	accumVelocity.y += newVelocity.y;
+	accumVelocity.z += newVelocity.z;
+	accumVelocity.w = PxMax(accumVelocity.w, fieldW);
+}
+
+APEX_CUDA_CALLABLE PX_INLINE void accumVELOCITY_DRAG(const FieldSamplerExecuteArgs& args, float dragCoeff,
+	const physx::PxVec3& field, float fieldW,
+	physx::PxVec4& accumAccel, physx::PxVec4& accumVelocity)
+{
+#if 1
+	const float dragFieldW = PxMin(fieldW * dragCoeff * args.elapsedTime / args.mass, 1.0f);
+	accumVELOCITY_DIRECT(args, field, dragFieldW, accumAccel, accumVelocity);
+#else
+	const physx::PxVec3 dragForce = (field - args.velocity) * dragCoeff;
+	accumFORCE(args, dragForce, fieldW, accumAccel, accumVelocity);
+#endif
+}
+
+
+} // namespace apex
+}
+#ifdef __CUDACC__
+
+template <int queryType>
+struct FieldSamplerExecutor;
+/*
+{
+	INPLACE_TEMPL_ARGS_DEF
+	static inline __device__ physx::PxVec3 func(const nvidia::fieldsampler::FieldSamplerParams* params, const nvidia::fieldsampler::FieldSamplerExecuteArgs& args, float& fieldWeight);
+};
+*/
+
+template <int queryType>
+struct FieldSamplerIncludeWeightEvaluator
+{
+	INPLACE_TEMPL_ARGS_DEF
+	static inline __device__ float func(const nvidia::fieldsampler::FieldSamplerParams* params, const physx::PxVec3& position, const physx::PxVec3& cellSize)
+	{
+		return nvidia::fieldsampler::evalWeightInShape(params->includeShape, position);
+	}
+};
+
+#endif
+
+#endif
diff --git a/APEX_1.4/module/fieldsampler/include/FieldSamplerManager.h b/APEX_1.4/module/fieldsampler/include/FieldSamplerManager.h
new file mode 100644
index 00000000..887a50c7
--- /dev/null
+++ b/APEX_1.4/module/fieldsampler/include/FieldSamplerManager.h
@@ -0,0 +1,133 @@
+/*
+ * Copyright (c) 2008-2015, 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 __FIELD_SAMPLER_MANAGER_H__
+#define __FIELD_SAMPLER_MANAGER_H__
+
+#include "FieldSamplerManagerIntl.h"
+#include "FieldSamplerScene.h"
+#include "ApexGroupsFiltering.h"
+
+#if APEX_CUDA_SUPPORT
+#define CUDA_OBJ(name) SCENE_CUDA_OBJ(mManager->getScene(), name)
+#endif
+
+namespace nvidia
+{
+namespace fieldsampler
+{
+
+PX_INLINE bool operator != (const PxFilterData& d1, const PxFilterData& d2)
+{
+	return d1.word0 != d2.word0 || d1.word1 != d2.word1 || d1.word2 != d2.word2 || d1.word3 != d2.word3;
+}
+
+class FieldSamplerScene;
+
+class FieldSamplerQuery;
+class FieldSamplerSceneWrapper;
+class FieldSamplerWrapper;
+class FieldBoundaryWrapper;
+
+class FieldSamplerManager : public FieldSamplerManagerIntl,  public UserAllocated
+{
+public:
+	FieldSamplerManager(FieldSamplerScene* scene);
+	virtual ~FieldSamplerManager() {}
+
+	void		submitTasks();
+	void		setTaskDependencies();
+	void		fetchResults();
+
+	/* FieldSamplerManager */
+	FieldSamplerQueryIntl* createFieldSamplerQuery(const FieldSamplerQueryDescIntl&);
+
+	void registerFieldSampler(FieldSamplerIntl* , const FieldSamplerDescIntl& , FieldSamplerSceneIntl*);
+	void unregisterFieldSampler(FieldSamplerIntl*);
+
+	void registerFieldBoundary(FieldBoundaryIntl* , const FieldBoundaryDescIntl&);
+	void unregisterFieldBoundary(FieldBoundaryIntl*);
+
+	void registerUnhandledParticleSystem(physx::PxActor*);
+	void unregisterUnhandledParticleSystem(physx::PxActor*);
+	bool isUnhandledParticleSystem(physx::PxActor*);
+
+	FieldSamplerScene*	getScene() const
+	{
+		return mScene;
+	}
+	SceneIntl&		getApexScene() const
+	{
+		return *mScene->mApexScene;
+	}
+
+	bool getFieldSamplerGroupsFiltering(const PxFilterData &o1,const PxFilterData &o2,float &weight) const;
+	void setFieldSamplerWeightedCollisionFilterCallback(FieldSamplerWeightedCollisionFilterCallback *callback)
+	{
+		mWeightedCallback = callback;
+	}
+
+	virtual bool doFieldSamplerFiltering(const PxFilterData &o1, const PxFilterData &o2, float &weight) const
+	{
+		return getFieldSamplerGroupsFiltering(o1, o2, weight);
+	}
+
+protected:
+	virtual FieldSamplerQuery* allocateFieldSamplerQuery(const FieldSamplerQueryDescIntl&) = 0;
+	virtual FieldSamplerSceneWrapper* allocateFieldSamplerSceneWrapper(FieldSamplerSceneIntl*) = 0;
+	virtual FieldSamplerWrapper* allocateFieldSamplerWrapper(FieldSamplerIntl* , const FieldSamplerDescIntl& , FieldSamplerSceneWrapper*) = 0;
+
+	static PX_INLINE void addFieldSamplerToQuery(FieldSamplerWrapper* fieldSamplerWrapper, FieldSamplerQuery* query);
+	void addAllFieldSamplersToQuery(FieldSamplerQuery*) const;
+
+	FieldSamplerScene*	mScene;
+
+	ResourceList		mFieldSamplerQueryList;
+	ResourceList		mFieldSamplerSceneWrapperList;
+	ResourceList		mFieldSamplerWrapperList;
+	ResourceList		mFieldBoundaryWrapperList;
+
+	nvidia::Array<physx::PxActor*>		mUnhandledParticleSystems;
+	FieldSamplerWeightedCollisionFilterCallback *mWeightedCallback;
+
+	bool								mFieldSamplerGroupsFilteringChanged;
+
+	friend class FieldSamplerSceneWrapperGPU;
+};
+
+class FieldSamplerManagerCPU : public FieldSamplerManager
+{
+public:
+	FieldSamplerManagerCPU(FieldSamplerScene* scene) : FieldSamplerManager(scene) {}
+
+protected:
+	FieldSamplerQuery* allocateFieldSamplerQuery(const FieldSamplerQueryDescIntl&);
+	FieldSamplerSceneWrapper* allocateFieldSamplerSceneWrapper(FieldSamplerSceneIntl*);
+	FieldSamplerWrapper* allocateFieldSamplerWrapper(FieldSamplerIntl* , const FieldSamplerDescIntl& , FieldSamplerSceneWrapper*);
+};
+
+#if APEX_CUDA_SUPPORT
+class FieldSamplerManagerGPU : public FieldSamplerManager
+{
+public:
+	FieldSamplerManagerGPU(FieldSamplerScene* scene) : FieldSamplerManager(scene) {}
+
+protected:
+	FieldSamplerQuery* allocateFieldSamplerQuery(const FieldSamplerQueryDescIntl&);
+	FieldSamplerSceneWrapper* allocateFieldSamplerSceneWrapper(FieldSamplerSceneIntl*);
+	FieldSamplerWrapper* allocateFieldSamplerWrapper(FieldSamplerIntl* , const FieldSamplerDescIntl& , FieldSamplerSceneWrapper*);
+};
+#endif // APEX_CUDA_SUPPORT
+
+
+} // end namespace nvidia::apex
+}
+#endif // __FIELD_SAMPLER_MANAGER_H__
diff --git a/APEX_1.4/module/fieldsampler/include/FieldSamplerPhysXMonitor.h b/APEX_1.4/module/fieldsampler/include/FieldSamplerPhysXMonitor.h
new file mode 100644
index 00000000..30ff79c4
--- /dev/null
+++ b/APEX_1.4/module/fieldsampler/include/FieldSamplerPhysXMonitor.h
@@ -0,0 +1,171 @@
+/*
+ * Copyright (c) 2008-2015, 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 __FIELD__SAMPLER_PHYSX_MONITOR_H___
+#define __FIELD__SAMPLER_PHYSX_MONITOR_H___
+
+#include "Apex.h"
+#include <PsArray.h>
+
+#include "FieldSamplerSceneIntl.h"
+#include "FieldSamplerPhysXMonitorParams.h"
+
+#if PX_PHYSICS_VERSION_MAJOR == 3
+#include <PxScene.h>
+#endif
+
+namespace nvidia
+{
+namespace apex
+{
+class FieldSamplerQueryIntl;
+}
+
+namespace fieldsampler
+{
+
+class FieldSamplerScene;
+class FieldSamplerManager;
+
+
+struct ShapeData : public UserAllocated
+{
+	uint32_t	fdIndex;
+	uint32_t	rbIndex;
+	float	mass;
+	PxVec3	pos;
+	PxVec3	vel;
+
+	static bool sortPredicate (ShapeData* sd1, ShapeData* sd2)
+	{
+		if (sd1 == 0) return false;
+		if (sd2 == 0) return true;
+		return sd1->fdIndex < sd2->fdIndex;
+	}
+};
+
+
+class FieldSamplerPhysXMonitor : public UserAllocated
+{
+private:
+
+
+public:
+	FieldSamplerPhysXMonitor(FieldSamplerScene& scene);
+	virtual ~FieldSamplerPhysXMonitor();
+
+#if PX_PHYSICS_VERSION_MAJOR == 3
+	virtual void	update();
+	virtual void	updatePhysX();
+
+	/* PhysX scene management */
+	void	setPhysXScene(PxScene* scene);
+	PX_INLINE PxScene*	getPhysXScene() const
+	{
+		return mScene;
+	}
+
+	/* Toggle PhysX Monitor on/off */
+	PX_INLINE void enablePhysXMonitor(bool enable)
+	{
+		mEnable = enable;
+	}
+
+	/* Is PhysX Monitor enabled */
+	PX_INLINE bool isEnable()
+	{
+		return mEnable;
+	}
+
+	PX_INLINE void FieldSamplerPhysXMonitor::setPhysXFilterData(physx::PxFilterData filterData)
+	{
+		mFilterData = filterData;
+	}
+
+private:
+	FieldSamplerPhysXMonitor& operator=(const FieldSamplerPhysXMonitor&);
+
+	void getParticles(uint32_t taskId);
+	void updateParticles();
+	void getRigidBodies(uint32_t taskId);
+	void updateRigidBodies();
+	//void getCloth(PxTask& task, bool isDataOnDevice);
+	//void updateCloth();
+
+protected:
+
+	void commonInitArray();
+
+	FieldSamplerScene*				mFieldSamplerScene;
+	FieldSamplerManager*			mFieldSamplerManager;
+
+	PxScene*						mScene;
+	
+	FieldSamplerPhysXMonitorParams*	mParams;
+	
+	PxFilterData					mFilterData;
+
+	//Particles
+	uint32_t					mNumPS;  //Number of particle systems
+	uint32_t					mPCount; //Number of particles in buffer
+	Array<PxActor*>					mParticleSystems;
+	Array<float>					mPSMass;
+	Array<PxVec4>					mPSOutField;
+	Array<PxVec3>					mOutVelocities;
+	Array<uint32_t>					mOutIndices;
+	Array<physx::PxParticleReadData*>		mParticleReadData;
+	Array<FieldSamplerQueryIntl*>		mPSFieldSamplerQuery;	
+	Array<PxTaskID>			mPSFieldSamplerTaskID;
+
+
+	//Rigid bodies
+	uint32_t					mNumRB; //Number of rigid bodies
+	Array<PxActor*>					mRBActors;
+	Array<ShapeData*>				mRBIndex;
+	Array<PxVec4>					mRBInPosition;
+	Array<PxVec4>					mRBInVelocity;
+	Array<PxVec4>					mRBOutField;
+	Array<PxFilterData>				mRBFilterData;
+	Array<FieldSamplerQueryIntl*>		mRBFieldSamplerQuery;	
+
+	//Enable or disable PhysX Monitor
+	bool mEnable;
+
+public:
+	class RunAfterActorUpdateTask : public PxTask
+	{
+	public:
+		RunAfterActorUpdateTask(FieldSamplerPhysXMonitor& owner) : mOwner(owner) {}
+		const char* getName() const
+		{
+			return FSST_PHYSX_MONITOR_UPDATE;
+		}
+		void run()
+		{
+			mOwner.updatePhysX();
+		}
+
+	protected:
+		FieldSamplerPhysXMonitor& mOwner;
+
+	private:
+		RunAfterActorUpdateTask operator=(const RunAfterActorUpdateTask&);
+	};
+	RunAfterActorUpdateTask				mTaskRunAfterActorUpdate;
+#endif
+};
+
+
+}
+} // end namespace nvidia::apex
+
+#endif
+
diff --git a/APEX_1.4/module/fieldsampler/include/FieldSamplerQuery.h b/APEX_1.4/module/fieldsampler/include/FieldSamplerQuery.h
new file mode 100644
index 00000000..65d2ba58
--- /dev/null
+++ b/APEX_1.4/module/fieldsampler/include/FieldSamplerQuery.h
@@ -0,0 +1,308 @@
+/*
+ * Copyright (c) 2008-2015, 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 __FIELD_SAMPLER_QUERY_H__
+#define __FIELD_SAMPLER_QUERY_H__
+
+#include "Apex.h"
+#include "ApexSDKHelpers.h"
+#include "FieldSamplerQueryIntl.h"
+#include "FieldSamplerIntl.h"
+
+#if APEX_CUDA_SUPPORT
+#include "ApexCudaWrapper.h"
+#endif
+
+#include "FieldSamplerCommon.h"
+
+namespace nvidia
+{
+namespace fieldsampler
+{
+
+class FieldSamplerManager;
+class FieldSamplerSceneWrapper;
+class FieldSamplerWrapper;
+
+class FieldSamplerQuery : public FieldSamplerQueryIntl, public ApexResourceInterface, public ApexResource
+{
+protected:
+
+	class SceneInfo : public ApexResourceInterface, public ApexResource
+	{
+	protected:
+		FieldSamplerQuery* mQuery;
+		FieldSamplerSceneWrapper* mSceneWrapper;
+		nvidia::Array<FieldSamplerInfo> mFieldSamplerArray;
+		bool mFieldSamplerArrayChanged;
+
+		uint32_t mEnabledFieldSamplerCount;
+
+		SceneInfo(ResourceList& list, FieldSamplerQuery* query, FieldSamplerSceneWrapper* sceneWrapper)
+			: mQuery(query), mSceneWrapper(sceneWrapper), mFieldSamplerArrayChanged(false), mEnabledFieldSamplerCount(0)
+		{
+			list.add(*this);
+		}
+
+	public:
+		FieldSamplerSceneWrapper* getSceneWrapper() const
+		{
+			return mSceneWrapper;
+		}
+
+		void addFieldSampler(FieldSamplerWrapper* fieldSamplerWrapper, float multiplier)
+		{
+			FieldSamplerInfo fsInfo;
+			fsInfo.mFieldSamplerWrapper = fieldSamplerWrapper;
+			fsInfo.mMultiplier = multiplier;
+			mFieldSamplerArray.pushBack(fsInfo);
+			mFieldSamplerArrayChanged = true;
+		}
+		bool removeFieldSampler(FieldSamplerWrapper* fieldSamplerWrapper)
+		{
+			const uint32_t size = mFieldSamplerArray.size();
+			for (uint32_t index = 0; index < size; ++index)
+			{
+				if (mFieldSamplerArray[index].mFieldSamplerWrapper == fieldSamplerWrapper)
+				{
+					mFieldSamplerArray.replaceWithLast(index);
+					mFieldSamplerArrayChanged = true;
+					return true;
+				}
+			}
+			return false;
+		}
+		void clearAllFieldSamplers()
+		{
+			mFieldSamplerArray.clear();
+			mFieldSamplerArrayChanged = true;
+		}
+		uint32_t getEnabledFieldSamplerCount() const
+		{
+			return mEnabledFieldSamplerCount;
+		}
+		const nvidia::Array<FieldSamplerInfo>& getFieldSamplerArray() const
+		{
+			return mFieldSamplerArray;
+		}
+
+		virtual bool update();
+
+		// ApexResourceInterface methods
+		void						release()
+		{
+			delete this;
+		}
+		void						setListIndex(ResourceList& list, uint32_t index)
+		{
+			m_listIndex = index;
+			m_list = &list;
+		}
+		uint32_t				getListIndex() const
+		{
+			return m_listIndex;
+		}
+	};
+
+	SceneInfo* findSceneInfo(FieldSamplerSceneWrapper* sceneWrapper) const;
+
+	virtual SceneInfo* createSceneInfo(FieldSamplerSceneWrapper* sceneWrapper) = 0;
+
+	FieldSamplerQuery(const FieldSamplerQueryDescIntl& desc, ResourceList& list, FieldSamplerManager* manager);
+
+public:
+	virtual void update();
+	
+	PX_INLINE void setOnStartCallback(FieldSamplerCallbackIntl* callback)
+	{
+		if (mOnStartCallback) 
+		{
+			PX_DELETE(mOnStartCallback);
+		}
+		mOnStartCallback = callback;
+	}
+	PX_INLINE void setOnFinishCallback(FieldSamplerCallbackIntl* callback)
+	{
+		if (mOnFinishCallback) 
+		{
+			PX_DELETE(mOnFinishCallback);
+		}
+		mOnFinishCallback = callback;
+	}
+
+	virtual void		submitTasks() {}
+	virtual void		setTaskDependencies() {}
+	virtual void		fetchResults() {}
+
+	void	release();
+	void	destroy();
+
+	// ApexResourceInterface methods
+	void						setListIndex(ResourceList& list, uint32_t index)
+	{
+		m_listIndex = index;
+		m_list = &list;
+	}
+	uint32_t				getListIndex() const
+	{
+		return m_listIndex;
+	}
+
+	void		setPhysXScene(PxScene*)	{}
+	PxScene*	getPhysXScene() const
+	{
+		return NULL;
+	}
+
+	bool addFieldSampler(FieldSamplerWrapper*);
+	bool removeFieldSampler(FieldSamplerWrapper*);
+	void clearAllFieldSamplers();
+
+	const FieldSamplerQueryDescIntl& getQueryDesc() const
+	{
+		return mQueryDesc;
+	}
+
+protected:
+	void submitFieldSamplerQuery(const FieldSamplerQueryDataIntl& data, PxTask*, PxTask* );
+
+	FieldSamplerManager* 				mManager;
+
+	FieldSamplerQueryDescIntl				mQueryDesc;
+	FieldSamplerQueryDataIntl				mQueryData;
+
+	ResourceList						mSceneList;
+
+	nvidia::Array<SceneInfo*>			mPrimarySceneList;
+	nvidia::Array<SceneInfo*>			mSecondarySceneList;
+
+	ApexMirroredArray<PxVec4>	mAccumVelocity;
+
+	FieldSamplerCallbackIntl*				mOnStartCallback;
+	FieldSamplerCallbackIntl*				mOnFinishCallback;
+
+	friend class FieldSamplerManager;
+};
+
+
+class FieldSamplerQueryCPU : public FieldSamplerQuery
+{
+public:
+	FieldSamplerQueryCPU(const FieldSamplerQueryDescIntl& desc, ResourceList& list, FieldSamplerManager* manager);
+	~FieldSamplerQueryCPU();
+
+	// FieldSamplerQueryIntl methods
+	PxTaskID submitFieldSamplerQuery(const FieldSamplerQueryDataIntl& data, PxTaskID taskID);
+
+protected:
+	void		execute();
+	void		executeScene(	const SceneInfo* sceneInfo, 
+								const FieldSamplerIntl::ExecuteData& executeData, 
+								PxVec4* accumField, 
+								PxVec4* accumVelocity, 
+								uint32_t positionStride, 
+								uint32_t velocityStride, 
+								uint32_t massStride);
+
+	uint32_t				mExecuteCount;
+	nvidia::Array<PxVec3>	mResultField;
+	nvidia::Array<float>	mWeights;
+
+	PxTask*				mTaskExecute;
+
+	friend class TaskExecute;
+
+	class SceneInfoCPU : public SceneInfo
+	{
+	public:
+		SceneInfoCPU(ResourceList& list, FieldSamplerQuery* query, FieldSamplerSceneWrapper* sceneWrapper)
+			: SceneInfo(list, query, sceneWrapper)
+		{
+		}
+	};
+
+	virtual SceneInfo* createSceneInfo(FieldSamplerSceneWrapper* sceneWrapper)
+	{
+		return PX_NEW(SceneInfoCPU)(mSceneList, this, sceneWrapper);
+	}
+
+};
+
+#if APEX_CUDA_SUPPORT
+
+class FieldSamplerQueryGPU : public FieldSamplerQueryCPU
+{
+public:
+	FieldSamplerQueryGPU(const FieldSamplerQueryDescIntl& desc, ResourceList& list, FieldSamplerManager* manager);
+	~FieldSamplerQueryGPU();
+
+	// FieldSamplerQueryIntl methods
+	PxTaskID submitFieldSamplerQuery(const FieldSamplerQueryDataIntl& data, PxTaskID taskID);
+
+	PxVec3 executeFieldSamplerQueryOnGrid(const FieldSamplerQueryGridDataIntl&);
+
+protected:
+	bool		launch(CUstream stream, int kernelIndex);
+	void		prepare();
+	bool		copy(CUstream stream, int kernelIndex);
+	void		fetch();
+
+	PxTask*		mTaskLaunch;
+	PxTask*		mTaskPrepare;
+	PxTask*		mTaskCopy;
+	PxTask*		mTaskFetch;
+
+	ApexMirroredArray<PxVec4>	mPositionMass;
+	ApexMirroredArray<PxVec4>	mVelocity;
+	ApexMirroredArray<PxVec4>	mAccumField;
+	PxGpuCopyDescQueue			mCopyQueue;
+
+	friend class FieldSamplerQueryLaunchTask;
+	friend class FieldSamplerQueryPrepareTask;
+	friend class FieldSamplerQueryCopyTask;
+	friend class FieldSamplerQueryFetchTask;
+
+	class SceneInfoGPU : public SceneInfo
+	{
+	public:
+		SceneInfoGPU(ResourceList& list, FieldSamplerQuery* query, FieldSamplerSceneWrapper* sceneWrapper);
+
+		virtual bool update();
+
+		PX_INLINE InplaceHandle<nvidia::fieldsampler::FieldSamplerParamsExArray> getParamsHandle() const
+		{
+			PX_ASSERT(mParamsExArrayHandle.isNull() == false);
+			return mParamsExArrayHandle;
+		}
+		PX_INLINE InplaceHandle<nvidia::fieldsampler::FieldSamplerQueryParams> getQueryParamsHandle() const
+		{
+			PX_ASSERT(mQueryParamsHandle.isNull() == false);
+			return mQueryParamsHandle;
+		}
+
+	private:
+		ApexCudaConstMemGroup                           mConstMemGroup;
+		InplaceHandle<nvidia::fieldsampler::FieldSamplerParamsExArray>    mParamsExArrayHandle;
+		InplaceHandle<nvidia::fieldsampler::FieldSamplerQueryParams>      mQueryParamsHandle;
+	};
+
+	virtual SceneInfo* createSceneInfo(FieldSamplerSceneWrapper* sceneWrapper)
+	{
+		return PX_NEW(SceneInfoGPU)(mSceneList, this, sceneWrapper);
+	}
+};
+
+#endif
+
+}
+} // end namespace nvidia::apex
+
+#endif
diff --git a/APEX_1.4/module/fieldsampler/include/FieldSamplerScene.h b/APEX_1.4/module/fieldsampler/include/FieldSamplerScene.h
new file mode 100644
index 00000000..1c63501f
--- /dev/null
+++ b/APEX_1.4/module/fieldsampler/include/FieldSamplerScene.h
@@ -0,0 +1,218 @@
+/*
+ * Copyright (c) 2008-2015, 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 __FIELD_SAMPLER_SCENE_H__
+#define __FIELD_SAMPLER_SCENE_H__
+
+#include "Apex.h"
+
+#include "ModuleFieldSamplerImpl.h"
+#include "FieldSamplerSceneIntl.h"
+#include "FieldSamplerQueryIntl.h"
+#include "ApexSDKIntl.h"
+#include "ModuleIntl.h"
+#include "ApexContext.h"
+#include "ApexSDKHelpers.h"
+#include "PsArray.h"
+
+#include "PxTask.h"
+
+#if APEX_CUDA_SUPPORT
+#include "ApexCudaWrapper.h"
+#include "ApexCuda.h"
+#include "CudaModuleScene.h"
+
+#include "../cuda/include/common.h"
+
+#define SCENE_CUDA_OBJ(scene, name) static_cast<FieldSamplerSceneGPU*>(scene)->APEX_CUDA_OBJ_NAME(name)
+#endif
+
+
+namespace nvidia
+{
+namespace apex
+{
+class SceneIntl;
+}
+namespace fieldsampler
+{
+
+class ModuleFieldSamplerImpl;
+class FieldSamplerPhysXMonitor;
+class FieldSamplerManager;
+class FieldSamplerQuery;
+
+class FieldSamplerScene : public ModuleSceneIntl, public ApexContext, public ApexResourceInterface, public ApexResource
+{
+public:
+	FieldSamplerScene(ModuleFieldSamplerImpl& module, SceneIntl& scene, RenderDebugInterface* debugRender, ResourceList& list);
+	~FieldSamplerScene();
+
+	/* ModuleSceneIntl */
+	void				visualize();
+
+	PxScene*			getModulePhysXScene() const
+	{
+		return mPhysXScene;
+	}
+	void				setModulePhysXScene(PxScene*);
+	PxScene* 			mPhysXScene;
+
+	void				submitTasks(float elapsedTime, float substepSize, uint32_t numSubSteps);
+	void				setTaskDependencies();
+	void				fetchResults();
+
+	virtual Module*	getModule()
+	{
+		return mModule;
+	}
+
+	virtual SceneStats* getStats()
+	{
+		return 0;
+	}
+
+	bool							lockRenderResources()
+	{
+		renderLockAllActors();	// Lock options not implemented yet
+		return true;
+	}
+
+	bool							unlockRenderResources()
+	{
+		renderUnLockAllActors();	// Lock options not implemented yet
+		return true;
+	}
+
+	/* ApexResourceInterface */
+	uint32_t				getListIndex() const
+	{
+		return m_listIndex;
+	}
+	void				setListIndex(ResourceList& list, uint32_t index)
+	{
+		m_listIndex = index;
+		m_list = &list;
+	}
+	void				release()
+	{
+		mModule->releaseModuleSceneIntl(*this);
+	}
+	SceneIntl& getApexScene() const
+	{
+		return *mApexScene;
+	}
+
+	FieldSamplerManagerIntl* getManager();
+
+	uint32_t createForceSampleBatch(uint32_t maxCount, const physx::PxFilterData filterData);
+	void releaseForceSampleBatch(uint32_t batchId);
+	void submitForceSampleBatch(	uint32_t batchId, PxVec4* forces, const uint32_t forcesStride,
+									const PxVec3* positions, const uint32_t positionsStride,
+									const PxVec3* velocities, const uint32_t velocitiesStride,
+									const float* mass, const uint32_t massStride,
+									const uint32_t* indices, const uint32_t numIndices);
+
+	/* Toggle PhysX Monitor on/off */
+	void enablePhysXMonitor(bool enable);
+
+	void setPhysXFilterData(physx::PxFilterData filterData);
+
+protected:
+	void                destroy();
+
+	virtual FieldSamplerManager* createManager() = 0;
+
+	class TaskPhysXMonitorLoad : public PxTask
+	{
+	public:
+		TaskPhysXMonitorLoad() {}
+		const char* getName() const
+		{
+			return FSST_PHYSX_MONITOR_LOAD;
+		}		
+		void run() {/* void task */};
+	};
+	TaskPhysXMonitorLoad	mPhysXMonitorLoadTask;
+	class TaskPhysXMonitorFetch : public PxTask
+	{
+	public:
+		TaskPhysXMonitorFetch() {}
+		const char* getName() const
+		{
+			return FSST_PHYSX_MONITOR_FETCH;
+		}		
+		void run() {/* void task */};
+	};
+	TaskPhysXMonitorFetch	mPhysXMonitorFetchTask;
+
+	ModuleFieldSamplerImpl*		mModule;
+	SceneIntl*			mApexScene;
+	RenderDebugInterface*		mDebugRender;
+	FieldSamplerPhysXMonitor* mPhysXMonitor;
+
+	FieldSamplerManager*	mManager;
+
+	uint32_t							mForceSampleBatchBufferSize;
+	uint32_t							mForceSampleBatchBufferPos;
+	Array <FieldSamplerQuery*>		mForceSampleBatchQuery;
+	Array <FieldSamplerQueryDataIntl> mForceSampleBatchQueryData;
+	Array <PxVec4>					mForceSampleBatchPosition;
+	Array <PxVec4>					mForceSampleBatchVelocity;
+	Array <float>					mForceSampleBatchMass;
+
+	friend class ModuleFieldSamplerImpl;
+	friend class FieldSamplerManager;
+};
+
+class FieldSamplerSceneCPU : public FieldSamplerScene
+{
+public:
+	FieldSamplerSceneCPU(ModuleFieldSamplerImpl& module, SceneIntl& scene, RenderDebugInterface* debugRender, ResourceList& list);
+	~FieldSamplerSceneCPU();
+
+protected:
+	virtual FieldSamplerManager* createManager();
+
+};
+
+#if APEX_CUDA_SUPPORT
+class FieldSamplerSceneGPU : public FieldSamplerScene, public CudaModuleScene
+{
+public:
+	FieldSamplerSceneGPU(ModuleFieldSamplerImpl& module, SceneIntl& scene, RenderDebugInterface* debugRender, ResourceList& list);
+	~FieldSamplerSceneGPU();
+
+	void* getHeadCudaObj()
+	{
+		return CudaModuleScene::getHeadCudaObj();
+	}
+
+//CUDA module objects
+#include "../cuda/include/fieldsampler.h"
+
+	PxCudaContextManager* getCudaContext() const
+	{
+		return mCtxMgr;
+	}
+
+protected:
+	virtual FieldSamplerManager* createManager();
+
+	/* keep a convenience pointer to the cuda context manager */
+	PxCudaContextManager* mCtxMgr;
+};
+#endif
+
+}
+} // end namespace nvidia::apex
+
+#endif
diff --git a/APEX_1.4/module/fieldsampler/include/FieldSamplerSceneWrapper.h b/APEX_1.4/module/fieldsampler/include/FieldSamplerSceneWrapper.h
new file mode 100644
index 00000000..dce7f529
--- /dev/null
+++ b/APEX_1.4/module/fieldsampler/include/FieldSamplerSceneWrapper.h
@@ -0,0 +1,206 @@
+/*
+ * Copyright (c) 2008-2015, 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 __FIELD_SAMPLER_SCENE_WRAPPER_H__
+#define __FIELD_SAMPLER_SCENE_WRAPPER_H__
+
+#include "Apex.h"
+#include "ApexSDKHelpers.h"
+#include "ApexActor.h"
+#include "FieldSamplerSceneIntl.h"
+
+#if APEX_CUDA_SUPPORT
+#include "ApexCudaWrapper.h"
+#endif
+
+#include "FieldSamplerCommon.h"
+
+namespace nvidia
+{
+namespace fieldsampler
+{
+
+class FieldSamplerManager;
+class FieldSamplerWrapper;
+class FieldBoundaryWrapper;
+
+class FieldSamplerSceneWrapper : public ApexResourceInterface, public ApexResource
+{
+public:
+	// ApexResourceInterface methods
+	void			release();
+	void			setListIndex(ResourceList& list, uint32_t index)
+	{
+		m_listIndex = index;
+		m_list = &list;
+	}
+	uint32_t	getListIndex() const
+	{
+		return m_listIndex;
+	}
+
+	FieldSamplerSceneWrapper(ResourceList& list, FieldSamplerManager* manager, FieldSamplerSceneIntl* fieldSamplerScene);
+
+	FieldSamplerSceneIntl* getInternalFieldSamplerScene() const
+	{
+		return mFieldSamplerScene;
+	}
+	const FieldSamplerSceneDescIntl& getInternalFieldSamplerSceneDesc() const
+	{
+		return mFieldSamplerSceneDesc;
+	}
+
+	virtual void update();
+	virtual void postUpdate() = 0;
+
+	class FieldBoundaryInfo;
+
+	FieldBoundaryInfo*	addFieldBoundary(FieldBoundaryWrapper* fieldBoundaryWrapper);
+	void				removeFieldBoundary(FieldBoundaryInfo* fieldBoundaryInfo);
+
+protected:
+	virtual FieldBoundaryInfo* createFieldBoundaryInfo(FieldBoundaryWrapper* fieldBoundaryWrapper) = 0;
+
+protected:
+	FieldSamplerManager*	mManager;
+	FieldSamplerSceneIntl*	mFieldSamplerScene;
+	FieldSamplerSceneDescIntl	mFieldSamplerSceneDesc;
+
+	ResourceList			mFieldBoundaryList;
+	bool					mFieldBoundaryListChanged;
+};
+
+class FieldSamplerSceneWrapper::FieldBoundaryInfo : public ApexResourceInterface, public ApexResource
+{
+protected:
+	FieldBoundaryWrapper*	mFieldBoundaryWrapper;
+	uint32_t			mRefCount;
+
+	FieldBoundaryInfo(ResourceList& list, FieldBoundaryWrapper* fieldBoundaryWrapper)
+		: mFieldBoundaryWrapper(fieldBoundaryWrapper), mRefCount(0)
+	{
+		list.add(*this);
+	}
+
+public:
+	FieldBoundaryWrapper* getFieldBoundaryWrapper() const
+	{
+		return mFieldBoundaryWrapper;
+	}
+
+	virtual void update() {}
+
+	void addRef()
+	{
+		++mRefCount;
+	}
+	bool releaseRef()
+	{
+		if (--mRefCount == 0)
+		{
+			release();
+			return true;
+		}
+		return false;
+	}
+
+	// ApexResourceInterface methods
+	void						release()
+	{
+		delete this;
+	}
+	void						setListIndex(ResourceList& list, uint32_t index)
+	{
+		m_listIndex = index;
+		m_list = &list;
+	}
+	uint32_t				getListIndex() const
+	{
+		return m_listIndex;
+	}
+};
+
+/******************************** CPU Version ********************************/
+
+class FieldSamplerSceneWrapperCPU : public FieldSamplerSceneWrapper
+{
+public:
+	FieldSamplerSceneWrapperCPU(ResourceList& list, FieldSamplerManager* manager, FieldSamplerSceneIntl* fieldSamplerScene);
+
+	virtual void postUpdate() {}
+
+	class FieldBoundaryInfoCPU : public FieldBoundaryInfo
+	{
+	public:
+		FieldBoundaryInfoCPU(ResourceList& list, FieldBoundaryWrapper* fieldBoundaryWrapper)
+			: FieldBoundaryInfo(list, fieldBoundaryWrapper)
+		{
+		}
+	};
+
+protected:
+	virtual FieldBoundaryInfo* createFieldBoundaryInfo(FieldBoundaryWrapper* fieldBoundaryWrapper)
+	{
+		return PX_NEW(FieldBoundaryInfoCPU)(mFieldBoundaryList, fieldBoundaryWrapper);
+	}
+
+};
+
+/******************************** GPU Version ********************************/
+
+#if APEX_CUDA_SUPPORT
+class FieldSamplerSceneWrapperGPU : public FieldSamplerSceneWrapper
+{
+public:
+	FieldSamplerSceneWrapperGPU(ResourceList& list, FieldSamplerManager* manager, FieldSamplerSceneIntl* fieldSamplerScene);
+
+	virtual void postUpdate();
+
+	PX_INLINE ApexCudaConstStorage& getConstStorage()
+	{
+		return mConstStorage;
+	}
+
+	class FieldBoundaryInfoGPU : public FieldBoundaryInfo
+	{
+	public:
+		FieldBoundaryInfoGPU(ResourceList& list, FieldBoundaryWrapper* fieldBoundaryWrapper, ApexCudaConstStorage& constStorage);
+
+		virtual void update();
+
+		PX_INLINE InplaceHandle<FieldShapeGroupParams> getShapeGroupParamsHandle() const
+		{
+			PX_ASSERT(mFieldShapeGroupParamsHandle.isNull() == false);
+			return mFieldShapeGroupParamsHandle;
+		}
+
+	private:
+		ApexCudaConstMemGroup                   mConstMemGroup;
+		InplaceHandle<FieldShapeGroupParams>    mFieldShapeGroupParamsHandle;
+	};
+
+protected:
+	virtual FieldBoundaryInfo* createFieldBoundaryInfo(FieldBoundaryWrapper* fieldBoundaryWrapper)
+	{
+		return PX_NEW(FieldBoundaryInfoGPU)(mFieldBoundaryList, fieldBoundaryWrapper, getConstStorage());
+	}
+
+private:
+	FieldSamplerSceneWrapperGPU& operator=(const FieldSamplerSceneWrapperGPU&);
+
+	ApexCudaConstStorage& mConstStorage;
+};
+#endif // APEX_CUDA_SUPPORT
+
+}
+} // end namespace nvidia::apex
+
+#endif
diff --git a/APEX_1.4/module/fieldsampler/include/FieldSamplerWrapper.h b/APEX_1.4/module/fieldsampler/include/FieldSamplerWrapper.h
new file mode 100644
index 00000000..4eebc10b
--- /dev/null
+++ b/APEX_1.4/module/fieldsampler/include/FieldSamplerWrapper.h
@@ -0,0 +1,149 @@
+/*
+ * Copyright (c) 2008-2015, 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 __FIELD_SAMPLER_WRAPPER_H__
+#define __FIELD_SAMPLER_WRAPPER_H__
+
+#include "Apex.h"
+#include "ApexSDKHelpers.h"
+#include "ApexActor.h"
+#include "FieldSamplerIntl.h"
+
+#if APEX_CUDA_SUPPORT
+#include "ApexCudaWrapper.h"
+#endif
+
+#include "FieldSamplerSceneWrapper.h"
+#include "FieldSamplerCommon.h"
+
+namespace nvidia
+{
+namespace fieldsampler
+{
+
+class FieldSamplerManager;
+
+class FieldBoundaryWrapper;
+
+class FieldSamplerWrapper : public ApexResourceInterface, public ApexResource
+{
+public:
+	// ApexResourceInterface methods
+	void			release();
+	void			setListIndex(ResourceList& list, uint32_t index)
+	{
+		m_listIndex = index;
+		m_list = &list;
+	}
+	uint32_t	getListIndex() const
+	{
+		return m_listIndex;
+	}
+
+	FieldSamplerWrapper(ResourceList& list, FieldSamplerManager* manager, FieldSamplerIntl* fieldSampler, const FieldSamplerDescIntl& fieldSamplerDesc, FieldSamplerSceneWrapper* fieldSamplerSceneWrapper);
+
+	virtual void update();
+
+	PX_INLINE FieldSamplerIntl* getInternalFieldSampler() const
+	{
+		return mFieldSampler;
+	}
+	PX_INLINE const FieldSamplerDescIntl& getInternalFieldSamplerDesc() const
+	{
+		return mFieldSamplerDesc;
+	}
+	PX_INLINE FieldSamplerSceneWrapper* getFieldSamplerSceneWrapper() const
+	{
+		return mSceneWrapper;
+	}
+
+	bool addFieldBoundary(FieldBoundaryWrapper* wrapper);
+	bool removeFieldBoundary(FieldBoundaryWrapper* wrapper);
+
+	uint32_t getFieldBoundaryCount() const
+	{
+		return mFieldBoundaryInfoArray.size();
+	}
+	FieldBoundaryWrapper* getFieldBoundaryWrapper(uint32_t index) const
+	{
+		return mFieldBoundaryInfoArray[index]->getFieldBoundaryWrapper();
+	}
+
+	PX_INLINE const FieldShapeDescIntl&   getInternalFieldSamplerShape() const
+	{
+		return mFieldSamplerShape;
+	}
+	PX_INLINE bool                      isFieldSamplerChanged() const
+	{
+		return mFieldSamplerShapeChanged;
+	}
+	PX_INLINE bool                      isEnabled() const
+	{
+		return mIsEnabled;
+	}
+	PX_INLINE bool                      isEnabledChanged() const
+	{
+		return (mIsEnabled != mIsEnabledLast);
+	}
+
+protected:
+	FieldSamplerManager*		mManager;
+	FieldSamplerIntl*				mFieldSampler;
+	FieldSamplerDescIntl			mFieldSamplerDesc;
+
+	FieldShapeDescIntl			mFieldSamplerShape;
+	bool						mFieldSamplerShapeChanged;
+
+	FieldSamplerSceneWrapper*	mSceneWrapper;
+	uint32_t				mQueryRefCount;
+
+	nvidia::Array<FieldSamplerSceneWrapper::FieldBoundaryInfo*>	mFieldBoundaryInfoArray;
+	bool														mFieldBoundaryInfoArrayChanged;
+
+	bool						mIsEnabled;
+	bool						mIsEnabledLast;
+
+	friend class FieldSamplerManager;
+};
+
+
+class FieldSamplerWrapperCPU : public FieldSamplerWrapper
+{
+public:
+	FieldSamplerWrapperCPU(ResourceList& list, FieldSamplerManager* manager, FieldSamplerIntl* fieldSampler, const FieldSamplerDescIntl& fieldSamplerDesc, FieldSamplerSceneWrapper* fieldSamplerSceneWrapper);
+
+private:
+};
+
+#if APEX_CUDA_SUPPORT
+class FieldSamplerWrapperGPU : public FieldSamplerWrapperCPU
+{
+public:
+	FieldSamplerWrapperGPU(ResourceList& list, FieldSamplerManager* manager, FieldSamplerIntl* fieldSampler, const FieldSamplerDescIntl& fieldSamplerDesc, FieldSamplerSceneWrapper* fieldSamplerSceneWrapper);
+
+	virtual void update();
+
+	PX_INLINE InplaceHandle<FieldSamplerParams>   getParamsHandle() const
+	{
+		PX_ASSERT(mFieldSamplerParamsHandle.isNull() == false);
+		return mFieldSamplerParamsHandle;
+	}
+
+private:
+	ApexCudaConstMemGroup               mConstMemGroup;
+	InplaceHandle<FieldSamplerParams>   mFieldSamplerParamsHandle;
+};
+#endif
+
+}
+} // end namespace nvidia::apex
+
+#endif
diff --git a/APEX_1.4/module/fieldsampler/include/ModuleEventDefs.h b/APEX_1.4/module/fieldsampler/include/ModuleEventDefs.h
new file mode 100644
index 00000000..583b50d1
--- /dev/null
+++ b/APEX_1.4/module/fieldsampler/include/ModuleEventDefs.h
@@ -0,0 +1,26 @@
+/*
+ * Copyright (c) 2008-2015, 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.
+ */
+
+
+// This file is used to define a list of AgPerfMon events.
+//
+// This file is included exclusively by ModulePerfScope.h
+// and by ModuleFoo.cpp, for the purpose of building
+// an enumeration (enum xx) and an array of strings ()
+// that contain the list of events.
+//
+// This file should only contain event definitions, using the
+// DEFINE_EVENT macro.  E.g.:
+//
+//     DEFINE_EVENT(sample_name_1)
+//     DEFINE_EVENT(sample_name_2)
+//     DEFINE_EVENT(sample_name_3)
+
+DEFINE_EVENT(FieldSamplerSceneFetchResults)
diff --git a/APEX_1.4/module/fieldsampler/include/ModuleFieldSamplerImpl.h b/APEX_1.4/module/fieldsampler/include/ModuleFieldSamplerImpl.h
new file mode 100644
index 00000000..b15f5189
--- /dev/null
+++ b/APEX_1.4/module/fieldsampler/include/ModuleFieldSamplerImpl.h
@@ -0,0 +1,98 @@
+/*
+ * Copyright (c) 2008-2015, 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 __MODULE_FIELD_SAMPLER_IMPL_H__
+#define __MODULE_FIELD_SAMPLER_IMPL_H__
+
+#include "Apex.h"
+#include "ModuleFieldSampler.h"
+#include "ApexSDKIntl.h"
+#include "ModuleIntl.h"
+#include "ModuleBase.h"
+
+#include "ApexSDKHelpers.h"
+#include "ApexRWLockable.h"
+#include "ModuleFieldSamplerRegistration.h"
+
+#include "ModuleFieldSamplerIntl.h"
+
+namespace nvidia
+{
+namespace apex
+{
+class SceneIntl;
+}
+
+namespace fieldsampler
+{
+
+class FieldSamplerScene;
+
+class ModuleFieldSamplerImpl : public ModuleFieldSampler, public ModuleFieldSamplerIntl, public ModuleBase, public ApexRWLockable
+{
+public:
+	APEX_RW_LOCKABLE_BOILERPLATE
+
+	ModuleFieldSamplerImpl(ApexSDKIntl* sdk);
+	~ModuleFieldSamplerImpl();
+
+	FieldSamplerManagerIntl*		getInternalFieldSamplerManager(const Scene& apexScene);
+
+	bool setFieldSamplerWeightedCollisionFilterCallback(const Scene& apexScene,FieldSamplerWeightedCollisionFilterCallback *callback);
+
+#if PX_PHYSICS_VERSION_MAJOR == 3
+	void						enablePhysXMonitor(const Scene& apexScene, bool enable);
+	void						setPhysXMonitorFilterData(const Scene& apexScene, physx::PxFilterData filterData);
+#endif
+
+	uint32_t				createForceSampleBatch(const Scene& apexScene, uint32_t maxCount, const physx::PxFilterData filterData);
+	void						releaseForceSampleBatch(const Scene& apexScene, uint32_t batchId);
+	void						submitForceSampleBatch(const Scene& apexScene, uint32_t batchId,
+													   PxVec4* forces, const uint32_t forcesStride,
+													   const PxVec3* positions, const uint32_t positionsStride,
+													   const PxVec3* velocities, const uint32_t velocitiesStride,
+													   const float* mass, const uint32_t massStride,
+													   const uint32_t* indices, const uint32_t numIndices);
+
+
+	// base class methods
+	void						init(NvParameterized::Interface&);
+	NvParameterized::Interface* getDefaultModuleDesc();
+	void release()
+	{
+		ModuleBase::release();
+	}
+	void destroy();
+	const char* getName() const
+	{
+		return ModuleBase::getName();
+	}
+	
+	ModuleSceneIntl* 				createInternalModuleScene(SceneIntl&, RenderDebugInterface*);
+	void						releaseModuleSceneIntl(ModuleSceneIntl&);
+	AuthObjTypeID				getModuleID() const;
+	RenderableIterator* 	createRenderableIterator(const Scene&);
+
+
+protected:
+	FieldSamplerScene* 	getFieldSamplerScene(const Scene& apexScene) const;
+
+	ResourceList				mFieldSamplerScenes;
+
+private:
+
+	FieldSamplerModuleParameters* 			mModuleParams;
+};
+
+}
+} // end namespace nvidia::apex
+
+#endif // __MODULE_FIELD_SAMPLER_IMPL_H__
diff --git a/APEX_1.4/module/fieldsampler/include/ModuleGpuEventDefs.h b/APEX_1.4/module/fieldsampler/include/ModuleGpuEventDefs.h
new file mode 100644
index 00000000..faa6b9d4
--- /dev/null
+++ b/APEX_1.4/module/fieldsampler/include/ModuleGpuEventDefs.h
@@ -0,0 +1,26 @@
+/*
+ * Copyright (c) 2008-2015, 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.
+ */
+
+
+// This file is used to define a list of AgPerfMon events.
+//
+// This file is included exclusively by ModulePerfScope.h
+// and by ModuleFoo.cpp, for the purpose of building
+// an enumeration (enum xx) and an array of strings ()
+// that contain the list of events.
+//
+// This file should only contain event definitions, using the
+// DEFINE_EVENT macro.  E.g.:
+//
+//     DEFINE_GPU_EVENT(kernel_name_1)
+//     DEFINE_GPU_EVENT(kernel_name_2)
+//     DEFINE_GPU_EVENT(copy_name_1)
+
+DEFINE_GPU_EVENT(FieldSamplerKernel)
diff --git a/APEX_1.4/module/fieldsampler/include/ModulePerfScope.h b/APEX_1.4/module/fieldsampler/include/ModulePerfScope.h
new file mode 100644
index 00000000..2692f4ba
--- /dev/null
+++ b/APEX_1.4/module/fieldsampler/include/ModulePerfScope.h
@@ -0,0 +1,18 @@
+/*
+ * Copyright (c) 2008-2015, 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 __MODULE_PERF_SCOPE_H___
+#define __MODULE_PERF_SCOPE_H___
+
+#define MODULE_NAMESPACE fieldsampler
+#include "ModuleProfileCommon.h"
+
+#endif
\ No newline at end of file
diff --git a/APEX_1.4/module/fieldsampler/include/autogen/FieldSamplerModuleParameters.h b/APEX_1.4/module/fieldsampler/include/autogen/FieldSamplerModuleParameters.h
new file mode 100644
index 00000000..4267524c
--- /dev/null
+++ b/APEX_1.4/module/fieldsampler/include/autogen/FieldSamplerModuleParameters.h
@@ -0,0 +1,230 @@
+// This code contains NVIDIA Confidential Information and is disclosed to you
+// under a form of NVIDIA software license agreement provided separately to you.
+//
+// Notice
+// NVIDIA Corporation and its licensors retain all intellectual property and
+// proprietary rights in and to this software and 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.
+//
+// ALL NVIDIA DESIGN SPECIFICATIONS, CODE ARE PROVIDED "AS IS.". NVIDIA MAKES
+// NO WARRANTIES, EXPRESSED, IMPLIED, STATUTORY, OR OTHERWISE WITH RESPECT TO
+// THE MATERIALS, AND EXPRESSLY DISCLAIMS ALL IMPLIED WARRANTIES OF NONINFRINGEMENT,
+// MERCHANTABILITY, AND FITNESS FOR A PARTICULAR PURPOSE.
+//
+// Information and code furnished is believed to be accurate and reliable.
+// However, NVIDIA Corporation assumes no responsibility for the consequences of use of such
+// information or for any infringement of patents or other rights of third parties that may
+// result from its use. No license is granted by implication or otherwise under any patent
+// or patent rights of NVIDIA Corporation. Details are subject to change without notice.
+// This code supersedes and replaces all information previously supplied.
+// NVIDIA Corporation products are not authorized for use as critical
+// components in life support devices or systems without express written approval of
+// NVIDIA Corporation.
+//
+// Copyright (c) 2008-2015 NVIDIA Corporation. All rights reserved.
+
+// This file was generated by NvParameterized/scripts/GenParameterized.pl
+
+
+#ifndef HEADER_FieldSamplerModuleParameters_h
+#define HEADER_FieldSamplerModuleParameters_h
+
+#include "NvParametersTypes.h"
+
+#ifndef NV_PARAMETERIZED_ONLY_LAYOUTS
+#include "nvparameterized/NvParameterized.h"
+#include "nvparameterized/NvParameterizedTraits.h"
+#include "NvParameters.h"
+#include "NvTraitsInternal.h"
+#endif
+
+namespace nvidia
+{
+namespace fieldsampler
+{
+
+#if PX_VC
+#pragma warning(push)
+#pragma warning(disable: 4324) // structure was padded due to __declspec(align())
+#endif
+
+namespace FieldSamplerModuleParametersNS
+{
+
+
+
+struct ParametersStruct
+{
+
+	uint32_t unused;
+
+};
+
+static const uint32_t checksum[] = { 0xe3bea6f1, 0xd89e4c54, 0x40faae3e, 0x35941431, };
+
+} // namespace FieldSamplerModuleParametersNS
+
+#ifndef NV_PARAMETERIZED_ONLY_LAYOUTS
+class FieldSamplerModuleParameters : public NvParameterized::NvParameters, public FieldSamplerModuleParametersNS::ParametersStruct
+{
+public:
+	FieldSamplerModuleParameters(NvParameterized::Traits* traits, void* buf = 0, int32_t* refCount = 0);
+
+	virtual ~FieldSamplerModuleParameters();
+
+	virtual void destroy();
+
+	static const char* staticClassName(void)
+	{
+		return("FieldSamplerModuleParameters");
+	}
+
+	const char* className(void) const
+	{
+		return(staticClassName());
+	}
+
+	static const uint32_t ClassVersion = ((uint32_t)0 << 16) + (uint32_t)0;
+
+	static uint32_t staticVersion(void)
+	{
+		return ClassVersion;
+	}
+
+	uint32_t version(void) const
+	{
+		return(staticVersion());
+	}
+
+	static const uint32_t ClassAlignment = 8;
+
+	static const uint32_t* staticChecksum(uint32_t& bits)
+	{
+		bits = 8 * sizeof(FieldSamplerModuleParametersNS::checksum);
+		return FieldSamplerModuleParametersNS::checksum;
+	}
+
+	static void freeParameterDefinitionTable(NvParameterized::Traits* traits);
+
+	const uint32_t* checksum(uint32_t& bits) const
+	{
+		return staticChecksum(bits);
+	}
+
+	const FieldSamplerModuleParametersNS::ParametersStruct& parameters(void) const
+	{
+		FieldSamplerModuleParameters* tmpThis = const_cast<FieldSamplerModuleParameters*>(this);
+		return *(static_cast<FieldSamplerModuleParametersNS::ParametersStruct*>(tmpThis));
+	}
+
+	FieldSamplerModuleParametersNS::ParametersStruct& parameters(void)
+	{
+		return *(static_cast<FieldSamplerModuleParametersNS::ParametersStruct*>(this));
+	}
+
+	virtual NvParameterized::ErrorType getParameterHandle(const char* long_name, NvParameterized::Handle& handle) const;
+	virtual NvParameterized::ErrorType getParameterHandle(const char* long_name, NvParameterized::Handle& handle);
+
+	void initDefaults(void);
+
+protected:
+
+	virtual const NvParameterized::DefinitionImpl* getParameterDefinitionTree(void);
+	virtual const NvParameterized::DefinitionImpl* getParameterDefinitionTree(void) const;
+
+
+	virtual void getVarPtr(const NvParameterized::Handle& handle, void*& ptr, size_t& offset) const;
+
+private:
+
+	void buildTree(void);
+	void initDynamicArrays(void);
+	void initStrings(void);
+	void initReferences(void);
+	void freeDynamicArrays(void);
+	void freeStrings(void);
+	void freeReferences(void);
+
+	static bool mBuiltFlag;
+	static NvParameterized::MutexType mBuiltFlagMutex;
+};
+
+class FieldSamplerModuleParametersFactory : public NvParameterized::Factory
+{
+	static const char* const vptr;
+
+public:
+
+	virtual void freeParameterDefinitionTable(NvParameterized::Traits* traits)
+	{
+		FieldSamplerModuleParameters::freeParameterDefinitionTable(traits);
+	}
+
+	virtual NvParameterized::Interface* create(NvParameterized::Traits* paramTraits)
+	{
+		// placement new on this class using mParameterizedTraits
+
+		void* newPtr = paramTraits->alloc(sizeof(FieldSamplerModuleParameters), FieldSamplerModuleParameters::ClassAlignment);
+		if (!NvParameterized::IsAligned(newPtr, FieldSamplerModuleParameters::ClassAlignment))
+		{
+			NV_PARAM_TRAITS_WARNING(paramTraits, "Unaligned memory allocation for class FieldSamplerModuleParameters");
+			paramTraits->free(newPtr);
+			return 0;
+		}
+
+		memset(newPtr, 0, sizeof(FieldSamplerModuleParameters)); // always initialize memory allocated to zero for default values
+		return NV_PARAM_PLACEMENT_NEW(newPtr, FieldSamplerModuleParameters)(paramTraits);
+	}
+
+	virtual NvParameterized::Interface* finish(NvParameterized::Traits* paramTraits, void* bufObj, void* bufStart, int32_t* refCount)
+	{
+		if (!NvParameterized::IsAligned(bufObj, FieldSamplerModuleParameters::ClassAlignment)
+		        || !NvParameterized::IsAligned(bufStart, FieldSamplerModuleParameters::ClassAlignment))
+		{
+			NV_PARAM_TRAITS_WARNING(paramTraits, "Unaligned memory allocation for class FieldSamplerModuleParameters");
+			return 0;
+		}
+
+		// Init NvParameters-part
+		// We used to call empty constructor of FieldSamplerModuleParameters here
+		// but it may call default constructors of members and spoil the data
+		NV_PARAM_PLACEMENT_NEW(bufObj, NvParameterized::NvParameters)(paramTraits, bufStart, refCount);
+
+		// Init vtable (everything else is already initialized)
+		*(const char**)bufObj = vptr;
+
+		return (FieldSamplerModuleParameters*)bufObj;
+	}
+
+	virtual const char* getClassName()
+	{
+		return (FieldSamplerModuleParameters::staticClassName());
+	}
+
+	virtual uint32_t getVersion()
+	{
+		return (FieldSamplerModuleParameters::staticVersion());
+	}
+
+	virtual uint32_t getAlignment()
+	{
+		return (FieldSamplerModuleParameters::ClassAlignment);
+	}
+
+	virtual const uint32_t* getChecksum(uint32_t& bits)
+	{
+		return (FieldSamplerModuleParameters::staticChecksum(bits));
+	}
+};
+#endif // NV_PARAMETERIZED_ONLY_LAYOUTS
+
+} // namespace fieldsampler
+} // namespace nvidia
+
+#if PX_VC
+#pragma warning(pop)
+#endif
+
+#endif
diff --git a/APEX_1.4/module/fieldsampler/include/autogen/FieldSamplerPhysXMonitorParams.h b/APEX_1.4/module/fieldsampler/include/autogen/FieldSamplerPhysXMonitorParams.h
new file mode 100644
index 00000000..037d7b3b
--- /dev/null
+++ b/APEX_1.4/module/fieldsampler/include/autogen/FieldSamplerPhysXMonitorParams.h
@@ -0,0 +1,232 @@
+// This code contains NVIDIA Confidential Information and is disclosed to you
+// under a form of NVIDIA software license agreement provided separately to you.
+//
+// Notice
+// NVIDIA Corporation and its licensors retain all intellectual property and
+// proprietary rights in and to this software and 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.
+//
+// ALL NVIDIA DESIGN SPECIFICATIONS, CODE ARE PROVIDED "AS IS.". NVIDIA MAKES
+// NO WARRANTIES, EXPRESSED, IMPLIED, STATUTORY, OR OTHERWISE WITH RESPECT TO
+// THE MATERIALS, AND EXPRESSLY DISCLAIMS ALL IMPLIED WARRANTIES OF NONINFRINGEMENT,
+// MERCHANTABILITY, AND FITNESS FOR A PARTICULAR PURPOSE.
+//
+// Information and code furnished is believed to be accurate and reliable.
+// However, NVIDIA Corporation assumes no responsibility for the consequences of use of such
+// information or for any infringement of patents or other rights of third parties that may
+// result from its use. No license is granted by implication or otherwise under any patent
+// or patent rights of NVIDIA Corporation. Details are subject to change without notice.
+// This code supersedes and replaces all information previously supplied.
+// NVIDIA Corporation products are not authorized for use as critical
+// components in life support devices or systems without express written approval of
+// NVIDIA Corporation.
+//
+// Copyright (c) 2008-2015 NVIDIA Corporation. All rights reserved.
+
+// This file was generated by NvParameterized/scripts/GenParameterized.pl
+
+
+#ifndef HEADER_FieldSamplerPhysXMonitorParams_h
+#define HEADER_FieldSamplerPhysXMonitorParams_h
+
+#include "NvParametersTypes.h"
+
+#ifndef NV_PARAMETERIZED_ONLY_LAYOUTS
+#include "nvparameterized/NvParameterized.h"
+#include "nvparameterized/NvParameterizedTraits.h"
+#include "NvParameters.h"
+#include "NvTraitsInternal.h"
+#endif
+
+namespace nvidia
+{
+namespace fieldsampler
+{
+
+#if PX_VC
+#pragma warning(push)
+#pragma warning(disable: 4324) // structure was padded due to __declspec(align())
+#endif
+
+namespace FieldSamplerPhysXMonitorParamsNS
+{
+
+
+
+struct ParametersStruct
+{
+
+	uint32_t maxRBCount;
+	uint32_t maxPSCount;
+	uint32_t maxParticleCount;
+
+};
+
+static const uint32_t checksum[] = { 0x357a67e3, 0x9ed0cf4f, 0xd9fe5e4e, 0x532febb8, };
+
+} // namespace FieldSamplerPhysXMonitorParamsNS
+
+#ifndef NV_PARAMETERIZED_ONLY_LAYOUTS
+class FieldSamplerPhysXMonitorParams : public NvParameterized::NvParameters, public FieldSamplerPhysXMonitorParamsNS::ParametersStruct
+{
+public:
+	FieldSamplerPhysXMonitorParams(NvParameterized::Traits* traits, void* buf = 0, int32_t* refCount = 0);
+
+	virtual ~FieldSamplerPhysXMonitorParams();
+
+	virtual void destroy();
+
+	static const char* staticClassName(void)
+	{
+		return("FieldSamplerPhysXMonitorParams");
+	}
+
+	const char* className(void) const
+	{
+		return(staticClassName());
+	}
+
+	static const uint32_t ClassVersion = ((uint32_t)0 << 16) + (uint32_t)0;
+
+	static uint32_t staticVersion(void)
+	{
+		return ClassVersion;
+	}
+
+	uint32_t version(void) const
+	{
+		return(staticVersion());
+	}
+
+	static const uint32_t ClassAlignment = 8;
+
+	static const uint32_t* staticChecksum(uint32_t& bits)
+	{
+		bits = 8 * sizeof(FieldSamplerPhysXMonitorParamsNS::checksum);
+		return FieldSamplerPhysXMonitorParamsNS::checksum;
+	}
+
+	static void freeParameterDefinitionTable(NvParameterized::Traits* traits);
+
+	const uint32_t* checksum(uint32_t& bits) const
+	{
+		return staticChecksum(bits);
+	}
+
+	const FieldSamplerPhysXMonitorParamsNS::ParametersStruct& parameters(void) const
+	{
+		FieldSamplerPhysXMonitorParams* tmpThis = const_cast<FieldSamplerPhysXMonitorParams*>(this);
+		return *(static_cast<FieldSamplerPhysXMonitorParamsNS::ParametersStruct*>(tmpThis));
+	}
+
+	FieldSamplerPhysXMonitorParamsNS::ParametersStruct& parameters(void)
+	{
+		return *(static_cast<FieldSamplerPhysXMonitorParamsNS::ParametersStruct*>(this));
+	}
+
+	virtual NvParameterized::ErrorType getParameterHandle(const char* long_name, NvParameterized::Handle& handle) const;
+	virtual NvParameterized::ErrorType getParameterHandle(const char* long_name, NvParameterized::Handle& handle);
+
+	void initDefaults(void);
+
+protected:
+
+	virtual const NvParameterized::DefinitionImpl* getParameterDefinitionTree(void);
+	virtual const NvParameterized::DefinitionImpl* getParameterDefinitionTree(void) const;
+
+
+	virtual void getVarPtr(const NvParameterized::Handle& handle, void*& ptr, size_t& offset) const;
+
+private:
+
+	void buildTree(void);
+	void initDynamicArrays(void);
+	void initStrings(void);
+	void initReferences(void);
+	void freeDynamicArrays(void);
+	void freeStrings(void);
+	void freeReferences(void);
+
+	static bool mBuiltFlag;
+	static NvParameterized::MutexType mBuiltFlagMutex;
+};
+
+class FieldSamplerPhysXMonitorParamsFactory : public NvParameterized::Factory
+{
+	static const char* const vptr;
+
+public:
+
+	virtual void freeParameterDefinitionTable(NvParameterized::Traits* traits)
+	{
+		FieldSamplerPhysXMonitorParams::freeParameterDefinitionTable(traits);
+	}
+
+	virtual NvParameterized::Interface* create(NvParameterized::Traits* paramTraits)
+	{
+		// placement new on this class using mParameterizedTraits
+
+		void* newPtr = paramTraits->alloc(sizeof(FieldSamplerPhysXMonitorParams), FieldSamplerPhysXMonitorParams::ClassAlignment);
+		if (!NvParameterized::IsAligned(newPtr, FieldSamplerPhysXMonitorParams::ClassAlignment))
+		{
+			NV_PARAM_TRAITS_WARNING(paramTraits, "Unaligned memory allocation for class FieldSamplerPhysXMonitorParams");
+			paramTraits->free(newPtr);
+			return 0;
+		}
+
+		memset(newPtr, 0, sizeof(FieldSamplerPhysXMonitorParams)); // always initialize memory allocated to zero for default values
+		return NV_PARAM_PLACEMENT_NEW(newPtr, FieldSamplerPhysXMonitorParams)(paramTraits);
+	}
+
+	virtual NvParameterized::Interface* finish(NvParameterized::Traits* paramTraits, void* bufObj, void* bufStart, int32_t* refCount)
+	{
+		if (!NvParameterized::IsAligned(bufObj, FieldSamplerPhysXMonitorParams::ClassAlignment)
+		        || !NvParameterized::IsAligned(bufStart, FieldSamplerPhysXMonitorParams::ClassAlignment))
+		{
+			NV_PARAM_TRAITS_WARNING(paramTraits, "Unaligned memory allocation for class FieldSamplerPhysXMonitorParams");
+			return 0;
+		}
+
+		// Init NvParameters-part
+		// We used to call empty constructor of FieldSamplerPhysXMonitorParams here
+		// but it may call default constructors of members and spoil the data
+		NV_PARAM_PLACEMENT_NEW(bufObj, NvParameterized::NvParameters)(paramTraits, bufStart, refCount);
+
+		// Init vtable (everything else is already initialized)
+		*(const char**)bufObj = vptr;
+
+		return (FieldSamplerPhysXMonitorParams*)bufObj;
+	}
+
+	virtual const char* getClassName()
+	{
+		return (FieldSamplerPhysXMonitorParams::staticClassName());
+	}
+
+	virtual uint32_t getVersion()
+	{
+		return (FieldSamplerPhysXMonitorParams::staticVersion());
+	}
+
+	virtual uint32_t getAlignment()
+	{
+		return (FieldSamplerPhysXMonitorParams::ClassAlignment);
+	}
+
+	virtual const uint32_t* getChecksum(uint32_t& bits)
+	{
+		return (FieldSamplerPhysXMonitorParams::staticChecksum(bits));
+	}
+};
+#endif // NV_PARAMETERIZED_ONLY_LAYOUTS
+
+} // namespace fieldsampler
+} // namespace nvidia
+
+#if PX_VC
+#pragma warning(pop)
+#endif
+
+#endif
diff --git a/APEX_1.4/module/fieldsampler/include/autogen/ModuleFieldSamplerRegistration.h b/APEX_1.4/module/fieldsampler/include/autogen/ModuleFieldSamplerRegistration.h
new file mode 100644
index 00000000..62baddf6
--- /dev/null
+++ b/APEX_1.4/module/fieldsampler/include/autogen/ModuleFieldSamplerRegistration.h
@@ -0,0 +1,117 @@
+/*
+ * Copyright (c) 2008-2015, 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 MODULE_MODULEFIELDSAMPLERREGISTRATIONH_H
+#define MODULE_MODULEFIELDSAMPLERREGISTRATIONH_H
+
+#include "PsAllocator.h"
+#include "NvRegistrationsForTraitsBase.h"
+#include "nvparameterized/NvParameterizedTraits.h"
+#include "PxAssert.h"
+#include <stdint.h>
+
+// INCLUDE GENERATED FACTORIES
+#include "FieldSamplerPhysXMonitorParams.h"
+#include "FieldSamplerModuleParameters.h"
+
+
+// INCLUDE GENERATED CONVERSION
+
+
+namespace nvidia {
+namespace fieldsampler {
+
+
+class ModuleFieldSamplerRegistration : public NvParameterized::RegistrationsForTraitsBase
+{
+public:
+	static void invokeRegistration(NvParameterized::Traits* parameterizedTraits)
+	{
+		if (parameterizedTraits)
+		{
+			ModuleFieldSamplerRegistration().registerAll(*parameterizedTraits);
+		}
+	}
+
+	static void invokeUnregistration(NvParameterized::Traits* parameterizedTraits)
+	{
+		if (parameterizedTraits)
+		{
+			ModuleFieldSamplerRegistration().unregisterAll(*parameterizedTraits);
+		}
+	}
+
+	void registerAvailableFactories(NvParameterized::Traits& parameterizedTraits)
+	{
+		::NvParameterized::Factory* factoriesToRegister[] = {
+// REGISTER GENERATED FACTORIES
+			new nvidia::fieldsampler::FieldSamplerPhysXMonitorParamsFactory(),
+			new nvidia::fieldsampler::FieldSamplerModuleParametersFactory(),
+
+		};
+
+		for (size_t i = 0; i < sizeof(factoriesToRegister)/sizeof(factoriesToRegister[0]); ++i)
+		{
+			parameterizedTraits.registerFactory(*factoriesToRegister[i]);
+		}
+	}
+
+	virtual void registerAvailableConverters(NvParameterized::Traits& parameterizedTraits)
+	{
+// REGISTER GENERATED CONVERSION
+PX_UNUSED(parameterizedTraits);
+
+	}
+
+	void unregisterAvailableFactories(NvParameterized::Traits& parameterizedTraits)
+	{
+		struct FactoryDesc
+		{
+			const char* name;
+			uint32_t version;
+		};
+
+		::NvParameterized::Factory* factoriesToUnregister[] = {
+// UNREGISTER GENERATED FACTORIES
+			new nvidia::fieldsampler::FieldSamplerPhysXMonitorParamsFactory(),
+			new nvidia::fieldsampler::FieldSamplerModuleParametersFactory(),
+
+		};
+
+		for (size_t i = 0; i < sizeof(factoriesToUnregister)/sizeof(factoriesToUnregister[0]); ++i)
+		{
+			::NvParameterized::Factory* removedFactory = parameterizedTraits.removeFactory(factoriesToUnregister[i]->getClassName(), factoriesToUnregister[i]->getVersion());
+			if (!removedFactory) 
+			{
+				PX_ASSERT_WITH_MESSAGE(0, "Factory can not be removed!");
+			}
+			else
+			{
+				removedFactory->freeParameterDefinitionTable(&parameterizedTraits);
+				delete removedFactory;
+				delete factoriesToUnregister[i];
+			}
+		}
+	}
+
+	virtual void unregisterAvailableConverters(NvParameterized::Traits& parameterizedTraits)
+	{
+// UNREGISTER GENERATED CONVERSION
+PX_UNUSED(parameterizedTraits);
+
+	}
+
+};
+
+
+}
+} //nvidia::fieldsampler
+
+#endif
diff --git a/APEX_1.4/module/fieldsampler/src/FieldBoundaryWrapper.cpp b/APEX_1.4/module/fieldsampler/src/FieldBoundaryWrapper.cpp
new file mode 100644
index 00000000..ebeedc80
--- /dev/null
+++ b/APEX_1.4/module/fieldsampler/src/FieldBoundaryWrapper.cpp
@@ -0,0 +1,45 @@
+/*
+ * Copyright (c) 2008-2015, 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 "ApexDefs.h"
+#include "Apex.h"
+#include "FieldBoundaryWrapper.h"
+#include "FieldSamplerManager.h"
+
+
+namespace nvidia
+{
+namespace fieldsampler
+{
+
+FieldBoundaryWrapper::FieldBoundaryWrapper(ResourceList& list, FieldSamplerManager* manager, FieldBoundaryIntl* fieldBoundary, const FieldBoundaryDescIntl& fieldBoundaryDesc)
+	: mManager(manager)
+	, mFieldBoundary(fieldBoundary)
+	, mFieldBoundaryDesc(fieldBoundaryDesc)
+	, mFieldShapesChanged(false)
+{
+	list.add(*this);
+
+}
+
+void FieldBoundaryWrapper::release()
+{
+	delete this;
+}
+
+void FieldBoundaryWrapper::update()
+{
+	mFieldShapesChanged = mFieldBoundary->updateFieldBoundary(mFieldShapes);
+}
+
+}
+} // end namespace nvidia::apex
+
diff --git a/APEX_1.4/module/fieldsampler/src/FieldSamplerManager.cpp b/APEX_1.4/module/fieldsampler/src/FieldSamplerManager.cpp
new file mode 100644
index 00000000..a7c83f88
--- /dev/null
+++ b/APEX_1.4/module/fieldsampler/src/FieldSamplerManager.cpp
@@ -0,0 +1,329 @@
+/*
+ * Copyright (c) 2008-2015, 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 "ApexDefs.h"
+#include "Apex.h"
+
+#include "FieldSamplerManager.h"
+
+#include "FieldSamplerQuery.h"
+#include "FieldSamplerSceneWrapper.h"
+#include "FieldSamplerWrapper.h"
+#include "FieldBoundaryWrapper.h"
+
+#include "SceneIntl.h"
+
+namespace nvidia
+{
+namespace fieldsampler
+{
+
+FieldSamplerManager::FieldSamplerManager(FieldSamplerScene* scene)
+	: mScene(scene)
+	, mFieldSamplerGroupsFilteringChanged(true)
+{
+	mWeightedCallback = NULL;
+}
+
+
+PX_INLINE void FieldSamplerManager::addFieldSamplerToQuery(FieldSamplerWrapper* fieldSamplerWrapper, FieldSamplerQuery* query)
+{
+	if (query->addFieldSampler(fieldSamplerWrapper))
+	{
+		fieldSamplerWrapper->mQueryRefCount += 1;
+	}
+}
+
+void FieldSamplerManager::addAllFieldSamplersToQuery(FieldSamplerQuery* query) const
+{
+	for (uint32_t i = 0; i < mFieldSamplerWrapperList.getSize(); ++i)
+	{
+		addFieldSamplerToQuery(static_cast<FieldSamplerWrapper*>(mFieldSamplerWrapperList.getResource(i)), query);
+	}
+}
+
+void FieldSamplerManager::submitTasks()
+{
+	if (mFieldSamplerGroupsFilteringChanged)
+	{
+		mFieldSamplerGroupsFilteringChanged = false;
+
+		//clear queryRefCounts
+		for (uint32_t i = 0; i < mFieldSamplerWrapperList.getSize(); ++i)
+		{
+			FieldSamplerWrapper* fieldSamplerWrapper = static_cast<FieldSamplerWrapper*>(mFieldSamplerWrapperList.getResource(i));
+			fieldSamplerWrapper->mQueryRefCount = 0;
+		}
+
+		//rebuild all connection based on changed collision checking
+		for (uint32_t i = 0; i < mFieldSamplerQueryList.getSize(); ++i)
+		{
+			FieldSamplerQuery* query = DYNAMIC_CAST(FieldSamplerQuery*)(mFieldSamplerQueryList.getResource(i));
+			query->clearAllFieldSamplers();
+			addAllFieldSamplersToQuery(query);
+		}
+	}
+
+	for (uint32_t i = 0; i < mFieldSamplerQueryList.getSize(); ++i)
+	{
+		FieldSamplerQuery* query = DYNAMIC_CAST(FieldSamplerQuery*)(mFieldSamplerQueryList.getResource(i));
+		query->submitTasks();
+	}
+}
+
+void FieldSamplerManager::setTaskDependencies()
+{
+	for (uint32_t i = 0; i < mFieldSamplerQueryList.getSize(); ++i)
+	{
+		FieldSamplerQuery* query = DYNAMIC_CAST(FieldSamplerQuery*)(mFieldSamplerQueryList.getResource(i));
+		query->setTaskDependencies();
+	}
+
+	//update
+	for (uint32_t i = 0; i < mFieldBoundaryWrapperList.getSize(); ++i)
+	{
+		FieldBoundaryWrapper* wrapper = DYNAMIC_CAST(FieldBoundaryWrapper*)(mFieldBoundaryWrapperList.getResource(i));
+		wrapper->update();
+	}
+	for (uint32_t i = 0; i < mFieldSamplerWrapperList.getSize(); ++i)
+	{
+		FieldSamplerWrapper* wrapper = DYNAMIC_CAST(FieldSamplerWrapper*)(mFieldSamplerWrapperList.getResource(i));
+		wrapper->update();
+	}
+	for (uint32_t i = 0; i < mFieldSamplerSceneWrapperList.getSize(); ++i)
+	{
+		FieldSamplerSceneWrapper* wrapper = DYNAMIC_CAST(FieldSamplerSceneWrapper*)(mFieldSamplerSceneWrapperList.getResource(i));
+		wrapper->update();
+	}
+	for (uint32_t i = 0; i < mFieldSamplerQueryList.getSize(); ++i)
+	{
+		FieldSamplerQuery* query = DYNAMIC_CAST(FieldSamplerQuery*)(mFieldSamplerQueryList.getResource(i));
+		query->update();
+	}
+
+	//postUpdate
+	for (uint32_t i = 0; i < mFieldSamplerSceneWrapperList.getSize(); ++i)
+	{
+		FieldSamplerSceneWrapper* wrapper = DYNAMIC_CAST(FieldSamplerSceneWrapper*)(mFieldSamplerSceneWrapperList.getResource(i));
+		wrapper->postUpdate();
+	}
+}
+
+void FieldSamplerManager::fetchResults()
+{
+	for (uint32_t i = 0; i < mFieldSamplerQueryList.getSize(); ++i)
+	{
+		FieldSamplerQuery* query = DYNAMIC_CAST(FieldSamplerQuery*)(mFieldSamplerQueryList.getResource(i));
+		query->fetchResults();
+	}
+}
+
+
+FieldSamplerQueryIntl* FieldSamplerManager::createFieldSamplerQuery(const FieldSamplerQueryDescIntl& desc)
+{
+	FieldSamplerQuery* query = allocateFieldSamplerQuery(desc);
+	if (query)
+	{
+		addAllFieldSamplersToQuery(query);
+	}
+	return query;
+}
+
+void FieldSamplerManager::registerFieldSampler(FieldSamplerIntl* fieldSampler, const FieldSamplerDescIntl& fieldSamplerDesc, FieldSamplerSceneIntl* fieldSamplerScene)
+{
+	FieldSamplerSceneWrapper* fieldSamplerSceneWrapper = NULL;
+	//find FieldSamplerSceneWrapper
+	for (uint32_t i = 0; i < mFieldSamplerSceneWrapperList.getSize(); ++i)
+	{
+		FieldSamplerSceneWrapper* wrapper = DYNAMIC_CAST(FieldSamplerSceneWrapper*)(mFieldSamplerSceneWrapperList.getResource(i));
+		if (wrapper->getInternalFieldSamplerScene() == fieldSamplerScene)
+		{
+			fieldSamplerSceneWrapper = wrapper;
+			break;
+		}
+	}
+	if (fieldSamplerSceneWrapper == NULL)
+	{
+		fieldSamplerSceneWrapper = allocateFieldSamplerSceneWrapper(fieldSamplerScene);
+	}
+	PX_ASSERT(fieldSamplerSceneWrapper != NULL);
+
+	FieldSamplerWrapper* fieldSamplerWrapper = allocateFieldSamplerWrapper(fieldSampler, fieldSamplerDesc, fieldSamplerSceneWrapper);
+	PX_ASSERT(fieldSamplerWrapper != NULL);
+
+	// add all mFieldBoundaryWrapperList
+	for (uint32_t i = 0; i < mFieldBoundaryWrapperList.getSize(); ++i)
+	{
+		FieldBoundaryWrapper* wrapper = DYNAMIC_CAST(FieldBoundaryWrapper*)(mFieldBoundaryWrapperList.getResource(i));
+		fieldSamplerWrapper->addFieldBoundary(wrapper);
+	}
+
+	for (uint32_t i = 0; i < mFieldSamplerQueryList.getSize(); ++i)
+	{
+		FieldSamplerQuery* query = DYNAMIC_CAST(FieldSamplerQuery*)(mFieldSamplerQueryList.getResource(i));
+		addFieldSamplerToQuery(fieldSamplerWrapper, query);
+	}
+}
+
+void FieldSamplerManager::unregisterFieldSampler(FieldSamplerIntl* fieldSampler)
+{
+	FieldSamplerWrapper* fieldSamplerWrapper = NULL;
+	//find FieldSamplerWrapper
+	for (uint32_t i = 0; i < mFieldSamplerWrapperList.getSize(); ++i)
+	{
+		FieldSamplerWrapper* wrapper = static_cast<FieldSamplerWrapper*>(mFieldSamplerWrapperList.getResource(i));
+		if (wrapper->getInternalFieldSampler() == fieldSampler)
+		{
+			fieldSamplerWrapper = wrapper;
+			break;
+		}
+	}
+	if (fieldSamplerWrapper != NULL)
+	{
+		for (uint32_t i = 0; i < mFieldSamplerQueryList.getSize(); ++i)
+		{
+			FieldSamplerQuery* query = DYNAMIC_CAST(FieldSamplerQuery*)(mFieldSamplerQueryList.getResource(i));
+			query->removeFieldSampler(fieldSamplerWrapper);
+		}
+		fieldSamplerWrapper->release();
+	}
+}
+
+void FieldSamplerManager::registerFieldBoundary(FieldBoundaryIntl* fieldBoundary, const FieldBoundaryDescIntl& fieldBoundaryDesc)
+{
+	FieldBoundaryWrapper* fieldBoundaryWrapper = PX_NEW(FieldBoundaryWrapper)(mFieldBoundaryWrapperList, this, fieldBoundary, fieldBoundaryDesc);
+	if (fieldBoundaryWrapper)
+	{
+		for (uint32_t i = 0; i < mFieldSamplerWrapperList.getSize(); ++i)
+		{
+			static_cast<FieldSamplerWrapper*>(mFieldSamplerWrapperList.getResource(i))->addFieldBoundary(fieldBoundaryWrapper);
+		}
+	}
+}
+void FieldSamplerManager::unregisterFieldBoundary(FieldBoundaryIntl* fieldBoundary)
+{
+	FieldBoundaryWrapper* fieldBoundaryWrapper = 0;
+	for (uint32_t i = 0; i < mFieldBoundaryWrapperList.getSize(); ++i)
+	{
+		FieldBoundaryWrapper* wrapper = static_cast<FieldBoundaryWrapper*>(mFieldBoundaryWrapperList.getResource(i));
+		if (wrapper->getInternalFieldBoundary() == fieldBoundary)
+		{
+			fieldBoundaryWrapper = wrapper;
+			break;
+		}
+	}
+	if (fieldBoundaryWrapper != 0)
+	{
+		for (uint32_t i = 0; i < mFieldSamplerWrapperList.getSize(); ++i)
+		{
+			static_cast<FieldSamplerWrapper*>(mFieldSamplerWrapperList.getResource(i))->removeFieldBoundary(fieldBoundaryWrapper);
+		}
+		fieldBoundaryWrapper->release();
+	}
+}
+
+void FieldSamplerManager::registerUnhandledParticleSystem(physx::PxActor* actor)
+{
+	if (!isUnhandledParticleSystem(actor))
+	{
+		mUnhandledParticleSystems.pushBack(actor);
+	}
+}
+
+void FieldSamplerManager::unregisterUnhandledParticleSystem(physx::PxActor* actor)
+{
+	mUnhandledParticleSystems.findAndReplaceWithLast(actor);
+}
+
+bool FieldSamplerManager::isUnhandledParticleSystem(physx::PxActor* actor)
+{
+	for (uint32_t i = 0; i < mUnhandledParticleSystems.size(); i++)
+	if (mUnhandledParticleSystems[i] == actor)
+	{
+		return true;
+	}
+	return false;
+}
+
+bool FieldSamplerManager::getFieldSamplerGroupsFiltering(const PxFilterData &o1,const PxFilterData &o2,float &weight) const
+{
+	bool ret = true;
+
+	weight = 1.0f;
+	if ( mScene )
+	{
+		if ( mWeightedCallback )
+		{
+			ret = mWeightedCallback->fieldSamplerWeightedCollisionFilter(o1,o2,weight);
+		}
+		else
+		{
+			physx::PxScene *scene = mScene->getModulePhysXScene();
+			scene->lockRead(__FILE__,__LINE__);
+			physx::PxSimulationFilterShader shader = scene->getFilterShader();
+			scene->unlockRead();
+			if ( shader )
+			{
+				physx::PxFilterObjectAttributes atr0 = 0;
+				physx::PxFilterObjectAttributes atr1 = 0;
+				physx::PxPairFlags pairFlags;
+				physx::PxFilterFlags result = (*shader)(atr0,o1,atr1,o2,pairFlags,NULL,0);
+				if ( result & (physx::PxFilterFlag::eKILL | physx::PxFilterFlag::eSUPPRESS) )
+				{
+					ret = false;
+				}
+			}
+		}
+	}
+	return ret;
+}
+
+
+
+
+/******************************** CPU Version ********************************/
+
+FieldSamplerQuery* FieldSamplerManagerCPU::allocateFieldSamplerQuery(const FieldSamplerQueryDescIntl& desc)
+{
+	return PX_NEW(FieldSamplerQueryCPU)(desc, mFieldSamplerQueryList, this);
+}
+FieldSamplerSceneWrapper* FieldSamplerManagerCPU::allocateFieldSamplerSceneWrapper(FieldSamplerSceneIntl* fieldSamplerScene)
+{
+	return PX_NEW(FieldSamplerSceneWrapperCPU)(mFieldSamplerSceneWrapperList, this, fieldSamplerScene);
+}
+FieldSamplerWrapper* FieldSamplerManagerCPU::allocateFieldSamplerWrapper(FieldSamplerIntl* fieldSampler, const FieldSamplerDescIntl& fieldSamplerDesc, FieldSamplerSceneWrapper* fieldSamplerSceneWrapper)
+{
+	return PX_NEW(FieldSamplerWrapperCPU)(mFieldSamplerWrapperList, this, fieldSampler, fieldSamplerDesc, fieldSamplerSceneWrapper);
+}
+
+/******************************** GPU Version ********************************/
+
+#if APEX_CUDA_SUPPORT
+
+FieldSamplerQuery* FieldSamplerManagerGPU::allocateFieldSamplerQuery(const FieldSamplerQueryDescIntl& desc)
+{
+	return PX_NEW(FieldSamplerQueryGPU)(desc, mFieldSamplerQueryList, this);
+}
+FieldSamplerSceneWrapper* FieldSamplerManagerGPU::allocateFieldSamplerSceneWrapper(FieldSamplerSceneIntl* fieldSamplerScene)
+{
+	return PX_NEW(FieldSamplerSceneWrapperGPU)(mFieldSamplerSceneWrapperList, this, fieldSamplerScene);
+}
+FieldSamplerWrapper* FieldSamplerManagerGPU::allocateFieldSamplerWrapper(FieldSamplerIntl* fieldSampler, const FieldSamplerDescIntl& fieldSamplerDesc, FieldSamplerSceneWrapper* fieldSamplerSceneWrapper)
+{
+	return PX_NEW(FieldSamplerWrapperGPU)(mFieldSamplerWrapperList, this, fieldSampler, fieldSamplerDesc, fieldSamplerSceneWrapper);
+}
+
+#endif // APEX_CUDA_SUPPORT
+
+}
+} // end namespace nvidia::apex
+
+
diff --git a/APEX_1.4/module/fieldsampler/src/FieldSamplerPhysXMonitor.cpp b/APEX_1.4/module/fieldsampler/src/FieldSamplerPhysXMonitor.cpp
new file mode 100644
index 00000000..e75c8700
--- /dev/null
+++ b/APEX_1.4/module/fieldsampler/src/FieldSamplerPhysXMonitor.cpp
@@ -0,0 +1,354 @@
+/*
+ * Copyright (c) 2008-2015, 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 "ApexDefs.h"
+
+#if PX_PHYSICS_VERSION_MAJOR == 3
+
+#include "ScopedPhysXLock.h"
+
+#include "PsSort.h"
+#include "FieldSamplerPhysXMonitor.h"
+#include "FieldSamplerScene.h"
+#include "FieldSamplerManager.h"
+#include "FieldSamplerQueryIntl.h"
+
+#include "extensions/PxShapeExt.h"
+
+namespace nvidia
+{
+namespace fieldsampler
+{
+
+	using namespace physx;
+
+#pragma warning(disable: 4355) // 'this' : used in base member initializer list
+
+FieldSamplerPhysXMonitor::FieldSamplerPhysXMonitor(FieldSamplerScene& scene)
+	: mFieldSamplerScene(&scene)
+	, mNumPS(0)
+	, mNumRB(0)
+	, mEnable(false)
+	, mTaskRunAfterActorUpdate(*this)
+{
+	mFilterData.setToDefault();
+
+	mScene = mFieldSamplerScene->getModulePhysXScene();
+	mParams = static_cast<FieldSamplerPhysXMonitorParams*>(GetInternalApexSDK()->getParameterizedTraits()->createNvParameterized(FieldSamplerPhysXMonitorParams::staticClassName()));
+	mFieldSamplerManager = DYNAMIC_CAST(FieldSamplerManager*)(mFieldSamplerScene->getManager());
+
+	mRBIndex.reserve(mParams->maxRBCount);
+	mParticleSystems.resize(mParams->maxPSCount);
+
+	mPSOutField.resize(mParams->maxParticleCount);
+	mOutVelocities.resize(mParams->maxParticleCount);
+	mOutIndices.resize(mParams->maxParticleCount);
+
+	mRBOutField.resize(mParams->maxRBCount);
+
+	mRBActors.resize(mParams->maxRBCount);
+}
+
+
+FieldSamplerPhysXMonitor::~FieldSamplerPhysXMonitor()
+{
+	for (uint32_t i = 0; i < mPSFieldSamplerQuery.size(); i++)
+	{
+		if (mPSFieldSamplerQuery[i])
+		{
+			mPSFieldSamplerQuery[i]->release();
+		}
+	}
+	for (uint32_t i = 0; i < mRBFieldSamplerQuery.size(); i++)
+	{
+		if (mRBFieldSamplerQuery[i])
+		{
+			mRBFieldSamplerQuery[i]->release();
+		}
+	}
+	if(mParams)
+	{
+		mParams->destroy();
+	}
+}
+
+
+void FieldSamplerPhysXMonitor::setPhysXScene(PxScene* scene)
+{
+	mScene = scene;
+}
+
+
+void FieldSamplerPhysXMonitor::getParticles(uint32_t taskId)
+{
+	SCOPED_PHYSX_LOCK_READ(&mFieldSamplerScene->getApexScene());
+	float deltaTime = mFieldSamplerScene->getApexScene().getPhysXSimulateTime();
+	mPCount = 0;
+	mNumPS = mScene->getActors(physx::PxActorTypeFlag::ePARTICLE_SYSTEM, &mParticleSystems[0], mParams->maxPSCount);
+	for(uint32_t i = 0; i < mNumPS; i++)
+	if (!mFieldSamplerManager->isUnhandledParticleSystem(mParticleSystems[i]) && mParticleSystems[i]->is<PxParticleBase>())
+	{
+		if (mPSFieldSamplerQuery.size() == i)
+		{
+			FieldSamplerQueryDescIntl queryDesc;
+			queryDesc.maxCount = mParams->maxParticleCount;
+			queryDesc.samplerFilterData = mFilterData;
+			mPSFieldSamplerQuery.pushBack( mFieldSamplerManager->createFieldSamplerQuery(queryDesc) );
+			mPSFieldSamplerTaskID.pushBack(0);
+			mParticleReadData.pushBack(0);
+			mPSMass.pushBack(0.f);
+		}
+
+		PxParticleSystem* particleSystem = DYNAMIC_CAST(PxParticleSystem*)((mParticleSystems[i]));
+		mParticleReadData[i] = particleSystem->lockParticleReadData();
+		uint32_t numParticles;
+		if (mParticleReadData[i])
+		{
+			numParticles = mParticleReadData[i]->validParticleRange;
+			if(mPCount + numParticles >= mParams->maxParticleCount) break;
+		
+			FieldSamplerQueryDataIntl queryData;
+			queryData.timeStep = deltaTime;
+			queryData.count = numParticles;
+			queryData.isDataOnDevice = false;
+			
+			//hack for PhysX particle stride calculation
+			physx::PxStrideIterator<const PxVec3> positionIt(mParticleReadData[i]->positionBuffer);
+#ifdef WIN64
+			queryData.positionStrideBytes = (uint32_t)(-(int64_t)&*positionIt + (int64_t)&*(++positionIt));
+#else
+			queryData.positionStrideBytes = (uint32_t)(-(int32_t)&*positionIt + (int32_t)&*(++positionIt));
+#endif
+			queryData.velocityStrideBytes = queryData.positionStrideBytes;
+			queryData.massStrideBytes = 0;
+			queryData.pmaInPosition = (float*)&*(mParticleReadData[i]->positionBuffer);
+			queryData.pmaInVelocity = (float*)&*(mParticleReadData[i]->velocityBuffer);
+			queryData.pmaInIndices = 0;
+			mPSMass[i] = particleSystem->getParticleMass();
+			queryData.pmaInMass = &mPSMass[i];
+			queryData.pmaOutField = &mPSOutField[mPCount];
+			mPSFieldSamplerTaskID[i] = mPSFieldSamplerQuery[i]->submitFieldSamplerQuery(queryData, taskId);
+
+			mPCount += numParticles;
+		}
+	}
+}
+
+
+void FieldSamplerPhysXMonitor::updateParticles()
+{
+	uint32_t pCount = 0;
+	SCOPED_PHYSX_LOCK_WRITE(mFieldSamplerScene->getApexScene().getPhysXScene());
+	for(uint32_t i = 0; i < mNumPS; i++)
+	if (!mFieldSamplerManager->isUnhandledParticleSystem(mParticleSystems[i]))
+	{
+		PxParticleSystem* particleSystem = DYNAMIC_CAST(PxParticleSystem*)((mParticleSystems[i]));
+		uint32_t numParticles = PxMin(mParticleReadData[i]->validParticleRange, mParams->maxParticleCount);
+
+		uint32_t numUpdates = 0;
+
+		if (numParticles > 0)
+		{
+			for (uint32_t w = 0; w <= (mParticleReadData[i]->validParticleRange-1) >> 5; w++)
+			{
+				for (uint32_t b = mParticleReadData[i]->validParticleBitmap[w]; b; b &= b-1) 
+				{
+					uint32_t index = (w << 5 | shdfnd::lowestSetBit(b));
+
+					PxVec3 diffVel = mPSOutField[pCount + index].getXYZ();
+					if (!diffVel.isZero())
+					{
+						const PxVec3& sourceVelocity = mParticleReadData[i]->velocityBuffer[index];
+						mOutVelocities[numUpdates] = sourceVelocity + diffVel;
+						mOutIndices[numUpdates] = index;
+						numUpdates++;
+					}
+				}
+			}
+		}
+		pCount += numParticles;
+		// return ownership of the buffers back to the SDK
+		mParticleReadData[i]->unlock();
+	
+		if(pCount <= mParams->maxParticleCount && numUpdates > 0)
+		{
+			physx::PxStrideIterator<uint32_t> indices(&mOutIndices[0]);
+			physx::PxStrideIterator<PxVec3> outVelocities(&mOutVelocities[0]);
+			particleSystem->setVelocities(numUpdates, indices, outVelocities);
+		}		
+	}
+}
+
+
+void FieldSamplerPhysXMonitor::getRigidBodies(uint32_t taskId)
+{
+	SCOPED_PHYSX_LOCK_READ(&mFieldSamplerScene->getApexScene());
+	float deltaTime = mFieldSamplerScene->getApexScene().getPhysXSimulateTime();
+
+	FieldSamplerQueryDataIntl queryData;
+	queryData.timeStep = deltaTime;
+	queryData.pmaInIndices = 0;
+
+	uint32_t rbCount = mScene->getActors(physx::PxActorTypeFlag::eRIGID_DYNAMIC, &mRBActors[0], mParams->maxRBCount);
+	Array<PxShape*> shapes;
+	mNumRB = 0;
+	mRBIndex.clear();
+	float weight = 1.f;
+
+	for (uint32_t i = 0; i < rbCount; i++)
+	{
+		physx::PxRigidDynamic* rb = (physx::PxRigidDynamic*)mRBActors[i];
+		if (rb->getRigidBodyFlags() == PxRigidBodyFlag::eKINEMATIC) 
+		{
+			continue;
+		}
+		const PxVec3& cmPos = rb->getGlobalPose().p;
+		const PxVec3& velocity = rb->getLinearVelocity();
+		const PxVec3& rotation = rb->getAngularVelocity();
+		float mass = rb->getMass();
+
+		const uint32_t numShapes = rb->getNbShapes();
+		shapes.resize(numShapes);
+		if (numShapes == 0) 
+		{
+			continue;
+		}
+		rb->getShapes(&shapes[0], numShapes);
+		for (uint32_t j = 0; j < numShapes && mNumRB < mParams->maxRBCount; j++)
+		{
+			PxFilterData filterData = shapes[j]->getQueryFilterData();
+			if (mFieldSamplerManager->getFieldSamplerGroupsFiltering(mFilterData, filterData, weight))
+			{
+				PxFilterData* current = mRBFilterData.find(filterData);
+				if (current == mRBFilterData.end())
+				{
+					mRBFilterData.pushBack(filterData);
+					current = &mRBFilterData.back();
+					FieldSamplerQueryDescIntl queryDesc;
+					queryDesc.maxCount = mParams->maxParticleCount;
+					queryDesc.samplerFilterData = filterData;
+					mRBFieldSamplerQuery.pushBack( mFieldSamplerManager->createFieldSamplerQuery(queryDesc) );
+				}
+
+				ShapeData* sd = PX_NEW(ShapeData)();
+				sd->fdIndex = (uint32_t)(current - &mRBFilterData[0]);
+				sd->rbIndex = i;
+				sd->mass = mass / numShapes;
+				sd->pos = PxShapeExt::getWorldBounds(*shapes[j], *rb).getCenter();
+				sd->vel = velocity + rotation.cross(sd->pos - cmPos);
+				mRBIndex.pushBack(sd);
+				++mNumRB;
+			}
+		}		
+	}
+
+	if (mNumRB == 0) 
+	{
+		return;
+	}
+
+	
+	sort(&mRBIndex[0], mNumRB, ShapeData::sortPredicate);
+
+	mRBInPosition.resize(mNumRB);
+	mRBInVelocity.resize(mNumRB);
+
+	uint32_t current(mRBIndex[0]->fdIndex);
+	uint32_t currentCount = 0;
+	uint32_t fdCount = 0;
+	for (uint32_t i = 0; i <= mNumRB; i++)
+	{
+		if (i == mNumRB || current != mRBIndex[i]->fdIndex)
+		{
+			queryData.count = currentCount;
+			queryData.isDataOnDevice = false;
+			queryData.massStrideBytes = sizeof(PxVec4);
+			queryData.positionStrideBytes = sizeof(PxVec4);
+			queryData.velocityStrideBytes = queryData.positionStrideBytes;
+
+			queryData.pmaInPosition = (float*)(&mRBInPosition[fdCount]);
+			queryData.pmaInVelocity = (float*)(&mRBInVelocity[fdCount]);
+			queryData.pmaInMass = &(mRBInPosition[fdCount].w);
+			queryData.pmaOutField = &mRBOutField[fdCount];
+	
+			mRBFieldSamplerQuery[current]->submitFieldSamplerQuery(queryData, taskId);
+			
+			fdCount += currentCount;
+			if (i != mNumRB) 
+			{
+				current = mRBIndex[i]->fdIndex;
+				currentCount = 1;
+			}
+		}
+		else
+		{
+			currentCount++;
+		}
+		if (i < mNumRB)
+		{
+			mRBInPosition[i] = PxVec4 (mRBIndex[i]->pos, mRBIndex[i]->mass);
+			mRBInVelocity[i] = PxVec4 (mRBIndex[i]->vel, 0.f);
+		}
+	}
+}
+
+void FieldSamplerPhysXMonitor::updateRigidBodies()
+{
+	SCOPED_PHYSX_LOCK_WRITE(&mFieldSamplerScene->getApexScene());
+
+	for(uint32_t i = 0; i < mNumRB; i++)
+	{
+		physx::PxRigidDynamic* rb = (physx::PxRigidDynamic*)mRBActors[mRBIndex[i]->rbIndex];
+		const PxVec3 velocity = mRBOutField[i].getXYZ();
+		const PxVec3 rotation = mRBOutField[i].getXYZ().cross(mRBInPosition[i].getXYZ() - rb->getGlobalPose().p);
+		if (!velocity.isZero() || !rotation.isZero())
+		{
+			rb->setLinearVelocity(rb->getLinearVelocity() + velocity);
+			rb->setAngularVelocity(rb->getAngularVelocity() + rotation);
+		}
+
+		PX_DELETE(mRBIndex[i]);
+	}
+}
+
+
+void FieldSamplerPhysXMonitor::update()
+{	
+	PxTaskManager* tm = mFieldSamplerScene->getApexScene().getTaskManager();
+	uint32_t taskId = tm->getNamedTask(FSST_PHYSX_MONITOR_LOAD);
+	if(mScene)
+	{
+		getParticles(taskId);
+
+		getRigidBodies(taskId);
+
+		//getCloth(task);
+	}
+	if(mNumPS > 0 || mNumRB > 0)
+	{	
+		tm->submitNamedTask(&mTaskRunAfterActorUpdate, FSST_PHYSX_MONITOR_UPDATE);
+		mTaskRunAfterActorUpdate.startAfter(tm->getNamedTask(FSST_PHYSX_MONITOR_FETCH));
+		mTaskRunAfterActorUpdate.finishBefore(tm->getNamedTask(AST_PHYSX_SIMULATE));
+	}	
+}
+
+void FieldSamplerPhysXMonitor::updatePhysX()
+{
+	updateParticles();
+	updateRigidBodies();
+}
+
+
+}
+} // end namespace nvidia::apex
+
+#endif
+
diff --git a/APEX_1.4/module/fieldsampler/src/FieldSamplerQuery.cpp b/APEX_1.4/module/fieldsampler/src/FieldSamplerQuery.cpp
new file mode 100644
index 00000000..bada10fc
--- /dev/null
+++ b/APEX_1.4/module/fieldsampler/src/FieldSamplerQuery.cpp
@@ -0,0 +1,876 @@
+/*
+ * Copyright (c) 2008-2015, 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 "ApexDefs.h"
+#include "Apex.h"
+#include "FieldSamplerQuery.h"
+#include "FieldSamplerManager.h"
+#include "FieldSamplerWrapper.h"
+#include "FieldSamplerSceneWrapper.h"
+#include "FieldBoundaryWrapper.h"
+
+#include "SceneIntl.h"
+
+#if APEX_CUDA_SUPPORT
+#include "PxGpuTask.h"
+#endif
+
+#include "FieldSamplerCommon.h"
+
+
+namespace nvidia
+{
+namespace fieldsampler
+{
+
+
+FieldSamplerQuery::FieldSamplerQuery(const FieldSamplerQueryDescIntl& desc, ResourceList& list, FieldSamplerManager* manager)
+	: mManager(manager)
+	, mQueryDesc(desc)
+	, mAccumVelocity(manager->getApexScene(), PX_ALLOC_INFO("mAccumVelocity", PARTICLES))
+	, mOnStartCallback(NULL)
+	, mOnFinishCallback(NULL)
+{
+	list.add(*this);
+}
+
+void FieldSamplerQuery::release()
+{
+	if (mInRelease)
+	{
+		return;
+	}
+	mInRelease = true;
+	destroy();
+}
+
+void FieldSamplerQuery::destroy()
+{
+	delete this;
+}
+
+
+FieldSamplerQuery::SceneInfo* FieldSamplerQuery::findSceneInfo(FieldSamplerSceneWrapper* sceneWrapper) const
+{
+	for (uint32_t i = 0; i < mSceneList.getSize(); ++i)
+	{
+		SceneInfo* sceneInfo = DYNAMIC_CAST(SceneInfo*)(mSceneList.getResource(i));
+		if (sceneInfo->getSceneWrapper() == sceneWrapper)
+		{
+			return sceneInfo;
+		}
+	}
+	return NULL;
+}
+
+
+bool FieldSamplerQuery::addFieldSampler(FieldSamplerWrapper* fieldSamplerWrapper)
+{
+	const FieldSamplerDescIntl& fieldSamplerDesc = fieldSamplerWrapper->getInternalFieldSamplerDesc();
+	float multiplier = 1.0f;
+	bool result = mManager->getFieldSamplerGroupsFiltering(mQueryDesc.samplerFilterData, fieldSamplerDesc.samplerFilterData, multiplier);
+	if (result)
+	{
+		FieldSamplerSceneWrapper* sceneWrapper = fieldSamplerWrapper->getFieldSamplerSceneWrapper();
+		SceneInfo* sceneInfo = findSceneInfo(sceneWrapper);
+		if (sceneInfo == NULL)
+		{
+			sceneInfo = createSceneInfo(sceneWrapper);
+		}
+		sceneInfo->addFieldSampler(fieldSamplerWrapper, multiplier);
+	}
+	return result;
+}
+
+bool FieldSamplerQuery::removeFieldSampler(FieldSamplerWrapper* fieldSamplerWrapper)
+{
+	FieldSamplerSceneWrapper* sceneWrapper = fieldSamplerWrapper->getFieldSamplerSceneWrapper();
+	SceneInfo* sceneInfo = findSceneInfo(sceneWrapper);
+	return (sceneInfo != NULL) ? sceneInfo->removeFieldSampler(fieldSamplerWrapper) : false;
+}
+
+void FieldSamplerQuery::clearAllFieldSamplers()
+{
+	for (uint32_t i = 0; i < mSceneList.getSize(); ++i)
+	{
+		SceneInfo* sceneInfo = DYNAMIC_CAST(SceneInfo*)(mSceneList.getResource(i));
+		sceneInfo->clearAllFieldSamplers();
+	}
+}
+
+void FieldSamplerQuery::submitFieldSamplerQuery(const FieldSamplerQueryDataIntl& data, PxTask* task, PxTask* readyTask)
+{
+	PX_UNUSED(readyTask);
+	for (uint32_t i = 0; i < mSceneList.getSize(); ++i)
+	{
+		SceneInfo* sceneInfo = DYNAMIC_CAST(SceneInfo*)(mSceneList.getResource(i));
+		FieldSamplerSceneIntl* niFieldSamplerScene = sceneInfo->getSceneWrapper()->getInternalFieldSamplerScene();
+		const PxTask* fieldSamplerReadyTask = niFieldSamplerScene->onSubmitFieldSamplerQuery(data, readyTask);
+		if (fieldSamplerReadyTask != 0)
+		{
+			task->startAfter(fieldSamplerReadyTask->getTaskID());
+		}
+	}
+}
+
+void FieldSamplerQuery::update()
+{
+	mPrimarySceneList.clear();
+	mSecondarySceneList.clear();
+
+	for (uint32_t i = 0; i < mSceneList.getSize(); ++i)
+	{
+		SceneInfo* sceneInfo = DYNAMIC_CAST(SceneInfo*)(mSceneList.getResource(i));
+		sceneInfo->update();
+
+		if (sceneInfo->getEnabledFieldSamplerCount() > 0 && (sceneInfo->getSceneWrapper()->getInternalFieldSamplerScene() != mQueryDesc.ownerFieldSamplerScene))
+		{
+			((sceneInfo->getSceneWrapper()->getInternalFieldSamplerSceneDesc().isPrimary) ? mPrimarySceneList : mSecondarySceneList).pushBack(sceneInfo);
+		}
+	}
+}
+
+bool FieldSamplerQuery::SceneInfo::update()
+{
+	mEnabledFieldSamplerCount = 0;
+	for (uint32_t i = 0; i < mFieldSamplerArray.size(); ++i)
+	{
+		if (mFieldSamplerArray[i].mFieldSamplerWrapper->isEnabled())
+		{
+			++mEnabledFieldSamplerCount;
+		}
+		if (mFieldSamplerArray[i].mFieldSamplerWrapper->isEnabledChanged())
+		{
+			mFieldSamplerArrayChanged = true;
+		}
+	}
+
+	if (mFieldSamplerArrayChanged)
+	{
+		mFieldSamplerArrayChanged = false;
+		return true;
+	}
+	return false;
+}
+/******************************** CPU Version ********************************/
+class TaskExecute : public PxTask, public UserAllocated
+{
+public:
+	TaskExecute(FieldSamplerQueryCPU* query) : mQuery(query) {}
+
+	const char* getName() const
+	{
+		return "FieldSamplerQueryCPU::TaskExecute";
+	}
+	void run()
+	{
+		mQuery->execute();
+	}
+
+protected:
+	FieldSamplerQueryCPU* mQuery;
+};
+
+FieldSamplerQueryCPU::FieldSamplerQueryCPU(const FieldSamplerQueryDescIntl& desc, ResourceList& list, FieldSamplerManager* manager)
+	: FieldSamplerQuery(desc, list, manager)
+{
+	mTaskExecute = PX_NEW(TaskExecute)(this);
+
+	mExecuteCount = 256;
+}
+
+FieldSamplerQueryCPU::~FieldSamplerQueryCPU()
+{
+	delete mTaskExecute;
+}
+
+PxTaskID FieldSamplerQueryCPU::submitFieldSamplerQuery(const FieldSamplerQueryDataIntl& data, PxTaskID taskID)
+{
+	PX_ASSERT(data.isDataOnDevice == false);
+	PX_ASSERT(data.count <= mQueryDesc.maxCount);
+	if (data.count == 0)
+	{
+		return taskID;
+	}
+	mQueryData = data;
+
+	mResultField.resize(mExecuteCount);
+	mWeights.resize(mExecuteCount);
+	mAccumVelocity.reserve(mQueryDesc.maxCount);
+
+	PxTaskManager* tm = mManager->getApexScene().getTaskManager();
+	tm->submitUnnamedTask(*mTaskExecute);
+
+	FieldSamplerQuery::submitFieldSamplerQuery(data, mTaskExecute, NULL);
+
+	mTaskExecute->finishBefore(taskID);
+	return mTaskExecute->getTaskID();
+}
+
+void FieldSamplerQueryCPU::execute()
+{
+	if (mOnStartCallback)
+	{
+		(*mOnStartCallback)(NULL);
+	}
+
+	FieldSamplerIntl::ExecuteData executeData;
+	
+	executeData.position		= mQueryData.pmaInPosition;
+	executeData.velocity		= mQueryData.pmaInVelocity;
+	executeData.mass			= mQueryData.pmaInMass;//  + massOffset;
+	executeData.resultField		= mResultField.begin();
+	executeData.positionStride	= mQueryData.positionStrideBytes;
+	executeData.velocityStride	= mQueryData.velocityStrideBytes;
+	executeData.massStride		= mQueryData.massStrideBytes;
+	executeData.indicesMask		= 0;
+
+	uint32_t beginIndex;
+	uint32_t* indices = &beginIndex;
+	if (mQueryData.pmaInIndices)
+	{
+		indices = mQueryData.pmaInIndices;
+		executeData.indicesMask = ~executeData.indicesMask;
+	}
+
+	for (uint32_t executeOffset = 0; executeOffset < mQueryData.count; executeOffset += mExecuteCount)
+	{
+		const uint32_t positionStride = mQueryData.positionStrideBytes / 4;
+		const uint32_t velocityStride = mQueryData.velocityStrideBytes / 4;
+		const uint32_t massStride = mQueryData.massStrideBytes / 4;
+		//const uint32_t offset = executeOffset * stride;
+		//const uint32_t massOffset = executeOffset * massStride;
+
+		beginIndex = executeOffset;
+		executeData.count        = PxMin(mExecuteCount, mQueryData.count - executeOffset);
+		executeData.indices		 = indices + (executeOffset & executeData.indicesMask);
+
+		PxVec4* accumField = (PxVec4*)(mQueryData.pmaOutField);
+		PxVec4* accumVelocity = mAccumVelocity.getPtr() + executeOffset;
+		//clear accum
+		for (uint32_t i = 0; i < executeData.count; ++i)
+		{
+			uint32_t j = executeData.indices[i & executeData.indicesMask] + (i & ~executeData.indicesMask);
+			accumField[j] = PxVec4(0.0f);
+			accumVelocity[i] = PxVec4(0.0f);
+		}
+		for (uint32_t sceneIdx = 0; sceneIdx < mPrimarySceneList.size(); ++sceneIdx)
+		{
+			executeScene(mPrimarySceneList[sceneIdx], executeData, accumField, accumVelocity, positionStride, velocityStride, massStride);
+		}
+
+		//setup weights for secondary scenes
+		for (uint32_t i = 0; i < executeData.count; ++i)
+		{
+			uint32_t j = executeData.indices[i & executeData.indicesMask] + (i & ~executeData.indicesMask);
+			accumField[j].w = accumVelocity[i].w;
+			accumVelocity[i].w = 0.0f;
+		}
+		for (uint32_t sceneIdx = 0; sceneIdx < mSecondarySceneList.size(); ++sceneIdx)
+		{
+			executeScene(mSecondarySceneList[sceneIdx], executeData, accumField, accumVelocity, positionStride, velocityStride, massStride);
+		}
+
+		//compose accum field
+		for (uint32_t i = 0; i < executeData.count; ++i)
+		{
+			uint32_t j = executeData.indices[i & executeData.indicesMask] + (i & ~executeData.indicesMask);
+			float blend = accumField[j].w;
+			float velW = accumVelocity[i].w;
+			float weight = blend + velW * (1 - blend);
+			if (weight >= VELOCITY_WEIGHT_THRESHOLD)
+			{
+				PxVec3 result = accumField[j].getXYZ();
+				const PxVec3& velocity = *(PxVec3*)(executeData.velocity + j * velocityStride);
+				result += (accumVelocity[i].getXYZ() - weight * velocity);
+				accumField[j] = PxVec4(result, 0);
+			}
+		}
+	}
+
+	if (mOnFinishCallback)
+	{
+		(*mOnFinishCallback)(NULL);
+	}
+}
+
+void FieldSamplerQueryCPU::executeScene(const SceneInfo* sceneInfo, 
+										const FieldSamplerIntl::ExecuteData& executeData, 
+										PxVec4* accumField, 
+										PxVec4* accumVelocity, 
+										uint32_t positionStride, 
+										uint32_t velocityStride, 
+										uint32_t massStride)
+{
+	FieldSamplerExecuteArgs execArgs;
+	execArgs.elapsedTime = mQueryData.timeStep;
+	execArgs.totalElapsedMS = mManager->getApexScene().getTotalElapsedMS();
+
+	const nvidia::Array<FieldSamplerInfo>& fieldSamplerArray = sceneInfo->getFieldSamplerArray();
+	for (uint32_t fieldSamplerIdx = 0; fieldSamplerIdx < fieldSamplerArray.size(); ++fieldSamplerIdx)
+	{
+		const FieldSamplerWrapperCPU* fieldSampler = DYNAMIC_CAST(FieldSamplerWrapperCPU*)(fieldSamplerArray[fieldSamplerIdx].mFieldSamplerWrapper);
+		if (fieldSampler->isEnabled())
+		{
+			const float multiplier = fieldSamplerArray[fieldSamplerIdx].mMultiplier;
+			PX_UNUSED(multiplier);
+
+			const FieldSamplerDescIntl& desc = fieldSampler->getInternalFieldSamplerDesc();
+			if (desc.cpuSimulationSupport)
+			{
+				const FieldShapeDescIntl& shapeDesc = fieldSampler->getInternalFieldSamplerShape();
+				PX_ASSERT(shapeDesc.weight >= 0.0f && shapeDesc.weight <= 1.0f);
+
+				for (uint32_t i = 0; i < executeData.count; ++i)
+				{
+					mWeights[i] = 0;
+				}
+
+				uint32_t boundaryCount = fieldSampler->getFieldBoundaryCount();
+				for (uint32_t boundaryIndex = 0; boundaryIndex < boundaryCount; ++boundaryIndex)
+				{
+					FieldBoundaryWrapper* fieldBoundaryWrapper = fieldSampler->getFieldBoundaryWrapper(boundaryIndex);
+
+					const nvidia::Array<FieldShapeDescIntl>& fieldShapes = fieldBoundaryWrapper->getFieldShapes();
+					for (uint32_t shapeIndex = 0; shapeIndex < fieldShapes.size(); ++shapeIndex)
+					{
+						const FieldShapeDescIntl& boundaryShapeDesc = fieldShapes[shapeIndex];
+						PX_ASSERT(boundaryShapeDesc.weight >= 0.0f && boundaryShapeDesc.weight <= 1.0f);
+
+						for (uint32_t i = 0; i < executeData.count; ++i)
+						{
+							uint32_t j = executeData.indices[i & executeData.indicesMask] + (i & ~executeData.indicesMask);
+							PxVec3* pos = (PxVec3*)(executeData.position + j * positionStride);
+							const float excludeWeight = evalFade(evalDistInShape(boundaryShapeDesc, *pos), 0.0f) * boundaryShapeDesc.weight;
+							mWeights[i] = PxMax(mWeights[i], excludeWeight);
+						}
+					}
+				}
+
+				for (uint32_t i = 0; i < executeData.count; ++i)
+				{
+					uint32_t j = executeData.indices[i & executeData.indicesMask] + (i & ~executeData.indicesMask);
+					PxVec3* pos = (PxVec3*)(executeData.position + j * positionStride);
+					const float includeWeight = evalFade(evalDistInShape(shapeDesc, *pos), desc.boundaryFadePercentage) * shapeDesc.weight;
+					const float excludeWeight = mWeights[i];
+					mWeights[i] = includeWeight * (1.0f - excludeWeight);
+#if FIELD_SAMPLER_MULTIPLIER == FIELD_SAMPLER_MULTIPLIER_WEIGHT
+					mWeights[i] *= multiplier;
+#endif
+				}
+
+				//execute field
+				fieldSampler->getInternalFieldSampler()->executeFieldSampler(executeData);
+
+#if FIELD_SAMPLER_MULTIPLIER == FIELD_SAMPLER_MULTIPLIER_VALUE
+				const float multiplier = fieldSamplerArray[fieldSamplerIdx].mMultiplier;
+				for (uint32_t i = 0; i < executeData.count; ++i)
+				{
+					executeData.resultField[i] *= multiplier;
+				}
+#endif
+
+				//accum field
+				switch (desc.type)
+				{
+				case FieldSamplerTypeIntl::FORCE:
+					for (uint32_t i = 0; i < executeData.count; ++i)
+					{
+						uint32_t j = executeData.indices[i & executeData.indicesMask] + (i & ~executeData.indicesMask);
+						execArgs.position = *(PxVec3*)(executeData.position + j * positionStride);
+						execArgs.velocity = *(PxVec3*)(executeData.velocity + j * velocityStride);
+						execArgs.mass     = *(executeData.mass + massStride * j);
+
+						accumFORCE(execArgs, executeData.resultField[i], mWeights[i], accumField[j], accumVelocity[i]);
+					}
+					break;
+				case FieldSamplerTypeIntl::ACCELERATION:
+					for (uint32_t i = 0; i < executeData.count; ++i)
+					{
+						uint32_t j = executeData.indices[i & executeData.indicesMask] + (i & ~executeData.indicesMask);
+						execArgs.position = *(PxVec3*)(executeData.position + j * positionStride);
+						execArgs.velocity = *(PxVec3*)(executeData.velocity + j * velocityStride);
+						execArgs.mass     = *(executeData.mass + massStride * j);
+
+						accumACCELERATION(execArgs, executeData.resultField[i], mWeights[i], accumField[j], accumVelocity[i]);
+					}
+					break;
+				case FieldSamplerTypeIntl::VELOCITY_DRAG:
+					for (uint32_t i = 0; i < executeData.count; ++i)
+					{
+						uint32_t j = executeData.indices[i & executeData.indicesMask] + (i & ~executeData.indicesMask);
+						execArgs.position = *(PxVec3*)(executeData.position + j * positionStride);
+						execArgs.velocity = *(PxVec3*)(executeData.velocity + j * velocityStride);
+						execArgs.mass     = *(executeData.mass + massStride * j);
+
+						accumVELOCITY_DRAG(execArgs, desc.dragCoeff, executeData.resultField[i], mWeights[i], accumField[j], accumVelocity[i]);
+					}
+					break;
+				case FieldSamplerTypeIntl::VELOCITY_DIRECT:
+					for (uint32_t i = 0; i < executeData.count; ++i)
+					{
+						uint32_t j = executeData.indices[i & executeData.indicesMask] + (i & ~executeData.indicesMask);
+						execArgs.position = *(PxVec3*)(executeData.position + j * positionStride);
+						execArgs.velocity = *(PxVec3*)(executeData.velocity + j * velocityStride);
+						execArgs.mass     = *(executeData.mass + massStride * j);
+
+						accumVELOCITY_DIRECT(execArgs, executeData.resultField[i], mWeights[i], accumField[j], accumVelocity[i]);
+					}
+					break;
+				};
+			}
+		}
+	}
+
+}
+
+
+/******************************** GPU Version ********************************/
+#if APEX_CUDA_SUPPORT
+
+class FieldSamplerQueryLaunchTask : public PxGpuTask, public UserAllocated
+{
+public:
+	FieldSamplerQueryLaunchTask(FieldSamplerQueryGPU* query) : mQuery(query) {}
+	const char* getName() const
+	{
+		return "FieldSamplerQueryLaunchTask";
+	}
+	void         run()
+	{
+		PX_ALWAYS_ASSERT();
+	}
+	bool         launchInstance(CUstream stream, int kernelIndex)
+	{
+		return mQuery->launch(stream, kernelIndex);
+	}
+	PxGpuTaskHint::Enum getTaskHint() const
+	{
+		return PxGpuTaskHint::Kernel;
+	}
+
+protected:
+	FieldSamplerQueryGPU* mQuery;
+};
+
+class FieldSamplerQueryPrepareTask : public PxTask, public UserAllocated
+{
+public:
+	FieldSamplerQueryPrepareTask(FieldSamplerQueryGPU* query) : mQuery(query) {}
+
+	const char* getName() const
+	{
+		return "FieldSamplerQueryPrepareTask";
+	}
+	void run()
+	{
+		mQuery->prepare();
+	}
+
+protected:
+	FieldSamplerQueryGPU* mQuery;
+};
+
+class FieldSamplerQueryCopyTask : public PxGpuTask, public UserAllocated
+{
+public:
+	FieldSamplerQueryCopyTask(FieldSamplerQueryGPU* query) : mQuery(query) {}
+	const char* getName() const
+	{
+		return "FieldSamplerQueryCopyTask";
+	}
+	void         run()
+	{
+		PX_ALWAYS_ASSERT();
+	}
+	bool         launchInstance(CUstream stream, int kernelIndex)
+	{
+		return mQuery->copy(stream, kernelIndex);
+	}
+	PxGpuTaskHint::Enum getTaskHint() const
+	{
+		return PxGpuTaskHint::Kernel;
+	}
+
+protected:
+	FieldSamplerQueryGPU* mQuery;
+};
+
+class FieldSamplerQueryFetchTask : public PxTask, public UserAllocated
+{
+public:
+	FieldSamplerQueryFetchTask(FieldSamplerQueryGPU* query) : mQuery(query) {}
+
+	const char* getName() const
+	{
+		return "FieldSamplerQueryFetchTask";
+	}
+	void run()
+	{
+		mQuery->fetch();
+	}
+
+protected:
+	FieldSamplerQueryGPU* mQuery;
+};
+
+
+FieldSamplerQueryGPU::FieldSamplerQueryGPU(const FieldSamplerQueryDescIntl& desc, ResourceList& list, FieldSamplerManager* manager)
+	: FieldSamplerQueryCPU(desc, list, manager)
+	, mPositionMass(manager->getApexScene(), PX_ALLOC_INFO("mPositionMass", PARTICLES))
+	, mVelocity(manager->getApexScene(), PX_ALLOC_INFO("mVelocity", PARTICLES))
+	, mAccumField(manager->getApexScene(), PX_ALLOC_INFO("mAccumField", PARTICLES))
+	, mCopyQueue(*manager->getApexScene().getTaskManager()->getGpuDispatcher())
+{
+	mTaskLaunch   = PX_NEW(FieldSamplerQueryLaunchTask)(this);
+	mTaskPrepare  = PX_NEW(FieldSamplerQueryPrepareTask)(this);
+	mTaskCopy     = PX_NEW(FieldSamplerQueryCopyTask)(this);
+	mTaskFetch    = PX_NEW(FieldSamplerQueryFetchTask)(this);
+}
+
+FieldSamplerQueryGPU::~FieldSamplerQueryGPU()
+{
+	PX_DELETE(mTaskFetch);
+	PX_DELETE(mTaskCopy);
+	PX_DELETE(mTaskPrepare);
+	PX_DELETE(mTaskLaunch);
+}
+
+PxTaskID FieldSamplerQueryGPU::submitFieldSamplerQuery(const FieldSamplerQueryDataIntl& data, PxTaskID taskID)
+{
+	PX_ASSERT(data.count <= mQueryDesc.maxCount);
+	if (data.count == 0)
+	{
+		return taskID;
+	}
+	mQueryData = data;
+
+	if (!data.isDataOnDevice)
+	{
+		bool isWorkOnCPU = true;
+		// try to find FieldSampler which has no CPU implemntation (Turbulence for example)
+		for (uint32_t sceneIdx = 0; (sceneIdx < mPrimarySceneList.size() + mSecondarySceneList.size()) && isWorkOnCPU; ++sceneIdx)
+		{
+			const nvidia::Array<FieldSamplerInfo>& fsArray = sceneIdx < mPrimarySceneList.size() 
+				? mPrimarySceneList[sceneIdx]->getFieldSamplerArray() 
+				: mSecondarySceneList[sceneIdx-mPrimarySceneList.size()]->getFieldSamplerArray();
+			for (uint32_t fsIdx = 0; fsIdx < fsArray.size() && isWorkOnCPU; fsIdx++)
+			{
+				if (fsArray[fsIdx].mFieldSamplerWrapper->isEnabled())
+				{
+					isWorkOnCPU = fsArray[fsIdx].mFieldSamplerWrapper->getInternalFieldSamplerDesc().cpuSimulationSupport;
+				}
+			}
+		}
+
+		// if all FSs can work on CPU we will execute FieldSamplerQuery on CPU
+		if (isWorkOnCPU)
+		{
+			return FieldSamplerQueryCPU::submitFieldSamplerQuery(data, taskID);
+		}
+
+		mPositionMass.reserve(mQueryDesc.maxCount, ApexMirroredPlace::CPU_GPU);
+		mVelocity.reserve(mQueryDesc.maxCount, ApexMirroredPlace::CPU_GPU);
+		mAccumField.reserve(mQueryDesc.maxCount, ApexMirroredPlace::CPU_GPU);
+	}
+	mAccumVelocity.reserve(mQueryDesc.maxCount, ApexMirroredPlace::CPU_GPU);
+	
+	// if data on device or some FS can't work on CPU we will launch FieldSamplerQuery on GPU
+	PxTaskManager* tm = mManager->getApexScene().getTaskManager();
+	tm->submitUnnamedTask(*mTaskLaunch, PxTaskType::TT_GPU);
+
+	if (data.isDataOnDevice)
+	{
+		FieldSamplerQuery::submitFieldSamplerQuery(data, mTaskLaunch, NULL);
+
+		mTaskLaunch->finishBefore(taskID);
+		return mTaskLaunch->getTaskID();
+	}
+	else
+	{
+		FieldSamplerQueryDataIntl data4Device;
+		data4Device.timeStep = data.timeStep;
+		data4Device.count = data.count;
+		data4Device.isDataOnDevice = true;
+		data4Device.positionStrideBytes = sizeof(PxVec4);
+		data4Device.velocityStrideBytes = sizeof(PxVec4);
+		data4Device.massStrideBytes = sizeof(PxVec4);
+		data4Device.pmaInPosition = (float*)mPositionMass.getGpuPtr();
+		data4Device.pmaInVelocity = (float*)mVelocity.getGpuPtr();
+		data4Device.pmaInMass = &mPositionMass.getGpuPtr()->w;
+		data4Device.pmaOutField = mAccumField.getGpuPtr();
+		data4Device.pmaInIndices = 0;
+
+		FieldSamplerQuery::submitFieldSamplerQuery(data4Device, mTaskLaunch, mTaskCopy);
+
+		tm->submitUnnamedTask(*mTaskPrepare);
+		tm->submitUnnamedTask(*mTaskCopy, PxTaskType::TT_GPU);
+		tm->submitUnnamedTask(*mTaskFetch);
+
+		mTaskPrepare->finishBefore(mTaskCopy->getTaskID());
+		mTaskCopy->finishBefore(mTaskLaunch->getTaskID());
+		mTaskLaunch->finishBefore(mTaskFetch->getTaskID());
+		mTaskFetch->finishBefore(taskID);
+		return mTaskPrepare->getTaskID();
+	}
+}
+
+void FieldSamplerQueryGPU::prepare()
+{
+	const uint32_t positionStride = mQueryData.positionStrideBytes / sizeof(float);
+	const uint32_t velocityStride = mQueryData.velocityStrideBytes / sizeof(float);
+	const uint32_t massStride = mQueryData.massStrideBytes / sizeof(float);
+	for (uint32_t idx = 0; idx < mQueryData.count; idx++)
+	{
+		mPositionMass[idx] = PxVec4(*(PxVec3*)(mQueryData.pmaInPosition + idx * positionStride), *(mQueryData.pmaInMass + idx * massStride));
+		mVelocity[idx] = PxVec4(*(PxVec3*)(mQueryData.pmaInVelocity + idx * velocityStride), 0.f);
+	}
+}
+
+void FieldSamplerQueryGPU::fetch()
+{
+	for (uint32_t idx = 0; idx < mQueryData.count; idx++)
+	{
+		mQueryData.pmaOutField[idx] = mAccumField[idx];
+	}
+}
+
+bool FieldSamplerQueryGPU::copy(CUstream stream, int kernelIndex)
+{
+	if (kernelIndex == 0)
+	{
+		mCopyQueue.reset(stream, 4);
+		mPositionMass.copyHostToDeviceQ(mCopyQueue, mQueryData.count);
+		mVelocity.copyHostToDeviceQ(mCopyQueue, mQueryData.count);
+		mCopyQueue.flushEnqueued();
+	}
+	return false;
+}
+
+bool FieldSamplerQueryGPU::launch(CUstream stream, int kernelIndex)
+{
+	FieldSamplerPointsKernelArgs args;
+	args.elapsedTime        = mQueryData.timeStep;
+	args.totalElapsedMS     = mManager->getApexScene().getTotalElapsedMS();
+	if (mQueryData.isDataOnDevice)
+	{
+		args.positionMass   = (float4*)mQueryData.pmaInPosition;
+		args.velocity       = (float4*)mQueryData.pmaInVelocity;
+		args.accumField     = (float4*)mQueryData.pmaOutField;
+	}
+	else
+	{
+		args.positionMass   = (float4*)mPositionMass.getGpuPtr();
+		args.velocity       = (float4*)mVelocity.getGpuPtr();
+		args.accumField     = (float4*)mAccumField.getGpuPtr();
+	}
+	args.accumVelocity      = (float4*)mAccumVelocity.getGpuPtr();
+
+	FieldSamplerPointsKernelLaunchDataIntl launchData;
+	launchData.stream       = stream;
+	launchData.kernelType   = FieldSamplerKernelType::POINTS;
+	launchData.kernelArgs   = &args;
+	launchData.threadCount  = mQueryData.count;
+	launchData.memRefSize   = mQueryData.count;
+
+	if (kernelIndex == 0 && mOnStartCallback)
+	{
+		(*mOnStartCallback)(stream);
+	}
+
+	if (kernelIndex == 0)
+	{
+		CUDA_OBJ(clearKernel)(stream, mQueryData.count,
+		                      createApexCudaMemRef(args.accumField, launchData.memRefSize, ApexCudaMemFlags::OUT),
+		                      createApexCudaMemRef(args.accumVelocity, launchData.memRefSize, ApexCudaMemFlags::OUT));
+		return true;
+	}
+	--kernelIndex;
+
+	const uint32_t bothSceneCount = mPrimarySceneList.size() + mSecondarySceneList.size();
+	if (kernelIndex < (int) bothSceneCount)
+	{
+		SceneInfo* sceneInfo = (kernelIndex < (int) mPrimarySceneList.size()) 
+			? mPrimarySceneList[(uint32_t)kernelIndex] 
+			: mSecondarySceneList[(uint32_t)kernelIndex - mPrimarySceneList.size()];
+		SceneInfoGPU* sceneInfoGPU = DYNAMIC_CAST(SceneInfoGPU*)(sceneInfo);
+
+		launchData.kernelMode = FieldSamplerKernelMode::DEFAULT;
+		if (kernelIndex == (int) mPrimarySceneList.size() - 1)
+		{
+			launchData.kernelMode = FieldSamplerKernelMode::FINISH_PRIMARY;
+		}
+		if ((kernelIndex == (int) bothSceneCount - 1))
+		{
+			launchData.kernelMode = FieldSamplerKernelMode::FINISH_SECONDARY;
+		}
+
+		FieldSamplerSceneWrapperGPU* sceneWrapper = DYNAMIC_CAST(FieldSamplerSceneWrapperGPU*)(sceneInfo->getSceneWrapper());
+
+		launchData.queryParamsHandle = sceneInfoGPU->getQueryParamsHandle();
+		launchData.paramsExArrayHandle = sceneInfoGPU->getParamsHandle();
+		launchData.fieldSamplerArray = &sceneInfo->getFieldSamplerArray();
+		launchData.activeFieldSamplerCount = sceneInfo->getEnabledFieldSamplerCount();
+
+		sceneWrapper->getInternalFieldSamplerScene()->launchFieldSamplerCudaKernel(launchData);
+		return true;
+	}
+	kernelIndex -= bothSceneCount;
+
+	if (kernelIndex == 0)
+	{
+		CUDA_OBJ(composeKernel)(stream, mQueryData.count,
+								createApexCudaMemRef(args.accumField, launchData.memRefSize, ApexCudaMemFlags::IN_OUT),
+								createApexCudaMemRef((const float4*)args.accumVelocity, launchData.memRefSize, ApexCudaMemFlags::IN),
+								createApexCudaMemRef(args.velocity, launchData.memRefSize, ApexCudaMemFlags::IN),
+								args.elapsedTime);
+		return true;
+	}
+	--kernelIndex;
+
+	if (!mQueryData.isDataOnDevice)
+	{
+		mAccumField.copyDeviceToHostQ(mCopyQueue, mQueryData.count);
+		mCopyQueue.flushEnqueued();
+
+		PxTaskManager* tm = mManager->getApexScene().getTaskManager();
+		tm->getGpuDispatcher()->addCompletionPrereq(*mTaskFetch);
+	}
+
+	if (mOnFinishCallback)
+	{
+		(*mOnFinishCallback)(stream);
+	}
+	return false;
+}
+
+FieldSamplerQueryGPU::SceneInfoGPU::SceneInfoGPU(ResourceList& list, FieldSamplerQuery* query, FieldSamplerSceneWrapper* sceneWrapper)
+	: SceneInfo(list, query, sceneWrapper)
+	, mConstMemGroup(DYNAMIC_CAST(FieldSamplerSceneWrapperGPU*)(sceneWrapper)->getConstStorage())
+{
+	APEX_CUDA_CONST_MEM_GROUP_SCOPE(mConstMemGroup);
+
+	mQueryParamsHandle.alloc(_storage_);
+}
+
+bool FieldSamplerQueryGPU::SceneInfoGPU::update()
+{
+	if (FieldSamplerQuery::SceneInfo::update())
+	{
+		APEX_CUDA_CONST_MEM_GROUP_SCOPE(mConstMemGroup);
+
+		FieldSamplerParamsExArray paramsExArray;
+		mParamsExArrayHandle.allocOrFetch(_storage_, paramsExArray);
+		if (paramsExArray.resize(_storage_, mEnabledFieldSamplerCount))
+		{
+			for (uint32_t i = 0, enabledIdx = 0; i < mFieldSamplerArray.size(); ++i)
+			{
+				FieldSamplerWrapperGPU* fieldSamplerWrapper = DYNAMIC_CAST(FieldSamplerWrapperGPU*)(mFieldSamplerArray[i].mFieldSamplerWrapper);
+				if (fieldSamplerWrapper->isEnabled())
+				{
+					FieldSamplerParamsEx fsParamsEx;
+					fsParamsEx.paramsHandle = fieldSamplerWrapper->getParamsHandle();
+					fsParamsEx.multiplier = mFieldSamplerArray[i].mMultiplier;
+					PX_ASSERT(enabledIdx < mEnabledFieldSamplerCount);
+					paramsExArray.updateElem(_storage_, fsParamsEx, enabledIdx++);
+				}
+			}
+			mParamsExArrayHandle.update(_storage_, paramsExArray);
+		}
+		return true;
+	}
+	return false;
+}
+
+PxVec3 FieldSamplerQueryGPU::executeFieldSamplerQueryOnGrid(const FieldSamplerQueryGridDataIntl& data)
+{
+	FieldSamplerGridKernelArgs args;
+
+	args.numX           = data.numX;
+	args.numY           = data.numY;
+	args.numZ           = data.numZ;
+
+	args.gridToWorld    = data.gridToWorld;
+
+	args.mass           = data.mass;
+	args.elapsedTime    = data.timeStep;
+	args.cellSize		= data.cellSize;
+	args.totalElapsedMS = mManager->getApexScene().getTotalElapsedMS();
+
+	FieldSamplerGridKernelLaunchDataIntl launchData;
+	launchData.stream       = data.stream;
+	launchData.kernelType   = FieldSamplerKernelType::GRID;
+	launchData.kernelArgs   = &args;
+	launchData.threadCountX	= data.numX;
+	launchData.threadCountY	= data.numY;
+	launchData.threadCountZ	= data.numZ;
+	launchData.accumArray	= data.resultVelocity;
+
+
+	{
+		APEX_CUDA_SURFACE_SCOPE_BIND(surfRefGridAccum, *launchData.accumArray, ApexCudaMemFlags::OUT);
+
+		CUDA_OBJ(clearGridKernel)(data.stream, launchData.threadCountX, launchData.threadCountY, launchData.threadCountZ,
+								  args.numX, args.numY, args.numZ);
+	}
+
+	PxVec3 velocity(0.0f);
+	for (uint32_t i = 0; i < mSecondarySceneList.size(); ++i)
+	{
+		SceneInfoGPU* sceneInfo = DYNAMIC_CAST(SceneInfoGPU*)(mSecondarySceneList[i]);
+		FieldSamplerSceneWrapperGPU* sceneWrapper = DYNAMIC_CAST(FieldSamplerSceneWrapperGPU*)(sceneInfo->getSceneWrapper());
+
+		launchData.activeFieldSamplerCount = 0;
+
+		const nvidia::Array<FieldSamplerInfo>& fieldSamplerArray = sceneInfo->getFieldSamplerArray();
+		for (uint32_t fieldSamplerIdx = 0; fieldSamplerIdx < fieldSamplerArray.size(); ++fieldSamplerIdx)
+		{
+			const FieldSamplerWrapperGPU* wrapper = static_cast<const FieldSamplerWrapperGPU* >( fieldSamplerArray[fieldSamplerIdx].mFieldSamplerWrapper );
+			if (wrapper->isEnabled())
+			{
+				switch (wrapper->getInternalFieldSamplerDesc().gridSupportType)
+				{
+					case FieldSamplerGridSupportTypeIntl::SINGLE_VELOCITY:
+					{
+						const FieldSamplerIntl* fieldSampler = wrapper->getInternalFieldSampler();
+						velocity += fieldSampler->queryFieldSamplerVelocity();
+					}
+					break;
+					case FieldSamplerGridSupportTypeIntl::VELOCITY_PER_CELL:
+					{
+						launchData.activeFieldSamplerCount += 1;
+					}
+					break;
+					default:
+						break;
+				}
+			}
+		}
+
+		if (launchData.activeFieldSamplerCount > 0)
+		{
+			launchData.queryParamsHandle = sceneInfo->getQueryParamsHandle();
+			launchData.paramsExArrayHandle = sceneInfo->getParamsHandle();
+			launchData.fieldSamplerArray = &sceneInfo->getFieldSamplerArray();
+			launchData.kernelMode = FieldSamplerKernelMode::DEFAULT;
+
+			sceneWrapper->getInternalFieldSamplerScene()->launchFieldSamplerCudaKernel(launchData);
+		}
+	}
+	return velocity;
+}
+
+
+#endif
+
+}
+} // end namespace nvidia::apex
+
diff --git a/APEX_1.4/module/fieldsampler/src/FieldSamplerScene.cpp b/APEX_1.4/module/fieldsampler/src/FieldSamplerScene.cpp
new file mode 100644
index 00000000..58c74ed7
--- /dev/null
+++ b/APEX_1.4/module/fieldsampler/src/FieldSamplerScene.cpp
@@ -0,0 +1,304 @@
+/*
+ * Copyright (c) 2008-2015, 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 "ApexDefs.h"
+#include "Apex.h"
+
+#include "FieldSamplerScene.h"
+#include "FieldSamplerManager.h"
+#include "FieldSamplerQuery.h"
+#include "FieldSamplerPhysXMonitor.h"
+#include "SceneIntl.h"
+#include "RenderDebugInterface.h"
+#include "ModulePerfScope.h"
+
+#if APEX_CUDA_SUPPORT
+#include "PxGpuTask.h"
+#include "ApexCudaSource.h"
+#endif
+
+#include "Lock.h"
+
+namespace nvidia
+{
+namespace fieldsampler
+{
+
+FieldSamplerScene::FieldSamplerScene(ModuleFieldSamplerImpl& module, SceneIntl& scene, RenderDebugInterface* debugRender, ResourceList& list)
+	: mModule(&module)
+	, mApexScene(&scene)
+	, mDebugRender(debugRender)
+	, mManager(NULL)
+	, mPhysXScene(NULL)
+	, mForceSampleBatchBufferPos(0)
+	, mForceSampleBatchBufferSize(0)
+{
+	list.add(*this);		// Add self to module's list of FieldSamplerScenes
+}
+
+FieldSamplerScene::~FieldSamplerScene()
+{
+}
+
+void FieldSamplerScene::visualize()
+{
+#ifndef WITHOUT_DEBUG_VISUALIZE
+#endif
+}
+
+void FieldSamplerScene::destroy()
+{
+#if PX_PHYSICS_VERSION_MAJOR == 3
+	PX_DELETE(mPhysXMonitor);
+#endif
+	PX_DELETE(mManager);
+
+	removeAllActors();
+	mApexScene->moduleReleased(*this);
+	delete this;
+}
+
+FieldSamplerManagerIntl* FieldSamplerScene::getManager()
+{
+	if (mManager == NULL)
+	{
+		mManager = createManager();
+		PX_ASSERT(mManager != NULL);
+	}
+	return mManager;
+}
+
+
+
+void FieldSamplerScene::setModulePhysXScene(PxScene* s)
+{
+	if (s)
+	{
+		mPhysXMonitor->setPhysXScene(s);
+	}
+	mPhysXScene = s;
+}
+
+
+void FieldSamplerScene::submitTasks(float /*elapsedTime*/, float /*substepSize*/, uint32_t /*numSubSteps*/)
+{
+#if PX_PHYSICS_VERSION_MAJOR == 3
+	PxTaskManager* tm;
+	{
+		READ_LOCK(*mApexScene);
+		tm = mApexScene->getTaskManager();
+	}
+	tm->submitNamedTask(&mPhysXMonitorFetchTask, FSST_PHYSX_MONITOR_FETCH);
+	tm->submitNamedTask(&mPhysXMonitorLoadTask, FSST_PHYSX_MONITOR_LOAD);
+#endif
+	if (mManager != NULL)
+	{
+		mManager->submitTasks();
+	}
+}
+
+void FieldSamplerScene::setTaskDependencies()
+{
+#if PX_PHYSICS_VERSION_MAJOR == 3
+	for (uint32_t i = 0; i < mForceSampleBatchQuery.size(); i++)
+	{
+		if (mForceSampleBatchQuery[i] && mForceSampleBatchQueryData[i].count > 0)
+		{
+			static_cast<FieldSamplerQueryIntl*>(mForceSampleBatchQuery[i])->submitFieldSamplerQuery(mForceSampleBatchQueryData[i], mApexScene->getTaskManager()->getNamedTask(FSST_PHYSX_MONITOR_LOAD));
+			mForceSampleBatchQueryData[i].count = 0;
+		}
+		mForceSampleBatchBufferPos = 0;
+	}
+	if (mPhysXMonitor->isEnable())
+	{
+		mPhysXMonitor->update();
+	}
+#endif
+	if (mManager != NULL)
+	{
+		mManager->setTaskDependencies();
+	}
+
+#if PX_PHYSICS_VERSION_MAJOR == 3
+	// Just in case one of the scene conditions doesn't set a bounding dependency, let's not let these dangle
+	PxTaskManager* tm;
+	{
+		READ_LOCK(*mApexScene);
+		tm = mApexScene->getTaskManager();
+	}
+	mPhysXMonitorFetchTask.finishBefore(tm->getNamedTask(AST_PHYSX_FETCH_RESULTS));
+	mPhysXMonitorLoadTask.finishBefore(tm->getNamedTask(AST_PHYSX_FETCH_RESULTS));
+#endif
+}
+
+void FieldSamplerScene::fetchResults()
+{
+	if (mManager != NULL)
+	{
+		mManager->fetchResults();
+	}
+}
+
+#if PX_PHYSICS_VERSION_MAJOR == 3
+void FieldSamplerScene::enablePhysXMonitor(bool enable)
+{
+	PX_UNUSED(enable);
+	mPhysXMonitor->enablePhysXMonitor(enable);
+}
+
+void FieldSamplerScene::setPhysXFilterData(physx::PxFilterData filterData)
+{
+	mPhysXMonitor->setPhysXFilterData(filterData);
+}
+
+
+uint32_t FieldSamplerScene::createForceSampleBatch(uint32_t maxCount, const physx::PxFilterData filterData)
+{
+	mForceSampleBatchBufferSize += maxCount;
+	mForceSampleBatchPosition.resize(mForceSampleBatchBufferSize);
+	mForceSampleBatchVelocity.resize(mForceSampleBatchBufferSize);
+	mForceSampleBatchMass.resize(mForceSampleBatchBufferSize);
+
+	FieldSamplerQueryDescIntl desc;
+	desc.maxCount = maxCount;
+	desc.samplerFilterData = filterData;
+	//SceneInfo* sceneInfo = DYNAMIC_CAST(SceneInfo*)(mSceneList.getResource(i));
+	//InternalFieldSamplerScene* niFieldSamplerScene = sceneInfo->getSceneWrapper()->getInternalFieldSamplerScene();
+	//desc.ownerFieldSamplerScene = this;
+	
+
+	uint32_t id = 0;
+	while (id < mForceSampleBatchQuery.size() && mForceSampleBatchQuery[id]) 
+	{
+		id++;
+	}
+	if (id == mForceSampleBatchQuery.size())
+	{
+		mForceSampleBatchQuery.pushBack(0);
+		FieldSamplerQueryDataIntl data;
+		data.count = 0;
+		mForceSampleBatchQueryData.pushBack(data);
+	}
+	mForceSampleBatchQuery[id] = static_cast<FieldSamplerQuery*>(mManager->createFieldSamplerQuery(desc));
+	return id;
+}
+
+
+void FieldSamplerScene::releaseForceSampleBatch(uint32_t batchId)
+{
+	if (batchId < mForceSampleBatchQuery.size() && mForceSampleBatchQuery[batchId])
+	{
+		mForceSampleBatchBufferSize -= mForceSampleBatchQuery[batchId]->getQueryDesc().maxCount;
+		mForceSampleBatchPosition.resize(mForceSampleBatchBufferSize);
+		mForceSampleBatchVelocity.resize(mForceSampleBatchBufferSize);
+		mForceSampleBatchMass.resize(mForceSampleBatchBufferSize);
+
+		mForceSampleBatchQuery[batchId]->release();
+		mForceSampleBatchQuery[batchId] = 0;
+	}
+}
+
+
+void FieldSamplerScene::submitForceSampleBatch(	uint32_t batchId, PxVec4* forces, const uint32_t forcesStride,
+								const PxVec3* positions, const uint32_t positionsStride,
+								const PxVec3* velocities, const uint32_t velocitiesStride,
+								const float* mass, const uint32_t massStride,
+								const uint32_t* indices, const uint32_t numIndices)
+{
+	PX_UNUSED(forcesStride);
+	PX_ASSERT(forcesStride == sizeof(PxVec4));
+	PX_ASSERT(indices);
+	if (batchId >= mForceSampleBatchQuery.size() || mForceSampleBatchQuery[batchId] == 0) return;
+
+	uint32_t maxIndices = indices[numIndices - 1] + 1; //supposed that indices are sorted
+	for (uint32_t i = 0; i < maxIndices; i++)
+	{
+		mForceSampleBatchPosition[mForceSampleBatchBufferPos + i] = *(PxVec4*)((uint8_t*)positions + i * positionsStride);
+		mForceSampleBatchVelocity[mForceSampleBatchBufferPos + i] = *(PxVec4*)((uint8_t*)velocities + i * velocitiesStride);
+		mForceSampleBatchMass[mForceSampleBatchBufferPos + i] = *(float*)((uint8_t*)mass + i * massStride);
+	}
+
+	FieldSamplerQueryDataIntl& data = mForceSampleBatchQueryData[batchId];
+	data.count = numIndices;
+	data.isDataOnDevice = false;
+	data.positionStrideBytes = sizeof(PxVec4);
+	data.velocityStrideBytes = sizeof(PxVec4);
+	data.massStrideBytes = massStride ? sizeof(float) : 0;
+	data.pmaInMass = (float*)&mForceSampleBatchMass[mForceSampleBatchBufferPos];
+	data.pmaInPosition = (float*)&mForceSampleBatchPosition[mForceSampleBatchBufferPos];
+	data.pmaInVelocity = (float*)&mForceSampleBatchVelocity[mForceSampleBatchBufferPos];
+	data.pmaInIndices = (uint32_t*)indices;
+	data.pmaOutField = forces;
+	data.timeStep = getApexScene().getPhysXSimulateTime();
+
+	mForceSampleBatchBufferPos += maxIndices;
+}
+
+#endif
+
+/******************************** CPU Version ********************************/
+
+
+FieldSamplerSceneCPU::FieldSamplerSceneCPU(ModuleFieldSamplerImpl& module, SceneIntl& scene, RenderDebugInterface* debugRender, ResourceList& list) :
+	FieldSamplerScene(module, scene, debugRender, list)
+{
+#if PX_PHYSICS_VERSION_MAJOR == 3
+	mPhysXMonitor = PX_NEW(FieldSamplerPhysXMonitor)(*this);
+#endif
+}
+
+FieldSamplerSceneCPU::~FieldSamplerSceneCPU()
+{
+}
+
+FieldSamplerManager* FieldSamplerSceneCPU::createManager()
+{
+	return PX_NEW(FieldSamplerManagerCPU)(this);
+}
+
+/******************************** GPU Version ********************************/
+
+#if APEX_CUDA_SUPPORT
+
+FieldSamplerSceneGPU::FieldSamplerSceneGPU(ModuleFieldSamplerImpl& module, SceneIntl& scene, RenderDebugInterface* debugRender, ResourceList& list)
+	: FieldSamplerScene(module, scene, debugRender, list)
+	, CudaModuleScene(scene, *mModule, APEX_CUDA_TO_STR(APEX_CUDA_MODULE_PREFIX))
+{
+#if PX_PHYSICS_VERSION_MAJOR == 3
+	mPhysXMonitor = PX_NEW(FieldSamplerPhysXMonitor)(*this);
+#endif
+	{
+		PxGpuDispatcher* gd = mApexScene->getTaskManager()->getGpuDispatcher();
+		PX_ASSERT(gd != NULL);
+		mCtxMgr = gd->getCudaContextManager();
+		PxScopedCudaLock _lock_(*mCtxMgr);
+
+//CUDA module objects
+#include "../cuda/include/fieldsampler.h"
+	}
+}
+
+FieldSamplerSceneGPU::~FieldSamplerSceneGPU()
+{
+	CudaModuleScene::destroy(*mApexScene);
+}
+
+FieldSamplerManager* FieldSamplerSceneGPU::createManager()
+{
+	return PX_NEW(FieldSamplerManagerGPU)(this);
+}
+
+#endif
+
+}
+} // end namespace nvidia::apex
+
+
diff --git a/APEX_1.4/module/fieldsampler/src/FieldSamplerSceneWrapper.cpp b/APEX_1.4/module/fieldsampler/src/FieldSamplerSceneWrapper.cpp
new file mode 100644
index 00000000..405462e7
--- /dev/null
+++ b/APEX_1.4/module/fieldsampler/src/FieldSamplerSceneWrapper.cpp
@@ -0,0 +1,143 @@
+/*
+ * Copyright (c) 2008-2015, 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 "ApexDefs.h"
+#include "Apex.h"
+#include "FieldSamplerSceneWrapper.h"
+#include "FieldSamplerManager.h"
+#include "FieldBoundaryWrapper.h"
+#include "FieldSamplerWrapper.h"
+
+
+namespace nvidia
+{
+namespace fieldsampler
+{
+
+FieldSamplerSceneWrapper::FieldSamplerSceneWrapper(ResourceList& list, FieldSamplerManager* manager, FieldSamplerSceneIntl* fieldSamplerScene)
+	: mManager(manager)
+	, mFieldSamplerScene(fieldSamplerScene)
+	, mFieldBoundaryListChanged(false)
+{
+	mFieldSamplerScene->getFieldSamplerSceneDesc(mFieldSamplerSceneDesc);
+
+	list.add(*this);
+}
+
+void FieldSamplerSceneWrapper::release()
+{
+	delete this;
+}
+
+FieldSamplerSceneWrapper::FieldBoundaryInfo* FieldSamplerSceneWrapper::addFieldBoundary(FieldBoundaryWrapper* fieldBoundaryWrapper)
+{
+	FieldBoundaryInfo* fieldBoundaryInfo = NULL;
+	for (uint32_t i = 0; i < mFieldBoundaryList.getSize(); ++i)
+	{
+		FieldBoundaryInfo* info = DYNAMIC_CAST(FieldBoundaryInfo*)(mFieldBoundaryList.getResource(i));
+		if (info->getFieldBoundaryWrapper() == fieldBoundaryWrapper)
+		{
+			fieldBoundaryInfo = info;
+			break;
+		}
+	}
+	if (fieldBoundaryInfo == NULL)
+	{
+		fieldBoundaryInfo = createFieldBoundaryInfo(fieldBoundaryWrapper);
+		mFieldBoundaryListChanged = true;
+	}
+	fieldBoundaryInfo->addRef();
+	return fieldBoundaryInfo;
+}
+
+void FieldSamplerSceneWrapper::removeFieldBoundary(FieldBoundaryInfo* fieldBoundaryInfo)
+{
+	if (fieldBoundaryInfo->releaseRef())
+	{
+		mFieldBoundaryListChanged = true;
+	}
+}
+
+void FieldSamplerSceneWrapper::update()
+{
+	for (uint32_t i = 0; i < mFieldBoundaryList.getSize(); ++i)
+	{
+		FieldBoundaryInfo* info = DYNAMIC_CAST(FieldBoundaryInfo*)(mFieldBoundaryList.getResource(i));
+		info->update();
+	}
+}
+
+/******************************** CPU Version ********************************/
+FieldSamplerSceneWrapperCPU::FieldSamplerSceneWrapperCPU(ResourceList& list, FieldSamplerManager* manager, FieldSamplerSceneIntl* fieldSamplerScene)
+	: FieldSamplerSceneWrapper(list, manager, fieldSamplerScene)
+{
+}
+
+/******************************** GPU Version ********************************/
+#if APEX_CUDA_SUPPORT
+
+FieldSamplerSceneWrapperGPU::FieldSamplerSceneWrapperGPU(ResourceList& list, FieldSamplerManager* manager, FieldSamplerSceneIntl* fieldSamplerScene)
+	: FieldSamplerSceneWrapper(list, manager, fieldSamplerScene)
+	, mConstStorage(*fieldSamplerScene->getFieldSamplerCudaConstStorage())
+{
+}
+
+void FieldSamplerSceneWrapperGPU::postUpdate()
+{
+	PxCudaContextManager* ctx = DYNAMIC_CAST(FieldSamplerSceneGPU*)(mManager->getScene())->getCudaContext();
+	PxScopedCudaLock _lock_(*ctx);
+
+	getConstStorage().copyToDevice(ctx, 0);
+}
+
+
+FieldSamplerSceneWrapperGPU::FieldBoundaryInfoGPU::FieldBoundaryInfoGPU(ResourceList& list, FieldBoundaryWrapper* fieldBoundaryWrapper, ApexCudaConstStorage& constStorage)
+	: FieldBoundaryInfo(list, fieldBoundaryWrapper)
+	, mConstMemGroup(constStorage)
+{
+	APEX_CUDA_CONST_MEM_GROUP_SCOPE(mConstMemGroup);
+	mFieldShapeGroupParamsHandle.alloc(_storage_);
+}
+
+void FieldSamplerSceneWrapperGPU::FieldBoundaryInfoGPU::update()
+{
+	if (mFieldBoundaryWrapper->getFieldShapesChanged())
+	{
+		APEX_CUDA_CONST_MEM_GROUP_SCOPE(mConstMemGroup);
+
+		FieldShapeGroupParams shapeGroupParams;
+		mFieldShapeGroupParamsHandle.fetch(_storage_, shapeGroupParams);
+
+		const nvidia::Array<FieldShapeDescIntl>& shapes = mFieldBoundaryWrapper->getFieldShapes();
+		uint32_t shapeCount = shapes.size();
+		shapeGroupParams.shapeArray.resize(_storage_, shapeCount);
+		for (uint32_t i = 0; i < shapeCount; ++i)
+		{
+			FieldShapeParams elem;
+			elem.type          = shapes[i].type;
+			elem.dimensions    = shapes[i].dimensions;
+			elem.worldToShape  = shapes[i].worldToShape;
+			PX_ASSERT(shapes[i].weight >= 0.0f && shapes[i].weight <= 1.0f);
+			elem.weight        = PxClamp(shapes[i].weight, 0.0f, 1.0f);
+			elem.fade          = 0;
+
+			shapeGroupParams.shapeArray.updateElem(_storage_, elem, i);
+		}
+
+		mFieldShapeGroupParamsHandle.update(_storage_, shapeGroupParams);
+	}
+}
+
+#endif // APEX_CUDA_SUPPORT
+
+}
+} // end namespace nvidia::apex
+
diff --git a/APEX_1.4/module/fieldsampler/src/FieldSamplerWrapper.cpp b/APEX_1.4/module/fieldsampler/src/FieldSamplerWrapper.cpp
new file mode 100644
index 00000000..94d18a76
--- /dev/null
+++ b/APEX_1.4/module/fieldsampler/src/FieldSamplerWrapper.cpp
@@ -0,0 +1,177 @@
+/*
+ * Copyright (c) 2008-2015, 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 "ApexDefs.h"
+#include "Apex.h"
+#include "FieldSamplerWrapper.h"
+#include "FieldBoundaryWrapper.h"
+#include "FieldSamplerManager.h"
+#include "FieldSamplerSceneWrapper.h"
+
+
+namespace nvidia
+{
+namespace fieldsampler
+{
+
+
+FieldSamplerWrapper::FieldSamplerWrapper(ResourceList& list, FieldSamplerManager* manager, FieldSamplerIntl* fieldSampler, const FieldSamplerDescIntl& fieldSamplerDesc, FieldSamplerSceneWrapper* fieldSamplerSceneWrapper)
+	: mManager(manager)
+	, mFieldSampler(fieldSampler)
+	, mFieldSamplerDesc(fieldSamplerDesc)
+	, mFieldSamplerShapeChanged(false)
+	, mSceneWrapper(fieldSamplerSceneWrapper)
+	, mQueryRefCount(0)
+	, mFieldBoundaryInfoArrayChanged(true)
+	, mIsEnabled(false)
+	, mIsEnabledLast(false)
+{
+	list.add(*this);
+
+	//set default shape weight to 1
+	mFieldSamplerShape.weight = 1;
+}
+
+void FieldSamplerWrapper::release()
+{
+	for (uint32_t i = 0; i < mFieldBoundaryInfoArray.size(); ++i)
+	{
+		FieldSamplerSceneWrapper::FieldBoundaryInfo* fieldBoundaryInfo = mFieldBoundaryInfoArray[i];
+		mSceneWrapper->removeFieldBoundary(fieldBoundaryInfo);
+	}
+
+	delete this;
+}
+
+bool FieldSamplerWrapper::addFieldBoundary(FieldBoundaryWrapper* fieldBoundaryWrapper)
+{
+	const FieldBoundaryDescIntl& fieldBoundaryDesc = fieldBoundaryWrapper->getInternalFieldBoundaryDesc();
+
+	float weight;
+	if (mManager->getFieldSamplerGroupsFiltering(fieldBoundaryDesc.boundaryFilterData, mFieldSamplerDesc.boundaryFilterData,weight))
+	{
+		FieldSamplerSceneWrapper::FieldBoundaryInfo* fieldBoundaryInfo =
+		    mSceneWrapper->addFieldBoundary(fieldBoundaryWrapper);
+
+		mFieldBoundaryInfoArray.pushBack(fieldBoundaryInfo);
+		mFieldBoundaryInfoArrayChanged = true;
+		return true;
+	}
+	return false;
+}
+
+bool FieldSamplerWrapper::removeFieldBoundary(FieldBoundaryWrapper* fieldBoundaryWrapper)
+{
+	for (uint32_t i = 0; i < mFieldBoundaryInfoArray.size(); ++i)
+	{
+		FieldSamplerSceneWrapper::FieldBoundaryInfo* fieldBoundaryInfo = mFieldBoundaryInfoArray[i];
+		if (fieldBoundaryInfo->getFieldBoundaryWrapper() == fieldBoundaryWrapper)
+		{
+			mSceneWrapper->removeFieldBoundary(fieldBoundaryInfo);
+
+			mFieldBoundaryInfoArray.replaceWithLast(i);
+			mFieldBoundaryInfoArrayChanged = true;
+			return true;
+		}
+	}
+	return false;
+}
+
+void FieldSamplerWrapper::update()
+{
+	mIsEnabledLast = mIsEnabled;
+	mFieldSamplerShapeChanged = mFieldSampler->updateFieldSampler(mFieldSamplerShape, mIsEnabled);
+	if (!mIsEnabledLast && mIsEnabled)
+	{
+		mFieldSamplerShapeChanged = true;
+	}
+	else if (!mIsEnabled)
+	{
+		mFieldSamplerShapeChanged = false;
+	}
+}
+
+/******************************** CPU Version ********************************/
+
+FieldSamplerWrapperCPU::FieldSamplerWrapperCPU(ResourceList& list, FieldSamplerManager* manager, FieldSamplerIntl* fieldSampler, const FieldSamplerDescIntl& fieldSamplerDesc, FieldSamplerSceneWrapper* fieldSamplerSceneWrapper)
+	: FieldSamplerWrapper(list, manager, fieldSampler, fieldSamplerDesc, fieldSamplerSceneWrapper)
+{
+}
+
+/******************************** GPU Version ********************************/
+#if APEX_CUDA_SUPPORT
+
+FieldSamplerWrapperGPU::FieldSamplerWrapperGPU(ResourceList& list, FieldSamplerManager* manager, FieldSamplerIntl* fieldSampler, const FieldSamplerDescIntl& fieldSamplerDesc, FieldSamplerSceneWrapper* fieldSamplerSceneWrapper)
+	: FieldSamplerWrapperCPU(list, manager, fieldSampler, fieldSamplerDesc, fieldSamplerSceneWrapper)
+	, mConstMemGroup(DYNAMIC_CAST(FieldSamplerSceneWrapperGPU*)(fieldSamplerSceneWrapper)->getConstStorage())
+{
+}
+
+void FieldSamplerWrapperGPU::update()
+{
+	FieldSamplerWrapper::update();
+	if (mFieldSamplerShapeChanged || mFieldBoundaryInfoArrayChanged)
+	{
+		APEX_CUDA_CONST_MEM_GROUP_SCOPE(mConstMemGroup);
+
+		FieldSamplerParams params;
+		if (mFieldSamplerParamsHandle.allocOrFetch(_storage_, params))
+		{
+			//only on alloc
+			params.type = mFieldSamplerDesc.type;
+			params.gridSupportType = mFieldSamplerDesc.gridSupportType;
+			params.dragCoeff = mFieldSamplerDesc.dragCoeff;
+
+			params.includeShape.fade           = PxClamp(mFieldSamplerDesc.boundaryFadePercentage, 0.0f, 1.0f);
+		}
+		
+		if (mFieldSamplerShapeChanged)
+		{
+			FieldSamplerIntl::CudaExecuteInfo executeInfo;
+			mFieldSampler->getFieldSamplerCudaExecuteInfo(executeInfo);
+
+			params.executeType = executeInfo.executeType;
+			params.executeParamsHandle = executeInfo.executeParamsHandle;
+
+			params.includeShape.type           = mFieldSamplerShape.type;
+			params.includeShape.dimensions     = mFieldSamplerShape.dimensions;
+			params.includeShape.worldToShape   = mFieldSamplerShape.worldToShape;
+			PX_ASSERT(mFieldSamplerShape.weight >= 0.0f && mFieldSamplerShape.weight <= 1.0f);
+			params.includeShape.weight         = PxClamp(mFieldSamplerShape.weight, 0.0f, 1.0f);
+		}
+
+		if (mFieldBoundaryInfoArrayChanged)
+		{
+			uint32_t shapeGroupCount = mFieldBoundaryInfoArray.size();
+			if (params.excludeShapeGroupHandleArray.resize(_storage_, shapeGroupCount))
+			{
+				for (uint32_t shapeGroupIndex = 0; shapeGroupIndex < shapeGroupCount; ++shapeGroupIndex)
+				{
+					FieldSamplerSceneWrapperGPU::FieldBoundaryInfoGPU* fieldBoundaryInfo =
+						static_cast<FieldSamplerSceneWrapperGPU::FieldBoundaryInfoGPU*>(mFieldBoundaryInfoArray[shapeGroupIndex]);
+
+					InplaceHandle<FieldShapeGroupParams> elem = fieldBoundaryInfo->getShapeGroupParamsHandle();
+					params.excludeShapeGroupHandleArray.updateElem(_storage_, elem, shapeGroupIndex);
+				}
+			}
+
+			mFieldBoundaryInfoArrayChanged = false;
+		}
+
+		mFieldSamplerParamsHandle.update(_storage_, params);
+	}
+}
+
+#endif
+
+}
+} // end namespace nvidia::apex
+
diff --git a/APEX_1.4/module/fieldsampler/src/ModuleFieldSamplerImpl.cpp b/APEX_1.4/module/fieldsampler/src/ModuleFieldSamplerImpl.cpp
new file mode 100644
index 00000000..7634c180
--- /dev/null
+++ b/APEX_1.4/module/fieldsampler/src/ModuleFieldSamplerImpl.cpp
@@ -0,0 +1,281 @@
+/*
+ * Copyright (c) 2008-2015, 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 "ApexDefs.h"
+#include "Apex.h"
+#include "ModuleFieldSamplerImpl.h"
+#include "ModuleFieldSamplerRegistration.h"
+#include "FieldSamplerScene.h"
+#include "FieldSamplerManager.h"
+#include "SceneIntl.h"
+#include "PsMemoryBuffer.h"
+#include "ModulePerfScope.h"
+using namespace fieldsampler;
+
+#include "ApexSDKIntl.h"
+#include "ApexUsingNamespace.h"
+
+#include "Lock.h"
+
+#include "ReadCheck.h"
+#include "WriteCheck.h"
+
+namespace nvidia
+{
+namespace apex
+{
+
+#if defined(_USRDLL)
+
+/* Modules don't have to link against the framework, they keep their own */
+ApexSDKIntl* gApexSdk = 0;
+ApexSDK* GetApexSDK()
+{
+	return gApexSdk;
+}
+ApexSDKIntl* GetInternalApexSDK()
+{
+	return gApexSdk;
+}
+
+APEX_API Module*  CALL_CONV createModule(
+    ApexSDKIntl* inSdk,
+    ModuleIntl** niRef,
+    uint32_t APEXsdkVersion,
+    uint32_t PhysXsdkVersion,
+    ApexCreateError* errorCode)
+{
+	if (APEXsdkVersion != APEX_SDK_VERSION)
+	{
+		if (errorCode)
+		{
+			*errorCode = APEX_CE_WRONG_VERSION;
+		}
+		return NULL;
+	}
+
+	if (PhysXsdkVersion != PX_PHYSICS_VERSION)
+	{
+		if (errorCode)
+		{
+			*errorCode = APEX_CE_WRONG_VERSION;
+		}
+		return NULL;
+	}
+
+	gApexSdk = inSdk;
+	ModuleFieldSamplerImpl* impl = PX_NEW(ModuleFieldSampler)(inSdk);
+	*niRef  = (ModuleIntl*) impl;
+	return (Module*) impl;
+}
+
+#else
+/* Statically linking entry function */
+void instantiateModuleFieldSampler()
+{
+	ApexSDKIntl* sdk = GetInternalApexSDK();
+	nvidia::fieldsampler::ModuleFieldSamplerImpl* impl = PX_NEW(fieldsampler::ModuleFieldSamplerImpl)(sdk);
+	sdk->registerExternalModule((Module*) impl, (ModuleIntl*) impl);
+}
+#endif // `defined(_USRDLL)
+}
+
+namespace fieldsampler
+{
+/* === ModuleFieldSamplerImpl Implementation === */
+ModuleFieldSamplerImpl::ModuleFieldSamplerImpl(ApexSDKIntl* sdk)
+{
+	mName = "FieldSampler";
+	mSdk = sdk;
+	mApiProxy = this;
+	mModuleParams = NULL;
+
+	/* Register the NvParameterized factories */
+	NvParameterized::Traits* traits = mSdk->getParameterizedTraits();
+	ModuleFieldSamplerRegistration::invokeRegistration(traits);
+}
+
+ModuleFieldSamplerImpl::~ModuleFieldSamplerImpl()
+{
+}
+
+void ModuleFieldSamplerImpl::destroy()
+{
+	NvParameterized::Traits* traits = mSdk->getParameterizedTraits();
+
+	if (mModuleParams)
+	{
+		mModuleParams->destroy();
+		mModuleParams = NULL;
+	}
+
+	ModuleBase::destroy();
+
+	if (traits)
+	{
+		/* Remove the NvParameterized factories */
+		ModuleFieldSamplerRegistration::invokeUnregistration(traits);
+	}
+
+	delete this;
+}
+
+void ModuleFieldSamplerImpl::init(NvParameterized::Interface&)
+{
+}
+
+NvParameterized::Interface* ModuleFieldSamplerImpl::getDefaultModuleDesc()
+{
+	NvParameterized::Traits* traits = mSdk->getParameterizedTraits();
+
+	if (!mModuleParams)
+	{
+		mModuleParams = DYNAMIC_CAST(FieldSamplerModuleParameters*)
+		                (traits->createNvParameterized("FieldSamplerModuleParameters"));
+		PX_ASSERT(mModuleParams);
+	}
+	else
+	{
+		mModuleParams->initDefaults();
+	}
+
+	return mModuleParams;
+}
+
+FieldSamplerManagerIntl* ModuleFieldSamplerImpl::getInternalFieldSamplerManager(const Scene& apexScene)
+{
+	FieldSamplerScene* scene = ModuleFieldSamplerImpl::getFieldSamplerScene(apexScene);
+	return scene->getManager();
+}
+
+
+AuthObjTypeID ModuleFieldSamplerImpl::getModuleID() const
+{
+	return 0;
+}
+
+
+/* == Example Scene methods == */
+ModuleSceneIntl* ModuleFieldSamplerImpl::createInternalModuleScene(SceneIntl& scene, RenderDebugInterface* debugRender)
+{
+#if APEX_CUDA_SUPPORT
+	READ_LOCK(scene);
+	if (scene.getTaskManager()->getGpuDispatcher() && scene.isUsingCuda())
+	{
+		return PX_NEW(FieldSamplerSceneGPU)(*this, scene, debugRender, mFieldSamplerScenes);
+	}
+	else
+#endif
+		return PX_NEW(FieldSamplerSceneCPU)(*this, scene, debugRender, mFieldSamplerScenes);
+}
+
+void ModuleFieldSamplerImpl::releaseModuleSceneIntl(ModuleSceneIntl& scene)
+{
+	FieldSamplerScene* es = DYNAMIC_CAST(FieldSamplerScene*)(&scene);
+	es->destroy();
+}
+
+nvidia::fieldsampler::FieldSamplerScene* ModuleFieldSamplerImpl::getFieldSamplerScene(const Scene& apexScene) const
+{
+	for (uint32_t i = 0 ; i < mFieldSamplerScenes.getSize() ; i++)
+	{
+		FieldSamplerScene* es = DYNAMIC_CAST(FieldSamplerScene*)(mFieldSamplerScenes.getResource(i));
+		if (es->mApexScene == &apexScene)
+		{
+			return es;
+		}
+	}
+
+	PX_ASSERT(!"Unable to locate an appropriate FieldSamplerScene");
+	return NULL;
+}
+
+RenderableIterator* ModuleFieldSamplerImpl::createRenderableIterator(const Scene& apexScene)
+{
+	FieldSamplerScene* es = getFieldSamplerScene(apexScene);
+	if (es)
+	{
+		return es->createRenderableIterator();
+	}
+
+	return NULL;
+}
+
+
+
+bool ModuleFieldSamplerImpl::setFieldSamplerWeightedCollisionFilterCallback(const Scene& apexScene,FieldSamplerWeightedCollisionFilterCallback *callback)
+{
+	WRITE_ZONE();
+	FieldSamplerScene* scene = getFieldSamplerScene(apexScene);
+	if (scene != NULL)
+	{
+		DYNAMIC_CAST(FieldSamplerManager*)(scene->getManager())->setFieldSamplerWeightedCollisionFilterCallback(callback);
+		return true;
+	}
+	return false;
+
+}
+
+#if PX_PHYSICS_VERSION_MAJOR == 3
+void ModuleFieldSamplerImpl::enablePhysXMonitor(const Scene& apexScene, bool enable)
+{
+	WRITE_ZONE();
+	getFieldSamplerScene(apexScene)->enablePhysXMonitor(enable);	
+}
+
+void ModuleFieldSamplerImpl::setPhysXMonitorFilterData(const Scene& apexScene, physx::PxFilterData filterData)
+{
+	WRITE_ZONE();
+	getFieldSamplerScene(apexScene)->setPhysXFilterData(filterData);
+}
+#endif
+
+uint32_t ModuleFieldSamplerImpl::createForceSampleBatch(const Scene& apexScene, uint32_t maxCount, const physx::PxFilterData filterData)
+{
+	WRITE_ZONE();
+	FieldSamplerScene* fsScene = getFieldSamplerScene(apexScene);
+	if (fsScene)
+	{
+		return fsScene->createForceSampleBatch(maxCount, filterData);
+	}
+	return (uint32_t)~0;
+}
+
+
+void ModuleFieldSamplerImpl::releaseForceSampleBatch(const Scene& apexScene, uint32_t batchId)
+{
+	WRITE_ZONE();
+	FieldSamplerScene* fsScene = getFieldSamplerScene(apexScene);
+	if (fsScene)
+	{
+		fsScene->releaseForceSampleBatch(batchId);
+	}
+}
+
+
+void ModuleFieldSamplerImpl::submitForceSampleBatch(const Scene& apexScene, uint32_t batchId,
+												PxVec4* forces, const uint32_t forcesStride,
+												const PxVec3* positions, const uint32_t positionsStride, 
+												const PxVec3* velocities, const uint32_t velocitiesStride,
+												const float* mass, const uint32_t massStride,
+												const uint32_t* indices, const uint32_t numIndices)
+{
+	WRITE_ZONE();
+	FieldSamplerScene* fsScene = getFieldSamplerScene(apexScene);
+	if (fsScene)
+	{
+		fsScene->submitForceSampleBatch(batchId, forces, forcesStride, positions, positionsStride, velocities, velocitiesStride, mass, massStride, indices, numIndices);
+	}
+}
+
+
+}
+} // end namespace nvidia::apex
diff --git a/APEX_1.4/module/fieldsampler/src/autogen/FieldSamplerModuleParameters.cpp b/APEX_1.4/module/fieldsampler/src/autogen/FieldSamplerModuleParameters.cpp
new file mode 100644
index 00000000..2b76ff70
--- /dev/null
+++ b/APEX_1.4/module/fieldsampler/src/autogen/FieldSamplerModuleParameters.cpp
@@ -0,0 +1,318 @@
+// This code contains NVIDIA Confidential Information and is disclosed to you
+// under a form of NVIDIA software license agreement provided separately to you.
+//
+// Notice
+// NVIDIA Corporation and its licensors retain all intellectual property and
+// proprietary rights in and to this software and 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.
+//
+// ALL NVIDIA DESIGN SPECIFICATIONS, CODE ARE PROVIDED "AS IS.". NVIDIA MAKES
+// NO WARRANTIES, EXPRESSED, IMPLIED, STATUTORY, OR OTHERWISE WITH RESPECT TO
+// THE MATERIALS, AND EXPRESSLY DISCLAIMS ALL IMPLIED WARRANTIES OF NONINFRINGEMENT,
+// MERCHANTABILITY, AND FITNESS FOR A PARTICULAR PURPOSE.
+//
+// Information and code furnished is believed to be accurate and reliable.
+// However, NVIDIA Corporation assumes no responsibility for the consequences of use of such
+// information or for any infringement of patents or other rights of third parties that may
+// result from its use. No license is granted by implication or otherwise under any patent
+// or patent rights of NVIDIA Corporation. Details are subject to change without notice.
+// This code supersedes and replaces all information previously supplied.
+// NVIDIA Corporation products are not authorized for use as critical
+// components in life support devices or systems without express written approval of
+// NVIDIA Corporation.
+//
+// Copyright (c) 2008-2015 NVIDIA Corporation. All rights reserved.
+
+// This file was generated by NvParameterized/scripts/GenParameterized.pl
+
+
+#include "FieldSamplerModuleParameters.h"
+#include <string.h>
+#include <stdlib.h>
+
+using namespace NvParameterized;
+
+namespace nvidia
+{
+namespace fieldsampler
+{
+
+using namespace FieldSamplerModuleParametersNS;
+
+const char* const FieldSamplerModuleParametersFactory::vptr =
+    NvParameterized::getVptr<FieldSamplerModuleParameters, FieldSamplerModuleParameters::ClassAlignment>();
+
+const uint32_t NumParamDefs = 2;
+static NvParameterized::DefinitionImpl* ParamDefTable; // now allocated in buildTree [NumParamDefs];
+
+
+static const size_t ParamLookupChildrenTable[] =
+{
+	1,
+};
+
+#define TENUM(type) nvidia::##type
+#define CHILDREN(index) &ParamLookupChildrenTable[index]
+static const NvParameterized::ParamLookupNode ParamLookupTable[NumParamDefs] =
+{
+	{ TYPE_STRUCT, false, 0, CHILDREN(0), 1 },
+	{ TYPE_U32, false, (size_t)(&((ParametersStruct*)0)->unused), NULL, 0 }, // unused
+};
+
+
+bool FieldSamplerModuleParameters::mBuiltFlag = false;
+NvParameterized::MutexType FieldSamplerModuleParameters::mBuiltFlagMutex;
+
+FieldSamplerModuleParameters::FieldSamplerModuleParameters(NvParameterized::Traits* traits, void* buf, int32_t* refCount) :
+	NvParameters(traits, buf, refCount)
+{
+	//mParameterizedTraits->registerFactory(className(), &FieldSamplerModuleParametersFactoryInst);
+
+	if (!buf) //Do not init data if it is inplace-deserialized
+	{
+		initDynamicArrays();
+		initStrings();
+		initReferences();
+		initDefaults();
+	}
+}
+
+FieldSamplerModuleParameters::~FieldSamplerModuleParameters()
+{
+	freeStrings();
+	freeReferences();
+	freeDynamicArrays();
+}
+
+void FieldSamplerModuleParameters::destroy()
+{
+	// We cache these fields here to avoid overwrite in destructor
+	bool doDeallocateSelf = mDoDeallocateSelf;
+	NvParameterized::Traits* traits = mParameterizedTraits;
+	int32_t* refCount = mRefCount;
+	void* buf = mBuffer;
+
+	this->~FieldSamplerModuleParameters();
+
+	NvParameters::destroy(this, traits, doDeallocateSelf, refCount, buf);
+}
+
+const NvParameterized::DefinitionImpl* FieldSamplerModuleParameters::getParameterDefinitionTree(void)
+{
+	if (!mBuiltFlag) // Double-checked lock
+	{
+		NvParameterized::MutexType::ScopedLock lock(mBuiltFlagMutex);
+		if (!mBuiltFlag)
+		{
+			buildTree();
+		}
+	}
+
+	return(&ParamDefTable[0]);
+}
+
+const NvParameterized::DefinitionImpl* FieldSamplerModuleParameters::getParameterDefinitionTree(void) const
+{
+	FieldSamplerModuleParameters* tmpParam = const_cast<FieldSamplerModuleParameters*>(this);
+
+	if (!mBuiltFlag) // Double-checked lock
+	{
+		NvParameterized::MutexType::ScopedLock lock(mBuiltFlagMutex);
+		if (!mBuiltFlag)
+		{
+			tmpParam->buildTree();
+		}
+	}
+
+	return(&ParamDefTable[0]);
+}
+
+NvParameterized::ErrorType FieldSamplerModuleParameters::getParameterHandle(const char* long_name, Handle& handle) const
+{
+	ErrorType Ret = NvParameters::getParameterHandle(long_name, handle);
+	if (Ret != ERROR_NONE)
+	{
+		return(Ret);
+	}
+
+	size_t offset;
+	void* ptr;
+
+	getVarPtr(handle, ptr, offset);
+
+	if (ptr == NULL)
+	{
+		return(ERROR_INDEX_OUT_OF_RANGE);
+	}
+
+	return(ERROR_NONE);
+}
+
+NvParameterized::ErrorType FieldSamplerModuleParameters::getParameterHandle(const char* long_name, Handle& handle)
+{
+	ErrorType Ret = NvParameters::getParameterHandle(long_name, handle);
+	if (Ret != ERROR_NONE)
+	{
+		return(Ret);
+	}
+
+	size_t offset;
+	void* ptr;
+
+	getVarPtr(handle, ptr, offset);
+
+	if (ptr == NULL)
+	{
+		return(ERROR_INDEX_OUT_OF_RANGE);
+	}
+
+	return(ERROR_NONE);
+}
+
+void FieldSamplerModuleParameters::getVarPtr(const Handle& handle, void*& ptr, size_t& offset) const
+{
+	ptr = getVarPtrHelper(&ParamLookupTable[0], const_cast<FieldSamplerModuleParameters::ParametersStruct*>(&parameters()), handle, offset);
+}
+
+
+/* Dynamic Handle Indices */
+
+void FieldSamplerModuleParameters::freeParameterDefinitionTable(NvParameterized::Traits* traits)
+{
+	if (!traits)
+	{
+		return;
+	}
+
+	if (!mBuiltFlag) // Double-checked lock
+	{
+		return;
+	}
+
+	NvParameterized::MutexType::ScopedLock lock(mBuiltFlagMutex);
+
+	if (!mBuiltFlag)
+	{
+		return;
+	}
+
+	for (uint32_t i = 0; i < NumParamDefs; ++i)
+	{
+		ParamDefTable[i].~DefinitionImpl();
+	}
+
+	traits->free(ParamDefTable);
+
+	mBuiltFlag = false;
+}
+
+#define PDEF_PTR(index) (&ParamDefTable[index])
+
+void FieldSamplerModuleParameters::buildTree(void)
+{
+
+	uint32_t allocSize = sizeof(NvParameterized::DefinitionImpl) * NumParamDefs;
+	ParamDefTable = (NvParameterized::DefinitionImpl*)(mParameterizedTraits->alloc(allocSize));
+	memset(ParamDefTable, 0, allocSize);
+
+	for (uint32_t i = 0; i < NumParamDefs; ++i)
+	{
+		NV_PARAM_PLACEMENT_NEW(ParamDefTable + i, NvParameterized::DefinitionImpl)(*mParameterizedTraits);
+	}
+
+	// Initialize DefinitionImpl node: nodeIndex=0, longName=""
+	{
+		NvParameterized::DefinitionImpl* ParamDef = &ParamDefTable[0];
+		ParamDef->init("", TYPE_STRUCT, "STRUCT", true);
+
+#ifdef NV_PARAMETERIZED_HIDE_DESCRIPTIONS
+
+#else
+
+		static HintImpl HintTable[1];
+		static Hint* HintPtrTable[1] = { &HintTable[0], };
+		HintTable[0].init("shortDescription", "This class is used for initializing the ModuleFieldSampler.", true);
+		ParamDefTable[0].setHints((const NvParameterized::Hint**)HintPtrTable, 1);
+
+#endif /* NV_PARAMETERIZED_HIDE_DESCRIPTIONS */
+
+
+
+
+
+	}
+
+	// Initialize DefinitionImpl node: nodeIndex=1, longName="unused"
+	{
+		NvParameterized::DefinitionImpl* ParamDef = &ParamDefTable[1];
+		ParamDef->init("unused", TYPE_U32, NULL, true);
+
+#ifdef NV_PARAMETERIZED_HIDE_DESCRIPTIONS
+
+#else
+
+		static HintImpl HintTable[1];
+		static Hint* HintPtrTable[1] = { &HintTable[0], };
+		HintTable[0].init("shortDescription", "No parameters necessary", true);
+		ParamDefTable[1].setHints((const NvParameterized::Hint**)HintPtrTable, 1);
+
+#endif /* NV_PARAMETERIZED_HIDE_DESCRIPTIONS */
+
+
+
+
+
+	}
+
+	// SetChildren for: nodeIndex=0, longName=""
+	{
+		static Definition* Children[1];
+		Children[0] = PDEF_PTR(1);
+
+		ParamDefTable[0].setChildren(Children, 1);
+	}
+
+	mBuiltFlag = true;
+
+}
+void FieldSamplerModuleParameters::initStrings(void)
+{
+}
+
+void FieldSamplerModuleParameters::initDynamicArrays(void)
+{
+}
+
+void FieldSamplerModuleParameters::initDefaults(void)
+{
+
+	freeStrings();
+	freeReferences();
+	freeDynamicArrays();
+	unused = uint32_t(0);
+
+	initDynamicArrays();
+	initStrings();
+	initReferences();
+}
+
+void FieldSamplerModuleParameters::initReferences(void)
+{
+}
+
+void FieldSamplerModuleParameters::freeDynamicArrays(void)
+{
+}
+
+void FieldSamplerModuleParameters::freeStrings(void)
+{
+}
+
+void FieldSamplerModuleParameters::freeReferences(void)
+{
+}
+
+} // namespace fieldsampler
+} // namespace nvidia
diff --git a/APEX_1.4/module/fieldsampler/src/autogen/FieldSamplerPhysXMonitorParams.cpp b/APEX_1.4/module/fieldsampler/src/autogen/FieldSamplerPhysXMonitorParams.cpp
new file mode 100644
index 00000000..06762a63
--- /dev/null
+++ b/APEX_1.4/module/fieldsampler/src/autogen/FieldSamplerPhysXMonitorParams.cpp
@@ -0,0 +1,358 @@
+// This code contains NVIDIA Confidential Information and is disclosed to you
+// under a form of NVIDIA software license agreement provided separately to you.
+//
+// Notice
+// NVIDIA Corporation and its licensors retain all intellectual property and
+// proprietary rights in and to this software and 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.
+//
+// ALL NVIDIA DESIGN SPECIFICATIONS, CODE ARE PROVIDED "AS IS.". NVIDIA MAKES
+// NO WARRANTIES, EXPRESSED, IMPLIED, STATUTORY, OR OTHERWISE WITH RESPECT TO
+// THE MATERIALS, AND EXPRESSLY DISCLAIMS ALL IMPLIED WARRANTIES OF NONINFRINGEMENT,
+// MERCHANTABILITY, AND FITNESS FOR A PARTICULAR PURPOSE.
+//
+// Information and code furnished is believed to be accurate and reliable.
+// However, NVIDIA Corporation assumes no responsibility for the consequences of use of such
+// information or for any infringement of patents or other rights of third parties that may
+// result from its use. No license is granted by implication or otherwise under any patent
+// or patent rights of NVIDIA Corporation. Details are subject to change without notice.
+// This code supersedes and replaces all information previously supplied.
+// NVIDIA Corporation products are not authorized for use as critical
+// components in life support devices or systems without express written approval of
+// NVIDIA Corporation.
+//
+// Copyright (c) 2008-2015 NVIDIA Corporation. All rights reserved.
+
+// This file was generated by NvParameterized/scripts/GenParameterized.pl
+
+
+#include "FieldSamplerPhysXMonitorParams.h"
+#include <string.h>
+#include <stdlib.h>
+
+using namespace NvParameterized;
+
+namespace nvidia
+{
+namespace fieldsampler
+{
+
+using namespace FieldSamplerPhysXMonitorParamsNS;
+
+const char* const FieldSamplerPhysXMonitorParamsFactory::vptr =
+    NvParameterized::getVptr<FieldSamplerPhysXMonitorParams, FieldSamplerPhysXMonitorParams::ClassAlignment>();
+
+const uint32_t NumParamDefs = 4;
+static NvParameterized::DefinitionImpl* ParamDefTable; // now allocated in buildTree [NumParamDefs];
+
+
+static const size_t ParamLookupChildrenTable[] =
+{
+	1, 2, 3,
+};
+
+#define TENUM(type) nvidia::##type
+#define CHILDREN(index) &ParamLookupChildrenTable[index]
+static const NvParameterized::ParamLookupNode ParamLookupTable[NumParamDefs] =
+{
+	{ TYPE_STRUCT, false, 0, CHILDREN(0), 3 },
+	{ TYPE_U32, false, (size_t)(&((ParametersStruct*)0)->maxRBCount), NULL, 0 }, // maxRBCount
+	{ TYPE_U32, false, (size_t)(&((ParametersStruct*)0)->maxPSCount), NULL, 0 }, // maxPSCount
+	{ TYPE_U32, false, (size_t)(&((ParametersStruct*)0)->maxParticleCount), NULL, 0 }, // maxParticleCount
+};
+
+
+bool FieldSamplerPhysXMonitorParams::mBuiltFlag = false;
+NvParameterized::MutexType FieldSamplerPhysXMonitorParams::mBuiltFlagMutex;
+
+FieldSamplerPhysXMonitorParams::FieldSamplerPhysXMonitorParams(NvParameterized::Traits* traits, void* buf, int32_t* refCount) :
+	NvParameters(traits, buf, refCount)
+{
+	//mParameterizedTraits->registerFactory(className(), &FieldSamplerPhysXMonitorParamsFactoryInst);
+
+	if (!buf) //Do not init data if it is inplace-deserialized
+	{
+		initDynamicArrays();
+		initStrings();
+		initReferences();
+		initDefaults();
+	}
+}
+
+FieldSamplerPhysXMonitorParams::~FieldSamplerPhysXMonitorParams()
+{
+	freeStrings();
+	freeReferences();
+	freeDynamicArrays();
+}
+
+void FieldSamplerPhysXMonitorParams::destroy()
+{
+	// We cache these fields here to avoid overwrite in destructor
+	bool doDeallocateSelf = mDoDeallocateSelf;
+	NvParameterized::Traits* traits = mParameterizedTraits;
+	int32_t* refCount = mRefCount;
+	void* buf = mBuffer;
+
+	this->~FieldSamplerPhysXMonitorParams();
+
+	NvParameters::destroy(this, traits, doDeallocateSelf, refCount, buf);
+}
+
+const NvParameterized::DefinitionImpl* FieldSamplerPhysXMonitorParams::getParameterDefinitionTree(void)
+{
+	if (!mBuiltFlag) // Double-checked lock
+	{
+		NvParameterized::MutexType::ScopedLock lock(mBuiltFlagMutex);
+		if (!mBuiltFlag)
+		{
+			buildTree();
+		}
+	}
+
+	return(&ParamDefTable[0]);
+}
+
+const NvParameterized::DefinitionImpl* FieldSamplerPhysXMonitorParams::getParameterDefinitionTree(void) const
+{
+	FieldSamplerPhysXMonitorParams* tmpParam = const_cast<FieldSamplerPhysXMonitorParams*>(this);
+
+	if (!mBuiltFlag) // Double-checked lock
+	{
+		NvParameterized::MutexType::ScopedLock lock(mBuiltFlagMutex);
+		if (!mBuiltFlag)
+		{
+			tmpParam->buildTree();
+		}
+	}
+
+	return(&ParamDefTable[0]);
+}
+
+NvParameterized::ErrorType FieldSamplerPhysXMonitorParams::getParameterHandle(const char* long_name, Handle& handle) const
+{
+	ErrorType Ret = NvParameters::getParameterHandle(long_name, handle);
+	if (Ret != ERROR_NONE)
+	{
+		return(Ret);
+	}
+
+	size_t offset;
+	void* ptr;
+
+	getVarPtr(handle, ptr, offset);
+
+	if (ptr == NULL)
+	{
+		return(ERROR_INDEX_OUT_OF_RANGE);
+	}
+
+	return(ERROR_NONE);
+}
+
+NvParameterized::ErrorType FieldSamplerPhysXMonitorParams::getParameterHandle(const char* long_name, Handle& handle)
+{
+	ErrorType Ret = NvParameters::getParameterHandle(long_name, handle);
+	if (Ret != ERROR_NONE)
+	{
+		return(Ret);
+	}
+
+	size_t offset;
+	void* ptr;
+
+	getVarPtr(handle, ptr, offset);
+
+	if (ptr == NULL)
+	{
+		return(ERROR_INDEX_OUT_OF_RANGE);
+	}
+
+	return(ERROR_NONE);
+}
+
+void FieldSamplerPhysXMonitorParams::getVarPtr(const Handle& handle, void*& ptr, size_t& offset) const
+{
+	ptr = getVarPtrHelper(&ParamLookupTable[0], const_cast<FieldSamplerPhysXMonitorParams::ParametersStruct*>(&parameters()), handle, offset);
+}
+
+
+/* Dynamic Handle Indices */
+
+void FieldSamplerPhysXMonitorParams::freeParameterDefinitionTable(NvParameterized::Traits* traits)
+{
+	if (!traits)
+	{
+		return;
+	}
+
+	if (!mBuiltFlag) // Double-checked lock
+	{
+		return;
+	}
+
+	NvParameterized::MutexType::ScopedLock lock(mBuiltFlagMutex);
+
+	if (!mBuiltFlag)
+	{
+		return;
+	}
+
+	for (uint32_t i = 0; i < NumParamDefs; ++i)
+	{
+		ParamDefTable[i].~DefinitionImpl();
+	}
+
+	traits->free(ParamDefTable);
+
+	mBuiltFlag = false;
+}
+
+#define PDEF_PTR(index) (&ParamDefTable[index])
+
+void FieldSamplerPhysXMonitorParams::buildTree(void)
+{
+
+	uint32_t allocSize = sizeof(NvParameterized::DefinitionImpl) * NumParamDefs;
+	ParamDefTable = (NvParameterized::DefinitionImpl*)(mParameterizedTraits->alloc(allocSize));
+	memset(ParamDefTable, 0, allocSize);
+
+	for (uint32_t i = 0; i < NumParamDefs; ++i)
+	{
+		NV_PARAM_PLACEMENT_NEW(ParamDefTable + i, NvParameterized::DefinitionImpl)(*mParameterizedTraits);
+	}
+
+	// Initialize DefinitionImpl node: nodeIndex=0, longName=""
+	{
+		NvParameterized::DefinitionImpl* ParamDef = &ParamDefTable[0];
+		ParamDef->init("", TYPE_STRUCT, "STRUCT", true);
+
+
+
+
+
+
+	}
+
+	// Initialize DefinitionImpl node: nodeIndex=1, longName="maxRBCount"
+	{
+		NvParameterized::DefinitionImpl* ParamDef = &ParamDefTable[1];
+		ParamDef->init("maxRBCount", TYPE_U32, NULL, true);
+
+#ifdef NV_PARAMETERIZED_HIDE_DESCRIPTIONS
+
+#else
+
+		static HintImpl HintTable[1];
+		static Hint* HintPtrTable[1] = { &HintTable[0], };
+		HintTable[0].init("shortDescription", "Maximum number of rigid bodies in scene", true);
+		ParamDefTable[1].setHints((const NvParameterized::Hint**)HintPtrTable, 1);
+
+#endif /* NV_PARAMETERIZED_HIDE_DESCRIPTIONS */
+
+
+
+
+
+	}
+
+	// Initialize DefinitionImpl node: nodeIndex=2, longName="maxPSCount"
+	{
+		NvParameterized::DefinitionImpl* ParamDef = &ParamDefTable[2];
+		ParamDef->init("maxPSCount", TYPE_U32, NULL, true);
+
+#ifdef NV_PARAMETERIZED_HIDE_DESCRIPTIONS
+
+#else
+
+		static HintImpl HintTable[1];
+		static Hint* HintPtrTable[1] = { &HintTable[0], };
+		HintTable[0].init("shortDescription", "Maximum number of particle systems in scene", true);
+		ParamDefTable[2].setHints((const NvParameterized::Hint**)HintPtrTable, 1);
+
+#endif /* NV_PARAMETERIZED_HIDE_DESCRIPTIONS */
+
+
+
+
+
+	}
+
+	// Initialize DefinitionImpl node: nodeIndex=3, longName="maxParticleCount"
+	{
+		NvParameterized::DefinitionImpl* ParamDef = &ParamDefTable[3];
+		ParamDef->init("maxParticleCount", TYPE_U32, NULL, true);
+
+#ifdef NV_PARAMETERIZED_HIDE_DESCRIPTIONS
+
+#else
+
+		static HintImpl HintTable[1];
+		static Hint* HintPtrTable[1] = { &HintTable[0], };
+		HintTable[0].init("shortDescription", "Maximum number of particles in scene", true);
+		ParamDefTable[3].setHints((const NvParameterized::Hint**)HintPtrTable, 1);
+
+#endif /* NV_PARAMETERIZED_HIDE_DESCRIPTIONS */
+
+
+
+
+
+	}
+
+	// SetChildren for: nodeIndex=0, longName=""
+	{
+		static Definition* Children[3];
+		Children[0] = PDEF_PTR(1);
+		Children[1] = PDEF_PTR(2);
+		Children[2] = PDEF_PTR(3);
+
+		ParamDefTable[0].setChildren(Children, 3);
+	}
+
+	mBuiltFlag = true;
+
+}
+void FieldSamplerPhysXMonitorParams::initStrings(void)
+{
+}
+
+void FieldSamplerPhysXMonitorParams::initDynamicArrays(void)
+{
+}
+
+void FieldSamplerPhysXMonitorParams::initDefaults(void)
+{
+
+	freeStrings();
+	freeReferences();
+	freeDynamicArrays();
+	maxRBCount = uint32_t(10000);
+	maxPSCount = uint32_t(128);
+	maxParticleCount = uint32_t(100000);
+
+	initDynamicArrays();
+	initStrings();
+	initReferences();
+}
+
+void FieldSamplerPhysXMonitorParams::initReferences(void)
+{
+}
+
+void FieldSamplerPhysXMonitorParams::freeDynamicArrays(void)
+{
+}
+
+void FieldSamplerPhysXMonitorParams::freeStrings(void)
+{
+}
+
+void FieldSamplerPhysXMonitorParams::freeReferences(void)
+{
+}
+
+} // namespace fieldsampler
+} // namespace nvidia