You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by GitBox <gi...@apache.org> on 2020/04/23 18:42:33 UTC
[GitHub] [incubator-tvm] wpan11nv opened a new pull request #5424: [CodeGen] Cleanup generated code
wpan11nv opened a new pull request #5424:
URL: https://github.com/apache/incubator-tvm/pull/5424
- remove unnecessary white spaces from storage kind
- do not start a new scope for vectorization as temporary
variables are alll uniquely generated.
The above two changes make vectorized code much cleaner.
Signed-off-by: Wei Pan <we...@nvidia.com>
Thanks for contributing to TVM! Please refer to guideline https://tvm.apache.org/docs/contribute/ for useful information and tips. After the pull request is submitted, please request code reviews from [Reviewers](https://github.com/apache/incubator-tvm/blob/master/CONTRIBUTORS.md#reviewers) by @ them in the pull request thread.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] wpan11nv commented on pull request #5424: [CodeGen] Cleanup generated code
Posted by GitBox <gi...@apache.org>.
wpan11nv commented on pull request #5424:
URL: https://github.com/apache/incubator-tvm/pull/5424#issuecomment-618591146
Notably extra white spaces or scopes are removed. When the loop nesting or expression is deep, the emitted becomes unreadable.
Before:
`extern "C" __global__ void test_kernel0(void* __restrict__ B, void* __restrict__ A) {
float4 _1;
{
float4 _2 = (( float4*)(( float*)A + ((((int)blockIdx.x) * 4))))[0];
float4 _3 = make_float4(1.000000e+00f, 1.000000e+00f, 1.000000e+00f, 1.000000e+00f);
_1.x = (_2.x+_3.x);
_1.y = (_2.y+_3.y);
_1.z = (_2.z+_3.z);
_1.w = (_2.w+_3.w);
}
(( float4*)(( float*)B + ((((int)blockIdx.x) * 4))))[0] = _1;
}`
After:
`extern "C" __global__ void test_kernel0(void* __restrict__ B, void* __restrict__ A) {
float4 _1;
float4 _2 = ((float4*)((float*)A + ((((int)blockIdx.x) * 4))))[0];
float4 _3 = make_float4(1.000000e+00f, 1.000000e+00f, 1.000000e+00f, 1.000000e+00f);
_1.x = (_2.x+_3.x);
_1.y = (_2.y+_3.y);
_1.z = (_2.z+_3.z);
_1.w = (_2.w+_3.w);
((float4*)((float*)B + ((((int)blockIdx.x) * 4))))[0] = _1;
}`
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] wpan11nv edited a comment on pull request #5424: [CodeGen] Cleanup generated code
Posted by GitBox <gi...@apache.org>.
wpan11nv edited a comment on pull request #5424:
URL: https://github.com/apache/incubator-tvm/pull/5424#issuecomment-618591146
Notably extra white spaces or scopes are removed. When the loop nesting or expression is deep, the emitted becomes unreadable.
Before:
`extern "C" __global__ void test_kernel0(void* __restrict__ B, void* __restrict__ A) {
float4 _1;
{
float4 _2 = (( float4*)(( float*)A + ((((int)blockIdx.x) * 4))))[0];
float4 _3 = make_float4(1.000000e+00f, 1.000000e+00f, 1.000000e+00f, 1.000000e+00f);
_1.x = (_2.x+_3.x);
_1.y = (_2.y+_3.y);
_1.z = (_2.z+_3.z);
_1.w = (_2.w+_3.w);
}
(( float4*)(( float*)B + ((((int)blockIdx.x) * 4))))[0] = _1;
}`
After:
```extern "C" __global__ void test_kernel0(void* __restrict__ B, void* __restrict__ A) {
float4 _1;
float4 _2 = ((float4*)((float*)A + ((((int)blockIdx.x) * 4))))[0];
float4 _3 = make_float4(1.000000e+00f, 1.000000e+00f, 1.000000e+00f, 1.000000e+00f);
_1.x = (_2.x+_3.x);
_1.y = (_2.y+_3.y);
_1.z = (_2.z+_3.z);
_1.w = (_2.w+_3.w);
((float4*)((float*)B + ((((int)blockIdx.x) * 4))))[0] = _1;
}```
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] wpan11nv edited a comment on pull request #5424: [CodeGen] Cleanup generated code
Posted by GitBox <gi...@apache.org>.
wpan11nv edited a comment on pull request #5424:
URL: https://github.com/apache/incubator-tvm/pull/5424#issuecomment-618591146
Notably extra white spaces or scopes are removed. When the loop nesting or expression is deep, the emitted becomes unreadable.
Before:
``` c++
extern "C" __global__ void test_kernel0(void* __restrict__ B, void* __restrict__ A) {
float4 _1;
{
float4 _2 = (( float4*)(( float*)A + ((((int)blockIdx.x) * 4))))[0];
float4 _3 = make_float4(1.000000e+00f, 1.000000e+00f, 1.000000e+00f, 1.000000e+00f);
_1.x = (_2.x+_3.x);
_1.y = (_2.y+_3.y);
_1.z = (_2.z+_3.z);
_1.w = (_2.w+_3.w);
}
(( float4*)(( float*)B + ((((int)blockIdx.x) * 4))))[0] = _1;
}
```
After:
```c++
extern "C" __global__ void test_kernel0(void* __restrict__ B, void* __restrict__ A) {
float4 _1;
float4 _2 = ((float4*)((float*)A + ((((int)blockIdx.x) * 4))))[0];
float4 _3 = make_float4(1.000000e+00f, 1.000000e+00f, 1.000000e+00f, 1.000000e+00f);
_1.x = (_2.x+_3.x);
_1.y = (_2.y+_3.y);
_1.z = (_2.z+_3.z);
_1.w = (_2.w+_3.w);
((float4*)((float*)B + ((((int)blockIdx.x) * 4))))[0] = _1;
}
```
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] wpan11nv edited a comment on pull request #5424: [CodeGen] Cleanup generated code
Posted by GitBox <gi...@apache.org>.
wpan11nv edited a comment on pull request #5424:
URL: https://github.com/apache/incubator-tvm/pull/5424#issuecomment-618591146
Notably extra white spaces or scopes are removed. When the loop nesting or expression is deep, the emitted becomes unreadable.
Before:
`extern "C" __global__ void test_kernel0(void* __restrict__ B, void* __restrict__ A) {
float4 _1;
{
float4 _2 = (( float4*)(( float*)A + ((((int)blockIdx.x) * 4))))[0];
float4 _3 = make_float4(1.000000e+00f, 1.000000e+00f, 1.000000e+00f, 1.000000e+00f);
_1.x = (_2.x+_3.x);
_1.y = (_2.y+_3.y);
_1.z = (_2.z+_3.z);
_1.w = (_2.w+_3.w);
}
(( float4*)(( float*)B + ((((int)blockIdx.x) * 4))))[0] = _1;
}`
After:
``extern "C" __global__ void test_kernel0(void* __restrict__ B, void* __restrict__ A) {
float4 _1;
float4 _2 = ((float4*)((float*)A + ((((int)blockIdx.x) * 4))))[0];
float4 _3 = make_float4(1.000000e+00f, 1.000000e+00f, 1.000000e+00f, 1.000000e+00f);
_1.x = (_2.x+_3.x);
_1.y = (_2.y+_3.y);
_1.z = (_2.z+_3.z);
_1.w = (_2.w+_3.w);
((float4*)((float*)B + ((((int)blockIdx.x) * 4))))[0] = _1;
}
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] wpan11nv edited a comment on pull request #5424: [CodeGen] Cleanup generated code
Posted by GitBox <gi...@apache.org>.
wpan11nv edited a comment on pull request #5424:
URL: https://github.com/apache/incubator-tvm/pull/5424#issuecomment-618591146
Notably extra white spaces or scopes are removed. When the loop nesting or expression is deep, the emitted becomes unreadable.
Before:
`extern "C" __global__ void test_kernel0(void* __restrict__ B, void* __restrict__ A) {
float4 _1;
{
float4 _2 = (( float4*)(( float*)A + ((((int)blockIdx.x) * 4))))[0];
float4 _3 = make_float4(1.000000e+00f, 1.000000e+00f, 1.000000e+00f, 1.000000e+00f);
_1.x = (_2.x+_3.x);
_1.y = (_2.y+_3.y);
_1.z = (_2.z+_3.z);
_1.w = (_2.w+_3.w);
}
(( float4*)(( float*)B + ((((int)blockIdx.x) * 4))))[0] = _1;
}`
After:
extern "C" __global__ void test_kernel0(void* __restrict__ B, void* __restrict__ A) {
float4 _1;
float4 _2 = ((float4*)((float*)A + ((((int)blockIdx.x) * 4))))[0];
float4 _3 = make_float4(1.000000e+00f, 1.000000e+00f, 1.000000e+00f, 1.000000e+00f);
_1.x = (_2.x+_3.x);
_1.y = (_2.y+_3.y);
_1.z = (_2.z+_3.z);
_1.w = (_2.w+_3.w);
((float4*)((float*)B + ((((int)blockIdx.x) * 4))))[0] = _1;
}
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org