[update]update ver:3.0.7

Author: yyzhang

doc/tutorial/release.md

@@ -17,4 +17,7 @@
 # V3.0.6  2026.03.26
 * 解决linger.init之后加载浮点模型dict不匹配问题
 * 解决avgpool算子一致性问题
-* 解决gru算子一致性问题
+* 解决gru算子一致性问题
+# V3.0.7  2026.04.30
+* 解决venusA和arcs平台gru一致性问题；
+* 重构导图后处理代码，增加图优化和常量折叠功能；

linger/__version.py

@@ -8,5 +8,5 @@ def _to_int(s):
         return s
 
 
-__version__ = "3.0.6"
+__version__ = "3.0.7"
 version_info = tuple(_to_int(s) for s in __version__.split("."))

linger/constrain/clayernorm.py

@@ -20,7 +20,6 @@ def ccreate(
             module.normalized_shape,
             module.eps,
             module.elementwise_affine,
-            None if module.bias is None else True,
             dtype=module.weight.dtype,
             device=device,
             constrain=constrain,

linger/initialize.py

@@ -85,29 +85,31 @@ def quant_module(module: nn.Module, c_activation_val: float = 8.0, c_weight_val:
             m.output_quantizer.data_bits = out_bits
             m.output_quantizer.clamp_activation_value = c_activation_val
 
-def constrain(model: nn.Module, config_file: str = None, disable_module=None, disable_submodel=[]):
+def constrain(model: nn.Module, config_file: str = None, disable_module=None, disable_submodel=None):
     c_configs = QUANT_CONFIGS
     if config_file is not None:
         c_configs._load_from_yaml(config_file)
+    disabled_submodels = [] if disable_submodel is None else list(disable_submodel)
 
     if disable_module is not None:
         for name in disable_module:
             if _CMODULE_TABLE.get(name, None) is not None:
                 _CMODULE_TABLE.pop(name)
 
     for name, m in model.named_modules():
-        if any(name.startswith(p) for p in disable_submodel): # disable_submodel 直接按照module的名称进行屏蔽
+        if any(name.startswith(p) for p in disabled_submodels): # disable_submodel 直接按照module的名称进行屏蔽
             continue
         _constrain_submodule(model, name, m, c_configs.clamp_info.to_dict())
 
     model.to(c_configs.device)
     return model
 
-def init(model: nn.Module, config_file: str = None, disable_module=None, disable_submodel=[]):
+def init(model: nn.Module, config_file: str = None, disable_module=None, disable_submodel=None):
 
     q_configs = QUANT_CONFIGS
     if config_file is not None:
         q_configs._load_from_yaml(config_file)
+    disabled_submodels = [] if disable_submodel is None else list(disable_submodel)
 
     if disable_module is not None:
         for name in disable_module:
@@ -118,15 +120,15 @@ def init(model: nn.Module, config_file: str = None, disable_module=None, disable
     # model = _replace_ops(traced_model, q_configs)
 
     for name, m in model.named_modules():
-        if any(name.startswith(p) for p in disable_submodel): # disable_submodel 直接按照module的名称进行量化屏蔽
+        if any(name.startswith(p) for p in disabled_submodels): # disable_submodel 直接按照module的名称进行量化屏蔽
             continue
 
         m.register_forward_pre_hook(hook_pre_forward)
         m.register_forward_hook(hook_forward)
 
         is_replaced = _quantize_submodule(model, name, m, weights_cfg=q_configs.quant_info.to_dict(), activations_cfg=q_configs.quant_info.to_dict(), bias_cfg=q_configs.quant_info.to_dict(), constrain =  q_configs.clamp_info.to_dict())
         if is_replaced:
-            disable_submodel.append(name)
+            disabled_submodels.append(name)
 
     def quant_tensor_pre_hook(state_dict, prefix, local_metadata, strict, missing_keys, unexpected_keys, error_msgs):
         quantizer_tags = (
@@ -192,6 +194,14 @@ def quant_tensor_layer(module, prefix=''):
                                 iq_layer.training = model.training
                                 # iq_layer = iq_layer.to(device)
                                 setattr(module, input_name, iq_layer)
+                            elif '_qgelu_' in input_name:
+                                iq_layer = QGelu(activate_config=activate_cfg, num_input=1)
+                                iq_layer.training = model.training
+                                setattr(module, input_name, iq_layer)
+                            elif '_qswish_' in input_name:
+                                iq_layer = QSwish(activate_config=activate_cfg, num_input=1)
+                                iq_layer.training = model.training
+                                setattr(module, input_name, iq_layer)
                             elif '_qsoftmax_' in input_name:
                                 iq_layer = QSoftmax(activate_config=activate_cfg, num_input=1)
                                 iq_layer.training = model.training

linger/kernel/cpu/venusa_qsigmoid_kernel.cpp

@@ -11,6 +11,21 @@
   ((x) > MAX_BITS(bits) ? MAX_BITS(bits) \
                         : ((x) < MIN_BITS(bits) ? MIN_BITS(bits) : (x)))
 
+static int64_t shfit_floor_x05_int64(int64_t x, int32_t shift)
+{
+	int64_t val = x;
+
+	if (shift >= 64) {
+		return 0;
+	}
+	if (shift > 0) {
+		val = val >> (shift - 1);
+		val = (val & 0x1) + (val >> 1);
+	}
+
+	return val;
+}
+
 torch::Tensor venusa_qsigmoid_cpu(torch::Tensor a)
 {
 	int32_t N = a.numel();
@@ -29,7 +44,6 @@ torch::Tensor venusa_qsigmoid_cpu(torch::Tensor a)
 	int64_t absx = 0;
 	int64_t slope = 0;
 	int64_t bias = 0;
-	int32_t shift = 0;
 	int64_t tmp = 0;
 	int64_t out = 0;
 
@@ -73,8 +87,8 @@ torch::Tensor venusa_qsigmoid_cpu(torch::Tensor a)
 				bias = 0;
 			}
 		}
-        bias = bias << 3;
-		out = ((slope * tmp) >> 27) + bias;
+		out = slope * tmp + (bias << 30);
+		out = shfit_floor_x05_int64(out, 27);
 		c_ptr[i] = SATURATE(out, 32);
 	}
 

linger/kernel/cpu/venusa_qsoftmax_kernel.cpp

@@ -21,7 +21,7 @@ static int32_t shift_pure(int64_t v, int32_t s)
 	} else {
         v = v >> (-s);
     }
-    return SATURATE(v, 32);
+    return v;
 }
 
 static int32_t sub32s(int32_t a, int32_t b)
@@ -150,7 +150,7 @@ torch::Tensor venusa_qsoftmax_cpu(const torch::Tensor& in, int64_t dim)
 		for (int i = 0; i < N; i++)
 		{
 			data = sub32s(x_ptr[i], max_value);
-            X = shfit_floor_x05_int64((int64_t)X * (int64_t)774541002, 31);
+            X = shfit_floor_x05_int64((int64_t)data * (int64_t)774541002, 31);
 			// X = shift_rasyms((int64_t)data * (int64_t)774541002, -31);//exp=>2xp，Q6.25=>Q8.23
 			E = X >> 23;
 			E = E + 1;//与118行对应
@@ -201,4 +201,3 @@ torch::Tensor venusa_qsoftmax_cpu(const torch::Tensor& in, int64_t dim)
 
 	return out;
 }
-

linger/kernel/cpu/venusa_qtanh_kernel.cpp

@@ -11,6 +11,21 @@
   ((x) > MAX_BITS(bits) ? MAX_BITS(bits) \
                         : ((x) < MIN_BITS(bits) ? MIN_BITS(bits) : (x)))
 
+static int64_t shfit_floor_x05_int64(int64_t x, int32_t shift)
+{
+	int64_t val = x;
+
+	if (shift >= 64) {
+		return 0;
+	}
+	if (shift > 0) {
+		val = val >> (shift - 1);
+		val = (val & 0x1) + (val >> 1);
+	}
+
+	return val;
+}
+
 torch::Tensor venusa_qtanh_cpu(torch::Tensor a)
 {
 	int32_t N = a.numel();
@@ -28,7 +43,6 @@ torch::Tensor venusa_qtanh_cpu(torch::Tensor a)
 	int64_t absx = 0;
 	int64_t slope = 0;
 	int64_t bias = 0;
-	int32_t shift = 0;
 	int64_t tmp = 0;
 	int64_t out = 0;
 
@@ -68,16 +82,15 @@ torch::Tensor venusa_qtanh_cpu(torch::Tensor a)
 				bias = 0;
 			}
 		}
-
-        bias = bias << 3;
 		if (1 == sign)
 		{
-			out = ((-1 * slope * absx) >> 27) - bias;
+			out = (-slope * absx) - (bias << 30);
 		}
 		else
 		{
-			out = ((slope * absx) >> 27) + bias;
+			out = (slope * absx) + (bias << 30);
 		}
+		out = shfit_floor_x05_int64(out, 27);
 
 		c_ptr[i] = SATURATE(out, 32);
 	}

linger/kernel/gpu/fake_quant_kernel.cu

@@ -44,7 +44,7 @@ std::tuple<torch::Tensor, torch::Tensor, float> fake_quant_cuda(
     // printf("到位置 2 了 \n");
     // 计算 scale
     float f = (float)(bit - 1) - factor;
-    float scale = powf(2.0f, roundf(f));
+    float scale = powf(2.0f, fminf(fmaxf(roundf(f), 0.0f), 23.0f));
     // scale = fminf(fmaxf(scale, 1e-6f), powf(2.0f, 32));
     // printf("到位置 1 了,scale:%f \n", scale);
     auto output = torch::empty_like(input);
@@ -84,7 +84,7 @@ std::tuple<torch::Tensor, torch::Tensor, float> fake_quant_cuda(
 std::tuple<torch::Tensor, torch::Tensor> bias_quant_cuda(
     torch::Tensor input,
     int bit,
-    float scale,
+    float scale, // 这里scale是在python代码里进行clamp到0-31的
     float scale_min,
     float quant_min,
     float quant_max) {
@@ -175,7 +175,7 @@ std::tuple<torch::Tensor, torch::Tensor, torch::Tensor, float> fake_quant_cuda_w
     // printf("到位置 2 了 \n");
     // 计算 scale
     float f = (float)(bit - 1) - factor;
-    float scale = powf(2.0f, roundf(f));
+    float scale = powf(2.0f, fminf(fmaxf(roundf(f), 0.0f), 23.0f));
     // scale = fminf(fmaxf(scale, 1e-6f), powf(2.0f, 32));
     // printf("到位置 1 了,scale:%f \n", scale);
     auto output = torch::empty_like(input);
@@ -218,7 +218,7 @@ std::tuple<torch::Tensor, torch::Tensor, torch::Tensor, float> fake_quant_cuda_w
 std::tuple<torch::Tensor, torch::Tensor, torch::Tensor> bias_quant_cuda_with_grad_scale(
     torch::Tensor input,
     int bit,
-    float scale,
+    float scale, // 这里的scale是在python里clamp的
     float scale_min,
     float quant_min,
     float quant_max) {

linger/kernel/gpu/venusa_qsigmoid_kernel.cu

@@ -10,6 +10,21 @@
   ((x) > MAX_BITS(bits) ? MAX_BITS(bits) \
                         : ((x) < MIN_BITS(bits) ? MIN_BITS(bits) : (x)))
 
+static __device__ int64_t shfit_floor_x05_int64(int64_t x, int32_t shift)
+{
+	int64_t val = x;
+
+	if (shift >= 64) {
+		return 0;
+	}
+	if (shift > 0) {
+		val = val >> (shift - 1);
+		val = (val & 0x1) + (val >> 1);
+	}
+
+	return val;
+}
+
 __global__ void venusa_qsigmoid_gpu_kernel(const int* __restrict__ a,
                              int* __restrict__ c, 
                             int32_t len, uint32_t* bands , uint32_t* slopes,
@@ -21,7 +36,6 @@ __global__ void venusa_qsigmoid_gpu_kernel(const int* __restrict__ a,
     int64_t absx = 0;
     int64_t slope = 0;
     int64_t bias = 0;
-    int32_t shift = 0;
     int64_t tmp = 0;
     int64_t out = 0;
 
@@ -65,12 +79,11 @@ __global__ void venusa_qsigmoid_gpu_kernel(const int* __restrict__ a,
 			{
 				slope = 0;
 				bias = 0;
-				shift = 0;
 			}
 		}
 
-        bias = bias << 3;
-		out = ((slope * tmp) >> 27) + bias;
+		out = slope * tmp + (bias << 30);
+		out = shfit_floor_x05_int64(out, 27);
 		c[idx] = SATURATE(out, 32);
     }
 }
@@ -111,4 +124,3 @@ torch::Tensor venusa_qsigmoid_gpu(torch::Tensor a)
 
     return c;
 }
-

linger/kernel/gpu/venusa_qsoftmax_kernel.cu

@@ -22,7 +22,7 @@ static __device__ int32_t shift_pure(int64_t v, int32_t s)
 	} else {
 		v = v >> (-s);
 	}
-	return SATURATE(v, 32);
+	return v;
 }
 
 static __device__ int32_t sub32s(int32_t a, int32_t b)
@@ -98,6 +98,7 @@ static __device__ int64_t shfit_floor_x05_int64(int64_t x, int32_t shift)
 static __device__ void venusa_qsoftmax_c(const int* __restrict__ x_ptr, int* __restrict__ y_ptr, int32_t N, int32_t* p23)
 {
 	int32_t max_value = x_ptr[0];
+	int32_t data = 0;
 	int32_t X = 0;
 	int64_t Y = 0;
 	int32_t E = 0;
@@ -112,8 +113,8 @@ static __device__ void venusa_qsoftmax_c(const int* __restrict__ x_ptr, int* __r
 
 	for (int i = 0; i < N; i++)
 	{
-		X = sub32s(x_ptr[i] ,max_value);
-        X = shfit_floor_x05_int64((int64_t)X * (int64_t)774541002, 31);
+		data = sub32s(x_ptr[i] ,max_value);
+        X = shfit_floor_x05_int64((int64_t)data * (int64_t)774541002, 31);
 		// X = shift_rasyms((int64_t)X * (int64_t)774541002,-31);
 		E = X >> 23;
 		E = E + 1;
@@ -207,4 +208,3 @@ torch::Tensor venusa_qsoftmax_gpu(const torch::Tensor& in, int64_t dim)
 
 	return out;
 }
-