diff --git a/src/core/tnlCuda.cu b/src/core/tnlCuda.cu index 113cba35227a0c461f1a32dc1148cd1e9646ac0f..158e29f4fa24b699b3825fe178b9726b23eeca86 100644 --- a/src/core/tnlCuda.cu +++ b/src/core/tnlCuda.cu @@ -41,6 +41,15 @@ bool tnlCuda::setup( const tnlParameterContainer& parameters, } */ +int tnlCuda::getDeviceId() +{ + int id( 0 ); +#ifdef HAVE_CUDA + cudaGetDevice( &id ); +#endif + return id; +} + bool tnlCuda::checkDevice( const char* file_name, int line ) { cudaError error = cudaGetLastError(); diff --git a/src/core/tnlCuda.h b/src/core/tnlCuda.h index 3d5e6eb6ebfd391e634a2bb9ad2979d425b5353a..4de1a2d8acb9edc2144a788f08e6026d43642fc4 100644 --- a/src/core/tnlCuda.h +++ b/src/core/tnlCuda.h @@ -42,7 +42,7 @@ class tnlCuda static tnlString getDeviceType(); - __cuda_callable__ static inline tnlDeviceEnum getDevice(); + __cuda_callable__ static inline tnlDeviceEnum getDevice(); __cuda_callable__ static inline int getMaxGridSize(); @@ -51,6 +51,8 @@ class tnlCuda __cuda_callable__ static inline int getWarpSize(); #ifdef HAVE_CUDA + static int getDeviceId(); + template< typename Index > __device__ static Index getGlobalThreadIdx( const Index gridIdx = 0 ); #endif diff --git a/src/core/tnlDeviceObject.h b/src/core/tnlDeviceObject.h new file mode 100644 index 0000000000000000000000000000000000000000..68b050f43c989e705e00bd92742d3e4786bf5f6f --- /dev/null +++ b/src/core/tnlDeviceObject.h @@ -0,0 +1,137 @@ +/*************************************************************************** + * * + * This program is free software; you can redistribute it and/or modify * + * it under the terms of the GNU General Public License as published by * + * the Free Software Foundation; either version 2 of the License, or * + * (at your option) any later version. * + * * + ***************************************************************************/ + +/*************************************************************************** + tnlDeviceObject.h - description + ------------------- + begin : Apr 29, 2016 + copyright : (C) 2016 by Tomas Oberhuber + email : tomas.oberhuber@fjfi.cvut.cz + ***************************************************************************/ + +#pragma once + +#include <core/tnlHost.h> +#include <core/tnlCuda.h> +#include <core/tnlDeviceObjectBase.h> + +template< typename Object, + typename Device > +class tnlDeviceObject +{ +}; + +template< typename Object > +class tnlDeviceObject< Object, tnlHost > : public tnlDeviceObjectBase +{ + public: + + tnlDeviceObject( Object& object ) + { + this->pointer = &object; + } + + Object* operator->() + { + return this->pointer; + } + + const Object* operator->() const + { + return this->pointer; + } + + const Object& get() const + { + return *this->pointer; + } + + Object& modify() + { + return *this->pointer; + } + + Object* getDevicePointer() + { + return this->pointer; + } + + const Object* getDevicePointer() const + { + return this->pointer; + } + + bool synchronize() + { + return true; + } + + protected: + + Object* pointer; +}; + +template< typename Object > +class tnlDeviceObject< Object, tnlCuda > : public tnlDeviceObjectBase +{ + public: + + tnlDeviceObject( Object& object ) + { + this->host_pointer = &object; +#ifdef HAVE_CUDA + cudaMalloc( ( void** ) &this->device_pointer, sizeof( Object ) ); + deviceId = tnlCuda::getDeviceId(); + tnlCuda::getDeviceObjectsContainer().enregister( this ); +#endif + } + + + Object* operator->() + { + return host_pointer; + } + + const Object* operator->() const + { + return host_pointer; + } + + const Object& get() const + { + return *host_pointer; + } + + Object& modify() + { + return *host_pointer; + } + + Object* getDevicePointer() + { + return device_pointer; + } + + const Object* getDevicePointer() const + { + return device_pointer; + } + + bool synchronize() + { + + } + + protected: + + Object *host_pointer, *device_pointer; + + int deviceId; +}; + diff --git a/src/core/tnlDeviceObjectBase.h b/src/core/tnlDeviceObjectBase.h new file mode 100644 index 0000000000000000000000000000000000000000..468bfaac8365883ae3465636d70af1ae63e66371 --- /dev/null +++ b/src/core/tnlDeviceObjectBase.h @@ -0,0 +1,27 @@ +/*************************************************************************** + * * + * This program is free software; you can redistribute it and/or modify * + * it under the terms of the GNU General Public License as published by * + * the Free Software Foundation; either version 2 of the License, or * + * (at your option) any later version. * + * * + ***************************************************************************/ + +/*************************************************************************** + tnlDeviceObjectBase.h - description + ------------------- + begin : Apr 29, 2016 + copyright : (C) 2016 by Tomas Oberhuber + email : tomas.oberhuber@fjfi.cvut.cz + ***************************************************************************/ + +#pragma once + +class tnlDeviceObjectBase +{ + public: + + virtual bool synchronize() = 0; +}; + + diff --git a/src/core/tnlDeviceObjectsContainer.h b/src/core/tnlDeviceObjectsContainer.h new file mode 100644 index 0000000000000000000000000000000000000000..4d4defb51b086a690b5bb3e4b0b479fe44b319e3 --- /dev/null +++ b/src/core/tnlDeviceObjectsContainer.h @@ -0,0 +1,57 @@ +/*************************************************************************** + * * + * This program is free software; you can redistribute it and/or modify * + * it under the terms of the GNU General Public License as published by * + * the Free Software Foundation; either version 2 of the License, or * + * (at your option) any later version. * + * * + ***************************************************************************/ + +/*************************************************************************** + tnlDeviceObjectsContainer.h - description + ------------------- + begin : Apr 29, 2016 + copyright : (C) 2016 by Tomas Oberhuber + email : tomas.oberhuber@fjfi.cvut.cz + ***************************************************************************/ + +#pragma once + +#include <vector> +#include <list> +#include <core/tnlDeviceObjectBase.h> +#include <core/tnlAssert.h> + +class tnlDeviceObjectsContainer +{ + + public: + + tnlDeviceObjectsContainer( int devicesCount = 1 ) + { + tnlAssert( deviceCount > 0, std::cerr << "deviceCount = " << deviceCount ); + objectsOnDevices.resize( devicesCount ); + this->devicesCount = devicesCount; + } + + void push( tnlDeviceObjectBase* object, int deviceId ) + { + tnlAssert( deviceId >= 0 && deviceId < this->devicesCount, + std::cerr << "deviceId = " << deviceId << " devicesCount = " << this->devicesCount ); + objectsOnDevices[ deviceId ].push( object ); + } + + bool synchronize() + { + .... + } + + protected: + + typedef std::list< tnlDeviceObjectBase* > ListType; + + std::vector< ListType > objectsOnDevices; + + int devicesCount; +}; +