Add HIP backend for AMD RDNA2-RDNA4 to the existing wheels#74
Open
0xDELUXA wants to merge 1 commit into
Open
Conversation
📝 WalkthroughWalkthroughThe pull request adds an AMD HIP backend with ROCm builds, native quantization and tensor kernels, architecture-aware dispatch, Python integration, validation tests, documentation, and combined CUDA+HIP wheel workflows for Linux and Windows. ChangesHIP backend and packaging
Native HIP implementation
Validation
Sequence Diagram(s)sequenceDiagram
participant PythonDispatch
participant HIPBackend
participant NanobindBindings
participant HIPKernel
PythonDispatch->>HIPBackend: Check architecture, dtype, shape, and WMMA eligibility
HIPBackend->>NanobindBindings: Pass validated tensors and dtype codes
NanobindBindings->>HIPKernel: Launch HIP operation on the selected stream
HIPKernel-->>NanobindBindings: Produce typed output
NanobindBindings-->>PythonDispatch: Return HIP tensor or raise validation error
Possibly related PRs
🚥 Pre-merge checks | ✅ 2✅ Passed checks (2 passed)
✨ Finishing Touches🧪 Generate unit tests (beta)
✨ Simplify code
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. Comment |
|
Review the following changes in direct dependencies. Learn more about Socket for GitHub.
|
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
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 @.github/workflows/build-wheels.yml:
- Line 37: Pin the workflow actions to immutable full commit SHAs instead of
mutable version tags: update pypa/cibuildwheel in
.github/workflows/build-wheels.yml at lines 37-37, actions/upload-artifact at
lines 161-161, and actions/download-artifact at lines 576-576, 594-594, and
600-600; retain readable version comments where appropriate, and also pin the
unchanged actions/checkout@v4 and actions/setup-python@v5 references in the same
file for consistency.
In `@comfy_kitchen/backends/hip/dlpack_bindings.cpp`:
- Around line 16-29: Update map_dtype_to_code to explicitly recognize DLPack’s
float8_e4m3fn and float8_e5m2 dtype enum values, assigning each the intended
distinct mapping used by this backend. If those fp8 types are not supported at
this boundary, remove their mapping from the surrounding comment instead of
implying support.
In `@comfy_kitchen/backends/hip/fp8_utils.h`:
- Around line 42-45: Standardize the BF16 type used by the
to_float<hip_bfloat16> specialization across the HIP backend. Use the preferred
__hip_bfloat16 spelling from <hip/hip_bf16.h> consistently, updating related
declarations and call sites as needed so overload resolution uses one ROCm BF16
type.
In `@README.md`:
- Around line 84-90: Update the README build command section to add a PowerShell
example for setting COMFY_HIP_ARCHS before running pip install, while preserving
the existing POSIX command and clearly labeling the Windows syntax.
🪄 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: cebcd438-4d82-4143-ad8f-a6295972d2a5
📒 Files selected for processing (25)
.github/workflows/build-wheels.ymlREADME.mdcomfy_kitchen/__init__.pycomfy_kitchen/backends/hip/CMakeLists.txtcomfy_kitchen/backends/hip/__init__.pycomfy_kitchen/backends/hip/dlpack_bindings.cppcomfy_kitchen/backends/hip/epilogue.hcomfy_kitchen/backends/hip/fp8_utils.hcomfy_kitchen/backends/hip/gemm_wmma.hcomfy_kitchen/backends/hip/hadamard.hcomfy_kitchen/backends/hip/mma.hcomfy_kitchen/backends/hip/ops/adaln.hipcomfy_kitchen/backends/hip/ops/apply_rope.hipcomfy_kitchen/backends/hip/ops/convrot_w4a4.hipcomfy_kitchen/backends/hip/ops/gemm_fp8.hipcomfy_kitchen/backends/hip/ops/gemm_int8.hipcomfy_kitchen/backends/hip/ops/gemv_awq.hipcomfy_kitchen/backends/hip/ops/per_tensor_fp8.hipcomfy_kitchen/backends/hip/ops/quantize_int8.hipcomfy_kitchen/backends/hip/ops/stochastic_round_fp8.hipcomfy_kitchen/backends/hip/ops/svdquant_w4a4.hipcomfy_kitchen/scaled_mm_v2.pypyproject.tomlsetup.pytests/test_hip_wmma.py
0xDELUXA
force-pushed
the
amd/hip-backend
branch
from
July 13, 2026 16:47
4ab4c89 to
2dba0e3
Compare
There was a problem hiding this comment.
Actionable comments posted: 3
🤖 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/__init__.py`:
- Around line 29-31: Update the comment above the hip priority configuration to
accurately state that the HIP backend registers for all supported AMD
architectures, rather than only gfx12. Leave the registry.is_available("hip")
condition and priority ordering unchanged.
In `@comfy_kitchen/backends/hip/CMakeLists.txt`:
- Line 1: Update the CMake minimum version in cmake_minimum_required to 3.26 so
it matches the Development.SABIModule requirement used by find_package; keep the
existing SABIModule dependency unchanged.
In `@README.md`:
- Line 35: Update the three README.md references to include RDNA3.5: revise the
HIP backend heading at lines 35-35, the Custom HIP WMMA kernels heading at lines
224-224, and the no-visible-AMD-GPU wording at lines 96-99 to mention
RDNA2/3/3.5/4 consistently.
🪄 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: 05e6bded-3ce8-42c1-8a90-8752063c44f4
📒 Files selected for processing (25)
.github/workflows/build-wheels.ymlREADME.mdcomfy_kitchen/__init__.pycomfy_kitchen/backends/hip/CMakeLists.txtcomfy_kitchen/backends/hip/__init__.pycomfy_kitchen/backends/hip/dlpack_bindings.cppcomfy_kitchen/backends/hip/epilogue.hcomfy_kitchen/backends/hip/fp8_utils.hcomfy_kitchen/backends/hip/gemm_wmma.hcomfy_kitchen/backends/hip/hadamard.hcomfy_kitchen/backends/hip/mma.hcomfy_kitchen/backends/hip/ops/adaln.hipcomfy_kitchen/backends/hip/ops/apply_rope.hipcomfy_kitchen/backends/hip/ops/convrot_w4a4.hipcomfy_kitchen/backends/hip/ops/gemm_fp8.hipcomfy_kitchen/backends/hip/ops/gemm_int8.hipcomfy_kitchen/backends/hip/ops/gemv_awq.hipcomfy_kitchen/backends/hip/ops/per_tensor_fp8.hipcomfy_kitchen/backends/hip/ops/quantize_int8.hipcomfy_kitchen/backends/hip/ops/stochastic_round_fp8.hipcomfy_kitchen/backends/hip/ops/svdquant_w4a4.hipcomfy_kitchen/scaled_mm_v2.pypyproject.tomlsetup.pytests/test_hip_wmma.py
0xDELUXA
force-pushed
the
amd/hip-backend
branch
from
July 13, 2026 17:20
2dba0e3 to
0453863
Compare
There was a problem hiding this comment.
Actionable comments posted: 9
🤖 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 @.github/workflows/build-wheels.yml:
- Around line 54-68: Update the ROCm repository setup in the workflow to
install/import AMD’s signing key and set gpgcheck=1 for both [rocm] and
[amdgraphics] before installing ROCm packages. Also update the nightly pip
install command for rocm[devel] to specify an explicit pinned version while
retaining the existing nightly index.
- Around line 409-422: Pin the ROCm package in the “Install ROCm HIP SDK
(Windows)” workflow step to a known-good version instead of installing the
unversioned rocm[devel] nightly package. Update the pip requirement while
preserving the existing multi-arch index and devel-only extra; add a package
hash if the workflow’s dependency-pinning conventions support it.
In `@comfy_kitchen/backends/hip/__init__.py`:
- Around line 136-147: Update has_wmma() to obtain _visible_gfx_arches() once,
then use that snapshot for both the availability/unsupported-architecture check
and _has_wmma(). Preserve the existing boolean behavior while avoiding the
duplicate traversal currently caused by calling is_available().
- Around line 347-362: Update _convrot_supported so the _C.convrot_max_k()
result is cached per current HIP device ID rather than in the process-wide
_convrot_max_k value. Track the device ID alongside its limit, refresh both when
the active device changes, and preserve the existing group-size, divisibility,
and limit checks.
In `@comfy_kitchen/backends/hip/dlpack_bindings.cpp`:
- Around line 92-127: Add contract validation to quantize_per_tensor_fp8,
dequantize_per_tensor_fp8, and stochastic_round_fp8 before launching kernels:
require nonempty float32 scale tensors, validate the expected input/output
transport dtypes, ensure RNG storage is uint8 where applicable, and verify numel
matches all relevant input/output/RNG element counts. Reject invalid arguments
with runtime errors and only launch after validation.
In `@comfy_kitchen/backends/hip/fp8_utils.h`:
- Around line 128-155: Update fp8_to_float and bf8_to_float to preserve the
FP8/BF8 special-value encodings handled by unpack_fp8: return the corresponding
infinity or NaN for e4m3 and e5m2 inputs instead of converting those exponent
patterns to finite values. Keep the existing normal and subnormal conversion
behavior unchanged for all non-special bytes.
In `@comfy_kitchen/backends/hip/ops/gemv_awq.hip`:
- Around line 74-80: Add host-side validation before the launch setup in the
GEMV AWQ entry point containing the M/N zero-dimension guard: reject group_size
== 0, require group_size and K to satisfy the kernel’s eight-weight
packing/alignment constraints, and ensure the AWQ dimensions are compatible with
those groups. Return through the existing invalid-argument/error path before
constructing block/grid when validation fails.
In `@setup.py`:
- Around line 510-522: Restrict the fallback in the CUDA setup block around
setup_cuda_extension() to the dedicated “CUDA unavailable” exception type only.
Re-raise other RuntimeError failures, such as invalid architectures or broken
configurations, and preserve the existing ROCm availability and BUILD_NO_HIP
checks for genuine CUDA-unavailable cases.
In `@tests/test_hip_wmma.py`:
- Around line 13-15: Replace the module-level pytestmark in
tests/test_hip_wmma.py with targeted skips on tests that execute HIP native
kernels. Leave CPU-testable routing checks such as _unsupported_arch_reason,
_has_wmma, and _build_constraints runnable without HIP, while preserving the
unavailable-HIP skip behavior for native-kernel coverage.
🪄 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: a37d2823-0260-4adf-837f-03f2909ec3c9
📒 Files selected for processing (25)
.github/workflows/build-wheels.ymlREADME.mdcomfy_kitchen/__init__.pycomfy_kitchen/backends/hip/CMakeLists.txtcomfy_kitchen/backends/hip/__init__.pycomfy_kitchen/backends/hip/dlpack_bindings.cppcomfy_kitchen/backends/hip/epilogue.hcomfy_kitchen/backends/hip/fp8_utils.hcomfy_kitchen/backends/hip/gemm_wmma.hcomfy_kitchen/backends/hip/hadamard.hcomfy_kitchen/backends/hip/mma.hcomfy_kitchen/backends/hip/ops/adaln.hipcomfy_kitchen/backends/hip/ops/apply_rope.hipcomfy_kitchen/backends/hip/ops/convrot_w4a4.hipcomfy_kitchen/backends/hip/ops/gemm_fp8.hipcomfy_kitchen/backends/hip/ops/gemm_int8.hipcomfy_kitchen/backends/hip/ops/gemv_awq.hipcomfy_kitchen/backends/hip/ops/per_tensor_fp8.hipcomfy_kitchen/backends/hip/ops/quantize_int8.hipcomfy_kitchen/backends/hip/ops/stochastic_round_fp8.hipcomfy_kitchen/backends/hip/ops/svdquant_w4a4.hipcomfy_kitchen/scaled_mm_v2.pypyproject.tomlsetup.pytests/test_hip_wmma.py
Author
|
Thanks for triggering the CI and for being open to adding a HIP backend to |
0xDELUXA
force-pushed
the
amd/hip-backend
branch
from
July 13, 2026 19:00
0453863 to
388d9f3
Compare
There was a problem hiding this comment.
Actionable comments posted: 3
🤖 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/hip/__init__.py`:
- Around line 681-683: Update the smooth operand preparation in
svdquant_quantize to cast smooth to x.dtype before passing it to _operand.
Preserve the existing reshape, device, and shape handling so smooth and ascales
use the same kernel dtype code.
In `@comfy_kitchen/backends/hip/CMakeLists.txt`:
- Around line 14-20: Remove the leading semicolons from the continuation strings
in the default COMFY_HIP_ARCHS set() block, while preserving all listed gfx
targets and their RDNA grouping comments. Ensure CMake stores exactly the 17
architectures without empty list elements.
In `@comfy_kitchen/backends/hip/ops/gemm_fp8.hip`:
- Around line 41-42: Create a shared wave_reduce_sum helper in the existing
mma.h/gemm_wmma.h utility layer that initializes its reduction offset from kWave
/ 2, then replace the duplicated inline reductions in
comfy_kitchen/backends/hip/ops/gemm_fp8.hip lines 41-42,
comfy_kitchen/backends/hip/ops/gemm_int8.hip lines 36-37, and
comfy_kitchen/backends/hip/ops/gemv_awq.hip lines 57-58 with calls to that
helper, preserving each accumulator’s type.
🪄 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: 4f12fd1c-3c71-46f3-8eb0-fb42c159f30f
📒 Files selected for processing (26)
.github/workflows/build-wheels.ymlREADME.mdcomfy_kitchen/__init__.pycomfy_kitchen/backends/hip/CMakeLists.txtcomfy_kitchen/backends/hip/__init__.pycomfy_kitchen/backends/hip/dlpack_bindings.cppcomfy_kitchen/backends/hip/epilogue.hcomfy_kitchen/backends/hip/fp8_utils.hcomfy_kitchen/backends/hip/gemm_wmma.hcomfy_kitchen/backends/hip/hadamard.hcomfy_kitchen/backends/hip/mma.hcomfy_kitchen/backends/hip/ops/adaln.hipcomfy_kitchen/backends/hip/ops/apply_rope.hipcomfy_kitchen/backends/hip/ops/convrot_w4a4.hipcomfy_kitchen/backends/hip/ops/gemm_fp8.hipcomfy_kitchen/backends/hip/ops/gemm_int8.hipcomfy_kitchen/backends/hip/ops/gemv_awq.hipcomfy_kitchen/backends/hip/ops/per_tensor_fp8.hipcomfy_kitchen/backends/hip/ops/quantize_int8.hipcomfy_kitchen/backends/hip/ops/stochastic_round_fp8.hipcomfy_kitchen/backends/hip/ops/svdquant_w4a4.hipcomfy_kitchen/scaled_mm_v2.pypyproject.tomlsetup.pytests/test_hip_dispatch.pytests/test_hip_wmma.py
0xDELUXA
force-pushed
the
amd/hip-backend
branch
from
July 13, 2026 19:31
388d9f3 to
0609de8
Compare
There was a problem hiding this comment.
Actionable comments posted: 3
🤖 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 @.github/workflows/build-wheels.yml:
- Around line 428-433: Update the ROCm dependency setup near the pinned install
to avoid relying on an aging nightly-only artifact: either mirror the exact
required wheel for continued reproducibility or switch to a stable ROCm
channel/version that remains available. Preserve the Windows wheel job’s
toolchain consistency and ensure the subsequent rocm_sdk path command uses the
resolved installation.
In `@comfy_kitchen/backends/hip/dlpack_bindings.cpp`:
- Around line 177-184: Validate scale_a and scale_b before launching the kernel
in scaled_mm_fp8, and add the same guard to int8_gemm and convrot_w4a4_gemm.
Reject missing or invalid scale tensors with a clean error and do not call the
corresponding launch function when validation fails.
In `@comfy_kitchen/backends/hip/ops/quantize_int8.hip`:
- Around line 97-111: Mirror the immediate post-launch hipGetLastError()
handling used by launch_stochastic_round_fp8_kernel after every listed launcher:
add checks after convrot_quant_kernel<false>, quantize_int8_rowwise_kernel, the
absmax/scale_from_absmax/quantize_int8_tensorwise launches in
comfy_kitchen/backends/hip/ops/quantize_int8.hip (97-111, 86-95, and 114-136),
CK_ADALN_LAUNCH in comfy_kitchen/backends/hip/ops/adaln.hip (98-124),
CK_ROPE_LAUNCH in comfy_kitchen/backends/hip/ops/apply_rope.hip (89-119), and
launch_int8 in comfy_kitchen/backends/hip/ops/gemm_int8.hip (78-111). Ensure
each check occurs immediately after its corresponding launch or dispatch.
🪄 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: 5ec4702a-37e9-4291-9e8a-b4075aad730f
📒 Files selected for processing (26)
.github/workflows/build-wheels.ymlREADME.mdcomfy_kitchen/__init__.pycomfy_kitchen/backends/hip/CMakeLists.txtcomfy_kitchen/backends/hip/__init__.pycomfy_kitchen/backends/hip/dlpack_bindings.cppcomfy_kitchen/backends/hip/epilogue.hcomfy_kitchen/backends/hip/fp8_utils.hcomfy_kitchen/backends/hip/gemm_wmma.hcomfy_kitchen/backends/hip/hadamard.hcomfy_kitchen/backends/hip/mma.hcomfy_kitchen/backends/hip/ops/adaln.hipcomfy_kitchen/backends/hip/ops/apply_rope.hipcomfy_kitchen/backends/hip/ops/convrot_w4a4.hipcomfy_kitchen/backends/hip/ops/gemm_fp8.hipcomfy_kitchen/backends/hip/ops/gemm_int8.hipcomfy_kitchen/backends/hip/ops/gemv_awq.hipcomfy_kitchen/backends/hip/ops/per_tensor_fp8.hipcomfy_kitchen/backends/hip/ops/quantize_int8.hipcomfy_kitchen/backends/hip/ops/stochastic_round_fp8.hipcomfy_kitchen/backends/hip/ops/svdquant_w4a4.hipcomfy_kitchen/scaled_mm_v2.pypyproject.tomlsetup.pytests/test_hip_dispatch.pytests/test_hip_wmma.py
0xDELUXA
force-pushed
the
amd/hip-backend
branch
from
July 13, 2026 20:28
0609de8 to
8a65b1a
Compare
Author
|
Switched from the TheRock nightlies (https://rocm.nightlies.amd.com/) to the "stable" ROCm repo (https://repo.radeon.com/rocm/windows/) in 8a65b1a. The CI was already green with TheRock; this change was made in response to a CodeRabbit suggestion. |
Author
|
@coderabbitai resume |
✅ Action performedReviews resumed. |
There was a problem hiding this comment.
Actionable comments posted: 2
🤖 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/hip/__init__.py`:
- Around line 365-379: Update _convrot_supported to accept the operand device
(or its device index) and use that value instead of torch.cuda.current_device()
when selecting and caching _convrot_max_k. Update every caller to pass x.device
consistently with _stream(x), preserving the existing group-size, divisibility,
and max_k checks.
In `@setup.py`:
- Around line 137-140: Update the HIP-specific CMake argument handling around
ext.backend == "hip" so forcing the Ninja generator is limited to Windows, or
ensure Ninja is installed and available for Linux HIP builds before configuring.
Preserve the existing Windows workaround while preventing Linux builds from
failing due to a missing Ninja executable.
🪄 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: 76ea145b-4646-4ec9-b4ed-5ffa349e86bc
📒 Files selected for processing (26)
.github/workflows/build-wheels.ymlREADME.mdcomfy_kitchen/__init__.pycomfy_kitchen/backends/hip/CMakeLists.txtcomfy_kitchen/backends/hip/__init__.pycomfy_kitchen/backends/hip/dlpack_bindings.cppcomfy_kitchen/backends/hip/epilogue.hcomfy_kitchen/backends/hip/fp8_utils.hcomfy_kitchen/backends/hip/gemm_wmma.hcomfy_kitchen/backends/hip/hadamard.hcomfy_kitchen/backends/hip/mma.hcomfy_kitchen/backends/hip/ops/adaln.hipcomfy_kitchen/backends/hip/ops/apply_rope.hipcomfy_kitchen/backends/hip/ops/convrot_w4a4.hipcomfy_kitchen/backends/hip/ops/gemm_fp8.hipcomfy_kitchen/backends/hip/ops/gemm_int8.hipcomfy_kitchen/backends/hip/ops/gemv_awq.hipcomfy_kitchen/backends/hip/ops/per_tensor_fp8.hipcomfy_kitchen/backends/hip/ops/quantize_int8.hipcomfy_kitchen/backends/hip/ops/stochastic_round_fp8.hipcomfy_kitchen/backends/hip/ops/svdquant_w4a4.hipcomfy_kitchen/scaled_mm_v2.pypyproject.tomlsetup.pytests/test_hip_dispatch.pytests/test_hip_wmma.py
0xDELUXA
force-pushed
the
amd/hip-backend
branch
from
July 13, 2026 20:56
8a65b1a to
8e826da
Compare
There was a problem hiding this comment.
Actionable comments posted: 1
🤖 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/hip/dlpack_bindings.cpp`:
- Around line 392-403: The HIP DLPack entry points need validate tensor lengths
and dtypes before passing raw buffers to kernels. In
comfy_kitchen/backends/hip/dlpack_bindings.cpp lines 392-403, update
svdquant_gemm to validate ascales, wscales, lora_act, lora_up, bias, and c
against the kernel indexing and expected types; in lines 363-372, apply
equivalent checks in the sibling entry point for wscales, wzeros, bias, and out.
Raise a clean exception before launch for any mismatch, while preserving
valid-call behavior.
🪄 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: c9deffe0-d745-4a6d-8bb9-fccaaa66b144
📒 Files selected for processing (26)
.github/workflows/build-wheels.ymlREADME.mdcomfy_kitchen/__init__.pycomfy_kitchen/backends/hip/CMakeLists.txtcomfy_kitchen/backends/hip/__init__.pycomfy_kitchen/backends/hip/dlpack_bindings.cppcomfy_kitchen/backends/hip/epilogue.hcomfy_kitchen/backends/hip/fp8_utils.hcomfy_kitchen/backends/hip/gemm_wmma.hcomfy_kitchen/backends/hip/hadamard.hcomfy_kitchen/backends/hip/mma.hcomfy_kitchen/backends/hip/ops/adaln.hipcomfy_kitchen/backends/hip/ops/apply_rope.hipcomfy_kitchen/backends/hip/ops/convrot_w4a4.hipcomfy_kitchen/backends/hip/ops/gemm_fp8.hipcomfy_kitchen/backends/hip/ops/gemm_int8.hipcomfy_kitchen/backends/hip/ops/gemv_awq.hipcomfy_kitchen/backends/hip/ops/per_tensor_fp8.hipcomfy_kitchen/backends/hip/ops/quantize_int8.hipcomfy_kitchen/backends/hip/ops/stochastic_round_fp8.hipcomfy_kitchen/backends/hip/ops/svdquant_w4a4.hipcomfy_kitchen/scaled_mm_v2.pypyproject.tomlsetup.pytests/test_hip_dispatch.pytests/test_hip_wmma.py
0xDELUXA
force-pushed
the
amd/hip-backend
branch
from
July 13, 2026 21:23
8e826da to
a7f78b5
Compare
✅ Action performedReviews resumed. |
There was a problem hiding this comment.
Actionable comments posted: 2
🤖 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/hip/dlpack_bindings.cpp`:
- Around line 249-267: Add operand extent and dtype validation to int8_gemm,
scaled_mm_fp8, and convrot_w4a4_gemm before their HIP launcher calls, matching
the existing svdquant_gemm validation pattern. Use require_len with each
operation’s M/N/K-derived expected sizes and require_dtype for every raw operand
a, b, and c, while preserving the existing scale and bias checks.
In `@comfy_kitchen/scaled_mm_v2.py`:
- Around line 83-98: Update the `hip.scaled_mm_fp8` call to preserve
`out_dtype=None` instead of converting it to `torch.bfloat16`, so the HIP path
matches `torch._scaled_mm` semantics and defaults to the input matrix dtype.
Keep explicit output dtype handling unchanged.
🪄 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: 153b97dc-0f8b-4028-820e-df55867221f8
📒 Files selected for processing (26)
.github/workflows/build-wheels.ymlREADME.mdcomfy_kitchen/__init__.pycomfy_kitchen/backends/hip/CMakeLists.txtcomfy_kitchen/backends/hip/__init__.pycomfy_kitchen/backends/hip/dlpack_bindings.cppcomfy_kitchen/backends/hip/epilogue.hcomfy_kitchen/backends/hip/fp8_utils.hcomfy_kitchen/backends/hip/gemm_wmma.hcomfy_kitchen/backends/hip/hadamard.hcomfy_kitchen/backends/hip/mma.hcomfy_kitchen/backends/hip/ops/adaln.hipcomfy_kitchen/backends/hip/ops/apply_rope.hipcomfy_kitchen/backends/hip/ops/convrot_w4a4.hipcomfy_kitchen/backends/hip/ops/gemm_fp8.hipcomfy_kitchen/backends/hip/ops/gemm_int8.hipcomfy_kitchen/backends/hip/ops/gemv_awq.hipcomfy_kitchen/backends/hip/ops/per_tensor_fp8.hipcomfy_kitchen/backends/hip/ops/quantize_int8.hipcomfy_kitchen/backends/hip/ops/stochastic_round_fp8.hipcomfy_kitchen/backends/hip/ops/svdquant_w4a4.hipcomfy_kitchen/scaled_mm_v2.pypyproject.tomlsetup.pytests/test_hip_dispatch.pytests/test_hip_wmma.py
0xDELUXA
force-pushed
the
amd/hip-backend
branch
from
July 13, 2026 22:10
a7f78b5 to
2bd3a2b
Compare
Author
|
@comfyanonymous CodeRabbit finally seems happy. If possible, I'd appreciate it if you could test the backend on your AMD systems and also rerun the CI. |
|
Tested the whl file from https://github.com/Comfy-Org/comfy-kitchen/actions/runs/29288989671 on gfx1103, comfy is running but get error output like this: Windows 11 24H2 |
Author
Will look into what might be causing the issue. Could you run
Triton with the native nodes? Some custom nodes have their own Triton backend, for example https://github.com/patientx/ComfyUI-INT8-Fast-ROCM. |
|
Tried this on Linux on a 7900xtx MIOPEN_FIND_MODE=2 TORCH_BLAS_PREFER_HIPBLASLT=1 python main.py --disable-api-nodes --disable-smart-memory --disable-dynamic-vram --disable-pinned-memory --use-flash-attention Anima and Flux2-K-9B worked, performance was within single digit % of using --enable-triton-backend. Krea2 gives garbage noise output similar to picture above. Krea2 does not work when using -enable-triton-backend, has nothing to do with this PR. |
1 task
I am using portable comfy. I try to run pytest tests/test_hip_wmma.py, but all tests skipped. It seems that pytest did not found hip. I am using native triton nodes. ComfyUI-INT8-Fast-ROCM is slower than native on my gfx1103. triton autotune configs in \comfy_kitchen\backends\triton are replaced to get better performance. |
0xDELUXA
force-pushed
the
amd/hip-backend
branch
from
July 17, 2026 08:43
36e72f2 to
5888a1b
Compare
There was a problem hiding this comment.
Actionable comments posted: 1
♻️ Duplicate comments (2)
comfy_kitchen/backends/hip/fp8_utils.h (1)
51-58: 🎯 Functional Correctness | 🟡 Minor | ⚡ Quick winPreserve e5m2 infinities when packing.
fp8_is_nan()excludes infinity, so Line 58 clamps e5m2±infto±57344instead of emitting0x7c/0xfc. Don’t finite-ish infinity again.Proposed fix
if (fp8_is_nan(value)) { return static_cast<uint8_t>(sign_bit | 0x7f); } + if (!e4m3 && + (__float_as_uint(value) & 0x7fffffffu) == 0x7f800000u) { + return static_cast<uint8_t>(sign_bit | 0x7c); + } abs_value = fp8_clamp(abs_value, 0.0f, fp8_max);🤖 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/hip/fp8_utils.h` around lines 51 - 58, Update the FP8 packing logic around fp8_is_nan() to detect infinities before fp8_clamp, preserving e5m2 positive and negative infinity as 0x7c and 0xfc respectively. Keep NaN encoding unchanged and continue clamping finite values to fp8_max.comfy_kitchen/backends/hip/dlpack_bindings.cpp (1)
196-213: 🎯 Functional Correctness | 🟠 Major | ⚡ Quick winRequire uint8 transport buffers for FP8 entry points.
quantize_per_tensor_fp8writes raw FP8 bytes without checkingoutput.dtype(), whiledequantize_per_tensor_fp8reads raw bytes without checkinginput.dtype(). Direct_Ccalls with float storage silently produce corrupt results.Proposed fix
void quantize_per_tensor_fp8(...) { constexpr const char* kFn = "quantize_per_tensor_fp8"; require_code(input_dtype_code, 0, 2, kFn, "input dtype"); require_code(output_dtype_code, 5, 6, kFn, "output dtype"); + require_dtype(output, 3, 3, kFn, "output"); if (map_dtype_to_code(input.dtype()) != input_dtype_code) { throw std::runtime_error(std::string(kFn) + ": input dtype does not match its code"); } void dequantize_per_tensor_fp8(...) { constexpr const char* kFn = "dequantize_per_tensor_fp8"; require_code(input_dtype_code, 5, 6, kFn, "input dtype"); require_code(output_dtype_code, 0, 2, kFn, "output dtype"); + require_dtype(input, 3, 3, kFn, "input"); if (map_dtype_to_code(output.dtype()) != output_dtype_code) {Also applies to: 215-232
🤖 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/hip/dlpack_bindings.cpp` around lines 196 - 213, Update the FP8 entry points quantize_per_tensor_fp8 and dequantize_per_tensor_fp8 to require uint8 transport storage: validate output.dtype() is uint8 before quantization and input.dtype() is uint8 before dequantization, using the existing validation/error conventions. Preserve the current dtype-code, scale, numel, and kernel-launch behavior after these checks.
🤖 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/hip/ops/gemm_fp8.hip`:
- Around line 89-97: Move the M == 0 || N == 0 early return before the K % 16
alignment validation in the GEMM launcher in
comfy_kitchen/backends/hip/ops/gemm_fp8.hip at lines 89-97. Apply the same
ordering change in comfy_kitchen/backends/hip/ops/gemm_int8.hip at lines 86-94,
so empty GEMMs return without kernel-specific K validation.
---
Duplicate comments:
In `@comfy_kitchen/backends/hip/dlpack_bindings.cpp`:
- Around line 196-213: Update the FP8 entry points quantize_per_tensor_fp8 and
dequantize_per_tensor_fp8 to require uint8 transport storage: validate
output.dtype() is uint8 before quantization and input.dtype() is uint8 before
dequantization, using the existing validation/error conventions. Preserve the
current dtype-code, scale, numel, and kernel-launch behavior after these checks.
In `@comfy_kitchen/backends/hip/fp8_utils.h`:
- Around line 51-58: Update the FP8 packing logic around fp8_is_nan() to detect
infinities before fp8_clamp, preserving e5m2 positive and negative infinity as
0x7c and 0xfc respectively. Keep NaN encoding unchanged and continue clamping
finite values to fp8_max.
🪄 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: cd87ace5-12bf-4ebb-b4e1-fbaa6e15783d
📒 Files selected for processing (26)
.github/workflows/build-wheels.ymlREADME.mdcomfy_kitchen/__init__.pycomfy_kitchen/backends/hip/CMakeLists.txtcomfy_kitchen/backends/hip/__init__.pycomfy_kitchen/backends/hip/dlpack_bindings.cppcomfy_kitchen/backends/hip/epilogue.hcomfy_kitchen/backends/hip/fp8_utils.hcomfy_kitchen/backends/hip/gemm_wmma.hcomfy_kitchen/backends/hip/hadamard.hcomfy_kitchen/backends/hip/mma.hcomfy_kitchen/backends/hip/ops/adaln.hipcomfy_kitchen/backends/hip/ops/apply_rope.hipcomfy_kitchen/backends/hip/ops/convrot_w4a4.hipcomfy_kitchen/backends/hip/ops/gemm_fp8.hipcomfy_kitchen/backends/hip/ops/gemm_int8.hipcomfy_kitchen/backends/hip/ops/gemv_awq.hipcomfy_kitchen/backends/hip/ops/per_tensor_fp8.hipcomfy_kitchen/backends/hip/ops/quantize_int8.hipcomfy_kitchen/backends/hip/ops/stochastic_round_fp8.hipcomfy_kitchen/backends/hip/ops/svdquant_w4a4.hipcomfy_kitchen/scaled_mm_v2.pypyproject.tomlsetup.pytests/test_hip_dispatch.pytests/test_hip_wmma.py
0xDELUXA
force-pushed
the
amd/hip-backend
branch
2 times, most recently
from
July 17, 2026 08:48
59ea2aa to
5e110a9
Compare
✅ Action performedReview finished.
|
✅ Action performedReview finished.
|
There was a problem hiding this comment.
Actionable comments posted: 3
🤖 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 @.github/workflows/build-wheels.yml:
- Around line 163-188: Add a CPU-runner smoke test after the ZIP-content
validation that identifies the matching combined Linux wheel, installs it, and
imports both CUDA and HIP backend extensions. Keep the existing archive checks,
and ensure the smoke test executes the installed wheel so it catches ELF
dependency failures or backend withdrawal before publication.
In `@comfy_kitchen/backends/hip/__init__.py`:
- Around line 70-77: Update the exception handler around the HIP extension
loading flow to remove the partially initialized "comfy_kitchen.backends.hip._C"
entry from sys.modules when _spec.loader.exec_module(_C) fails. Keep the
existing _EXT_ERROR assignment, _C reset, and import-safe behavior unchanged.
In `@comfy_kitchen/backends/hip/ops/convrot_w4a4.hip`:
- Around line 63-71: In the convrot_w4a4 operation, move the M/N zero-dimension
guard before the K alignment validation. Return immediately when M == 0 or N ==
0, and retain the existing multiple-of-32 validation for non-empty GEMMs.
🪄 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: 4f0f4726-2260-4cf0-a919-e50b29bbb76b
📒 Files selected for processing (26)
.github/workflows/build-wheels.ymlREADME.mdcomfy_kitchen/__init__.pycomfy_kitchen/backends/hip/CMakeLists.txtcomfy_kitchen/backends/hip/__init__.pycomfy_kitchen/backends/hip/dlpack_bindings.cppcomfy_kitchen/backends/hip/epilogue.hcomfy_kitchen/backends/hip/fp8_utils.hcomfy_kitchen/backends/hip/gemm_wmma.hcomfy_kitchen/backends/hip/hadamard.hcomfy_kitchen/backends/hip/mma.hcomfy_kitchen/backends/hip/ops/adaln.hipcomfy_kitchen/backends/hip/ops/apply_rope.hipcomfy_kitchen/backends/hip/ops/convrot_w4a4.hipcomfy_kitchen/backends/hip/ops/gemm_fp8.hipcomfy_kitchen/backends/hip/ops/gemm_int8.hipcomfy_kitchen/backends/hip/ops/gemv_awq.hipcomfy_kitchen/backends/hip/ops/per_tensor_fp8.hipcomfy_kitchen/backends/hip/ops/quantize_int8.hipcomfy_kitchen/backends/hip/ops/stochastic_round_fp8.hipcomfy_kitchen/backends/hip/ops/svdquant_w4a4.hipcomfy_kitchen/scaled_mm_v2.pypyproject.tomlsetup.pytests/test_hip_dispatch.pytests/test_hip_wmma.py
✅ Action performedReviews resumed. |
0xDELUXA
force-pushed
the
amd/hip-backend
branch
from
July 17, 2026 10:03
5e110a9 to
1e2790a
Compare
There was a problem hiding this comment.
Actionable comments posted: 2
♻️ Duplicate comments (2)
comfy_kitchen/backends/hip/fp8_utils.h (1)
51-58: 🎯 Functional Correctness | 🟡 Minor | ⚡ Quick winRestore e5m2 infinity encoding.
pack_fp8()still clamps e5m2±infto±57344, whileunpack_fp8()supports the infinity encoding. Infinity is being finite-ished again.Proposed fix
if (fp8_is_nan(value)) { return static_cast<uint8_t>(sign_bit | 0x7f); } + if (!e4m3 && (__float_as_uint(value) & 0x7fffffffu) == 0x7f800000u) { + return static_cast<uint8_t>(sign_bit | 0x7c); + } abs_value = fp8_clamp(abs_value, 0.0f, fp8_max);🤖 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/hip/fp8_utils.h` around lines 51 - 58, Update pack_fp8() to detect e5m2 positive and negative infinity before fp8_clamp and return the corresponding signed infinity encoding. Preserve the existing NaN handling and finite-value clamping, while leaving e4m3 behavior unchanged.comfy_kitchen/backends/hip/dlpack_bindings.cpp (1)
196-232: 🗄️ Data Integrity & Integration | 🟠 Major | ⚡ Quick winRequire uint8 storage for FP8 transport tensors.
The format code is separate, but the kernels still read/write raw FP8 bytes. Validate
outputin quantization andinputin dequantization as uint8; otherwise direct_Ccalls silently reinterpret unrelated storage.Proposed fix
require_code(output_dtype_code, 5, 6, kFn, "output dtype"); + require_dtype(output, 3, 3, kFn, "output"); if (map_dtype_to_code(input.dtype()) != input_dtype_code) {require_code(input_dtype_code, 5, 6, kFn, "input dtype"); + require_dtype(input, 3, 3, kFn, "input"); require_code(output_dtype_code, 0, 2, kFn, "output dtype");🤖 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/hip/dlpack_bindings.cpp` around lines 196 - 232, Validate the raw storage dtype for FP8 transport tensors: in quantize_per_tensor_fp8, require output to be uint8, and in dequantize_per_tensor_fp8, require input to be uint8. Keep the existing format-code checks and kernel calls unchanged, and reject mismatched storage before launching either kernel.
🤖 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 `@README.md`:
- Around line 84-103: Update the architecture/build behavior description near
“Architectures default” and the extension-skip statement to distinguish zero
detected GPUs from detected-but-unsupported GPUs: no visible GPUs should fall
back to the full supported target list, while skipping applies only when visible
AMD GPUs are present and none are RDNA2/3/3.5/4. Preserve the existing
hard-error and COMFY_KITCHEN_BUILD_NO_HIP behavior.
In `@setup.py`:
- Around line 167-171: Update the HIP architecture selection in the setup
extension configuration so the raw self.hip_archs CLI value is passed through
normalize_archs() before being selected for CMake, while preserving
ext.hip_archs as the fallback. Ensure comma-separated overrides such as
gfx1100,gfx1200 become separate normalized architectures before appending
COMFY_HIP_ARCHS.
---
Duplicate comments:
In `@comfy_kitchen/backends/hip/dlpack_bindings.cpp`:
- Around line 196-232: Validate the raw storage dtype for FP8 transport tensors:
in quantize_per_tensor_fp8, require output to be uint8, and in
dequantize_per_tensor_fp8, require input to be uint8. Keep the existing
format-code checks and kernel calls unchanged, and reject mismatched storage
before launching either kernel.
In `@comfy_kitchen/backends/hip/fp8_utils.h`:
- Around line 51-58: Update pack_fp8() to detect e5m2 positive and negative
infinity before fp8_clamp and return the corresponding signed infinity encoding.
Preserve the existing NaN handling and finite-value clamping, while leaving e4m3
behavior unchanged.
🪄 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: 5b5bb0c8-4f77-4794-b247-033e93163606
📒 Files selected for processing (26)
.github/workflows/build-wheels.ymlREADME.mdcomfy_kitchen/__init__.pycomfy_kitchen/backends/hip/CMakeLists.txtcomfy_kitchen/backends/hip/__init__.pycomfy_kitchen/backends/hip/dlpack_bindings.cppcomfy_kitchen/backends/hip/epilogue.hcomfy_kitchen/backends/hip/fp8_utils.hcomfy_kitchen/backends/hip/gemm_wmma.hcomfy_kitchen/backends/hip/hadamard.hcomfy_kitchen/backends/hip/mma.hcomfy_kitchen/backends/hip/ops/adaln.hipcomfy_kitchen/backends/hip/ops/apply_rope.hipcomfy_kitchen/backends/hip/ops/convrot_w4a4.hipcomfy_kitchen/backends/hip/ops/gemm_fp8.hipcomfy_kitchen/backends/hip/ops/gemm_int8.hipcomfy_kitchen/backends/hip/ops/gemv_awq.hipcomfy_kitchen/backends/hip/ops/per_tensor_fp8.hipcomfy_kitchen/backends/hip/ops/quantize_int8.hipcomfy_kitchen/backends/hip/ops/stochastic_round_fp8.hipcomfy_kitchen/backends/hip/ops/svdquant_w4a4.hipcomfy_kitchen/scaled_mm_v2.pypyproject.tomlsetup.pytests/test_hip_dispatch.pytests/test_hip_wmma.py
0xDELUXA
force-pushed
the
amd/hip-backend
branch
from
July 17, 2026 10:29
1e2790a to
88b2c35
Compare
There was a problem hiding this comment.
Actionable comments posted: 7
♻️ Duplicate comments (1)
comfy_kitchen/backends/hip/dlpack_bindings.cpp (1)
196-226: 🎯 Functional Correctness | 🟠 Major | ⚡ Quick winEnforce the uint8 FP8 transport buffers.
quantize_per_tensor_fp8never checksoutput.dtype(), anddequantize_per_tensor_fp8never checksinput.dtype(). A direct_Ccall with wider storage is interpreted as a packed byte stream, silently reading or writing the wrong elements.Proposed fix
void quantize_per_tensor_fp8(...) { ... + require_dtype(output, 3, 3, kFn, "output"); require_scale(scale, kFn); void dequantize_per_tensor_fp8(...) { ... + require_dtype(input, 3, 3, kFn, "input"); require_scale(scale, kFn);🤖 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/hip/dlpack_bindings.cpp` around lines 196 - 226, Update quantize_per_tensor_fp8 to validate output.dtype() against output_dtype_code and dequantize_per_tensor_fp8 to validate input.dtype() against input_dtype_code, matching the existing opposite-buffer checks. Reject mismatches before validation or kernel launch using the same error style and dtype mapping already used in these functions.
🤖 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/hip/__init__.py`:
- Around line 107-167: Restrict HIP architecture checks to exact documented
target membership: update _ARCH_SUPPORTED and _ARCH_WMMA and replace
prefix-based checks in _unsupported_arch_reason and _has_wmma with exact
membership tests. In tests/test_hip_dispatch.py lines 13-32, add near-neighbor
unsupported cases including gfx1037, gfx1104, gfx1154, and gfx1202. In
tests/test_hip_wmma.py lines 409-420, assert exact WMMA target membership rather
than prefix matching.
In `@comfy_kitchen/backends/hip/dlpack_bindings.cpp`:
- Around line 341-349: Update require_convrot_group to reject negative k before
the divisibility check, while preserving the existing validation for allowed
group_size values and non-divisible K values. Ensure both ConvRot bindings
cannot proceed to their launchers when K is negative.
In `@comfy_kitchen/backends/hip/fp8_utils.h`:
- Around line 63-81: Update the FP8 packing helper’s rounding in both the
subnormal and normal mantissa calculations to use round-to-nearest-even instead
of floorf(value + 0.5f). Preserve the existing exponent carry handling while
ensuring exact halfway values select the even mantissa code point.
In `@comfy_kitchen/backends/hip/gemm_wmma.h`:
- Around line 51-58: Update the chunk index calculation from tid-major ordering
to lane-interleaved ordering by using i * THREADS + tid in both affected
methods, including the corresponding code around the second occurrence. Preserve
the existing bounds checks and loads, then benchmark the shared path on gfx11
and gfx12.
In `@comfy_kitchen/backends/hip/ops/svdquant_w4a4.hip`:
- Around line 277-285: Move the M == 0 || N == 0 early return in the SVD
quantized GEMM launcher before the K % kSvdGroup validation, so empty workloads
exit without requiring aligned K. Keep the existing invalid-K exception
unchanged for non-empty GEMMs.
In `@setup.py`:
- Around line 338-376: Update the compiler candidate list in get_rocm_path() to
also check the HIP SDK Windows location root / "bin" / "clang++.exe", while
preserving the existing candidates and first-existing-file selection behavior.
In `@tests/test_hip_wmma.py`:
- Around line 358-406: Extend test_no_hipblaslt_on_the_quantized_paths to also
trap eager matmul fallbacks by monkeypatching the relevant torch.matmul entry
point and recording its invocation in called. Keep the existing _scaled_mm and
_int_mm traps, then ensure the final assertion fails when any regular matmul
path is used.
---
Duplicate comments:
In `@comfy_kitchen/backends/hip/dlpack_bindings.cpp`:
- Around line 196-226: Update quantize_per_tensor_fp8 to validate output.dtype()
against output_dtype_code and dequantize_per_tensor_fp8 to validate
input.dtype() against input_dtype_code, matching the existing opposite-buffer
checks. Reject mismatches before validation or kernel launch using the same
error style and dtype mapping already used in these functions.
🪄 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: abb37ea8-8666-419b-a7ee-a149a0433305
📒 Files selected for processing (26)
.github/workflows/build-wheels.ymlREADME.mdcomfy_kitchen/__init__.pycomfy_kitchen/backends/hip/CMakeLists.txtcomfy_kitchen/backends/hip/__init__.pycomfy_kitchen/backends/hip/dlpack_bindings.cppcomfy_kitchen/backends/hip/epilogue.hcomfy_kitchen/backends/hip/fp8_utils.hcomfy_kitchen/backends/hip/gemm_wmma.hcomfy_kitchen/backends/hip/hadamard.hcomfy_kitchen/backends/hip/mma.hcomfy_kitchen/backends/hip/ops/adaln.hipcomfy_kitchen/backends/hip/ops/apply_rope.hipcomfy_kitchen/backends/hip/ops/convrot_w4a4.hipcomfy_kitchen/backends/hip/ops/gemm_fp8.hipcomfy_kitchen/backends/hip/ops/gemm_int8.hipcomfy_kitchen/backends/hip/ops/gemv_awq.hipcomfy_kitchen/backends/hip/ops/per_tensor_fp8.hipcomfy_kitchen/backends/hip/ops/quantize_int8.hipcomfy_kitchen/backends/hip/ops/stochastic_round_fp8.hipcomfy_kitchen/backends/hip/ops/svdquant_w4a4.hipcomfy_kitchen/scaled_mm_v2.pypyproject.tomlsetup.pytests/test_hip_dispatch.pytests/test_hip_wmma.py
0xDELUXA
force-pushed
the
amd/hip-backend
branch
from
July 17, 2026 13:37
88b2c35 to
ce30a43
Compare
There was a problem hiding this comment.
Actionable comments posted: 3
♻️ Duplicate comments (2)
comfy_kitchen/backends/hip/fp8_utils.h (1)
67-74: 🎯 Functional Correctness | 🟡 Minor | ⚡ Quick winPreserve e5m2 infinity when packing.
±infbypassesfp8_is_nan()and is clamped to±57344, althoughunpack_fp8()supports the e5m2 infinity code. Restore the explicit bit check; infinity should not be finite-ished again.Proposed fix
if (fp8_is_nan(value)) { return static_cast<uint8_t>(sign_bit | 0x7f); } + if (!e4m3 && + (__float_as_uint(value) & 0x7fffffffu) == 0x7f800000u) { + return static_cast<uint8_t>(sign_bit | 0x7c); + } abs_value = fp8_clamp(abs_value, 0.0f, fp8_max);🤖 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/hip/fp8_utils.h` around lines 67 - 74, Update the packing logic near fp8_is_nan() to detect positive and negative infinity before fp8_clamp(), and return the correct e5m2 infinity encoding while preserving the sign bit. Keep NaN handling unchanged and ensure finite values continue through the existing clamp path.comfy_kitchen/backends/hip/dlpack_bindings.cpp (1)
196-207: 🎯 Functional Correctness | 🟠 Major | ⚡ Quick winEnforce uint8 FP8 transport storage.
Quantization never verifies that
outputis uint8, and dequantization never verifies thatinputis uint8. Direct_Ccalls can therefore reinterpret multi-byte tensors and silently produce corrupted values.Proposed fix
void quantize_per_tensor_fp8(...) { constexpr const char* kFn = "quantize_per_tensor_fp8"; + require_dtype(output, 3, 3, kFn, "output"); require_code(input_dtype_code, 0, 2, kFn, "input dtype");void dequantize_per_tensor_fp8(...) { constexpr const char* kFn = "dequantize_per_tensor_fp8"; + require_dtype(input, 3, 3, kFn, "input"); require_code(input_dtype_code, 5, 6, kFn, "input dtype");The earlier transport-validation fix still leaves these bytes incognito.
Also applies to: 215-226
🤖 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/hip/dlpack_bindings.cpp` around lines 196 - 207, Update quantize_per_tensor_fp8 to validate that output uses uint8 transport storage, and update the corresponding dequantize_per_tensor_fp8 path to validate that input uses uint8 storage before processing. Use the existing dtype-mapping and error-reporting conventions around these functions, while preserving their current validation flow for logical input/output dtype codes.
🤖 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/hip/ops/svdquant_w4a4.hip`:
- Around line 251-265: Update launch_svdquant_quant_kernel to reject nonzero K
values that are not multiples of 64, placing the validation after the existing M
== 0 early return and before either kernel launch. Preserve the zero-work
behavior and prevent quantization from proceeding with partial SVDQuant groups.
In `@README.md`:
- Around line 37-39: Update the HIP backend descriptions in the README,
including the repeated occurrence, to describe them as custom kernels rather
than universally WMMA-based. Preserve the statement that HIP does not link or
call hipBLAS/hipBLASLt and that matmuls are compiled from sources under
comfy_kitchen/backends/hip/.
In `@setup.py`:
- Around line 268-284: Update hip_arch_supported() to validate architectures
against an exact allowlist matching the WMMA-enabled targets in mma.h and
DEFAULT_HIP_ARCHS, rather than accepting gfx11/gfx12 prefixes. Reuse a shared
target collection for the allowlist and default architecture string so newly
recognized GPU targets are rejected until they have a compiled MMA policy.
---
Duplicate comments:
In `@comfy_kitchen/backends/hip/dlpack_bindings.cpp`:
- Around line 196-207: Update quantize_per_tensor_fp8 to validate that output
uses uint8 transport storage, and update the corresponding
dequantize_per_tensor_fp8 path to validate that input uses uint8 storage before
processing. Use the existing dtype-mapping and error-reporting conventions
around these functions, while preserving their current validation flow for
logical input/output dtype codes.
In `@comfy_kitchen/backends/hip/fp8_utils.h`:
- Around line 67-74: Update the packing logic near fp8_is_nan() to detect
positive and negative infinity before fp8_clamp(), and return the correct e5m2
infinity encoding while preserving the sign bit. Keep NaN handling unchanged and
ensure finite values continue through the existing clamp path.
🪄 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: 44b8f404-0bff-4fd5-801b-b14bc06172fa
📒 Files selected for processing (26)
.github/workflows/build-wheels.ymlREADME.mdcomfy_kitchen/__init__.pycomfy_kitchen/backends/hip/CMakeLists.txtcomfy_kitchen/backends/hip/__init__.pycomfy_kitchen/backends/hip/dlpack_bindings.cppcomfy_kitchen/backends/hip/epilogue.hcomfy_kitchen/backends/hip/fp8_utils.hcomfy_kitchen/backends/hip/gemm_wmma.hcomfy_kitchen/backends/hip/hadamard.hcomfy_kitchen/backends/hip/mma.hcomfy_kitchen/backends/hip/ops/adaln.hipcomfy_kitchen/backends/hip/ops/apply_rope.hipcomfy_kitchen/backends/hip/ops/convrot_w4a4.hipcomfy_kitchen/backends/hip/ops/gemm_fp8.hipcomfy_kitchen/backends/hip/ops/gemm_int8.hipcomfy_kitchen/backends/hip/ops/gemv_awq.hipcomfy_kitchen/backends/hip/ops/per_tensor_fp8.hipcomfy_kitchen/backends/hip/ops/quantize_int8.hipcomfy_kitchen/backends/hip/ops/stochastic_round_fp8.hipcomfy_kitchen/backends/hip/ops/svdquant_w4a4.hipcomfy_kitchen/scaled_mm_v2.pypyproject.tomlsetup.pytests/test_hip_dispatch.pytests/test_hip_wmma.py
0xDELUXA
force-pushed
the
amd/hip-backend
branch
from
July 17, 2026 13:52
ce30a43 to
1e9b6fa
Compare
✅ Action performedReviews resumed. |
There was a problem hiding this comment.
Actionable comments posted: 5
🤖 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 @.github/workflows/build-wheels.yml:
- Around line 163-188: Update the Linux wheel verification step around the wheel
collection and per-wheel loop to assert that every Python 3.12 wheel filename
retains the abi3 tag before checking CUDA and HIP members. Add the equivalent
abi3 assertion to the Windows wheel validation block covering the same wheel
checks, while preserving the existing extension and linkage validations.
In `@comfy_kitchen/backends/hip/__init__.py`:
- Around line 579-590: Update the validation flow in convrot_w4a4_linear so
zero-work inputs return an empty ConvRot result immediately after the existing
inner-shape check, before the K % 32 and convrot_groupsize alignment checks.
Preserve the current validation and fallback behavior for non-empty inputs.
In `@comfy_kitchen/scaled_mm_v2.py`:
- Around line 53-58: Update the HIP fast paths around the shown availability
checks and the corresponding sections at the other referenced locations to honor
explicit backend priority before executing HIP. Use registry backend selection,
or a registry query confirming HIP is selected, instead of relying only on
registry.is_available("hip"); preserve the existing WMMA capability check after
selection is confirmed.
In `@README.md`:
- Around line 46-48: Update the RDNA3.5 and RDNA3 capability descriptions in the
README table to replace “Everything” with “All HIP-supported kernels,” while
preserving the existing fp8 GEMM clarification and other table content.
In `@setup.py`:
- Around line 60-65: Update the source-distribution packaging configuration so
the complete comfy_kitchen/backends/hip tree is included, including
CMakeLists.txt, headers, and .hip sources. Add the necessary
manifest/package-data rules near the existing packaging configuration; do not
rely on implicit inclusion or change the build logic in the visible __init__
method.
🪄 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: 39c20fe6-c1f8-4da7-9a55-ee189fc093fc
📒 Files selected for processing (26)
.github/workflows/build-wheels.ymlREADME.mdcomfy_kitchen/__init__.pycomfy_kitchen/backends/hip/CMakeLists.txtcomfy_kitchen/backends/hip/__init__.pycomfy_kitchen/backends/hip/dlpack_bindings.cppcomfy_kitchen/backends/hip/epilogue.hcomfy_kitchen/backends/hip/fp8_utils.hcomfy_kitchen/backends/hip/gemm_wmma.hcomfy_kitchen/backends/hip/hadamard.hcomfy_kitchen/backends/hip/mma.hcomfy_kitchen/backends/hip/ops/adaln.hipcomfy_kitchen/backends/hip/ops/apply_rope.hipcomfy_kitchen/backends/hip/ops/convrot_w4a4.hipcomfy_kitchen/backends/hip/ops/gemm_fp8.hipcomfy_kitchen/backends/hip/ops/gemm_int8.hipcomfy_kitchen/backends/hip/ops/gemv_awq.hipcomfy_kitchen/backends/hip/ops/per_tensor_fp8.hipcomfy_kitchen/backends/hip/ops/quantize_int8.hipcomfy_kitchen/backends/hip/ops/stochastic_round_fp8.hipcomfy_kitchen/backends/hip/ops/svdquant_w4a4.hipcomfy_kitchen/scaled_mm_v2.pypyproject.tomlsetup.pytests/test_hip_dispatch.pytests/test_hip_wmma.py
0xDELUXA
force-pushed
the
amd/hip-backend
branch
from
July 17, 2026 14:18
1e9b6fa to
d6747a0
Compare
✅ Action performedReviews resumed. |
There was a problem hiding this comment.
Actionable comments posted: 1
🤖 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/hip/__init__.py`:
- Around line 579-582: Replace the broad x.numel() zero-input guard with a check
that only matches zero-row inputs (M == 0), so zero-K inputs continue through
the GEMM path and produce valid zero accumulations plus bias. Preserve the
existing empty-result shape and dtype/device behavior for genuinely zero-row
inputs.
🪄 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: 2be23dfd-3a7d-4f60-987c-bb6a5c33a64d
📒 Files selected for processing (26)
.github/workflows/build-wheels.ymlREADME.mdcomfy_kitchen/__init__.pycomfy_kitchen/backends/hip/CMakeLists.txtcomfy_kitchen/backends/hip/__init__.pycomfy_kitchen/backends/hip/dlpack_bindings.cppcomfy_kitchen/backends/hip/epilogue.hcomfy_kitchen/backends/hip/fp8_utils.hcomfy_kitchen/backends/hip/gemm_wmma.hcomfy_kitchen/backends/hip/hadamard.hcomfy_kitchen/backends/hip/mma.hcomfy_kitchen/backends/hip/ops/adaln.hipcomfy_kitchen/backends/hip/ops/apply_rope.hipcomfy_kitchen/backends/hip/ops/convrot_w4a4.hipcomfy_kitchen/backends/hip/ops/gemm_fp8.hipcomfy_kitchen/backends/hip/ops/gemm_int8.hipcomfy_kitchen/backends/hip/ops/gemv_awq.hipcomfy_kitchen/backends/hip/ops/per_tensor_fp8.hipcomfy_kitchen/backends/hip/ops/quantize_int8.hipcomfy_kitchen/backends/hip/ops/stochastic_round_fp8.hipcomfy_kitchen/backends/hip/ops/svdquant_w4a4.hipcomfy_kitchen/scaled_mm_v2.pypyproject.tomlsetup.pytests/test_hip_dispatch.pytests/test_hip_wmma.py
0xDELUXA
force-pushed
the
amd/hip-backend
branch
from
July 17, 2026 19:39
d6747a0 to
6d8a2f3
Compare
Author
|
Thanks @comfyanonymous for all the CI reruns! The latest run is green. I'm still making my way through CodeRabbit's reviews, and I'm starting to wonder whether I'll ever reach the end, or hear from a maintainer about the direction for |
There was a problem hiding this comment.
Actionable comments posted: 1
♻️ Duplicate comments (1)
comfy_kitchen/backends/hip/dlpack_bindings.cpp (1)
196-232: 🎯 Functional Correctness | 🟠 Major | ⚡ Quick winRequire uint8 storage for both FP8 transport buffers.
The format codes do not validate actual ndarray storage. A float output is partially byte-written during quantization, while a float input is byte-reinterpreted during dequantization—silent corruption on direct
_Ccalls. Don’t let FP8 wear a four-byte disguise, imp.Proposed fix
void quantize_per_tensor_fp8(...) { ... require_code(output_dtype_code, 5, 6, kFn, "output dtype"); + require_dtype(output, 3, 3, kFn, "output"); if (map_dtype_to_code(input.dtype()) != input_dtype_code) { ... } } void dequantize_per_tensor_fp8(...) { ... require_code(output_dtype_code, 0, 2, kFn, "output dtype"); + require_dtype(input, 3, 3, kFn, "input"); if (map_dtype_to_code(output.dtype()) != output_dtype_code) { ... } }🤖 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/hip/dlpack_bindings.cpp` around lines 196 - 232, Add explicit uint8 storage validation for both FP8 transport buffers in quantize_per_tensor_fp8 and dequantize_per_tensor_fp8: require input and output ndarrays to have one-byte uint8 dtype before launching their kernels, while retaining the existing format-code checks for logical FP8 types and dequantized outputs.
🤖 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/hip/hadamard.h`:
- Around line 38-65: Update convrot_max_k() and its caller check_convrot_k() to
query the LDS limit for the device associated with the kernel launch stream
rather than relying on hipGetDevice() from the caller thread. Thread the
stream/device context through the validation path, using the stream’s device API
where available, and preserve the existing zero/failure handling when the device
or attribute query fails.
---
Duplicate comments:
In `@comfy_kitchen/backends/hip/dlpack_bindings.cpp`:
- Around line 196-232: Add explicit uint8 storage validation for both FP8
transport buffers in quantize_per_tensor_fp8 and dequantize_per_tensor_fp8:
require input and output ndarrays to have one-byte uint8 dtype before launching
their kernels, while retaining the existing format-code checks for logical FP8
types and dequantized outputs.
🪄 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: 0929eb11-44e3-4851-a6dc-c97b088d3465
📒 Files selected for processing (26)
.github/workflows/build-wheels.ymlREADME.mdcomfy_kitchen/__init__.pycomfy_kitchen/backends/hip/CMakeLists.txtcomfy_kitchen/backends/hip/__init__.pycomfy_kitchen/backends/hip/dlpack_bindings.cppcomfy_kitchen/backends/hip/epilogue.hcomfy_kitchen/backends/hip/fp8_utils.hcomfy_kitchen/backends/hip/gemm_wmma.hcomfy_kitchen/backends/hip/hadamard.hcomfy_kitchen/backends/hip/mma.hcomfy_kitchen/backends/hip/ops/adaln.hipcomfy_kitchen/backends/hip/ops/apply_rope.hipcomfy_kitchen/backends/hip/ops/convrot_w4a4.hipcomfy_kitchen/backends/hip/ops/gemm_fp8.hipcomfy_kitchen/backends/hip/ops/gemm_int8.hipcomfy_kitchen/backends/hip/ops/gemv_awq.hipcomfy_kitchen/backends/hip/ops/per_tensor_fp8.hipcomfy_kitchen/backends/hip/ops/quantize_int8.hipcomfy_kitchen/backends/hip/ops/stochastic_round_fp8.hipcomfy_kitchen/backends/hip/ops/svdquant_w4a4.hipcomfy_kitchen/scaled_mm_v2.pypyproject.tomlsetup.pytests/test_hip_dispatch.pytests/test_hip_wmma.py
0xDELUXA
force-pushed
the
amd/hip-backend
branch
from
July 17, 2026 19:58
6d8a2f3 to
3e66d59
Compare
Member
|
The goal is to eventually merge this PR, I asked some people from AMD to take a look at it. |
Author
Great news, thanks! |
There was a problem hiding this comment.
Actionable comments posted: 3
♻️ Duplicate comments (2)
comfy_kitchen/backends/hip/fp8_utils.h (1)
67-74: 🎯 Functional Correctness | 🟡 Minor | ⚡ Quick winRestore e5m2 infinity encoding.
±infis not NaN, so it reachesfp8_clamp()and becomes finite±57344instead of e5m20x7c/0xfc. The decoder already supports infinity; don’t let it get finite-ished.Proposed fix
if (fp8_is_nan(value)) { return static_cast<uint8_t>(sign_bit | 0x7f); } + if (!e4m3 && (__float_as_uint(value) & 0x7fffffffu) == 0x7f800000u) { + return static_cast<uint8_t>(sign_bit | 0x7c); + } abs_value = fp8_clamp(abs_value, 0.0f, fp8_max);🤖 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/hip/fp8_utils.h` around lines 67 - 74, Update the FP8 conversion logic around fp8_is_nan() to detect positive and negative infinity before fp8_clamp(). Preserve the sign bit and return the e5m2 infinity encoding (0x7c for positive, 0xfc for negative), while leaving NaN and finite-value handling unchanged.comfy_kitchen/backends/hip/dlpack_bindings.cpp (1)
196-213: 🎯 Functional Correctness | 🟠 Major | ⚡ Quick winRestore the FP8 transport-storage checks.
The format codes are validated, but quantization still accepts a non-uint8
output, and dequantization accepts a non-uint8input. The kernels process one raw byte per element, so wider storage produces corrupted direct_Cresults. Require uint8 at both boundaries—a byte-sized fix with a big bite.Proposed fix
void quantize_per_tensor_fp8(...) { constexpr const char* kFn = "quantize_per_tensor_fp8"; require_code(input_dtype_code, 0, 2, kFn, "input dtype"); require_code(output_dtype_code, 5, 6, kFn, "output dtype"); + require_dtype(output, 3, 3, kFn, "output"); if (map_dtype_to_code(input.dtype()) != input_dtype_code) {void dequantize_per_tensor_fp8(...) { constexpr const char* kFn = "dequantize_per_tensor_fp8"; require_code(input_dtype_code, 5, 6, kFn, "input dtype"); require_code(output_dtype_code, 0, 2, kFn, "output dtype"); + require_dtype(input, 3, 3, kFn, "input"); if (map_dtype_to_code(output.dtype()) != output_dtype_code) {Also applies to: 215-232
🤖 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/hip/dlpack_bindings.cpp` around lines 196 - 213, Update quantize_per_tensor_fp8 to require output storage to be uint8 after validating output_dtype_code, and apply the corresponding check in the dequantization function covering the referenced lines so its input storage is uint8. Reuse the existing dtype-validation helpers and preserve the current format-code checks and kernel launch flow.
🤖 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/hip/__init__.py`:
- Around line 485-491: In both ConvRot validation paths, update the checks
around convrot_groupsize before the modulo operation so zero is rejected first
with the existing ValueError validation. Preserve the allowed-size validation
for 16, 64, and 256, and apply the same ordering in the second path identified
by the duplicate logic.
In `@README.md`:
- Line 231: Update the README backend description for “hip” to say “Custom HIP
kernels” without claiming they are universally WMMA, while preserving the listed
AMD RDNA2 through RDNA4 architectures.
In `@setup.py`:
- Around line 324-360: Ensure the PEP 517 isolated build environment can resolve
rocm-sdk by adding it to build-system.requires, or remove the claim that plain
installs automatically detect caller-environment ROCm and document
--no-build-isolation/ROCM_HOME instead. Keep rocm_sdk_root() and get_rocm_path()
behavior aligned with the chosen installation contract.
---
Duplicate comments:
In `@comfy_kitchen/backends/hip/dlpack_bindings.cpp`:
- Around line 196-213: Update quantize_per_tensor_fp8 to require output storage
to be uint8 after validating output_dtype_code, and apply the corresponding
check in the dequantization function covering the referenced lines so its input
storage is uint8. Reuse the existing dtype-validation helpers and preserve the
current format-code checks and kernel launch flow.
In `@comfy_kitchen/backends/hip/fp8_utils.h`:
- Around line 67-74: Update the FP8 conversion logic around fp8_is_nan() to
detect positive and negative infinity before fp8_clamp(). Preserve the sign bit
and return the e5m2 infinity encoding (0x7c for positive, 0xfc for negative),
while leaving NaN and finite-value handling unchanged.
🪄 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: cef584f2-9a57-427b-87a3-bfced9d5b1e6
📒 Files selected for processing (26)
.github/workflows/build-wheels.ymlREADME.mdcomfy_kitchen/__init__.pycomfy_kitchen/backends/hip/CMakeLists.txtcomfy_kitchen/backends/hip/__init__.pycomfy_kitchen/backends/hip/dlpack_bindings.cppcomfy_kitchen/backends/hip/epilogue.hcomfy_kitchen/backends/hip/fp8_utils.hcomfy_kitchen/backends/hip/gemm_wmma.hcomfy_kitchen/backends/hip/hadamard.hcomfy_kitchen/backends/hip/mma.hcomfy_kitchen/backends/hip/ops/adaln.hipcomfy_kitchen/backends/hip/ops/apply_rope.hipcomfy_kitchen/backends/hip/ops/convrot_w4a4.hipcomfy_kitchen/backends/hip/ops/gemm_fp8.hipcomfy_kitchen/backends/hip/ops/gemm_int8.hipcomfy_kitchen/backends/hip/ops/gemv_awq.hipcomfy_kitchen/backends/hip/ops/per_tensor_fp8.hipcomfy_kitchen/backends/hip/ops/quantize_int8.hipcomfy_kitchen/backends/hip/ops/stochastic_round_fp8.hipcomfy_kitchen/backends/hip/ops/svdquant_w4a4.hipcomfy_kitchen/scaled_mm_v2.pypyproject.tomlsetup.pytests/test_hip_dispatch.pytests/test_hip_wmma.py
Comment on lines
+485
to
+491
| if convrot: | ||
| if k % convrot_groupsize != 0: | ||
| raise ValueError( | ||
| f"ConvRot group size {convrot_groupsize} does not divide input features {k}" | ||
| ) | ||
| if convrot_groupsize not in (16, 64, 256): | ||
| raise ValueError(f"ConvRot group size must be 16, 64 or 256, got {convrot_groupsize}") |
There was a problem hiding this comment.
🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win
Validate convrot_groupsize before using it as a divisor.
Both paths evaluate % convrot_groupsize before excluding zero, so convrot_groupsize=0 escapes the intended validation/fallback as ZeroDivisionError. Defuse this modulo landmine first.
Proposed fix
if convrot:
+ if convrot_groupsize not in (16, 64, 256):
+ raise ValueError(
+ f"ConvRot group size must be 16, 64 or 256, got {convrot_groupsize}"
+ )
if k % convrot_groupsize != 0:
raise ValueError(
f"ConvRot group size {convrot_groupsize} does not divide input features {k}"
)
- if convrot_groupsize not in (16, 64, 256):
- raise ValueError(f"ConvRot group size must be 16, 64 or 256, got {convrot_groupsize}")+ if convrot_groupsize == 0:
+ raise ValueError("convrot_groupsize must be non-zero")
if x.shape[-1] % convrot_groupsize != 0:Also applies to: 596-604
🤖 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/hip/__init__.py` around lines 485 - 491, In both
ConvRot validation paths, update the checks around convrot_groupsize before the
modulo operation so zero is rejected first with the existing ValueError
validation. Preserve the allowed-size validation for 16, 64, and 256, and apply
the same ordering in the second path identified by the duplicate logic.
| The library supports multiple backends: | ||
| - **eager**: Pure PyTorch implementation | ||
| - **cuda**: Custom CUDA C kernels (CUDA only) | ||
| - **hip**: Custom HIP WMMA kernels (AMD RDNA2 / RDNA3 / RDNA3.5 / RDNA4) |
There was a problem hiding this comment.
📐 Maintainability & Code Quality | 🟡 Minor | ⚡ Quick win
Describe the backend as custom HIP kernels, not universally WMMA.
RDNA2 only exposes non-WMMA kernels, and Lines 37-40 already use the accurate broader wording.
Proposed fix
-- **hip**: Custom HIP WMMA kernels (AMD RDNA2 / RDNA3 / RDNA3.5 / RDNA4)
+- **hip**: Custom HIP kernels (AMD RDNA2 / RDNA3 / RDNA3.5 / RDNA4)🤖 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 `@README.md` at line 231, Update the README backend description for “hip” to
say “Custom HIP kernels” without claiming they are universally WMMA, while
preserving the listed AMD RDNA2 through RDNA4 architectures.
Comment on lines
+324
to
+360
| def rocm_sdk_root() -> str | None: | ||
| """Ask the pip rocm-sdk wheel for its root, if it is installed.""" | ||
| try: | ||
| root = subprocess.run( | ||
| [sys.executable, "-m", "rocm_sdk", "path", "--root"], | ||
| capture_output=True, | ||
| check=True, | ||
| text=True, | ||
| ).stdout.strip() | ||
| except (subprocess.CalledProcessError, OSError): | ||
| return None | ||
| return root if root and pathlib.Path(root).exists() else None | ||
|
|
||
|
|
||
| def get_rocm_path() -> tuple[str | None, pathlib.Path | None]: | ||
| """Locate a ROCm root and its clang driver. | ||
|
|
||
| Handles the pip ``rocm-sdk`` layout (site-packages/_rocm_sdk_devel) as well | ||
| as a system ROCm install. | ||
| """ | ||
| rocm_home = os.getenv("ROCM_HOME") or os.getenv("ROCM_PATH") or os.getenv("HIP_PATH") | ||
|
|
||
| if rocm_home is None: | ||
| rocm_home = rocm_sdk_root() | ||
|
|
||
| if rocm_home is None: | ||
| sdk = pathlib.Path(sys.prefix) / "Lib" / "site-packages" / "_rocm_sdk_devel" | ||
| if not sdk.exists(): | ||
| sdk = ( | ||
| pathlib.Path(sys.prefix) | ||
| / "lib" | ||
| / f"python{sys.version_info.major}.{sys.version_info.minor}" | ||
| / "site-packages" | ||
| / "_rocm_sdk_devel" | ||
| ) | ||
| if sdk.exists(): | ||
| rocm_home = str(sdk) |
There was a problem hiding this comment.
🩺 Stability & Availability | 🟠 Major | ⚡ Quick win
🧩 Analysis chain
🏁 Script executed:
#!/bin/bash
set -euo pipefail
echo "Build-system requirements:"
sed -n '1,80p' pyproject.toml
echo "ROCm discovery and documented install paths:"
rg -n -C3 'rocm-sdk|rocm_sdk|_rocm_sdk_devel|no-build-isolation|pip install \.' \
setup.py pyproject.toml README.mdRepository: Comfy-Org/comfy-kitchen
Length of output: 6490
🏁 Script executed:
#!/bin/bash
set -euo pipefail
echo "pyproject build requirements:"
sed -n '1,40p' pyproject.toml
echo
echo "ROCm-related code paths:"
sed -n '324,410p' setup.py
echo
echo "Docs mentioning isolation / ROCm:"
rg -n -C2 'rocm-sdk|_rocm_sdk_devel|no-build-isolation|PEP 517|pip install \.' README.md setup.py pyproject.tomlRepository: Comfy-Org/comfy-kitchen
Length of output: 7362
Include rocm-sdk in the build env or drop the plain-install claim
Under PEP 517 isolation, rocm_sdk_root() and the _rocm_sdk_devel fallback only see the temporary build prefix, so a rocm-sdk wheel installed in the caller env can be missed and HIP gets skipped. Add rocm-sdk to build-system.requires, or document --no-build-isolation/ROCM_HOME for ROCm installs.
🧰 Tools
🪛 ast-grep (0.44.1)
[error] 326-331: Command coming from incoming request
Context: subprocess.run(
[sys.executable, "-m", "rocm_sdk", "path", "--root"],
capture_output=True,
check=True,
text=True,
)
Note: [CWE-78] Improper Neutralization of Special Elements used in an OS Command ('OS Command Injection').
(subprocess-from-request)
🤖 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 `@setup.py` around lines 324 - 360, Ensure the PEP 517 isolated build
environment can resolve rocm-sdk by adding it to build-system.requires, or
remove the claim that plain installs automatically detect caller-environment
ROCm and document --no-build-isolation/ROCM_HOME instead. Keep rocm_sdk_root()
and get_rocm_path() behavior aligned with the chosen installation contract.
Author
|
I think I'll stop the back-and-forth with CodeRabbit and rebasing to main, given that we're in a situation of waiting for AMD. 👀 |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
6 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Supersedes #72. That PR is sitting at ~100 review comments and is gfx12-only, so rather than force-push another rewrite on top of it I opened a fresh one. This picks up the two things that actually blocked it:
This builds one wheel with both backends in it, and keeps the
abi3tag while doing it.This adds RDNA3, RDNA3.5 and RDNA2.
The kernels are #72's, rebased on current main. The packaging follows the approach in dev/hip1 by @rattus128.
Packaging
setup.pydetects both toolchains and emitsbackends/cuda/_C*andbackends/hip/_C*into a single wheel. The Linux job installs CUDA and ROCm into the same manylinux container; the Windows job installs the ROCm toolchain from AMD's release directory (https://repo.radeon.com/), the same host the Linux job takes its RPMs from.auditwheelexcludes the ROCm runtime libs the same way it already excludes the CUDA ones, and both jobs fail if the wheel does not contain both extensions. Each backend withdraws itself at import when its runtime is missing, so a CUDA-only or ROCm-only machine gets the same wheel and just sees one backend register.The HIP extension is built against the Python limited API on 3.12+, like the CUDA one. That matters for the packaging goal: an extension that is not abi3 drags the whole wheel back to version-specific builds, so adding AMD would otherwise have multiplied the wheel matrix. It does not. Nothing outside
backends/hip/changes in the build, and the wheel matrix does not grow: one wheel, both backends, sameabi3tag as today.Architectures
gfx1200,gfx1201gfx1150-gfx1153gfx1100-gfx1103gfx1030-gfx1036v_wmma_*_w32_gfx12has no gfx11 encoding, and the difference is not only the intrinsic name: gfx12 splits a K-step across the half-waves and takes 8-byte operands, while gfx11 gives each lane the whole K-step and duplicates it across the half-waves (rocWMMA calls this input duplication). The fragment type and the bytes of a K-row one MMA consumes therefore moved into anMmapolicy in the newmma.h, and the tile GEMM stays byte-addressed and shared between the two.gfx11 has no fp8 WMMA at all, so
MmaFp8/MmaBf8widen the operands to bf16 there. The widening is exact (e4m3 and e5m2 mantissas both fit bf16's 7 bits) and both paths accumulate in f32, so RDNA3 fp8 should be as accurate as RDNA4's native path, just slower.RDNA2 has no matrix cores. It runs the kernels that never needed them (RoPE, AdaLN, the fp8/int8 quantizers, stochastic rounding, the AWQ GEMV) and does not advertise the four GEMMs, which fall through to triton/eager. The GEMM kernels still compile for gfx103x, where they trap, so the registry gates them by architecture instead: with a mix of GPUs in one process the capability set is the intersection, since kernels launch on the tensor's own device.
scaled_mm_fp8is reached directly rather than through the registry, so it testship.has_wmma()explicitly.The default arch list is every RDNA2/3/3.5/4 target ROCm supports.
COMFY_HIP_ARCHSstill overrides it, and a local build for the one GPU it can see is much faster.I've added RDNA2 support as well: ROCm ships the device bitcode for
gfx103xand the toolchain builds it, and in my experience there are many RDNA2 users in the AMD community.Testing
tests/test_hip_wmma.pyon a gfx1200 (RX 9060 XT), Windows, ROCm 7.15, torch 2.12. The suite is arch-aware: the GEMM tests skip themselves on a device without matrix cores.v_wmmainstructions and each of the 7 gfx103x targets emits zero plus the traps. A missing arch macro would not fail the build, it would silently drop that GPU into the no-WMMA stub, which is how I caught thatgfx1103was missing.bdist_wheelon Windows producescomfy_kitchen-0.2.21-cp312-abi3-win_amd64.whlcontainingbackends/hip/_C.abi3.pyd, so theabi3tag does survive the HIP extension.ruffclean.v_wmma, the 7 gfx103x emit none, as intended.backends/cuda/_C*.soandbackends/hip/_C*.so, and the "Verify Linux wheels contain CUDA and HIP extensions" step - which hard-fails if either is missing - passes; the Windows jobs do the same for.pyd. The CPU-only test jobs then install that wheel on runners with neither GPU runtime and pass, which exercises the "each backend withdraws itself at import" path.Kernel changes to improve RDNA3 performance can only be properly validated on an actual RDNA3 GPU. If you have access to RDNA3 hardware and would like to help, feel free to fork the amd/hip-backend branch, implement and test RDNA3-specific improvements there, and open a PR against that branch. Once those changes are merged, they'll automatically become part of this PR.
RDNA2 users should still use Triton for int8, as rocBLAS doesn't contain int kernels, and hipBLASLt doesn't support RDNA2. As things are today, the INT8-Fast-ROCM node is the best option.