Comfy-Org/comfy-kitchen

Fork 0

mirror of https://github.com/Comfy-Org/comfy-kitchen.git synced 2026-11-04 23:37:42 +00:00

[GH-ISSUE #214] int8 attention returns NaN for prepared dense masks containing -inf (gfx1103, ROCm) #62

New issue

Closed

opened 2026-10-26 22:19:50 +00:00 by bsakaguchi · 4 comments

bsakaguchi commented

2026-10-26 22:19:50 +00:00

Owner

Originally created by @LuXuxue on GitHub (Sep 29, 2026).
Original GitHub issue: https://github.com/Comfy-Org/comfy-kitchen/issues/214

Summary

int8_attention returns NaN for the output rows whenever an attn_mask drops keys
via -inf, on the prepared dense mask path introduced in #208. A large finite
negative bias produces correct output; only -inf fails. The same mask expressed as
torch.bool also NaNs, because the bit-packed path synthesizes -inf as the bias
internally.

This is reproducible on an unmodified checkout of main at 888b13e (v0.2.36):

tests/test_int8_attention.py: 541 failed, 445 passed, 5 skipped

failing test	count
`test_hip_compact_dense_bool_matches_float_bias`	192
`test_hip_dense_16_accuracy_and_snapshot`	192
`test_hip_fused_dense_mask_matches_snapshot`	96
`test_hip_dense_16_preserves_constant_values`	36
`test_int8_attention_long_masked_sequence`	12
`test_int8_attention_mask_gqa_broadcast_and_fully_masked_row`	9
`test_hip_fused_dense_mask_graph_reads_updated_bias`	4

Environment


GPU	AMD Radeon 780M Graphics, `gfx1103` (RDNA3, 6 WGP)
ROCm	7.15.26333
PyTorch	2.13.0+rocm10.0.0
checkout	`888b13e` (`git describe`: v0.2.36), unmodified
build	`HIP_PLATFORM=amd HIP_ARCHS=gfx1103 python setup.py build_ext --inplace`

Reproducer

import torch
import comfy_kitchen as ck

torch.manual_seed(183)
q = torch.randn(2, 4, 129, 64, device="cuda", dtype=torch.float16)
k = torch.randn(2, 2, 257, 64, device="cuda", dtype=torch.float16)
v = torch.randn(2, 2, 257, 64, device="cuda", dtype=torch.float16)

# keep ~85% of keys; -inf is the idiomatic way to say "drop this key"
keep = torch.rand(1, 1, 129, 257, device="cuda") > 0.15
keep[..., 0, :] = False

inf_bias    = torch.where(keep, 0.0, -torch.inf)
finite_bias = torch.where(keep, 0.0, -60000.0)
bool_mask   = keep

for name, m in [("additive -inf", inf_bias),
                ("additive finite", finite_bias),
                ("bool", bool_mask)]:
    out = ck.int8_attention(q, k, v, scale=0.0, attn_mask=m)
    print(f"{name:>15}: nonfinite {int((~torch.isfinite(out)).sum()):>6}/{out.numel()}")

Output on 888b13e:

additive -inf : nonfinite  10240/66048
additive finite: nonfinite      0/66048
bool          : nonfinite  10240/66048

What I have established

1. It is -inf specifically, not a large-magnitude bias. Sweeping the additive
bias value on a fixed mask, every finite value is fine down to -1e9:

bias	nonfinite
`-inf`	9216/66048
`-1e9`	0
`-1e6`	0
`-72135` (= `-kMaskedScoreBase2`)	0
`-50000` (= `kMaskedScore`)	0
`-1e4`	0

So there is no magnitude threshold to tune around; the finite sentinel already used
elsewhere in the kernel is sufficient.

2. It is the prepared path, not the short-key path. kv_length <= 64 skips mask
packing and takes MaskMode::kCustom, which is clean. The first failing length is 65:

`kv_length`	nonfinite
16, 32, 64	0 (ok)
77	65536/66048
128	65536/66048
257	9216/66048
1153	11264/70144

Affects head_dim 64, 128 and 256 alike, and both the fused int8_attention and
prequantize_int8_attention + int8_attention_from_prequantized.

3. Any -inf in a row poisons the entire row, not just the dropped positions.
With exactly one key kept, at any position, the whole output is NaN:

only key 0 kept    : nonfinite 66048/66048
only key 128 kept  : nonfinite 66048/66048
only key 256 kept  : nonfinite 66048/66048
first 256 kept     : nonfinite 66048/66048
all 257 kept       : nonfinite      0/66048   <- control

Rows that are fully dropped are NaN rather than the zeros the
row_valid bookkeeping implies they should be.

4. The -inf reaches the score as a bias. In int8_attn.hip the packed mask is
read and fma'd into the score:

prob[fk][e] = fmaf(static_cast<float>(s_acc[fk][e]), tile_scale_k[fk], bias);

prepare_dense_mask_kernel stores __uint_as_float(0xff800000u) (i.e. -inf) for a
dropped key, and prepare_key_mask_kernel likewise. The bit-packed path does not even
read a bias: bool_mask_bias constructs one,

return __uint_as_float(static_cast<uint32_t>(dropped) & 0xff800000u);

which explains why the bool mask fails identically.

5. -ffast-math is not folding the guard. I expected the kDense validity test to
be optimized away and initially reported that as the cause; that was wrong. Compiling
int8_attn.hip for gfx1103 with and without -ffast-math and counting the emitted
instructions:

build	kernels	`0x7f800000` compares	`v_cmp_*_f32`	`v_exp_f32`
`-ffast-math` on	60	396	2308	2176
`-ffast-math` off	60	396	2344	2176

The raw-bits test at int8_attn.hip:692,

const bool valid_tile = (__float_as_uint(tile_m) & 0x7f800000u) != 0x7f800000u;
row_valid[fq] = row_valid[fq] || valid_tile;
tile_m = valid_tile ? tile_m : kMaskedScore;

survives in both builds, and the guard runs after tile_m has been reduced, so it
can only repair a tile whose maximum is already -inf. It does not sanitize the
individual -inf bias values inside an otherwise-valid tile, and tile_m for such a
tile is finite, so the guard does not fire. I do not currently see where the -inf
bias turns into a NaN instead of the expected exp2(-inf - tile_m) == 0, and I would
rather ask than guess.

One observation that may be the actual bug

The prepared mask is packed in the base-two domain (prepare_dense_mask_kernel
stores bias * kLog2e), and the kernel accounts for that asymmetrically:

// int8_attn.hip:699
if constexpr (kPreparedDense16) tile_m *= kLog2e;

kLog2e is applied to tile_m only for the 16-bit prepared path, and not for the
float32 prepared path — yet the float32 bias it is added to was multiplied by
kLog2e at packing time. The sentinel substituted by the validity test,
kMaskedScore = -50000, is likewise never rescaled.

This is consistent with the observed split: the clamp fixes everything routed through
the 16-bit and bit-packed read sites, and leaves the float32 dense path failing. I have
not confirmed it is the NaN source, and it is a correctness question independent of the
NaN, so I am flagging it rather than claiming it.

Suggested direction

Clamping the bias at the read sites — replacing any non-finite packed value with the
existing finite sentinel, whose sign is preserved so +inf still means "strongly
attend" — fixes the bit-packed and 16-bit paths completely:

// Uses kMaskedScore, which already exists in sage_common.h.
template <bool kBF16>
__forceinline__ __device__ float dense16_mask_bias(uint16_t bits) {
    const uint16_t inf = kBF16 ? 0x7f80u : 0x7c00u;
    const float finite = kBF16 ? __uint_as_float(static_cast<uint32_t>(bits) << 16)
                               : __half2float(*reinterpret_cast<const __half*>(&bits));
    const float clamped = (bits & 0x8000u) != 0 ? kMaskedScore : -kMaskedScore;
    return (bits & inf) == inf ? clamped : finite;
}

__forceinline__ __device__ float dense32_mask_bias(float bias) {
    const uint32_t bits = __float_as_uint(bias);
    return (bits & 0x7f800000u) == 0x7f800000u
        ? ((bits & 0x80000000u) != 0 ? kMaskedScore : -kMaskedScore)
        : bias;
}

plus changing bool_mask_bias to return -kMaskedScore * (float)dropped instead of
the synthesized 0xff800000u. Note the base-two-domain question in the previous
section applies to whichever sentinel is chosen, since kMaskedScore is in the
natural domain.

Measured on 888b13e with that clamp applied to the three dense read sites and
bool_mask_bias returning a finite sentinel: 541 -> 423 failed, 0 newly broken.
Per test:

test	before	after
`test_hip_fused_dense_mask_matches_snapshot`	96	0
`test_hip_fused_dense_mask_graph_reads_updated_bias`	4	0
`test_hip_dense_16_preserves_constant_values`	36	30
`test_int8_attention_long_masked_sequence`	12	6
`test_int8_attention_mask_gqa_broadcast_and_fully_masked_row`	9	3
`test_hip_compact_dense_bool_matches_float_bias`	192	192
`test_hip_dense_16_accuracy_and_snapshot`	192	192

Note that test_hip_compact_dense_bool_matches_float_bias does not improve even
though the bool path itself now returns finite values — the test also compares against
expected, which is the additive -inf path, so it keeps failing. The two 192-count
tests that remain are the ones exercising the additive path.

The remainder are the additive -inf path, which the read-site clamp does not
reach, consistent with point 4: that -inf is already in the packed buffer, and
something between the packed read and exp2 turns it into NaN. That is the part I
have not solved.

Questions

Is this AMD-specific?
Is the kLog2e asymmetry at int8_attn.hip:699 intentional for the float32 dense
path, given the bias arrives pre-scaled?
For the additive path, where between the packed read and exp2 does -inf become
NaN rather than the expected exp2(-inf - tile_m) == 0? Clamping the read sites
does not change the outcome there, so the -inf must be arriving by another route
that I have not identified.

Originally created by @LuXuxue on GitHub (Sep 29, 2026). Original GitHub issue: https://github.com/Comfy-Org/comfy-kitchen/issues/214 ## Summary `int8_attention` returns NaN for the output rows whenever an `attn_mask` drops keys via `-inf`, on the **prepared dense** mask path introduced in #208. A large finite negative bias produces correct output; only `-inf` fails. The same mask expressed as `torch.bool` also NaNs, because the bit-packed path synthesizes `-inf` as the bias internally. This is reproducible on an unmodified checkout of `main` at `888b13e` (v0.2.36): ``` tests/test_int8_attention.py: 541 failed, 445 passed, 5 skipped ``` | failing test | count | |---|---| | `test_hip_compact_dense_bool_matches_float_bias` | 192 | | `test_hip_dense_16_accuracy_and_snapshot` | 192 | | `test_hip_fused_dense_mask_matches_snapshot` | 96 | | `test_hip_dense_16_preserves_constant_values` | 36 | | `test_int8_attention_long_masked_sequence` | 12 | | `test_int8_attention_mask_gqa_broadcast_and_fully_masked_row` | 9 | | `test_hip_fused_dense_mask_graph_reads_updated_bias` | 4 | ## Environment | | | |---|---| | GPU | AMD Radeon 780M Graphics, `gfx1103` (RDNA3, 6 WGP) | | ROCm | 7.15.26333 | | PyTorch | 2.13.0+rocm10.0.0 | | checkout | `888b13e` (`git describe`: v0.2.36), unmodified | | build | `HIP_PLATFORM=amd HIP_ARCHS=gfx1103 python setup.py build_ext --inplace` | ## Reproducer ```python import torch import comfy_kitchen as ck torch.manual_seed(183) q = torch.randn(2, 4, 129, 64, device="cuda", dtype=torch.float16) k = torch.randn(2, 2, 257, 64, device="cuda", dtype=torch.float16) v = torch.randn(2, 2, 257, 64, device="cuda", dtype=torch.float16) # keep ~85% of keys; -inf is the idiomatic way to say "drop this key" keep = torch.rand(1, 1, 129, 257, device="cuda") > 0.15 keep[..., 0, :] = False inf_bias = torch.where(keep, 0.0, -torch.inf) finite_bias = torch.where(keep, 0.0, -60000.0) bool_mask = keep for name, m in [("additive -inf", inf_bias), ("additive finite", finite_bias), ("bool", bool_mask)]: out = ck.int8_attention(q, k, v, scale=0.0, attn_mask=m) print(f"{name:>15}: nonfinite {int((~torch.isfinite(out)).sum()):>6}/{out.numel()}") ``` Output on `888b13e`: ``` additive -inf : nonfinite 10240/66048 additive finite: nonfinite 0/66048 bool : nonfinite 10240/66048 ``` ## What I have established **1. It is `-inf` specifically, not a large-magnitude bias.** Sweeping the additive bias value on a fixed mask, every finite value is fine down to `-1e9`: | bias | nonfinite | |---|---| | `-inf` | 9216/66048 | | `-1e9` | 0 | | `-1e6` | 0 | | `-72135` (= `-kMaskedScoreBase2`) | 0 | | `-50000` (= `kMaskedScore`) | 0 | | `-1e4` | 0 | So there is no magnitude threshold to tune around; the finite sentinel already used elsewhere in the kernel is sufficient. **2. It is the prepared path, not the short-key path.** `kv_length <= 64` skips mask packing and takes `MaskMode::kCustom`, which is clean. The first failing length is 65: | `kv_length` | nonfinite | |---|---| | 16, 32, 64 | 0 (ok) | | 77 | 65536/66048 | | 128 | 65536/66048 | | 257 | 9216/66048 | | 1153 | 11264/70144 | Affects `head_dim` 64, 128 and 256 alike, and both the fused `int8_attention` and `prequantize_int8_attention` + `int8_attention_from_prequantized`. **3. Any `-inf` in a row poisons the entire row, not just the dropped positions.** With exactly one key kept, at any position, the whole output is NaN: ``` only key 0 kept : nonfinite 66048/66048 only key 128 kept : nonfinite 66048/66048 only key 256 kept : nonfinite 66048/66048 first 256 kept : nonfinite 66048/66048 all 257 kept : nonfinite 0/66048 <- control ``` Rows that are *fully* dropped are NaN rather than the zeros the `row_valid` bookkeeping implies they should be. **4. The `-inf` reaches the score as a bias.** In `int8_attn.hip` the packed mask is read and fma'd into the score: ``` prob[fk][e] = fmaf(static_cast<float>(s_acc[fk][e]), tile_scale_k[fk], bias); ``` `prepare_dense_mask_kernel` stores `__uint_as_float(0xff800000u)` (i.e. `-inf`) for a dropped key, and `prepare_key_mask_kernel` likewise. The bit-packed path does not even read a bias: `bool_mask_bias` *constructs* one, ``` return __uint_as_float(static_cast<uint32_t>(dropped) & 0xff800000u); ``` which explains why the bool mask fails identically. **5. `-ffast-math` is not folding the guard.** I expected the `kDense` validity test to be optimized away and initially reported that as the cause; that was wrong. Compiling `int8_attn.hip` for `gfx1103` with and without `-ffast-math` and counting the emitted instructions: | build | kernels | `0x7f800000` compares | `v_cmp_*_f32` | `v_exp_f32` | |---|---|---|---|---| | `-ffast-math` on | 60 | 396 | 2308 | 2176 | | `-ffast-math` off | 60 | 396 | 2344 | 2176 | The raw-bits test at `int8_attn.hip:692`, ``` const bool valid_tile = (__float_as_uint(tile_m) & 0x7f800000u) != 0x7f800000u; row_valid[fq] = row_valid[fq] || valid_tile; tile_m = valid_tile ? tile_m : kMaskedScore; ``` survives in both builds, and the guard runs *after* `tile_m` has been reduced, so it can only repair a tile whose maximum is already `-inf`. It does not sanitize the individual `-inf` bias values inside an otherwise-valid tile, and `tile_m` for such a tile is finite, so the guard does not fire. I do not currently see where the `-inf` bias turns into a NaN instead of the expected `exp2(-inf - tile_m) == 0`, and I would rather ask than guess. ## One observation that may be the actual bug The prepared mask is packed in the **base-two** domain (`prepare_dense_mask_kernel` stores `bias * kLog2e`), and the kernel accounts for that asymmetrically: ``` // int8_attn.hip:699 if constexpr (kPreparedDense16) tile_m *= kLog2e; ``` `kLog2e` is applied to `tile_m` only for the **16-bit** prepared path, and not for the **float32** prepared path — yet the float32 bias it is added to was multiplied by `kLog2e` at packing time. The sentinel substituted by the validity test, `kMaskedScore = -50000`, is likewise never rescaled. This is consistent with the observed split: the clamp fixes everything routed through the 16-bit and bit-packed read sites, and leaves the float32 dense path failing. I have not confirmed it is the NaN source, and it is a correctness question independent of the NaN, so I am flagging it rather than claiming it. ## Suggested direction Clamping the bias at the read sites — replacing any non-finite packed value with the existing finite sentinel, whose sign is preserved so `+inf` still means "strongly attend" — fixes the bit-packed and 16-bit paths completely: ```cpp // Uses kMaskedScore, which already exists in sage_common.h. template <bool kBF16> __forceinline__ __device__ float dense16_mask_bias(uint16_t bits) { const uint16_t inf = kBF16 ? 0x7f80u : 0x7c00u; const float finite = kBF16 ? __uint_as_float(static_cast<uint32_t>(bits) << 16) : __half2float(*reinterpret_cast<const __half*>(&bits)); const float clamped = (bits & 0x8000u) != 0 ? kMaskedScore : -kMaskedScore; return (bits & inf) == inf ? clamped : finite; } __forceinline__ __device__ float dense32_mask_bias(float bias) { const uint32_t bits = __float_as_uint(bias); return (bits & 0x7f800000u) == 0x7f800000u ? ((bits & 0x80000000u) != 0 ? kMaskedScore : -kMaskedScore) : bias; } ``` plus changing `bool_mask_bias` to return `-kMaskedScore * (float)dropped` instead of the synthesized `0xff800000u`. Note the base-two-domain question in the previous section applies to whichever sentinel is chosen, since `kMaskedScore` is in the natural domain. Measured on `888b13e` with that clamp applied to the three dense read sites and `bool_mask_bias` returning a finite sentinel: **541 -> 423 failed, 0 newly broken.** Per test: | test | before | after | |---|---|---| | `test_hip_fused_dense_mask_matches_snapshot` | 96 | **0** | | `test_hip_fused_dense_mask_graph_reads_updated_bias` | 4 | **0** | | `test_hip_dense_16_preserves_constant_values` | 36 | 30 | | `test_int8_attention_long_masked_sequence` | 12 | 6 | | `test_int8_attention_mask_gqa_broadcast_and_fully_masked_row` | 9 | 3 | | `test_hip_compact_dense_bool_matches_float_bias` | 192 | 192 | | `test_hip_dense_16_accuracy_and_snapshot` | 192 | 192 | Note that `test_hip_compact_dense_bool_matches_float_bias` does **not** improve even though the bool path itself now returns finite values — the test also compares against `expected`, which is the additive `-inf` path, so it keeps failing. The two 192-count tests that remain are the ones exercising the additive path. The remainder are the **additive** `-inf` path, which the read-site clamp does not reach, consistent with point 4: that `-inf` is already in the packed buffer, and something between the packed read and `exp2` turns it into NaN. That is the part I have not solved. ## Questions 1. Is this AMD-specific? 2. Is the `kLog2e` asymmetry at `int8_attn.hip:699` intentional for the float32 dense path, given the bias arrives pre-scaled? 3. For the additive path, where between the packed read and `exp2` does `-inf` become NaN rather than the expected `exp2(-inf - tile_m) == 0`? Clamping the read sites does not change the outcome there, so the `-inf` must be arriving by another route that I have not identified.

bsakaguchi closed this issue

2026-10-26 22:19:51 +00:00

bsakaguchi commented

2026-10-26 22:19:51 +00:00

Author

Owner

@tangzzycc commented on GitHub (Sep 29, 2026):

I can reproduce this issue on gfx1151 (AMD Radeon 8060S) as well, so it does not appear to be limited to gfx1103. On main@888b13e, the same input produces:

additive -inf: 11264 NaNs
additive -60000: 0 NaNs
bool mask: 11264 NaNs

I tried fixing this in the HIP kernel. My current approach determines mask validity directly from the Boolean bit or the raw FP16/BF16/FP32 exponent bits, tracks valid keys in a lane-local bitset, and prevents masked -inf values from entering FMA, max reduction, or exp2. Masked probabilities are explicitly set to zero, while fully masked rows still produce exact zero outputs.

After the change, all three mask variants produce zero NaNs. Regression tests also pass for Boolean, FP16, BF16, and FP32 masks across head dimensions 64, 128, and 256, covering both fused and prequantized paths.

The current fix has a measurable performance cost in preliminary gfx1151 A/B tests:

Q=129, KV=257, HD=64: effectively unchanged;
Q=129, KV=257, HD=256: approximately 3% slower;
Q=256, KV=2048, HD=128: approximately 9–12% slower for the prequantized attention-kernel path, and 8–14% slower for the complete ck.int8_attention operator call including quantization, mask preparation, and the attention kernel.

The regression appears to come from the additional validity tracking and branch/register scheduling pressure.

@tangzzycc commented on GitHub (Sep 29, 2026): I can reproduce this issue on gfx1151 (AMD Radeon 8060S) as well, so it does not appear to be limited to gfx1103. On `main@888b13e`, the same input produces: ```text additive -inf: 11264 NaNs additive -60000: 0 NaNs bool mask: 11264 NaNs ``` I tried fixing this in the HIP kernel. My current approach determines mask validity directly from the Boolean bit or the raw FP16/BF16/FP32 exponent bits, tracks valid keys in a lane-local bitset, and prevents masked `-inf` values from entering FMA, max reduction, or `exp2`. Masked probabilities are explicitly set to zero, while fully masked rows still produce exact zero outputs. After the change, all three mask variants produce zero NaNs. Regression tests also pass for Boolean, FP16, BF16, and FP32 masks across head dimensions 64, 128, and 256, covering both fused and prequantized paths. The current fix has a measurable performance cost in preliminary gfx1151 A/B tests: - `Q=129, KV=257, HD=64`: effectively unchanged; - `Q=129, KV=257, HD=256`: approximately 3% slower; - `Q=256, KV=2048, HD=128`: approximately 9–12% slower for the prequantized attention-kernel path, and 8–14% slower for the complete `ck.int8_attention` operator call including quantization, mask preparation, and the attention kernel. The regression appears to come from the additional validity tracking and branch/register scheduling pressure.

bsakaguchi commented

2026-10-26 22:19:52 +00:00

Author

Owner

@lightzl commented on GitHub (Sep 29, 2026):

Real-world confirmation on gfx1151 (AMD Radeon 8060S, ROCm 10.1) — attached is a 2.1 MB minimal reproducer.

This is not only a unit-test issue: with comfy-kitchen 0.2.36 a Qwen-Image 2.1 (bf16) workflow in
ComfyUI 0.37.0 produces garbage images, while 0.2.35 produces correct images for the same
workflow/seed. Instrumenting all comfy_kitchen calls of the real run (4252 calls):

0.2.35: 0 calls with NaN/Inf
0.2.36: 1119 calls with NaN/Inf; first at an int8_attention call whose input is finite and
identical in both versions (0.2.35 -> finite, 0.2.36 -> NaN).
The NaN then propagates through the model and the image is corrupted (pixel correlation with the
0.2.35 output: 0.13).

Minimal case in the bundle: q=(1,1,26,128) bf16, k=v=(1,1,4162,128) bf16,
attn_mask=(1,1,26,4162) bool with 325 False entries, scale=None.
0.2.36 -> 3072/3328 NaN; 0.2.35 -> 0 NaN
Trigger follows the MASK, not the data (random q/k/v + captured mask -> NaN; captured q/k/v +
all-True mask -> finite; all-False mask -> NaN; no mask -> finite).
Run: python repro.py (also attached: README.md, cross_test.py)

Workaround for users on ROCm: pin comfy-kitchen==0.2.35 until a fixed release.

comfy-kitchen-036-nan-repro.tar.gz

@lightzl commented on GitHub (Sep 29, 2026): Real-world confirmation on gfx1151 (AMD Radeon 8060S, ROCm 10.1) — attached is a 2.1 MB minimal reproducer. This is not only a unit-test issue: with comfy-kitchen 0.2.36 a Qwen-Image 2.1 (bf16) workflow in ComfyUI 0.37.0 produces garbage images, while 0.2.35 produces correct images for the same workflow/seed. Instrumenting all comfy_kitchen calls of the real run (4252 calls): - 0.2.35: 0 calls with NaN/Inf - 0.2.36: 1119 calls with NaN/Inf; first at an int8_attention call whose input is finite and identical in both versions (0.2.35 -> finite, 0.2.36 -> NaN). The NaN then propagates through the model and the image is corrupted (pixel correlation with the 0.2.35 output: 0.13). Minimal case in the bundle: q=(1,1,26,128) bf16, k=v=(1,1,4162,128) bf16, attn_mask=(1,1,26,4162) bool with 325 False entries, scale=None. 0.2.36 -> 3072/3328 NaN; 0.2.35 -> 0 NaN Trigger follows the MASK, not the data (random q/k/v + captured mask -> NaN; captured q/k/v + all-True mask -> finite; all-False mask -> NaN; no mask -> finite). Run: python repro.py (also attached: README.md, cross_test.py) Workaround for users on ROCm: pin comfy-kitchen==0.2.35 until a fixed release. [comfy-kitchen-036-nan-repro.tar.gz](https://github.com/user-attachments/files/32812412/comfy-kitchen-036-nan-repro.tar.gz)

bsakaguchi commented

2026-10-26 22:19:52 +00:00

Author

Owner

@comfyanonymous commented on GitHub (Sep 29, 2026):

This should fix it: github.com/Comfy-Org/comfy-kitchen@19ea55b9eb

@comfyanonymous commented on GitHub (Sep 29, 2026): This should fix it: https://github.com/Comfy-Org/comfy-kitchen/commit/19ea55b9ebdaf77942dab36223e1222009d3ce11

bsakaguchi commented

2026-10-26 22:19:52 +00:00

Author

Owner

@tangzzycc commented on GitHub (Sep 30, 2026):

@comfyanonymous Thanks, I tested commit 19ea55b on gfx1151, and it fixes the issue for me. The original reproducer now reports zero non-finite values for additive -inf, additive -60000, and Boolean masks. The new test_hip_dense_mask_empty_tiles cases also all pass (24 passed).

I also compared kernel-only performance before and after the fix across head dimensions 64/128/256 and Boolean/FP16/FP32 masks, and did not observe any meaningful performance impact. Representative FP16 results (median of 9 runs):

Shape	Before (`888b13e`)	After (`19ea55b`)
`Q=129, KV=257, HD=64`	19.310 us	19.405 us
`Q=256, KV=2048, HD=128`	102.286 us	102.463 us
`Q=129, KV=257, HD=256`	42.012 us	42.057 us

@tangzzycc commented on GitHub (Sep 30, 2026): @comfyanonymous Thanks, I tested commit [`19ea55b`](https://github.com/Comfy-Org/comfy-kitchen/commit/19ea55b9ebdaf77942dab36223e1222009d3ce11) on gfx1151, and it fixes the issue for me. The original reproducer now reports zero non-finite values for additive `-inf`, additive `-60000`, and Boolean masks. The new `test_hip_dense_mask_empty_tiles` cases also all pass (`24 passed`). I also compared kernel-only performance before and after the fix across head dimensions 64/128/256 and Boolean/FP16/FP32 masks, and did not observe any meaningful performance impact. Representative FP16 results (median of 9 runs): | Shape | Before (`888b13e`) | After (`19ea55b`) | |---|---:|---:| | `Q=129, KV=257, HD=64` | 19.310 us | 19.405 us | | `Q=256, KV=2048, HD=128` | 102.286 us | 102.463 us | | `Q=129, KV=257, HD=256` | 42.012 us | 42.057 us |

bsakaguchi referenced this issue

2026-10-26 22:20:31 +00:00

[PR #62] [MERGED] Enhance CLA workflow with author-only allowlist #115

No labels

pull-request

No milestone

No project

No assignees

1 participant

Notifications

Due date

The due date is invalid or out of range. Please use the format "yyyy-mm-dd".

No due date set.

Dependencies

No dependencies set.

Reference

Comfy-Org/comfy-kitchen#62

No description provided.

Rows
Columns