You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by "wrongtest-intellif (via GitHub)" <gi...@apache.org> on 2025/11/05 06:55:56 UTC
[PR] [TIR] Add step attribute to ForNode (Initial codes) [tvm]
wrongtest-intellif opened a new pull request, #18421:
URL: https://github.com/apache/tvm/pull/18421
A prototyping change to add `ForNode::step`
It try add minimal codes to run following naive test cases.
```python
import tvm
import numpy as np
from tvm.script import tir as T
@T.prim_func
def function(A: T.Buffer[(1024)], B: T.Buffer[(1024)], C: T.Buffer[(1024)]):
for i in range(0, 100, 3):
C[i] = A[i] + B[i]
print(function)
lib = tvm.compile(function, target="c")
print(lib.mod.inspect_source())
lib2 = tvm.compile(function, target="llvm")
a = np.random.uniform(1, 100, [1024]).astype("float32")
b = np.random.uniform(1, 100, [1024]).astype("float32")
c = np.zeros([1024]).astype("float32")
lib(a, b, c)
c[:] = 0
print(c[:])
lib2(a, b, c)
print(c[:])
```
The aspects to check for a real roadmap may be
1. Roundtrip support for TIR tvmscript grammar
2. Correctness of TIR lowering pipeline
- For **all transformations and analysis tools**, either it make adaptions to non-consecutive loop iteration indices, or loop canonicalization required.
- Ensure the original `ForNode::step` is not dropped by mutations on `ForNode`.
3. Correctness of TensorIR schedule and MetaSchedule
- Since many primitives depend on affine bindings. Loop canonicalization is required.
4. CodeGen support
- Check mainstream targets could support the loop step.
5. Compatibility issues
- Try to argue that the change would not affect existing works, since `ForNode` is an important construction in TVM.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] [TIR] Add step attribute to ForNode (Initial codes) [tvm]
Posted by "gemini-code-assist[bot] (via GitHub)" <gi...@apache.org>.
gemini-code-assist[bot] commented on PR #18421:
URL: https://github.com/apache/tvm/pull/18421#issuecomment-3489649231
## Summary of Changes
Hello @wrongtest-intellif, I'm Gemini Code Assist[^1]! I'm currently reviewing this pull request and will post my feedback shortly. In the meantime, here's a summary to help you and other reviewers quickly get up to speed!
This pull request introduces a `step` attribute to the `ForNode` in TVM's Tensor IR (TIR), enabling the representation and generation of loops with non-unit iteration steps. This foundational change impacts various components, from the core IR definition and Python scripting interface to code generation backends and scheduling primitives, ensuring that loops with custom steps are correctly parsed, lowered, and compiled across TVM.
### Highlights
* **New `step` attribute for `ForNode`**: Introduced an optional `step` attribute to the `ForNode` in TVM's Tensor IR (TIR), allowing loops to iterate with non-unit steps (e.g., `for i in range(0, 100, 3)`).
* **Python API and Scripting Support**: Updated the Python `tvm.script.ir_builder.tir` module to accept `step` arguments for `serial`, `parallel`, `vectorized`, and `unroll` loop types. The `tvm.script.parser.tir` now handles Python's `range` sugar with a `step` argument, mapping it to the new TIR `step` attribute.
* **Code Generation Updates**: Modified various code generation backends (LLVM, C, WebGPU, SPIRV) to correctly interpret and generate code for loops that include a custom `step` value. This ensures that the new loop semantics are respected during compilation.
* **TIR Schedule and Transform Compatibility**: Adjusted numerous TIR scheduling primitives and transformation passes to either propagate the new `step` attribute or utilize the newly introduced `For::ForSimple` constructor, which implicitly assumes a trivial (unit) step. This maintains compatibility and correctness across the TIR pipeline.
* **New Utility and Tests**: Added a `HasTrivialStep()` method to `ForNode` for easier checking of unit steps. New test cases were introduced in `test_tir_nodes.py` and `test_tvmscript_parser_tir.py` to validate the functionality of loops with custom steps.
<details>
<summary><b>Using Gemini Code Assist</b></summary>
<br>
The full guide for Gemini Code Assist can be found on our [documentation page](https://developers.google.com/gemini-code-assist/docs/review-github-code), here are some quick tips.
<b>Invoking Gemini</b>
You can request assistance from Gemini at any point by creating a comment using either `/gemini <command>` or `@gemini-code-assist <command>`. Below is a summary of the supported commands on the current page.
Feature | Command | Description
--- | --- | ---
Code Review | `/gemini review` | Performs a code review for the current pull request in its current state.
Pull Request Summary | `/gemini summary` | Provides a summary of the current pull request in its current state.
Comment | @gemini-code-assist | Responds in comments when explicitly tagged, both in pull request comments and review comments.
Help | `/gemini help` | Displays a list of available commands.
<b>Customization</b>
To customize Gemini Code Assist for GitHub experience, repository maintainers can create a configuration file and/or provide a custom code review style guide (such as PEP-8 for Python) by creating and adding files to a `.gemini/` folder in the base of the repository. Detailed instructions can be found [here](https://developers.google.com/gemini-code-assist/docs/customize-gemini-behavior-github).
<b>Limitations & Feedback</b>
Gemini Code Assist may make mistakes. Please leave feedback on any instances where its feedback is incorrect or counter productive. You can react with :thumbsup: and :thumbsdown: on @gemini-code-assist comments. If you're interested in giving your feedback about your experience with Gemini Code Assist for Github and other Google products, sign up [here](https://google.qualtrics.com/jfe/form/SV_2cyuGuTWsEw84yG).
<b>You can also get AI-powered code generation, chat, as well as code reviews directly in the IDE at no cost with the [Gemini Code Assist IDE Extension](https://cloud.google.com/products/gemini/code-assist).</b>
</details>
[^1]: Review the [Privacy Notices](https://policies.google.com/privacy), [Generative AI Prohibited Use Policy](https://policies.google.com/terms/generative-ai/use-policy), [Terms of Service](https://policies.google.com/terms), and learn how to configure Gemini Code Assist in GitHub [here](https://developers.google.com/gemini-code-assist/docs/customize-gemini-behavior-github). Gemini can make mistakes, so double check it and [use code with caution](https://support.google.com/legal/answer/13505487).
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] [TIR] Add step attribute to ForNode (Initial codes) [tvm]
Posted by "wrongtest-intellif (via GitHub)" <gi...@apache.org>.
wrongtest-intellif commented on PR #18421:
URL: https://github.com/apache/tvm/pull/18421#issuecomment-3541687566
/gemini review
--
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: discuss-archive-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
---------------------------------------------------------------------
To unsubscribe, e-mail: discuss-archive-unsubscribe@tvm.apache.org
For additional commands, e-mail: discuss-archive-help@tvm.apache.org
Re: [PR] [TIR] Add step attribute to ForNode (Initial codes) [tvm]
Posted by "wrongtest-intellif (via GitHub)" <gi...@apache.org>.
wrongtest-intellif commented on code in PR #18421:
URL: https://github.com/apache/tvm/pull/18421#discussion_r2532756406
##########
src/target/source/codegen_c.cc:
##########
@@ -1121,12 +1121,19 @@ void CodeGenC::VisitStmt_(const AssertStmtNode* op) {
void CodeGenC::VisitStmt_(const ForNode* op) {
std::string extent = PrintExpr(op->extent);
+ std::string step = op->step.has_value()? PrintExpr(*op->step) : "";
PrintIndent();
std::string vid = AllocVarID(op->loop_var.get());
ICHECK(is_zero(op->min));
stream << "for (";
PrintType(op->loop_var.dtype(), stream);
- stream << ' ' << vid << " = 0; " << vid << " < " << extent << "; ++" << vid << ") {\n";
+ stream << ' ' << vid << " = 0; " << vid << " < " << extent << "; ";
+ if (step.empty()) {
+ stream << "++" << vid;
+ } else {
+ stream << vid << " += " << step;
+ }
+ stream << ") {\n";
int for_scope = BeginScope();
PrintStmt(op->body);
this->EndScope(for_scope);
Review Comment:
/gemini review
##########
src/target/source/codegen_webgpu.cc:
##########
@@ -668,12 +668,19 @@ void CodeGenWebGPU::VisitStmt_(const AllocateNode* op) {
void CodeGenWebGPU::VisitStmt_(const ForNode* op) {
std::string extent = PrintExpr(op->extent);
+ std::string step = op->step.has_value()? PrintExpr(*op->step) : "";
std::string vid = AllocVarID(op->loop_var.get());
ICHECK(is_zero(op->min));
PrintIndent();
stream << "for (var " << vid << " : ";
PrintType(op->loop_var.dtype(), stream);
- stream << " = 0; " << vid << " < " << extent << "; " << vid << "++) {\n";
+ stream << " = 0; " << vid << " < " << extent << "; " << vid;
+ if (step.empty()) {
+ stream << "++";
+ } else {
+ stream << " += " << step;
+ }
+ stream << ") {\n";
int for_scope = BeginScope();
PrintStmt(op->body);
this->EndScope(for_scope);
Review Comment:
/gemini review
--
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: discuss-archive-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
---------------------------------------------------------------------
To unsubscribe, e-mail: discuss-archive-unsubscribe@tvm.apache.org
For additional commands, e-mail: discuss-archive-help@tvm.apache.org
Re: [PR] [TIR] Add step attribute to ForNode (Initial codes) [tvm]
Posted by "gemini-code-assist[bot] (via GitHub)" <gi...@apache.org>.
gemini-code-assist[bot] commented on code in PR #18421:
URL: https://github.com/apache/tvm/pull/18421#discussion_r2533965816
##########
python/tvm/script/parser/tir/parser.py:
##########
@@ -168,6 +168,28 @@ def find_decorator_annotation(node: doc.FunctionDef, annotation: str, default: b
return default
+def range_sugar(
+ start: PrimExpr,
+ stop: PrimExpr = None,
+ step: Optional[PrimExpr] = None,
+ *,
+ annotations: Dict[str, Any] = None,
+) -> T.frame.ForFrame:
+ """The sugar for python range builtin."""
+
+ # Since `tir.For` do not support reversed iteration semantic,
+ # the step must be checked to be positive integer when use range sugar
+ if step is not None:
+ try:
+ step = int(step)
+ if step < 0:
+ raise ValueError(f"Only support positive step in range(), get {step}")
Review Comment:
Python's `range()` function does not allow a step of zero. It's better to check for this case here and provide a more specific error message, similar to Python's behavior. The `CanonicalizeLoop` pass would catch this later, but failing early with a more informative message would improve user experience.
```suggestion
if step == 0:
raise ValueError("range() arg 3 must not be zero")
if step < 0:
raise ValueError(f"Only support positive step in range(), get {step}")
```
--
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: discuss-archive-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
---------------------------------------------------------------------
To unsubscribe, e-mail: discuss-archive-unsubscribe@tvm.apache.org
For additional commands, e-mail: discuss-archive-help@tvm.apache.org
Re: [PR] [TIR] Add step attribute to ForNode (Initial codes) [tvm]
Posted by "tqchen (via GitHub)" <gi...@apache.org>.
tqchen commented on code in PR #18421:
URL: https://github.com/apache/tvm/pull/18421#discussion_r2531956520
##########
src/target/spirv/codegen_spirv.cc:
##########
@@ -672,10 +672,19 @@ void CodeGenSPIRV::VisitStmt_(const BufferStoreNode* op) {
}
void CodeGenSPIRV::VisitStmt_(const ForNode* op) {
- ICHECK(is_zero(op->min));
analyzer_->Bind(op->loop_var, Range::FromMinExtent(op->min, op->extent));
spirv::Value init_value = MakeValue(op->min);
spirv::Value extent_value = MakeValue(op->extent);
+
+ // loop step
+ spirv::Value step;
+ if (op->HasTrivialStep()) {
+ step = op->loop_var.dtype().is_int() ? builder_->IntImm(loop_var.stype, 1)
+ : builder_->UIntImm(loop_var.stype, 1);
+ } else {
+ step = MakeValue(tir::cast(op->extent->dtype, *op->step));
+ }
Review Comment:
@wrongtest-intellif seems this is a valid comment
--
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: discuss-archive-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
---------------------------------------------------------------------
To unsubscribe, e-mail: discuss-archive-unsubscribe@tvm.apache.org
For additional commands, e-mail: discuss-archive-help@tvm.apache.org
Re: [PR] [TIR] Add step attribute to ForNode (Initial codes) [tvm]
Posted by "tqchen (via GitHub)" <gi...@apache.org>.
tqchen commented on code in PR #18421:
URL: https://github.com/apache/tvm/pull/18421#discussion_r2532522702
##########
src/target/spirv/codegen_spirv.cc:
##########
@@ -672,10 +672,19 @@ void CodeGenSPIRV::VisitStmt_(const BufferStoreNode* op) {
}
void CodeGenSPIRV::VisitStmt_(const ForNode* op) {
- ICHECK(is_zero(op->min));
analyzer_->Bind(op->loop_var, Range::FromMinExtent(op->min, op->extent));
spirv::Value init_value = MakeValue(op->min);
spirv::Value extent_value = MakeValue(op->extent);
+
+ // loop step
+ spirv::Value step;
+ if (op->HasTrivialStep()) {
+ step = op->loop_var.dtype().is_int() ? builder_->IntImm(loop_var.stype, 1)
+ : builder_->UIntImm(loop_var.stype, 1);
+ } else {
+ step = MakeValue(tir::cast(op->extent->dtype, *op->step));
+ }
Review Comment:
you can do `/gemini review` in comment to trigger
--
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: discuss-archive-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
---------------------------------------------------------------------
To unsubscribe, e-mail: discuss-archive-unsubscribe@tvm.apache.org
For additional commands, e-mail: discuss-archive-help@tvm.apache.org
Re: [PR] [TIR] Add step attribute to ForNode (Initial codes) [tvm]
Posted by "wrongtest-intellif (via GitHub)" <gi...@apache.org>.
wrongtest-intellif commented on PR #18421:
URL: https://github.com/apache/tvm/pull/18421#issuecomment-3540961235
/gemini review
--
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: discuss-archive-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
---------------------------------------------------------------------
To unsubscribe, e-mail: discuss-archive-unsubscribe@tvm.apache.org
For additional commands, e-mail: discuss-archive-help@tvm.apache.org
Re: [PR] [TIR] Add step attribute to ForNode (Initial codes) [tvm]
Posted by "wrongtest-intellif (via GitHub)" <gi...@apache.org>.
wrongtest-intellif commented on code in PR #18421:
URL: https://github.com/apache/tvm/pull/18421#discussion_r2532513549
##########
src/target/spirv/codegen_spirv.cc:
##########
@@ -672,10 +672,19 @@ void CodeGenSPIRV::VisitStmt_(const BufferStoreNode* op) {
}
void CodeGenSPIRV::VisitStmt_(const ForNode* op) {
- ICHECK(is_zero(op->min));
analyzer_->Bind(op->loop_var, Range::FromMinExtent(op->min, op->extent));
spirv::Value init_value = MakeValue(op->min);
spirv::Value extent_value = MakeValue(op->extent);
+
+ // loop step
+ spirv::Value step;
+ if (op->HasTrivialStep()) {
+ step = op->loop_var.dtype().is_int() ? builder_->IntImm(loop_var.stype, 1)
+ : builder_->UIntImm(loop_var.stype, 1);
+ } else {
+ step = MakeValue(tir::cast(op->extent->dtype, *op->step));
+ }
Review Comment:
> @wrongtest-intellif seems this is a valid comment
The review comment seems outdated, could gemini make another round review?
--
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: discuss-archive-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
---------------------------------------------------------------------
To unsubscribe, e-mail: discuss-archive-unsubscribe@tvm.apache.org
For additional commands, e-mail: discuss-archive-help@tvm.apache.org
Re: [PR] [TIR] Add step attribute to ForNode (Initial codes) [tvm]
Posted by "tqchen (via GitHub)" <gi...@apache.org>.
tqchen commented on PR #18421:
URL: https://github.com/apache/tvm/pull/18421#issuecomment-3538754487
@wrongtest-intellif @Hzfengsy would be great to check some of the gemini reviews, some seems to be valid
--
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: discuss-archive-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
---------------------------------------------------------------------
To unsubscribe, e-mail: discuss-archive-unsubscribe@tvm.apache.org
For additional commands, e-mail: discuss-archive-help@tvm.apache.org
Re: [PR] [TIR] Add step attribute to ForNode (Initial codes) [tvm]
Posted by "gemini-code-assist[bot] (via GitHub)" <gi...@apache.org>.
gemini-code-assist[bot] commented on code in PR #18421:
URL: https://github.com/apache/tvm/pull/18421#discussion_r2532632309
##########
src/target/source/codegen_webgpu.cc:
##########
@@ -668,12 +668,19 @@ void CodeGenWebGPU::VisitStmt_(const AllocateNode* op) {
void CodeGenWebGPU::VisitStmt_(const ForNode* op) {
std::string extent = PrintExpr(op->extent);
+ std::string step = op->step.has_value()? PrintExpr(*op->step) : "";
std::string vid = AllocVarID(op->loop_var.get());
ICHECK(is_zero(op->min));
PrintIndent();
stream << "for (var " << vid << " : ";
PrintType(op->loop_var.dtype(), stream);
- stream << " = 0; " << vid << " < " << extent << "; " << vid << "++) {\n";
+ stream << " = 0; " << vid << " < " << extent << "; " << vid;
+ if (step.empty()) {
+ stream << "++";
+ } else {
+ stream << " += " << step;
+ }
+ stream << ") {\n";
int for_scope = BeginScope();
PrintStmt(op->body);
this->EndScope(for_scope);
Review Comment:
It looks like you've made good progress on supporting non-zero `min` and non-trivial `step` in the WebGPU codegen! However, I noticed that the `ICHECK(is_zero(op->min));` on line 675 is still present. This check will prevent the codegen from handling loops where `min` is not zero, effectively negating the new functionality.
Please remove this `ICHECK` to fully enable support for non-zero `min` values.
```suggestion
void CodeGenWebGPU::VisitStmt_(const ForNode* op) {
std::string begin_str = PrintExpr(op->min);
PrimExpr end = is_zero(op->min) ? op->extent : arith::Analyzer().Simplify(op->min + op->extent);
std::string end_str = PrintExpr(end);
std::string step_str = op->step.has_value() ? PrintExpr(*op->step) : "";
std::string vid = AllocVarID(op->loop_var.get());
PrintIndent();
stream << "for (var " << vid << " : ";
PrintType(op->loop_var.dtype(), stream);
stream << " = " << begin_str << "; " << vid << " < " << end_str << "; " << vid;
if (step_str.empty()) {
stream << "++";
} else {
stream << " += " << step_str;
}
stream << ") {\n";
int for_scope = BeginScope();
```
```suggestion
void CodeGenWebGPU::VisitStmt_(const ForNode* op) {
std::string begin_str = PrintExpr(op->min);
PrimExpr end = is_zero(op->min) ? op->extent : arith::Analyzer().Simplify(op->min + op->extent);
std::string end_str = PrintExpr(end);
std::string step_str = op->step.has_value() ? PrintExpr(*op->step) : "";
std::string vid = AllocVarID(op->loop_var.get());
PrintIndent();
stream << "for (var " << vid << " : ";
PrintType(op->loop_var.dtype(), stream);
stream << " = " << begin_str << "; " << vid << " < " << end_str << "; " << vid;
if (step_str.empty()) {
stream << "++";
} else {
stream << " += " << step_str;
}
stream << ") {\n";
int for_scope = BeginScope();
```
--
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: discuss-archive-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
---------------------------------------------------------------------
To unsubscribe, e-mail: discuss-archive-unsubscribe@tvm.apache.org
For additional commands, e-mail: discuss-archive-help@tvm.apache.org
Re: [PR] [TIR] Add step attribute to ForNode (Initial codes) [tvm]
Posted by "wrongtest-intellif (via GitHub)" <gi...@apache.org>.
wrongtest-intellif commented on code in PR #18421:
URL: https://github.com/apache/tvm/pull/18421#discussion_r2532630691
##########
src/target/source/codegen_c.cc:
##########
@@ -1121,12 +1121,19 @@ void CodeGenC::VisitStmt_(const AssertStmtNode* op) {
void CodeGenC::VisitStmt_(const ForNode* op) {
std::string extent = PrintExpr(op->extent);
+ std::string step = op->step.has_value()? PrintExpr(*op->step) : "";
PrintIndent();
std::string vid = AllocVarID(op->loop_var.get());
ICHECK(is_zero(op->min));
stream << "for (";
PrintType(op->loop_var.dtype(), stream);
- stream << ' ' << vid << " = 0; " << vid << " < " << extent << "; ++" << vid << ") {\n";
+ stream << ' ' << vid << " = 0; " << vid << " < " << extent << "; ";
+ if (step.empty()) {
+ stream << "++" << vid;
+ } else {
+ stream << vid << " += " << step;
+ }
+ stream << ") {\n";
int for_scope = BeginScope();
PrintStmt(op->body);
this->EndScope(for_scope);
Review Comment:
/gemini review
##########
src/target/source/codegen_webgpu.cc:
##########
@@ -668,12 +668,19 @@ void CodeGenWebGPU::VisitStmt_(const AllocateNode* op) {
void CodeGenWebGPU::VisitStmt_(const ForNode* op) {
std::string extent = PrintExpr(op->extent);
+ std::string step = op->step.has_value()? PrintExpr(*op->step) : "";
std::string vid = AllocVarID(op->loop_var.get());
ICHECK(is_zero(op->min));
PrintIndent();
stream << "for (var " << vid << " : ";
PrintType(op->loop_var.dtype(), stream);
- stream << " = 0; " << vid << " < " << extent << "; " << vid << "++) {\n";
+ stream << " = 0; " << vid << " < " << extent << "; " << vid;
+ if (step.empty()) {
+ stream << "++";
+ } else {
+ stream << " += " << step;
+ }
+ stream << ") {\n";
int for_scope = BeginScope();
PrintStmt(op->body);
this->EndScope(for_scope);
Review Comment:
/gemini review
--
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: discuss-archive-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
---------------------------------------------------------------------
To unsubscribe, e-mail: discuss-archive-unsubscribe@tvm.apache.org
For additional commands, e-mail: discuss-archive-help@tvm.apache.org
Re: [PR] [TIR] Add step attribute to ForNode (Initial codes) [tvm]
Posted by "tqchen (via GitHub)" <gi...@apache.org>.
tqchen merged PR #18421:
URL: https://github.com/apache/tvm/pull/18421
--
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: discuss-archive-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
---------------------------------------------------------------------
To unsubscribe, e-mail: discuss-archive-unsubscribe@tvm.apache.org
For additional commands, e-mail: discuss-archive-help@tvm.apache.org
Re: [PR] [TIR] Add step attribute to ForNode (Initial codes) [tvm]
Posted by "gemini-code-assist[bot] (via GitHub)" <gi...@apache.org>.
gemini-code-assist[bot] commented on code in PR #18421:
URL: https://github.com/apache/tvm/pull/18421#discussion_r2532757292
##########
src/target/source/codegen_c.cc:
##########
@@ -1121,12 +1121,19 @@ void CodeGenC::VisitStmt_(const AssertStmtNode* op) {
void CodeGenC::VisitStmt_(const ForNode* op) {
std::string extent = PrintExpr(op->extent);
+ std::string step = op->step.has_value()? PrintExpr(*op->step) : "";
PrintIndent();
std::string vid = AllocVarID(op->loop_var.get());
ICHECK(is_zero(op->min));
stream << "for (";
PrintType(op->loop_var.dtype(), stream);
- stream << ' ' << vid << " = 0; " << vid << " < " << extent << "; ++" << vid << ") {\n";
+ stream << ' ' << vid << " = 0; " << vid << " < " << extent << "; ";
+ if (step.empty()) {
+ stream << "++" << vid;
+ } else {
+ stream << vid << " += " << step;
+ }
+ stream << ") {\n";
int for_scope = BeginScope();
PrintStmt(op->body);
this->EndScope(for_scope);
Review Comment:
Great job! The `ICHECK(is_zero(op->min));` has been successfully removed, and the C codegen for `ForNode` now correctly handles non-zero `min` values and non-trivial steps as per the generalized loop behavior. This resolves the `high` severity issue I previously highlighted. The code now aligns with the suggested implementation. Thank you for addressing this!
--
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: discuss-archive-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
---------------------------------------------------------------------
To unsubscribe, e-mail: discuss-archive-unsubscribe@tvm.apache.org
For additional commands, e-mail: discuss-archive-help@tvm.apache.org
Re: [PR] [TIR] Add step attribute to ForNode (Initial codes) [tvm]
Posted by "gemini-code-assist[bot] (via GitHub)" <gi...@apache.org>.
gemini-code-assist[bot] commented on code in PR #18421:
URL: https://github.com/apache/tvm/pull/18421#discussion_r2493249865
##########
src/target/source/codegen_webgpu.cc:
##########
@@ -668,12 +668,19 @@ void CodeGenWebGPU::VisitStmt_(const AllocateNode* op) {
void CodeGenWebGPU::VisitStmt_(const ForNode* op) {
std::string extent = PrintExpr(op->extent);
+ std::string step = op->step.has_value()? PrintExpr(*op->step) : "";
std::string vid = AllocVarID(op->loop_var.get());
ICHECK(is_zero(op->min));
PrintIndent();
stream << "for (var " << vid << " : ";
PrintType(op->loop_var.dtype(), stream);
- stream << " = 0; " << vid << " < " << extent << "; " << vid << "++) {\n";
+ stream << " = 0; " << vid << " < " << extent << "; " << vid;
+ if (step.empty()) {
+ stream << "++";
+ } else {
+ stream << " += " << step;
+ }
+ stream << ") {\n";
int for_scope = BeginScope();
PrintStmt(op->body);
this->EndScope(for_scope);
Review Comment:
Similar to the C codegen, the WebGPU codegen for `ForNode` still assumes `min` is zero. This should be updated to support non-zero `min` for consistency with other backends like LLVM and SPIR-V.
```c
std::string min_str = PrintExpr(op->min);
std::string extent_str = PrintExpr(op->extent);
std::string step_str = op->step.has_value() ? PrintExpr(*op->step) : "";
std::string vid = AllocVarID(op->loop_var.get());
PrintIndent();
stream << "for (var " << vid << " : ";
PrintType(op->loop_var.dtype(), stream);
stream << " = " << min_str << "; " << vid << " < (" << min_str << " + " << extent_str << "); " << vid;
if (step_str.empty()) {
stream << "++";
} else {
stream << " += " << step_str;
}
stream << ") {\n";
int for_scope = BeginScope();
PrintStmt(op->body);
this->EndScope(for_scope);
```
##########
src/target/spirv/codegen_spirv.cc:
##########
@@ -672,10 +672,19 @@ void CodeGenSPIRV::VisitStmt_(const BufferStoreNode* op) {
}
void CodeGenSPIRV::VisitStmt_(const ForNode* op) {
- ICHECK(is_zero(op->min));
analyzer_->Bind(op->loop_var, Range::FromMinExtent(op->min, op->extent));
spirv::Value init_value = MakeValue(op->min);
spirv::Value extent_value = MakeValue(op->extent);
+
+ // loop step
+ spirv::Value step;
+ if (op->HasTrivialStep()) {
+ step = op->loop_var.dtype().is_int() ? builder_->IntImm(loop_var.stype, 1)
+ : builder_->UIntImm(loop_var.stype, 1);
+ } else {
+ step = MakeValue(tir::cast(op->extent->dtype, *op->step));
+ }
Review Comment:
There seems to be a bug in the step calculation. The code `builder_->IntImm(loop_var.stype, 1)` uses `loop_var`, but the local `loop_var` of `spirv::PhiValue` type is defined later in the function. This means it's likely using the member variable `this->loop_var` from an outer loop, which is incorrect. The type of the current loop variable `op->loop_var` should be used instead.
```suggestion
if (op->HasTrivialStep()) {
spirv::SType stype = builder_->GetSType(op->loop_var.dtype());
step = op->loop_var.dtype().is_int() ? builder_->IntImm(stype, 1)
: builder_->UIntImm(stype, 1);
} else {
step = MakeValue(tir::cast(op->extent->dtype, *op->step));
}
```
##########
src/target/source/codegen_c.cc:
##########
@@ -1121,12 +1121,19 @@ void CodeGenC::VisitStmt_(const AssertStmtNode* op) {
void CodeGenC::VisitStmt_(const ForNode* op) {
std::string extent = PrintExpr(op->extent);
+ std::string step = op->step.has_value()? PrintExpr(*op->step) : "";
PrintIndent();
std::string vid = AllocVarID(op->loop_var.get());
ICHECK(is_zero(op->min));
stream << "for (";
PrintType(op->loop_var.dtype(), stream);
- stream << ' ' << vid << " = 0; " << vid << " < " << extent << "; ++" << vid << ") {\n";
+ stream << ' ' << vid << " = 0; " << vid << " < " << extent << "; ";
+ if (step.empty()) {
+ stream << "++" << vid;
+ } else {
+ stream << vid << " += " << step;
+ }
+ stream << ") {\n";
int for_scope = BeginScope();
PrintStmt(op->body);
this->EndScope(for_scope);
Review Comment:
The C codegen for `ForNode` still assumes that the loop starts from 0, as indicated by the `ICHECK(is_zero(op->min))` and the hardcoded `vid = 0`. Other backends like LLVM and SPIR-V have been updated to support non-zero `min`. For consistency and to fully support generalized loops, this backend should also be updated to handle non-zero `min`. This would involve removing the `ICHECK` and using `op->min` for loop initialization and in the loop condition.
```c
std::string min_str = PrintExpr(op->min);
std::string extent_str = PrintExpr(op->extent);
std::string step_str = op->step.has_value() ? PrintExpr(*op->step) : "";
PrintIndent();
std::string vid = AllocVarID(op->loop_var.get());
stream << "for (";
PrintType(op->loop_var.dtype(), stream);
stream << ' ' << vid << " = " << min_str << "; " << vid << " < (" << min_str << " + " << extent_str << "); ";
if (step_str.empty()) {
stream << "++" << vid;
} else {
stream << vid << " += " << step_str;
}
stream << ") {\n";
int for_scope = BeginScope();
PrintStmt(op->body);
this->EndScope(for_scope);
```
##########
src/te/operation/create_primfunc.cc:
##########
@@ -725,6 +725,8 @@ void RewriteStageToBlock(const te::Operation& op, CreateFuncInfo* info,
} else if (const auto extern_op = op.as<te::ExternOp>()) {
// Case 3. ExternOp (te.extern)
root_stmts->push_back(GenerateStmtFromExternOp(extern_op.value(), info));
+ } else if (const auto scan_op = op.as<te::ScanOp>()) {
+
Review Comment:
This `if` branch for `te::ScanOp` is empty. If this is a placeholder for future work, it would be better to add a `// TODO` comment explaining the intent or remove it for now to avoid dead code.
##########
src/relax/distributed/transform/lower_global_view_to_local_view.cc:
##########
@@ -330,8 +330,8 @@ class DistributedBufferCompactor : StmtExprMutator {
if (shard > 1) {
arith::Analyzer analyzer;
ICHECK(analyzer.CanProve(floormod(new_loop->extent, shard) == 0));
- return For(new_loop->loop_var, new_loop->min, floordiv(new_loop->extent, shard),
- new_loop->kind, new_loop->body, new_loop->thread_binding, new_loop->annotations);
+ new_loop.CopyOnWrite()->extent = floordiv(new_loop->extent, shard);
+ return new_loop;
Review Comment:
Using `CopyOnWrite` here is a great way to preserve the new `step` field and any other future fields on the `ForNode`. This pattern has been applied consistently across the PR, which is excellent.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] [TIR] Add step attribute to ForNode (Initial codes) [tvm]
Posted by "Hzfengsy (via GitHub)" <gi...@apache.org>.
Hzfengsy commented on PR #18421:
URL: https://github.com/apache/tvm/pull/18421#issuecomment-3500342887
cc @LeiWang1999
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] [TIR] Add step attribute to ForNode (Initial codes) [tvm]
Posted by "gemini-code-assist[bot] (via GitHub)" <gi...@apache.org>.
gemini-code-assist[bot] commented on code in PR #18421:
URL: https://github.com/apache/tvm/pull/18421#discussion_r2532523707
##########
src/target/spirv/codegen_spirv.cc:
##########
@@ -672,10 +672,19 @@ void CodeGenSPIRV::VisitStmt_(const BufferStoreNode* op) {
}
void CodeGenSPIRV::VisitStmt_(const ForNode* op) {
- ICHECK(is_zero(op->min));
analyzer_->Bind(op->loop_var, Range::FromMinExtent(op->min, op->extent));
spirv::Value init_value = MakeValue(op->min);
spirv::Value extent_value = MakeValue(op->extent);
+
+ // loop step
+ spirv::Value step;
+ if (op->HasTrivialStep()) {
+ step = op->loop_var.dtype().is_int() ? builder_->IntImm(loop_var.stype, 1)
+ : builder_->UIntImm(loop_var.stype, 1);
+ } else {
+ step = MakeValue(tir::cast(op->extent->dtype, *op->step));
+ }
Review Comment:
You are absolutely right, @wrongtest-intellif! My apologies for the outdated comment. The issue I highlighted regarding the `loop_var` being used before its local definition in `src/target/spirv/codegen_spirv.cc` has been resolved by the changes in this pull request. The `spirv::PhiValue loop_var` is now correctly declared before its usage in the step calculation, ensuring it refers to the local loop variable's type. Thank you for pointing this out!
--
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: discuss-archive-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
---------------------------------------------------------------------
To unsubscribe, e-mail: discuss-archive-unsubscribe@tvm.apache.org
For additional commands, e-mail: discuss-archive-help@tvm.apache.org
Re: [PR] [TIR] Add step attribute to ForNode (Initial codes) [tvm]
Posted by "gemini-code-assist[bot] (via GitHub)" <gi...@apache.org>.
gemini-code-assist[bot] commented on code in PR #18421:
URL: https://github.com/apache/tvm/pull/18421#discussion_r2533514663
##########
src/tir/transforms/canonicalize_loop.cc:
##########
@@ -0,0 +1,102 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+/*!
+ * \file tir/transforms/canonicalize_loop.cc
+ * \brief Canonicalize all loops to start from zero and step one.
+ */
+#include <tvm/arith/analyzer.h>
+#include <tvm/ffi/reflection/registry.h>
+#include <tvm/runtime/device_api.h>
+#include <tvm/tir/function.h>
+#include <tvm/tir/op.h>
+#include <tvm/tir/stmt_functor.h>
+#include <tvm/tir/transform.h>
+
+#include <utility>
+
+namespace tvm {
+namespace tir {
+
+class LoopCanonicalizer : public StmtExprMutator {
+ public:
+ LoopCanonicalizer() = default;
+
+ private:
+ Stmt VisitStmt_(const ForNode* op) final {
+ if (is_zero(op->min) && op->HasTrivialStep()) {
+ return StmtExprMutator::VisitStmt_(op);
+ }
+ arith::Analyzer analyzer;
+ const auto* loop_var = op->loop_var.get();
+ PrimExpr step = op->step.value_or(make_const(loop_var->dtype, 1));
+
+ // report warning for negative step, since it would be a forever loop
+ if (!analyzer.CanProveGreaterEqual(step, 1)) {
+ // TODO(tvm): prove dynamic shaped step
+ LOG(FATAL) << "Loop step for " << op->loop_var << " may not be positive: " << step;
+ }
+
+ new_iter_info_[loop_var] = std::make_pair(step, op->min);
+ auto n = CopyOnWrite(op);
+ n->body = VisitStmt(op->body);
+ n->min = make_zero(loop_var->dtype);
+ n->extent = analyzer.Simplify(ceildiv(op->extent - op->min + 1, step));
Review Comment:
The calculation for the new loop extent seems incorrect. The `ForNode` documentation states that a loop runs from `min` to `min + extent`. With a step, the number of iterations should be `ceildiv(extent, step)`.
The current formula `ceildiv(op->extent - op->min + 1, step)` appears to be incorrect. For example, for a loop `for i in range(2, 12, 3)`, we have `min=2` and `extent=10`. The loop should have 4 iterations (2, 5, 8, 11). The correct formula `ceildiv(10, 3)` gives 4. However, the formula in the code gives `ceildiv(10 - 2 + 1, 3) = ceildiv(9, 3) = 3`, which is incorrect.
The correct formula for the number of iterations should be `ceildiv(op->extent, step)`.
```suggestion
n->extent = analyzer.Simplify(ceildiv(op->extent, step));
```
##########
tests/python/codegen/test_target_codegen_cuda.py:
##########
@@ -877,5 +877,37 @@ def main(A: T.Buffer((16, 16), "float32"), B: T.Buffer((16, 16), "float32")):
assert "return;" in cuda_code
+@tvm.testing.requires_gpu
+@tvm.testing.requires_cuda
+def test_cuda_loop_step():
+ @T.prim_func
+ def cuda_loop_step(
+ A: T.Buffer((1024,), "float32"),
+ B: T.Buffer((1024,), "float32"),
+ C: T.Buffer((1024,), "float32"),
+ ):
+ # Each thread computes a strided subset of the i loop: start = tx*3, step = 96 (3 * 32 threads)
+ for bx in T.thread_binding(1, "blockIdx.x"):
+ for tx in T.thread_binding(96, "threadIdx.x"):
+ for i in T.serial(tx, 4096, step=96):
Review Comment:
The loop's stop condition `4096` will lead to out-of-bounds memory access, as the buffers `A`, `B`, and `C` are all defined with a size of 1024. The test expects all elements to be computed, which means the loop should cover indices from 0 to 1023. To achieve this with the given threading scheme, the stop value should be 1024.
```suggestion
for i in T.serial(tx, 1024, step=96):
```
--
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: discuss-archive-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
---------------------------------------------------------------------
To unsubscribe, e-mail: discuss-archive-unsubscribe@tvm.apache.org
For additional commands, e-mail: discuss-archive-help@tvm.apache.org
Re: [PR] [TIR] Add step attribute to ForNode (Initial codes) [tvm]
Posted by "gemini-code-assist[bot] (via GitHub)" <gi...@apache.org>.
gemini-code-assist[bot] commented on code in PR #18421:
URL: https://github.com/apache/tvm/pull/18421#discussion_r2532631337
##########
src/target/source/codegen_c.cc:
##########
@@ -1121,12 +1121,19 @@ void CodeGenC::VisitStmt_(const AssertStmtNode* op) {
void CodeGenC::VisitStmt_(const ForNode* op) {
std::string extent = PrintExpr(op->extent);
+ std::string step = op->step.has_value()? PrintExpr(*op->step) : "";
PrintIndent();
std::string vid = AllocVarID(op->loop_var.get());
ICHECK(is_zero(op->min));
stream << "for (";
PrintType(op->loop_var.dtype(), stream);
- stream << ' ' << vid << " = 0; " << vid << " < " << extent << "; ++" << vid << ") {\n";
+ stream << ' ' << vid << " = 0; " << vid << " < " << extent << "; ";
+ if (step.empty()) {
+ stream << "++" << vid;
+ } else {
+ stream << vid << " += " << step;
+ }
+ stream << ") {\n";
int for_scope = BeginScope();
PrintStmt(op->body);
this->EndScope(for_scope);
Review Comment:
It looks like you've made good progress on generalizing the C codegen for `ForNode` to handle non-zero `min` and non-trivial `step`, which aligns with my previous suggestion. However, I noticed that the `ICHECK(is_zero(op->min));` on line 1129 is still present. This check will prevent the code from handling non-zero `min` values, contradicting the new logic you've introduced. Please remove this `ICHECK` to fully enable the generalized loop behavior.
This is a `high` severity issue as it prevents the intended functionality of supporting non-zero `min` values in loops.
```suggestion
std::string begin_str = PrintExpr(op->min);
PrimExpr end = is_zero(op->min) ? op->extent : arith::Analyzer().Simplify(op->min + op->extent);
std::string end_str = PrintExpr(end);
std::string step_str = op->step.has_value() ? PrintExpr(*op->step) : "";
PrintIndent();
std::string vid = AllocVarID(op->loop_var.get());
stream << "for (";
PrintType(op->loop_var.dtype(), stream);
stream << ' ' << vid << " = " << begin_str << "; " << vid << " < " << end_str << "; ";
if (step_str.empty()) {
stream << "++" << vid;
} else {
stream << vid << " += " << step_str;
}
stream << ") {\n";
int for_scope = BeginScope();
PrintStmt(op->body);
this->EndScope(for_scope);
```
--
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: discuss-archive-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
---------------------------------------------------------------------
To unsubscribe, e-mail: discuss-archive-unsubscribe@tvm.apache.org
For additional commands, e-mail: discuss-archive-help@tvm.apache.org
Re: [PR] [TIR] Add step attribute to ForNode (Initial codes) [tvm]
Posted by "wrongtest-intellif (via GitHub)" <gi...@apache.org>.
wrongtest-intellif commented on code in PR #18421:
URL: https://github.com/apache/tvm/pull/18421#discussion_r2533751212
##########
tests/python/codegen/test_target_codegen_cuda.py:
##########
@@ -877,5 +877,37 @@ def main(A: T.Buffer((16, 16), "float32"), B: T.Buffer((16, 16), "float32")):
assert "return;" in cuda_code
+@tvm.testing.requires_gpu
+@tvm.testing.requires_cuda
+def test_cuda_loop_step():
+ @T.prim_func
+ def cuda_loop_step(
+ A: T.Buffer((1024,), "float32"),
+ B: T.Buffer((1024,), "float32"),
+ C: T.Buffer((1024,), "float32"),
+ ):
+ # Each thread computes a strided subset of the i loop: start = tx*3, step = 96 (3 * 32 threads)
+ for bx in T.thread_binding(1, "blockIdx.x"):
+ for tx in T.thread_binding(96, "threadIdx.x"):
+ for i in T.serial(tx, 4096, step=96):
Review Comment:
/gemini review
--
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: discuss-archive-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
---------------------------------------------------------------------
To unsubscribe, e-mail: discuss-archive-unsubscribe@tvm.apache.org
For additional commands, e-mail: discuss-archive-help@tvm.apache.org
Re: [PR] [TIR] Add step attribute to ForNode (Initial codes) [tvm]
Posted by "gemini-code-assist[bot] (via GitHub)" <gi...@apache.org>.
gemini-code-assist[bot] commented on code in PR #18421:
URL: https://github.com/apache/tvm/pull/18421#discussion_r2533752398
##########
tests/python/codegen/test_target_codegen_cuda.py:
##########
@@ -877,5 +877,37 @@ def main(A: T.Buffer((16, 16), "float32"), B: T.Buffer((16, 16), "float32")):
assert "return;" in cuda_code
+@tvm.testing.requires_gpu
+@tvm.testing.requires_cuda
+def test_cuda_loop_step():
+ @T.prim_func
+ def cuda_loop_step(
+ A: T.Buffer((1024,), "float32"),
+ B: T.Buffer((1024,), "float32"),
+ C: T.Buffer((1024,), "float32"),
+ ):
+ # Each thread computes a strided subset of the i loop: start = tx*3, step = 96 (3 * 32 threads)
+ for bx in T.thread_binding(1, "blockIdx.x"):
+ for tx in T.thread_binding(96, "threadIdx.x"):
+ for i in T.serial(tx, 4096, step=96):
Review Comment:
It seems the previous suggestion was not applied. The loop's stop condition `4096` will still lead to out-of-bounds memory access, as the buffers `A`, `B`, and `C` are all defined with a size of 1024. The test expects all elements to be computed, which means the loop should cover indices from 0 to 1023. To achieve this with the given threading scheme, the stop value should be 1024.
This is a `critical` issue as it can lead to incorrect behavior or crashes due to memory corruption.
```suggestion
for i in T.serial(tx, 1024, step=96):
```
--
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: discuss-archive-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
---------------------------------------------------------------------
To unsubscribe, e-mail: discuss-archive-unsubscribe@tvm.apache.org
For additional commands, e-mail: discuss-archive-help@tvm.apache.org
Re: [PR] [TIR] Add step attribute to ForNode (Initial codes) [tvm]
Posted by "tqchen (via GitHub)" <gi...@apache.org>.
tqchen commented on PR #18421:
URL: https://github.com/apache/tvm/pull/18421#issuecomment-3570804705
Thanks @wrongtest-intellif !
--
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: discuss-archive-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
---------------------------------------------------------------------
To unsubscribe, e-mail: discuss-archive-unsubscribe@tvm.apache.org
For additional commands, e-mail: discuss-archive-help@tvm.apache.org
Re: [PR] [TIR] Add step attribute to ForNode (Initial codes) [tvm]
Posted by "wrongtest-intellif (via GitHub)" <gi...@apache.org>.
wrongtest-intellif commented on PR #18421:
URL: https://github.com/apache/tvm/pull/18421#issuecomment-3550508822
> @wrongtest-intellif @Hzfengsy would be great to check some of the gemini reviews, some seems to be valid
gemini review issues resolved cc @tqchen
--
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: discuss-archive-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
---------------------------------------------------------------------
To unsubscribe, e-mail: discuss-archive-unsubscribe@tvm.apache.org
For additional commands, e-mail: discuss-archive-help@tvm.apache.org