Repository navigation
Add thread include - #212
Merged
bstefanuk merged 4 commits intoJun 18, 2025
Merged
Conversation
…evelop/bstefanuk_Tensile/add-thread-include
|
/azp run Tensile |
|
Azure Pipelines successfully started running 1 pipeline(s). |
davidd-amd
approved these changes
Jun 18, 2025
ibrahimw1
approved these changes
Jun 18, 2025
bstefanuk
deleted the
import/develop/bstefanuk_Tensile/add-thread-include
branch
June 18, 2025 17:37
assistant-librarian Bot
pushed a commit
to ROCm/Tensile
that referenced
this pull request
Jun 18, 2025
Add thread include MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 🔁 Imported from [#2151](#2151) 🧑💻 Originally authored by @bstefanuk
bstefanuk
added a commit
to bstefanuk/rocm-libraries
that referenced
this pull request
Jun 19, 2025
🔁 Imported from [ROCm/Tensile#2151](ROCm/Tensile#2151) 🧑💻 Originally authored by @bstefanuk --------- Co-authored-by: Braden Stefanuk <121893577+bstefanuk@users.noreply.github.com> Co-authored-by: assistant-librarian[bot] <assistant-librarian[bot]@users.noreply.github.com>
ammallya
pushed a commit
that referenced
this pull request
Sep 24, 2025
* MIOpen convolution forward integration
ammallya
pushed a commit
that referenced
this pull request
Sep 24, 2025
* MIOpen convolution forward integration [ROCm/hipDNN commit: a4e2c82]
amd-alavin
added a commit
that referenced
this pull request
Oct 9, 2026
JIRA ID : ALMIOPEN-2361 Takes the vendored `src/hipconv` snapshot from v0.3.2 to v0.4.0 and fixes ConvHipConv's perf-db round trip, which broke SD2.1 fp16 batch 4 on gfx950 in the 10-02 nightly (ROCM-32185). ## ConvHipConv perf-db records #9680 made ConvHipConv's perf config the hipconv kernel label, e.g. `direct[tile_size_k=256,tile_size_n=1,...]`. `PerfConfigBase` deserializes field by field and splits on `,`, so a stored label reloaded as `direct[tile_size_k=256`, matched no kernel, and the solver fell back to config 0. Under MIGraphX's `--exhaustive-tune`, execution then built config 0 against the workspace Find had reported for the tuned config and threw `ConvHipConv: not enough workspace for direct kernel` (`conv_hipconv.cpp:702`). - `fix(miopen): serialize the ConvHipConv perf config whole`: `PerformanceConfigConvHipConv` derives from `PerfConfig` and stores the label as one string. - `fix(miopen): key ConvHipConv perf-db records to the hipconv minor version`: a record is `<hipconv config_version>:family[full description]`, e.g. `v0.4:direct[tile_size_k=64,...,direction=fprop,type=fp16,...]`. hipconv guarantees that a full description selects the same config only within a minor version, so a record from another minor version, an unstamped record, and any record in a build whose `config_version()` is `unknown` select nothing. The solver then falls back as for any invalid record (search, the AI picker, or hipconv's top-ranked config). Every ConvHipConv record written before this change is unstamped and is discarded once. The performance-log kernel name keeps the short label. #9680's perf-config picker catalog (`lgbm_pcfg.bin`) holds 166 ConvHipConv labels in the old form. They never resolved, since every one contains a comma, and they will not match stamped records; they need regenerating against v0.4.0. ## hipconv v0.4.0 Per-version detail is in the shipped `src/hipconv/CHANGELOG.md`. - Every kernel family describes its configs, and `describe_config(kernel, full)` renders every field, which is what the stamp stores. `version()` and `config_version()` report the snapshot's version from the shipped `VERSION` file. - gfx1250 `direct_wgrad`: dense backward-weights, fp16/bf16 with an fp32 dW, 2x2, 3x3, 4x4, 5x5, 3x1 and 1x3 filters, stride 1 (hipconv #212, #297). - `ConvParams::is_valid()` rejects a layer whose dilated filter is wider than its padded input (hipconv #288). - Explicit-GEMM launch fixes (hipconv #281); MIOpen leaves those paths off. Data types are unchanged: fp16, bf16 and TF32. `projects/miopen/CHANGELOG.md` gains the gfx1250 dense backward-weights kernel and the perf-db fix under MIOpen 3.6.2. ## Downstream patches None. On develop, `projects/miopen/src/hipconv` was byte-identical to hipconv v0.3.2. ## Validation - hipconv's internal suite at v0.4.0: 65/65 on gfx950 (MI355X); 65/65 on MI455 A0 (gfx1250) with explicit GEMM off, including about 14,000 cases that run every kernel config against hipconv's test layers. - `test_unit_conv_solver_ConvHipConv` on gfx950 (MI355X), MIOpen built with clang 24 (ROCm 10.1) from this branch: 377 pass and 24 skip (the gfx125X-only suites), the same in two repeats. On MI455 A0 (gfx1250), built for `gfx1250;gfx1250-strict`: 210 pass and 191 skip (the gfx950-only suites), the same in two runs. New: `CPU_UnitTestConvSolverConvHipConvPerfConfig_NONE.SerializeRoundTrip` and `GPU_UnitTestConvSolverConvHipConvPerfConfig_FP16.VersionStamp`, which pass on both. - ROCM-32185 reproduction on gfx950, the 18 stride-1 SD2.1 batch-8 layers through Find 2.0 with ConvHipConv forced: exhaustive Find, save and load the solution, run on a second handle with the workspace Find reported. On develop it throws at `conv_hipconv.cpp:702`. With this branch, every layer passes on a fresh user perf db (one search each), again reloading those records (no search, no warning), and with the records restamped `v0.3` (each rejected once with the "Invalid config" warning, re-searched, then passing). - clang-tidy over the consumer TU is clean (LLVM 21.1.8). A find-db record also caches a workspace size and carries no hipconv version; that is unchanged here. --------- Co-authored-by: Boris Prokopenko <boris.prokopenko@amd.com> Co-authored-by: Carsten Uphoff <carsten.uphoff@amd.com> Co-authored-by: Jin Zhou <jin.zhou@amd.com> Co-authored-by: Petar Jovanovic <petar.jovanovic@amd.com> Co-authored-by: Yingmao Lu <yingmao.lu@amd.com>
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.
Add to resolve compile time errors.
🔁 Imported from ROCm/Tensile#2151
🧑💻 Originally authored by @bstefanuk