         // Copy input array to output array
         auto it_in = arg.in_indexer.It(i + reduce_offset, reduce_stride);
         auto it_out = arg.out_indexer.It(i + reduce_offset, reduce_stride);
-        for (int64_t j = reduce_block_offset; j < reduce_dim; j += reduce_block_size, ++it_in, ++it_out) {
+        for (int64_t j = reduce_block_offset; j < reduction_dim; j += reduce_block_size, ++it_in, ++it_out) {
             auto value = cuda_internal::StorageToDataType<const In>(arg.in[it_in]);
             arg.out[it_out] = cuda_internal::DataToStorageType<Out>(impl.MapIn(value, j));
         }

            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) {
                reduce(it_from, it_to);
            }
            __syncthreads();

    }

    ReductionArg arg{in, Axes{axis}, out};
    int64_t reduction_dim = in.shape()[axis];

    // TODO(niboshi): Calculate kMaxBlockSize per device
    std::lock_guard<std::mutex> lock{*cuda_internal::g_mutex};
