Skip to content
Merged
Show file tree
Hide file tree
Changes from 1 commit
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Prev Previous commit
hipblaslt: build v0 kernels for gfx1250 and pick them by ASIC revision
A v0 part reports gfx1250-strict only while HSA_DISABLE_GFX12_STRICT=0.
Unset or =1 it reports plain gfx1250 and loads only gfx1250 (0x49) code
objects, so a gfx1250 package hands it v1 kernels, some of which use
instructions v0 does not have: the fp4 32x16 WMMA kernel returns NaN on
v0, and v0 has no TDM multicast.

A build that requests gfx1250 now also builds gfx1250v0: gfx1250-strict's
logic and capability overrides, assembled and compiled for the gfx1250
target and installed as library/gfx1250v0/ under gfx1250 file names. Its
kernels match gfx1250-strict's instruction for instruction; only the
target differs. gfx1250v0 is a build alias, not an architecture: it never
appears in GPU_TARGETS and is rejected when named directly.
TensileCreateLibrary adds it only under --gfx1250v0, which the
HIPBLASLT_BUILD_GFX1250V0 CMake option (default ON, shared with
hipSPARSELt) passes, so dropping v0 from the gfx1250 target later is that
one switch. The alias joins before the ISA partition, so it builds in a
process and scratch directory of its own, and its assembler and helper
compile get strict's target features, so v1-only instructions are still
rejected at build time.

At runtime hipBLASLt and hipSPARSELt choose the Tensile subtree by the
reported name first: gfx1250-strict keeps loading library/gfx1250-strict/.
A device reporting gfx1250 with hipDeviceProp_t::asicRevision 0 loads
library/gfx1250v0/, and fails to initialize if it is missing rather than
fall back to v1 kernels that give wrong results. Every other device is
unchanged. ExtOp and matrix-transform still load from
library/<reported name>/.

validate_library_layout.py accepts the gfx1250v0 subtree, which carries
gfx1250 file names and no ExtOp or Transform files, and row_uniformity
reads the subtree the runtime loads.

test_MatrixInstructionConversion updated the shared defaultSolution in
place, which leaked into whichever test xdist ran next on that worker;
it now works on a copy.

The gfx1250-strict copy of spmm_f16_sia4.yaml picks up #12647, which
moved the non-TDM group of the gfx1250 copy to ScheduleIterAlg 3; the two
now differ only in their skip mark and Architecture.

Co-authored-by: Cursor <cursoragent@cursor.com>
  • Loading branch information
geotseng-amd and cursoragent committed Sep 30, 2026
commit 6fdee38bedbf9e56c3554155c761d9dc55385a76
1 change: 1 addition & 0 deletions projects/hipblaslt/clients/tests/src/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -7,6 +7,7 @@ target_sources(hipblaslt-test
${CMAKE_CURRENT_SOURCE_DIR}/hipblaslt_test.cpp
${CMAKE_CURRENT_SOURCE_DIR}/caching_library_gtest.cpp
${CMAKE_CURRENT_SOURCE_DIR}/secure_env_gtest.cpp
${CMAKE_CURRENT_SOURCE_DIR}/arch_revision_gtest.cpp
${CMAKE_CURRENT_SOURCE_DIR}/efficiency_monitor_gtest.cpp
${CMAKE_CURRENT_SOURCE_DIR}/matmul_gtest.cpp
${CMAKE_CURRENT_SOURCE_DIR}/auxiliary_gtest.cpp
Expand Down
48 changes: 48 additions & 0 deletions projects/hipblaslt/clients/tests/src/arch_revision_gtest.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,48 @@
// Copyright Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT

// Host-only smoke test for the ASIC-revision -> library-arch mapping, a pure
// function with no HIP dependency (see rocblaslt_arch_revision.hpp). Included by
// relative path so the white-box test stays self-contained without adding the
// internal rocblaslt include dir to the whole test target.

#include <gtest/gtest.h>

#include "../../../library/src/amd_detail/rocblaslt/src/include/rocblaslt_arch_revision.hpp"

namespace
{
TEST(arch_revision_smoke, gfx1250_revision0_is_the_v0_subtree)
{
// An A0 part in non-strict mode reports gfx1250 and loads its own tree.
EXPECT_EQ(rocblaslt_revisioned_arch_name("gfx1250", 0), "gfx1250v0");
}

TEST(arch_revision_smoke, gfx1250_strict_keeps_its_reported_name)
{
// An A0 part in strict mode reports gfx1250-strict (asicRevision is still 0)
// and keeps loading library/gfx1250-strict/.
EXPECT_EQ(rocblaslt_revisioned_arch_name("gfx1250-strict", 0), "gfx1250-strict");
}

TEST(arch_revision_smoke, gfx1250_nonzero_revision_keeps_the_base_name)
{
// Assumes B0 reports a non-zero revision; it maps to the plain gfx1250 tree.
EXPECT_EQ(rocblaslt_revisioned_arch_name("gfx1250", 1), "gfx1250");
EXPECT_EQ(rocblaslt_revisioned_arch_name("gfx1250", 2), "gfx1250");
}

TEST(arch_revision_smoke, gfx1250_unknown_revision_keeps_the_base_name)
{
// Only a revision known to be A0 selects the A0 tree.
EXPECT_EQ(rocblaslt_revisioned_arch_name("gfx1250", -1), "gfx1250");
}

TEST(arch_revision_smoke, other_arches_are_unaffected_by_revision)
{
EXPECT_EQ(rocblaslt_revisioned_arch_name("gfx942", 0), "gfx942");
EXPECT_EQ(rocblaslt_revisioned_arch_name("gfx950", 0), "gfx950");
EXPECT_EQ(rocblaslt_revisioned_arch_name("gfx1250v0", 0), "gfx1250v0");
EXPECT_EQ(rocblaslt_revisioned_arch_name("gfx1250-strict", 1), "gfx1250-strict");
}
}
13 changes: 12 additions & 1 deletion projects/hipblaslt/clients/tests/src/row_uniformity_gtest.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -37,6 +37,8 @@
#include <origami/hardware.hpp>
#include <origami/streamk.hpp>

#include "../../../library/src/amd_detail/rocblaslt/src/include/rocblaslt_arch_revision.hpp"

#ifdef _WIN32
#ifndef WIN32_LEAN_AND_MEAN
#define WIN32_LEAN_AND_MEAN
Expand Down Expand Up @@ -169,14 +171,23 @@ namespace
// that never appear in the library filename.
const std::string arch = gpuArchName();
const std::string processor = arch.substr(0, arch.find(':'));
// The subtree the runtime loads, which for an A0 part reporting
// gfx1250 is not the one named for processor.
hipDeviceProp_t props{};
int device = 0;
const int rev = hipGetDevice(&device) == hipSuccess
&& hipGetDeviceProperties(&props, device) == hipSuccess
? props.asicRevision
: -1;
const std::string libArch = rocblaslt_revisioned_arch_name(processor, rev);

std::string tried;
for(const auto& root : libraryRootCandidates())
{
// Always the logical single-extension name: the loader resolves
// the shipped ".dat.zlib" by appending the suffix itself.
const std::filesystem::path logical
= root / processor / ("TensileLibrary_lazy_" + processor + ".dat");
= root / libArch / ("TensileLibrary_lazy_" + processor + ".dat");
tried += (tried.empty() ? "" : ", ") + logical.string();

if(!std::filesystem::exists(logical)
Expand Down
15 changes: 14 additions & 1 deletion projects/hipblaslt/cmake/hipblaslt_codegen.cmake
Original file line number Diff line number Diff line change
Expand Up @@ -130,6 +130,10 @@ function(hipblaslt_detect_sanitizer_runtime out_options out_lib_dirs)
set(${out_lib_dirs} "${_lib_dirs}" PARENT_SCOPE)
endfunction()

# When ON, a build that requests gfx1250 also produces library/gfx1250v0/
# (gfx1250-strict's kernels built for gfx1250). Never add gfx1250v0 to GPU_TARGETS.
option(HIPBLASLT_BUILD_GFX1250V0 "Build library/gfx1250v0/ alongside gfx1250 for A0 parts reporting gfx1250." ON)

function(create_device_library)
set(_opts HOST_ASAN HOST_TSAN)
set(_one
Expand Down Expand Up @@ -293,11 +297,20 @@ function(create_device_library)
if(NOT _cdl_ASM_COMMENTS)
list(APPEND _opts_list "--disable-asm-comments")
endif()
set(_logic_arches ${_cdl_ARCHES})
if(HIPBLASLT_BUILD_GFX1250V0)
list(APPEND _opts_list "--gfx1250v0")
# gfx1250v0 is built from gfx1250-strict's logic, so validate that too.
if("gfx1250" IN_LIST _cdl_ARCHES AND NOT "gfx1250-strict" IN_LIST _cdl_ARCHES)
list(APPEND _logic_arches "gfx1250-strict")
endif()
endif()
list(JOIN _logic_arches "$<SEMICOLON>" _logic_arches_semi)

set(_tensile_logic_args
"${_cdl_LOGIC_PATH}"
--architecture
"${_arches_semi}"
"${_logic_arches_semi}"
--use-bundled-known-bugs
--check-all
)
Expand Down
Original file line number Diff line number Diff line change
@@ -0,0 +1,18 @@
// Copyright Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT

#pragma once

#include <string>

// The GEMM library subtree a device loads. With HSA_DISABLE_GFX12_STRICT unset/1 an
// A0 part reports the same name as B0 (gfx1250), so only hipDeviceProp_t::asicRevision
// tells them apart (A0 -> 0): an A0 part loads library/gfx1250v0/ only (no fallback to
// gfx1250 -- B0 kernels are not valid on A0). gfx1250-strict and every other reported
// name are unchanged. Dependency-free so it can be unit-tested GPU-free.
inline std::string rocblaslt_revisioned_arch_name(const std::string& baseArch, int asicRevision)
{
if(baseArch == "gfx1250" && asicRevision == 0)
return "gfx1250v0";
return baseArch;
}
Original file line number Diff line number Diff line change
Expand Up @@ -36,6 +36,7 @@
#include "Debug.hpp"
#include "include/check_numerics_matrix.hpp"
#include "rocblaslt-types.h"
#include "rocblaslt_arch_revision.hpp"
#include "rocblaslt_fused_a2a_peers.hpp"
#include "rocblaslt_mat_utils.hpp"
#include "rocblaslt_secure_env.hpp"
Expand Down Expand Up @@ -3004,12 +3005,20 @@ namespace
// path. Only use the subdir if a Tensile mapping file is actually present
// there; otherwise the directory may have been created by ExtOp/Transform
// installs without a corresponding Tensile library (multi-arch non-TheRock
// builds). Both the subdir and the mapping filenames carry `processor`,
// the name the runtime reports for this device -- which is the compiler
// target the kernels in it were built for, including a silicon-revision
// variant such as gfx1250-strict.
// builds). The mapping filenames carry `processor`, the name the runtime
// reports for this device -- which is the compiler target the kernels in
// it were built for, including gfx1250-strict. The subdir is revisioned:
// an A0 part reporting gfx1250 loads library/gfx1250v0/, with no fallback.
{
auto processor_path = path / processor;
int asicRevision = -1;
if(processor == "gfx1250")
{
hipDeviceProp_t deviceProperties;
HIP_CHECK_EXC(hipGetDeviceProperties(&deviceProperties, deviceId));
asicRevision = deviceProperties.asicRevision;
}
const auto libArch = rocblaslt_revisioned_arch_name(processor, asicRevision);
auto processor_path = path / libArch;
auto mapping_msgpack = processor_path / ("TensileLibrary_lazy_" + processor + ".dat");
auto mapping_msgpack_gz = processor_path / ("TensileLibrary_lazy_" + processor + ".dat.zlib");
auto mapping_yaml = processor_path / ("TensileLibrary_lazy_" + processor + ".yaml");
Expand All @@ -3018,6 +3027,23 @@ namespace
{
path = std::move(processor_path);
}
else if(libArch != processor)
{
auto master = processor_path / ("TensileLibrary_" + processor);
if(!std::filesystem::exists(master.string() + ".dat")
&& !std::filesystem::exists(master.string() + ".dat.zlib")
&& !std::filesystem::exists(master.string() + ".yaml"))
std::cerr << "\nrocblaslt error: " << processor
<< " device with asicRevision 0 (A0) needs the " << libArch
<< " GEMM library, but " << processor_path
<< " holds no Tensile library. Build with the "
<< libArch << " subtree (the default for -a " << processor
<< "), or run with HSA_DISABLE_GFX12_STRICT=0 and a "
<< processor << "-strict build. Not falling back to "
<< (path / processor)
<< ": its kernels give wrong results on A0." << std::endl;
path = std::move(processor_path);
}
}

if(get_logger_layer_mode() & rocblaslt_layer_mode_log_info)
Expand Down
69 changes: 67 additions & 2 deletions projects/hipblaslt/tensilelite/Tensile/Common/Architectures.py
Original file line number Diff line number Diff line change
Expand Up @@ -72,8 +72,69 @@
# Spelled as clang and ROCr both spell it. Shares gfx1250's ISA, so
# SUPPORTED_ISA cannot name it; this entry is where `all` picks it up.
"gfx1250-strict": "gfx1250-strict",
"gfx1250v0": "gfx1250v0",
}

# Names that build another architecture's codegen for an existing target,
# shipped under that target's file names in a subtree of their own. gfx1250v0
# is gfx1250-strict's logic and capabilities built for gfx1250, for A0 parts
# the runtime reports as gfx1250. No compiler knows the alias, so it stays out
# of `all`, archMacroNames and GPU_TARGETS, and is only built on request
# (TensileCreateLibrary --gfx1250v0).
ARCH_BUILD_ALIASES = {
"gfx1250v0": {
# Both the compiler target and the name the library files carry.
"target": "gfx1250",
"tuningArch": "gfx1250-strict",
# gfx1250-strict's features on a gfx1250 compile or assembly.
"deviceTargetFeatures": [
"+needs-aligned-2addr-lds",
"-wmma-f4-insts",
"-block16-cvt-scale-insts",
],
},
}


def _buildAlias(spec: str) -> Optional[dict]:
return ARCH_BUILD_ALIASES.get(baseArchName(spec))


def compilerTargetOf(spec: str) -> str:
"""The compiler target a spec is built for, and the name its library files
carry; the spec itself unless aliased."""
alias = _buildAlias(spec)
return alias["target"] if alias else spec


def tuningArchOf(spec: str) -> str:
"""The architecture whose logic and capability overrides a spec builds; the spec itself unless aliased."""
alias = _buildAlias(spec)
return alias["tuningArch"] if alias else spec


def deviceTargetFeaturesOf(spec: str) -> List[str]:
"""Target features added when compiling or assembling a spec; none unless aliased."""
alias = _buildAlias(spec)
return list(alias["deviceTargetFeatures"]) if alias else []


def withArchBuildAliases(archs: List[str]) -> List[str]:
"""``archs`` plus the alias of every target it requests.

gfx1250 gains gfx1250v0; gfx1250-strict names a different target and does
not expand.

Call before ``isaCollisionFreeGroups``: the alias shares its ISA with what
it expands from and has to reach the partitioner that splits them.
"""
named = {baseArchName(a) for a in archs}
return archs + [
alias
for alias, fields in ARCH_BUILD_ALIASES.items()
if fields["target"] in named and alias not in named
]

gfxVariantMap = {
"gfx906": ["gfx906:xnack+", "gfx906:xnack-"],
"gfx908": ["gfx908:xnack+", "gfx908:xnack-"],
Expand Down Expand Up @@ -206,7 +267,7 @@ def supportedSteppings() -> List[str]:
return [
name
for name in architectureMap
if steppingArchOf(name) in covered
if steppingArchOf(name) in covered and name not in ARCH_BUILD_ALIASES
]


Expand Down Expand Up @@ -268,7 +329,11 @@ def archMacroNames(isa: IsaVersion) -> List[str]:
if isa is None:
return []
names = sorted(
{baseArchName(name) for name in architectureMap if gfxToIsa(name) == isa}
{
baseArchName(name)
for name in architectureMap
if gfxToIsa(name) == isa and name not in ARCH_BUILD_ALIASES
}
)
return ["__" + name.replace("-", "_") + "__" for name in names]

Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -26,7 +26,7 @@

from typing import List, Dict

from .Architectures import ARCH_CAP_OVERRIDES, baseArchName, gfxToIsa
from .Architectures import ARCH_CAP_OVERRIDES, baseArchName, gfxToIsa, tuningArchOf
from .Types import IsaVersion, IsaInfo


Expand Down Expand Up @@ -62,7 +62,7 @@ def applyArchCapOverrides(isaInfoMap: Dict[IsaVersion, IsaInfo], archNames: List
# Keyed on the bare name: --gpu-targets forwards a requested spec verbatim,
# predicates and all, and a lookup that missed gfx1250-strict[cu=64] would
# silently build it with gfx1250's capabilities.
overrides = ARCH_CAP_OVERRIDES.get(baseArchName(name))
overrides = ARCH_CAP_OVERRIDES.get(tuningArchOf(baseArchName(name)))
if not overrides:
continue
isa = gfxToIsa(name)
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -849,6 +849,7 @@ def _assertGlobalParametersAreValid(config, ignoreKeys):
"Experimental", # --experimental logic-dir toggle in ParseArguments
"EnableGemmA2AFusion", # --enable-gemm-a2a-fusion toggle in ParseArguments
"GenSolTable", # --gen-sol-table toggle in ParseArguments
"BuildGfx1250v0", # --gfx1250v0 toggle in ParseArguments
# Keys with a sanctioned opt-out from the strict gate:
# - Live but read via DebugConfig (makeDebugConfig in
# Tensile/Common/Types.py) directly from the raw config dict
Expand Down
17 changes: 13 additions & 4 deletions projects/hipblaslt/tensilelite/Tensile/LibraryIO.py
Original file line number Diff line number Diff line change
Expand Up @@ -549,7 +549,8 @@ def parseLibraryLogicFile(
printSolutionRejectionReason: bool,
printIndexAssignmentInfo: bool,
isaInfoMap: Dict[str, IsaInfo],
lazyLibraryLoading: bool
lazyLibraryLoading: bool,
archRenames: Optional[Dict[str, str]] = None,
):
"""Wrapper function to read and parse a library logic file."""
return parseLibraryLogicData(
Expand All @@ -560,7 +561,8 @@ def parseLibraryLogicFile(
printSolutionRejectionReason,
printIndexAssignmentInfo,
isaInfoMap,
lazyLibraryLoading
lazyLibraryLoading,
archRenames,
)


Expand Down Expand Up @@ -672,9 +674,14 @@ def parseLibraryLogicData(
printSolutionRejectionReason: bool,
printIndexAssignmentInfo: bool,
isaInfoMap: Dict[str, IsaInfo],
lazyLibraryLoading: bool
lazyLibraryLoading: bool,
archRenames: Optional[Dict[str, str]] = None,
):
"""Parses the data of a library logic file."""
"""Parses the data of a library logic file.

``archRenames`` maps a declared ArchitectureName to the name the library is
keyed and its files are named by, for a build alias (see ARCH_BUILD_ALIASES).
"""
# Reset the type mismatch collector at the start to capture all type
# mismatches from both ProblemType and Solution constructors
resetTypeMismatchCollector()
Expand All @@ -684,6 +691,8 @@ def parseLibraryLogicData(
elif isinstance(data, dict):
prepareLibraryLogicDict(data)

if archRenames:
data["ArchitectureName"] = archRenames.get(data["ArchitectureName"], data["ArchitectureName"])
if "CUCount" not in data:
data["CUCount"] = None
if 'MacDataTypeA' not in data["ProblemType"]: #it will either be set as d['MacDataType'] or a specified input
Expand Down
8 changes: 5 additions & 3 deletions projects/hipblaslt/tensilelite/Tensile/Tensile.py
Original file line number Diff line number Diff line change
Expand Up @@ -42,7 +42,7 @@
from Tensile.Common import print1, printExit, printWarning, ensurePath, HR, isRhel8, \
LIBRARY_LOGIC_DIR, setVerbosity, IsaInfo, makeDebugConfig, \
DebugConfig, IsaVersion, coVersionMap
from Tensile.Common.Architectures import architectureMap, archNameForIsa, \
from Tensile.Common.Architectures import ARCH_BUILD_ALIASES, architectureMap, archNameForIsa, \
baseArchName, detectGlobalCurrentArch, \
gfxToIsa, steppingArchOf
from Tensile.Common.Capabilities import applyArchCapOverrides, makeIsaInfoMap
Expand Down Expand Up @@ -709,7 +709,8 @@ def Tensile(userArgs):
# (gfx1250v1, gfx1250v) still resolves to (12,5,0) and would silently
# build the shipping stepping. Check the base name, so target IDs and
# predicates (gfx942:xnack-, gfx950[cu=64]) keep working as before.
if isa is None or baseArchName(arch) not in architectureMap:
if isa is None or baseArchName(arch) not in architectureMap \
or baseArchName(arch) in ARCH_BUILD_ALIASES:
raise ValueError(f"Unrecognized GPU target: '{arch}'")
isaList.append(isa)
archNames.append(arch)
Expand Down Expand Up @@ -738,7 +739,8 @@ def Tensile(userArgs):
if not arch:
continue
isa = gfxToIsa(arch)
if isa is None or baseArchName(arch) not in architectureMap:
if isa is None or baseArchName(arch) not in architectureMap \
or baseArchName(arch) in ARCH_BUILD_ALIASES:
raise ValueError(f"Unrecognized Architecture in config: '{arch}'")
if nameCameFromDetection:
continue
Expand Down
Loading
Loading