[Feature request] Support strides for `T.serial`

[LyricZhao]

For lots of CUDA kernels, we conventionaly write as:

for (int i = thread_idx; i < numel; i += num_threads)
    out[i] = 0;

But in tilelang:

for i in T.serial(0, T.ceildiv(numel - thread_idx, num_threads)):
    j = thread_idx + i * num_threads
    out[j] = -1

which is annoying to manually compute the range and index transformation, also introduces more registers and lower index computation performance.

So I purpose supporting strides in T.serial.

[LyricZhao]

Also negative strides are wanted, e.g. reversed loops.

[LeiWang1999]

now we use a workaround to implement for with steps, because a lot of pass assume step to be 1:

__global__ void vecAdd(const float *A, const float *B, float *C, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = blockDim.x * gridDim.x;
    for (int i = tid; i < n; i += stride) {
        C[i] = A[i] + B[i];
    }
}

=>

__global__ void vecAdd2(const float *A, const float *B, float *C, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = blockDim.x * gridDim.x;
    for (int j = 0; j < (n + stride - 1) / stride; ++j) {
        int i = tid + j * stride;
        C[i] = A[i] + B[i];
    }
}

However, it might be better for us to implement the loop with actual steps, as the Sass output from those implementations is not identical. We can join the TVM threads to discuss and address this together:
https://discuss.tvm.apache.org/t/do-we-have-plan-to-introduce-step-attribute-to-fornode/18685/3?u=leiwang1999