         int64_t stride = 1;
 
         // Up-Sweep Phase
-        for (stride = 1; stride * 2 <= reduce_dim; stride <<= 1) {
+        for (stride = 1; stride * 2 <= reduction_dim; stride <<= 1) {
             int64_t index_from = reduce_block_offset * stride * 2 + stride - 1;
             int64_t index_to = index_from + stride;
             auto it_from = arg.out_indexer.It(i + index_from * len, reduce_stride * stride * 2);
             auto it_to = arg.out_indexer.It(i + index_to * len, reduce_stride * stride * 2);
-            for (int64_t j = index_to; j < reduce_dim; j += reduce_block_size * stride * 2) {
+            for (int64_t j = index_to; j < reduction_dim; j += reduce_block_size * stride * 2) {
                 reduce(it_from, it_to);
             }
             __syncthreads();

    // Currently, only contiguous output arrays are optimized.
    if (arg.in_strides().ndim() == 1 && arg.out_strides().ndim() == 1) {
        reduce_detail::ScanKernel<<<grid_size, block_size, shared_mem_size>>>(
                MakeReductionKernelArg<In, Out, 1, 1>(arg), out_block_size, reduce_block_size, impl, reduce_len);
        return;
    }
#endif  // NDEBUG

    reduce_detail::ScanKernel<<<grid_size, block_size, shared_mem_size>>>(
            MakeReductionKernelArg<In, Out>(arg), out_block_size, reduce_block_size, impl, reduce_len);
}

}  // namespace cuda

        int out_block_size,
        int reduce_block_size,
        ReductionImpl impl,
        int64_t reduction_dim) {
    int tid = threadIdx.x;

    int64_t len = arg.out_indexer.total_size() / reduction_dim;
    int64_t reduce_block_offset = tid / out_block_size;
    int64_t reduce_offset = reduce_block_offset * len;
    int64_t reduce_stride = reduce_block_size * len;
