__global__
 void __launch_bounds__(128) fn_reduce_max_subtract_exp_reduce_sum_divide_kernel(const float* __restrict__ var_66069008, float* __restrict__ var_65126352)
 {
   float _var_64979360_temp_buffer [ 1 ];
   __shared__ float _var_65147536_temp_buffer [ 1 ];
   float _var_65147536_tmp_temp_buffer [ 1 ];
   float _var_65430496_temp_buffer [ 1 ];
   __shared__ float _var_65460480_temp_buffer [ 1 ];
   float _var_65460480_tmp_temp_buffer [ 1 ];
   float* var_64979360 = _var_64979360_temp_buffer;
   float* var_65147536 = _var_65147536_temp_buffer;
   float* var_65147536_tmp = _var_65147536_tmp_temp_buffer;
   float* var_65430496 = _var_65430496_temp_buffer;
   float* var_65460480 = _var_65460480_temp_buffer;
   float* var_65460480_tmp = _var_65460480_tmp_temp_buffer;
   if (((int)blockIdx.x < 64)) {
     if (((int)threadIdx.x < 128)) {
       var_65460480_tmp[0] = cinn_block_reduce_max_fp32_internal(var_66069008[((128 * (int)blockIdx.x) + (int)threadIdx.x)]);
     };
     if (((int)threadIdx.x < 1)) {
       var_65460480[0] = var_65460480_tmp[0];
     };
     __syncthreads();
     if (((int)threadIdx.x < 128)) {
       var_65430496[0] = (var_66069008[((128 * (int)blockIdx.x) + (int)threadIdx.x)] - var_65460480[0]);
       var_64979360[0] = cinn_nvgpu_exp_fp32(var_65430496[0]);
       var_65147536_tmp[0] = cinn_block_reduce_sum_fp32_internal(var_64979360[0]);
     };
     if (((int)threadIdx.x < 1)) {
       var_65147536[0] = var_65147536_tmp[0];
     };
     __syncthreads();
     if (((int)threadIdx.x < 128)) {
       var_65126352[((128 * (int)blockIdx.x) + (int)threadIdx.x)] = (var_64979360[0] / var_65147536[0]);
     };
   };
 }