[CINN+PIR]Support softmax graph for PIRCompiler

[Aurelius84]

PR types

Others

PR changes

Others

Description

Pcard-67164

验证了softmax 子图的生成和执行，机制上主要包括：

• 添加了OpMapper机制，处理可变Attribute对Operation的Oprand和Attribute的影响。后续将考虑升级为 Pass + cinn::dialect
• 删除了 cinn/utlils/attribute_utils.h，统一到hlir/framework/pir/utils.h
• 规范化 new_ir 目录、newir namespace 为 pir

附 softmax 子图生成的 CUDA C code:

#include <cstdint>

#define CINN_WITH_CUDA
#include "bfloat16.h"
#include "float16.h"
using cinn::common::bfloat16;
using cinn::common::float16;
using cinn::common::half4;
using cinn::common::half8;
using cinn::common::float8;

#include "cinn_cuda_runtime_source.cuh"

extern "C" {

__global__
void __launch_bounds__(1024) fn_fill_constant_kernel(float* __restrict__ var_72428416)
{
  if (((int)blockIdx.x < 8)) {
    if (((int)threadIdx.x < 1024)) {
      var_72428416[((1024 * (int)blockIdx.x) + (int)threadIdx.x)] = 1.00000000f;
    };
  };
}__global__
void __launch_bounds__(1024) fn_reduce_max_subtract_exp_reduce_sum_divide_kernel(const float* __restrict__ var_72428416, float* __restrict__ var_71976912, float* __restrict__ var_71946928, float* __restrict__ var_71269344, float* __restrict__ var_71256320, float* __restrict__ var_71249104)
{
  float _var_71256320_tmp_temp_buffer [ 1 ];
  float _var_71976912_tmp_temp_buffer [ 1 ];
  float* var_71256320_tmp = _var_71256320_tmp_temp_buffer;
  float* var_71976912_tmp = _var_71976912_tmp_temp_buffer;
  if (((int)blockIdx.x < 64)) {
    if (((int)threadIdx.x < 128)) {
      var_71976912_tmp[0] = cinn_block_reduce_max_fp32_internal(var_72428416[((128 * (int)blockIdx.x) + (int)threadIdx.x)]);
    };
  };
  if (((int)blockIdx.x < 64)) {
    if (((int)threadIdx.x < 1)) {
      var_71976912[(int)blockIdx.x] = var_71976912_tmp[0];
    };
  };
  if (((int)blockIdx.x < 8)) {
    if (((int)threadIdx.x < 1024)) {
      var_71946928[((1024 * (int)blockIdx.x) + (int)threadIdx.x)] = (var_72428416[((1024 * (int)blockIdx.x) + (int)threadIdx.x)] - var_71976912[(((int)threadIdx.x / 128) + (8 * (int)blockIdx.x))]);
    };
  };
  if (((int)blockIdx.x < 8)) {
    if (((int)threadIdx.x < 1024)) {
      var_71269344[((1024 * (int)blockIdx.x) + (int)threadIdx.x)] = cinn_nvgpu_exp_fp32(var_71946928[((1024 * (int)blockIdx.x) + (int)threadIdx.x)]);
    };
  };
  if (((int)blockIdx.x < 64)) {
    if (((int)threadIdx.x < 128)) {
      var_71256320_tmp[0] = cinn_block_reduce_sum_fp32_internal(var_71269344[((128 * (int)blockIdx.x) + (int)threadIdx.x)]);
    };
  };
  if (((int)blockIdx.x < 64)) {
    if (((int)threadIdx.x < 1)) {
      var_71256320[(int)blockIdx.x] = var_71256320_tmp[0];
    };
  };
  if (((int)blockIdx.x < 8)) {
    if (((int)threadIdx.x < 1024)) {
      var_71249104[((1024 * (int)blockIdx.x) + (int)threadIdx.x)] = (var_71269344[((1024 * (int)blockIdx.x) + (int)threadIdx.x)] / var_71256320[(((int)threadIdx.x / 128) + (8 * (int)blockIdx.x))]);
    };
  };
}

}

[paddle-bot]

你的PR提交成功，感谢你对开源项目的贡献!
请关注后续CI自动化测试结果，详情请参考Paddle-CI手册。
Your PR has been submitted. Thanks for your contribution!
Please wait for the result of CI firstly. See Paddle CI Manual for details.

[XiaoguangHu01]

LGTM

[XiaoguangHu01]

LGTM