Repository navigation
feat(tensilelite): build a solution record in C++ - #13573
jaopaulolc wants to merge 7 commits into
Conversation
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
8666d31 to
942ee97
Compare
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
6cb295f to
75035b4
Compare
jaopaulolc
added
the
rocm:gemm algos
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #13573 +/- ##
===========================================
- Coverage 78.09% 78.09% -0.00%
===========================================
Files 3005 3005
Lines 521953 522055 +102
Branches 81452 81485 +33
===========================================
+ Hits 407579 407656 +77
- Misses 85459 85475 +16
- Partials 28915 28924 +9
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
c5dae21 to
0b08d00
Compare
70bc297 to
fd9cdb2
Compare
15f1e26 to
2832364
Compare
2832364 to
1e2d140
Compare
1e2d140 to
0d1cffb
Compare
0d1cffb to
1fc8739
Compare
ThanHenderson
left a comment
ThanHenderson
left a comment
There was a problem hiding this comment.
Approving, but consider addressing my concerns with the architecture IDs. Could be a follow up though too.
1fc8739 to
533ce1a
Compare
Alex-Vasile
left a comment
Alex-Vasile
left a comment
There was a problem hiding this comment.
Thanks for the fixes since the last round. The gaps I raised in the builder are closed in the code. I am requesting changes for three reasons, each with an inline comment:
- The two new
SolutionRecordgtests fail in the TheRockTest tensilelitejob (gfx94X, job 128031777598: 840 of 852 gtests pass, these 2 fail). SolutionState.SizeMappingOptionalsFollowTheStatereads freed memory.- The record check compares predicate type names only, so it does not yet catch most drift between the Python and C++ builders.
The description also still cites ec66cfb1f05 with 14 tests, says the automated checks are pending, and says the PR is based on the MessagePack writer branch. Could it be refreshed with the final commit, the current test counts (15 SolutionState gtests, 2 SolutionRecord gtests, 16 pytests) and which lane ran each?
| "${CMAKE_CURRENT_SOURCE_DIR}/SolutionRecord_test.cpp" | ||
| ) | ||
| target_compile_definitions(tensilelite-tests PRIVATE | ||
| "TENSILELITE_SOLUTION_RECORD_DIR=\"${SOLUTION_STATE_RECORD_DIR}\"" |
There was a problem hiding this comment.
These two definitions put the build-tree path into the tensilelite-tests binary, and the TheRock Test tensilelite job runs that binary on a different runner from the one that built it. In job 128031777598 the path /__w/rocm-libraries/rocm-libraries/build/math-libs/BLAS/hipBLASLt/build/tensilelite/tests/data/solution-state-catalog does not exist, so CatalogMatchesSolutionState compares the C++ lists with {} and MatchesThePythonBuilder stops at ASSERT_FALSE(bytes.empty()) for all three fixtures. Those are the 2 failures in that job. Math CI passes the same tests only because it builds and runs in one tree.
The other tests in this directory find their data through TestData::Instance(), which resolves <directory of the test binary>/data (tests/TestData.cpp:119). The rocm_install(DIRECTORY "${TEST_DATA_DIR}/" ...) step copies that directory next to the installed binary, so files placed under TEST_DATA_DIR are found there. Could these tests use the same lookup instead of a compile-time path?
Also, fileLines() in SolutionRecord_test.cpp returns an empty list when a file is missing, which is why the catalog failure reads as "lists differ". Could it fail with the missing path instead?
There was a problem hiding this comment.
Fixed in 4806b97c35a. The compile-time path definitions are gone. SolutionRecord_test.cpp now finds the fixtures through TestData::Instance(), in the data directory next to the test binary, like the neighboring tests. A build target, tensilelite-solution-state-fixtures, copies the fixtures into TEST_DATA_DIR on every build, so the rocm_install(DIRECTORY ...) step carries them next to an installed binary. The readers no longer return an empty list for a missing file: they fail with the path, for example cannot open test fixture <dir>/data/solution-state/catalog/catalog_chip_ids.txt. I checked this the way the TheRock job runs it: I copied the tensilelite-tests binary and its data directory to another directory and ran SolutionRecord.* and SolutionState.* from there, and they passed; with the catalog file deleted from the copy, the test failed with the message above.
| StateOptions options; | ||
| options.oneLdsBuffer = true; | ||
| options.bufferLoadOff = true; | ||
| auto const& size = build(options)->sizeMapping; |
There was a problem hiding this comment.
build(options) returns a shared_ptr<ContractionSolution>, and ->sizeMapping is a member of the object it points to. The temporary shared_ptr is destroyed at the end of this statement, so size refers to freed memory on lines 730-733. Binding a reference to a member does not keep the temporary alive. An AddressSanitizer build of this test (a local run, not CI) reports a heap-use-after-free at line 730. CI passes because the freed memory still holds the old values.
Could the result be kept in a named variable (auto solution = build(options);) and read through solution->sizeMapping? Line 692 is fine, because build(options)->problemType.stochasticRounding is read inside one statement.
There was a problem hiding this comment.
Fixed in 4806b97c35a. SizeMappingOptionalsFollowTheState now keeps auto solution = build(options); in a named variable and takes size from solution->sizeMapping. I went through SolutionState_test.cpp and SolutionRecord_test.cpp for other references or views bound to a member of a temporary returned by value and found none left. I reproduced your report with an AddressSanitizer run of the old test (heap-use-after-free at line 730); the same run on the fixed test, and on the SolutionRecord and SolutionState tests, reports nothing.
| add("problemType.stochasticRounding", jsonBool(problem.stochasticRounding)); | ||
|
|
||
| std::vector<std::string> tags; | ||
| collect(*solution.problemPredicate, tags); |
There was a problem hiding this comment.
This compares predicate type names only: collect() (line 117) pushes predicate.type(), so no value inside a predicate is compared in C++. The values are in the checked-in *.predicates.yaml files, but only test_solution_state_record.py reads them. The three fixtures are also one fp16 kernel on three ISAs (MacroTile0 = MacroTile1 = DepthU = 16, GlobalReadVectorWidthA = GlobalReadVectorWidthB = 4), so swapped A/B or M/N values would give the same output.
In a local copy of SolutionState.cpp, each of these changes left SolutionRecord.MatchesThePythonBuilder and SolutionState.* passing: GlobalSplitUCheckMinK 32 changed to 64; the A and B types in TypesEqual swapped; the ReuseAcrossPersistent block deleted; the MXBlockFreeA predicate dropped; occupancyProbe ignored; Int4EncodingA::UnsignedBias8 read as Signed. That is the kind of divergence this check was added to catch, and the last four are the same class of gap that came up in the earlier threads.
Could the C++ side compare predicate values as well, for example by writing the predicates in the {type, index, value} form that predicates.yaml uses, and could the fixtures include a kernel with different tile sizes and at least one non-default feature (MX, int4, StreamK, ReuseAcrossPersistent or an occupancy probe)? The C++ builder and the Python builder are both maintained, so this check is the only thing that fails when they differ.
There was a problem hiding this comment.
Done in 4806b97c35a. The C++ side now compares predicate values, not only type names: SolutionRecord.PredicateValuesMatchThePythonBuilder projects the problem, task and hardware predicate trees into the {type, index, value} form of the checked-in *.predicates.yaml and compares them as a deep comparison, reporting the path of the first mismatch. SolutionRecord.FixturesAreComplete checks every fixture has all its files, and the record test discovers fixtures from the data directory instead of a fixed list. The fixtures are now ten derived states: the three plain fp16 kernels, plus gfx942 stochastic rounding, gfx942 StreamK fp8, gfx950 MX, gfx950 MX with free blocks, gfx1151 int4 (several scale block sizes, UnsignedBias8), gfx1250 StreamK cluster, and gfx1250 StreamK cluster with ReuseAcrossPersistent and an occupancy probe; they have different tile sizes, depth and vector widths, so swapped A/B or M/N values show up.
I applied your six mutations one at a time to SolutionState.cpp; each now fails at least one test: GlobalSplitUCheckMinK 32 to 64 and swapped A/B types in TypesEqual and Int4EncodingA::UnsignedBias8 read as Signed fail PredicateValuesMatchThePythonBuilder; dropping the ReuseAcrossPersistent block and dropping MXBlockFreeA fail both MatchesThePythonBuilder and PredicateValuesMatchThePythonBuilder; ignoring occupancyProbe fails MatchesThePythonBuilder. The stricter comparison also found real differences, which are fixed: the MX predicate order in the C++ builder differed from Python, and the test projection was missing customKernel.args padding and the occupancyProbe line. I also refreshed the description with the final commit and the test counts.
TheRock runs this test from build/share/hipblaslt/tensilelite, which has no src/msgpack. Look up the checkout that produced that copy.
solutionFromState used the test constructor, which freezes a compute-unit count of 0 and never queries the GPU, so every XCC kernel was rejected.
Building the record still leaves the compute-unit count at 0. Evaluating the check reads the device, so a frozen count of 0 cannot reject every kernel.
solutionFromState and FromOriginalState are kept by a checked-in record of the three plain kernels. A Python-side field change fails the pytest until that record is regenerated, and the regenerated lines fail the gtest until the C++ builder matches them.
…g walk The installed pytest tree has no SolutionState.cpp, and the plain-kernel YAML was collected as a tensilelite config. The catalog now compares Python to a checked-in list, and the fixtures live under test_data so findConfigs skips them.
Read MX block-free and occupancy-probe fields, align the ReuseAcrossPersistent predicates with develop, and make the catalog test compare exported tables while loading fixture data from the test data dir.
533ce1a to
4806b97
Compare
JIRA ID
Motivation
solutionFromStatebuilds aContractionSolutionin C++ from the stateSolution.FromOriginalStatereads. This pull request is the builder, its tests and fixtures, and the CMake wiring. It is based ondevelop; the MessagePack writer it was stacked on has merged.Technical Details
TensileLite::solutionFromStatereads one MessagePack map and returns aContractionSolution: name, kernel name, problem type, predicates, size mapping, custom kernel, internal-argument support, and the other fieldsLoadLibraryFilestores. It does not buildoriginalSolution.The record follows the Python builder for the fields that had drifted:
PersistentLoopArgsVersion0, 1, and 2. Version 2 requires StreamK with a static grid, aClusterDimaxis above 1, andKernArgsVersion3.DataTypeAandDataTypeBas the compute inputs when both MAC inputs are still theDataTypedefault.ScaleBlockSizeA(anOrwhen several of 32, 64, and 128 are listed),ScaleZeroPointA, andInt4EncodingA.StochasticRounding,CustomKernelSizePredicate(SizeGreaterThan,SizeLessThan,SizeEqual), and the size-mapping fields the performance model reads:WavefrontSize,1LDSBuffer,SourceSwap,UnrollMajorLDSA/B,enableLDSTrA/B,BufferLoad, andBufferStore. AbsentBufferLoadandBufferStorestay true. AbsentUnrollMajorLDS*stay -1.Idealskey isstd::runtime_errornaming the key.WorkgroupMappingXCCis stored with a compute-unit count of 0. Building the record does not query a GPU. A gfx950DeviceNamesentry such asDevice 75a0becomes a PCI chip-id predicate.test_solution_state_catalog.pycompares the hand-copied data-type names and build chip ids withrocisa.enum.DataTypeEnumandGFX_CHIP_IDS. The chip ids live inTensile/Architectures.hpp, so other C++ code can use the same list.The C++ record is checked against the Python builder in two ways.
SolutionRecord.MatchesThePythonBuildercompares the record lines, andSolutionRecord.PredicateValuesMatchThePythonBuildercompares the problem, task and hardware predicates, values included, with the{type, index, value}form of the checked-in*.predicates.yaml.SolutionRecord.FixturesAreCompletefails when a fixture file is missing. The fixtures undertensilelite/Tests/unit/test_data/solution_recordsare ten derived states: three plain fp16 kernels (gfx90a, gfx942, gfx950) and gfx942 stochastic rounding, gfx942 StreamK fp8, gfx950 MX, gfx950 MX with free blocks, gfx1151 int4, gfx1250 StreamK cluster, and gfx1250 StreamK cluster withReuseAcrossPersistentand an occupancy probe.python tensilelite/Tests/unit/test_solution_state_record.pyregenerates them, andderive_solution_records.pyderives states from library logic files and benchmark configs. Changing, for example, a predicate value, the A and B types ofTypesEqual,MXBlockFreeA,occupancyProbeor the int4 encoding inSolutionState.cppmakes at least one of these tests fail.The record tests read their fixtures through
TestData::Instance(), from thedatadirectory next to the test binary. A build step copies the fixtures there on every build, so a binary run on another machine from an installed copy finds them. A missing fixture fails with its path.Test Plan
Host build, gfx950, Release.
tensilelite-tests --gtest_filter='SolutionState.*:SolutionRecord.*'pytest tensilelite/Tests/unit/test_solution_state_record.py tensilelite/Tests/unit/test_solution_state_catalog.pyTest Result
On
4806b97c35a, on a gfx950 host build (Release,build-13573):tensilelite-testsran 854 tests, 848 passed, 6 were skipped (they need a device feature this host lacks) and none failed. That includes 15SolutionStateand 4SolutionRecordgtests.pytestontest_solution_state_record.pyandtest_solution_state_catalog.pypassed 48 tests. Thetensilelite-testsbinary and itsdatadirectory also passed when copied to another directory and run from there.The gtests are part of
tensilelite-tests, which the TheRockTest tensilelitejob runs on gfx94X. The pytests run in theTensileLite coverage / Linuxlane, in the TheRockTest tensilelitejob, and in the Math CIpreliminaryandtensilelite-unit-codecovjobs. Check results for this head are on the pull request.Submission Checklist
Risk level
Medium. A record built here is what a later library write stores. The new cases are the ones the Python builder already accepts: cluster StreamK, mixed MX operands, int4 group scales, stochastic rounding, exact-size bounds, and the size-mapping fields. Rollback is reverting this pull request.