[CodeGen] Cleanup generated code

[wpan11nv]

• remove unnecessary white spaces from storage kind
• do not start a new scope for vectorization as temporary
  variables are all uniquely generated.

The above two changes make vectorized code much cleaner.

Signed-off-by: Wei Pan [email protected]

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 by @ them in the pull request thread.

[wpan11nv]

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;
}