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 \
}
}
top_data[index] = maxval;
- } // (if index < nthreads)
+ }
}
template <typename Dtype>
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;
}
}
top_data[index] = aveval / (hend - hstart) / (wend - wstart);
- } // (if index < nthreads)
+ }
}
template <typename Dtype>
}
}
}
- } // (if index < nthreads)
+ }
}
}
}
top_data[index] = cumvalues / cumsum;
- } // (if index < nthreads)
+ }
}
}
}
bottom_diff[index] = gradient;
- } // (if index < nthreads)
+ }
}
}
}
bottom_diff[index] = gradient;
- } // (if index < nthreads)
+ }
}
}
}
bottom_diff[index] = gradient;
- } // (if index < nthreads)
+ }
}