[Apache TVM Discuss] [Questions] Intrinsic Function with Constant Input Array

[moderato via Apache TVM Discuss <no...@discuss.tvm.ai>]

Hello, to get started I have an example code of intrinsic function like this:

```python
from __future__ import absolute_import, print_function

import tvm
from tvm import te
import numpy as np
from tvm.topi.utils import get_const_tuple

ctx = tvm.context("cpu", 0)
M = 8
factor = 4
A = te.placeholder((M,), name='A')
B = te.placeholder((M,), name='B')
C = te.compute((M,), lambda i: A[i] + B[i], name='C')
s = te.create_schedule(C.op)
x, = C.op.axis
xo, xi = s[C].split(x, factor=factor)
print(tvm.lower(s, [A, B, C], simple_mode=True))
dtype = A.dtype

def intrin_add(m):
    a = te.placeholder((m,), name='a')
    b = te.placeholder((m,), name='b')
    c = te.compute((m,), lambda i: a[i] + b[i], name='c')
    d = tvm.nd.array(np.zeros(get_const_tuple(C.shape), dtype=dtype), ctx)
    Ab = tvm.tir.decl_buffer(a.shape, a.dtype,
                         name="A",
                         offset_factor=1,
                         strides=[1])
    Bb = tvm.tir.decl_buffer(b.shape, b.dtype,
                         name="B",
                         offset_factor=1,
                         strides=[1])
    Cb = tvm.tir.decl_buffer(c.shape, c.dtype,
                         name="C",
                         offset_factor=1,
                         strides=[1])
    def intrin_func(ins, outs):
        aa, bb = ins
        cc = outs[0]
        return tvm.tir.call_extern('int32', 'add', aa.access_ptr('r'), bb.access_ptr('r'), cc.access_ptr('w'), m)

    return te.decl_tensor_intrin(c.op, intrin_func, binds={a: Ab, b: Bb, c: Cb})

add = intrin_add(factor)
s[C].tensorize(xi, add)
print(tvm.lower(s, [A, B, C], simple_mode=True))
func = tvm.build(s, [A, B, C], target="llvm -mcpu=core-avx2", name="add")

a = np.random.uniform(size=get_const_tuple(A.shape)).astype(dtype)
b = np.random.uniform(size=get_const_tuple(B.shape)).astype(dtype)
c = tvm.nd.array(np.zeros(get_const_tuple(C.shape), dtype=dtype), ctx)
func(tvm.nd.array(a, ctx), tvm.nd.array(b, ctx), c)
np.testing.assert_allclose(c.asnumpy(), np.add(a, b), rtol=1e-3)
```
while the extern C function looks like:
```C
extern "C" int add(const float* a, const float* b, float* c, int M) {
    for (int i = 0; i < M; i++) {
        c[i] = a[i] + b[i];
    }
    return 0;
}
```

Now I wanna pass a constant input array `dd` to the intrinsic function so that it looks like:
```python
def intrin_func(ins, outs):
    aa, bb = ins
    cc = outs[0]
    ##
    # compute dd here
    ##
    return tvm.tir.call_extern('int32', 'add', 
                                aa.access_ptr('r'), 
                                bb.access_ptr('r'), 
                                cc.access_ptr('w'), 
                                dd,
                                m)
```
and correspondingly the C function's API becomes
```C
extern "C" int add(const float* a, const float* b, float* c, const float* d, int M)
```

I have tried multiple ways, e.g. passing a list / ndarray to the function, etc, but they all failed. It seems like the intrinsic function requires inputs to be `string` or `primexpr` and I didn't see any types that is compatible with a list or an array.

Any one can help? Thanks!





---
[Visit Topic](https://discuss.tvm.apache.org/t/intrinsic-function-with-constant-input-array/9955/1) to respond.

You are receiving this because you enabled mailing list mode.

To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/aa9ca6fd668d7fb2c77ed376b05de562ba4beebf3cb952ad8e5ec6663a73161d).

[Apache TVM Discuss] [Questions] Intrinsic Function with Constant Input Array

[tqchen via Apache TVM Discuss <no...@discuss.tvm.ai>]

In this case we should make it a runtime array. This is also how relay handles constants





---
[Visit Topic](https://discuss.tvm.apache.org/t/intrinsic-function-with-constant-input-array/9955/4) to respond.

You are receiving this because you enabled mailing list mode.

To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/0a98b156f258fc5a941618f621531ed7246ef67b3ccee071492af764e0fec1ce).

[Apache TVM Discuss] [Questions] Intrinsic Function with Constant Input Array

[moderato via Apache TVM Discuss <no...@discuss.tvm.ai>]

Thanks for the reply! 
> If you want to inline a constant array(that is really small), you can attempt to use select to dispatch each index to the value.

Do you mean putting all the values I wanna pass in the C API individually?

In my case, the array size would be around 100 so I think I have to go for the second. I'm fine if I have to make it a normal array. It was being computed inside the C function during the runtime, but later I found it can be computed during the compile time so I tried to bring it out and that's the whole point I asked this.

What should I do here?





---
[Visit Topic](https://discuss.tvm.apache.org/t/intrinsic-function-with-constant-input-array/9955/3) to respond.

You are receiving this because you enabled mailing list mode.

To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/90770212b05093ec7fe6d16a2d868d56373d93c4b8e75ff2f93cb30e698b0a0e).

[Apache TVM Discuss] [Questions] Intrinsic Function with Constant Input Array

[tqchen via Apache TVM Discuss <no...@discuss.tvm.ai>]

right now constant arrays are still being processed as normal array in memory. this requirement is due to the fact that the intrinsic function have to follow C API convention. If you want to inline a constant array(that is really small), you can attempt to use select to dispatch each index to the value.

For larger constant array, all systems will have to implement them as in-memory arrays





---
[Visit Topic](https://discuss.tvm.apache.org/t/intrinsic-function-with-constant-input-array/9955/2) to respond.

You are receiving this because you enabled mailing list mode.

To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/53c421f22f1525e0223e16da6af94306c1b8de689ba63cacaf02457a104f3439).