diff options
Diffstat (limited to 'PxShared/src/cudamanager/include')
| -rw-r--r-- | PxShared/src/cudamanager/include/CudaContextManager.h | 51 | ||||
| -rw-r--r-- | PxShared/src/cudamanager/include/CudaKernelWrangler.h | 331 | ||||
| -rw-r--r-- | PxShared/src/cudamanager/include/GpuDispatcher.h | 332 | ||||
| -rw-r--r-- | PxShared/src/cudamanager/include/PhysXDeviceSettings.h | 56 |
4 files changed, 0 insertions, 770 deletions
diff --git a/PxShared/src/cudamanager/include/CudaContextManager.h b/PxShared/src/cudamanager/include/CudaContextManager.h deleted file mode 100644 index 3d68f82..0000000 --- a/PxShared/src/cudamanager/include/CudaContextManager.h +++ /dev/null @@ -1,51 +0,0 @@ -// 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-2017 NVIDIA Corporation. All rights reserved. - -#ifndef PXCUDACONTEXTMANAGER_CUDACONTEXTMANAGER_H -#define PXCUDACONTEXTMANAGER_CUDACONTEXTMANAGER_H - -#include "task/PxTaskDefine.h" - -#if PX_SUPPORT_GPU_PHYSX - -namespace physx -{ - -class PxCudaContextManager; -class PxCudaContextManagerDesc; -class PxErrorCallback; - -/** -Creates cuda context manager for PhysX and APEX. -*/ -PxCudaContextManager* createCudaContextManager(const PxCudaContextManagerDesc& desc, PxErrorCallback& errorCallback); - -} - -#endif - -#endif // PXCUDACONTEXTMANAGER_CUDACONTEXTMANAGER_H diff --git a/PxShared/src/cudamanager/include/CudaKernelWrangler.h b/PxShared/src/cudamanager/include/CudaKernelWrangler.h deleted file mode 100644 index 36a2cc8..0000000 --- a/PxShared/src/cudamanager/include/CudaKernelWrangler.h +++ /dev/null @@ -1,331 +0,0 @@ -// 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-2017 NVIDIA Corporation. All rights reserved. - -#ifndef __CUDA_KERNEL_WRANGLER__ -#define __CUDA_KERNEL_WRANGLER__ - -// Make this header is safe for inclusion in headers that are shared with device code. -#if !defined(__CUDACC__) - -#include "task/PxTaskDefine.h" -#include "task/PxGpuDispatcher.h" - -#include "PsUserAllocated.h" -#include "PsArray.h" - -#include <cuda.h> - -namespace physx -{ - -class KernelWrangler : public shdfnd::UserAllocated -{ - PX_NOCOPY(KernelWrangler) -public: - KernelWrangler(PxGpuDispatcher& gd, PxErrorCallback& errorCallback, const char** funcNames, uint16_t numFuncs); - ~KernelWrangler(); - - CUfunction getCuFunction(uint16_t funcIndex) const - { - return mCuFunctions[ funcIndex ]; - } - - CUmodule getCuModule(uint16_t funcIndex) const - { - uint16_t modIndex = mCuFuncModIndex[ funcIndex ]; - return mCuModules[ modIndex ]; - } - - static void const* const* getImages(); - static int getNumImages(); - - bool hadError() const { return mError; } - -protected: - bool mError; - shdfnd::Array<CUfunction> mCuFunctions; - shdfnd::Array<uint16_t> mCuFuncModIndex; - shdfnd::Array<CUmodule> mCuModules; - PxGpuDispatcher& mGpuDispatcher; - PxErrorCallback& mErrorCallback; -}; - -/* SJB - These were "borrowed" from an Ignacio Llamas email to devtech-compute. - * If we feel this is too clumsy, we can steal the boost based bits from APEX - */ - -class ExplicitCudaFlush -{ -public: - ExplicitCudaFlush(int cudaFlushCount) : mCudaFlushCount(cudaFlushCount), mDefaultCudaFlushCount(mCudaFlushCount) {} - ~ExplicitCudaFlush() {} - - void setCudaFlushCount(int value) { mCudaFlushCount = mDefaultCudaFlushCount = value; } - unsigned int getCudaFlushCount() const { return (unsigned int)mCudaFlushCount; } - void resetCudaFlushCount() { mCudaFlushCount = mDefaultCudaFlushCount; } - - void decrementFlushCount() - { - if (mCudaFlushCount == 0) return; - - if (--mCudaFlushCount == 0) - { - CUresult ret = cuStreamQuery(0); // flushes current push buffer - PX_UNUSED(ret); - PX_ASSERT(ret == CUDA_SUCCESS || ret == CUDA_ERROR_NOT_READY); - - // For current implementation, disable resetting of cuda flush count - // reset cuda flush count - // mCudaFlushCount = mDefaultCudaFlushCount; - } - } - -private: - int mCudaFlushCount; - int mDefaultCudaFlushCount; -}; - -} - -template <typename T0> -PX_NOINLINE CUresult launchKernel(CUfunction func, unsigned int numBlocks, unsigned int numThreads, unsigned int sharedMem, CUstream stream, - T0 v0) -{ - void* kernelParams[] = - { - &v0, - }; - return cuLaunchKernel(func, numBlocks, 1, 1, numThreads, 1, 1, sharedMem, stream, kernelParams, NULL); -} - -template <typename T0, typename T1> -PX_NOINLINE CUresult launchKernel(CUfunction func, unsigned int numBlocks, unsigned int numThreads, unsigned int sharedMem, CUstream stream, - T0 v0, T1 v1) -{ - void* kernelParams[] = - { - &v0, &v1, - }; - return cuLaunchKernel(func, numBlocks, 1, 1, numThreads, 1, 1, sharedMem, stream, kernelParams, NULL); -} - -template <typename T0, typename T1, typename T2> -PX_NOINLINE CUresult launchKernel(CUfunction func, unsigned int numBlocks, unsigned int numThreads, unsigned int sharedMem, CUstream stream, - T0 v0, T1 v1, T2 v2) -{ - void* kernelParams[] = - { - &v0, &v1, &v2, - }; - return cuLaunchKernel(func, numBlocks, 1, 1, numThreads, 1, 1, sharedMem, stream, kernelParams, NULL); -} - -template <typename T0, typename T1, typename T2, typename T3> -PX_NOINLINE CUresult launchKernel(CUfunction func, unsigned int numBlocks, unsigned int numThreads, unsigned int sharedMem, CUstream stream, - T0 v0, T1 v1, T2 v2, T3 v3) -{ - void* kernelParams[] = - { - &v0, &v1, &v2, &v3, - }; - return cuLaunchKernel(func, numBlocks, 1, 1, numThreads, 1, 1, sharedMem, stream, kernelParams, NULL); -} - -template <typename T0, typename T1, typename T2, typename T3, typename T4> -PX_NOINLINE CUresult launchKernel(CUfunction func, unsigned int numBlocks, unsigned int numThreads, unsigned int sharedMem, CUstream stream, - T0 v0, T1 v1, T2 v2, T3 v3, T4 v4) -{ - void* kernelParams[] = - { - &v0, &v1, &v2, &v3, &v4, - }; - return cuLaunchKernel(func, numBlocks, 1, 1, numThreads, 1, 1, sharedMem, stream, kernelParams, NULL); -} - -template <typename T0, typename T1, typename T2, typename T3, typename T4, typename T5> -PX_NOINLINE CUresult launchKernel(CUfunction func, unsigned int numBlocks, unsigned int numThreads, unsigned int sharedMem, CUstream stream, - T0 v0, T1 v1, T2 v2, T3 v3, T4 v4, T5 v5) -{ - void* kernelParams[] = - { - &v0, &v1, &v2, &v3, &v4, &v5, - }; - return cuLaunchKernel(func, numBlocks, 1, 1, numThreads, 1, 1, sharedMem, stream, kernelParams, NULL); -} - -template <typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6> -PX_NOINLINE CUresult launchKernel(CUfunction func, unsigned int numBlocks, unsigned int numThreads, unsigned int sharedMem, CUstream stream, - T0 v0, T1 v1, T2 v2, T3 v3, T4 v4, T5 v5, T6 v6) -{ - void* kernelParams[] = - { - &v0, &v1, &v2, &v3, &v4, &v5, &v6, - }; - return cuLaunchKernel(func, numBlocks, 1, 1, numThreads, 1, 1, sharedMem, stream, kernelParams, NULL); -} - -template <typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7> -PX_NOINLINE CUresult launchKernel(CUfunction func, unsigned int numBlocks, unsigned int numThreads, unsigned int sharedMem, CUstream stream, - T0 v0, T1 v1, T2 v2, T3 v3, T4 v4, T5 v5, T6 v6, T7 v7) -{ - void* kernelParams[] = - { - &v0, &v1, &v2, &v3, &v4, &v5, &v6, &v7, - }; - return cuLaunchKernel(func, numBlocks, 1, 1, numThreads, 1, 1, sharedMem, stream, kernelParams, NULL); -} - -template <typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, - typename T8> -PX_NOINLINE CUresult launchKernel(CUfunction func, unsigned int numBlocks, unsigned int numThreads, unsigned int sharedMem, CUstream stream, - T0 v0, T1 v1, T2 v2, T3 v3, T4 v4, T5 v5, T6 v6, T7 v7, T8 v8) -{ - void* kernelParams[] = - { - &v0, &v1, &v2, &v3, &v4, &v5, &v6, &v7, &v8, - }; - return cuLaunchKernel(func, numBlocks, 1, 1, numThreads, 1, 1, sharedMem, stream, kernelParams, NULL); -} - -template <typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, - typename T8, typename T9> -PX_NOINLINE CUresult launchKernel(CUfunction func, unsigned int numBlocks, unsigned int numThreads, unsigned int sharedMem, CUstream stream, - T0 v0, T1 v1, T2 v2, T3 v3, T4 v4, T5 v5, T6 v6, T7 v7, T8 v8, T9 v9) -{ - void* kernelParams[] = - { - &v0, &v1, &v2, &v3, &v4, &v5, &v6, &v7, &v8, &v9, - }; - return cuLaunchKernel(func, numBlocks, 1, 1, numThreads, 1, 1, sharedMem, stream, kernelParams, NULL); -} - -template <typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, - typename T8, typename T9, typename T10> -PX_NOINLINE CUresult launchKernel(CUfunction func, unsigned int numBlocks, unsigned int numThreads, unsigned int sharedMem, CUstream stream, - T0 v0, T1 v1, T2 v2, T3 v3, T4 v4, T5 v5, T6 v6, T7 v7, T8 v8, T9 v9, T10 v10) -{ - void* kernelParams[] = - { - &v0, &v1, &v2, &v3, &v4, &v5, &v6, &v7, &v8, &v9, &v10, - }; - return cuLaunchKernel(func, numBlocks, 1, 1, numThreads, 1, 1, sharedMem, stream, kernelParams, NULL); -} - -template <typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, - typename T8, typename T9, typename T10, typename T11> -PX_NOINLINE CUresult launchKernel(CUfunction func, unsigned int numBlocks, unsigned int numThreads, unsigned int sharedMem, CUstream stream, - T0 v0, T1 v1, T2 v2, T3 v3, T4 v4, T5 v5, T6 v6, T7 v7, T8 v8, T9 v9, T10 v10, T11 v11) -{ - void* kernelParams[] = - { - &v0, &v1, &v2, &v3, &v4, &v5, &v6, &v7, &v8, &v9, &v10, &v11, - }; - return cuLaunchKernel(func, numBlocks, 1, 1, numThreads, 1, 1, sharedMem, stream, kernelParams, NULL); -} - -template <typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, - typename T8, typename T9, typename T10, typename T11, typename T12> -PX_NOINLINE CUresult launchKernel(CUfunction func, unsigned int numBlocks, unsigned int numThreads, unsigned int sharedMem, CUstream stream, - T0 v0, T1 v1, T2 v2, T3 v3, T4 v4, T5 v5, T6 v6, T7 v7, T8 v8, T9 v9, T10 v10, T11 v11, T12 v12) -{ - void* kernelParams[] = - { - &v0, &v1, &v2, &v3, &v4, &v5, &v6, &v7, &v8, &v9, &v10, &v11, &v12, - }; - return cuLaunchKernel(func, numBlocks, 1, 1, numThreads, 1, 1, sharedMem, stream, kernelParams, NULL); -} - -template <typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, - typename T8, typename T9, typename T10, typename T11, typename T12, typename T13> -PX_NOINLINE CUresult launchKernel(CUfunction func, unsigned int numBlocks, unsigned int numThreads, unsigned int sharedMem, CUstream stream, - T0 v0, T1 v1, T2 v2, T3 v3, T4 v4, T5 v5, T6 v6, T7 v7, T8 v8, T9 v9, T10 v10, T11 v11, T12 v12, - T13 v13) -{ - void* kernelParams[] = - { - &v0, &v1, &v2, &v3, &v4, &v5, &v6, &v7, &v8, &v9, &v10, &v11, &v12, &v13, - }; - return cuLaunchKernel(func, numBlocks, 1, 1, numThreads, 1, 1, sharedMem, stream, kernelParams, NULL); -} - -template <typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, - typename T8, typename T9, typename T10, typename T11, typename T12, typename T13, typename T14> -PX_NOINLINE CUresult launchKernel(CUfunction func, unsigned int numBlocks, unsigned int numThreads, unsigned int sharedMem, CUstream stream, - T0 v0, T1 v1, T2 v2, T3 v3, T4 v4, T5 v5, T6 v6, T7 v7, T8 v8, T9 v9, T10 v10, T11 v11, T12 v12, - T13 v13, T14 v14) -{ - void* kernelParams[] = - { - &v0, &v1, &v2, &v3, &v4, &v5, &v6, &v7, &v8, &v9, &v10, &v11, &v12, &v13, &v14, - }; - return cuLaunchKernel(func, numBlocks, 1, 1, numThreads, 1, 1, sharedMem, stream, kernelParams, NULL); -} - -template <typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, - typename T8, typename T9, typename T10, typename T11, typename T12, typename T13, typename T14, typename T15> -PX_NOINLINE CUresult launchKernel(CUfunction func, unsigned int numBlocks, unsigned int numThreads, unsigned int sharedMem, CUstream stream, - T0 v0, T1 v1, T2 v2, T3 v3, T4 v4, T5 v5, T6 v6, T7 v7, T8 v8, T9 v9, T10 v10, T11 v11, T12 v12, - T13 v13, T14 v14, T15 v15) -{ - void* kernelParams[] = - { - &v0, &v1, &v2, &v3, &v4, &v5, &v6, &v7, &v8, &v9, &v10, &v11, &v12, &v13, &v14, &v15, - }; - return cuLaunchKernel(func, numBlocks, 1, 1, numThreads, 1, 1, sharedMem, stream, kernelParams, NULL); -} - -template <typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, - typename T8, typename T9, typename T10, typename T11, typename T12, typename T13, typename T14, typename T15, - typename T16> -PX_NOINLINE CUresult launchKernel(CUfunction func, unsigned int numBlocks, unsigned int numThreads, unsigned int sharedMem, CUstream stream, - T0 v0, T1 v1, T2 v2, T3 v3, T4 v4, T5 v5, T6 v6, T7 v7, T8 v8, T9 v9, T10 v10, T11 v11, T12 v12, - T13 v13, T14 v14, T15 v15, T16 v16) -{ - void* kernelParams[] = - { - &v0, &v1, &v2, &v3, &v4, &v5, &v6, &v7, &v8, &v9, &v10, &v11, &v12, &v13, &v14, &v15, &v16, - }; - return cuLaunchKernel(func, numBlocks, 1, 1, numThreads, 1, 1, sharedMem, stream, kernelParams, NULL); -} - -template <typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, - typename T8, typename T9, typename T10, typename T11, typename T12, typename T13, typename T14, typename T15, - typename T16, typename T17> -PX_NOINLINE CUresult launchKernel(CUfunction func, unsigned int numBlocks, unsigned int numThreads, unsigned int sharedMem, CUstream stream, - T0 v0, T1 v1, T2 v2, T3 v3, T4 v4, T5 v5, T6 v6, T7 v7, T8 v8, T9 v9, T10 v10, T11 v11, T12 v12, - T13 v13, T14 v14, T15 v15, T16 v16, T17 v17) -{ - void* kernelParams[] = - { - &v0, &v1, &v2, &v3, &v4, &v5, &v6, &v7, &v8, &v9, &v10, &v11, &v12, &v13, &v14, &v15, &v16, &v17, - }; - return cuLaunchKernel(func, numBlocks, 1, 1, numThreads, 1, 1, sharedMem, stream, kernelParams, NULL); -} - -#endif - -#endif diff --git a/PxShared/src/cudamanager/include/GpuDispatcher.h b/PxShared/src/cudamanager/include/GpuDispatcher.h deleted file mode 100644 index 10c412f..0000000 --- a/PxShared/src/cudamanager/include/GpuDispatcher.h +++ /dev/null @@ -1,332 +0,0 @@ -// 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-2017 NVIDIA Corporation. All rights reserved. - -#ifndef PXTASK_GPUDISPATCHER_H -#define PXTASK_GPUDISPATCHER_H - -#include "task/PxTask.h" -#include "task/PxTaskDefine.h" -#include "task/PxGpuTask.h" -#include "task/PxTaskManager.h" -#include "task/PxGpuDispatcher.h" -#include "foundation/PxProfiler.h" - -#include "PsUserAllocated.h" -#include "PsThread.h" -#include "PsAtomic.h" -#include "PsMutex.h" -#include "PsSync.h" -#include "PsArray.h" - -#include <cuda.h> - -namespace physx { - -typedef uint16_t EventID; - -void releaseGpuDispatcher(PxGpuDispatcher&); - -class KernelWrangler; -class BlockingWaitThread; -class FanoutTask; -class LaunchTask; -class BlockTask; -class PxGpuWorkerThread; - -class GpuDispatcherImpl : public PxGpuDispatcher, public shdfnd::UserAllocated -{ -public: - GpuDispatcherImpl(PxErrorCallback& errorCallback, PxCudaContextManager& ctx); - virtual ~GpuDispatcherImpl(); - - void start(); - void startSimulation(); - void startGroup(); - void submitTask(PxTask& task); - void finishGroup(); - void addCompletionPrereq(PxBaseTask& task); - bool failureDetected() const; - void forceFailureMode(); - void stopSimulation(); - void launchCopyKernel(PxGpuCopyDesc* desc, uint32_t count, CUstream stream); - - PxBaseTask& getPreLaunchTask(); - void addPreLaunchDependent(PxBaseTask& dependent); - - PxBaseTask& getPostLaunchTask(); - void addPostLaunchDependent(PxBaseTask& dependent); - - PxCudaContextManager* getCudaContextManager(); - - PxGpuWorkerThread* mDispatcher; - BlockingWaitThread* mBlockingThread; - LaunchTask* mLaunchTask; // predecessor of tasks launching kernels - BlockTask* mBlockTask; // continuation of tasks launching kernels - FanoutTask* mSyncTask; // predecessor of tasks waiting for cuda context synchronize -}; - -class JobQueue -{ - PX_NOCOPY(JobQueue) -public: - JobQueue() : taskarray(PX_DEBUG_EXP("PxTask*")) {} - void push(PxTask* t) - { - access.lock(); - taskarray.pushBack(t); - access.unlock(); - } - PxTask* popBack() - { - access.lock(); - PxTask* t = NULL; - if (taskarray.size()) - { - t = taskarray.popBack(); - } - access.unlock(); - return t; - } - uint32_t size() - { - return taskarray.size(); - } - bool empty() - { - return taskarray.size() == 0; - } - -private: - shdfnd::Array<PxTask*> taskarray; - shdfnd::Mutex access; -}; - -class EventPool -{ - PX_NOCOPY(EventPool) -public: - EventPool(uint32_t inflags) : flags(inflags), evarray(PX_DEBUG_EXP("CUevent")) {} - void add(CUevent ev) - { - access.lock(); - evarray.pushBack(ev); - access.unlock(); - } - CUevent get() - { - access.lock(); - CUevent ev; - if (evarray.size()) - { - ev = evarray.popBack(); - } - else - { - cuEventCreate(&ev, flags); - } - access.unlock(); - return ev; - } - bool empty() const - { - return evarray.size() == 0; - } - void clear() - { - access.lock(); - for (uint32_t i = 0; i < evarray.size(); i++) - { - cuEventDestroy(evarray[i]); - } - access.unlock(); - } - -private: - uint32_t flags; - shdfnd::Array<CUevent> evarray; - shdfnd::Mutex access; -}; - -class StreamCache -{ -public: - StreamCache() : sarray(PX_DEBUG_EXP("CUstream")), freeIndices(PX_DEBUG_EXP("freeIndices")) - { - } - CUstream get(uint32_t s) - { - PX_ASSERT(s); - return sarray[ s - 1 ]; - } - void push(uint32_t s) - { - freeIndices.pushBack(s); - } - uint32_t popBack() - { - if (freeIndices.size()) - { - return freeIndices.popBack(); - } - else - { - CUstream s; - cuStreamCreate(&s, 0); - sarray.pushBack(s); - return sarray.size(); - } - } - void reset() - { - freeIndices.resize(sarray.size()); - for (uint32_t i = 0 ; i < sarray.size() ; i++) - { - freeIndices[i] = i + 1; - } - } - bool empty() - { - return freeIndices.size() == 0; - } - -private: - shdfnd::Array<CUstream> sarray; - shdfnd::Array<uint32_t> freeIndices; -}; - -class KernelBar -{ -public: - KernelBar() - { - reset(); - } - void reset() - { - start = 0xffffffff; - stop = 0; - } - - uint32_t start; - uint32_t stop; -}; - -const int SIZE_COMPLETION_RING = 1024; - -struct CudaBatch -{ - CUevent blockingEvent; - CUstream blockingStream; // sync on stream instead of event if lsb is zero (faster) - PxBaseTask* continuationTask; -}; - -struct ReadyTask -{ - PxGpuTask* task; - uint32_t iteration; -}; - -class PxGpuWorkerThread : public shdfnd::Thread -{ - PX_NOCOPY(PxGpuWorkerThread) -public: - PxGpuWorkerThread(); - ~PxGpuWorkerThread(); - - void setCudaContext(PxCudaContextManager& ctx); - - /* API to TaskManager */ - void startSimulation(); - void stopSimulation(); - - /* API to GPU tasks */ - void addCompletionPrereq(PxBaseTask& task); - - /* PxGpuTask execution thread */ - void execute(); - void pollSubmitted(shdfnd::Array<ReadyTask> *ready); - void processActiveTasks(); - void flushBatch(CUevent endEvent, CUstream, PxBaseTask* task); - void launchCopyKernel(PxGpuCopyDesc* desc, uint32_t count, CUstream stream); - - /* Blocking wait thread */ - void blockingWaitFunc(); - - StreamCache mCachedStreams; - shdfnd::Array<PxBaseTask*> mCompletionTasks; - JobQueue mSubmittedTaskList; - volatile int mActiveGroups; - shdfnd::Sync mInputReady; - shdfnd::Sync mRecordEventQueued; - PxCudaContextManager* mCtxMgr; - bool mNewTasksSubmitted; - bool mFailureDetected; - - bool mUsingConcurrentStreams; - - CudaBatch mCompletionRing[ SIZE_COMPLETION_RING ]; - volatile int mCompletionRingPush; - volatile int mCompletionRingPop; - - EventPool mCachedBlockingEvents; - EventPool mCachedNonBlockingEvents; - - volatile int mCountActiveScenes; - - uint32_t* mSmStartTimes; - uint32_t mSmClockFreq; - - shdfnd::Array<ReadyTask> mReady[ PxGpuTaskHint::NUM_GPU_TASK_HINTS ]; - - KernelWrangler* mUtilKernelWrapper; - - CUevent mStartEvent; - - shdfnd::Mutex mMutex; -}; - -class BlockingWaitThread : public shdfnd::Thread -{ -public: - BlockingWaitThread(PxGpuWorkerThread& worker) : mWorker(worker) {} - ~BlockingWaitThread() {} - - void execute(); - -protected: - PxGpuWorkerThread& mWorker; - -private: - BlockingWaitThread& operator=(const BlockingWaitThread&); -}; - -#define GD_CHECK_CALL(call) { CUresult ret = call; \ - if( CUDA_SUCCESS != ret ) { mFailureDetected=true; PX_ASSERT(!ret); } } - -} - -#endif // PXTASK_GPUDISPATCHER_H diff --git a/PxShared/src/cudamanager/include/PhysXDeviceSettings.h b/PxShared/src/cudamanager/include/PhysXDeviceSettings.h deleted file mode 100644 index 5358915..0000000 --- a/PxShared/src/cudamanager/include/PhysXDeviceSettings.h +++ /dev/null @@ -1,56 +0,0 @@ -// 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-2017 NVIDIA Corporation. All rights reserved. - -#ifndef PXCUDACONTEXTMANAGER_PHYSXDEVICESETTINGS_H -#define PXCUDACONTEXTMANAGER_PHYSXDEVICESETTINGS_H - -#include "task/PxTaskDefine.h" - -#if PX_SUPPORT_GPU_PHYSX - -namespace physx -{ - class PxErrorCallback; - - /** - Helper functions to expose control panel functionality - */ - class PhysXDeviceSettings - { - private: - PhysXDeviceSettings() {} - - public: - static int getSuggestedCudaDeviceOrdinal(PxErrorCallback& errc); - static int isUsingDedicatedGPU(); - static bool isSLIEnabled(void* graphicsDevice); - }; -} - -#endif - -#endif // PXCUDACONTEXTMANAGER_PHYSXDEVICESETTINGS_H |