Skip to content

Commit e151543

Browse files
unamedkrclaude
unamedkr
and
claude
committed
v0.5: Bin-centered quantization (floor + midpoint reconstruction)
Switch from round-to-nearest/edge-aligned to floor/bin-centered scheme, matching the Triton reference (refs/PolarQuant/models/kernel4group.py). Algorithm change: - Quantize: scale = range/levels (was range/(levels-1)), use floor (was round) - Dequantize: val = min + (q + 0.5) * scale (was min + q * scale) - PolarQuant: angle range shifted to [0, 2pi] (was [-pi, pi]) Applied uniformly across all backends: - Core C reference (tq_polar.c, tq_uniform.c) - ARM NEON (tq_neon.c) - x86 AVX2 (tq_avx2.c) - CUDA kernels (tq_polar.cu, tq_value.cu, tq_turbo.cu, tq_fused_cache.cu) - Metal shaders (tq_polar.metal, tq_value.metal, tq_turbo.metal, tq_fused_cache.metal) A/B test results on real model data: - uniform_2b MSE: 0.069 → 0.040 (-42% improvement) - uniform_4b MSE: 0.00255 → 0.00247 (-3% improvement) - uniform_2b cosine: 0.827 → 0.855 (+3.4% improvement) - Synthetic roundtrip_mse: 0.0014 → 0.0013 (improved) 13/13 C++ tests pass, score 99.7% Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
1 parent 86a451e commit e151543

15 files changed

Lines changed: 375 additions & 121 deletions

‎docs/prd_v0.5.md‎

Lines changed: 117 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,117 @@
1+
# TurboQuant.cpp — Product Requirements Document v0.5
2+
3+
**Version**: 0.5
4+
**Date**: 2026-03-29
5+
**Focus**: 극한 품질 최적화 — refs 코드와의 차이를 0으로 만든다
6+
7+
---
8+
9+
## 1. v0.5 Goal
10+
11+
v0.4까지 프로덕션 안정성을 확보했다. v0.5는 **논문 구현과의 알고리즘 차이를 좁혀 품질을 극한까지 끌어올린다**. 모든 개선은 A/B 테스트로 측정하고, 측정 불가능한 변경은 하지 않는다.
12+
13+
### 발견된 알고리즘 차이 (refs 코드 대비)
14+
15+
| # | 차이 | 파일 | 예상 임팩트 | 난이도 |
16+
|---|------|------|-----------|--------|
17+
| ALG-1 | PolarQuant 각도 범위 [-π,π] → [0,2π] 미적용 | tq_polar.c | MSE -5% | 1줄 |
18+
| ALG-2 | 역양자화 시 +0.5 중심 보정 미적용 | tq_polar.c | MSE -15~20% | 1줄 |
19+
| ALG-3 | QJL inlier/outlier 이중 스트림 스코어 미적용 | tq_qjl.c | Cosine +10~15% | 중 |
20+
| ALG-4 | QJL 그룹 기반 아웃라이어 (개별→그룹) | tq_qjl.c | Cosine +5% | 중 |
21+
| ALG-5 | Value 양자화 per-tile 스케일링 (vLLM 패턴) | tq_uniform.c | Value MSE -10% | 중 |
22+
| ALG-6 | Key/Value 분리 양자화 전략 | tq_types.h | 아키텍처 정확성 | 소 |
23+
24+
**핵심 원칙**: ALG-1, ALG-2는 각각 1줄 수정으로 즉시 품질 개선이 가능. 먼저 적용하고 측정한다.
25+
26+
---
27+
28+
## 2. Functional Requirements
29+
30+
### FR-V5-1: PolarQuant 정밀도 개선 (ALG-1 + ALG-2)
31+
32+
**각도 범위 수정** (ALG-1):
33+
```c
34+
// Before (tq_polar.c):
35+
float t = atan2f(y, x); // [-π, π]
36+
37+
// After:
38+
float t = atan2f(y, x);
39+
if (t < 0.0f) t += 2.0f * TQ_PI; // [0, 2π]
40+
```
41+
이유: [-π,π]에서 min-max 양자화하면 π 근처에서 wrap-around 오차 발생. [0,2π]로 이동하면 연속적.
42+
43+
**중심 보정** (ALG-2):
44+
```c
45+
// Before (dequantize):
46+
float theta_f = tscale * (float)theta_idx + tmn;
47+
48+
// After:
49+
float theta_f = tscale * ((float)theta_idx + 0.5f) + tmn;
50+
float radius_f = rscale * ((float)rho_idx + 0.5f) + rmn;
51+
```
52+
이유: 양자화 bin의 가장자리가 아닌 중심으로 복원하면 평균 오차가 절반.
53+
54+
**검증**: MSE 측정 — before vs after on 실제 모델 데이터.
55+
56+
### FR-V5-2: QJL 이중 스트림 스코어 (ALG-3)
57+
58+
현재: `score = sqrt(π/2)/S × norm × (S - 2×hamming)`
59+
목표: `score = sqrt(π/2)/S × norm_inlier × hamming_inlier + sqrt(π/2)/S_out × norm_outlier × hamming_outlier`
60+
61+
핵심 변경:
62+
1. 양자화 시: inlier와 outlier를 분리하여 각각 별도 hash 생성
63+
2. 쿼리 투영 시: 전체 sketch에서 outlier 기여분을 빼서 inlier sketch 계산
64+
3. 점수 계산: 두 스트림의 점수를 각각 norm-weighted하여 합산
65+
66+
참조: `refs/QJL/qjl_kernel/csrc/qjl_score_kernel.cu` line 130
67+
68+
**검증**: Cosine similarity — QJL attention on 실제 모델 데이터.
69+
70+
### FR-V5-3: Value Per-Tile Scaling (ALG-5)
71+
72+
현재: 128개 요소에 단일 scale
73+
목표: 128개 요소를 32개씩 4개 타일로 나눠 각각 scale
74+
75+
```c
76+
typedef struct {
77+
uint16_t scale[4]; // 4 tiles × fp16 scale
78+
uint16_t zero_point[4]; // 4 tiles × fp16 zero
79+
uint8_t qs[TQ_BK / 2]; // 4-bit packed values
80+
} block_tq_value_tiled;
81+
```
82+
83+
참조: `refs/vllm/csrc/cache_kernels.cu` line 400-418
84+
85+
**검증**: Value roundtrip MSE — per-tile vs per-block on 실제 모델 데이터.
86+
87+
### FR-V5-4: Key/Value 분리 전략 (ALG-6)
88+
89+
Progressive config에 별도 key_type, value_type 설정:
90+
```c
91+
typedef struct {
92+
tq_type key_warm_type; // Keys: PolarQuant or TurboQuant
93+
tq_type value_warm_type; // Values: Uniform (amplitude preservation)
94+
tq_type key_cold_type;
95+
tq_type value_cold_type;
96+
...
97+
} tq_progressive_config_t;
98+
```
99+
100+
이유: 키는 방향 보존이 중요 (PolarQuant/QJL), 값은 진폭 보존이 중요 (Uniform).
101+
102+
---
103+
104+
## 3. Success Criteria
105+
106+
모든 측정은 `build/real_model_validation` 기준:
107+
108+
| 지표 | v0.4 현재 | v0.5 목표 |
109+
|------|----------|----------|
110+
| uniform_4b real cosine | 0.991 | > 0.993 |
111+
| polar_4b real cosine | 0.786 | > 0.85 (ALG-1+2 적용) |
112+
| turbo_3b real cosine | 0.939 | > 0.95 |
113+
| QJL real cosine | 0.857 | > 0.90 (ALG-3 적용) |
114+
| uniform_4b real MSE | 0.0025 | < 0.002 |
115+
| polar_4b real MSE | 0.053 | < 0.03 (ALG-1+2 적용) |
116+
117+
**모든 개선은 A/B 테스트 증거가 있어야 머지한다.**

‎docs/wbs_v0.5.md‎

Lines changed: 124 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,124 @@
1+
# TurboQuant.cpp — Work Breakdown Structure v0.5
2+
3+
**Version**: 0.5
4+
**Date**: 2026-03-29
5+
**Focus**: 극한 품질 최적화 — 모든 변경을 A/B 테스트로 측정
6+
7+
---
8+
9+
## Phase 1: 1줄 수정 고임팩트 (ALG-1 + ALG-2)
10+
11+
### 1.1 PolarQuant 각도 범위 [0, 2π]
12+
13+
- [ ] `src/core/tq_polar.c` — quantize 함수에서 atan2 결과 보정
14+
- [ ] `if (t < 0.0f) t += 2.0f * TQ_PI;` 추가
15+
- [ ] A/B 측정: `build/real_model_validation` 실행
16+
- [ ] Before polar_4b MSE: ____
17+
- [ ] After polar_4b MSE: ____
18+
- [ ] 개선율: ____%
19+
20+
### 1.2 중심 보정 (+0.5 offset)
21+
22+
- [ ] `src/core/tq_polar.c` — dequantize 함수에서 +0.5 추가
23+
- [ ] theta: `tscale * ((float)tq + 0.5f) + tmn`
24+
- [ ] radius: `rscale * ((float)rq + 0.5f) + rmn`
25+
- [ ] `src/core/tq_polar.c` — attention 함수의 LUT에도 +0.5 반영
26+
- [ ] `cos_lut[q] = cosf(tscale * ((float)q + 0.5f) + tmn)`
27+
- [ ] `src/core/tq_uniform.c` — uniform dequantize에도 +0.5 추가
28+
- [ ] `val = scale * ((float)q + 0.5f) + zero_point`
29+
- [ ] A/B 측정
30+
- [ ] Before uniform_4b MSE: ____
31+
- [ ] After uniform_4b MSE: ____
32+
33+
### 1.3 검증
34+
35+
- [ ] 모든 기존 테스트 통과 (테스트 tolerance 조정 필요할 수 있음)
36+
- [ ] `build/tq_quality` 결과 개선 확인
37+
- [ ] `build/ab_test` 결과 개선 확인
38+
39+
---
40+
41+
## Phase 2: QJL 이중 스트림 (ALG-3)
42+
43+
### 2.1 양자화 분리
44+
45+
- [ ] `src/core/tq_qjl.c` — `tq_qjl_quantize_ref()` 수정
46+
- [ ] 아웃라이어 차원에 대해 별도 hash 계산 (outlier_hash)
47+
- [ ] 인라이어 = 전체 투영에서 아웃라이어 투영을 뺀 것
48+
- [ ] `block_tq_qjl`에 outlier hash 추가 (또는 기존 hash와 분리)
49+
50+
### 2.2 어텐션 이중 스트림
51+
52+
- [ ] `src/core/tq_qjl.c` — `tq_qjl_attention_ref()` 수정
53+
- [ ] Query sketch 계산
54+
- [ ] Query outlier sketch 계산 (outlier 차원만)
55+
- [ ] Inlier sketch = query_sketch - query_outlier_sketch
56+
- [ ] Inlier score = sqrt(π/2)/S × norm_inlier × (S - 2×hamming_inlier)
57+
- [ ] Outlier score = sqrt(π/2)/S_out × norm_outlier × (S_out - 2×hamming_outlier)
58+
- [ ] Total = inlier_score + outlier_score
59+
60+
### 2.3 검증
61+
62+
- [ ] A/B 측정: QJL cosine on real model data
63+
- [ ] Before: ____
64+
- [ ] After: ____
65+
- [ ] 모든 QJL 테스트 통과
66+
67+
---
68+
69+
## Phase 3: Value Per-Tile Scaling (ALG-5)
70+
71+
### 3.1 타일 구조 추가
72+
73+
- [ ] `include/turboquant/tq_types.h` — `block_tq_value_tiled` 구조체 정의
74+
- [ ] 4개 타일 × (scale + zero_point) = 16 bytes metadata
75+
- [ ] qs[TQ_BK/2] = 64 bytes data
76+
- [ ] 총 80 bytes (기존 68 bytes 대비 12 bytes 증가)
77+
78+
### 3.2 타일 양자화 구현
79+
80+
- [ ] `src/core/tq_value_quant.c` — `tq_value_quantize_tiled()` 구현
81+
- [ ] 128 요소를 32개씩 4개 타일로 분할
82+
- [ ] 각 타일마다 독립 min/max → scale/zero_point
83+
- [ ] 4-bit 양자화 + 패킹
84+
85+
### 3.3 검증
86+
87+
- [ ] A/B 측정: Value roundtrip MSE
88+
- [ ] Per-block MSE: ____
89+
- [ ] Per-tile MSE: ____
90+
- [ ] 개선율: ____%
91+
92+
---
93+
94+
## Phase 4: 통합 + 최종 벤치마크
95+
96+
### 4.1 Key/Value 분리 설정
97+
98+
- [ ] `include/turboquant/tq_types.h` — progressive config 확장
99+
- [ ] `key_warm_type`, `value_warm_type` 분리
100+
- [ ] `tq_progressive_default_config()` 업데이트
101+
102+
### 4.2 최종 A/B 비교 벤치마크
103+
104+
- [ ] `bench/ab_comparison_v05.cpp` — v0.4 vs v0.5 비교 벤치마크
105+
- [ ] 모든 타입 × 실제 모델 데이터
106+
- [ ] Before/After MSE, Cosine 비교 테이블
107+
- [ ] 개선율 요약
108+
109+
### 4.3 문서 업데이트
110+
111+
- [ ] `docs/real_model_results.md` — v0.5 결과 추가
112+
- [ ] README — 최신 수치 반영
113+
114+
---
115+
116+
## 완료 기준
117+
118+
- [ ] polar_4b real cosine > 0.85 (현재 0.786)
119+
- [ ] QJL real cosine > 0.90 (현재 0.857)
120+
- [ ] turbo_3b real cosine > 0.95 (현재 0.939)
121+
- [ ] uniform_4b real MSE < 0.002 (현재 0.0025)
122+
- [ ] 모든 개선에 A/B 측정 수치 기록
123+
- [ ] 13+ C++ 테스트 전체 통과
124+
- [ ] score.sh ≥ 0.99 유지

‎src/backend/cpu/tq_avx2.c‎

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -98,7 +98,7 @@ void tq_uniform_4b_quantize_avx2(const float* src, void* dst, int n) {
9898

9999
float range = mx - mn;
100100
if (range < 1e-8f) range = 1e-8f;
101-
float scale = range / 15.0f;
101+
float scale = range / 16.0f; /* 16 bins of width range/16 */
102102
float inv_scale = 1.0f / scale;
103103

104104
block->scale = avx_fp32_to_fp16(scale);
@@ -116,8 +116,8 @@ void tq_uniform_4b_quantize_avx2(const float* src, void* dst, int n) {
116116
__m256 v = _mm256_loadu_ps(src + i);
117117
__m256 shifted = _mm256_sub_ps(v, v_mn);
118118
__m256 scaled = _mm256_mul_ps(shifted, v_invs);
119-
/* Round to nearest: _mm256_round_ps with _MM_FROUND_TO_NEAREST_INT */
120-
__m256 rounded = _mm256_round_ps(scaled, _MM_FROUND_TO_NEAREST_INT | _MM_FROUND_NO_EXC);
119+
/* Floor: _mm256_round_ps with _MM_FROUND_TO_NEG_INF */
120+
__m256 rounded = _mm256_round_ps(scaled, _MM_FROUND_TO_NEG_INF | _MM_FROUND_NO_EXC);
121121
/* Clamp to [0, 15] */
122122
rounded = _mm256_max_ps(rounded, v_zero);
123123
rounded = _mm256_min_ps(rounded, v_15);
@@ -139,7 +139,7 @@ void tq_uniform_4b_quantize_avx2(const float* src, void* dst, int n) {
139139

140140
/* Scalar tail */
141141
for (; i < count; i++) {
142-
int q = (int)roundf((src[i] - mn) * inv_scale);
142+
int q = (int)floorf((src[i] - mn) * inv_scale);
143143
if (q < 0) q = 0;
144144
if (q > 15) q = 15;
145145
if (i % 2 == 0) {
@@ -172,11 +172,11 @@ void tq_uniform_4b_dequantize_avx2(const void* src, float* dst, int n) {
172172
for (int k = 0; k < 8; k++) {
173173
int idx = i + k;
174174
uint8_t byte = block->qs[idx / 2];
175-
q_arr[k] = (float)((idx % 2 == 0) ? (byte & 0x0F) : (byte >> 4));
175+
q_arr[k] = (float)((idx % 2 == 0) ? (byte & 0x0F) : (byte >> 4)) + 0.5f;
176176
}
177177

178178
__m256 q = _mm256_loadu_ps(q_arr);
179-
/* dst = mn + q * scale using FMA */
179+
/* dst = mn + (q + 0.5) * scale using FMA (0.5 already added above) */
180180
__m256 result = _mm256_fmadd_ps(q, v_scale, v_mn);
181181
_mm256_storeu_ps(dst + i, result);
182182
}
@@ -185,7 +185,7 @@ void tq_uniform_4b_dequantize_avx2(const void* src, float* dst, int n) {
185185
for (; i < count; i++) {
186186
uint8_t byte = block->qs[i / 2];
187187
int q = (i % 2 == 0) ? (byte & 0x0F) : (byte >> 4);
188-
dst[i] = mn + q * scale;
188+
dst[i] = mn + ((float)q + 0.5f) * scale;
189189
}
190190
}
191191

0 commit comments

Comments
 (0)