   dim4 tgt_stride(0);
   dim4 size(1);
 
-  uint ndim = THCudaTensor_nDimension(src);
-  for(uint dim=0; dim < ndim; dim++) {
-    uint odim = ndim - 1 - dim;
+  unsigned ndim = THCudaTensor_nDimension(src);
+  for(unsigned dim=0; dim < ndim; dim++) {
+    unsigned odim = ndim - 1 - dim;
     src_stride[odim] = THCudaTensor_stride(src, dim);
     tgt_stride[odim] = THCudaTensor_stride(tgt, dim);
     size[odim]       = THCudaTensor_size(src, dim);
   }
 
   dim3 threads(32, 16);
-  uint nBlockPerRow = (size[1] + threads.y - 1) / threads.y;
+  unsigned nBlockPerRow = (size[1] + threads.y - 1) / threads.y;
   dim3 grid(size[2], nBlockPerRow, size[3]);
 
   THCudaTensor_kernel_reduceInnermostDim<<<grid, threads>>>(THCudaTensor_data(tgt),


  T acc = init;

  uint row = blockIdx.y * blockDim.y + threadIdx.y;
  float *src = src_ + blockIdx.z * src_stride[3] + blockIdx.x * src_stride[2] + row * src_stride[1];
  bool reducing = threadIdx.x < blockDim.y
                && blockIdx.y * blockDim.y + threadIdx.x < size[1]
                && threadIdx.y == 0;

  for(uint bCol=0; bCol < size[0]; bCol += blockDim.x) {

    sbuf[threadIdx.y][threadIdx.x] = init;
    uint col = bCol + threadIdx.x;
    if(row < size[1] && col < size[0]) {
      sbuf[threadIdx.y][threadIdx.x] = src[col];
    }
    __syncthreads();

    if(reducing) {
      for(uint x=0; x < blockDim.x; ++x) {
        /* Could eliminate shared memory bank conflicst hare by using modulo 32
         * addressing: (x + threadIdx) % 32. However, due to the memory-bound
         * nature of the kernel this makes no difference to performance */

  }

  if(reducing) {
    unsigned row = blockIdx.y * blockDim.y + threadIdx.x;
    unsigned tgt_offset = blockIdx.z * tgt_stride[3] + blockIdx.x * tgt_stride[2];
    tgt[tgt_offset + row] = float(acc);
  }
}
