[SVE] Change the dtype of Ramp and Broadcast lanes to PrimExpr by ekalda · Pull Request #16523 · apache/tvm

ekalda · 2024-02-05T16:07:26Z

This change will allow us to express scalable vectors through Ramp and Broadcast nodes, e.g.

vec = tvm.tir.expr.Ramp(0, 1, 4 * tvm.tir.vscale())

We will use negative values for runtime::DataType the encode the scalable lane values, e.g. the above example would result in lanes = -4. That's because the lanes in runtime::DataType are tied to DLPack standard which uses uint16_t for the lanes. The conversion happens in the node definitions and runtime::DataType, so the int and uint16_t values should never be exposed to the API user, especially after the string support has been added.

Also include the TVMScript support for scalable Ramp and Broadcasts.

Note that this patch doesn't include lowering to the appropriate LLVM vectors, support for data type string representation or LoopVectorizer support. All of these will be part of future patches.

ekalda · 2024-02-05T16:32:45Z

cc @neildhickey @lhutton1 @tqchen @cbalint13 @Lunderberg @Anndrey24

ekalda · 2024-02-06T08:49:07Z

@tvm-bot rerun

lhutton1

Thanks for the great work @ekalda! I had a look over and noticed a few nitpicks :)

tqchen · 2024-02-06T14:34:11Z

if it is not high effort, consider add https://github.com/apache/tvm/blob/main/python/tvm/ir/json_compact.py so previously serialized node can be loaded

tqchen · 2024-02-06T14:37:16Z

 */
 inline int GetVectorBytes(DataType dtype) {
+ if (dtype.is_scalable()) {
+ LOG(FATAL) << "Cannot get vector bytes of scalable vector";


runtime Data type cannot be scalable vector

ekalda · 2024-02-13T16:20:48Z

Thanks @tqchen, @Lunderberg and @lhutton1 for your feedback, I uploaded a reworked version of the patch. Here's what's changed:

Separation of vscale multiplier and fixed length vector lanes APIs in runtime::DataType - now we access these constants via vscale_factor() and lanes() methods
is_vector() is retierd now and replaced with is_scalable_vector(), is_fixed_length_vector() and is_scalable_or_fixed_length_vector()
Refactor of the function that extracts the integer multiplier from a lanes expression as per @Lunderberg's suggestions
Removed the ScalableLanes function. Actually, the new form of ExtractVscaleFactor does the job there, so I used that. I reckon that
```
if (!arith::ExtractVscaleFactor(lanes.Eval())):
... 
```
reads somewhat weird, so LMK if you think it would be better to wrap it into something more self-documenting.
Now that an attempt to fetch lanes() on a scalable vector results in an error, I decided to change the pattern in codegens
```
ICHECK(!op->dtype.is_scalable()) << "Scalable vectors are not supported in codegen_c_host";
int lanes = static_cast<int>(Downcast<IntImm>(op->lanes)->value);
```
into
```
int lanes = op->dtype.lanes();
```
Which essentially pushes the error into runtime::Datatype. This has the advantage of reducing the logic in codegens that is not really related to these codegens.
Implemented JSON serialisation support such that graphs serialised with older versions of TVM can be correctly loaded in versions that include the changes is this patch. Unfortunately, in case of a strategic choice of lanes value, a graph serialised with an older version of TVM can be loaded as an incorrect graph without triggrering an error that would then trigger the upgrade_json function, so now we have to force the json upgrade every time we try to load a serialised graph.

…imExpr This change will allow us to express scalable vectors through Ramp and Broadcast nodes, e.g. ``` vec = tvm.tir.expr.Ramp(0, 1, 4 * tvm.tir.vscale()) ``` We will use negative values for `runtime::DataType` the encode the scalable lane values, e.g. the above example would result in `lanes` = -4. That's because the lanes in `runtime::DataType` are tied to DLPack standard which uses `uint16_t` for the lanes. The conversion happens in the node definitions and `runtime::DataType`, so the `int` and `uint16_t` values should never be exposed to the API user, especially after the string support has been added. Also include the TVMScript support for scalable Ramp and Broadcasts. Note that this patch doesn't include lowering to the appropriate LLVM vectors, support for data type string representation or `LoopVectorizer` support. All of these will be part of future patches. Co-authored-by: Luke Hutton <luke.hutton@arm.com> Co-authored-by: Neil Hickey <neil.hickey@arm.com>

Change-Id: I8eb77ce5632359b6e4a2e63c4da490e3abab3ee9

* Separate APIs for fixed length and scalable vectors * Improve the function that extracts the multiplier from scalable lanes expression * Update json_compact.py

Fix failures in test_arith_intset.py and test_tvmscript_printer_tir.py

lhutton1 · 2024-02-20T09:13:45Z

Thanks @ekalda @tqchen @Lunderberg

ekalda · 2024-02-05T16:07:26Z

This change will allow us to express scalable vectors through Ramp and Broadcast nodes, e.g.

vec = tvm.tir.expr.Ramp(0, 1, 4 * tvm.tir.vscale())

We will use negative values for runtime::DataType the encode the scalable lane values, e.g. the above example would result in lanes = -4. That's because the lanes in runtime::DataType are tied to DLPack standard which uses uint16_t for the lanes. The conversion happens in the node definitions and runtime::DataType, so the int and uint16_t values should never be exposed to the API user, especially after the string support has been added.

Also include the TVMScript support for scalable Ramp and Broadcasts.

Note that this patch doesn't include lowering to the appropriate LLVM vectors, support for data type string representation or LoopVectorizer support. All of these will be part of future patches.

ekalda · 2024-02-05T16:32:45Z

cc @neildhickey @lhutton1 @tqchen @cbalint13 @Lunderberg @Anndrey24

ekalda · 2024-02-06T08:49:07Z

@tvm-bot rerun

lhutton1

Thanks for the great work @ekalda! I had a look over and noticed a few nitpicks :)

tqchen · 2024-02-06T14:34:11Z

if it is not high effort, consider add https://github.com/apache/tvm/blob/main/python/tvm/ir/json_compact.py so previously serialized node can be loaded

tqchen · 2024-02-06T14:37:16Z

 */
 inline int GetVectorBytes(DataType dtype) {
+ if (dtype.is_scalable()) {
+ LOG(FATAL) << "Cannot get vector bytes of scalable vector";


runtime Data type cannot be scalable vector

ekalda · 2024-02-13T16:20:48Z

Thanks @tqchen, @Lunderberg and @lhutton1 for your feedback, I uploaded a reworked version of the patch. Here's what's changed:

Separation of vscale multiplier and fixed length vector lanes APIs in runtime::DataType - now we access these constants via vscale_factor() and lanes() methods
is_vector() is retierd now and replaced with is_scalable_vector(), is_fixed_length_vector() and is_scalable_or_fixed_length_vector()
Refactor of the function that extracts the integer multiplier from a lanes expression as per @Lunderberg's suggestions
Removed the ScalableLanes function. Actually, the new form of ExtractVscaleFactor does the job there, so I used that. I reckon that
```
if (!arith::ExtractVscaleFactor(lanes.Eval())):
... 
```
reads somewhat weird, so LMK if you think it would be better to wrap it into something more self-documenting.
Now that an attempt to fetch lanes() on a scalable vector results in an error, I decided to change the pattern in codegens
```
ICHECK(!op->dtype.is_scalable()) << "Scalable vectors are not supported in codegen_c_host";
int lanes = static_cast<int>(Downcast<IntImm>(op->lanes)->value);
```
into
```
int lanes = op->dtype.lanes();
```
Which essentially pushes the error into runtime::Datatype. This has the advantage of reducing the logic in codegens that is not really related to these codegens.
Implemented JSON serialisation support such that graphs serialised with older versions of TVM can be correctly loaded in versions that include the changes is this patch. Unfortunately, in case of a strategic choice of lanes value, a graph serialised with an older version of TVM can be loaded as an incorrect graph without triggrering an error that would then trigger the upgrade_json function, so now we have to force the json upgrade every time we try to load a serialised graph.

…imExpr This change will allow us to express scalable vectors through Ramp and Broadcast nodes, e.g. ``` vec = tvm.tir.expr.Ramp(0, 1, 4 * tvm.tir.vscale()) ``` We will use negative values for `runtime::DataType` the encode the scalable lane values, e.g. the above example would result in `lanes` = -4. That's because the lanes in `runtime::DataType` are tied to DLPack standard which uses `uint16_t` for the lanes. The conversion happens in the node definitions and `runtime::DataType`, so the `int` and `uint16_t` values should never be exposed to the API user, especially after the string support has been added. Also include the TVMScript support for scalable Ramp and Broadcasts. Note that this patch doesn't include lowering to the appropriate LLVM vectors, support for data type string representation or `LoopVectorizer` support. All of these will be part of future patches. Co-authored-by: Luke Hutton <luke.hutton@arm.com> Co-authored-by: Neil Hickey <neil.hickey@arm.com>

Change-Id: I8eb77ce5632359b6e4a2e63c4da490e3abab3ee9

* Separate APIs for fixed length and scalable vectors * Improve the function that extracts the multiplier from scalable lanes expression * Update json_compact.py

Fix failures in test_arith_intset.py and test_tvmscript_printer_tir.py

lhutton1 · 2024-02-20T09:13:45Z

Thanks @ekalda @tqchen @Lunderberg

ekalda · 2024-02-05T16:07:26Z

This change will allow us to express scalable vectors through Ramp and Broadcast nodes, e.g.

vec = tvm.tir.expr.Ramp(0, 1, 4 * tvm.tir.vscale())

We will use negative values for runtime::DataType the encode the scalable lane values, e.g. the above example would result in lanes = -4. That's because the lanes in runtime::DataType are tied to DLPack standard which uses uint16_t for the lanes. The conversion happens in the node definitions and runtime::DataType, so the int and uint16_t values should never be exposed to the API user, especially after the string support has been added.

Also include the TVMScript support for scalable Ramp and Broadcasts.

Note that this patch doesn't include lowering to the appropriate LLVM vectors, support for data type string representation or LoopVectorizer support. All of these will be part of future patches.

ekalda · 2024-02-05T16:32:45Z

cc @neildhickey @lhutton1 @tqchen @cbalint13 @Lunderberg @Anndrey24

ekalda · 2024-02-06T08:49:07Z

@tvm-bot rerun

lhutton1

Thanks for the great work @ekalda! I had a look over and noticed a few nitpicks :)

tqchen · 2024-02-06T14:34:11Z

if it is not high effort, consider add https://github.com/apache/tvm/blob/main/python/tvm/ir/json_compact.py so previously serialized node can be loaded

tqchen · 2024-02-06T14:37:16Z

 */
 inline int GetVectorBytes(DataType dtype) {
+ if (dtype.is_scalable()) {
+ LOG(FATAL) << "Cannot get vector bytes of scalable vector";


runtime Data type cannot be scalable vector

ekalda · 2024-02-13T16:20:48Z

Thanks @tqchen, @Lunderberg and @lhutton1 for your feedback, I uploaded a reworked version of the patch. Here's what's changed:

Separation of vscale multiplier and fixed length vector lanes APIs in runtime::DataType - now we access these constants via vscale_factor() and lanes() methods
is_vector() is retierd now and replaced with is_scalable_vector(), is_fixed_length_vector() and is_scalable_or_fixed_length_vector()
Refactor of the function that extracts the integer multiplier from a lanes expression as per @Lunderberg's suggestions
Removed the ScalableLanes function. Actually, the new form of ExtractVscaleFactor does the job there, so I used that. I reckon that
```
if (!arith::ExtractVscaleFactor(lanes.Eval())):
... 
```
reads somewhat weird, so LMK if you think it would be better to wrap it into something more self-documenting.
Now that an attempt to fetch lanes() on a scalable vector results in an error, I decided to change the pattern in codegens
```
ICHECK(!op->dtype.is_scalable()) << "Scalable vectors are not supported in codegen_c_host";
int lanes = static_cast<int>(Downcast<IntImm>(op->lanes)->value);
```
into
```
int lanes = op->dtype.lanes();
```
Which essentially pushes the error into runtime::Datatype. This has the advantage of reducing the logic in codegens that is not really related to these codegens.
Implemented JSON serialisation support such that graphs serialised with older versions of TVM can be correctly loaded in versions that include the changes is this patch. Unfortunately, in case of a strategic choice of lanes value, a graph serialised with an older version of TVM can be loaded as an incorrect graph without triggrering an error that would then trigger the upgrade_json function, so now we have to force the json upgrade every time we try to load a serialised graph.

…imExpr This change will allow us to express scalable vectors through Ramp and Broadcast nodes, e.g. ``` vec = tvm.tir.expr.Ramp(0, 1, 4 * tvm.tir.vscale()) ``` We will use negative values for `runtime::DataType` the encode the scalable lane values, e.g. the above example would result in `lanes` = -4. That's because the lanes in `runtime::DataType` are tied to DLPack standard which uses `uint16_t` for the lanes. The conversion happens in the node definitions and `runtime::DataType`, so the `int` and `uint16_t` values should never be exposed to the API user, especially after the string support has been added. Also include the TVMScript support for scalable Ramp and Broadcasts. Note that this patch doesn't include lowering to the appropriate LLVM vectors, support for data type string representation or `LoopVectorizer` support. All of these will be part of future patches. Co-authored-by: Luke Hutton <luke.hutton@arm.com> Co-authored-by: Neil Hickey <neil.hickey@arm.com>

Change-Id: I8eb77ce5632359b6e4a2e63c4da490e3abab3ee9

* Separate APIs for fixed length and scalable vectors * Improve the function that extracts the multiplier from scalable lanes expression * Update json_compact.py

Fix failures in test_arith_intset.py and test_tvmscript_printer_tir.py

lhutton1 · 2024-02-20T09:13:45Z

Thanks @ekalda @tqchen @Lunderberg

ekalda · 2024-02-05T16:07:26Z

This change will allow us to express scalable vectors through Ramp and Broadcast nodes, e.g.

vec = tvm.tir.expr.Ramp(0, 1, 4 * tvm.tir.vscale())

We will use negative values for runtime::DataType the encode the scalable lane values, e.g. the above example would result in lanes = -4. That's because the lanes in runtime::DataType are tied to DLPack standard which uses uint16_t for the lanes. The conversion happens in the node definitions and runtime::DataType, so the int and uint16_t values should never be exposed to the API user, especially after the string support has been added.

Also include the TVMScript support for scalable Ramp and Broadcasts.

Note that this patch doesn't include lowering to the appropriate LLVM vectors, support for data type string representation or LoopVectorizer support. All of these will be part of future patches.

ekalda · 2024-02-05T16:32:45Z

cc @neildhickey @lhutton1 @tqchen @cbalint13 @Lunderberg @Anndrey24

ekalda · 2024-02-06T08:49:07Z

@tvm-bot rerun

lhutton1

Thanks for the great work @ekalda! I had a look over and noticed a few nitpicks :)

tqchen · 2024-02-06T14:34:11Z

if it is not high effort, consider add https://github.com/apache/tvm/blob/main/python/tvm/ir/json_compact.py so previously serialized node can be loaded

tqchen · 2024-02-06T14:37:16Z

 */
 inline int GetVectorBytes(DataType dtype) {
+ if (dtype.is_scalable()) {
+ LOG(FATAL) << "Cannot get vector bytes of scalable vector";


runtime Data type cannot be scalable vector

ekalda · 2024-02-13T16:20:48Z

Thanks @tqchen, @Lunderberg and @lhutton1 for your feedback, I uploaded a reworked version of the patch. Here's what's changed:

Separation of vscale multiplier and fixed length vector lanes APIs in runtime::DataType - now we access these constants via vscale_factor() and lanes() methods
is_vector() is retierd now and replaced with is_scalable_vector(), is_fixed_length_vector() and is_scalable_or_fixed_length_vector()
Refactor of the function that extracts the integer multiplier from a lanes expression as per @Lunderberg's suggestions
Removed the ScalableLanes function. Actually, the new form of ExtractVscaleFactor does the job there, so I used that. I reckon that
```
if (!arith::ExtractVscaleFactor(lanes.Eval())):
... 
```
reads somewhat weird, so LMK if you think it would be better to wrap it into something more self-documenting.
Now that an attempt to fetch lanes() on a scalable vector results in an error, I decided to change the pattern in codegens
```
ICHECK(!op->dtype.is_scalable()) << "Scalable vectors are not supported in codegen_c_host";
int lanes = static_cast<int>(Downcast<IntImm>(op->lanes)->value);
```
into
```
int lanes = op->dtype.lanes();
```
Which essentially pushes the error into runtime::Datatype. This has the advantage of reducing the logic in codegens that is not really related to these codegens.
Implemented JSON serialisation support such that graphs serialised with older versions of TVM can be correctly loaded in versions that include the changes is this patch. Unfortunately, in case of a strategic choice of lanes value, a graph serialised with an older version of TVM can be loaded as an incorrect graph without triggrering an error that would then trigger the upgrade_json function, so now we have to force the json upgrade every time we try to load a serialised graph.

…imExpr This change will allow us to express scalable vectors through Ramp and Broadcast nodes, e.g. ``` vec = tvm.tir.expr.Ramp(0, 1, 4 * tvm.tir.vscale()) ``` We will use negative values for `runtime::DataType` the encode the scalable lane values, e.g. the above example would result in `lanes` = -4. That's because the lanes in `runtime::DataType` are tied to DLPack standard which uses `uint16_t` for the lanes. The conversion happens in the node definitions and `runtime::DataType`, so the `int` and `uint16_t` values should never be exposed to the API user, especially after the string support has been added. Also include the TVMScript support for scalable Ramp and Broadcasts. Note that this patch doesn't include lowering to the appropriate LLVM vectors, support for data type string representation or `LoopVectorizer` support. All of these will be part of future patches. Co-authored-by: Luke Hutton <luke.hutton@arm.com> Co-authored-by: Neil Hickey <neil.hickey@arm.com>

Change-Id: I8eb77ce5632359b6e4a2e63c4da490e3abab3ee9

* Separate APIs for fixed length and scalable vectors * Improve the function that extracts the multiplier from scalable lanes expression * Update json_compact.py

Fix failures in test_arith_intset.py and test_tvmscript_printer_tir.py

lhutton1 · 2024-02-20T09:13:45Z

Thanks @ekalda @tqchen @Lunderberg

ekalda · 2024-02-05T16:07:26Z

This change will allow us to express scalable vectors through Ramp and Broadcast nodes, e.g.

vec = tvm.tir.expr.Ramp(0, 1, 4 * tvm.tir.vscale())

We will use negative values for runtime::DataType the encode the scalable lane values, e.g. the above example would result in lanes = -4. That's because the lanes in runtime::DataType are tied to DLPack standard which uses uint16_t for the lanes. The conversion happens in the node definitions and runtime::DataType, so the int and uint16_t values should never be exposed to the API user, especially after the string support has been added.

Also include the TVMScript support for scalable Ramp and Broadcasts.

Note that this patch doesn't include lowering to the appropriate LLVM vectors, support for data type string representation or LoopVectorizer support. All of these will be part of future patches.

ekalda · 2024-02-05T16:32:45Z

cc @neildhickey @lhutton1 @tqchen @cbalint13 @Lunderberg @Anndrey24

ekalda · 2024-02-06T08:49:07Z

@tvm-bot rerun

lhutton1

Thanks for the great work @ekalda! I had a look over and noticed a few nitpicks :)

tqchen · 2024-02-06T14:34:11Z

if it is not high effort, consider add https://github.com/apache/tvm/blob/main/python/tvm/ir/json_compact.py so previously serialized node can be loaded

tqchen · 2024-02-06T14:37:16Z

 */
 inline int GetVectorBytes(DataType dtype) {
+ if (dtype.is_scalable()) {
+ LOG(FATAL) << "Cannot get vector bytes of scalable vector";


runtime Data type cannot be scalable vector

ekalda · 2024-02-13T16:20:48Z

Thanks @tqchen, @Lunderberg and @lhutton1 for your feedback, I uploaded a reworked version of the patch. Here's what's changed:

Separation of vscale multiplier and fixed length vector lanes APIs in runtime::DataType - now we access these constants via vscale_factor() and lanes() methods
is_vector() is retierd now and replaced with is_scalable_vector(), is_fixed_length_vector() and is_scalable_or_fixed_length_vector()
Refactor of the function that extracts the integer multiplier from a lanes expression as per @Lunderberg's suggestions
Removed the ScalableLanes function. Actually, the new form of ExtractVscaleFactor does the job there, so I used that. I reckon that
```
if (!arith::ExtractVscaleFactor(lanes.Eval())):
... 
```
reads somewhat weird, so LMK if you think it would be better to wrap it into something more self-documenting.
Now that an attempt to fetch lanes() on a scalable vector results in an error, I decided to change the pattern in codegens
```
ICHECK(!op->dtype.is_scalable()) << "Scalable vectors are not supported in codegen_c_host";
int lanes = static_cast<int>(Downcast<IntImm>(op->lanes)->value);
```
into
```
int lanes = op->dtype.lanes();
```
Which essentially pushes the error into runtime::Datatype. This has the advantage of reducing the logic in codegens that is not really related to these codegens.
Implemented JSON serialisation support such that graphs serialised with older versions of TVM can be correctly loaded in versions that include the changes is this patch. Unfortunately, in case of a strategic choice of lanes value, a graph serialised with an older version of TVM can be loaded as an incorrect graph without triggrering an error that would then trigger the upgrade_json function, so now we have to force the json upgrade every time we try to load a serialised graph.

…imExpr This change will allow us to express scalable vectors through Ramp and Broadcast nodes, e.g. ``` vec = tvm.tir.expr.Ramp(0, 1, 4 * tvm.tir.vscale()) ``` We will use negative values for `runtime::DataType` the encode the scalable lane values, e.g. the above example would result in `lanes` = -4. That's because the lanes in `runtime::DataType` are tied to DLPack standard which uses `uint16_t` for the lanes. The conversion happens in the node definitions and `runtime::DataType`, so the `int` and `uint16_t` values should never be exposed to the API user, especially after the string support has been added. Also include the TVMScript support for scalable Ramp and Broadcasts. Note that this patch doesn't include lowering to the appropriate LLVM vectors, support for data type string representation or `LoopVectorizer` support. All of these will be part of future patches. Co-authored-by: Luke Hutton <luke.hutton@arm.com> Co-authored-by: Neil Hickey <neil.hickey@arm.com>

Change-Id: I8eb77ce5632359b6e4a2e63c4da490e3abab3ee9

* Separate APIs for fixed length and scalable vectors * Improve the function that extracts the multiplier from scalable lanes expression * Update json_compact.py

Fix failures in test_arith_intset.py and test_tvmscript_printer_tir.py

lhutton1 · 2024-02-20T09:13:45Z

Thanks @ekalda @tqchen @Lunderberg

ekalda · 2024-02-05T16:07:26Z

This change will allow us to express scalable vectors through Ramp and Broadcast nodes, e.g.

vec = tvm.tir.expr.Ramp(0, 1, 4 * tvm.tir.vscale())

We will use negative values for runtime::DataType the encode the scalable lane values, e.g. the above example would result in lanes = -4. That's because the lanes in runtime::DataType are tied to DLPack standard which uses uint16_t for the lanes. The conversion happens in the node definitions and runtime::DataType, so the int and uint16_t values should never be exposed to the API user, especially after the string support has been added.

Also include the TVMScript support for scalable Ramp and Broadcasts.

Note that this patch doesn't include lowering to the appropriate LLVM vectors, support for data type string representation or LoopVectorizer support. All of these will be part of future patches.

ekalda · 2024-02-05T16:32:45Z

cc @neildhickey @lhutton1 @tqchen @cbalint13 @Lunderberg @Anndrey24

ekalda · 2024-02-06T08:49:07Z

@tvm-bot rerun

lhutton1

Thanks for the great work @ekalda! I had a look over and noticed a few nitpicks :)

tqchen · 2024-02-06T14:34:11Z

if it is not high effort, consider add https://github.com/apache/tvm/blob/main/python/tvm/ir/json_compact.py so previously serialized node can be loaded

tqchen · 2024-02-06T14:37:16Z

 */
 inline int GetVectorBytes(DataType dtype) {
+ if (dtype.is_scalable()) {
+ LOG(FATAL) << "Cannot get vector bytes of scalable vector";


runtime Data type cannot be scalable vector

ekalda · 2024-02-13T16:20:48Z

Thanks @tqchen, @Lunderberg and @lhutton1 for your feedback, I uploaded a reworked version of the patch. Here's what's changed:

Separation of vscale multiplier and fixed length vector lanes APIs in runtime::DataType - now we access these constants via vscale_factor() and lanes() methods
is_vector() is retierd now and replaced with is_scalable_vector(), is_fixed_length_vector() and is_scalable_or_fixed_length_vector()
Refactor of the function that extracts the integer multiplier from a lanes expression as per @Lunderberg's suggestions
Removed the ScalableLanes function. Actually, the new form of ExtractVscaleFactor does the job there, so I used that. I reckon that
```
if (!arith::ExtractVscaleFactor(lanes.Eval())):
... 
```
reads somewhat weird, so LMK if you think it would be better to wrap it into something more self-documenting.
Now that an attempt to fetch lanes() on a scalable vector results in an error, I decided to change the pattern in codegens
```
ICHECK(!op->dtype.is_scalable()) << "Scalable vectors are not supported in codegen_c_host";
int lanes = static_cast<int>(Downcast<IntImm>(op->lanes)->value);
```
into
```
int lanes = op->dtype.lanes();
```
Which essentially pushes the error into runtime::Datatype. This has the advantage of reducing the logic in codegens that is not really related to these codegens.
Implemented JSON serialisation support such that graphs serialised with older versions of TVM can be correctly loaded in versions that include the changes is this patch. Unfortunately, in case of a strategic choice of lanes value, a graph serialised with an older version of TVM can be loaded as an incorrect graph without triggrering an error that would then trigger the upgrade_json function, so now we have to force the json upgrade every time we try to load a serialised graph.

…imExpr This change will allow us to express scalable vectors through Ramp and Broadcast nodes, e.g. ``` vec = tvm.tir.expr.Ramp(0, 1, 4 * tvm.tir.vscale()) ``` We will use negative values for `runtime::DataType` the encode the scalable lane values, e.g. the above example would result in `lanes` = -4. That's because the lanes in `runtime::DataType` are tied to DLPack standard which uses `uint16_t` for the lanes. The conversion happens in the node definitions and `runtime::DataType`, so the `int` and `uint16_t` values should never be exposed to the API user, especially after the string support has been added. Also include the TVMScript support for scalable Ramp and Broadcasts. Note that this patch doesn't include lowering to the appropriate LLVM vectors, support for data type string representation or `LoopVectorizer` support. All of these will be part of future patches. Co-authored-by: Luke Hutton <luke.hutton@arm.com> Co-authored-by: Neil Hickey <neil.hickey@arm.com>

Change-Id: I8eb77ce5632359b6e4a2e63c4da490e3abab3ee9

* Separate APIs for fixed length and scalable vectors * Improve the function that extracts the multiplier from scalable lanes expression * Update json_compact.py

Fix failures in test_arith_intset.py and test_tvmscript_printer_tir.py

lhutton1 · 2024-02-20T09:13:45Z

Thanks @ekalda @tqchen @Lunderberg

ekalda · 2024-02-05T16:07:26Z

This change will allow us to express scalable vectors through Ramp and Broadcast nodes, e.g.

vec = tvm.tir.expr.Ramp(0, 1, 4 * tvm.tir.vscale())

We will use negative values for runtime::DataType the encode the scalable lane values, e.g. the above example would result in lanes = -4. That's because the lanes in runtime::DataType are tied to DLPack standard which uses uint16_t for the lanes. The conversion happens in the node definitions and runtime::DataType, so the int and uint16_t values should never be exposed to the API user, especially after the string support has been added.

Also include the TVMScript support for scalable Ramp and Broadcasts.

Note that this patch doesn't include lowering to the appropriate LLVM vectors, support for data type string representation or LoopVectorizer support. All of these will be part of future patches.

ekalda · 2024-02-05T16:32:45Z

cc @neildhickey @lhutton1 @tqchen @cbalint13 @Lunderberg @Anndrey24

ekalda · 2024-02-06T08:49:07Z

@tvm-bot rerun

lhutton1

Thanks for the great work @ekalda! I had a look over and noticed a few nitpicks :)

tqchen · 2024-02-06T14:34:11Z

if it is not high effort, consider add https://github.com/apache/tvm/blob/main/python/tvm/ir/json_compact.py so previously serialized node can be loaded

tqchen · 2024-02-06T14:37:16Z

 */
 inline int GetVectorBytes(DataType dtype) {
+ if (dtype.is_scalable()) {
+ LOG(FATAL) << "Cannot get vector bytes of scalable vector";


runtime Data type cannot be scalable vector

ekalda · 2024-02-13T16:20:48Z

Thanks @tqchen, @Lunderberg and @lhutton1 for your feedback, I uploaded a reworked version of the patch. Here's what's changed:

Separation of vscale multiplier and fixed length vector lanes APIs in runtime::DataType - now we access these constants via vscale_factor() and lanes() methods
is_vector() is retierd now and replaced with is_scalable_vector(), is_fixed_length_vector() and is_scalable_or_fixed_length_vector()
Refactor of the function that extracts the integer multiplier from a lanes expression as per @Lunderberg's suggestions
Removed the ScalableLanes function. Actually, the new form of ExtractVscaleFactor does the job there, so I used that. I reckon that
```
if (!arith::ExtractVscaleFactor(lanes.Eval())):
... 
```
reads somewhat weird, so LMK if you think it would be better to wrap it into something more self-documenting.
Now that an attempt to fetch lanes() on a scalable vector results in an error, I decided to change the pattern in codegens
```
ICHECK(!op->dtype.is_scalable()) << "Scalable vectors are not supported in codegen_c_host";
int lanes = static_cast<int>(Downcast<IntImm>(op->lanes)->value);
```
into
```
int lanes = op->dtype.lanes();
```
Which essentially pushes the error into runtime::Datatype. This has the advantage of reducing the logic in codegens that is not really related to these codegens.
Implemented JSON serialisation support such that graphs serialised with older versions of TVM can be correctly loaded in versions that include the changes is this patch. Unfortunately, in case of a strategic choice of lanes value, a graph serialised with an older version of TVM can be loaded as an incorrect graph without triggrering an error that would then trigger the upgrade_json function, so now we have to force the json upgrade every time we try to load a serialised graph.

…imExpr This change will allow us to express scalable vectors through Ramp and Broadcast nodes, e.g. ``` vec = tvm.tir.expr.Ramp(0, 1, 4 * tvm.tir.vscale()) ``` We will use negative values for `runtime::DataType` the encode the scalable lane values, e.g. the above example would result in `lanes` = -4. That's because the lanes in `runtime::DataType` are tied to DLPack standard which uses `uint16_t` for the lanes. The conversion happens in the node definitions and `runtime::DataType`, so the `int` and `uint16_t` values should never be exposed to the API user, especially after the string support has been added. Also include the TVMScript support for scalable Ramp and Broadcasts. Note that this patch doesn't include lowering to the appropriate LLVM vectors, support for data type string representation or `LoopVectorizer` support. All of these will be part of future patches. Co-authored-by: Luke Hutton <luke.hutton@arm.com> Co-authored-by: Neil Hickey <neil.hickey@arm.com>

Change-Id: I8eb77ce5632359b6e4a2e63c4da490e3abab3ee9

* Separate APIs for fixed length and scalable vectors * Improve the function that extracts the multiplier from scalable lanes expression * Update json_compact.py

Fix failures in test_arith_intset.py and test_tvmscript_printer_tir.py

lhutton1 · 2024-02-20T09:13:45Z

Thanks @ekalda @tqchen @Lunderberg

ekalda · 2024-02-05T16:07:26Z

This change will allow us to express scalable vectors through Ramp and Broadcast nodes, e.g.

vec = tvm.tir.expr.Ramp(0, 1, 4 * tvm.tir.vscale())

We will use negative values for runtime::DataType the encode the scalable lane values, e.g. the above example would result in lanes = -4. That's because the lanes in runtime::DataType are tied to DLPack standard which uses uint16_t for the lanes. The conversion happens in the node definitions and runtime::DataType, so the int and uint16_t values should never be exposed to the API user, especially after the string support has been added.

Also include the TVMScript support for scalable Ramp and Broadcasts.

Note that this patch doesn't include lowering to the appropriate LLVM vectors, support for data type string representation or LoopVectorizer support. All of these will be part of future patches.

ekalda · 2024-02-05T16:32:45Z

cc @neildhickey @lhutton1 @tqchen @cbalint13 @Lunderberg @Anndrey24

ekalda · 2024-02-06T08:49:07Z

@tvm-bot rerun

lhutton1

Thanks for the great work @ekalda! I had a look over and noticed a few nitpicks :)

tqchen · 2024-02-06T14:34:11Z

if it is not high effort, consider add https://github.com/apache/tvm/blob/main/python/tvm/ir/json_compact.py so previously serialized node can be loaded

tqchen · 2024-02-06T14:37:16Z

 */
 inline int GetVectorBytes(DataType dtype) {
+ if (dtype.is_scalable()) {
+ LOG(FATAL) << "Cannot get vector bytes of scalable vector";


runtime Data type cannot be scalable vector

ekalda · 2024-02-13T16:20:48Z

Thanks @tqchen, @Lunderberg and @lhutton1 for your feedback, I uploaded a reworked version of the patch. Here's what's changed:

Separation of vscale multiplier and fixed length vector lanes APIs in runtime::DataType - now we access these constants via vscale_factor() and lanes() methods
is_vector() is retierd now and replaced with is_scalable_vector(), is_fixed_length_vector() and is_scalable_or_fixed_length_vector()
Refactor of the function that extracts the integer multiplier from a lanes expression as per @Lunderberg's suggestions
Removed the ScalableLanes function. Actually, the new form of ExtractVscaleFactor does the job there, so I used that. I reckon that
```
if (!arith::ExtractVscaleFactor(lanes.Eval())):
... 
```
reads somewhat weird, so LMK if you think it would be better to wrap it into something more self-documenting.
Now that an attempt to fetch lanes() on a scalable vector results in an error, I decided to change the pattern in codegens
```
ICHECK(!op->dtype.is_scalable()) << "Scalable vectors are not supported in codegen_c_host";
int lanes = static_cast<int>(Downcast<IntImm>(op->lanes)->value);
```
into
```
int lanes = op->dtype.lanes();
```
Which essentially pushes the error into runtime::Datatype. This has the advantage of reducing the logic in codegens that is not really related to these codegens.
Implemented JSON serialisation support such that graphs serialised with older versions of TVM can be correctly loaded in versions that include the changes is this patch. Unfortunately, in case of a strategic choice of lanes value, a graph serialised with an older version of TVM can be loaded as an incorrect graph without triggrering an error that would then trigger the upgrade_json function, so now we have to force the json upgrade every time we try to load a serialised graph.

…imExpr This change will allow us to express scalable vectors through Ramp and Broadcast nodes, e.g. ``` vec = tvm.tir.expr.Ramp(0, 1, 4 * tvm.tir.vscale()) ``` We will use negative values for `runtime::DataType` the encode the scalable lane values, e.g. the above example would result in `lanes` = -4. That's because the lanes in `runtime::DataType` are tied to DLPack standard which uses `uint16_t` for the lanes. The conversion happens in the node definitions and `runtime::DataType`, so the `int` and `uint16_t` values should never be exposed to the API user, especially after the string support has been added. Also include the TVMScript support for scalable Ramp and Broadcasts. Note that this patch doesn't include lowering to the appropriate LLVM vectors, support for data type string representation or `LoopVectorizer` support. All of these will be part of future patches. Co-authored-by: Luke Hutton <luke.hutton@arm.com> Co-authored-by: Neil Hickey <neil.hickey@arm.com>

Change-Id: I8eb77ce5632359b6e4a2e63c4da490e3abab3ee9

* Separate APIs for fixed length and scalable vectors * Improve the function that extracts the multiplier from scalable lanes expression * Update json_compact.py

Fix failures in test_arith_intset.py and test_tvmscript_printer_tir.py

lhutton1 · 2024-02-20T09:13:45Z

Thanks @ekalda @tqchen @Lunderberg

ekalda mentioned this pull request Feb 5, 2024
[Tracking Issue] Scalable Vector Extension (SVE) upstreaming #16455
Closed
12 tasks

lhutton1 reviewed Feb 6, 2024
View reviewed changes

Lunderberg reviewed Feb 6, 2024
View reviewed changes

Comment threadsrc/arith/scalable_expression.cc Outdated

tqchen requested changes Feb 6, 2024
View reviewed changes

tqchen reviewed Feb 6, 2024
View reviewed changes

Comment threadinclude/tvm/runtime/data_type.h

ekalda force-pushed the p2-scalable-ramps2 branch from cea202e to 888ad3bCompare February 13, 2024 16:17

ekaldaand others added 4 commits February 14, 2024 09:57

Fix the linting
f74c43a
Change-Id: I8eb77ce5632359b6e4a2e63c4da490e3abab3ee9

Respond to review comments
ab3914f
* Separate APIs for fixed length and scalable vectors * Improve the function that extracts the multiplier from scalable lanes expression * Update json_compact.py

Fix test failures
60e7d83
Fix failures in test_arith_intset.py and test_tvmscript_printer_tir.py

ekalda force-pushed the p2-scalable-ramps2 branch from 888ad3b to 60e7d83Compare February 14, 2024 10:25

tqchen approved these changes Feb 19, 2024
View reviewed changes

lhutton1 merged commit a6157a6 into apache:mainFeb 20, 2024

ysh329 mentioned this pull request Apr 21, 2024
[Release] v0.16.0 Release Candidate Notes #16911
Closed

This was referenced Jul 27, 2026
[Metal] Add 16-byte alignment padding to shared/threadgroup memory tile-ai/tilelang#2786
Merged
[TIR][Python] Trim redundant op proxy wrappers tile-ai/tilelang#2850
Merged