enhance reduce mapping: bind block for outer-fused-loop · Pull Request !26 · MindSpore/akg

Pull Requests

/

详情

已合并

yangsijia:enhance-auto-tiling MindSpore:master

创建于 2021-03-27 17:13

克隆/下载

<!--  Thanks for sending a pull request!  Here are some tips for you:

If this is your first time, please read our contributor guidelines: https://gitee.com/mindspore/akg/blob/master/CONTRIBUTING.md
-->

**What type of PR is this?**
> kind task

**What does this PR do / why do we need it**:
> This PR enhances auto tiling strategy for reduce op with more than three dimension: bind block for outer loop to activate more blocks to increase parallelism.

After applying this PR, we increases the performance from `8e-4 sec/op` to `4e-05 sec/op`, and the result cuda is 
```
#include "akg_reduce/reduce.cuh"
extern "C" __global__ void reduce_sum_float32_41_54_40_2_True_kernel0( float* __restrict__ input_1,  float* __restrict__ input_1_red) {
  __shared__ float input_1_red_shared[54];
   float acc_0[1];
  __shared__ float red_buf0[32];
  if (((int)threadIdx.x) == 0) {
    for (int cc2 = 0; cc2 < 54; ++cc2) {
      input_1_red_shared[cc2] = 0.000000e+00f;
    }
  }
  __syncthreads();
  for (int cc21 = 0; cc21 < 54; ++cc21) {
    acc_0[0] = 0.000000e+00f;
    for (int cc3 = 0; cc3 < 2; ++cc3) {
      if (((cc3 * 32) + ((int)threadIdx.x)) <= 39) {
        acc_0[0] = (acc_0[0] + input_1[((((((int)blockIdx.y) * 2160) + (cc21 * 40)) + (cc3 * 32)) + ((int)threadIdx.x))]);
      }
    }
    (void)akg_reduce::AkgReduce<float,akg_reduce::SumOp, 32, 1, akg_reduce::REDUCE2D_X>(akg_reduce::SumOp(), &(input_1_red_shared[cc21]), red_buf0, acc_0[0], 32);
  }
  __syncthreads();
  for (int cc22 = 0; cc22 < 2; ++cc22) {
    if (((cc22 * 32) + ((int)threadIdx.x)) <= 53) {
      input_1_red[(((((int)blockIdx.y) * 54) + (cc22 * 32)) + ((int)threadIdx.x))] = input_1_red_shared[((cc22 * 32) + ((int)threadIdx.x))];
    }
  }
  __syncthreads();
}

```
**Which issue(s) this PR fixes**:

Fixes https://gitee.com/mindspore/akg/issues/I3DV0V

**Special notes for your reviewers**: