         int out_block_size,
         int reduce_block_size,
         ReductionImpl impl,
-        int64_t reduce_dim) {
+        int64_t reduction_dim) {
     int tid = threadIdx.x;
 
-    int64_t len = arg.out_indexer.total_size() / reduce_dim;
+    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;

        // 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) {
            auto value = cuda_internal::StorageToDataType<const In>(arg.in[it_in]);
            arg.out[it_out] = cuda_internal::DataToStorageType<Out>(impl.MapIn(value, j));
        }

    }

    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};
