                     y_indexer,
                     kernel_indexer,
                     CudaStackVector{stride_},
-                    CudaStackVector{pad_},
-                    NdimIndex{x_iarray.ndim()});
+                    CudaStackVector{pad_});
         });
 
         return ggy;

                int64_t idx = it_y.index()[i] * stride.data[i - 2] - pad.data[i - 2] + it_kernel.index()[i - 2];
                idx = max(idx, int64_t{0});
                idx = min(idx, x_indexer.shape()[i] - 1);
                x_index.index()[i] = idx;
            }
            it_x.CopyIndex(x_index);
            if (y == x_iarray[it_x]) {
                ggy_iarray[it_y] = ggx_iarray[it_x];
            }

        Indexer<> y_indexer,
        Indexer<> kernel_indexer,
        CudaStackVector stride,
        CudaStackVector pad) {
    auto it_y = y_indexer.It(blockIdx.x * blockDim.x + threadIdx.x, blockDim.x * gridDim.x);
    auto it_kernel = kernel_indexer.It(kernel_indexer.total_size() - 1);
    auto it_x = x_indexer.It(0);

    for (it_y.Restart(); it_y; ++it_y) {
        it_x.index()[0] = it_y.index()[0];  // batch.
        it_x.index()[1] = it_y.index()[1];  // channel.

        T y = y_iarray[it_y];

