Remove unreachable test_direct_mapped_lxu_cache_lookup_large_grid#5900
Closed
q10 wants to merge 2 commits into
Closed
Remove unreachable test_direct_mapped_lxu_cache_lookup_large_grid#5900q10 wants to merge 2 commits into
q10 wants to merge 2 commits into
Conversation
pytorch#5899) Summary: `test_lru_cache_insert_large_grid` (added by D105282095) hardcodes the LXU cache associativity as `32`. The split-embeddings LXU cache is set-associative with associativity == warp size == `DEFAULT_ASSOC` (32 on NVIDIA, 64 on AMD). On AMD wavefront64 (gfx942 / MI300) `lru_cache_insert_kernel` strides cache rows by `kWarpSize = 64` and writes `lxu_cache_state` / `lxu_cache_weights` / `lru_state` for `slot` in `[0, 64)`, indexing past the 32-wide test allocations -> out-of-bounds -> non-deterministic memory corruption -> flaky `assertEqual(lru_state != time_stamp, 0)` failures in OSS ROCm CI (see P2378242263). On NVIDIA (32 == 32) the allocation matches the kernel, so the test passed. Fix (test-only; no kernel/production change): - Size the three cache tensors and assertions by `DEFAULT_ASSOC` instead of the literal `32`, matching the established pattern in `lxu_cache_test.py` and `nbit_cache_test.py`, so the allocation width matches the kernel's `kWarpSize` associativity on both platforms. - Fix `torch.accelerator.current_accelerator("cuda")` -> `current_accelerator()` (the string was silently coerced to `check_available=True`; flagged by ai_diff_reviewer). - Generalize the docstring's NVIDIA-specific (32) grid math. Differential Revision: D108540654
Summary: `test_direct_mapped_lxu_cache_lookup_large_grid` (added by D105286434) tests an unreachable code path and fails on ROCm in OSS CI (see P2378802042). `direct_mapped_lxu_cache_lookup_kernel` takes `linear_cache_indices` as a `PackedTensorAccessor32` (lxu_cache.cu:326,530). Building a 32-bit accessor runs `validate_tensor`, which TORCH_CHECKs `numel < INT32_MAX` (~2.1B). The test passes `N = 2**32 + 1` (~4.3B) to try to exceed HIP's 2**32 thread-per-launch limit, but the accessor rejects the tensor at ~2.1B elements, before the kernel launch is ever reached: RuntimeError: lxu_cache.hip(538): direct_mapped_lxu_cache_lookup_kernel: Tensor 'linear_cache_indices' numel needs to be smaller than int32_t max; otherwise, please use [packed_]accessor64 So the host-side grid cap added by D105286434 (`cap_grid_dim_x_from_workload`) is unreachable on this kernel: any N large enough to need it (>= 2**32) is rejected by the 32-bit accessor (>= 2**31) first. The regression test therefore cannot validate anything and only breaks ROCm CI. This mirrors the abandoned D105202994 (delinearize_unique_index_kernel), which had the identical PackedTensorAccessor32 INT32_MAX limitation. Removes the test and its now-unused `gpu_memory_lt_gb` import. No production change; the cap is left in place (harmless, and correct once the accessor is upgraded to 64-bit). Differential Revision: D108559251
Contributor
|
@q10 has exported this pull request. If you are a Meta employee, you can view the originating Diff in D108559251. |
q10
added a commit
to q10/FBGEMM
that referenced
this pull request
Jun 15, 2026
…torch#5900) Summary: X-link: facebookresearch/FBGEMM#2821 `test_direct_mapped_lxu_cache_lookup_large_grid` (added by D105286434) tests an unreachable code path and fails on ROCm in OSS CI (see P2378802042). `direct_mapped_lxu_cache_lookup_kernel` takes `linear_cache_indices` as a `PackedTensorAccessor32` (lxu_cache.cu:326,530). Building a 32-bit accessor runs `validate_tensor`, which TORCH_CHECKs `numel < INT32_MAX` (~2.1B). The test passes `N = 2**32 + 1` (~4.3B) to try to exceed HIP's 2**32 thread-per-launch limit, but the accessor rejects the tensor at ~2.1B elements, before the kernel launch is ever reached: RuntimeError: lxu_cache.hip(538): direct_mapped_lxu_cache_lookup_kernel: Tensor 'linear_cache_indices' numel needs to be smaller than int32_t max; otherwise, please use [packed_]accessor64 So the host-side grid cap added by D105286434 (`cap_grid_dim_x_from_workload`) is unreachable on this kernel: any N large enough to need it (>= 2**32) is rejected by the 32-bit accessor (>= 2**31) first. The regression test therefore cannot validate anything and only breaks ROCm CI. This mirrors the abandoned D105202994 (delinearize_unique_index_kernel), which had the identical PackedTensorAccessor32 INT32_MAX limitation. Removes the test and its now-unused `gpu_memory_lt_gb` import. No production change; the cap is left in place (harmless, and correct once the accessor is upgraded to 64-bit). Reviewed By: henrylhtsang Differential Revision: D108559251
Contributor
|
This pull request has been merged in 1864530. |
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
2 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.
Summary:
test_direct_mapped_lxu_cache_lookup_large_grid(added by D105286434)tests an unreachable code path and fails on ROCm in OSS CI
(see P2378802042).
direct_mapped_lxu_cache_lookup_kerneltakeslinear_cache_indicesas a
PackedTensorAccessor32(lxu_cache.cu:326,530). Building a 32-bitaccessor runs
validate_tensor, which TORCH_CHECKsnumel < INT32_MAX(~2.1B). The test passesN = 2**32 + 1(~4.3B) totry to exceed HIP's 2**32 thread-per-launch limit, but the accessor
rejects the tensor at ~2.1B elements, before the kernel launch is ever
reached:
RuntimeError: lxu_cache.hip(538): direct_mapped_lxu_cache_lookup_kernel:
Tensor 'linear_cache_indices' numel needs to be smaller than int32_t
max; otherwise, please use [packed_]accessor64
So the host-side grid cap added by D105286434
(
cap_grid_dim_x_from_workload) is unreachable on this kernel: any Nlarge enough to need it (>= 232) is rejected by the 32-bit accessor
(>= 231) first. The regression test therefore cannot validate
anything and only breaks ROCm CI.
This mirrors the abandoned D105202994 (delinearize_unique_index_kernel),
which had the identical PackedTensorAccessor32 INT32_MAX limitation.
Removes the test and its now-unused
gpu_memory_lt_gbimport. Noproduction change; the cap is left in place (harmless, and correct once
the accessor is upgraded to 64-bit).
Differential Revision: D108559251