[SreecharanGundaboluAMD]

Motivation

LSTM backward weights produced wrong bias gradients (~30% too small) on gfx1250 (mi450/RDNA4), causing test_variable_sequence_cuda_float32 to fail in PyTorch. The bug only manifested with variable-length (packed) sequences and was gfx1250-specific — gfx942 (MI300X) with identical inputs passed.

Root cause: the legacy ComposableKernel code vendored inside MIOpen (src/legacy_composable_kernel/) constructs AMD buffer resource descriptors using two fields that changed format on gfx1250:

1. CK_BUFFER_RESOURCE_3RD_DWORD (config.hpp): gfx1250 was grouped with gfx10/11 and assigned 0x31014000. gfx1250 (RDNA4/gfx12) requires 0 — confirmed against CK's implementation (composablekernel/include/ck/ck.hpp). With the wrong word-3, buffer_load treats in-range elements as out-of-bounds and returns 0, silently dropping reduction terms.

2. make_wave_buffer_resource (amd_buffer_addressing.hpp): gfx1250 changed the buffer resource descriptor layout — num_records became a 45-bit field split across dword1[31:25] and dword2[31:0], instead of fitting in dword2 alone. The old single-dword write puts bounds in the wrong bit positions, causing the hardware to see a garbage range and silently return 0 for out-of-bounds elements. Also confirmed by CK's implementation.

Technical Details

Two files changed, both in src/legacy_composable_kernel/composable_kernel/include/utility/:

config.hpp — split gfx1250 out of the gfx10/11 group and assign the correct word-3:

// Before: gfx1250 grouped with gfx10/11, gets 0x31014000 (wrong)
#elif defined(...GFX1201) || defined(CK_AMD_GPU_GFX1250)
#define CK_BUFFER_RESOURCE_3RD_DWORD 0x31014000

// After:
#elif defined(...GFX1201)
#define CK_BUFFER_RESOURCE_3RD_DWORD 0x31014000
#elif defined(CK_AMD_GPU_GFX1250)
#define CK_BUFFER_RESOURCE_3RD_DWORD 0          // matches ck.hpp for __gfx125__

amd_buffer_addressing.hpp — add gfx1250 branch in make_wave_buffer_resource with the correct 45-bit split num_records encoding, matching what modern CK's amd_buffer_addressing_builtins.hpp does for __gfx125__:

#if defined(CK_AMD_GPU_GFX1250)
    wave_buffer_resource.address(Number<0>{}) = const_cast<remove_cv_t<T>*>(p_wave);
    uint64_t num_records = static_cast<uint64_t>(element_space_size) * sizeof(T);
    wave_buffer_resource.range(Number<1>{}) |= static_cast<int32_t>((num_records & 0x7f) << 25);
    wave_buffer_resource.range(Number<2>{}) = static_cast<int32_t>(num_records >> 7);
    wave_buffer_resource.config(Number<3>{}) = CK_BUFFER_RESOURCE_3RD_DWORD;
#else
    // existing pre-gfx12 path unchanged

Why both changes are needed: word-3 controls the descriptor format/control bits; num_records controls the buffer bounds. A descriptor with correct word-3 but wrong range encoding still causes the hardware to misinterpret bounds. Both fields must be correct.

Why the legacy file needs the fix independently: Modern CK routes gfx1250 buffer loads through amd_buffer_addressing_builtins.hpp when __clang_major__ >= 20 (build environment uses clang 23). The legacy path in MIOpen has no such gate — it always uses amd_buffer_addressing.hpp — so the fix must live there directly.

Known follow-up (TODO in config.hpp): gfx11 (1100–1153) and gfx120 (1200/1201) are also assigned 0x31014000 in legacy config.hpp, but modern CK uses 0x31004000 for them (word-3 format changed RDNA2→RDNA3). Only gfx1250 has been verified here; the gfx11/gfx120 discrepancy is tracked in the TODO comment.

Risk Assessment

Risk Level: 🟢 Low

Impacted Components

• src/legacy_composable_kernel/composable_kernel/include/utility/config.hpp — buffer resource descriptor word-3 definition; affects all legacy-CK kernels compiled for gfx1250.
• src/legacy_composable_kernel/composable_kernel/include/utility/amd_buffer_addressing.hpp — make_wave_buffer_resource; affects all code using legacy-CK buffer loads on gfx1250.
• On all other architectures: no change — both fixes are behind #if defined(CK_AMD_GPU_GFX1250).

Potential Side Effects

• On gfx1250: all legacy-CK buffer load/store operations now use a correctly formed descriptor. This can only improve correctness — wrong descriptors cause silent data corruption or zeros, never correct results, so there is no risk of behavioral regression on gfx1250.
• On all other arches: changes are completely isolated by the preprocessor guard.

Mitigation Steps

• Both fixes are strictly additive for gfx1250 and leave all other arches unchanged.
• Verified with MIOPEN_RNNWRW_REDUCTION=0 and =2 as cross-checks (both use non-buffer-load reduction paths and pass, confirming the RNN logic itself is correct).
• Verified gfx942 is unaffected.

Test Plan

1. PyTorch end-to-end on gfx1250 where original failure was observed:

   python test/test_nn.py TestNNDeviceTypeCUDA.test_variable_sequence_cuda_float32
   

Test Result

• PyTorch test_variable_sequence_cuda_float32 on gfx1250: passes (was: AssertionError: 16/16 elements mismatched, max abs diff 0.988).
• gfx942 (MI300X): unaffected, continues to pass with default reduction algo (MIOPEN_RNNWRW_REDUCTION=1).

Submission Checklist

• Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

[JonathanLichtnerAMD]

LGTM, but could we add a test that reproduces this issue? (We could do that as a follow-on change since we don't want to delay this fix.)

[SreecharanGundaboluAMD]

LGTM, but could we add a test that reproduces this issue? (We could do that as a follow-on change since we don't want to delay this fix.)

Good point. We should run gtests on gfx1250 first, as this may already be covered by an existing test. We have not really run gtests on gfx1250 before due to limited hardware availability. We should also create a ticket to track this.

[adickin-amd]

Merging as change does not affect gfx1151 and windows.