From 98972fee2c9c91b6c495b10c79c99ca00a0ada04 Mon Sep 17 00:00:00 2001 From: amd-ysanchez Date: Wed, 16 Sep 2026 10:05:52 -0500 Subject: [PATCH 1/5] Enable fix from PR #11963 for gfx950 --- .../hipblaslt/tensilelite/client/src/DataInitialization.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hipblaslt/tensilelite/client/src/DataInitialization.cpp b/projects/hipblaslt/tensilelite/client/src/DataInitialization.cpp index 33c71ae24f3e..4746d9ec2b64 100644 --- a/projects/hipblaslt/tensilelite/client/src/DataInitialization.cpp +++ b/projects/hipblaslt/tensilelite/client/src/DataInitialization.cpp @@ -2137,7 +2137,7 @@ namespace TensileLite // gfx1250 only: its scale layout comes from problem geometry, so // the descriptors fully determine the generated bytes. GFX950 // preswizzles per solution matrix instruction and must not reuse. - bool const reuseSupported = (m_mxScaleLayout == MXScaleLayout::GFX1250); + bool const reuseSupported = (m_mxScaleLayout == MXScaleLayout::GFX1250 || m_mxScaleLayout == MXScaleLayout::GFX950); bool const descriptorsUnchanged = pristineData.initDescriptor.size() == 1 && pristineScale.initDescriptor.size() == 1 @@ -2871,7 +2871,7 @@ namespace TensileLite ptr = p.gpuInput.valid.get(); } else if(auto mxIt = m_mxSwizzledDescriptor.find(i); - m_mxScaleLayout == MXScaleLayout::GFX1250 + (m_mxScaleLayout == MXScaleLayout::GFX1250 || m_mxScaleLayout == MXScaleLayout::GFX950) && mxIt != m_mxSwizzledDescriptor.end() && mxIt->second == desc) { // Already swizzled, and initOneMXSide kept the source bytes From 4d12f307903bcd5b3fed123583e3856c0e71a1d0 Mon Sep 17 00:00:00 2001 From: amd-ysanchez Date: Thu, 17 Sep 2026 06:09:26 -0500 Subject: [PATCH 2/5] Fix gfx950 re-use and add tests for gfx950 and gfx1250 --- .../client/src/DataInitialization.cpp | 59 +++++++-- .../tensilelite/tests/DataInit_test.cpp | 118 ++++++++++++++++++ 2 files changed, 167 insertions(+), 10 deletions(-) diff --git a/projects/hipblaslt/tensilelite/client/src/DataInitialization.cpp b/projects/hipblaslt/tensilelite/client/src/DataInitialization.cpp index 4746d9ec2b64..2f808acd3b4f 100644 --- a/projects/hipblaslt/tensilelite/client/src/DataInitialization.cpp +++ b/projects/hipblaslt/tensilelite/client/src/DataInitialization.cpp @@ -2133,19 +2133,45 @@ namespace TensileLite // prepareGPUInputs() runs once per solution, so without this // the generator re-runs for every solution on identical input. // Same idea as the non-MX initDescriptor check below. - // - // gfx1250 only: its scale layout comes from problem geometry, so - // the descriptors fully determine the generated bytes. GFX950 - // preswizzles per solution matrix instruction and must not reuse. - bool const reuseSupported = (m_mxScaleLayout == MXScaleLayout::GFX1250 || m_mxScaleLayout == MXScaleLayout::GFX950); + bool const descriptorsUnchanged = pristineData.initDescriptor.size() == 1 && pristineScale.initDescriptor.size() == 1 && pristineData.initDescriptor[0] == dataDesc && pristineScale.initDescriptor[0] == scaleDesc; - if(reuseSupported && descriptorsUnchanged) + // gfx1250 swizzles depend only on the tensor descriptors, so an + // unchanged descriptor means both the canonical scale and the cached + // swizzled bytes are still valid. gfx950 is excluded because its + // preswizzle choice also depends on the current solution's MI shape. + if(m_mxScaleLayout == MXScaleLayout::GFX1250 && descriptorsUnchanged) + { return; + } + + auto const cachedSwizzled = m_mxSwizzledDescriptor.find(scaleTensorEnum); + bool const hasMatchingSwizzledScale + = swizzleLayout != MXScaleLayout::None + && cachedSwizzled != m_mxSwizzledDescriptor.end() + && cachedSwizzled->second == scaleDesc; + + // A canonical-only upload means gpuInput.valid no longer holds a + // trustworthy swizzled scale for this tensor, so clear the cache tag + // before any later reuse decision observes stale metadata. + if(swizzleLayout == MXScaleLayout::None) + m_mxSwizzledDescriptor.erase(scaleTensorEnum); + + // gfx950 can skip regeneration only when the descriptor is unchanged + // and any required preswizzled copy is already known to match this + // scale descriptor. When that cache hit happens, remember that + // gpuInput.valid already contains the solution-ready layout. + if(m_mxScaleLayout == MXScaleLayout::GFX950 && descriptorsUnchanged + && (swizzleLayout == MXScaleLayout::None || hasMatchingSwizzledScale)) + { + if(hasMatchingSwizzledScale) + *preswizzledFlag = true; + return; + } if(pristineData.initDescriptor.size() == 1) pristineData.initDescriptor[0] = dataDesc; @@ -2272,6 +2298,10 @@ namespace TensileLite gpuScaleBuf.data(), gpuScaleBytes, hipMemcpyHostToDevice)); + // Record that gpuInput.valid now contains a verified swizzled + // scale for this exact descriptor, so later solution passes can + // reuse it without re-permuting host data. + m_mxSwizzledDescriptor[scaleTensorEnum] = scaleDesc; *preswizzledFlag = true; } }; @@ -2860,6 +2890,10 @@ namespace TensileLite kernelMxScaleFormat = m_currentSolution->problemType.mxScaleFormat; if(kernelMxScaleFormat == 0) { + // The selected kernel explicitly wants canonical scales, so + // uploading cpuInput.valid must also invalidate any prior + // "gpuInput.valid is swizzled" cache state for this tensor. + m_mxSwizzledDescriptor.erase(i); ptr = copyInputBuffers(desc, p.gpuInput.valid.get(), p.cpuInput.valid.get(), @@ -2874,13 +2908,18 @@ namespace TensileLite (m_mxScaleLayout == MXScaleLayout::GFX1250 || m_mxScaleLayout == MXScaleLayout::GFX950) && mxIt != m_mxSwizzledDescriptor.end() && mxIt->second == desc) { - // Already swizzled, and initOneMXSide kept the source bytes - // unchanged, so skip the host-side permute. + // The cache says gpuInput.valid already holds the swizzled + // bytes for this descriptor, so reuse that device buffer as-is + // instead of rebuilding the layout from cpuInput.valid. ptr = p.gpuInput.valid.get(); } - else + + if(ptr == nullptr) { - m_mxSwizzledDescriptor[i] = desc; + // We are about to rebuild a swizzled copy from canonical host + // data, so any previous descriptor tag is stale until the new + // transform completes and is written back to gpuInput.valid. + m_mxSwizzledDescriptor.erase(i); // Fallback K-dimension swizzle for the case where // initializeMXData did NOT pre-produce a swizzled diff --git a/projects/hipblaslt/tensilelite/tests/DataInit_test.cpp b/projects/hipblaslt/tensilelite/tests/DataInit_test.cpp index 2677cb1c0405..5b413fc768a7 100644 --- a/projects/hipblaslt/tensilelite/tests/DataInit_test.cpp +++ b/projects/hipblaslt/tensilelite/tests/DataInit_test.cpp @@ -3,9 +3,14 @@ #include #include "DataInitialization.hpp" // isMXTensor / Problem +#include "ClientProblemFactory.hpp" #include "DataInitializationHelpers.hpp" // detail::* (MX-only, internally guarded) +#include "ProgramOptions.hpp" #include +#include #include +#include +#include #include #include #include @@ -30,7 +35,10 @@ #endif using TensileLite::ContractionProblemGemm; +using TensileLite::ContractionSolution; using TensileLite::DataTypeInfo; +using TensileLite::KernelLanguage; +using TensileLite::PerformanceMetric; using TensileLite::TensorDescriptor; using TensileLite::Client::isMXProblem; using TensileLite::Client::isMXTensor; @@ -41,6 +49,12 @@ namespace dt = TensileLite::Client::detail; #endif namespace { + template + void setOption(TensileLite::Client::po::variables_map& args, std::string const& name, T value) + { + args[name].value() = std::move(value); + } + // ----------------------------------------------------------------------- // Helper: build a ContractionProblemGemm with the requested A/B dtypes. // Mirrors tests/MXScalePadding_test.cpp::makeMXProblem so the geometry @@ -77,6 +91,63 @@ namespace if(mxBlockB > 0) problem.setMXScaleB(rocisa::DataType::E8, mxBlockB); return problem; } + +#if HIPBLASLT_ENABLE_MXDATAGENERATOR + TensileLite::Client::po::variables_map makeMinimalClientArgs(int mxScaleFormat) + { + using namespace TensileLite::Client; + + po::variables_map args; + setOption(args, "problem-identifier", std::string("Contraction_l_Ailk_Bljk_Cijk_Dijk")); + setOption(args, "problem-size", std::vector>{{128, 128, 1, 512}}); + setOption(args, "strided-batched", true); + setOption(args, "batch-mode", 0); + setOption(args, "grouped-gemm", false); + setOption(args, "sparse", 0); + setOption(args, "high-precision-accumulate", false); + setOption(args, "kernel-language", KernelLanguage::Any); + setOption(args, "performance-metric", PerformanceMetric::DeviceEfficiency); + setOption(args, "deterministic-mode", false); + setOption(args, "c-equal-d", false); + setOption(args, "type", rocisa::DataType::Float8); + setOption(args, "a-type", rocisa::DataType::Float8); + setOption(args, "b-type", rocisa::DataType::Float8); + setOption(args, "c-type", rocisa::DataType::Float); + setOption(args, "d-type", rocisa::DataType::Float); + setOption(args, "alpha-type", rocisa::DataType::Float); + setOption(args, "beta-type", rocisa::DataType::Float); + setOption(args, "compute-input-type-A", rocisa::DataType::Float8); + setOption(args, "compute-input-type-B", rocisa::DataType::Float8); + setOption(args, "f32-xdl-math-op", rocisa::DataType::Float); + setOption(args, "activation-compute-type", rocisa::DataType::Float); + setOption(args, "mx-a-block", 32); + setOption(args, "mx-b-block", 32); + setOption(args, "mx-a-type", rocisa::DataType::E8); + setOption(args, "mx-b-type", rocisa::DataType::E8); + setOption(args, "mx-scale-format", mxScaleFormat); + setOption(args, "fused-gemm-a2a", false); + setOption(args, "metadata-layout", 0); + setOption(args, "a-ops", TensileLite::TensorOps{}); + setOption(args, "b-ops", TensileLite::TensorOps{}); + setOption(args, "c-ops", TensileLite::TensorOps{}); + setOption(args, "d-ops", TensileLite::TensorOps{}); + setOption(args, "use-gradient", false); + setOption(args, "output-amaxD", false); + setOption(args, "use-scaleAB", std::string()); + setOption(args, "use-scaleCD", false); + setOption(args, "use-scaleAlphaVec", 0); + setOption(args, "device-idx", 0); + setOption(args, "num-elements-to-validate", 0); + setOption(args, "pristine-on-gpu", true); + setOption(args, "prune-mode", TensileLite::Client::PruneSparseMode::PruneRandom); + setOption(args, "rotating-buffer-size", 0); + setOption(args, "rotating-buffer-mode", 0); + setOption(args, "bounds-check", TensileLite::Client::BoundsCheckMode::Disable); + return args; + } + +#endif + } // namespace // ============================================================================= @@ -246,6 +317,53 @@ TEST(InitializeMXDataForFP4OrFP8_BatchStrideFormula, BFloat8_OneBytePerElement) EXPECT_EQ(bytes, kStrideElems); } +#if HIPBLASLT_ENABLE_MXDATAGENERATOR +TEST(DataInitializationReferenceRecompute, Gfx950RequiresPerSolutionRecomputeForSwizzledMX) +{ + hipDeviceProp_t prop{}; + ASSERT_EQ(hipGetDeviceProperties(&prop, 0), hipSuccess); + if(std::string(prop.gcnArchName).find("gfx950") == std::string::npos) + GTEST_SKIP() << "gfx950-specific recompute behavior is only observable on gfx950"; + + auto problem = makeProblem(rocisa::DataType::Float8, + rocisa::DataType::Float8, + /*mxBlockA=*/32, + /*mxBlockB=*/32); + auto args = makeMinimalClientArgs(/*mxScaleFormat=*/1); + TensileLite::Client::ClientProblemFactory factory(args); + TensileLite::Client::DataInitialization dataInit(args, factory); + + ContractionSolution noSwizzleSolution; + noSwizzleSolution.problemType.mxScaleFormat = 0; + ContractionSolution swizzleSolution; + swizzleSolution.problemType.mxScaleFormat = 1; + + EXPECT_FALSE(dataInit.referenceNeedsPerSolutionRecompute(problem, &noSwizzleSolution)); + EXPECT_TRUE(dataInit.referenceNeedsPerSolutionRecompute(problem, &swizzleSolution)); +} + +TEST(DataInitializationReferenceRecompute, Gfx1250ReusesPreswizzledMXAcrossSolutions) +{ + hipDeviceProp_t prop{}; + ASSERT_EQ(hipGetDeviceProperties(&prop, 0), hipSuccess); + if(std::string(prop.gcnArchName).find("gfx1250") == std::string::npos) + GTEST_SKIP() << "gfx1250-specific reuse behavior is only observable on gfx1250"; + + auto problem = makeProblem(rocisa::DataType::Float8, + rocisa::DataType::Float8, + /*mxBlockA=*/32, + /*mxBlockB=*/32); + auto args = makeMinimalClientArgs(/*mxScaleFormat=*/1); + TensileLite::Client::ClientProblemFactory factory(args); + TensileLite::Client::DataInitialization dataInit(args, factory); + + ContractionSolution swizzleSolution; + swizzleSolution.problemType.mxScaleFormat = 1; + + EXPECT_FALSE(dataInit.referenceNeedsPerSolutionRecompute(problem, &swizzleSolution)); +} +#endif + // ============================================================================= // Section 4 — direct calls into TensileLite::Client::detail (MX builds only) // ============================================================================= From 229f50e0ea72cedb8977046eb1e9d0c2ff404bcc Mon Sep 17 00:00:00 2001 From: amd-ysanchez Date: Thu, 17 Sep 2026 06:10:05 -0500 Subject: [PATCH 3/5] Fix gfx950 re-use and add tests for gfx950 and gfx1250 --- .../hipblaslt/tensilelite/client/src/DataInitialization.cpp | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/projects/hipblaslt/tensilelite/client/src/DataInitialization.cpp b/projects/hipblaslt/tensilelite/client/src/DataInitialization.cpp index 2f808acd3b4f..af119671d98f 100644 --- a/projects/hipblaslt/tensilelite/client/src/DataInitialization.cpp +++ b/projects/hipblaslt/tensilelite/client/src/DataInitialization.cpp @@ -2908,9 +2908,8 @@ namespace TensileLite (m_mxScaleLayout == MXScaleLayout::GFX1250 || m_mxScaleLayout == MXScaleLayout::GFX950) && mxIt != m_mxSwizzledDescriptor.end() && mxIt->second == desc) { - // The cache says gpuInput.valid already holds the swizzled - // bytes for this descriptor, so reuse that device buffer as-is - // instead of rebuilding the layout from cpuInput.valid. + // Already swizzled, and initOneMXSide kept the source bytes + // unchanged, so skip the host-side permute. ptr = p.gpuInput.valid.get(); } From 15b80f999d267ff743979210ec1b63eea4d9798a Mon Sep 17 00:00:00 2001 From: amd-ysanchez Date: Thu, 17 Sep 2026 09:36:46 -0500 Subject: [PATCH 4/5] Fix failing CI test --- .../tensilelite/tests/DataInit_test.cpp | 17 +++++++++++++++-- 1 file changed, 15 insertions(+), 2 deletions(-) diff --git a/projects/hipblaslt/tensilelite/tests/DataInit_test.cpp b/projects/hipblaslt/tensilelite/tests/DataInit_test.cpp index 5b413fc768a7..9f60c011e3a1 100644 --- a/projects/hipblaslt/tensilelite/tests/DataInit_test.cpp +++ b/projects/hipblaslt/tensilelite/tests/DataInit_test.cpp @@ -120,8 +120,21 @@ namespace setOption(args, "compute-input-type-B", rocisa::DataType::Float8); setOption(args, "f32-xdl-math-op", rocisa::DataType::Float); setOption(args, "activation-compute-type", rocisa::DataType::Float); - setOption(args, "mx-a-block", 32); - setOption(args, "mx-b-block", 32); + // Deliberately 0 here (not 32): these only control whether + // ClientProblemFactory calls setMXScaleA/B on the *internal dummy* + // problem it builds for sizing purposes (see ClientProblemFactory.cpp), + // which is unrelated to the `problem` object this test passes directly + // into referenceNeedsPerSolutionRecompute(). When non-zero here, the + // DataInitialization ctor's storage-geometry computation for the + // swizzled MX-scale tensor indexes problem.freeIndicesA()[0] / + // freeIndicesB()[0] on that dummy problem without a bounds check; for + // the minimal synthetic problem-size/identifier used here those + // indices can be empty, which segfaults (only reachable on real + // hardware, so it stays hidden behind GTEST_SKIP() on any other arch). + // m_mxScaleFormat (checked by referenceNeedsPerSolutionRecompute) is + // read directly from "mx-scale-format" below, independent of these. + setOption(args, "mx-a-block", 0); + setOption(args, "mx-b-block", 0); setOption(args, "mx-a-type", rocisa::DataType::E8); setOption(args, "mx-b-type", rocisa::DataType::E8); setOption(args, "mx-scale-format", mxScaleFormat); From ba23120d8b31e60b5286f31aa664293a2b8a0ca0 Mon Sep 17 00:00:00 2001 From: amd-ysanchez Date: Thu, 17 Sep 2026 10:28:06 -0500 Subject: [PATCH 5/5] Finally fix CI segfault --- .../tensilelite/tests/DataInit_test.cpp | 18 ++++++++++++++++++ 1 file changed, 18 insertions(+) diff --git a/projects/hipblaslt/tensilelite/tests/DataInit_test.cpp b/projects/hipblaslt/tensilelite/tests/DataInit_test.cpp index 9f60c011e3a1..53f00f5dc121 100644 --- a/projects/hipblaslt/tensilelite/tests/DataInit_test.cpp +++ b/projects/hipblaslt/tensilelite/tests/DataInit_test.cpp @@ -148,6 +148,24 @@ namespace setOption(args, "output-amaxD", false); setOption(args, "use-scaleAB", std::string()); setOption(args, "use-scaleCD", false); + // ClientProblemFactory::m_useBias / m_biasSrc are plain `int` members + // with NO default member initializer (see ClientProblemFactory.hpp); + // they're only assigned when args.count("use-bias") / + // args.count("bias-source") is true. Without these two options, + // m_useBias holds whatever garbage was left on the stack/heap when + // ClientProblemFactory was constructed. If that garbage is non-zero, + // ContractionProblemGemm::setBias() takes a branch that indexes + // m_tensors[m_biasSrc].sizes()[batchIdx] with batchIdx left at its + // default of 2 for this non-batched dummy problem (m_batchIndices is + // empty) -- an out-of-bounds read that segfaults. This is why the + // crash only showed up when the test ran in isolation (a fresh + // process, as CI's per-test ctest invocation does): the leftover + // stack/heap contents differ from running the whole 733-test binary + // in one process, where prior tests happened to leave zeroed memory + // behind and masked the bug. Setting these explicitly makes the + // behavior deterministic regardless of process/memory history. + setOption(args, "use-bias", 0); + setOption(args, "bias-source", static_cast(ContractionProblemGemm::TENSOR::D)); setOption(args, "use-scaleAlphaVec", 0); setOption(args, "device-idx", 0); setOption(args, "num-elements-to-validate", 0);