[XPU] Supports the different types of post dynamic quantization for conv and fc (#59307)

Travis-Lee · web-flow · commit b10fd56bbcb0 · 2023-11-30T18:32:31.000+08:00
diff --git a/paddle/fluid/framework/ir/xpu/conv2d_xpu_fuse_pass.cc b/paddle/fluid/framework/ir/xpu/conv2d_xpu_fuse_pass.cc
@@ -489,8 +489,9 @@ void Conv2dXPUFusePass::CreateTheReplicatedWeights(
       true,
       platform::errors::InvalidArgument("conv node ptr can not be null"));
   auto conv_filter_name = conv->Op()->Input("Filter")[0];
-  std::string replicated_filter_name =
-      conv_filter_name + "_copy_" + std::to_string(conv->id());
+  std::string replicated_filter_name = conv_filter_name + "_copy_" +
+                                       std::to_string(block->ID()) + "_" +
+                                       std::to_string(conv->id());
   auto* replicated_filter_var = scope->FindVar(replicated_filter_name);
   if (replicated_filter_var == nullptr) {
     auto* filter_tensor =
@@ -536,8 +537,9 @@ void Conv2dXPUFusePass::CreateFusionWeightsAndBias(
   auto conv_filter_name = conv->Op()->Input("Filter")[0];
   Node* conv_filter = FindNodeWithName(graph, conv_filter_name);
   CreateTheReplicatedWeights(graph, scope, block, nodes_map);
-  std::string replicated_filter_name =
-      conv_filter_name + "_copy_" + std::to_string(conv->id());
+  std::string replicated_filter_name = conv_filter_name + "_copy_" +
+                                       std::to_string(block->ID()) + "_" +
+                                       std::to_string(conv->id());
   auto* conv_filter_replicated_node =
       FindNodeWithName(graph, replicated_filter_name);
   auto* filter_t =
@@ -718,28 +720,69 @@ void Conv2dXPUFusePass::CreateFusionWeightsAndBias(
   Node* filter_intx = nullptr;
   Node* filter_max = nullptr;
   Node* scale_max = nullptr;
-  bool per_channel_quant = false;
+
+  std::map<std::string, int> default_type;
+  default_type.insert(std::make_pair("conv2d", -1));
+  auto quant_post_type =
+      Has("quant_post_dynamic_weight_methods")
+          ? Get<std::map<std::string, int>>("quant_post_dynamic_weight_methods")
+          : default_type;
+
+  for (auto it = quant_post_type.begin(); it != quant_post_type.end(); ++it) {
+    VLOG(5) << "Key:" << it->first;
+    VLOG(5) << "Value:" << it->second;
+  }
+
   if (op_weights_precision != "int8") {
-    PrepareWeight<float, int16_t>(graph,
-                                  scope,
-                                  block,
-                                  conv_filter_replicated_node,
-                                  &filter_intx,
-                                  &filter_max,
-                                  &scale_max,
-                                  false,
-                                  weight_scale,
-                                  per_channel_quant);
+    if (quant_post_type.find("conv2d") != quant_post_type.end() &&
+            quant_post_type.find("conv2d")->second == 2 ||
+        quant_post_type.find("conv2d") != quant_post_type.end() &&
+            quant_post_type.find("conv2d")->second == -1) {
+      VLOG(5) << "Use int16 per-tensor weight";
+      PrepareWeight<float, int16_t>(graph,
+                                    scope,
+                                    block,
+                                    conv_filter_replicated_node,
+                                    &filter_intx,
+                                    &filter_max,
+                                    &scale_max,
+                                    false,
+                                    weight_scale,
+                                    false);
+    } else if (quant_post_type.find("conv2d") != quant_post_type.end() &&
+               quant_post_type.find("conv2d")->second == 3) {
+      VLOG(5) << "Use int16 per-channel weight";
+      PrepareWeight<float, int16_t>(graph,
+                                    scope,
+                                    block,
+                                    conv_filter_replicated_node,
+                                    &filter_intx,
+                                    &filter_max,
+                                    &scale_max,
+                                    false,
+                                    weight_scale,
+                                    true);
+    } else {
+      VLOG(5) << "Unsupported type weight by non-int8!";
+    }
+
   } else {
-    PrepareWeight<int8_t, int8_t>(graph,
-                                  scope,
-                                  block,
-                                  conv_filter_replicated_node,
-                                  &filter_intx,
-                                  &filter_max,
-                                  &scale_max,
-                                  false,
-                                  weight_scale);
+    if (quant_post_type.find("conv2d") != quant_post_type.end() &&
+            quant_post_type.find("conv2d")->second == 0 ||
+        quant_post_type.find("conv2d") != quant_post_type.end() &&
+            quant_post_type.find("conv2d")->second == 1) {
+      PrepareWeight<int8_t, int8_t>(graph,
+                                    scope,
+                                    block,
+                                    conv_filter_replicated_node,
+                                    &filter_intx,
+                                    &filter_max,
+                                    &scale_max,
+                                    false,
+                                    weight_scale);
+    } else {
+      VLOG(5) << "Unsupported type weight!";
+    }
   }
 
   (*fusion_nodes_map)["filter"] = filter_intx;
diff --git a/paddle/fluid/framework/ir/xpu/fc_xpu_fuse_pass.cc b/paddle/fluid/framework/ir/xpu/fc_xpu_fuse_pass.cc
@@ -299,8 +299,9 @@ void FcXPUFusePass::CreateTheReplicatedWeights(
       true,
       platform::errors::InvalidArgument("mul node ptr can not be null"));
   auto mul_w_name = mul->Op()->Input("Y")[0];
-  std::string replicated_w_name =
-      mul_w_name + "_copy_" + std::to_string(mul->id());
+  std::string replicated_w_name = mul_w_name + "_copy_" +
+                                  std::to_string(block->ID()) + "_" +
+                                  std::to_string(mul->id());
   auto* replicated_w_var = scope->FindVar(replicated_w_name);
   if (replicated_w_var == nullptr) {
     auto* filter_tensor =
@@ -395,8 +396,9 @@ void FcXPUFusePass::CreateFusionWeightsAndBias(
   auto mul_w_name = mul->Op()->Input("Y")[0];
   Node* mul_w = FindNodeWithName(graph, mul_w_name);
   CreateTheReplicatedWeights(graph, scope, block, nodes_map);
-  std::string replicated_w_name =
-      mul_w_name + "_copy_" + std::to_string(mul->id());
+  std::string replicated_w_name = mul_w_name + "_copy_" +
+                                  std::to_string(block->ID()) + "_" +
+                                  std::to_string(mul->id());
   auto* mul_w_replicated_node = FindNodeWithName(graph, replicated_w_name);
   // transfilter fp16 --> fp32
   auto* filter_t = scope->FindVar(mul_w_replicated_node->Name())
@@ -527,31 +529,83 @@ void FcXPUFusePass::CreateFusionWeightsAndBias(
   Node* filter_intx = nullptr;
   Node* filter_max = nullptr;
   Node* scale_max = nullptr;
-  bool per_channel_quant =
-      std::getenv("FLAGS_fc_gemm_use_per_channel") == nullptr ? false : true;
+
+  std::map<std::string, int> default_type;
+  default_type.insert(std::make_pair("fc", -1));
+  auto quant_post_type =
+      Has("quant_post_dynamic_weight_methods")
+          ? Get<std::map<std::string, int>>("quant_post_dynamic_weight_methods")
+          : default_type;
+
+  for (auto it = quant_post_type.begin(); it != quant_post_type.end(); ++it) {
+    VLOG(5) << "Key:" << it->first;
+    VLOG(5) << "Value:" << it->second;
+  }
+
   if (op_weights_precision != "int8") {
-    PrepareWeight<float, int16_t>(graph,
-                                  scope,
-                                  block,
-                                  mul_w_replicated_node,
-                                  &filter_intx,
-                                  &filter_max,
-                                  &scale_max,
-                                  !transpose_w,
-                                  weight_scale,
-                                  per_channel_quant);
+    if (quant_post_type.find("fc") != quant_post_type.end() &&
+            quant_post_type.find("fc")->second == 2 ||
+        quant_post_type.find("fc") != quant_post_type.end() &&
+            quant_post_type.find("fc")->second == -1) {
+      VLOG(5) << "Use int16 per-tensor weight";
+      PrepareWeight<float, int16_t>(graph,
+                                    scope,
+                                    block,
+                                    mul_w_replicated_node,
+                                    &filter_intx,
+                                    &filter_max,
+                                    &scale_max,
+                                    !transpose_w,
+                                    weight_scale,
+                                    false);
+    } else if (quant_post_type.find("fc") != quant_post_type.end() &&
+               quant_post_type.find("fc")->second == 3) {
+      VLOG(5) << "Use int16 per-channel weight";
+      PrepareWeight<float, int16_t>(graph,
+                                    scope,
+                                    block,
+                                    mul_w_replicated_node,
+                                    &filter_intx,
+                                    &filter_max,
+                                    &scale_max,
+                                    !transpose_w,
+                                    weight_scale,
+                                    true);
+    } else {
+      VLOG(5) << "Unsupported type weight by non-int8!";
+    }
   } else {
-    PrepareWeight<int8_t, int8_t>(graph,
-                                  scope,
-                                  block,
-                                  mul_w_replicated_node,
-                                  &filter_intx,
-                                  &filter_max,
-                                  &scale_max,
-                                  !transpose_w,
-                                  weight_scale,
-                                  per_channel_quant);
+    if (quant_post_type.find("fc") != quant_post_type.end() &&
+        quant_post_type.find("fc")->second == 0) {
+      VLOG(5) << "Use int8  per-tensor weight";
+      PrepareWeight<int8_t, int8_t>(graph,
+                                    scope,
+                                    block,
+                                    mul_w_replicated_node,
+                                    &filter_intx,
+                                    &filter_max,
+                                    &scale_max,
+                                    !transpose_w,
+                                    weight_scale,
+                                    false);
+    } else if (quant_post_type.find("fc") != quant_post_type.end() &&
+               quant_post_type.find("fc")->second == 1) {
+      VLOG(5) << "Use int8  per-channel weight";
+      PrepareWeight<int8_t, int8_t>(graph,
+                                    scope,
+                                    block,
+                                    mul_w_replicated_node,
+                                    &filter_intx,
+                                    &filter_max,
+                                    &scale_max,
+                                    !transpose_w,
+                                    weight_scale,
+                                    true);
+    } else {
+      VLOG(5) << "Unsupported type weight!";
+    }
   }
+
   (*fusion_nodes_map)["w"] = filter_intx;
   (*fusion_nodes_map)["w_max"] = filter_max;
   (*fusion_nodes_map)["scale_max"] = scale_max;
diff --git a/paddle/fluid/inference/analysis/argument.h b/paddle/fluid/inference/analysis/argument.h
@@ -333,6 +333,10 @@ struct Argument {
                       XpuFcAutotuneFileWriteback,
                       bool);
   DECL_ARGUMENT_FIELD(xpu_gemm_compute_precision, XpuGemmComputePrecision, int);
+  using quant_post_type = std::map<std::string, int>;
+  DECL_ARGUMENT_FIELD(xpu_quant_post_dynamic_weight_methods,
+                      XpuQuantPostDynamicWeightMethods,
+                      quant_post_type);
   DECL_ARGUMENT_FIELD(xpu_transformer_softmax_optimize_level,
                       XpuTransformerSoftmaxOptimizeLevel,
                       int);
diff --git a/paddle/fluid/inference/analysis/ir_pass_manager.cc b/paddle/fluid/inference/analysis/ir_pass_manager.cc
@@ -359,6 +359,20 @@ void IRPassManager::CreatePasses(Argument *argument,
       if (quant_post_dynamic_weight_precision == 0) {
         pass->Set("quant_post_dynamic_weight_precision ", new int(0));
       }
+    } else if (pass_name == "fc_xpu_fuse_pass") {
+      std::map<std::string, int> quant_post_type =
+          argument->xpu_quant_post_dynamic_weight_methods();
+      if (!quant_post_type.empty()) {
+        pass->Set("quant_post_dynamic_weight_methods",
+                  new std::map<std::string, int>(quant_post_type));
+      }
+    } else if (pass_name == "conv2d_xpu_fuse_pass") {
+      std::map<std::string, int> quant_post_type =
+          argument->xpu_quant_post_dynamic_weight_methods();
+      if (!quant_post_type.empty()) {
+        pass->Set("quant_post_dynamic_weight_methods",
+                  new std::map<std::string, int>(quant_post_type));
+      }
     }
     pre_pass = pass_name;
 
diff --git a/paddle/fluid/inference/api/analysis_predictor.cc b/paddle/fluid/inference/api/analysis_predictor.cc
@@ -1675,6 +1675,8 @@ void AnalysisPredictor::PrepareArgument() {
       config_.xpu_config_.fc_autotune_file_writeback);
   argument_->SetXpuGemmComputePrecision(
       config_.xpu_config_.gemm_compute_precision);
+  argument_->SetXpuQuantPostDynamicWeightMethods(
+      config_.xpu_config_.quant_post_dynamic_weight_methods);
   argument_->SetXpuTransformerSoftmaxOptimizeLevel(
       config_.xpu_config_.transformer_softmax_optimize_level);
   argument_->SetXpuTransformerEncoderAdaptiveSeqlen(
diff --git a/paddle/fluid/inference/api/paddle_analysis_config.h b/paddle/fluid/inference/api/paddle_analysis_config.h
@@ -145,6 +145,10 @@ struct PD_INFER_DECL XpuConfig {
   // Note: PaddleInference only.
   int quant_post_dynamic_weight_precision{1};
   std::vector<std::string> quant_post_dynamic_op_types;
+  // fc, conv2d
+  // 0: int8 per tensor, 1: int8 per-channel, 2: int16 per-tensor(default), 3:
+  // int16 per-channel, 4: int31 per-tensor. Note: PaddleInference only.
+  std::map<std::string, int> quant_post_dynamic_weight_methods;
 };
 
 struct DistConfig {
diff --git a/paddle/phi/kernels/xpu/plugin/src/kernel/kunlun2cpp/fast_reduce.xpu b/paddle/phi/kernels/xpu/plugin/src/kernel/kunlun2cpp/fast_reduce.xpu
@@ -112,19 +112,18 @@ __global__ void fast_reduce_sum_tiny(const T* x, T* y, int m, int t) {
   for (int i = mstart; i < mend; i += block_cnt) {
     int readlen = min((mend - i) * t, block_cnt * t);
     GM2LM(x + i * t, (T*)xlm, readlen * sizeof(T));
+    primitive_cast<T, float>((T*)xlm, xlm, readlen);
     if (t % 16 == 0 && t >= 32) {
       for (int j = 0; j < readlen; j += t) {
         ylm[j / t] = do_sum_align16(xlm + j, t);
       }
-      LM2GM((T*)ylm, y + i, readlen / t * sizeof(T));
     } else {
-      primitive_cast<T, float>((T*)xlm, xlm, readlen);
       for (int j = 0; j < readlen; j += t) {
         ylm[j / t] = do_sum(xlm + j, t);
       }
-      primitive_cast<float, T>(ylm, (T*)ylm, readlen / t);
-      LM2GM((T*)ylm, y + i, readlen / t * sizeof(T));
     }
+    primitive_cast<float, T>(ylm, (T*)ylm, readlen / t);
+    LM2GM((T*)ylm, y + i, readlen / t * sizeof(T));
   }
   return;
 }
@@ -147,19 +146,18 @@ __global__ void fast_reduce_mean_tiny(const T* x, T* y, int m, int t) {
   for (int i = mstart; i < mend; i += block_cnt) {
     int readlen = min((mend - i) * t, block_cnt * t);
     GM2LM(x + i * t, (T*)xlm, readlen * sizeof(T));
+    primitive_cast<T, float>((T*)xlm, xlm, readlen);
     if (t % 16 == 0 && t >= 32) {
       for (int j = 0; j < readlen; j += t) {
         ylm[j / t] = do_sum_align16(xlm + j, t) / t;
       }
-      LM2GM((T*)ylm, y + i, readlen / t * sizeof(T));
     } else {
-      primitive_cast<T, float>((T*)xlm, xlm, readlen);
       for (int j = 0; j < readlen; j += t) {
         ylm[j / t] = do_sum(xlm + j, t) / t;
       }
-      primitive_cast<float, T>(ylm, (T*)ylm, readlen / t);
-      LM2GM((T*)ylm, y + i, readlen / t * sizeof(T));
     }
+    primitive_cast<float, T>(ylm, (T*)ylm, readlen / t);
+    LM2GM((T*)ylm, y + i, readlen / t * sizeof(T));
   }
   return;
 }
@@ -182,19 +180,18 @@ __global__ void fast_reduce_max_tiny(const T* x, T* y, int m, int t) {
   for (int i = mstart; i < mend; i += block_cnt) {
     int readlen = min((mend - i) * t, block_cnt * t);
     GM2LM(x + i * t, (T*)xlm, readlen * sizeof(T));
+    primitive_cast<T, float>((T*)xlm, xlm, readlen);
     if (t % 16 == 0 && t >= 32) {
       for (int j = 0; j < readlen; j += t) {
         ylm[j / t] = do_max_align16(xlm + j, t);
       }
-      LM2GM((T*)ylm, y + i, readlen / t * sizeof(T));
     } else {
-      primitive_cast<T, float>((T*)xlm, xlm, readlen);
       for (int j = 0; j < readlen; j += t) {
         ylm[j / t] = do_max(xlm + j, t);
       }
-      primitive_cast<float, T>(ylm, (T*)ylm, readlen / t);
-      LM2GM((T*)ylm, y + i, readlen / t * sizeof(T));
     }
+    primitive_cast<float, T>(ylm, (T*)ylm, readlen / t);
+    LM2GM((T*)ylm, y + i, readlen / t * sizeof(T));
   }
   return;
 }
@@ -217,19 +214,18 @@ __global__ void fast_reduce_min_tiny(const T* x, T* y, int m, int t) {
   for (int i = mstart; i < mend; i += block_cnt) {
     int readlen = min((mend - i) * t, block_cnt * t);
     GM2LM(x + i * t, (T*)xlm, readlen * sizeof(T));
+    primitive_cast<T, float>((T*)xlm, xlm, readlen);
     if (t % 16 == 0 && t >= 32) {
       for (int j = 0; j < readlen; j += t) {
         ylm[j / t] = do_min_align16(xlm + j, t);
       }
-      LM2GM((T*)ylm, y + i, readlen / t * sizeof(T));
     } else {
-      primitive_cast<T, float>((T*)xlm, xlm, readlen);
       for (int j = 0; j < readlen; j += t) {
         ylm[j / t] = do_min(xlm + j, t);
       }
-      primitive_cast<float, T>(ylm, (T*)ylm, readlen / t);
-      LM2GM((T*)ylm, y + i, readlen / t * sizeof(T));
     }
+    primitive_cast<float, T>(ylm, (T*)ylm, readlen / t);
+    LM2GM((T*)ylm, y + i, readlen / t * sizeof(T));
   }
   return;
 }

Original file line number	Diff line number	Diff line change
`@@ -112,19 +112,18 @@ __global__ void fast_reduce_sum_tiny(const T* x, T* y, int m, int t) {`
`112`	`112`	`for (int i = mstart; i < mend; i += block_cnt) {`
`113`	`113`	`int readlen = min((mend - i) * t, block_cnt * t);`
`114`	`114`	`GM2LM(x + i * t, (T)xlm, readlen sizeof(T));`
	`115`	`+ primitive_cast<T, float>((T*)xlm, xlm, readlen);`
`115`	`116`	`if (t % 16 == 0 && t >= 32) {`
`116`	`117`	`for (int j = 0; j < readlen; j += t) {`
`117`	`118`	`ylm[j / t] = do_sum_align16(xlm + j, t);`
`118`	`119`	`}`
`119`		`- LM2GM((T)ylm, y + i, readlen / t sizeof(T));`
`120`	`120`	`} else {`
`121`		`- primitive_cast<T, float>((T*)xlm, xlm, readlen);`
`122`	`121`	`for (int j = 0; j < readlen; j += t) {`
`123`	`122`	`ylm[j / t] = do_sum(xlm + j, t);`
`124`	`123`	`}`
`125`		`- primitive_cast<float, T>(ylm, (T*)ylm, readlen / t);`
`126`		`- LM2GM((T)ylm, y + i, readlen / t sizeof(T));`
`127`	`124`	`}`
	`125`	`+ primitive_cast<float, T>(ylm, (T*)ylm, readlen / t);`
	`126`	`+ LM2GM((T)ylm, y + i, readlen / t sizeof(T));`
`128`	`127`	`}`
`129`	`128`	`return;`
`130`	`129`	`}`
`@@ -147,19 +146,18 @@ __global__ void fast_reduce_mean_tiny(const T* x, T* y, int m, int t) {`
`147`	`146`	`for (int i = mstart; i < mend; i += block_cnt) {`
`148`	`147`	`int readlen = min((mend - i) * t, block_cnt * t);`
`149`	`148`	`GM2LM(x + i * t, (T)xlm, readlen sizeof(T));`
	`149`	`+ primitive_cast<T, float>((T*)xlm, xlm, readlen);`
`150`	`150`	`if (t % 16 == 0 && t >= 32) {`
`151`	`151`	`for (int j = 0; j < readlen; j += t) {`
`152`	`152`	`ylm[j / t] = do_sum_align16(xlm + j, t) / t;`
`153`	`153`	`}`
`154`		`- LM2GM((T)ylm, y + i, readlen / t sizeof(T));`
`155`	`154`	`} else {`
`156`		`- primitive_cast<T, float>((T*)xlm, xlm, readlen);`
`157`	`155`	`for (int j = 0; j < readlen; j += t) {`
`158`	`156`	`ylm[j / t] = do_sum(xlm + j, t) / t;`
`159`	`157`	`}`
`160`		`- primitive_cast<float, T>(ylm, (T*)ylm, readlen / t);`
`161`		`- LM2GM((T)ylm, y + i, readlen / t sizeof(T));`
`162`	`158`	`}`
	`159`	`+ primitive_cast<float, T>(ylm, (T*)ylm, readlen / t);`
	`160`	`+ LM2GM((T)ylm, y + i, readlen / t sizeof(T));`
`163`	`161`	`}`
`164`	`162`	`return;`
`165`	`163`	`}`
`@@ -182,19 +180,18 @@ __global__ void fast_reduce_max_tiny(const T* x, T* y, int m, int t) {`
`182`	`180`	`for (int i = mstart; i < mend; i += block_cnt) {`
`183`	`181`	`int readlen = min((mend - i) * t, block_cnt * t);`
`184`	`182`	`GM2LM(x + i * t, (T)xlm, readlen sizeof(T));`
	`183`	`+ primitive_cast<T, float>((T*)xlm, xlm, readlen);`
`185`	`184`	`if (t % 16 == 0 && t >= 32) {`
`186`	`185`	`for (int j = 0; j < readlen; j += t) {`
`187`	`186`	`ylm[j / t] = do_max_align16(xlm + j, t);`
`188`	`187`	`}`
`189`		`- LM2GM((T)ylm, y + i, readlen / t sizeof(T));`
`190`	`188`	`} else {`
`191`		`- primitive_cast<T, float>((T*)xlm, xlm, readlen);`
`192`	`189`	`for (int j = 0; j < readlen; j += t) {`
`193`	`190`	`ylm[j / t] = do_max(xlm + j, t);`
`194`	`191`	`}`
`195`		`- primitive_cast<float, T>(ylm, (T*)ylm, readlen / t);`
`196`		`- LM2GM((T)ylm, y + i, readlen / t sizeof(T));`
`197`	`192`	`}`
	`193`	`+ primitive_cast<float, T>(ylm, (T*)ylm, readlen / t);`
	`194`	`+ LM2GM((T)ylm, y + i, readlen / t sizeof(T));`
`198`	`195`	`}`
`199`	`196`	`return;`
`200`	`197`	`}`
`@@ -217,19 +214,18 @@ __global__ void fast_reduce_min_tiny(const T* x, T* y, int m, int t) {`
`217`	`214`	`for (int i = mstart; i < mend; i += block_cnt) {`
`218`	`215`	`int readlen = min((mend - i) * t, block_cnt * t);`
`219`	`216`	`GM2LM(x + i * t, (T)xlm, readlen sizeof(T));`
	`217`	`+ primitive_cast<T, float>((T*)xlm, xlm, readlen);`
`220`	`218`	`if (t % 16 == 0 && t >= 32) {`
`221`	`219`	`for (int j = 0; j < readlen; j += t) {`
`222`	`220`	`ylm[j / t] = do_min_align16(xlm + j, t);`
`223`	`221`	`}`
`224`		`- LM2GM((T)ylm, y + i, readlen / t sizeof(T));`
`225`	`222`	`} else {`
`226`		`- primitive_cast<T, float>((T*)xlm, xlm, readlen);`
`227`	`223`	`for (int j = 0; j < readlen; j += t) {`
`228`	`224`	`ylm[j / t] = do_min(xlm + j, t);`
`229`	`225`	`}`
`230`		`- primitive_cast<float, T>(ylm, (T*)ylm, readlen / t);`
`231`		`- LM2GM((T)ylm, y + i, readlen / t sizeof(T));`
`232`	`226`	`}`
	`227`	`+ primitive_cast<float, T>(ylm, (T*)ylm, readlen / t);`
	`228`	`+ LM2GM((T)ylm, y + i, readlen / t sizeof(T));`
`233`	`229`	`}`
`234`	`230`	`return;`
`235`	`231`	`}`