 // Use grid strided loops, descriped here:
 // https://devblogs.nvidia.com/cuda-pro-tip-write-flexible-kernels-grid-stride-loops/
+// This pattern ensures that all of the loop values are visited once, no matter
+// what grid parameters are used for the function.
 	
 extern "C" __global__
 void seq2col(float* output,

    //    d_seq[i] += d_cols[i+1, 1]
    //    d_seq[i] += d_cols[i+2, 0]

    for (int b = blockIdx.x * blockDim.x + threadIdx.x; b < B; b += blockDim.x * gridDim.x)
    {
        int nF = nW * 2 + 1;
        int seq_row = b * I;
        int col_feat = nF * I;
        for (int f=-nW; f < (nW+1); ++f)
        {
            int col_row = (b+f) * (I * nF);
            col_feat -= I;
            if ((col_row >= 0) && (col_row < (B*I*nF)))
            {
                int start = col_row + col_feat;
                if ((start >= 0) && ((start+I) < (B*I*nF)))
                {
                    for (int i=0; i < I; ++i)
            	    atomicAdd(&d_seqs[seq_row+i], d_cols[start+i]);
                }
            }
        }

    // * x_start=-6, x_end=9 : (0-2) * 3, (0+2+1) * 3
    // * x_start=-3, x_end=13 : (1-2) * 3, (1+2+1) * 3
    // * x_start=0, x_end=16 : (2-2) * 3, (2+2+1) * 3
    int _loop_start = blockIdx.x * blockDim.x + threadIdx.x;
    int _loop_stride = blockDim.x * gridDim.x;
    int nF = nW * 2 + 1;
    for (int b = _loop_start; b < B; b += _loop_stride)
    {
        int o_start = b * I * nF;
        // Let's say b=0, nW=1, I=10, B=20
        // x_start = (0-1) * 10 : -10
