echuraev commented on PR #15137:
URL: https://github.com/apache/tvm/pull/15137#issuecomment-1609406063
@tqchen During implementing a limitation on the number of function
arguments, I met a problem. In case of tensor with dynamic shape, it is
possible that in addition to the parameters of the kernel and indices by
dynamic dimensions, strides will be passed to the function. For example, below
an example of the OpenCL code which was generated for one of concats from the
test on split_args:
```opencl
__kernel void vm_mod_fused_concatenate_kernel(__global float* restrict
T_concat,
__global float* restrict p0,
__global float* restrict p1,
__global float* restrict p2,
__global float* restrict p3,
__global float* restrict p4,
int any_dim,
int any_dim_1,
int any_dim_2,
int any_dim_3,
int any_dim_4,
int stride,
int stride_1,
int stride_2,
int stride_3,
int stride_4,
int stride_5,
int stride_6,
int stride_7,
int stride_8,
int stride_9,
int stride_10,
int stride_11,
int stride_12,
int stride_13,
int stride_14,
int stride_15,
int stride_16) {
// ...
}
```
As you can see, OpenCL kernel has three types of arguments: buffers, indices
of dynamic dimensions and strides. Unfortunately, I'm not able to get the
number of strides in `SplitArgs` and `FuseOps` passes. Strides appear only
after the stage of `LowerTE`.
To handle this problem, I added a new attribute to target
`max_function_args_dyn` which is used in case if at least one of the tensors
has dynamic shape. This attribute should have less value than
`max_function_args` and it helps to run topologies, which previously didn't
work in TVM. But I have some doubts about this solution. Probably you know the
better way, how can I calculate the number of strides which will be passed to
the kernel in case of dynamic shapes?
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
