Writing documentation on DevicePointer.

Documentation/Tutorials/Pointers/CMakeLists.txt

@@ -3,6 +3,8 @@ IF( BUILD_CUDA )
    ADD_CUSTOM_COMMAND( COMMAND UniquePointerExample > ${TNL_DOCUMENTATION_OUTPUT_SNIPPETS_PATH}/UniquePointerExample.out OUTPUT UniquePointerExample.out )
    CUDA_ADD_EXECUTABLE( SharedPointerExample SharedPointerExample.cu )
    ADD_CUSTOM_COMMAND( COMMAND SharedPointerExample > ${TNL_DOCUMENTATION_OUTPUT_SNIPPETS_PATH}/SharedPointerExample.out OUTPUT SharedPointerExample.out )
+   CUDA_ADD_EXECUTABLE( DevicePointerExample DevicePointerExample.cu )
+   ADD_CUSTOM_COMMAND( COMMAND DevicePointerExample > ${TNL_DOCUMENTATION_OUTPUT_SNIPPETS_PATH}/DevicePointerExample.out OUTPUT DevicePointerExample.out )
 ELSE()
    ADD_EXECUTABLE( UniquePointerExample UniquePointerExample.cpp )
    ADD_CUSTOM_COMMAND( COMMAND UniquePointerExample > ${TNL_DOCUMENTATION_OUTPUT_SNIPPETS_PATH}/UniquePointerExample.out OUTPUT UniquePointerExample.out )
@@ -15,7 +17,8 @@ ADD_CUSTOM_COMMAND( COMMAND UniquePointerHostExample > ${TNL_DOCUMENTATION_OUTPU
 IF( BUILD_CUDA )
 ADD_CUSTOM_TARGET( TutorialsPointersCuda ALL DEPENDS
    UniquePointerExample.out
-   SharedPointerExample.out )
+   SharedPointerExample.out
+   DevicePointerExample.out )
 ENDIF()
 
 ADD_CUSTOM_TARGET( TutorialsPointers ALL DEPENDS

Documentation/Tutorials/Pointers/DevicePointerExample.cpp

+#include <iostream>
+#include <cstdlib>
+#include <TNL/Containers/Array.h>
+#include <TNL/Pointers/DevicePointer.h>
+
+using namespace TNL;
+
+using ArrayCuda = Containers::Array< int, Devices::Cuda >;
+
+struct Tuple
+{
+   Tuple( ArrayCuda& _a1, ArrayCuda& _a2 ):
+   a1( _a1 ), a2( _a2 ){};
+
+   Pointers::DevicePointer< ArrayCuda > a1, a2;
+};
+
+__global__ void printTuple( const Tuple t )
+{
+   printf( "Tuple size is: %d\n", t.a1->getSize() );
+   for( int i = 0; i < t.a1->getSize(); i++ )
+   {
+      printf( "a1[ %d ] = %d \n", i, ( *t.a1 )[ i ] );
+      printf( "a2[ %d ] = %d \n", i, ( *t.a2 )[ i ] );
+   }
+}
+
+int main( int argc, char* argv[] )
+{
+   /***
+    * Create a tuple of arrays and print them in CUDA kernel
+    */
+#ifdef HAVE_CUDA
+   ArrayCuda a1( 3 ), a2( 3 );
+   Tuple t( a1, a2 );
+   a1 = 1;
+   a2 = 2;
+   Pointers::synchronizeSmartPointersOnDevice< Devices::Cuda >();
+   printTuple<<< 1, 1 >>>( t );
+
+   /***
+    * Resize the arrays
+    */
+   a1.setSize( 5 );
+   a2.setSize( 5 );
+   a1 = 3;
+   a2 = 4;
+   Pointers::synchronizeSmartPointersOnDevice< Devices::Cuda >();
+   printTuple<<< 1, 1 >>>( t );
+#endif
+   return EXIT_SUCCESS;
+
+}
+

Documentation/Tutorials/Pointers/DevicePointerExample.cu

+DevicePointerExample.cpp
\ No newline at end of file

Documentation/Tutorials/Pointers/SharedPointerExample.cpp

@@ -34,7 +34,7 @@ __global__ void printTuple( const Tuple t )
 int main( int argc, char* argv[] )
 {
    /***
-    * Create a tuple of arrays and print the in CUDA kernel
+    * Create a tuple of arrays and print them in CUDA kernel
     */
 #ifdef HAVE_CUDA
    Tuple t( 3 );
@@ -44,7 +44,7 @@ int main( int argc, char* argv[] )
    printTuple<<< 1, 1 >>>( t );
 
    /***
-    * Resize the array
+    * Resize the arrays
     */
    t.setSize( 5 );
    *t.a1 = 3;

Documentation/Tutorials/Pointers/tutorial_Pointers.md

@@ -77,4 +77,10 @@ One of the differences between `UniquePointer` and `SmartPointer` is that the `S
 
 ## Device pointers <a name="device_pointers"></a>
 
-The last type of the smart pointer implemented in TNL is `DevicePointer`. It works the same way as `SharedPointer` but it does not create new object on the host system. `DevicePointer` is therefore useful in situation when there is already an object created in the host memory and we want to create its image even on the device. Both images are linked one with each other and so one can just manipulate the one on the host and then synchronize it on the device.
+The last type of the smart pointer implemented in TNL is `DevicePointer`. It works the same way as `SharedPointer` but it does not create new object on the host system. `DevicePointer` is therefore useful in situation when there is already an object created in the host memory and we want to create its image even on the device. Both images are linked one with each other and so one can just manipulate the one on the host and then synchronize it on the device. The following listing is a modification of the previous example with tuple:
+
+\include DevicePointerExample.cpp
+
+The result looks the same:
+
+\include DevicePointerExample.out

src/TNL/Pointers/DevicePointer.h

@@ -267,26 +267,46 @@ class DevicePointer< Object, Devices::Cuda > : public SmartPointer
          pointer.cuda_pointer = nullptr;
       }
 
+      __cuda_callable__
       const Object* operator->() const
       {
+#ifdef __CUDA_ARCH__
+         return this->cuda_pointer;
+#else
          return this->pointer;
+#endif
       }
 
+      __cuda_callable__
       Object* operator->()
       {
+#ifdef __CUDA_ARCH__
+         return this->cuda_pointer;
+#else
          this->pd->maybe_modified = true;
          return this->pointer;
+#endif
       }
 
+      __cuda_callable__
       const Object& operator *() const
       {
+#ifdef __CUDA_ARCH__
+         return *( this->cuda_pointer );
+#else
          return *( this->pointer );
+#endif
       }
 
+      __cuda_callable__
       Object& operator *()
       {
+#ifdef __CUDA_ARCH__
+         return *( this->cuda_pointer );
+#else
          this->pd->maybe_modified = true;
          return *( this->pointer );
+#endif
       }
 
       __cuda_callable__