From 5f4e38fd88ad14e5a3c461f864bef443b1689c8b Mon Sep 17 00:00:00 2001
From: Jeff Donahue <jeff.donahue@gmail.com>
Date: Tue, 18 Mar 2014 17:34:25 -0700
Subject: [PATCH] fix remaining issues related to CUDA_KERNEL_LOOP

---
 include/caffe/common.hpp          |  2 +-
 src/caffe/layers/pooling_layer.cu | 17 ++++++++---------
 2 files changed, 9 insertions(+), 10 deletions(-)

diff --git a/include/caffe/common.hpp b/include/caffe/common.hpp
index f0cae65..96ba58c 100644
--- a/include/caffe/common.hpp
+++ b/include/caffe/common.hpp
@@ -22,7 +22,7 @@
   for (int i = blockIdx.x * blockDim.x + threadIdx.x; \
        i < (n); \
        i += blockDim.x * gridDim.x)
-       
+
 // After a kernel is executed, this will check the error and if there is one,
 // exit loudly.
 #define CUDA_POST_KERNEL_CHECK \
diff --git a/src/caffe/layers/pooling_layer.cu b/src/caffe/layers/pooling_layer.cu
index 8735dde..357a392 100644
--- a/src/caffe/layers/pooling_layer.cu
+++ b/src/caffe/layers/pooling_layer.cu
@@ -35,7 +35,7 @@ __global__ void MaxPoolForward(const int nthreads, const Dtype* bottom_data,
       }
     }
     top_data[index] = maxval;
-  }  // (if index < nthreads)
+  }
 }
 
 template <typename Dtype>
@@ -43,8 +43,7 @@ __global__ void AvePoolForward(const int nthreads, const Dtype* bottom_data,
     const int num, const int channels, const int height,
     const int width, const int pooled_height, const int pooled_width,
     const int ksize, const int stride, Dtype* top_data) {
-  int index = threadIdx.x + blockIdx.x * blockDim.x;
-  if (index < nthreads) {
+  CUDA_KERNEL_LOOP(index, nthreads) {
     int pw = index % pooled_width;
     int ph = (index / pooled_width) % pooled_height;
     int c = (index / pooled_width / pooled_height) % channels;
@@ -61,7 +60,7 @@ __global__ void AvePoolForward(const int nthreads, const Dtype* bottom_data,
       }
     }
     top_data[index] = aveval / (hend - hstart) / (wend - wstart);
-  }  // (if index < nthreads)
+  }
 }
 
 template <typename Dtype>
@@ -100,7 +99,7 @@ __global__ void StoPoolForwardTrain(const int nthreads,
         }
       }
     }
-  }  // (if index < nthreads)
+  }
 }
 
 
@@ -131,7 +130,7 @@ __global__ void StoPoolForwardTest(const int nthreads,
       }
     }
     top_data[index] = cumvalues / cumsum;
-  }  // (if index < nthreads)
+  }
 }
 
 
@@ -211,7 +210,7 @@ __global__ void MaxPoolBackward(const int nthreads, const Dtype* bottom_data,
       }
     }
     bottom_diff[index] = gradient;
-  }  // (if index < nthreads)
+  }
 }
 
 
@@ -242,7 +241,7 @@ __global__ void AvePoolBackward(const int nthreads, const Dtype* top_diff,
       }
     }
     bottom_diff[index] = gradient;
-  }  // (if index < nthreads)
+  }
 }
 
 
@@ -273,7 +272,7 @@ __global__ void StoPoolBackward(const int nthreads,
       }
     }
     bottom_diff[index] = gradient;
-  }  // (if index < nthreads)
+  }
 }
 
 
-- 
2.7.4