Skip to content

Try to optimize sdvquant implementation a bit.#50

Open
comfyanonymous wants to merge 1 commit into
mainfrom
dev/optim_int4
Open

Try to optimize sdvquant implementation a bit.#50
comfyanonymous wants to merge 1 commit into
mainfrom
dev/optim_int4

Conversation

@comfyanonymous

Copy link
Copy Markdown
Member

No description provided.

@coderabbitai

coderabbitai Bot commented Jun 23, 2026

Copy link
Copy Markdown

Review Change Stack

📝 Walkthrough

Walkthrough

This PR adds an AWQ W4A16 weight-dequantization CUDA kernel with Nanobind bindings, env-gated Python caching infrastructure, and a direct-CUDA dispatch fast path. It also lowers the AWQ GEMV/GEMM threshold from 8 to 4, parameterizes the SVDQuant W4A4 GEMM kernel by warp-grid shape, and adds a transposed LoRA-down path and a grouped-fuse inference fast path in the Python tensor layer.

Changes

AWQ W4A16 Dequantization Kernel and Routing

Layer / File(s) Summary
AWQ dequant CUDA kernel and GEMV threshold
comfy_kitchen/backends/cuda/ops/awq_w4a16.cu
Lowers kGemvMThreshold from 8 to 4; adds store_fp32_pair bf16/fp16 vector-store helpers; implements awq_w4a16_dequant_kernel and launch_dequant for nibble-packed weight dequantization; exports launch_awq_dequant_w4a16_kernel with bf16/fp16 dispatch.
Nanobind binding for awq_dequant_w4a16
comfy_kitchen/backends/cuda/dlpack_bindings.cpp
Declares launch_awq_dequant_w4a16_kernel; implements the awq_dequant_w4a16 Nanobind wrapper with full 2D/dtype/shape/group-size validation; registers it in the Python module; updates M-routing threshold comments to M ≤ 4.
Python AWQ caching and gemv_awq_w4a16 routing
comfy_kitchen/backends/cuda/__init__.py
Adds MMA cutoff, env-var helpers, tensor-attached dequant-cache keying, small-M call-count thresholding, _awq_w4a16_dequant_weight, and bias-aware matmul helpers; rewires gemv_awq_w4a16 to route through cached/dequant/fused-small-M paths with conditional external bias application.
AWQ tensor dispatch direct-CUDA fast path
comfy_kitchen/tensor/awq_w4a16.py
Adds _direct_cuda_backend and related helpers; routes _awq_forward to cuda_backend.gemv_awq_w4a16 when a direct CUDA backend is available; removes redundant in-function imports from dispatch handlers.
AWQ W4A16 tests
tests/test_awq_w4a16.py
New test file covering dequant-cache opt-in, CUDA kernel correctness vs nibble-unpack reference, cache invalidation on scale/zero replacement, small-m cache reuse, forced-fused-small-M bypass, QuantizedTensor-wrapped linear cache reuse, and eager-backend bypass.

SVDQuant W4A4 Warp-Grid GEMM, LoRA-Down, and Grouped-Fuse Fast Path

Layer / File(s) Summary
SVDQuant W4A4 warp-grid kernel parameterization
comfy_kitchen/backends/cuda/ops/scaled_mm_svdquant_w4a4.cu
Adds kWarpsMParam/kWarpsNParam template parameters; derives local CTA/warp geometry; updates shared-memory shapes, B-tile/WS/LoRA cp.async addressing, ldmatrix indexing, and per-mi accumulation; introduces LAUNCH_GEMM_SHAPE macro and M-threshold dispatch (4×2 / 8×1 / 2×4).
SVDQuant LoRA-down transposed GEMM and lora-up cache
comfy_kitchen/backends/cuda/__init__.py
Adds _SVDQUANT_LORA_DOWN_TMP_CACHE; reworks _natural_lora_up_from_tile_packed structured cache key; adds _svdquant_use_transposed_lora_down policy; updates quantize_svdquant_w4a4 with r > 0 guard and transposed-GEMM workspace reuse; refines scaled_mm_svdquant_w4a4 auto-fusion threshold policy and rank-0 guard.
SVDQuant W4A4 Python tensor restructuring and grouped-fuse fast path
comfy_kitchen/tensor/svdquant_w4a4.py
Adds _cuda_backend_module, _cuda_quantize_pad_size, grouped-fuse cache helpers, and _svdquant_cached_fused_group; refactors _w4a4_forward_from_quantized_activation to read fields directly; adds CUDA grouped-fuse fast path in svdquant_w4a4_grouped_linear; extracts _w4a4_forward_prepared and _svdquant_split_fused_output; simplifies _w4a4_forward.
SVDQuant W4A4 policy, cache, and rank-0 tests
tests/test_svdquant_w4a4.py
Adds tests for LoRA-down transposed-GEMM policy, grouped-fuse cache default/env-toggle, CUDA grouped-linear auto-fuse and cache reuse, rank-0 LoRA edge cases, _cuda_quantize_pad_size policy, lora-up cache identity reuse under inference-mode, and auto-fuse policy flag sequence across (m, n) scenarios.

Sequence Diagrams

sequenceDiagram
  rect rgba(100, 180, 255, 0.5)
    note over _awq_forward,gemv_awq_w4a16: AWQ W4A16 dispatch (direct-CUDA path)
  end
  participant _awq_forward
  participant _direct_cuda_backend
  participant gemv_awq_w4a16 as cuda_backend.gemv_awq_w4a16
  participant _awq_w4a16_cached_matmul_with_optional_bias
  participant _awq_w4a16_dequant_weight
  participant _C_awq_w4a16 as _C.awq_w4a16 (fused small-M)

  _awq_forward->>_direct_cuda_backend: input_tensor, qdata, wscales, wzeros
  _direct_cuda_backend-->>_awq_forward: cuda_backend or None
  _awq_forward->>gemv_awq_w4a16: x, qweight, wscales, wzeros, bias, group_size
  gemv_awq_w4a16->>_awq_w4a16_cached_matmul_with_optional_bias: try cache hit
  alt cache hit
    _awq_w4a16_cached_matmul_with_optional_bias-->>gemv_awq_w4a16: out, bias_applied=True
  else small-m dequant path
    gemv_awq_w4a16->>_awq_w4a16_dequant_weight: qweight → fp16/bf16 weight (cached or temp)
    _awq_w4a16_dequant_weight-->>gemv_awq_w4a16: weight
    gemv_awq_w4a16-->>gemv_awq_w4a16: matmul → out, bias_applied
  else fused small-M kernel
    gemv_awq_w4a16->>_C_awq_w4a16: x2d, qweight, wscales, wzeros, group_size
    _C_awq_w4a16-->>gemv_awq_w4a16: out, bias_applied=False
  end
  gemv_awq_w4a16-->>_awq_forward: final output (+ bias if not applied)
Loading
sequenceDiagram
  rect rgba(180, 130, 255, 0.5)
    note over svdquant_w4a4_grouped_linear,_w4a4_forward_prepared: SVDQuant W4A4 grouped-fuse fast path
  end
  participant svdquant_w4a4_grouped_linear
  participant _svdquant_cached_fused_group
  participant _w4a4_forward_prepared
  participant quantize_svdquant_w4a4
  participant scaled_mm_svdquant_w4a4
  participant _svdquant_split_fused_output

  svdquant_w4a4_grouped_linear->>_svdquant_cached_fused_group: weights list + input m
  _svdquant_cached_fused_group-->>svdquant_w4a4_grouped_linear: fused_qdata, fused_params, splits
  svdquant_w4a4_grouped_linear->>_w4a4_forward_prepared: x, fused_qdata, fused_params, cuda_backend
  _w4a4_forward_prepared->>quantize_svdquant_w4a4: x, smooth_factor, proj_down, pad_size
  quantize_svdquant_w4a4-->>_w4a4_forward_prepared: q_act, ascales, lora_act
  _w4a4_forward_prepared->>scaled_mm_svdquant_w4a4: q_act, fused_qdata, ascales, lora_act
  scaled_mm_svdquant_w4a4-->>_w4a4_forward_prepared: fused_out
  _w4a4_forward_prepared-->>svdquant_w4a4_grouped_linear: fused_out
  svdquant_w4a4_grouped_linear->>_svdquant_split_fused_output: fused_out, splits
  _svdquant_split_fused_output-->>svdquant_w4a4_grouped_linear: [out_0, out_1, ..., out_n]
Loading
🚥 Pre-merge checks | ✅ 2
✅ Passed checks (2 passed)
Check name Status Explanation
Linked Issues check ✅ Passed Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check ✅ Passed Check skipped because no linked issues were found for this pull request.

✏️ Tip: You can configure your own custom pre-merge checks in the settings.

✨ Finishing Touches
🧪 Generate unit tests (beta)
  • Create PR with unit tests
  • Commit unit tests in branch dev/optim_int4
✨ Simplify code
  • Create PR with simplified code
  • Commit simplified code in branch dev/optim_int4

Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out.

❤️ Share

Comment @coderabbitai help to get the list of available commands.

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actionable comments posted: 4

🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Inline comments:
In `@comfy_kitchen/backends/cuda/__init__.py`:
- Around line 981-1005: The issue is a race condition where the weight is cached
in qweight._ck_awq_dequant_cache immediately after enqueueing the async
_C.awq_dequant_w4a16 kernel, but the cache can be accessed by another CUDA
stream before the kernel finishes executing. To fix this, store a CUDA event
with the cache tuple when it is created (capturing the event after the async
kernel is enqueued), then modify the code that reads from this cache (in
_awq_w4a16_cached_dequant_weight) to synchronize on that event before using the
cached weight. Alternatively, key the cache by both the cache_key and the
producer stream to prevent cross-stream reuse until the first fill is complete.

In `@comfy_kitchen/backends/cuda/dlpack_bindings.cpp`:
- Around line 728-749: The awq_dequant_w4a16 function validates the data types
of wscales, wzeros, and weight tensors but fails to validate the dtype of
qweight, which should be byte-packed (int8 or uint8). Add a validation check
after the existing dtype validation using svdquant_dtype_code to ensure
qweight.dtype() is int8 or uint8, throwing a runtime_error if it is not. This
check must occur before launch_awq_dequant_w4a16_kernel is called to prevent the
CUDA kernel from misinterpreting non-byte-packed data as packed uint4 values.

In `@comfy_kitchen/backends/cuda/ops/awq_w4a16.cu`:
- Around line 615-625: The group index `g` is computed once from `k` but is used
for both the low nibble at position `k` and the high nibble at position `k + 1`.
When `group_size` is odd, the high nibble may belong to a different quantization
group but incorrectly uses the previous group's scale and zero values. Fix this
by either validating that `group_size` is even before launching all AWQ
byte-pair kernels, or by computing separate group indices for both nibble
positions (compute `g` for `k` and a separate group index for `k + 1`) in the
dequantization logic in the store_fp32_pair call.

In `@comfy_kitchen/tensor/svdquant_w4a4.py`:
- Around line 295-306: The grouped CUDA fuse path does not validate whether
shared quantization is possible before calling _svdquant_cached_fused_group(),
which internally assumes shared quantization. Add a validation check using
svdquant_w4a4_can_share_quant() on the weights parameter before entering the if
block that checks cuda_backend and _svdquant_grouped_fuse_enabled(). This
ensures that all weights have compatible quantization parameters (same
smooth_factor and proj_down) before attempting the fused operation, preventing
silent incorrect fusion of incompatible weights.
🪄 Autofix (Beta)

Fix all unresolved CodeRabbit comments on this PR:

  • Push a commit to this branch (recommended)
  • Create a new PR with the fixes

ℹ️ Review info
⚙️ Run configuration

Configuration used: Organization UI

Review profile: ASSERTIVE

Plan: Pro Plus

Run ID: 133e9f5e-ea54-408c-a639-88c0b9cd570c

📥 Commits

Reviewing files that changed from the base of the PR and between 0570b47 and cfcfd0c.

📒 Files selected for processing (8)
  • comfy_kitchen/backends/cuda/__init__.py
  • comfy_kitchen/backends/cuda/dlpack_bindings.cpp
  • comfy_kitchen/backends/cuda/ops/awq_w4a16.cu
  • comfy_kitchen/backends/cuda/ops/scaled_mm_svdquant_w4a4.cu
  • comfy_kitchen/tensor/awq_w4a16.py
  • comfy_kitchen/tensor/svdquant_w4a4.py
  • tests/test_awq_w4a16.py
  • tests/test_svdquant_w4a4.py

Comment on lines +981 to +1005
if hasattr(_C, "awq_dequant_w4a16"):
weight = torch.empty(n, k, dtype=compute_dtype, device=qweight.device)
stream_ptr = torch.cuda.current_stream(qweight.device).cuda_stream
_C.awq_dequant_w4a16(
_wrap_for_dlpack(qweight.view(torch.uint8)),
_wrap_for_dlpack(wscales),
_wrap_for_dlpack(wzeros),
_wrap_for_dlpack(weight),
group_size,
stream_ptr,
)
else:
x16 = qweight.view(torch.uint8).to(torch.int16)
lo = (x16 & 0xF).to(torch.int8)
hi = ((x16 >> 4) & 0xF).to(torch.int8)
nibbles = torch.stack([lo, hi], dim=-1).reshape(n, k).to(compute_dtype)
weight = (
(nibbles.view(n, k // group_size, group_size) - 8.0)
* wscales.t().unsqueeze(-1)
+ wzeros.t().unsqueeze(-1)
).view(n, k)

if cache:
weight = weight.contiguous()
qweight._ck_awq_dequant_cache = (cache_key, weight)

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🩺 Stability & Availability | 🟠 Major | 🏗️ Heavy lift

Do not publish the AWQ dequant cache before the producer stream is safe.

_C.awq_dequant_w4a16 enqueues an async kernel on stream_ptr, but line 1005 exposes weight via qweight._ck_awq_dequant_cache immediately. A second request on another CUDA stream can hit _awq_w4a16_cached_dequant_weight and enqueue mm reading weight before the dequant kernel finishes. Store a CUDA event with the cache and wait on reuse, or key the cache by producer stream until the first fill is complete.

🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@comfy_kitchen/backends/cuda/__init__.py` around lines 981 - 1005, The issue
is a race condition where the weight is cached in qweight._ck_awq_dequant_cache
immediately after enqueueing the async _C.awq_dequant_w4a16 kernel, but the
cache can be accessed by another CUDA stream before the kernel finishes
executing. To fix this, store a CUDA event with the cache tuple when it is
created (capturing the event after the async kernel is enqueued), then modify
the code that reads from this cache (in _awq_w4a16_cached_dequant_weight) to
synchronize on that event before using the cached weight. Alternatively, key the
cache by both the cache_key and the producer stream to prevent cross-stream
reuse until the first fill is complete.

Comment on lines +728 to +749
if (qweight.ndim() != 2 || wscales.ndim() != 2 ||
wzeros.ndim() != 2 || weight.ndim() != 2) {
throw std::runtime_error("awq_dequant_w4a16: all tensors must be 2D");
}
const int N = static_cast<int>(qweight.shape(0));
const int K = static_cast<int>(qweight.shape(1)) * 2;
if (group_size <= 0 || K % group_size != 0) {
throw std::runtime_error("awq_dequant_w4a16: K must be divisible by group_size");
}
const int dtype_code = svdquant_dtype_code(wscales.dtype());
if (dtype_code != 1 && dtype_code != 2) {
throw std::runtime_error("awq_dequant_w4a16: only fp16 (1) and bf16 (2) outputs supported");
}
if (svdquant_dtype_code(wzeros.dtype()) != dtype_code ||
svdquant_dtype_code(weight.dtype()) != dtype_code) {
throw std::runtime_error("awq_dequant_w4a16: wscales, wzeros, and weight dtypes must match");
}
if (wscales.shape(0) != K / group_size || wscales.shape(1) != N ||
wzeros.shape(0) != K / group_size || wzeros.shape(1) != N ||
weight.shape(0) != N || weight.shape(1) != K) {
throw std::runtime_error("awq_dequant_w4a16: expected qweight (N,K/2), scales/zeros (K/G,N), weight (N,K)");
}

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🎯 Functional Correctness | 🟠 Major | ⚡ Quick win

Validate qweight is byte-packed before launching CUDA.

The wrapper validates the fp output tensors but never checks qweight.dtype(). A fp16/fp32 CUDA tensor with the expected 2D shape would pass these checks, then launch_awq_dequant_w4a16_kernel would interpret its raw bytes as packed uint4 and return garbage. Reject anything except int8/uint8 before the shape-only contract lets the nibble gremlin in.

🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@comfy_kitchen/backends/cuda/dlpack_bindings.cpp` around lines 728 - 749, The
awq_dequant_w4a16 function validates the data types of wscales, wzeros, and
weight tensors but fails to validate the dtype of qweight, which should be
byte-packed (int8 or uint8). Add a validation check after the existing dtype
validation using svdquant_dtype_code to ensure qweight.dtype() is int8 or uint8,
throwing a runtime_error if it is not. This check must occur before
launch_awq_dequant_w4a16_kernel is called to prevent the CUDA kernel from
misinterpreting non-byte-packed data as packed uint4 values.

Comment on lines +615 to +625
const int k = kh * 2;
const int g = k / G;
const int idx = n * K_half + kh;

const uint8_t byte = static_cast<uint8_t>(qweight[idx]);
const float scale = to_fp32(wscales[g * N + n]);
const float zero = to_fp32(wzeros[g * N + n]);
store_fp32_pair(
weight + n * K + k,
(float(byte & 0xF) - 8.0f) * scale + zero,
(float((byte >> 4) & 0xF) - 8.0f) * scale + zero);

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🎯 Functional Correctness | 🟠 Major | ⚡ Quick win

Keep byte-paired nibbles inside the correct quantization group.

g is computed from k once and then used for both k and k + 1. Since the public validation only requires K % group_size == 0, an odd group_size lets the high nibble at a group boundary use the previous group's scale/zero, so CUDA dequant can silently diverge from the reference path. Either reject odd group_size before launching all AWQ byte-pair kernels, or compute separate group indices here. Tiny nibble, big quibble.

🐛 Possible kernel-side fix if odd groups should be supported
-    const int g = k / G;
+    const int g_lo = k / G;
+    const int g_hi = (k + 1) / G;
     const int idx = n * K_half + kh;
 
     const uint8_t byte = static_cast<uint8_t>(qweight[idx]);
-    const float scale = to_fp32(wscales[g * N + n]);
-    const float zero = to_fp32(wzeros[g * N + n]);
+    const float scale_lo = to_fp32(wscales[g_lo * N + n]);
+    const float zero_lo = to_fp32(wzeros[g_lo * N + n]);
+    const float scale_hi = to_fp32(wscales[g_hi * N + n]);
+    const float zero_hi = to_fp32(wzeros[g_hi * N + n]);
     store_fp32_pair(
         weight + n * K + k,
-        (float(byte & 0xF) - 8.0f) * scale + zero,
-        (float((byte >> 4) & 0xF) - 8.0f) * scale + zero);
+        (float(byte & 0xF) - 8.0f) * scale_lo + zero_lo,
+        (float((byte >> 4) & 0xF) - 8.0f) * scale_hi + zero_hi);
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@comfy_kitchen/backends/cuda/ops/awq_w4a16.cu` around lines 615 - 625, The
group index `g` is computed once from `k` but is used for both the low nibble at
position `k` and the high nibble at position `k + 1`. When `group_size` is odd,
the high nibble may belong to a different quantization group but incorrectly
uses the previous group's scale and zero values. Fix this by either validating
that `group_size` is even before launching all AWQ byte-pair kernels, or by
computing separate group indices for both nibble positions (compute `g` for `k`
and a separate group index for `k + 1`) in the dequantization logic in the
store_fp32_pair call.

Comment on lines +295 to +306
if cuda_backend is not None and len(weights) > 1 and _svdquant_grouped_fuse_enabled():
fused_weight, fused_bias, splits = _svdquant_cached_fused_group(weights, biases)
out = _w4a4_forward_prepared(
fused_weight,
fused_bias,
orig_shape=orig_shape,
m=m,
x_main=x_main,
lora_x=lora_x,
cuda_backend=cuda_backend,
)
return _svdquant_split_fused_output(out, splits)

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🎯 Functional Correctness | 🟠 Major | ⚡ Quick win

Validate shared quantization before taking the grouped CUDA fuse path.

Line 295 enters the fast path before svdquant_w4a4_can_share_quant(...), then _svdquant_cached_fused_group() fuses with assume_shared_quant=True. That skips the caller’s validate_shared_quant/assume_shared_quant contract, so CUDA can silently fuse weights with different smooth_factor/proj_down using the first tensor’s params. Quite the sneaky gremlin: fast fuse, wrong muse.

🐛 Proposed fix
     cuda_backend = _direct_cuda_backend(x_main, qdata)
-    if cuda_backend is not None and len(weights) > 1 and _svdquant_grouped_fuse_enabled():
-        fused_weight, fused_bias, splits = _svdquant_cached_fused_group(weights, biases)
-        out = _w4a4_forward_prepared(
-            fused_weight,
-            fused_bias,
-            orig_shape=orig_shape,
-            m=m,
-            x_main=x_main,
-            lora_x=lora_x,
-            cuda_backend=cuda_backend,
-        )
-        return _svdquant_split_fused_output(out, splits)
-
     if not svdquant_w4a4_can_share_quant(
         weights, validate=validate_shared_quant, trust=assume_shared_quant,
     ):
         raise ValueError("SVDQuant weights do not share quantization parameters")
+
+    if cuda_backend is not None and len(weights) > 1 and _svdquant_grouped_fuse_enabled():
+        fused_weight, fused_bias, splits = _svdquant_cached_fused_group(weights, biases)
+        out = _w4a4_forward_prepared(
+            fused_weight,
+            fused_bias,
+            orig_shape=orig_shape,
+            m=m,
+            x_main=x_main,
+            lora_x=lora_x,
+            cuda_backend=cuda_backend,
+        )
+        return _svdquant_split_fused_output(out, splits)
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@comfy_kitchen/tensor/svdquant_w4a4.py` around lines 295 - 306, The grouped
CUDA fuse path does not validate whether shared quantization is possible before
calling _svdquant_cached_fused_group(), which internally assumes shared
quantization. Add a validation check using svdquant_w4a4_can_share_quant() on
the weights parameter before entering the if block that checks cuda_backend and
_svdquant_grouped_fuse_enabled(). This ensures that all weights have compatible
quantization parameters (same smooth_factor and proj_down) before attempting the
fused operation, preventing silent incorrect fusion of incompatible weights.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant