Initial commit:

PhysX 3.4.0 Update @ 21294896
APEX 1.4.0 Update @ 21275617

[CL 21300167]

9 files changed, 393 insertions, 0 deletions

diff --git a/APEX_1.4/module/iofx/cuda/include/actorRanges.h b/APEX_1.4/module/iofx/cuda/include/actorRanges.h
new file mode 100644
index 00000000..57271da6
--- /dev/null
+++ b/APEX_1.4/module/iofx/cuda/include/actorRanges.h
@@ -0,0 +1,16 @@
+/*
+ * 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_BOUND_KERNEL(ACTOR_RANGE_KERNEL_CONFIG, actorRangeKernel,
+                       ((const unsigned int*, sortedActorID))((unsigned int, maxActorID))
+                       ((unsigned int*, actorStart))((unsigned int*, actorEnd))((unsigned int*, actorVisibleEnd))
+                       ((const uint32_t*, sortedStateID))
+                      )
diff --git a/APEX_1.4/module/iofx/cuda/include/bbox.h b/APEX_1.4/module/iofx/cuda/include/bbox.h
new file mode 100644
index 00000000..e58df0d6
--- /dev/null
+++ b/APEX_1.4/module/iofx/cuda/include/bbox.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.
+ */
+
+
+APEX_CUDA_TEXTURE_1D(texRefBBoxPositions, float4)
+
+APEX_CUDA_SYNC_KERNEL(BBOX_KERNEL_CONFIG, bboxSyncKernel, ((unsigned int, count))
+                      ((unsigned int*, g_actorID))
+                      ((unsigned int*, stateToInput))
+                      ((const float4*, g_positionMass))
+                      ((float4*, g_outMinBounds))((float4*, g_outMaxBounds))
+                      ((unsigned int*, g_tmpActorID))
+                      ((float4*, g_tmpMinBounds))((float4*, g_tmpMaxBounds))
+                     )
+
+APEX_CUDA_BOUND_KERNEL(BBOX_KERNEL_CONFIG, bboxKernel,
+                       ((unsigned int*, g_actorID))
+                       ((unsigned int*, stateToInput))
+                       ((const float4*, g_positionMass))
+                       ((float4*, g_outMinBounds))((float4*, g_outMaxBounds))
+                       ((unsigned int*, g_tmpActorID))
+                       ((float4*, g_tmpMinBounds))((float4*, g_tmpMaxBounds))
+					   ((unsigned int, phase))((unsigned int, gridSize))
+                     )
diff --git a/APEX_1.4/module/iofx/cuda/include/common.h b/APEX_1.4/module/iofx/cuda/include/common.h
new file mode 100644
index 00000000..6e7d6a8d
--- /dev/null
+++ b/APEX_1.4/module/iofx/cuda/include/common.h
@@ -0,0 +1,156 @@
+/*
+ * 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 IOFX_
+
+#include "ApexCuda.h"
+#include "InplaceTypes.h"
+#include "IofxManagerIntl.h"
+#include "PxBounds3.h"
+#include <float.h>
+
+#if PX_WINDOWS_FAMILY
+#pragma warning(push)
+#pragma warning(disable:4201)
+#pragma warning(disable:4408)
+#endif
+
+#include <vector_types.h>
+
+#if PX_WINDOWS_FAMILY
+#pragma warning(pop)
+#endif
+
+/* Reduce functions assume warps per block <= 32 !! */
+
+
+#define ACTOR_RANGE_KERNEL_CONFIG (2, WARP_SIZE * 2)
+
+const unsigned int BBOX_BLOCK_BOUNDS_PAD = 1;
+const unsigned int BBOX_WARP_BOUNDS_PAD = 1;
+#define BBOX_KERNEL_CONFIG (BBOX_BLOCK_BOUNDS_PAD * 6, WARP_SIZE * 7, 0, 1, WARP_SIZE)
+
+const unsigned int SPRITE_MAX_DWORDS_PER_OUTPUT = 18;
+const unsigned int MESH_MAX_DWORDS_PER_OUTPUT = 22;
+
+#define SPRITE_TEXTURE_MODIFIER_KERNEL_CONFIG (0, WARP_SIZE * 4)
+
+const unsigned int RADIX_SORT_NBITS = 4;
+
+#define SORT_KERNEL_CONFIG (0, WARP_SIZE * 2)
+
+//+1 here is to avoid shared memory bank conflicts
+const unsigned int NEW_SORT_WARP_STRIDE_PAD = 1;
+const unsigned int NEW_SORT_KEY_DIGITS = (1U << RADIX_SORT_NBITS);
+const unsigned int NEW_SORT_VECTOR_SIZE = 4;
+
+/* step 1: (NEW_SORT_KEY_DIGITS * NEW_SORT_WARP_STRIDE_PAD, WARP_SIZE * 2 + NEW_SORT_KEY_DIGITS) */
+/* step 2: (0, WARP_SIZE * 2)
+/* step 3 & block: (1 + NEW_SORT_KEY_DIGITS * NEW_SORT_WARP_STRIDE_PAD, WARP_SIZE * 2 + WARP_SIZE * NEW_SORT_VECTOR_SIZE * 2 + NEW_SORT_KEY_DIGITS) */
+#define NEW_SORT_KERNEL_CONFIG (1 + NEW_SORT_KEY_DIGITS * NEW_SORT_WARP_STRIDE_PAD, WARP_SIZE * 2 + WARP_SIZE * NEW_SORT_VECTOR_SIZE * 2 + NEW_SORT_KEY_DIGITS, 0, (NEW_SORT_KEY_DIGITS * MAX_BOUND_BLOCKS / 4) / WARP_SIZE)
+
+
+const unsigned int STATE_ID_MASK = 0x7FFFFFFFu;
+const unsigned int STATE_ID_DIST_SIGN = 0x80000000u;
+
+
+namespace nvidia
+{
+namespace iofx
+{
+
+#ifdef __CUDACC__
+
+APEX_CUDA_CALLABLE PX_INLINE IofxSlice uint4_to_IofxSlice(uint4 v)
+{
+	IofxSlice ret;
+	ret.x = v.x;
+	ret.y = v.y;
+	ret.z = v.z;
+	ret.w = v.w;
+	return ret;
+}
+
+APEX_CUDA_CALLABLE PX_INLINE uint4 IofxSlice_to_uint4(IofxSlice s)
+{
+	return make_uint4(s.x, s.y, s.z, s.w);
+}
+#endif
+
+
+//struct VolumeParams
+#define INPLACE_TYPE_STRUCT_NAME VolumeParams
+#define INPLACE_TYPE_STRUCT_FIELDS \
+	INPLACE_TYPE_FIELD(PxBounds3,	bounds) \
+	INPLACE_TYPE_FIELD(uint32_t,		priority)
+#include INPLACE_TYPE_BUILD()
+
+typedef InplaceArray<VolumeParams> VolumeParamsArray;
+typedef InplaceArray<uint32_t> ActorClassIDBitmapArray;
+
+typedef InplaceArray<uint32_t> ActorIDRemapArray;
+
+//struct ModifierListElem
+#define INPLACE_TYPE_STRUCT_NAME ModifierListElem
+#define INPLACE_TYPE_STRUCT_FIELDS \
+	INPLACE_TYPE_FIELD(uint32_t,		type) \
+	INPLACE_TYPE_FIELD(InplaceHandleBase,	paramsHandle)
+#include INPLACE_TYPE_BUILD()
+
+typedef InplaceArray<ModifierListElem> ModifierList;
+
+//struct AssetParams
+#define INPLACE_TYPE_STRUCT_NAME AssetParams
+#define INPLACE_TYPE_STRUCT_FIELDS \
+	INPLACE_TYPE_FIELD(ModifierList,	spawnModifierList) \
+	INPLACE_TYPE_FIELD(ModifierList,	continuousModifierList)
+#include INPLACE_TYPE_BUILD()
+
+//struct ClientParams
+#define INPLACE_TYPE_STRUCT_NAME ClientParams
+#define INPLACE_TYPE_STRUCT_FIELDS \
+	INPLACE_TYPE_FIELD(float,				objectScale) \
+	INPLACE_TYPE_FIELD(InplaceHandle<AssetParams>,	assetParamsHandle)
+#include INPLACE_TYPE_BUILD()
+
+typedef InplaceArray< InplaceHandle<ClientParams> > ClientParamsHandleArray;
+
+struct SpritePrivateState;
+
+struct SpritePrivateStateArgs
+{
+	IofxSlice* g_state[1];
+
+#ifdef __CUDACC__
+	static __device__ void read(const SpritePrivateStateArgs& args, SpritePrivateState& state, unsigned int pos);
+	static __device__ void write(SpritePrivateStateArgs& args, const SpritePrivateState& state, unsigned int pos);
+#endif
+};
+
+struct MeshPrivateState;
+
+struct MeshPrivateStateArgs
+{
+	IofxSlice* g_state[3];
+
+#ifdef __CUDACC__
+	static __device__ void read(const MeshPrivateStateArgs& args, MeshPrivateState& state, unsigned int pos);
+	static __device__ void write(MeshPrivateStateArgs& args, const MeshPrivateState& state, unsigned int pos);
+#endif
+};
+
+}
+} // namespace nvidia
+
+#endif
diff --git a/APEX_1.4/module/iofx/cuda/include/migration.h b/APEX_1.4/module/iofx/cuda/include/migration.h
new file mode 100644
index 00000000..3204d79d
--- /dev/null
+++ b/APEX_1.4/module/iofx/cuda/include/migration.h
@@ -0,0 +1,22 @@
+/*
+ * 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_STORAGE(migrationStorage)
+
+
+APEX_CUDA_BOUND_KERNEL((), volumeMigrationKernel,
+                       ((InplaceHandle<VolumeParamsArray>, volumeParamsArrayHandle))
+                       ((InplaceHandle<ActorClassIDBitmapArray>, actorClassIDBitmapArrayHandle))
+                       ((uint32_t, numActorClasses))((uint32_t, numVolumes))((uint32_t, numActorIDValues))
+                       ((IofxActorIDIntl*, actorID))((uint32_t, maxInputID))
+                       ((const float4*, positionMass))
+                       ((uint32_t*, actorStart))((uint32_t*, actorEnd))((uint32_t*, actorVisibleEnd))
+                      )
diff --git a/APEX_1.4/module/iofx/cuda/include/modifier.h b/APEX_1.4/module/iofx/cuda/include/modifier.h
new file mode 100644
index 00000000..fc0127e8
--- /dev/null
+++ b/APEX_1.4/module/iofx/cuda/include/modifier.h
@@ -0,0 +1,70 @@
+/*
+ * 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_STORAGE_SIZE(modifierStorage, MAX_CONST_MEM_SIZE - 32)
+
+
+APEX_CUDA_TEXTURE_1D(texRefPositionMass,         float4)
+APEX_CUDA_TEXTURE_1D(texRefVelocityLife,         float4)
+APEX_CUDA_TEXTURE_1D(texRefCollisionNormalFlags, float4)
+APEX_CUDA_TEXTURE_1D(texRefDensity,              float)
+APEX_CUDA_TEXTURE_1D(texRefActorIDs,             unsigned int)
+APEX_CUDA_TEXTURE_1D(texRefUserData,             unsigned int)
+
+APEX_CUDA_TEXTURE_1D(texRefSpritePrivState0,     uint4)
+
+APEX_CUDA_TEXTURE_1D(texRefMeshPrivState0,       uint4)
+APEX_CUDA_TEXTURE_1D(texRefMeshPrivState1,       uint4)
+APEX_CUDA_TEXTURE_1D(texRefMeshPrivState2,       uint4)
+
+APEX_CUDA_TEXTURE_1D(texRefInStateToInput,       unsigned int)
+APEX_CUDA_TEXTURE_1D(texRefStateSpawnSeed,       unsigned int)
+APEX_CUDA_TEXTURE_1D(texRefStateSpawnScale,      float)
+
+APEX_CUDA_TEXTURE_2D(texRefCurveSamples,         float)
+
+#ifndef SKIP_OUTPUT_SURFACES_DEF
+APEX_CUDA_SURFACE_2D(surfRefOutput0)
+APEX_CUDA_SURFACE_2D(surfRefOutput1)
+APEX_CUDA_SURFACE_2D(surfRefOutput2)
+APEX_CUDA_SURFACE_2D(surfRefOutput3)
+#else
+#undef SKIP_OUTPUT_SURFACES_DEF
+#endif
+
+APEX_CUDA_BOUND_KERNEL((), spriteModifiersKernel,
+                       ((unsigned int, inStateOffset))((unsigned int, outStateOffset))
+                       ((InplaceHandle<ClientParamsHandleArray>, clientParamsHandleArrayHandle))
+                       ((ModifierCommonParams, commonParams))
+                       ((unsigned int*, g_sortedActorIDs))((unsigned int*, g_sortedStateIDs))((unsigned int*, g_outStateToInput))
+                       ((SpritePrivateStateArgs, privStateArgs))((float*, g_stateSpawnScale))
+                       ((PRNGInfo, rand))((unsigned int*, g_outputBuffer))
+					   ((InplaceHandle<SpriteOutputLayout>, outputLayoutHandle))
+                      )
+
+APEX_CUDA_BOUND_KERNEL(SPRITE_TEXTURE_MODIFIER_KERNEL_CONFIG, spriteTextureModifiersKernel,
+                       ((unsigned int, inStateOffset))((unsigned int, outStateOffset))
+                       ((InplaceHandle<ClientParamsHandleArray>, clientParamsHandleArrayHandle))
+                       ((ModifierCommonParams, commonParams))
+                       ((unsigned int*, g_sortedActorIDs))((unsigned int*, g_sortedStateIDs))((unsigned int*, g_outStateToInput))
+                       ((SpritePrivateStateArgs, privStateArgs))((float*, g_stateSpawnScale))
+                       ((PRNGInfo, rand))((SpriteTextureOutputLayout, outputLayout))
+                      )
+
+APEX_CUDA_BOUND_KERNEL((), meshModifiersKernel,
+                       ((unsigned int, inStateOffset))((unsigned int, outStateOffset))
+                       ((InplaceHandle<ClientParamsHandleArray>, clientParamsHandleArrayHandle))
+                       ((ModifierCommonParams, commonParams))
+                       ((unsigned int*, g_sortedActorIDs))((unsigned int*, g_sortedStateIDs))((unsigned int*, g_outStateToInput))
+                       ((MeshPrivateStateArgs, privStateArgs))((float*, g_stateSpawnScale))
+                       ((PRNGInfo, rand))((unsigned int*, g_outputBuffer))
+					   ((InplaceHandle<MeshOutputLayout>, outputLayoutHandle))
+                      )
diff --git a/APEX_1.4/module/iofx/cuda/include/moduleList.h b/APEX_1.4/module/iofx/cuda/include/moduleList.h
new file mode 100644
index 00000000..2bcceed1
--- /dev/null
+++ b/APEX_1.4/module/iofx/cuda/include/moduleList.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.
+ */
+
+
+#include "actorRanges.h"
+#include "bbox.h"
+#include "migration.h"
+#include "modifier.h"
+#include "remap.h"
+#include "sort.h"
+#include "sortNew.h"
diff --git a/APEX_1.4/module/iofx/cuda/include/remap.h b/APEX_1.4/module/iofx/cuda/include/remap.h
new file mode 100644
index 00000000..d57fb441
--- /dev/null
+++ b/APEX_1.4/module/iofx/cuda/include/remap.h
@@ -0,0 +1,34 @@
+/*
+ * 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_STORAGE(remapStorage)
+
+
+APEX_CUDA_TEXTURE_1D(texRefRemapPositions,      float4)
+APEX_CUDA_TEXTURE_1D(texRefRemapActorIDs,       unsigned int)
+APEX_CUDA_TEXTURE_1D(texRefRemapInStateToInput, unsigned int)
+
+
+APEX_CUDA_BOUND_KERNEL((), makeSortKeys,
+                       ((const uint32_t*, inStateToInput))((uint32_t, maxInputID))
+                       ((uint32_t, numActorsPerVolume))((uint32_t, numActorIDs))
+                       ((InplaceHandle<ActorIDRemapArray>, actorIDRemapArrayHandle))
+                       ((const float4*, positionMass))((bool, outputDensityKeys))
+                       ((PxVec3, eyePos))((PxVec3, eyeDir))((float, zNear))
+                       ((uint32_t*, sortKey))((uint32_t*, sortValue))
+                      )
+
+APEX_CUDA_BOUND_KERNEL((), remapKernel,
+                       ((const uint32_t*, inStateToInput))((uint32_t, maxInputID))
+                       ((uint32_t, numActorsPerVolume))((uint32_t, numActorIDs))
+                       ((InplaceHandle<ActorIDRemapArray>, actorIDRemapArrayHandle))
+                       ((const unsigned int*, inSortedValue))((unsigned int*, outSortKey))
+                      )
diff --git a/APEX_1.4/module/iofx/cuda/include/sort.h b/APEX_1.4/module/iofx/cuda/include/sort.h
new file mode 100644
index 00000000..283e2884
--- /dev/null
+++ b/APEX_1.4/module/iofx/cuda/include/sort.h
@@ -0,0 +1,23 @@
+/*
+ * 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_SYNC_KERNEL(SORT_KERNEL_CONFIG, radixSortSyncKernel, ((unsigned int, numElements))
+                      ((unsigned int*, keys))((unsigned int*, values))
+                      ((unsigned int*, tempKeys))((unsigned int*, tempValues))
+                      ((unsigned int*, g_temp))((unsigned int, keyBits))((unsigned int, startbit))
+                     )
+
+APEX_CUDA_BOUND_KERNEL(SORT_KERNEL_CONFIG, radixSortStepKernel,
+                       ((unsigned int*, keys))((unsigned int*, values))
+                       ((unsigned int*, tempKeys))((unsigned int*, tempValues))
+                       ((unsigned int*, g_temp))((unsigned int, startbit))
+					   ((unsigned int, phase))((unsigned int, gridSize))
+                      )
diff --git a/APEX_1.4/module/iofx/cuda/include/sortNew.h b/APEX_1.4/module/iofx/cuda/include/sortNew.h
new file mode 100644
index 00000000..70c56315
--- /dev/null
+++ b/APEX_1.4/module/iofx/cuda/include/sortNew.h
@@ -0,0 +1,23 @@
+/*
+ * 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_BOUND_KERNEL(NEW_SORT_KERNEL_CONFIG, newRadixSortBlockKernel,
+                       ((unsigned int, count))((unsigned int, bitCount))((unsigned int, startBit))
+                       ((unsigned int*, inpKeys))((unsigned int*, inpValues))
+                      )
+
+APEX_CUDA_BOUND_KERNEL(NEW_SORT_KERNEL_CONFIG, newRadixSortStepKernel,
+                       ((unsigned int, count))((unsigned int, startBit))
+                       ((unsigned int*, inpKeys))((unsigned int*, inpValues))
+                       ((unsigned int*, outKeys))((unsigned int*, outValues))
+                       ((unsigned int*, tempScan))
+					   ((unsigned int, phase))((unsigned int, gridSize))
+                      )