diff --git a/include/caffe/blob.hpp b/include/caffe/blob.hpp
index c04375a10e2..bbea86aea69 100644
--- a/include/caffe/blob.hpp
+++ b/include/caffe/blob.hpp
@@ -6,6 +6,7 @@
 #include "caffe/common.hpp"
 #include "caffe/syncedmem.hpp"
 #include "caffe/proto/caffe.pb.h"
+#include "caffe/util/math_functions.hpp"
 
 namespace caffe {
 
diff --git a/include/caffe/syncedmem.hpp b/include/caffe/syncedmem.hpp
index bed55c3806e..2b7f349025a 100644
--- a/include/caffe/syncedmem.hpp
+++ b/include/caffe/syncedmem.hpp
@@ -6,6 +6,7 @@
 #include <cstdlib>
 
 #include "caffe/common.hpp"
+#include "caffe/util/math_functions.hpp"
 
 namespace caffe {
 
diff --git a/include/caffe/util/math_functions.hpp b/include/caffe/util/math_functions.hpp
index 99519974243..97a057103db 100644
--- a/include/caffe/util/math_functions.hpp
+++ b/include/caffe/util/math_functions.hpp
@@ -59,15 +59,14 @@ void caffe_gpu_axpby(const int N, const Dtype alpha, const Dtype* X,
 template <typename Dtype>
 void caffe_copy(const int N, const Dtype *X, Dtype *Y);
 
+void caffe_memcpy(const size_t N, const void *X, void *Y);
+
 template <typename Dtype>
 void caffe_set(const int N, const Dtype alpha, Dtype *X);
 
 template <typename Dtype>
 void caffe_gpu_set(const int N, const Dtype alpha, Dtype *X);
 
-template <typename Dtype>
-void caffe_gpu_copy(const int N, const Dtype *X, Dtype *Y);
-
 template <typename Dtype>
 void caffe_add_scalar(const int N, const Dtype alpha, Dtype *X);
 
diff --git a/matlab/caffe/matcaffe.cpp b/matlab/caffe/matcaffe.cpp
index 21f51e83994..957ebea008e 100644
--- a/matlab/caffe/matcaffe.cpp
+++ b/matlab/caffe/matcaffe.cpp
@@ -54,12 +54,12 @@ static mxArray* do_forward(const mxArray* const bottom) {
         reinterpret_cast<const float* const>(mxGetPr(elem));
     switch (Caffe::mode()) {
     case Caffe::CPU:
-      memcpy(input_blobs[i]->mutable_cpu_data(), data_ptr,
-          sizeof(float) * input_blobs[i]->count());
+      caffe_copy(input_blobs[i]->count(), data_ptr,
+          input_blobs[i]->mutable_cpu_data());
       break;
     case Caffe::GPU:
-      cudaMemcpy(input_blobs[i]->mutable_gpu_data(), data_ptr,
-          sizeof(float) * input_blobs[i]->count(), cudaMemcpyHostToDevice);
+      caffe_copy(input_blobs[i]->count(), data_ptr,
+          input_blobs[i]->mutable_gpu_data());
       break;
     default:
       LOG(FATAL) << "Unknown Caffe mode.";
@@ -77,12 +77,12 @@ static mxArray* do_forward(const mxArray* const bottom) {
     float* data_ptr = reinterpret_cast<float*>(mxGetPr(mx_blob));
     switch (Caffe::mode()) {
     case Caffe::CPU:
-      memcpy(data_ptr, output_blobs[i]->cpu_data(),
-          sizeof(float) * output_blobs[i]->count());
+      caffe_copy(output_blobs[i]->count(), output_blobs[i]->cpu_data(),
+          data_ptr);
       break;
     case Caffe::GPU:
-      cudaMemcpy(data_ptr, output_blobs[i]->gpu_data(),
-          sizeof(float) * output_blobs[i]->count(), cudaMemcpyDeviceToHost);
+      caffe_copy(output_blobs[i]->count(), output_blobs[i]->gpu_data(),
+          data_ptr);
       break;
     default:
       LOG(FATAL) << "Unknown Caffe mode.";
@@ -104,12 +104,12 @@ static mxArray* do_backward(const mxArray* const top_diff) {
         reinterpret_cast<const float* const>(mxGetPr(elem));
     switch (Caffe::mode()) {
     case Caffe::CPU:
-      memcpy(output_blobs[i]->mutable_cpu_diff(), data_ptr,
-        sizeof(float) * output_blobs[i]->count());
+      caffe_copy(output_blobs[i]->count(), data_ptr,
+          output_blobs[i]->mutable_cpu_diff());
       break;
     case Caffe::GPU:
-      cudaMemcpy(output_blobs[i]->mutable_gpu_diff(), data_ptr,
-        sizeof(float) * output_blobs[i]->count(), cudaMemcpyHostToDevice);
+      caffe_copy(output_blobs[i]->count(), data_ptr,
+          output_blobs[i]->mutable_gpu_diff());
       break;
     default:
       LOG(FATAL) << "Unknown Caffe mode.";
@@ -129,12 +129,10 @@ static mxArray* do_backward(const mxArray* const top_diff) {
     float* data_ptr = reinterpret_cast<float*>(mxGetPr(mx_blob));
     switch (Caffe::mode()) {
     case Caffe::CPU:
-      memcpy(data_ptr, input_blobs[i]->cpu_diff(),
-          sizeof(float) * input_blobs[i]->count());
+      caffe_copy(input_blobs[i]->count(), input_blobs[i]->cpu_diff(), data_ptr);
       break;
     case Caffe::GPU:
-      cudaMemcpy(data_ptr, input_blobs[i]->gpu_diff(),
-          sizeof(float) * input_blobs[i]->count(), cudaMemcpyDeviceToHost);
+      caffe_copy(input_blobs[i]->count(), input_blobs[i]->gpu_diff(), data_ptr);
       break;
     default:
       LOG(FATAL) << "Unknown Caffe mode.";
@@ -206,12 +204,12 @@ static mxArray* do_get_weights() {
 
         switch (Caffe::mode()) {
         case Caffe::CPU:
-          memcpy(weights_ptr, layer_blobs[j]->cpu_data(),
-              sizeof(float) * layer_blobs[j]->count());
+          caffe_copy(layer_blobs[j]->count(), layer_blobs[j]->cpu_data(),
+              weights_ptr);
           break;
         case Caffe::GPU:
-          CUDA_CHECK(cudaMemcpy(weights_ptr, layer_blobs[j]->gpu_data(),
-              sizeof(float) * layer_blobs[j]->count(), cudaMemcpyDeviceToHost));
+          caffe_copy(layer_blobs[j]->count(), layer_blobs[j]->gpu_data(),
+              weights_ptr);
           break;
         default:
           LOG(FATAL) << "Unknown caffe mode: " << Caffe::mode();
diff --git a/src/caffe/blob.cpp b/src/caffe/blob.cpp
index e603712fd82..8df46323eee 100644
--- a/src/caffe/blob.cpp
+++ b/src/caffe/blob.cpp
@@ -75,25 +75,25 @@ const Dtype* Blob<Dtype>::gpu_diff() const {
 template <typename Dtype>
 Dtype* Blob<Dtype>::mutable_cpu_data() {
   CHECK(data_);
-  return reinterpret_cast<Dtype*>(data_->mutable_cpu_data());
+  return static_cast<Dtype*>(data_->mutable_cpu_data());
 }
 
 template <typename Dtype>
 Dtype* Blob<Dtype>::mutable_gpu_data() {
   CHECK(data_);
-  return reinterpret_cast<Dtype*>(data_->mutable_gpu_data());
+  return static_cast<Dtype*>(data_->mutable_gpu_data());
 }
 
 template <typename Dtype>
 Dtype* Blob<Dtype>::mutable_cpu_diff() {
   CHECK(diff_);
-  return reinterpret_cast<Dtype*>(diff_->mutable_cpu_data());
+  return static_cast<Dtype*>(diff_->mutable_cpu_data());
 }
 
 template <typename Dtype>
 Dtype* Blob<Dtype>::mutable_gpu_diff() {
   CHECK(diff_);
-  return reinterpret_cast<Dtype*>(diff_->mutable_gpu_data());
+  return static_cast<Dtype*>(diff_->mutable_gpu_data());
 }
 
 template <typename Dtype>
@@ -121,15 +121,15 @@ void Blob<Dtype>::Update() {
   case SyncedMemory::HEAD_AT_CPU:
     // perform computation on CPU
     caffe_axpy<Dtype>(count_, Dtype(-1),
-        reinterpret_cast<const Dtype*>(diff_->cpu_data()),
-        reinterpret_cast<Dtype*>(data_->mutable_cpu_data()));
+        static_cast<const Dtype*>(diff_->cpu_data()),
+        static_cast<Dtype*>(data_->mutable_cpu_data()));
     break;
   case SyncedMemory::HEAD_AT_GPU:
   case SyncedMemory::SYNCED:
     // perform computation on GPU
     caffe_gpu_axpy<Dtype>(count_, Dtype(-1),
-        reinterpret_cast<const Dtype*>(diff_->gpu_data()),
-        reinterpret_cast<Dtype*>(data_->mutable_gpu_data()));
+        static_cast<const Dtype*>(diff_->gpu_data()),
+        static_cast<Dtype*>(data_->mutable_gpu_data()));
     break;
   default:
     LOG(FATAL) << "Syncedmem not initialized.";
@@ -149,20 +149,20 @@ void Blob<Dtype>::CopyFrom(const Blob& source, bool copy_diff, bool reshape) {
   switch (Caffe::mode()) {
   case Caffe::GPU:
     if (copy_diff) {
-      CUDA_CHECK(cudaMemcpy(diff_->mutable_gpu_data(), source.gpu_diff(),
-          sizeof(Dtype) * count_, cudaMemcpyDeviceToDevice));
+      caffe_copy(count_, source.gpu_diff(),
+          static_cast<Dtype*>(diff_->mutable_gpu_data()));
     } else {
-      CUDA_CHECK(cudaMemcpy(data_->mutable_gpu_data(), source.gpu_data(),
-          sizeof(Dtype) * count_, cudaMemcpyDeviceToDevice));
+      caffe_copy(count_, source.gpu_data(),
+          static_cast<Dtype*>(data_->mutable_gpu_data()));
     }
     break;
   case Caffe::CPU:
     if (copy_diff) {
-      memcpy(diff_->mutable_cpu_data(), source.cpu_diff(),
-          sizeof(Dtype) * count_);
+      caffe_copy(count_, source.cpu_diff(),
+          static_cast<Dtype*>(diff_->mutable_cpu_data()));
     } else {
-      memcpy(data_->mutable_cpu_data(), source.cpu_data(),
-        sizeof(Dtype) * count_);
+      caffe_copy(count_, source.cpu_data(),
+          static_cast<Dtype*>(data_->mutable_cpu_data()));
     }
     break;
   default:
diff --git a/src/caffe/layers/concat_layer.cu b/src/caffe/layers/concat_layer.cu
index ca0cf0c1b5b..2643d7441c8 100644
--- a/src/caffe/layers/concat_layer.cu
+++ b/src/caffe/layers/concat_layer.cu
@@ -16,7 +16,7 @@ Dtype ConcatLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
     int offset_num = 0;
     for (int i = 0; i < bottom.size(); ++i) {
       const Dtype* bottom_data = bottom[i]->gpu_data();
-      caffe_gpu_copy(bottom[i]->count(), bottom_data,
+      caffe_copy(bottom[i]->count(), bottom_data,
         top_data + (*top)[0]->offset(offset_num));
       offset_num += bottom[i]->num();
     }
@@ -27,7 +27,7 @@ Dtype ConcatLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
       int num_elem =
         bottom[i]->channels() * bottom[i]->height() * bottom[i]->width();
       for (int n = 0; n < num_; ++n) {
-        caffe_gpu_copy(num_elem, bottom_data+bottom[i]->offset(n),
+        caffe_copy(num_elem, bottom_data+bottom[i]->offset(n),
           top_data + (*top)[0]->offset(n, offset_channel));
       }
       offset_channel += bottom[i]->channels();
@@ -49,7 +49,7 @@ void ConcatLayer<Dtype>::Backward_gpu(const vector<Blob<Dtype>*>& top,
       Blob<Dtype>* blob = (*bottom)[i];
       if (propagate_down[i]) {
         Dtype* bottom_diff = blob->mutable_gpu_diff();
-        caffe_gpu_copy(blob->count(), top_diff + top[0]->offset(offset_num),
+        caffe_copy(blob->count(), top_diff + top[0]->offset(offset_num),
                        bottom_diff);
       }
       offset_num += blob->num();
@@ -62,7 +62,7 @@ void ConcatLayer<Dtype>::Backward_gpu(const vector<Blob<Dtype>*>& top,
         Dtype* bottom_diff = blob->mutable_gpu_diff();
         int num_elem = blob->channels()*blob->height()*blob->width();
         for (int n = 0; n < num_; ++n) {
-          caffe_gpu_copy(num_elem, top_diff + top[0]->offset(n, offset_channel),
+          caffe_copy(num_elem, top_diff + top[0]->offset(n, offset_channel),
                          bottom_diff + blob->offset(n));
         }
       }
diff --git a/src/caffe/layers/conv_layer.cpp b/src/caffe/layers/conv_layer.cpp
index 67913bfa574..963dc688a0e 100644
--- a/src/caffe/layers/conv_layer.cpp
+++ b/src/caffe/layers/conv_layer.cpp
@@ -126,11 +126,11 @@ void ConvolutionLayer<Dtype>::Backward_cpu(const vector<Blob<Dtype>*>& top,
       const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom) {
   const Dtype* weight = this->blobs_[0]->cpu_data();
   Dtype* weight_diff = this->blobs_[0]->mutable_cpu_diff();
-  memset(weight_diff, 0, sizeof(Dtype) * this->blobs_[0]->count());
+  caffe_set(this->blobs_[0]->count(), Dtype(0), weight_diff);
   Dtype* bias_diff = NULL;
   if (bias_term_) {
     bias_diff = this->blobs_[1]->mutable_cpu_diff();
-    memset(bias_diff, 0, sizeof(Dtype) * this->blobs_[1]->count());
+    caffe_set(this->blobs_[1]->count(), Dtype(0), bias_diff);
   }
   const int weight_offset = M_ * K_;
   const int col_offset = K_ * N_;
diff --git a/src/caffe/layers/conv_layer.cu b/src/caffe/layers/conv_layer.cu
index 71b00c9566b..59ec58dfebe 100644
--- a/src/caffe/layers/conv_layer.cu
+++ b/src/caffe/layers/conv_layer.cu
@@ -48,15 +48,13 @@ void ConvolutionLayer<Dtype>::Backward_gpu(const vector<Blob<Dtype>*>& top,
       const vector<bool>& propagate_down, vector<Blob<Dtype>*>* bottom) {
   const Dtype* weight = this->blobs_[0]->gpu_data();
   Dtype* weight_diff = this->blobs_[0]->mutable_gpu_diff();
-  CUDA_CHECK(cudaMemset(weight_diff, 0,
-      sizeof(Dtype) * this->blobs_[0]->count()));
+  caffe_gpu_set(this->blobs_[0]->count(), Dtype(0), weight_diff);
   Dtype* col_data = col_buffer_.mutable_gpu_data();
   Dtype* col_diff = col_buffer_.mutable_gpu_diff();
   Dtype* bias_diff = NULL;
   if (bias_term_) {
     bias_diff = this->blobs_[1]->mutable_gpu_diff();
-    CUDA_CHECK(cudaMemset(bias_diff, 0,
-        sizeof(Dtype) * this->blobs_[1]->count()));
+    caffe_gpu_set(this->blobs_[1]->count(), Dtype(0), bias_diff);
   }
   const int weight_offset = M_ * K_;
   const int col_offset = K_ * N_;
diff --git a/src/caffe/layers/data_layer.cu b/src/caffe/layers/data_layer.cu
index 2ff9a292b3e..40316a16772 100644
--- a/src/caffe/layers/data_layer.cu
+++ b/src/caffe/layers/data_layer.cu
@@ -21,13 +21,11 @@ Dtype DataLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
   // First, join the thread
   JoinPrefetchThread();
   // Copy the data
-  CUDA_CHECK(cudaMemcpy((*top)[0]->mutable_gpu_data(),
-      prefetch_data_->cpu_data(), sizeof(Dtype) * prefetch_data_->count(),
-      cudaMemcpyHostToDevice));
+  caffe_copy(prefetch_data_->count(), prefetch_data_->cpu_data(),
+      (*top)[0]->mutable_gpu_data());
   if (output_labels_) {
-    CUDA_CHECK(cudaMemcpy((*top)[1]->mutable_gpu_data(),
-        prefetch_label_->cpu_data(), sizeof(Dtype) * prefetch_label_->count(),
-        cudaMemcpyHostToDevice));
+    caffe_copy(prefetch_label_->count(), prefetch_label_->cpu_data(),
+        (*top)[1]->mutable_gpu_data());
   }
   // Start a new prefetch thread
   CreatePrefetchThread();
diff --git a/src/caffe/layers/dropout_layer.cu b/src/caffe/layers/dropout_layer.cu
index 225e0919409..c9f3ecd2dd5 100644
--- a/src/caffe/layers/dropout_layer.cu
+++ b/src/caffe/layers/dropout_layer.cu
@@ -40,7 +40,7 @@ Dtype DropoutLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
         count, bottom_data, mask, uint_thres_, scale_, top_data);
     CUDA_POST_KERNEL_CHECK;
   } else {
-    caffe_gpu_copy(count, bottom_data, top_data);
+    caffe_copy(count, bottom_data, top_data);
   }
   return Dtype(0);
 }
@@ -71,7 +71,7 @@ void DropoutLayer<Dtype>::Backward_gpu(const vector<Blob<Dtype>*>& top,
           count, top_diff, mask, uint_thres_, scale_, bottom_diff);
       CUDA_POST_KERNEL_CHECK;
     } else {
-      caffe_gpu_copy(top[0]->count(), top_diff, bottom_diff);
+      caffe_copy(top[0]->count(), top_diff, bottom_diff);
     }
   }
 }
diff --git a/src/caffe/layers/eltwise_layer.cu b/src/caffe/layers/eltwise_layer.cu
index 3860944889c..99c14feace1 100644
--- a/src/caffe/layers/eltwise_layer.cu
+++ b/src/caffe/layers/eltwise_layer.cu
@@ -51,7 +51,7 @@ void EltwiseLayer<Dtype>::Backward_gpu(const vector<Blob<Dtype>*>& top,
         break;
       case EltwiseParameter_EltwiseOp_SUM:
         if (coeffs_[i] == Dtype(1.)) {
-          caffe_gpu_copy(count, top_diff, bottom_diff);
+          caffe_copy(count, top_diff, bottom_diff);
         } else {
           caffe_gpu_scale(count, coeffs_[i], top_diff, bottom_diff);
         }
diff --git a/src/caffe/layers/hdf5_data_layer.cu b/src/caffe/layers/hdf5_data_layer.cu
index b2b09ef7dd1..3c27f37e360 100644
--- a/src/caffe/layers/hdf5_data_layer.cu
+++ b/src/caffe/layers/hdf5_data_layer.cu
@@ -40,16 +40,12 @@ Dtype HDF5DataLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
       }
       current_row_ = 0;
     }
-    CUDA_CHECK(cudaMemcpy(
-            &(*top)[0]->mutable_gpu_data()[i * data_count],
-            &data_blob_.cpu_data()[current_row_ * data_count],
-            sizeof(Dtype) * data_count,
-            cudaMemcpyHostToDevice));
-    CUDA_CHECK(cudaMemcpy(
-            &(*top)[1]->mutable_gpu_data()[i * label_data_count],
-            &label_blob_.cpu_data()[current_row_ * label_data_count],
-            sizeof(Dtype) * label_data_count,
-            cudaMemcpyHostToDevice));
+    caffe_copy(data_count,
+        &data_blob_.cpu_data()[current_row_ * data_count],
+        &(*top)[0]->mutable_gpu_data()[i * data_count]);
+    caffe_copy(label_data_count,
+        &label_blob_.cpu_data()[current_row_ * label_data_count],
+        &(*top)[1]->mutable_gpu_data()[i * label_data_count]);
   }
   return Dtype(0.);
 }
diff --git a/src/caffe/layers/hdf5_output_layer.cpp b/src/caffe/layers/hdf5_output_layer.cpp
index 3a513b9c366..8307ad7c184 100644
--- a/src/caffe/layers/hdf5_output_layer.cpp
+++ b/src/caffe/layers/hdf5_output_layer.cpp
@@ -54,12 +54,10 @@ Dtype HDF5OutputLayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
   const int label_datum_dim = bottom[1]->count() / bottom[1]->num();
 
   for (int i = 0; i < bottom[0]->num(); ++i) {
-    memcpy(&data_blob_.mutable_cpu_data()[i * data_datum_dim],
-           &bottom[0]->cpu_data()[i * data_datum_dim],
-           sizeof(Dtype) * data_datum_dim);
-    memcpy(&label_blob_.mutable_cpu_data()[i * label_datum_dim],
-           &bottom[1]->cpu_data()[i * label_datum_dim],
-           sizeof(Dtype) * label_datum_dim);
+    caffe_copy(data_datum_dim, &bottom[0]->cpu_data()[i * data_datum_dim],
+        &data_blob_.mutable_cpu_data()[i * data_datum_dim]);
+    caffe_copy(label_datum_dim, &bottom[0]->cpu_data()[i * label_datum_dim],
+        &label_blob_.mutable_cpu_data()[i * label_datum_dim]);
   }
   SaveBlobs();
   return Dtype(0.);
diff --git a/src/caffe/layers/hdf5_output_layer.cu b/src/caffe/layers/hdf5_output_layer.cu
index 59505ee6acf..744b8fe1128 100644
--- a/src/caffe/layers/hdf5_output_layer.cu
+++ b/src/caffe/layers/hdf5_output_layer.cu
@@ -27,12 +27,10 @@ Dtype HDF5OutputLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
   const int label_datum_dim = bottom[1]->count() / bottom[1]->num();
 
   for (int i = 0; i < bottom[0]->num(); ++i) {
-    CUDA_CHECK(cudaMemcpy(&data_blob_.mutable_cpu_data()[i * data_datum_dim],
-           &bottom[0]->gpu_data()[i * data_datum_dim],
-           sizeof(Dtype) * data_datum_dim, cudaMemcpyDeviceToHost));
-    CUDA_CHECK(cudaMemcpy(&label_blob_.mutable_cpu_data()[i * label_datum_dim],
-           &bottom[1]->gpu_data()[i * label_datum_dim],
-           sizeof(Dtype) * label_datum_dim, cudaMemcpyDeviceToHost));
+    caffe_copy(data_datum_dim, &bottom[0]->gpu_data()[i * data_datum_dim],
+        &data_blob_.mutable_cpu_data()[i * data_datum_dim]);
+    caffe_copy(label_datum_dim, &bottom[0]->gpu_data()[i * label_datum_dim],
+        &label_blob_.mutable_cpu_data()[i * label_datum_dim]);
   }
   SaveBlobs();
   return Dtype(0.);
diff --git a/src/caffe/layers/image_data_layer.cu b/src/caffe/layers/image_data_layer.cu
index 98047297d80..dd5bdbc2068 100644
--- a/src/caffe/layers/image_data_layer.cu
+++ b/src/caffe/layers/image_data_layer.cu
@@ -27,12 +27,10 @@ Dtype ImageDataLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
   // First, join the thread
   JoinPrefetchThread();
   // Copy the data
-  CUDA_CHECK(cudaMemcpy((*top)[0]->mutable_gpu_data(),
-      prefetch_data_->cpu_data(), sizeof(Dtype) * prefetch_data_->count(),
-      cudaMemcpyHostToDevice));
-  CUDA_CHECK(cudaMemcpy((*top)[1]->mutable_gpu_data(),
-      prefetch_label_->cpu_data(), sizeof(Dtype) * prefetch_label_->count(),
-      cudaMemcpyHostToDevice));
+  caffe_copy(prefetch_data_->count(), prefetch_data_->cpu_data(),
+      (*top)[0]->mutable_gpu_data());
+  caffe_copy(prefetch_label_->count(), prefetch_label_->cpu_data(),
+      (*top)[1]->mutable_gpu_data());
   // Start a new prefetch thread
   CreatePrefetchThread();
   return Dtype(0.);
diff --git a/src/caffe/layers/lrn_layer.cpp b/src/caffe/layers/lrn_layer.cpp
index a86c1d4c59d..2bda0430c15 100644
--- a/src/caffe/layers/lrn_layer.cpp
+++ b/src/caffe/layers/lrn_layer.cpp
@@ -123,7 +123,7 @@ Dtype LRNLayer<Dtype>::CrossChannelForward_cpu(
   }
   Blob<Dtype> padded_square(1, channels_ + size_ - 1, height_, width_);
   Dtype* padded_square_data = padded_square.mutable_cpu_data();
-  memset(padded_square_data, 0, sizeof(Dtype) * padded_square.count());
+  caffe_set(padded_square.count(), Dtype(0), padded_square_data);
   Dtype alpha_over_size = alpha_ / size_;
   // go through the images
   for (int n = 0; n < num_; ++n) {
@@ -201,7 +201,7 @@ void LRNLayer<Dtype>::CrossChannelBackward_cpu(
   Dtype* accum_ratio_data = accum_ratio.mutable_cpu_data();
   // We hack a little bit by using the diff() to store an additional result
   Dtype* accum_ratio_times_bottom = accum_ratio.mutable_cpu_diff();
-  memset(padded_ratio_data, 0, sizeof(Dtype) * padded_ratio.count());
+  caffe_set(padded_ratio.count(), Dtype(0), padded_ratio_data);
   Dtype cache_ratio_value = 2. * alpha_ * beta_ / size_;
 
   caffe_powx<Dtype>(scale_.count(), scale_data, -beta_, bottom_diff);
@@ -220,7 +220,7 @@ void LRNLayer<Dtype>::CrossChannelBackward_cpu(
         scale_data + block_offset,
         padded_ratio_data + padded_ratio.offset(0, inverse_pre_pad));
     // Now, compute the accumulated ratios and the bottom diff
-    memset(accum_ratio_data, 0, sizeof(Dtype) * accum_ratio.count());
+    caffe_set(accum_ratio.count(), Dtype(0), accum_ratio_data);
     for (int c = 0; c < size_ - 1; ++c) {
       caffe_axpy<Dtype>(height_ * width_, 1.,
           padded_ratio_data + padded_ratio.offset(0, c), accum_ratio_data);
diff --git a/src/caffe/layers/multinomial_logistic_loss_layer.cpp b/src/caffe/layers/multinomial_logistic_loss_layer.cpp
index dd5cae44d6d..8687784d6ab 100644
--- a/src/caffe/layers/multinomial_logistic_loss_layer.cpp
+++ b/src/caffe/layers/multinomial_logistic_loss_layer.cpp
@@ -55,7 +55,7 @@ void MultinomialLogisticLossLayer<Dtype>::Backward_cpu(
     Dtype* bottom_diff = (*bottom)[0]->mutable_cpu_diff();
     int num = (*bottom)[0]->num();
     int dim = (*bottom)[0]->count() / (*bottom)[0]->num();
-    memset(bottom_diff, 0, sizeof(Dtype) * (*bottom)[0]->count());
+    caffe_set((*bottom)[0]->count(), Dtype(0), bottom_diff);
     for (int i = 0; i < num; ++i) {
       int label = static_cast<int>(bottom_label[i]);
       Dtype prob = max(bottom_data[i * dim + label], Dtype(kLOG_THRESHOLD));
diff --git a/src/caffe/layers/power_layer.cu b/src/caffe/layers/power_layer.cu
index 6d699636e21..e7f9831eb13 100644
--- a/src/caffe/layers/power_layer.cu
+++ b/src/caffe/layers/power_layer.cu
@@ -23,7 +23,7 @@ Dtype PowerLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
     return Dtype(0);
   }
   const Dtype* bottom_data = bottom[0]->gpu_data();
-  caffe_gpu_copy(count, bottom_data, top_data);
+  caffe_copy(count, bottom_data, top_data);
   if (scale_ != Dtype(1)) {
     caffe_gpu_scal(count, scale_, top_data);
   }
@@ -68,7 +68,7 @@ void PowerLayer<Dtype>::Backward_gpu(const vector<Blob<Dtype>*>& top,
         caffe_gpu_div(count, top_data, bottom_data, bottom_diff);
         caffe_gpu_scal(count, power_, bottom_diff);
       } else {
-        caffe_gpu_copy(count, bottom_data, bottom_diff);
+        caffe_copy(count, bottom_data, bottom_diff);
         if (scale_ != Dtype(1)) {
           caffe_gpu_scal(count, scale_, bottom_diff);
         }
diff --git a/src/caffe/layers/sigmoid_cross_entropy_loss_layer.cu b/src/caffe/layers/sigmoid_cross_entropy_loss_layer.cu
index 8f7275827e2..0c858cdfc19 100644
--- a/src/caffe/layers/sigmoid_cross_entropy_loss_layer.cu
+++ b/src/caffe/layers/sigmoid_cross_entropy_loss_layer.cu
@@ -50,7 +50,7 @@ void SigmoidCrossEntropyLossLayer<Dtype>::Backward_gpu(
     const Dtype* sigmoid_output_data = sigmoid_output_->gpu_data();
     const Dtype* target = (*bottom)[1]->gpu_data();
     Dtype* bottom_diff = (*bottom)[0]->mutable_gpu_diff();
-    caffe_gpu_copy(count, sigmoid_output_data, bottom_diff);
+    caffe_copy(count, sigmoid_output_data, bottom_diff);
     caffe_gpu_axpy(count, Dtype(-1), target, bottom_diff);
     // Scale down gradient
     caffe_gpu_scal(count, Dtype(1) / num, bottom_diff);
diff --git a/src/caffe/layers/softmax_layer.cpp b/src/caffe/layers/softmax_layer.cpp
index 57847d005f6..5d60d9df1ba 100644
--- a/src/caffe/layers/softmax_layer.cpp
+++ b/src/caffe/layers/softmax_layer.cpp
@@ -34,7 +34,7 @@ Dtype SoftmaxLayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
   Dtype* scale_data = scale_.mutable_cpu_data();
   int num = bottom[0]->num();
   int dim = bottom[0]->count() / bottom[0]->num();
-  memcpy(top_data, bottom_data, sizeof(Dtype) * bottom[0]->count());
+  caffe_copy(bottom[0]->count(), bottom_data, top_data);
   // we need to subtract the max to avoid numerical issues, compute the exp,
   // and then normalize.
   for (int i = 0; i < num; ++i) {
@@ -68,7 +68,7 @@ void SoftmaxLayer<Dtype>::Backward_cpu(const vector<Blob<Dtype>*>& top,
   Dtype* scale_data = scale_.mutable_cpu_data();
   int num = top[0]->num();
   int dim = top[0]->count() / top[0]->num();
-  memcpy(bottom_diff, top_diff, sizeof(Dtype) * top[0]->count());
+  caffe_copy(top[0]->count(), top_diff, bottom_diff);
   // Compute inner1d(top_diff, top_data) and subtract them from the bottom diff
   for (int i = 0; i < num; ++i) {
     scale_data[i] = caffe_cpu_dot<Dtype>(dim, top_diff + i * dim,
diff --git a/src/caffe/layers/softmax_layer.cu b/src/caffe/layers/softmax_layer.cu
index f53883c9206..ceeaff5b020 100644
--- a/src/caffe/layers/softmax_layer.cu
+++ b/src/caffe/layers/softmax_layer.cu
@@ -50,8 +50,7 @@ Dtype SoftmaxLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
   Dtype* scale_data = scale_.mutable_gpu_data();
   int num = bottom[0]->num();
   int dim = bottom[0]->count() / bottom[0]->num();
-  CUDA_CHECK(cudaMemcpy(top_data, bottom_data,
-      sizeof(Dtype) * bottom[0]->count(), cudaMemcpyDeviceToDevice));
+  caffe_copy(bottom[0]->count(), bottom_data, top_data);
   // we need to subtract the max to avoid numerical issues, compute the exp,
   // and then normalize.
   // Compute max
@@ -85,8 +84,7 @@ void SoftmaxLayer<Dtype>::Backward_gpu(const vector<Blob<Dtype>*>& top,
   Dtype* bottom_diff = (*bottom)[0]->mutable_gpu_diff();
   int num = top[0]->num();
   int dim = top[0]->count() / top[0]->num();
-  CUDA_CHECK(cudaMemcpy(bottom_diff, top_diff,
-      sizeof(Dtype) * top[0]->count(), cudaMemcpyDeviceToDevice));
+  caffe_copy(top[0]->count(), top_diff, bottom_diff);
   // Compute inner1d(top_diff, top_data) and subtract them from the bottom diff
   // cuda dot returns the result to cpu, so we temporarily change the pointer
   // mode
diff --git a/src/caffe/layers/softmax_loss_layer.cpp b/src/caffe/layers/softmax_loss_layer.cpp
index 1a3601aa9e6..37c5ebc45be 100644
--- a/src/caffe/layers/softmax_loss_layer.cpp
+++ b/src/caffe/layers/softmax_loss_layer.cpp
@@ -66,7 +66,7 @@ void SoftmaxWithLossLayer<Dtype>::Backward_cpu(const vector<Blob<Dtype>*>& top,
   if (propagate_down[0]) {
     Dtype* bottom_diff = (*bottom)[0]->mutable_cpu_diff();
     const Dtype* prob_data = prob_.cpu_data();
-    memcpy(bottom_diff, prob_data, sizeof(Dtype) * prob_.count());
+    caffe_copy(prob_.count(), prob_data, bottom_diff);
     const Dtype* label = (*bottom)[1]->cpu_data();
     int num = prob_.num();
     int dim = prob_.count() / num;
diff --git a/src/caffe/layers/window_data_layer.cpp b/src/caffe/layers/window_data_layer.cpp
index fd4860f98be..5dbdff330ff 100644
--- a/src/caffe/layers/window_data_layer.cpp
+++ b/src/caffe/layers/window_data_layer.cpp
@@ -59,7 +59,7 @@ void* WindowDataLayerPrefetch(void* layer_pointer) {
   bool use_square = (crop_mode == "square") ? true : false;
 
   // zero out batch
-  memset(top_data, 0, sizeof(Dtype)*layer->prefetch_data_->count());
+  caffe_set(layer->prefetch_data_->count(), Dtype(0), top_data);
 
   const int num_fg = static_cast<int>(static_cast<float>(batch_size)
       * fg_fraction);
diff --git a/src/caffe/layers/window_data_layer.cu b/src/caffe/layers/window_data_layer.cu
index bc49fef6545..ca664fcef41 100644
--- a/src/caffe/layers/window_data_layer.cu
+++ b/src/caffe/layers/window_data_layer.cu
@@ -28,12 +28,10 @@ Dtype WindowDataLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
   // First, join the thread
   JoinPrefetchThread();
   // Copy the data
-  CUDA_CHECK(cudaMemcpy((*top)[0]->mutable_gpu_data(),
-      prefetch_data_->cpu_data(), sizeof(Dtype) * prefetch_data_->count(),
-      cudaMemcpyHostToDevice));
-  CUDA_CHECK(cudaMemcpy((*top)[1]->mutable_gpu_data(),
-      prefetch_label_->cpu_data(), sizeof(Dtype) * prefetch_label_->count(),
-      cudaMemcpyHostToDevice));
+  caffe_copy(prefetch_data_->count(), prefetch_data_->cpu_data(),
+      (*top)[0]->mutable_gpu_data());
+  caffe_copy(prefetch_label_->count(), prefetch_label_->cpu_data(),
+      (*top)[1]->mutable_gpu_data());
   // Start a new prefetch thread
   CreatePrefetchThread();
   return Dtype(0.);
diff --git a/src/caffe/solver.cpp b/src/caffe/solver.cpp
index 769618175ac..ca1d92525c9 100644
--- a/src/caffe/solver.cpp
+++ b/src/caffe/solver.cpp
@@ -310,7 +310,7 @@ void SGDSolver<Dtype>::ComputeUpdateValue() {
             history_[param_id]->mutable_gpu_data());
       }
       // copy
-      caffe_gpu_copy(net_params[param_id]->count(),
+      caffe_copy(net_params[param_id]->count(),
           history_[param_id]->gpu_data(),
           net_params[param_id]->mutable_gpu_diff());
     }
diff --git a/src/caffe/syncedmem.cpp b/src/caffe/syncedmem.cpp
index fec37d6e9ec..9fe55280de9 100644
--- a/src/caffe/syncedmem.cpp
+++ b/src/caffe/syncedmem.cpp
@@ -6,6 +6,7 @@
 
 #include "caffe/common.hpp"
 #include "caffe/syncedmem.hpp"
+#include "caffe/util/math_functions.hpp"
 
 namespace caffe {
 
@@ -32,7 +33,7 @@ inline void SyncedMemory::to_cpu() {
       CaffeMallocHost(&cpu_ptr_, size_);
       own_cpu_data_ = true;
     }
-    CUDA_CHECK(cudaMemcpy(cpu_ptr_, gpu_ptr_, size_, cudaMemcpyDeviceToHost));
+    caffe_memcpy(size_, gpu_ptr_, cpu_ptr_);
     head_ = SYNCED;
     break;
   case HEAD_AT_CPU:
@@ -52,7 +53,7 @@ inline void SyncedMemory::to_gpu() {
     if (gpu_ptr_ == NULL) {
       CUDA_CHECK(cudaMalloc(&gpu_ptr_, size_));
     }
-    CUDA_CHECK(cudaMemcpy(gpu_ptr_, cpu_ptr_, size_, cudaMemcpyHostToDevice));
+    caffe_memcpy(size_, cpu_ptr_, gpu_ptr_);
     head_ = SYNCED;
     break;
   case HEAD_AT_GPU:
diff --git a/src/caffe/test/test_gradient_check_util.hpp b/src/caffe/test/test_gradient_check_util.hpp
index ff104b90eb9..2d551f82428 100644
--- a/src/caffe/test/test_gradient_check_util.hpp
+++ b/src/caffe/test/test_gradient_check_util.hpp
@@ -230,7 +230,7 @@ Dtype GradientChecker<Dtype>::GetObjAndGradient(vector<Blob<Dtype>*>* top,
         loss += top_blob_data[j] * top_blob_data[j];
       }
       // set the diff: simply the data.
-      memcpy(top_blob_diff, top_blob_data, sizeof(Dtype) * top_blob->count());
+      caffe_copy(top_blob->count(), top_blob_data, top_blob_diff);
     }
     loss /= 2.;
   } else {
@@ -238,7 +238,7 @@ Dtype GradientChecker<Dtype>::GetObjAndGradient(vector<Blob<Dtype>*>* top,
     for (int i = 0; i < top->size(); ++i) {
       Blob<Dtype>* top_blob = (*top)[i];
       Dtype* top_blob_diff = top_blob->mutable_cpu_diff();
-      memset(top_blob_diff, 0, sizeof(Dtype) * top_blob->count());
+      caffe_set(top_blob->count(), Dtype(0), top_blob_diff);
     }
     loss = (*top)[top_id]->cpu_data()[top_data_id];
     (*top)[top_id]->mutable_cpu_diff()[top_data_id] = 1.;
diff --git a/src/caffe/test/test_math_functions.cpp b/src/caffe/test/test_math_functions.cpp
index d0265767c07..941d8b9479a 100644
--- a/src/caffe/test/test_math_functions.cpp
+++ b/src/caffe/test/test_math_functions.cpp
@@ -219,7 +219,7 @@ TYPED_TEST(MathFunctionsTest, TestCopyGPU) {
   const int n = this->blob_bottom_->count();
   const TypeParam* bottom_data = this->blob_bottom_->gpu_data();
   TypeParam* top_data = this->blob_top_->mutable_gpu_data();
-  caffe_gpu_copy(n, bottom_data, top_data);
+  caffe_copy(n, bottom_data, top_data);
   bottom_data = this->blob_bottom_->cpu_data();
   top_data = this->blob_top_->mutable_cpu_data();
   for (int i = 0; i < n; ++i) {
diff --git a/src/caffe/test/test_platform.cpp b/src/caffe/test/test_platform.cpp
index c3868f34d9f..7cf8306e8c9 100644
--- a/src/caffe/test/test_platform.cpp
+++ b/src/caffe/test/test_platform.cpp
@@ -47,6 +47,8 @@ TEST_F(PlatformTest, TestInitialization) {
          CAFFE_TEST_CUDA_PROP.multiProcessorCount);
   printf("Kernel execution timeout:      %s\n",
          (CAFFE_TEST_CUDA_PROP.kernelExecTimeoutEnabled ? "Yes" : "No"));
+  printf("Unified virtual addressing:    %s\n",
+         (CAFFE_TEST_CUDA_PROP.unifiedAddressing ? "Yes" : "No"));
   EXPECT_TRUE(true);
 }
 
diff --git a/src/caffe/test/test_syncedmem.cpp b/src/caffe/test/test_syncedmem.cpp
index 7bbbbab6a57..3aaeafc353e 100644
--- a/src/caffe/test/test_syncedmem.cpp
+++ b/src/caffe/test/test_syncedmem.cpp
@@ -7,6 +7,7 @@
 #include "gtest/gtest.h"
 #include "caffe/common.hpp"
 #include "caffe/syncedmem.hpp"
+#include "caffe/util/math_functions.hpp"
 
 #include "caffe/test/test_caffe_main.hpp"
 
@@ -57,8 +58,7 @@ TEST_F(SyncedMemoryTest, TestGPURead) {
   EXPECT_EQ(mem.head(), SyncedMemory::SYNCED);
   // check if values are the same
   char* recovered_value = new char[10];
-  cudaMemcpy(reinterpret_cast<void*>(recovered_value), gpu_data, 10,
-             cudaMemcpyDeviceToHost);
+  caffe_memcpy(10, gpu_data, recovered_value);
   for (int i = 0; i < mem.size(); ++i) {
     EXPECT_EQ((reinterpret_cast<char*>(recovered_value))[i], 1);
   }
@@ -72,8 +72,7 @@ TEST_F(SyncedMemoryTest, TestGPURead) {
   gpu_data = mem.gpu_data();
   EXPECT_EQ(mem.head(), SyncedMemory::SYNCED);
   // check if values are the same
-  cudaMemcpy(reinterpret_cast<void*>(recovered_value), gpu_data, 10,
-             cudaMemcpyDeviceToHost);
+  caffe_memcpy(10, gpu_data, recovered_value);
   for (int i = 0; i < mem.size(); ++i) {
     EXPECT_EQ((reinterpret_cast<char*>(recovered_value))[i], 2);
   }
diff --git a/src/caffe/test/test_util_blas.cpp b/src/caffe/test/test_util_blas.cpp
index 2e4c6795952..5b4c48eaf80 100644
--- a/src/caffe/test/test_util_blas.cpp
+++ b/src/caffe/test/test_util_blas.cpp
@@ -30,8 +30,8 @@ TYPED_TEST(GemmTest, TestGemm) {
   TypeParam A_reshape_data[6] = {1, 4, 2, 5, 3, 6};
   TypeParam B_reshape_data[12] = {1, 5, 9, 2, 6, 10, 3, 7, 11, 4, 8, 12};
   TypeParam result[8] = {38, 44, 50, 56, 83, 98, 113, 128};
-  memcpy(A.mutable_cpu_data(), data, 6 * sizeof(TypeParam));
-  memcpy(B.mutable_cpu_data(), data, 12 * sizeof(TypeParam));
+  caffe_copy(6, data, A.mutable_cpu_data());
+  caffe_copy(12, data, B.mutable_cpu_data());
 
   if (sizeof(TypeParam) == 4 || CAFFE_TEST_CUDA_PROP.major >= 2) {
     // [1, 2, 3; 4 5 6] * [1, 2, 3, 4; 5, 6, 7, 8; 9, 10, 11, 12];
@@ -48,7 +48,7 @@ TYPED_TEST(GemmTest, TestGemm) {
 
     // Test when we have a transposed A
     A.Reshape(1, 1, 3, 2);
-    memcpy(A.mutable_cpu_data(), A_reshape_data, 6 * sizeof(TypeParam));
+    caffe_copy(6, A_reshape_data, A.mutable_cpu_data());
     caffe_cpu_gemm<TypeParam>(CblasTrans, CblasNoTrans, 2, 4, 3, 1.,
         A.cpu_data(), B.cpu_data(), 0., C.mutable_cpu_data());
     for (int i = 0; i < 8; ++i) {
@@ -62,7 +62,7 @@ TYPED_TEST(GemmTest, TestGemm) {
 
     // Test when we have a transposed A and a transposed B too
     B.Reshape(1, 1, 4, 3);
-    memcpy(B.mutable_cpu_data(), B_reshape_data, 12 * sizeof(TypeParam));
+    caffe_copy(12, B_reshape_data, B.mutable_cpu_data());
     caffe_cpu_gemm<TypeParam>(CblasTrans, CblasTrans, 2, 4, 3, 1.,
         A.cpu_data(), B.cpu_data(), 0., C.mutable_cpu_data());
     for (int i = 0; i < 8; ++i) {
@@ -76,7 +76,7 @@ TYPED_TEST(GemmTest, TestGemm) {
 
     // Test when we have a transposed B
     A.Reshape(1, 1, 2, 3);
-    memcpy(A.mutable_cpu_data(), data, 6 * sizeof(TypeParam));
+    caffe_copy(6, data, A.mutable_cpu_data());
     caffe_cpu_gemm<TypeParam>(CblasNoTrans, CblasTrans, 2, 4, 3, 1.,
         A.cpu_data(), B.cpu_data(), 0., C.mutable_cpu_data());
     for (int i = 0; i < 8; ++i) {
@@ -100,8 +100,8 @@ TYPED_TEST(GemmTest, TestGemv) {
   TypeParam data[6] = {1, 2, 3, 4, 5, 6};
   TypeParam result_2[2] = {14, 32};
   TypeParam result_3[3] = {9, 12, 15};
-  memcpy(A.mutable_cpu_data(), data, 6 * sizeof(TypeParam));
-  memcpy(x.mutable_cpu_data(), data, 3 * sizeof(TypeParam));
+  caffe_copy(6, data, A.mutable_cpu_data());
+  caffe_copy(3, data, x.mutable_cpu_data());
 
   if (sizeof(TypeParam) == 4 || CAFFE_TEST_CUDA_PROP.major >= 2) {
     caffe_cpu_gemv<TypeParam>(CblasNoTrans, 2, 3, 1., A.cpu_data(),
@@ -116,7 +116,7 @@ TYPED_TEST(GemmTest, TestGemv) {
     }
 
     // Test transpose case
-    memcpy(y.mutable_cpu_data(), data, 2 * sizeof(TypeParam));
+    caffe_copy(2, data, y.mutable_cpu_data());
     caffe_cpu_gemv<TypeParam>(CblasTrans, 2, 3, 1., A.cpu_data(),
         y.cpu_data(), 0., x.mutable_cpu_data());
     for (int i = 0; i < 3; ++i) {
diff --git a/src/caffe/util/im2col.cpp b/src/caffe/util/im2col.cpp
index 037410e29b7..ce4e18848bf 100644
--- a/src/caffe/util/im2col.cpp
+++ b/src/caffe/util/im2col.cpp
@@ -5,6 +5,7 @@
 #include <cstring>
 
 #include "caffe/util/im2col.hpp"
+#include "caffe/util/math_functions.hpp"
 
 namespace caffe {
 
@@ -45,7 +46,7 @@ template <typename Dtype>
 void col2im_cpu(const Dtype* data_col, const int channels,
     const int height, const int width, const int ksize, const int pad,
     const int stride, Dtype* data_im) {
-  memset(data_im, 0, sizeof(Dtype) * height * width * channels);
+  caffe_set(height * width * channels, Dtype(0), data_im);
   int height_col = (height + 2 * pad - ksize) / stride + 1;
   int width_col = (width + 2 * pad - ksize) / stride + 1;
   int channels_col = channels * ksize * ksize;
diff --git a/src/caffe/util/im2col.cu b/src/caffe/util/im2col.cu
index ec4465eff66..79faa6cb5c9 100644
--- a/src/caffe/util/im2col.cu
+++ b/src/caffe/util/im2col.cu
@@ -106,8 +106,6 @@ template <typename Dtype>
 void col2im_gpu(const Dtype* data_col, const int channels,
     const int height, const int width, const int ksize, const int pad,
     const int stride, Dtype* data_im) {
-  // CUDA_CHECK(cudaMemset(data_im, 0,
-  //            sizeof(Dtype) * height * width * channels));
   int height_col = (height + 2 * pad - ksize) / stride + 1;
   int width_col = (width + 2 * pad - ksize) / stride + 1;
   int num_kernels = channels * height * width;
diff --git a/src/caffe/util/math_functions.cpp b/src/caffe/util/math_functions.cpp
index 90df51248e2..918bb3c361c 100644
--- a/src/caffe/util/math_functions.cpp
+++ b/src/caffe/util/math_functions.cpp
@@ -149,31 +149,22 @@ void caffe_add_scalar(const int N, const double alpha, double* Y) {
   }
 }
 
-template <>
-void caffe_copy<float>(const int N, const float* X, float* Y) {
-  if (X != Y) {
-    cblas_scopy(N, X, 1, Y, 1);
-  }
-}
-
-template <>
-void caffe_copy<double>(const int N, const double* X, double* Y) {
+template <typename Dtype>
+void caffe_copy(const int N, const Dtype* X, Dtype* Y) {
   if (X != Y) {
-    cblas_dcopy(N, X, 1, Y, 1);
+    CUDA_CHECK(cudaMemcpy(Y, X, sizeof(Dtype) * N, cudaMemcpyDefault));
   }
 }
 
-template <>
-void caffe_gpu_copy<float>(const int N, const float* X, float* Y) {
-  if (X != Y) {
-    CUBLAS_CHECK(cublasScopy(Caffe::cublas_handle(), N, X, 1, Y, 1));
-  }
-}
+template void caffe_copy<int>(const int N, const int* X, int* Y);
+template void caffe_copy<unsigned int>(const int N, const unsigned int* X,
+    unsigned int* Y);
+template void caffe_copy<float>(const int N, const float* X, float* Y);
+template void caffe_copy<double>(const int N, const double* X, double* Y);
 
-template <>
-void caffe_gpu_copy<double>(const int N, const double* X, double* Y) {
+void caffe_memcpy(const size_t N, const void* X, void* Y) {
   if (X != Y) {
-    CUBLAS_CHECK(cublasDcopy(Caffe::cublas_handle(), N, X, 1, Y, 1));
+    CUDA_CHECK(cudaMemcpy(Y, X, N, cudaMemcpyDefault));
   }
 }
 
diff --git a/src/caffe/util/math_functions.cu b/src/caffe/util/math_functions.cu
index 63c8fac69c5..849e53b9ca4 100644
--- a/src/caffe/util/math_functions.cu
+++ b/src/caffe/util/math_functions.cu
@@ -20,27 +20,20 @@ __global__ void set_kernel(const int n, const Dtype alpha, Dtype* y) {
   }
 }
 
-template <>
-void caffe_gpu_set(const int N, const float alpha, float* Y) {
+template <typename Dtype>
+void caffe_gpu_set(const int N, const Dtype alpha, Dtype* Y) {
   if (alpha == 0) {
-    CUDA_CHECK(cudaMemset(Y, 0, sizeof(float) * N));
+    CUDA_CHECK(cudaMemset(Y, 0, sizeof(Dtype) * N));
     return;
   }
   // NOLINT_NEXT_LINE(whitespace/operators)
-  set_kernel<float><<<CAFFE_GET_BLOCKS(N), CAFFE_CUDA_NUM_THREADS>>>(
+  set_kernel<Dtype><<<CAFFE_GET_BLOCKS(N), CAFFE_CUDA_NUM_THREADS>>>(
       N, alpha, Y);
 }
 
-template <>
-void caffe_gpu_set(const int N, const double alpha, double* Y) {
-  if (alpha == 0) {
-    CUDA_CHECK(cudaMemset(Y, 0, sizeof(double) * N));
-    return;
-  }
-  // NOLINT_NEXT_LINE(whitespace/operators)
-  set_kernel<double><<<CAFFE_GET_BLOCKS(N), CAFFE_CUDA_NUM_THREADS>>>(
-      N, alpha, Y);
-}
+template void caffe_gpu_set<int>(const int N, const int alpha, int* Y);
+template void caffe_gpu_set<float>(const int N, const float alpha, float* Y);
+template void caffe_gpu_set<double>(const int N, const double alpha, double* Y);
 
 template <typename Dtype>
 __global__ void add_scalar_kernel(const int n, const Dtype alpha, Dtype* y) {