Skip to content

[MetaSchedule] Use Add-Unit-Loop in Auto-Bind - #11581

Merged
spectrometerHBH merged 1 commit into
apache:mainfrom
junrushao:bugfix/2022-06-04/auto-bind-with-no-spatial-loop
Jun 5, 2022
Merged

[MetaSchedule] Use Add-Unit-Loop in Auto-Bind#11581
spectrometerHBH merged 1 commit into
apache:mainfrom
junrushao:bugfix/2022-06-04/auto-bind-with-no-spatial-loop

Conversation

@junrushao

Copy link
Copy Markdown
Member

Following #11575, this PR allows CUDA thread binding for TIR programs
like

@T.prim_func
def zero_dim_add(
    A: T.Buffer[(), "float32"],
    B: T.Buffer[(), "float32"],
    C: T.Buffer[(), "float32"],
) -> None:
    with T.block("C"):
        vi = T.axis.spatial(1, 0)
        C[()] = A[()] + B[()]

where there is no loop available to be bound to threadIdx/blockIdx.

Following apache#11575, this PR allows CUDA thread binding for TIR programs
like

```python
@T.prim_func
def zero_dim_add(
    A: T.Buffer[(), "float32"],
    B: T.Buffer[(), "float32"],
    C: T.Buffer[(), "float32"],
) -> None:
    with T.block("C"):
        vi = T.axis.spatial(1, 0)
        C[()] = A[()] + B[()]
```

where there is no loop available to be bound to threadIdx/blockIdx.
@spectrometerHBH
spectrometerHBH merged commit ba60788 into apache:main Jun 5, 2022
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants