Repository navigation
Conversation
⏳ Waiting on Required Checks⏳ Required checks have not all concluded yet. This bot run is exiting without failing so those checks can report independently.
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
⏳ Waiting for required checks to start or finish:
|
newling
force-pushed
the
refactor/host-validation-pure-boundary
branch
from
August 11, 2026 00:09
dcafb2a to
b344d93
Compare
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #10553 +/- ##
============================================
- Coverage 77.16% 76.84% -0.32%
============================================
Files 2888 259 -2629
Lines 483904 99184 -384720
Branches 73029 18628 -54401
============================================
- Hits 373380 76212 -297168
+ Misses 85920 18461 -67459
+ Partials 24604 4511 -20093
🚀 New features to boost your workflow:
|
newling
force-pushed
the
refactor/host-validation-pure-boundary
branch
from
August 11, 2026 23:56
b344d93 to
b6d77a7
Compare
Generated Documentation |
Code Coverage Report for gfx950Summary
This PR adds/edits 53 newly uncovered lines. Artifacts
Commit Hashes |
newling
force-pushed
the
refactor/host-validation-pure-boundary
branch
from
August 12, 2026 19:42
3474775 to
518353f
Compare
newling
force-pushed
the
refactor/host-validation-pure-boundary
branch
from
August 12, 2026 21:30
518353f to
d639690
Compare
…switches Develop's MAF initialization modes reached both recipe switches unhandled (clang -Wswitch). MAF fills tensors directly, so name the modes and reject them with a specific message.
normalizeBlockScale only validated its argument's shape, so make it a validation function. The default scales are already optional tensors; pass them through without rebuilding the optional.
initializeMatrix took forceNaN, a one-special value, positiveOnly,
integer-exact options and a batch index positionally, so most callers passed
'false, std::nullopt, false, {}'. Group them in MatrixInitializationOptions
with defaults and designated initializers.
OpenMP is a required HostNumerics dependency, so remove the serial `#ifdef _OPENMP` fallbacks from the shared threading helpers and MX packing. Route every parallel operation loop through forEachParallelIndex, which rethrows the first exception from its workers. Generation and MX generation had their own OpenMP regions with copies of that partitioning and exception-forwarding code; the unbounded MX generator's loop had no forwarding at all although each element looked up the recipe index order, which allocates. MX generation now resolves that order once per call. Freivalds projection used a separate OpenMP threshold and the full runtime thread count; it now follows the shared thread-count policy. Its residue arithmetic is exact, so results do not depend on the thread count. Every caller passed 8 as the default maximum thread count, so it is now a constant. evenlyPartitionedRange asserts its internal partition index instead of throwing, returns a pair instead of a struct that shadowed the public IndexRange, and also replaces structured sparsity's own chunk arithmetic.
OpenMP is required, so the generation and MX thread-count determinism checks and the structured-sparsity benchmark no longer need serial alternatives.
The comparison test's "zero selection stride" case only constructs an IndexRange; a comma expression discards the order argument and no comparison runs. Core tests already require IndexRange to reject a zero stride.
The blocked, BLAS and allocation-failure tests include private sources that use the OpenMP threading helpers. Without OpenMP flags those tests compiled the helpers' serial branch while the library compiled the OpenMP branch, and without the serial fallback they need omp.h from OpenMP's include path.
…r10553-scan-combined # Conflicts: # shared/hostnumerics/tests/check_amd_gpu_layout_boundary.cmake
…r10553-scan-combined
1 task done
…cit seed FreivaldsProbeAlgorithm offered the unbiased rejection mapping and a modulo mapping kept for an earlier hipBLASLt sequence. No caller selected the modulo mapping, so remove it with its enumeration. With one mapping, FreivaldsProbeOptions only carried a seed that defaulted to zero; callers now pass the probe seed to prepareFreivalds directly.
…onents The prototype added identical 139-line modules to hipBLASLt and rocRoller that ran git archive at configure time to extract shared/hostnumerics when a sparse checkout omitted it. Other shared dependencies (origami, stinkytofu, formerly mxDataGenerator) are listed in each sparse checkout instead, and the READMEs already list shared/hostnumerics. Remove the modules: use the in-tree source when present, otherwise require the installed package. Replace the retired shared/mxdatagenerator entry with shared/hostnumerics in the two CI workflows that build hipBLASLt from a sparse checkout, and drop the unused TENSILELITE_ENABLE_MXDATAGENERATOR forwarding.
GenerationRecipeSettings defaulted its index order to the rank-zero
permutation and its seed to zero, and GenerationRecipe kept a hidden flag
that replaced the stored order with last-dimension-fastest. Add IndexOrder,
which holds either LastDimensionFastest or an explicit Permutation and
resolves it against the tensor rank when used; an explicit permutation must
match that rank. It has no default.
Recipe factories now take the seed and index order as required arguments:
GenerationRecipe::realOnly(component, seed, LastDimensionFastest{});
GenerationRecipe::realOnly(component, seed, Permutation{1, 0});
This replaces GenerationRecipeSettings, the three *LastDimensionFastest
factories and withLastDimensionFastestIndexOrder. The index offset stays a
withIndexOffset modifier whose initial value is zero. Python exposes
IndexOrder with implicit conversion from LastDimensionFastest and
Permutation. Generated values are unchanged for the same seed, order and
index. Migrate the hipBLASLt, TensileLite and rocRoller call sites.
…explicit
StructuredSparsityPattern defaulted to axis 0, a fixed 2:4 pattern retaining
positions {0, 1}, seed 0 and a rank-zero index order, and it carried both
fixedPositions and seed although each selection uses only one of them.
Represent the selection as a variant of FixedStructuredSparsitySelection
(positions) and RandomStructuredSparsitySelection (seed), use IndexOrder for
the line numbering, and remove every default member value:
StructuredSparsityPattern pattern{
.axis = 1,
.groupSize = 4,
.retainedElements = 2,
.selection = FixedStructuredSparsitySelection{{0, 2}},
.indexOrder = Permutation{0, 1},
};
Python constructs the pattern with all five fields. Results are unchanged;
TensileLite's pruning modes map to the same positions and seed.
The host reference took scale tensors and one shared block size, with block size 0 meaning unscaled, an empty scale vector meaning no runtime scale, DataType::None meaning no scale type, and a block size of K standing in for a single scale. A one-value buffer also silently meant a uniform scale. Pass each operand's scale as std::optional<MatmulBlockScale>, so the tensor and its block size travel together. generateGEMMInputs takes std::optional<BlockScaleFormat> per operand and returns MatmulBlockScales; hostBlockScale requires one value per block, and hostUniformScale names the single-scale case as one block spanning K. The client and tests translate their ScaleMode settings into these values. Numerical results are unchanged: the blocked accumulation splits K at the same boundaries.
The reference adapter took --num-elements-to-validate as a size_t, so 0 and -1 (cast to SIZE_MAX) both meant every element. Take a ReferenceElements value instead, either all() or sampled(count), and translate the command-line convention once in ReferenceValidator: N > 0 samples about N elements, and -1 or 0 computes every element (0 still disables checking). Selections are unchanged.
…dapter The adapter stored MX block size 0 to mean 'not MX', with placeholder E8M0 scale types and free-block size 1. Store std::optional<MxOperandScaling> per operand instead, bundling scale type, block size and free-block size, and translate TensileLite's zero block sizes once when the problem is read.
A zero all-close tolerance selected the unit check (exact, or 4 ULP for FP32 and FP64). Make MatmulValidationCase::allCloseTolerance optional and document that leaving it unset selects the unit check; the tolerance policy sets it only for the GFX11 low-precision and special-initialization cases.
Operation inputs defaulted to values a reader has to look up: MXFP4 data with E8M0 scales in 32-element blocks along axis 0 and derived scales; an MX represented range of [-1, 1]; a block-scale size of zero, which is invalid; Float32 epilogue and LayerNorm outputs; LayerNorm axis 0, Float32 accumulation and epsilon 1e-5; and distribution parameters such as a uniform integer range of [0, 1], a standard normal, one binomial trial, dimension 0, a remainder divisor of 1, a zero raw constant and even-parity sign alternation. Remove those default member values and state the values at every call site. Python constructors take the same required arguments. Defaults that remain are absent std::optional values, false opt-in flags, named Default enumerators, identity transforms (affine scale 1 and offset 0, remainder offset 0), resource limits, and zero-initialized result counters. The MAF adapters pass MxScaleGenerationMode::One, which MAF generation already forced. Results are unchanged.
Axis orders are Permutations throughout the core API, but copyWithPermutedDimensions accepted a span or braced list of axis numbers and constructed the Permutation internally. Take const Permutation& so callers state the order with the same type as every other API, and migrate the hipBLASLt client, samples and TensileLite callers. TensileLite's swizzle permutation constants become Permutations. The tensor transformation test keeps the rank check; duplicate and out-of-range entries are Permutation construction errors covered by the Permutation tests.
…itly ScalarConversionOptions defaulted four of its six policies, including the non-obvious FloatingSourcePrecision::Float32, so a call naming only the integer policy silently selected the encoded-storage behavior. Its constructor and the Python ScalarConversionOptions now require all six. convertScalar converts to native types and only uses the integer rounding and overflow policies, so it takes a new IntegerConversionOptions pair instead of the full storage options. Storage code passes options.integerConversion(). Callers that relied on the defaults now spell out the same values, so results are unchanged: the implicit native and storage conversions, exact-integer output, the Freivalds and codec tests, hipBLASLt's device validation test, rocRoller's host reference output conversion, and a Python precision test.
LayoutFailure stored 0 in dimension, value and limit when a failure had no
axis or no sizes, so a rank mismatch looked like a failure on axis 0. The
fields are now optional and every construction states which ones apply.
Layout() silently meant a scalar with zero offset. Remove it from C++ and
Python; callers write Layout(Shape{}, {}, 0). A moved-from Layout is still
that scalar layout.
…th explicit entry points Three Python factories gave None a hidden meaning: - hv.from_numpy(array, scalar_type=None) inferred the type when None and converted otherwise; - hv.to_numpy(tensor, dtype) required a dtype but treated None as the default decoded dtype; - Tensor.from_storage(type, shape, storage, strides, offset) treated strides=None as C-contiguous. hv.from_numpy(array) now always stores the exactly matching ScalarType and hv.from_numpy_converted(array, scalar_type) converts. hv.to_numpy(tensor) decodes into default_decoded_dtype and hv.to_numpy_as(tensor, dtype) into a named dtype. Tensor.from_storage takes an explicit Layout, matching C++ takeOwnershipOfEncodedBackingStorage. Tensor.from_numpy's optional scalar_type only asserted that it equaled the dtype's type; remove it. Callers are rewritten mechanically and reformatted with black.
…ayout helpers private
Layout::getElementBounds returned the pair {0, -1} for an empty shape, so
callers had to recognize upper < lower as "no element". It now returns
std::optional<ElementOffsetBounds> (named lowest and highest), with
std::nullopt for an empty shape; Python element_bounds returns None.
The public headers also carried helpers used only by HostNumerics sources:
detail::getElementBounds (a wrapper of the member function), elementBounds,
strideMagnitude and hasProvablyDistinctElementOffsets in layout.hpp, and
canBroadcastTo in shape.hpp. Move the last four to src/detail/layout_analysis.hpp,
include it where used, and remove the wrapper. TensileLite's adapter calls
Layout::getElementBounds directly.
…lementCount detail::ShapeAccess was an installed friend struct with two jobs. It exposed Shape's non-throwing element-count check, which TensileLite uses to reject an overflowing descriptor without exceptions. It also had its own index-order validator, duplicating Permutation's. Make Shape::getElementCount a public static fallible query, following the get-prefix convention for result-returning queries, and point TensileLite at it. validatedIndexOrder now validates by constructing a Permutation. Delete shape_access.hpp from the installed headers. runtime_tensor_io.hpp uses the shared strideMagnitude instead of a local copy.
MSVC builds select the LLVM OpenMP runtime because the MX packing loops use unsigned loop indices, which the default OpenMP 2.0 runtime rejects. OpenMP_RUNTIME_MSVC first appeared in CMake 3.30, and older versions ignore it silently and then fail to compile those loops. Stop at configure time with a message naming the requirement.
…licit Tensor(Source) was a converting constructor, so assigning 2.0 to a Tensor option or passing 2.0f where an operand Tensor was expected silently created a rank-zero Float64 or Float32 tensor whose type came from the C++ literal. Make it explicit and spell out the seven test call sites that relied on the conversion. Document that Tensor(type, shape|layout) zero-initializes storage.
…ng/pr10553-integration
…ng/pr10553-integration
Only the three MX packing loops used unsigned OpenMP loop indices, which MSVC's OpenMP 2.0 runtime rejects. Give them signed indices, as the shared threading helpers already use, and drop the LLVM OpenMP runtime selection and its CMake 3.30 requirement.
This branch has not been deployed
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
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.
JIRA ID : AIHPBLAS-4195
Motivation
hipBLASLt, TensileLite and rocRoller each perform host data generation, reference arithmetic and numerical comparison. Those implementations repeat scalar conversion and tensor storage rules. This prototype moves that work into a shared HostNumerics component, with each product translating its descriptors and buffers at the boundary.
Technical Details
HostNumerics supplies scalar types and codecs, logical shapes and affine layouts, owned or shared tensors, typed generation recipes, reference operations and structured comparison reports through C++ and Python APIs. Generic numerical operations run on the CPU. A separate CPU module handles AMD GPU storage layouts. The product adapters retain responsibility for GPU execution, uploads and diagnostics. MX generation moves into HostNumerics, replacing the shared mxDataGenerator component.
Public APIs avoid implicit meanings. Numerical types, tolerances, NaN policies, seeds, sampling limits, conversion policies and execution policies are required arguments or required fields, with no default that silently selects a behavior; absence is expressed with
std::optionalor a separate entry point rather than 0, -1 or an empty value. Index orders are explicit: generation recipes and structured sparsity take anIndexOrderthat is eitherLastDimensionFastestor aPermutationwhose size must equal the tensor rank, and an emptyPermutationmeans rank zero only. Freivalds verification uses a single probe mapping. The adapters state their existing choices, so numerical results are unchanged; product command-line and configuration values that use such sentinels are translated once where the adapter reads them.HostNumerics requires OpenMP; there is no serial-only build. Parallel work uses shared helpers with signed loop indices, which MSVC's OpenMP 2.0 runtime accepts. Exceptions are carried from worker threads to the caller only where the parallel work can actually throw; other parallel regions require
noexceptwork. hipBLASLt and rocRoller use the in-tree HostNumerics sources or an installed HostNumerics package, and their sparse checkouts listshared/hostnumerics.Test Plan
Build the standalone component with Clang and GCC, run its native and Python suites, and check ASan/UBSan and installed C++ and Python consumers. Build affected hipBLASLt and TensileLite clients, samples and tests, plus rocRoller clients and tests. Run the host adapter suites, a rocRoller GEMM generate-and-validate run on a local GPU, root pre-commit, component formatting and HostNumerics clang-tidy.
Test Result
Local Clang 23 and GCC 13 standalone builds pass all native and Python tests, including BLAS-backed GEMM and installed-package checks. Clang 18 ASan/UBSan passes the native suite; its two nested package tests need an OpenMP runtime for Clang 18, which CI installs. Host-only hipBLASLt and TensileLite builds pass
hipblaslt-test-hostunit, the hipBLASLt host suites andtensilelite-tests, excluding two device tests that also crash at the previous head on this machine's GPU. rocRoller host reference, data-generation and Catch tests pass, and rocRoller GEMM validation is correct on gfx1201. Windows/MSVC, hipBLASLt device libraries and MX-capable GPU validation remain pending.Submission Checklist
Risk level
Medium. The prototype changes public C++ and Python signatures and their product callers. Numerical behavior is covered by the existing host oracles and adapter tests.