Added host_write_only() and device_write_only() to the gpu_data and tensor

objects.  Then made the relevant parts of the code use these functions.

dlib/dnn/cudnn_dlibapi.cpp

@@ -240,7 +240,7 @@ namespace dlib
                 return;
             CHECK_CUDNN(cudnnSetTensor(context(),
                                  descriptor(t),
-                                 t.device(),
+                                 t.device_write_only(),
                                  &value));
         }
 

dlib/dnn/gpu_data.h

@@ -95,6 +95,13 @@ namespace dlib
             return data_host.get(); 
         }
 
+        float* host_write_only() 
+        {
+            host_current = true;
+            device_current = false;
+            return data_host.get(); 
+        }
+
         const float* device() const 
         { 
 #ifndef DLIB_USE_CUDA
@@ -116,6 +123,18 @@ namespace dlib
             return data_device.get(); 
         }
 
+        float* device_write_only()
+        {
+#ifndef DLIB_USE_CUDA
+            DLIB_CASSERT(false, "CUDA NOT ENABLED");
+#endif
+            wait_for_transfer_to_finish();
+            host_current = false;
+            device_current = true;
+            device_in_use = true;
+            return data_device.get(); 
+        }
+
         bool host_ready (
         ) const { return host_current; }
 

dlib/dnn/gpu_data_abstract.h

@@ -125,6 +125,20 @@ namespace dlib
                   calling host().
         !*/
 
+        float* host_write_only(
+        );
+        /*!
+            ensures
+                - This function returns the same pointer as host(), except that it never
+                  performs a device to host memory copy.  Instead, it immediately marks the
+                  device side data as out of date, effectively discarding it.  Therefore,
+                  the values in the data pointed to by host_write_only() are undefined and
+                  you should only call host_write_only() if you are going to assign to
+                  every memory location in the returned memory block.  
+                - #host_ready() == true
+                - #device_ready() == false 
+        !*/
+
         const float* device(
         ) const;
         /*!
@@ -154,6 +168,23 @@ namespace dlib
                 - #device_ready() == true
         !*/
 
+        float* device_write_only(
+        );
+        /*!
+            requires
+                - DLIB_USE_CUDA is #defined
+            ensures
+                - This function returns the same pointer as device(), except that it never
+                  performs a host to device memory copy.  Instead, it immediately marks the
+                  host side data as out of date, effectively discarding it.  Therefore, the
+                  values in the data pointed to by device_write_only() are undefined and
+                  you should only call device_write_only() if you are going to assign to
+                  every memory location in the returned memory block.  
+                - #host_ready() == false 
+                - #device_ready() == true 
+        !*/
+
+
         size_t size(
         ) const; 
         /*!

dlib/dnn/loss.h

@@ -68,7 +68,7 @@ namespace dlib
             const double scale = 1.0/output_tensor.num_samples();
             double loss = 0;
             const float* out_data = output_tensor.host();
-            float* g = grad.host();
+            float* g = grad.host_write_only();
             for (long i = 0; i < output_tensor.num_samples(); ++i)
             {
                 const float y = *truth++;
@@ -79,6 +79,10 @@ namespace dlib
                     loss += scale*temp;
                     g[i] = -scale*y;
                 }
+                else
+                {
+                    g[i] = 0;
+                }
             }
             return loss;
         }

dlib/dnn/tensor.h

@@ -46,8 +46,10 @@ namespace dlib
 
         virtual const float* host() const = 0;
         virtual float*       host() = 0; 
+        virtual float*       host_write_only() = 0;
         virtual const float* device() const = 0;
         virtual float*       device() = 0;
+        virtual float*       device_write_only() = 0;
 
         tensor& operator= (float val)
         {
@@ -62,8 +64,9 @@ namespace dlib
                 return *this;
             }
 #endif
-            for (auto& d : *this)
-                d = val;
+            auto d = host_write_only();
+            for (size_t i = 0; i < size(); ++i)
+                d[i] = val;
 
             return *this;
         }
@@ -95,7 +98,7 @@ namespace dlib
             static_assert((is_same_type<float, typename EXP::type>::value == true),
                 "To assign a matrix to a tensor the matrix must contain float values");
 
-            set_ptrm(host(), m_n, m_nr*m_nc*m_k) = item;
+            set_ptrm(host_write_only(), m_n, m_nr*m_nc*m_k) = item;
             return *this;
         }
 
@@ -279,8 +282,10 @@ namespace dlib
 
         virtual const float* host() const { return data_instance.host(); }
         virtual float*       host()       { return data_instance.host(); }
+        virtual float*       host_write_only() { return data_instance.host_write_only(); }
         virtual const float* device() const { return data_instance.device(); }
         virtual float*       device()       { return data_instance.device(); }
+        virtual float*       device_write_only() { return data_instance.device_write_only(); }
 
         void clear(
         )
@@ -441,8 +446,10 @@ namespace dlib
 
         virtual const float* host() const { return data_instance->host()+data_offset; }
         virtual float*       host()       { return data_instance->host()+data_offset; }
+        virtual float*       host_write_only()    { return data_instance->host()+data_offset; }
         virtual const float* device() const { return data_instance->device()+data_offset; }
         virtual float*       device()       { return data_instance->device()+data_offset; }
+        virtual float*       device_write_only()  { return data_instance->device()+data_offset; }
 
 
 #ifdef DLIB_USE_CUDA

dlib/dnn/tensor_abstract.h

@@ -133,6 +133,18 @@ namespace dlib
                   calling host().
         !*/
 
+        float float* host_write_only(
+        ) = 0;
+        /*!
+            ensures
+                - This function returns the same pointer as host(), except that it never
+                  performs a device to host memory copy.  Instead, it immediately marks the
+                  device side data as out of date, effectively discarding it.  Therefore,
+                  the values in the data pointed to by host_write_only() are undefined and
+                  you should only call host_write_only() if you are going to assign to
+                  every memory location in the returned memory block.  
+        !*/
+
         virtual const float* device(
         ) const = 0;
         /*!
@@ -161,6 +173,20 @@ namespace dlib
                   host() will perform a device to host transfer.
         !*/
 
+        float float* device_write_only(
+        ) = 0;
+        /*!
+            requires
+                - DLIB_USE_CUDA is #defined
+            ensures
+                - This function returns the same pointer as device(), except that it never
+                  performs a host to device memory copy.  Instead, it immediately marks the
+                  host side data as out of date, effectively discarding it.  Therefore, the
+                  values in the data pointed to by device_write_only() are undefined and
+                  you should only call device_write_only() if you are going to assign to
+                  every memory location in the returned memory block.  
+        !*/
+
         tensor& operator= (
             float val
         );