[Bug]: mixtureOfExpertsTest crashes with SEGV when a parameter set has no MoE GEMM tactics on the running SM

System Info

• Container: nvcr.io/nvidia/tensorrt-llm/devel:1.3.0rc26
• Commit: 63d217f2 on main (version 1.3.0rc27)
• GPU: NVIDIA B300 SXM6, compute capability 10.3
• CUDA 13.2, GCC 13, CMake 4.0.3

Built with the sanitizer option the repository already ships:

python3 scripts/build_wheel.py --skip_building_wheel --skip-stubs \
  --cuda_architectures 103-real --build_type RelWithDebInfo --fast_build \
  -D "SANITIZE=address" -D "BUILD_TESTS=ON"

Who can help?

No response

Information

• The official example scripts
• My own modified scripts

Tasks

• An officially supported task in the examples folder (such as GLUE/SQuAD, ...)
• My own task or dataset (give details below)

Reproduction

ASAN_OPTIONS=protect_shadow_gap=0 ./cpp/build_RelWithDebInfo/tests/unit_tests/kernels/mixtureOfExpertsTest

Expected behavior

Either the parameter set runs, or the test reports that no MoE GEMM tactic exists for
this SM and type combination.

actual behavior

The process dies with a null dereference:

WARNING: Could not find config for sm version 103          (printed 16 times)
[ RUN      ] MixtureOfExpertsTest/5.Permute

ERROR: AddressSanitizer: SEGV on unknown address 0x000000000040
SUMMARY: AddressSanitizer: SEGV /usr/include/c++/13/bits/stl_pair.h:687 in
  std::pair<tensorrt_llm::_v1::cutlass_extensions::CutlassGemmConfig,
            tensorrt_llm::_v1::cutlass_extensions::CutlassGemmConfig>::pair<...>

Stack, trimmed to this repository's frames:

#1 MixtureOfExpertsTest<WeightParams<__nv_fp8_e4m3, cutlass::integer_subbyte<4, false>,
     __nv_bfloat16, void, __nv_bfloat16>>::selectTacticsForArch(int, bool)
#2 ...::getAllTileConfigsToTest()
#3 ...::BasicPermuteTest(int, long, ...)
#4 MixtureOfExpertsTest_Permute_Test<...>::TestBody()

The crash is in the fallback that exists for exactly this situation,
cpp/tests/unit_tests/kernels/mixtureOfExpertsTest.cu:1295-1303:

if (it1 == tactics1.end() || it2 == tactics2.end())
{
    // Fall back to any tactic
    std::cout << "WARNING: Could not find config for sm version " << sm << std::endl;
    it1 = (it1 == tactics1.end()) ? tactics1.begin() : it1;
    it2 = (it2 == tactics2.end()) ? tactics2.begin() : it2;
}

return std::pair(*it1, *it2);

The guard handles "no tactic matched the filter". It does not handle "the tactic list is
empty", where begin() == end() and the dereference below is undefined. On this device
mMoERunner.getTactics() returns an empty list for the
<__nv_fp8_e4m3 weights, int4, __nv_bfloat16> parameter set (MixtureOfExpertsTest/5);
getFilteredConfigs does not filter anything here because its extra filter only applies to
sm == 89 || sm >= 120.

AddressSanitizer is what turns this into a clean SEGV with a stack. Without it the same
code reads past the end of an empty vector, so the behaviour is undefined either way.

Note on the build architecture

I first saw this with a 90-real build on the same device and assumed the architecture
mismatch was the cause. It is not: rebuilding with --cuda_architectures 103-real on the
matching device reproduces it identically, and the warning count rises from 6 to 16.

Possible fix

Four lines, right where the fallback already is: check for an empty list and report it
instead of dereferencing. I have a patch ready and will open it as a PR against this issue.

additional notes

Found while running the repository's own SANITIZE=address build, which appears not to be
exercised by CI.

Start in cpp/tests/unit_tests/kernels/mixtureOfExpertsTest.cu:1295-1303 and inspect selectTacticsForArch, then reproduce with the listed mixtureOfExpertsTest command under AddressSanitizer. Verify the empty-tactic parameter set no longer dereferences an empty list and that the test reports the missing configuration instead of crashing.