Repository navigation
[rocprofiler-sdk][ATT] Enable back ATT tests in rocprofiler-sdk - #1881
Merged
Merged
Conversation
Contributor
There was a problem hiding this comment.
Pull Request Overview
This PR re-enables ATT (Advanced Thread Trace) tests that were previously disabled due to rocr runtime changes. The tests are being enabled again following a rocr patch that fixed the underlying issues.
- Removes
DISABLED ${ROCPROFILER_DISABLE_UNSTABLE_CTESTS}property from three thread-trace API tests - Removes conditional logic that would disable tests based on
ROCPROFILER_DISABLE_UNSTABLE_CTESTSflag - Consolidates multi-line test property definitions to single-line format
Reviewed Changes
Copilot reviewed 3 out of 3 changed files in this pull request and generated no comments.
| File | Description |
|---|---|
| projects/rocprofiler-sdk/tests/thread-trace/CMakeLists.txt | Re-enables three thread-trace API tests (single, multi, and agent) by removing DISABLED property and consolidating formatting |
| projects/rocprofiler-sdk/tests/rocprofv3/advanced-thread-trace/CMakeLists.txt | Removes conditional check that would force disable ATT tests based on ROCPROFILER_DISABLE_UNSTABLE_CTESTS |
| projects/rocprofiler-sdk/samples/thread_trace/CMakeLists.txt | Removes conditional check that would force disable thread-trace sample based on ROCPROFILER_DISABLE_UNSTABLE_CTESTS |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
ApoKalipse-V
approved these changes
Nov 17, 2025
venkat1361
force-pushed
the
users/venkat1361/enable-att-tests
branch
from
November 19, 2025 17:02
67812a6 to
cc2c72c
Compare
ammarwa
approved these changes
Nov 24, 2025
jamessiddeley-amd
pushed a commit
that referenced
this pull request
Dec 11, 2025
jrmadsen
approved these changes
Feb 24, 2026
Contributor
Author
|
/AzurePipelines run rocm-ci-caller |
|
No pipelines are associated with this pull request. |
1 task done
venkat1361
pushed a commit
that referenced
this pull request
Mar 12, 2026
## Motivation <!-- Explain the purpose of this PR and the goals it aims to achieve. --> Fix failing PSDB job ## Technical Details <!-- Explain the changes along with any relevant GitHub links. --> #1881 accounted for hidden call to `hipGetDeviceProperties*` in the HIP summary check but did not account for it in `get_dims(total)` and `get_dims(hip_and_marker)`. ## JIRA ID <!-- If applicable, mention the JIRA ID resolved by this PR (Example: Resolves SWDEV-12345). --> <!-- Do not post any JIRA links here. --> N/A ## Test Plan <!-- Explain any relevant testing done to verify this PR. --> N/A. This is a testing fix. ## Test Result <!-- Briefly summarize test outcomes. --> N/A. This is a testing fix. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
dayatsin-amd
added a commit
that referenced
this pull request
Apr 7, 2026
…mpute-external-merge-new-800 (#755) * [rocprofiler-sdk][rocprofiler-register] add CPackComponent (#3499) * [rocprofiler-sdk][rocprofiler-regsiter] add CPackComponent * add includes to the top * Tuning: added AINIC-specific tuning (#2761) * projects/RCCL: automatic AINIC detection and gfx942-specific tuning for AINIC RoCE Improves AINIC OOB experience. Changes: - Added rcclDetectIBNics() to auto-detect IB devices and their rate. The RCCL_AINIC_ROCE environment variable can still be used to manually override the auto-detection behavior. - Add tuning_model_8 with AINIC-specific for gfx942 - Refactor architecture-specific tuning overrides into a centralized rcclApplyTuningOverrides() function, replacing duplicated inline checks in rome_models.cc - Update rcclGetTuningIndexForArch() to use AINIC-specific tuning map when AINIC is enabled * projects/rccl: net.cc fixed error handling and addressed review comments --------- Co-authored-by: alexander-sannikov <Aleksandr.Sannikov@amd.com> * [rocprof-trace-decoder] Fix missing test dependency (#3800) * [rocprofiler-sdk] Add firmware restrictions validation for ROCprofiler-SDK counter collection (#344) * Add firmware restrictions validation for ROCprofiler-SDK counter collection Implement comprehensive firmware version checking to ensure counter collection operates on devices with compatible firmware versions, preventing issues from outdated firmware that could cause device instability or incorrect data. Key features: • YAML-based firmware restrictions configuration with schema versioning • Support for CP/MEC and SDMA firmware type restrictions per architecture • Environment variable (ROCPROFILER_FIRMWARE_RESTRICTIONS_PATH) and install path discovery • Thread-safe single-execution validation using std::call_once • Comprehensive warning messages for firmware version violations • Graceful handling of missing or invalid restriction files Implementation: • New FirmwareRestriction struct with firmware_type, min_version, reason, and affected_architectures • parse_firmware_restrictions() for YAML parsing and validation • check_agent_firmware_restrictions() for agent firmware validation • check_installed_firmware_restrictions() for file-based validation with caching • Extensive test suite covering YAML parsing, agent validation, and edge cases • Complete documentation in counter_collection_services.rst * remove unnecessary test * [rocprofiler-sdk] Rename counter_defs.yaml to config.yaml and consolidate configuration (#357) * Rename counter_defs.yaml to config.yaml and consolidate configuration files │ - Remove separate firmware_restrictions.yaml file │ - Rename counter_defs.yaml to config.yaml for unified configuration │ - Update all code references from counter_defs.yaml to config.yaml │ - Update documentation to reflect new config.yaml filename │ - Consolidate firmware restrictions into main config file │ - Update CMakeLists.txt to remove firmware_restrictions.yaml from build │ - Update test cases to use unified schema format * Minor fixes + adding config.yaml * Add missing newline at end of firmware_restrictions files Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Address PR review comments: default initializers, ROCP_WARNING, current_version - Add default member initializers to FirmwareRestriction struct fields - Add current_version field to track actual firmware version on agent - Change ROCP_ERROR to ROCP_WARNING for firmware restriction messages (per jrmadsen: if it's truly an error, it should be fatal) Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Address PR review comments: default initializers, ROCP_WARNING, current_version - Empty affected_architectures list now applies restriction to all architectures - Add ROCP_WARNING when firmware restriction YAML parsing fails - Rename EmptyArchitecturesList test to EmptyArchitecturesListAppliesToAll with behavioral verification via check_agent_firmware_restrictions Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Apply clang-format-11 formatting fixes Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> --------- Co-authored-by: root <Benjamin Welton> Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com> * [rocprofiler-sdk] Agent topology updates (#3059) * [rocprofiler-sdk] Agent topology updates - support env variable pointing to topology - ROCPROFILER_KFD_TOPOLOGY - AMD_KFD_TOPOLOGY - HSA_MODEL_TOPOLOGY - handle APUs (treat as GPUs) - add static topology files (without caches) - AMD Ryzen Threadripper PRO 3945WX 12-Cores - gfx906 - gfx1102 - gfx1032 - support /sys/devices/virtual/kfd/kfd/topology/nodes * Added node for gfx942 * Added node for gfx950 * Added node for gfx1201 * The properties of every node/agent are verified against a mock agent. * Remove comparison against HSA - Setting HSA_MODEL_TOPOLOGY requires HSA_MODEL_LIB - we don't want to mock up a model library - also, formatting --------- Co-authored-by: Evan Phoukong <Evan.Phoukong@amd.com> * [rocprofiler-sdk] Set ROCPROFILER_REGISTER_LIBRARY env variable in initialization (#2335) * Set ROCPROFILER_REGISTER_LIBRARY in rocprofiler-sdk initialization - Setting this environment variable tells rocprofiler-register which specific rocprofiler-sdk library to load (if it hasn't loaded it already) - This ensures that all rocprofiler-register instances communicate with the same rocprofiler-sdk library and/or the rocprofiler-sdk library that a tool library called rocprofiler_force_configure in. * Apply suggestions from code review Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> * Apply code review suggestions - added unit tests for common::dl::get_linked_path and common::dl::get_symbol_path - tweaked logging * Fix get_this_library_path() in rocprofiler-sdk-rocattach - search for versioned library, soversioned library, non-versioned library * Fix build error - missing include for <fmt/ranges.h> --------- Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> * [rocprofiler-sdk] Fix rocprofv3 handling of roctxProfilerPause/Resume (#1496) * [rocprofiler-sdk] Fix rocprofv3 handling of roctxProfilerPause/Resume - Fix roctxProfilerPause / roctxProfilerResume - support control over counter collection, ATT, PC sampling - Add --selected-regions-ref-count which ensures nested regions of resume / pause are ignored - e.g. rocprofv3 without this option enabled will pause at step 3 below instead of step 4: 1. resume 2. resume 3. pause 4. pause . * Kernel trace and counter tests added * Added pc-sampling and nested tests, removed ATT support * Changed CI_LOG(WARNING) to WARNING to test selected-regions-ref-count option, fixed pc-sampling test skips, fixed nested test * Added nested call after final pause to ensure it still works * Addressing comments and minor logic changes * Modified pytest.ini and added comment * Updating for comments. Need to confirm testing passes * Fix test error --------- Co-authored-by: Ian Trowbridge <Ian.Trowbridge@amd.com> * Add float/double support to environment variable helpers (#3678) Co-authored-by: ammallya <ameyakeshava.mallya@amd.com> * Dont build examples if we dont have MPI (#3876) * [ROCM-19766] Remove references to deprecated amd-smi reset -r (#3751) * [SWDEV-553392] Remove references to deprecated amd-smi reset -r Signed-off-by: Maisam Arif <Maisam.Arif@amd.com> * Changelog update Signed-off-by: Maisam Arif <Maisam.Arif@amd.com> --------- Signed-off-by: Maisam Arif <Maisam.Arif@amd.com> * [AMD-SMI] Fix CLI breakage when amdgpu is not present (#3765) * [AMD-SMI] Fix CLI breakage when amdgpu is not present Signed-off-by: Maisam Arif <Maisam.Arif@amd.com> * Adjusted Init to better catch driver loading issues Signed-off-by: gabrpham_amdeng <Gabriel.Pham@amd.com> --------- Signed-off-by: Maisam Arif <Maisam.Arif@amd.com> Signed-off-by: gabrpham_amdeng <Gabriel.Pham@amd.com> Co-authored-by: gabrpham_amdeng <Gabriel.Pham@amd.com> * Update Dependabot setting for docs (#3602) * Centralize dependabot setting file * Remove duplicated entry * Update .github/dependabot.yml Co-authored-by: David Galiffi <David.Galiffi@amd.com> --------- Co-authored-by: David Galiffi <David.Galiffi@amd.com> * SWDEV-549711: Fix IsLeafNodeSyncRequired to detect leaf segments at all levels (#3867) Previously IsLeafNodeSyncRequired only checked segments at max_dependency_level_, missing leaf segments (no outgoing edges) at earlier levels. This caused a missing sync when a host node with no dependents ran on a parallel stream — hipStreamSynchronize on the launch stream would return before the host callback completed. Fix: iterate all segments and count those with empty segment_ids_edges to correctly identify all leaf segments across all dependency levels. * ROCM-1667 - Enable OpenCL build from rocm-systems for ROCr backend on Windows (#3857) * ROCM-1667 - Enable OpenCL build from rocm-systems for ROCr backend on Windows * Address copilot comments * [rocprofiler-sdk] HSA API changes for HSA_AMD_EXT_API_TABLE_STEP_VERSION 0x0A (#3823) * add new hsa APIs to rocprofiler-sdk table version 0x0A. * add fmt support for new hsa api. * Copilot suggestion. * take _val by const reference. * use internal fmt for ci builds. * address comments: remove default in hsa_amd_memory_copy_op_type_t and reorganize hsa_amd_memory_copy_op_t. * add pip install for cmake 3.24. required by rocprofiler-register. * also install cmake 3.24 for sanitizers. * bump patch version to 1. * add local path to GITHUB_ENV. --------- Co-authored-by: Venkateshwar Reddy Kandula <venkateshwar.kandula1306@gmail.com> * clr: Restructure hsa_amd_memory_async_batch_copy (#3824) - Remove Indirect function at HIP level - Expand hsa_amd_memory_copy_op_t struct * [rocprofiler-compute] Pin gersemi version in github workflow to match that of pre commit (#3875) * Pin gersemi version to match that of pre commit hook * Fix formatting * Async Logging Implementation to improve runtime logging performance (#3666) * clr: Add async logging support * Add a missing file. * Remove the lock * Fix ring buffer logic * reduce host side flush interval to 1ms * port over the exception handling from Saleel's branch * limit the host thread drain logic(FlushAsyncLogsInCurrentThread) to be windows only to fix a hang issue And call TruncateLogFileFlushPath() in async path only to replace truncate_log_file(), so that freopen call in the origin func is avoided, in order to resolve the crash by race condition * clean up the format for debug files so it has 100 symbols per line * reduce message size to 2048 and keep Windows path the same as linux for TruncateLogFileFlushPath * pick a different implementation of TruncateLogFileFlushPath, so that it should handle Windows as well * quick Windows build fix and remove unused line * remove the sleep in worker thread * address PR comment: 1. use static AsyncLogger instance; and disable async logging in destruct path 2. use const char * for source file pointer * switch message size back to 4096, as current implementation is still using string. --------- Co-authored-by: German Andryeyev <German.Andryeyev@amd.com> Co-authored-by: German Andryeyev <56892148+gandryey@users.noreply.github.com> * rocr: fix hsaKmtGetNodeWallclockFrequency ret status (#3558) Signed-off-by: Flora Cui <flora.cui@amd.com> * Combine surface vector types into single test instance (#2593) * [rocprofiler-systems] Fix python bindings for manylinux builds (#3859) * Prevent bug causing build failure when multiple python versions are set * Prevent requiring libpython * Unset python cache when ROCPROFSYS_BUILD_FOR_THEROCK=ON * Remove unused vars and update comment * Apply review suggestion * Apply review suggestion * Revert "clr: Restructure hsa_amd_memory_async_batch_copy (#3824)" (#3907) This reverts commit f68eacbac0fba6aa271df3f8227c4b3c050d0cd2. * [rocprofiler-sdk] Add --copy in rocpd package generation tests (#3815) * hip-tests: Improve rtc compiler option test (#3343) * hip-tests: Improve rtc compiler option test * Cache parsing of config file for faster runtime * Remove kernel launch when only IR data is required * Cleanups based in copilot review * Remove powf(2,-126) case to avoid failures * Minor fixes based on copilot --------- Co-authored-by: Rahul Manocha <rmanocha@amd.com> Co-authored-by: ammallya <ameyakeshava.mallya@amd.com> * Disable negative tests (#3896) * Disable negative tests * Add multiproc tests for ASAN Skip * Disable Examples by default (#3920) * [rocSHMEM] Disable RDMA NIC firmware check (#3738) * disable firmware check * apply robot's feedback * Apply PR feedback * typo * Update projects/rocshmem/docs/api/env_variables.rst typo in doc --------- Co-authored-by: Aurelien Bouteiller <abouteil@amd.com> * enable ROCclr githash in hip (#3758) * AIRUNTIME-172: Pass HIP flags to device-only builds (#3669) - Parse CMAKE_HIP_FLAGS and per-config flags into HIP_DEVICE_BUILD_FLAGS_NO_ARCH, filtering --offload-arch= - Propagate parsed flags to all custom device-only compilation commands across 16 CMakeLists.txt in catch tests and samples - Guard flag extraction for empty build types * rocrtst: AddressSanitizer fix series - fixes hsa_memory_free mismatches (#3660) * rocrtst: Address sanitizer fix series: skip negative alloc tests under ASAN (#3788) * rocrtst: ASAN fixes for memory_basic and memory_access tests (#3852) * Bumping TheRock commit hash to include rocrtst fix (#3904) * Bumping TheRock commit hash to include rocrtst fix * Adding the change to build tools checkout * AIRUNTIME-84 - Remove MT path from the ROCR's backend (#3481) * clr: Remove direct dispatch from ROCR path MT path for ROCR isn't tested and can be broken in Linux, and in Windows it doesn't work at all. * Update projects/clr/rocclr/device/rocm/rocvirtual.cpp Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> * Apply suggestion from @Copilot Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> * Add copilot suggestions --------- Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> Co-authored-by: Godavarthy Surya, Anusha <agodavar@amd.com> * [RCCL] Implicit launch order test for MPI unit test. (#3822) * Added ImplicitLaunchOrder Test * Added multiple communicators to implcit order test * using base class create communicator function * Added Implicit launch order test case to test_runner * Used MPI test framerwork APIs - to handle resources - used macros to get rid of if/else conditions - removed redundant hipDeviceSynchronize and MPI_Barrier() --------- Co-authored-by: Atul Kulkarni <atul.kulkarni@amd.com> * [ci] Adding CI nightly to rocm-systems (#3840) * Adding CI nightly * Adding testing comment * removing test --------- Co-authored-by: RahulC <168588218+rahulc-gh@users.noreply.github.com> * [RCCL] Tune message thresholds and nchannels for PAT algorithm (#3756) * [RCCL] Tune message thresholds and nchannels for PAT algorithm Tune PAT algorithm selection for RS and AG in 1ppn multi-node configurations: 1. Add message size thresholds for PAT selection based on node count: * 4 nodes: 512KB - 2MB * 8 nodes: 32KB - 4MB * >= 16 nodes: up to 32MB 2. Add PAT specific channel tuning: * <= 32KB msg size: 1ch * <= 64KB msg size: 2ch * <= 1MB msg size: 4ch * > 1MB msg size: 8ch 3. Skip channelThresholds tuning table for 1ppn to fix non-monotonic channel count in 1ppn RS * Address review requests * [ROCProfiler-SDK][ATT] Other SIMD data generation (#1350) * [rocprofiler-systems] Deprecation note for rocprof-sys user API (#3873) The rocprof-sys user API is superseded by ROCTx Markers from rocprofiler-sdk. This PR marks the user API as deprecated across documentation and runtime to guide users toward the modern replacement before eventual removal. * fix: resolve CPU device path lookup in CLI query-device and library (#3861) Authored-by: Koushik Billakanti <kbillaka@amd.com> * Add rocm-bootstrap package (#3897) * Add rocm-bootstrap package: GPU detection and target hierarchy Pure-Python package providing AMD GPU detection via Linux sysfs (no ROCm installation required), a canonical 3-level GFX target hierarchy (family/sub-family/target), dist-safe and module-safe naming APIs, and a WheelNext variant plugin entry point. Key components: - targets.py: 33 GFX targets across 9 sub-families and 4 families, with LLVM generic ISA and xnack metadata from GCNProcessors.td/AMDGPU.td - naming.py: Deterministic name generation for pip/wheel and Python module contexts, fully validated and tested for every bundle - detect.py: GPU detection via KFD topology and ip_discovery sysfs with env var overrides. CLI: rocm-bootstrap-detect --unique/--verbose/--hierarchy - _platform.py: All fakeable I/O in one module for test monkeypatching - variant_provider.py: WheelNext variant_plugins entry point reporting detected leaf targets (wheel rebuilder owns hierarchy fan-out) 671 tests covering all targets, naming contracts, detection paths, and CLI. Published to PyPI under AMD-OSS org to reserve the package name. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Remove ip_discovery fallback from GPU detection The ip_discovery path (DRM sysfs GC IP version) can misidentify GPUs because GC IP block versions don't always map 1:1 to gfx_target_version. KFD topology is the reliable source — it's what rocminfo uses internally. Remove the fallback and only detect via KFD. If KFD isn't loaded, no GPUs are reported. A PCI device ID fallback can be added later if needed. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> --------- Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com> * Revert "hip-tests: Improve rtc compiler option test (#3343)" (#3942) This reverts commit 6053278c6ec46e2d716dd2731554f4925f5265cf. * [hip-tests] ROCM-2280 - fix Unit_Coalesced_Group_Tiled_Partition_Sync… (#3793) * [hip-tests] ROCM-2280 - fix Unit_Coalesced_Group_Tiled_Partition_Sync_Positive_Basic Signed-off-by: Sebastian Luzynski <sebastian.luzynski@amd.com> * [rocprofiler-compute] Added unit tests for data imputation (#3270) * Added unit tests for data imputation * Added test for incorrect and corner cases * Removed unnecessary import * Review changes * Move tests * Add new test in ctest * Split tests * [rocprofiler-compute] Add workload-level aggregation to analysis database (#3427) * Add workload-level aggregation to analysis database Database Schema Changes: - Increment analysis database schema version to 1.3.0 - Rename kernel-scoped tables for clarity: MetricValue → KernelMetricValue, RooflineData → KernelRooflineData - Add workload-level tables: WorkloadMetricValue and WorkloadRooflineData for aggregate metrics across all dispatches - Add workload_metric_view database view for querying workload-level metrics Metrics Calculation Refactoring: - Refactor calc_expressions() to return both kernel-level and workload-level metric values - Refactor calc_roofline_data() to calculate roofline data at both kernel and workload granularity - Optimize calc_metrics_data() to generate metric expression templates once instead of duplicating per kernel - Optimize database creation by pre-grouping kernel data Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com> * Address review comments * Breakdown run_analysis into sub-functions * Fix test case --------- Co-authored-by: Claude Opus 4.5 <noreply@anthropic.com> * [rocprofiler-compute] Add statistics to torch operators listing (#3201) * Refactoring to place all processing in Analyze mode. * Update projects/rocprofiler-compute/src/rocprof_compute_analyze/analysis_base.py Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> * Ruff * Ruff * Type correction * Ruff * Added cli flags, hierarchial display. * Error message indicating the order in which flags should be used. * removing list_torch_operators() from "list and exit" flags. * ruff fixes * Ruff fixes * Removing spaces * Update projects/rocprofiler-compute/src/utils/tty.py Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> * Display corrections * ruff * Ruff fixes. * Showing context info and correlation ids after the hierarchy tree. * Logging format fixed. * fstring * [AIPROFCOMP-154] Refactor marker naming to include full pytorch operator hierarchy (#3011) * Change marker naming to include full hierarchy * Add C++ ops, domain specific libraries * Wrapping all class methods for common modules and torch submodules. * Adding instrument_all_torch_ops() as torch._C._dispatch_call is not used for torch versions >= 2.9 * Adding a check for torch._C._dispatch_call * Using a context stack to keep marker names intact * Ruff * Fixing sanitized_operator_name for the csv filename. * Trigger CI * Correcting the docs * Modified profiling test and added a message for roctx import failure. * Added analysis tests * Ruff fixes * logging correction * logging correction * Update projects/rocprofiler-compute/src/utils/inject_roctx.py Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> * Use thread-local stacks for ROCTX markers; update profiling docs * Ruff * Update projects/rocprofiler-compute/tests/test_profile_general.py Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> * Correcting the usage of thread local stacks * Moved hierarchy to "--list-torch-operators" * Format fix * Removed kernel name shortening * Modified display. * Updated --torch-operator option. * Allosing --torch-operator to take multiple, comma/space separated values. * ruff and table formating * Removing profiling from analyze pytests. * Fix pytests * pytest fixed * Removing differing columns (between rocpd and csv runs) * Ruff fixes * Further removal of differing columns. * Further column correction * Profile data is cached and accessed by analyze pytest. * Ruff fix. * Fixed pytest * Normalize column name - Correlation_ID, to handle case inconsistencies --------- Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> * Refactor profiling tests and improve logging consistency * 1) Retaining workload counter_collection and marker_trace files after creating torch_trace directory 2) Remove previous torch_trace directory each time list-torch-operators is used * Workload python version display and improved error message. * Updating docs. * Display correction * Ruff * Addressing PR reviews. * Update projects/rocprofiler-compute/src/utils/inject_roctx.py Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> * Rebase and conflict resolution. * Delete projects/rocprofiler-compute/docs/stress-test-validation-plan.md * Modifying --list-torch-operators show kernel IDs and default to short names * COrrecting docs regarding the retention of profiling data on running --list-torch-operators * Removing dead code. * Removing experimental wording * Update projects/rocprofiler-compute/src/utils/tty.py Co-authored-by: xuchen-amd <xuchen@amd.com> * Using the same kernel_shortening logic as in default analyze tables. * Making extract_kernel_name an additional step * Fix comments. * Indentation fix. * Ruff fix. * Fix spacing in comments. * Resolving comments, reordering operators according to duration. * removing executables from tracking * Undoing changes in gitignore. * ruff * Update projects/rocprofiler-compute/src/rocprof_compute_analyze/analysis_base.py Co-authored-by: xuchen-amd <xuchen@amd.com> * Extracting duplicated code into a function * Adding early checks and continues while parsing context_id. * Refactoring the if else blocks. * Removed unnecessary else block. * Move sys.exit out of the funtion list_torch_operators. * Removed outdated variable - filter_torch_operators * Ruff * Fixing kernel_name_to_id, by reading operator csv files and not relying on pmc_kernel_top.csv * Ruff * Cleaner comments * Moving all helper functions to utils.py * Remove duplicate work, caching kernel_to_id mapping for display. * Removed dead code. * Comment correction and dead code removal. * Remove redundant code. * Added checks for cli output from --list-torch-operators * Fixing the sorting order in the assertion * RUff * Comments explaining the checks * Comment correction * Removed redundant comments. * Update projects/rocprofiler-compute/src/rocprof_compute_analyze/analysis_cli.py Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> * Ruff * Removing dependency on maxcolwidths --------- Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> Co-authored-by: xuchen-amd <xuchen@amd.com> * [rocprofiler-compute][bugfix] fix duplicated generate_machine_specs() call (#3534) * fix duplicated generate_machine_specs() call * Update projects/rocprofiler-compute/src/rocprof_compute_base.py Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> --------- Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> * [rocprofiler-compute] Update installation section of documentation (#3618) * Update installation section of documentation * Add newly added cmake configuration arguments * Remove mention of rocprofiler-compute specific tarball release in favor of cloing source code * Remove package manager based installation which will no longer be updated with latest version of rocprofiler-compute, instead this will be released by TheRock build system * Improve build steps to add clarity * Clarify the dependency on rocprofiler-sdk library and optional dependence on rocprofv3 script * Add tarball and TheRock based install instructions * Fix formatting * Address review comments * [rocprofiler-compute][bugfix][TUI] fix data overwrite and add tui tests (#3638) * fix data overwrite and add tui tests * address comments and add test file --------- Co-authored-by: ywang103-amd <ywang103@amd.com> Co-authored-by: vedithal-amd <Vignesh.Edithal@amd.com> * [rocprofiler-compute] Optimize data imputation for iteration multiplexing (#3771) * Optimize data imputation for iteration multiplexing Problem: * The data imputation used nested python loops resulting in O(num_unique_kernels * num_subgroups * num_counters) time complexity * This made data imputation prohibitively slow for workloads with large number of dispatches and uniquer kernels * Data imputation took 12 minutes for workload with 12k dispatches and 9 unique kernels! Solution: * Eliminate per counter and per subgroup loop to reduce time complexity to O(num_unique_kernels) * Instead use chained call groupby().bfill().ffill() which levearages pandas vectorize CPython routines * Logic is preserved since for a given dispatch group and counter, only one valid non-null value exists * Last incomplete dispatch group edge case handling is also preserved Performance Impact: * Optimized: 0.0139 seconds per group * Baseline: 3.57 seconds per group * Speedup: ~260x Testing: * All existing unit tests pass without modification * Logic correctness verified via test_impute_counters_iteration_multiplex Logging: * Added debug logging to indicate data imputation is being performed on X dispatches across Y kernels * Address review comments * Remove special handling of imputing last subgroup * The bfill+ffill pattern naturally handles both complete and incomplete subgroups * Remove redundant double indexing when calculating cols_frozenset * Add more test cases for data imputation algorithm * [rocprofiler-compute] Remove csv output format fallback for torch tracing by using correct correlation id for rocpd format (#3783) * Use stack_id in liue of correlation_id for rocpd * Doc update * Added tests * Update CMake * [rocprofiler-compute] Fail dynamic attachment test quickly in case of hangs (#3819) * Reduce timeout for test_profile_live_attach_detach from 1800 to 300 for quicker fail in case of hangs * Apply suggestion from @Copilot Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> * Fix cmake formatting --------- Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> * [rocprofiler-compute] Remove hip-trace from rocprofiler-compute profiling (#3833) * Removed hip trace * Update profiling_config for tests * Removed hip-trace tests * Update CHANGELOG --------- Co-authored-by: Vignesh Edithal <Vignesh.Edithal@amd.com> * [rocprofiler-compute][tui] improve button response (#3640) * improve button response * add unit tests * address comments * Update projects/rocprofiler-compute/src/rocprof_compute_tui/widgets/instant_button.py Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> --------- Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> * [rocprofiler-compute][bugfix] fix filter for node id and gpu id (#3648) * fix filter for node id and gpu id * use helper function * add informative exit message * fix formatting * add test cases * [rocprofiler-compute] Handle rocflop child process crashes gracefully - reduce vram usage of rocflop based tests for robustness (#3564) * Extract pipe-read error handling into read_records_from_pipe() - Isolate pipe-read logic into a dedicated function for better structure - Report actual errno via std::strerror() when read() fails - Warn on partial/incomplete reads (bytes_read % sizeof(Result) != 0) - Default to empty results when read() returns error or zero bytes - Use std::cout for error messages for consistency with rest of file Made-with: Cursor * Adding device ids to error messages. Making rocflop exit non-zero when any child process fails. * Adding MI300A CPX/NPS1 specific HBM bandwidth value to the expected values. * Confining rocflop to the lightest workload. * [rocprofiler-compute] Roofline: Use gfx950 builtins for MFMA (#3886) * Use gfx950 builtin for MFMA FP16 * Use gfx950 builtin for MFMA BF16 * Use gfx950 builtin for MFMA I8 * Fix comments * Update copyright * Update CHANGELOG * Fix formatting * Review comments * Clear CHANGELOG indicating we resolved roofline peaks for MI350 * Fix typo in pre-processor guard preventing roofline from running on MI300 * Ruff formatting * Fix uninitialized variables --------- Co-authored-by: benrichard-amd <ben.richard@amd.com> * [rocprofiler-compute] Fix test coverage path handling for install mode (#3624) * Fix test coverage path handling for install mode Fix coverage report generation when TEST_FROM_INSTALL is enabled by dynamically setting the coverage output path based on test mode. Also fix conftest.py to use the correct rocprof-compute script path that matches the successful module import location. Changes: - Remove unused CTEST_COVERAGE_EXTRA_FLAGS variable - Add COVERAGE_XML_PATH variable that points to install dir when TEST_FROM_INSTALL=ON, or build dir otherwise - Update conftest.py to store and reuse the script path that successfully loaded the module, preventing path mismatches in subprocess execution * Pin clang-format version in formatting worfklow * Formatting changes for cmake * Fix ruff formatting (#3932) * [rocprofiler-sdk] [Documentation] correcting counter collection example (#3947) * correcting multi-pass counter collection example * Update PMC parameters in using-rocprofv3.rst * Modify PMC job parameters and profiling pass details Updated PMC job configuration and corrected profiling pass description. * Modify rocprofv3 command example for clarity Updated the command example to include 'SQ_WAVE_CYCLES' in the profiling passes. * Fix graph dependency logic broken by PR #3481 (#3946) PR https://github.com/ROCm/rocm-systems/pull/3481 moved the DepHwEvents barrier logic inside the profilingBegin for-loop scope, causing external dependency signals to be skipped when the loop condition is not entered. This breaks graph node dependencies in the new scheduling path. Move the DepHwEvents iteration and clearDepHwEvents() back outside the for-loop scope so dependency barriers are always applied. * memory: fix "contiguous_bytes" calculation in generic conversion (#3779) * [rocshmem]: CI: disable tests known to fail due to AIROCSHMEM-316 (#3893) * CI: disable tests known to fail due to AIROCSHMEM-316 * We need to modify jenkins config to subcase for gda-mlx, disable for all gda for now * Also match gda-mlx5 for when the jenkins script gets updated * [rocdecode][samples] - fix rocdecdecode sample issue in SLES (#3958) * [rocDecode] Fix seg faults with batch sample for mem_type=2 (#3845) * fix seg faults with mem type =2 in batch sample * clean up required for host backend * clean up * address copilot comments * address review comments * [rocprofiler-compute][bugfix] fix config key access for iteration multiplexing (#3892) * [TheRock CI] Add custom container support for rocgdb tests (#3926) Add per-component container_image and container_options support to therock-test-component.yml, matching TheRock upstream convention, and enable rocgdb in the debug_tools and all test project lists now that the required custom container (no_rocm_image_ubuntu24_04_rocgdb with --cap-add=SYS_PTRACE) can be specified. Co-authored-by: Claude Sonnet 4.6 <noreply@anthropic.com> * add dll details to hiprtc, hiprtc-builtins (#3620) * add dll details to hiprtc, hiprtc-builtins * update copy rights year to 2026 for all 3 dlls * [ci] Build therock-archives and therock-dist together (#3956) ## Motivation Incorporating https://github.com/ROCm/TheRock/pull/1523 into these forked workflows. This should fix "report" step errors like `ls: cannot access 'TheRock/build/artifacts/*.tar.xz': No such file or directory` on https://github.com/ROCm/rocm-systems/actions/runs/22652232146/job/65691391937?pr=1705 when the build fails, leading to more clear error attribution. ## Technical Details * Also adding `-k 0` to keep going through errors * We originally split these steps because archive generation was very slow on Windows (on the order of 1h for a 1h build). * The "report" step should still be moved to a Python script that gracefully handles missing directories/files. * Eventually these workflows will be de-dup'd with those in TheRock ## JIRA ID N/A ## Test Plan * CI on this PR should succeed, check the logs for archive generation and step timing ## Test Result CI passed as expected. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. * [rocprofiler-systems] Decode OMPT "flags" argument into human-readable form (#3494) * Transform "flags" arguments for certain OMPT callbacks into human readable form * clr: Eliminate hsa-tracing workaround for profiler (#3895) * hip-tests: Improve rtc compiler option test (#3944) * hip-tests: Improve rtc compiler option test * Cache parsing of config file for faster runtime * Remove kernel launch when only IR data is required * Cleanups based in copilot review * Remove powf(2,-126) case to avoid failures * Minor fixes based on copilot * hip-tests: Fix ComplrOptn test failure in CI --------- Co-authored-by: Rahul Manocha <rmanocha@amd.com> Co-authored-by: ammallya <ameyakeshava.mallya@amd.com> * [TheRock CI] Enable compilation for rocshmem PRs (#3955) * [rocprofiler-systems] Fix `transpose` runtime-instrument crash (#3949) Add `libclang-cpp.so` and `libLLVM.so` to the set of libraries that are not instrumented by default * [LRT][HIP] update readme in clr (#3664) * Add fast-path fatbin detection and fix PHDR in-place write regression ELF surgery: - Add ElfSurgery.has_fatbin_section() static method for lightweight .hip_fatbin detection by scanning section headers only - Fix PHDR in-place write: when spare PT_NULL slots exist (common after patchelf), write new segment in-place instead of failing. Prior code only handled the resize-and-shift path. - Move test helpers' inline imports to top-level in test_surgery.py COFF surgery: - Add CoffSurgery.has_fatbin_section() static method for .hip_fat section detection (8-char PE section name limit) kpack_transform: - Add is_fat_binary() cross-format dispatcher using the new fast-path static methods instead of loading full binary - Remove elf_has_hip_fatbin() from wheel_splitter (now uses is_fat_binary) binutils/ccob_parser: - Remove catch-all exception handlers that silently swallowed errors - Let subprocess and parsing failures propagate naturally Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Add database handlers for wheel splitting with arch-aware file detection New plugin architecture for detecting architecture-specific kernel database files during wheel splitting. Each handler knows its library's file layout and extracts the GPU architecture from filenames/paths. Handlers: - rocBLAS: lib/rocblas/library/*_gfx*.{co,hsaco,dat} - hipBLASLt: lib/hipblaslt/library/*_gfx*.{co,hsaco,dat} - hipSPARSELt: lib/hipsparselt/library/*_gfx*.{co,hsaco,dat} - AOTriton: lib/aotriton.images/amd-gfx*/... with bundle key mapping (gfx11xx->gfx11, gfx120x->gfx12_0). Updated to match actual torch wheel layout rather than abstract spec. - MIOpen: share/miopen/db/gfx*.{db.txt,fdb.txt,model} with dedicated regex for concatenated arch+CU filenames (e.g., gfx90878 = gfx908 + 78 CUs). Interim handler — MIOpen database format will be adapted for better splitting support soon. Base class enforces path validation contract: detect() raises ValueError with clear message if path is not under prefix_root, rather than letting incidental errors from relative_to() propagate. Wheel-type presets map preset names (e.g., "torch-fat") to the relevant set of handlers. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Add Python wheel splitter for host/device separation Core wheel splitting implementation that takes a fat torch wheel and produces: - Host wheel: shared libraries with device code zeroed out, plus a kpack overlay directory containing the kpack manifest - Device wheels: per-architecture kernel databases and kpack bundles - Variant wheel (optional): PEP 817 variant metadata for wheelnext-aware installers Key components: - WheelSplitter class orchestrating the full split pipeline - Fat binary scanning via is_fat_binary() fast path - Parallel kpack_offload_binary() transforms with ProcessPoolExecutor - Database handler integration for arch-specific file routing - RECORD/METADATA rewriting for valid wheel output - TransformArgs dataclass for clean parallel dispatch CLI tool (split_python_wheels.py) with ModuleNotFoundError guard for the rocm-bootstrap dependency. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Format C++ kpack runtime ISA target matching code Formatting-only changes to isa_target_match.h and its test file. No functional changes. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Specify THEROCK_ROCM_SYSTEMS_SOURCE_DIR (#3968) THEROCK_ROCM_SYSTEMS_SOURCE_DIR needs to be set to test the checked out rocm-systems repo instead of the rocm-systems submodule pulled via TheRock. Furthermore, don't fetch rocm-systems submodule via TheRock. * [rocprofiler-sdk] Fix rocprofv3 summary validation (#3739) ## Motivation <!-- Explain the purpose of this PR and the goals it aims to achieve. --> Fix failing PSDB job ## Technical Details <!-- Explain the changes along with any relevant GitHub links. --> https://github.com/ROCm/rocm-systems/pull/1881 accounted for hidden call to `hipGetDeviceProperties*` in the HIP summary check but did not account for it in `get_dims(total)` and `get_dims(hip_and_marker)`. ## JIRA ID <!-- If applicable, mention the JIRA ID resolved by this PR (Example: Resolves SWDEV-12345). --> <!-- Do not post any JIRA links here. --> N/A ## Test Plan <!-- Explain any relevant testing done to verify this PR. --> N/A. This is a testing fix. ## Test Result <!-- Briefly summarize test outcomes. --> N/A. This is a testing fix. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. * Fix VA memory leaks in graph memalloc nodes (VM mode) (#3878) Move vaddr_sub_obj release into submitVirtualMap (single release point) and remove redundant releases from all callers: - VmHeap::UncommitMemory - VirtualMemFreeNode::submit - ihipFree graph VM path - hipMemUnmap - ReleaseVaRefs (removed along with va_refs_ tracking) Leaks fixed: - vaddr_sub_obj leaked in auto-free graph path (no free node to release) - va_ leaked in auto-free graph path (parent ref never dropped) - vaddr_sub_obj leaked in non-graph hipFree with use_vm_heap_ enabled - vaddr_sub_obj leaked during hipMemPoolTrimTo VmHeap trim path" ## Motivation <!-- Explain the purpose of this PR and the goals it aims to achieve. --> Fix leaks in Graph VM path for ASAN reported leaks. ## Technical Details <!-- Explain the changes along with any relevant GitHub links. --> ## JIRA ID <!-- If applicable, mention the JIRA ID resolved by this PR (Example: Resolves SWDEV-12345). --> <!-- Do not post any JIRA links here. --> ## Test Plan <!-- Explain any relevant testing done to verify this PR. --> ## Test Result <!-- Briefly summarize test outcomes. --> ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. * wsl/librocdxg: Add support for ROCgdb (#3888) Add ROCm debugger (ROCgdb) support in hsakmt/dxg. - include/impl: extend WDDM headers for debugger use - debug.cpp, queues.cpp: add or adjust debugger-related handling - wkmi: interface and library updates * Make TRACE_ENTRY/TRACE_EXIT print time the function took This commit makes dbgapi function tracing include the time each call took. This is very useful to diagnose performance problems. In addition, it prints the name of the function on the exit side too, so that it's much easier to grep for the time of the call you are interested in. For example, before: [amd-dbgapi-lib] amd_dbgapi_read_memory (process_id=process_1, wave_id=wave_1, lane_id=0, address_space_id=address_space_70 <private_lane>, segment_address=0x8, value_size=4@0x7ffd4e7cdd68, value=0x63d1cd5ca580) { [amd-dbgapi-lib] } = STATUS_SUCCESS, *value_size=4, *value=[0x1,0x0,0x0,0x0] After: [amd-dbgapi-lib] amd_dbgapi_read_memory (process_id=process_1, wave_id=wave_1, lane_id=0, address_space_id=address_space_70 <private_lane>, segment_address=0x8, value_size=4@0x7ffc5f20bdf8, value=0x5ed7de1eb8b0) { [amd-dbgapi-lib] } amd_dbgapi_read_memory = STATUS_SUCCESS, *value_size=4, *value=[0x1,0x0,0x0,0x0] (took 0.071165 ms) Printing the name at the end also makes it much easier to find the matching exit brace of a call's entry, especially if there are nested calls in between, making the end brace appear much later in the log. I considered making printing the time optional with a new CMake option, to make it less noisy to e.g., diff two logs, but decided against it since it's trivial to sed away the "(took ... ms)" bit before comparing logs. Change-Id: I6e8f56dc936f089fafebeeb8eb0970beccd0cc91 commit-id:5b8454b2 * Add missing hipFree in test and hide check of mem allocate behind ASAN (#3988) * [rocprof-sys-run] Remove ROCPROFSYS_USE_ROCM flag (#3601) **BREAKING CHANGE**: ROCPROFSYS_USE_ROCM CMake option and runtime environment variable have been removed. ROCm support cannot be disabled. ## Motivation ROCm support is now always enabled. Since rocprofiler-systems is a profiler specifically designed for ROCm systems, the ability to disable ROCm support is unnecessary. ## Technical Details - Remove ROCPROFSYS_USE_ROCM CMake option and compile definition - Remove all #if ROCPROFSYS_USE_ROCM > 0 preprocessor guards - Remove ROCPROFSYS_USE_ROCM runtime config setting and env var - Simplify build-release.sh from 4 variants to 2 (--base, --mpi) - Update documentation to reflect ROCm is always enabled - Update tests to remove ROCPROFSYS_USE_ROCM=OFF configurations ## JIRA ID - AIPROFSYST-199 - AIPROFSYST-233 ## Test Plan <!-- Explain any relevant testing done to verify this PR. --> Workflow Build and Tests run. ## Test Result <!-- Briefly summarize test outcomes. --> ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Co-authored-by: David Galiffi <David.Galiffi@amd.com> * rocr: Remove ldrm_bo from AllocationRegion (#3772) ## Motivation Removal of AllocationRegion::ldrm_bo as it is not used. ## Technical Details It seems that ldrm_bo is no longer needed since amdgpu calls have moved out of Runtime. ## JIRA ID NA ## Test Plan Run roctst. ## Test Result Success. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> * SWDEV-567112 - Introduce new mechanism for tagging and disabling tests - Part 4 (#3870) ## Motivation Since the number of unit test cases is constantly growing, the execution time is also increasing. Specific tests are suited for non standard testing and there is no need to run them in the regular PSDB. Add a mechanism to easily disable/enable tests and easily categorize them. Part 4 covers test name changes for following groups: - graph tests - hip_specific tests - kernel tests - launchBounds tests - library tests - math tests - memory tests ## Technical Details Explained in the ticket in the document attached. ## JIRA ID SWDEV-567112 ## Test Plan Ensure build passes. All enabled tests pass, all disabled tests are not run ## Test Result Local build passed, local enabled tests pass, disabled are not run ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. * Add config header generation to standalone test build readme (#3795) ## Motivation Explain in the README how can standalone tests be compiled by manually generating hip_tests_config.hh ## Technical Details None ## JIRA ID None ## Test Plan Build should pass. Standalone tests can be compiled using the instructions ## Test Result Local build passes. Standalone tests can be compiled locally ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. * Drop env file (#3691) * Fix KFDTest KFDDBGTest.HitTrapEvent to support GFX12 (#3883) ## Motivation Fix KFDTest KFDDBGTest.HitTrapEvent to support GFX12 ## Technical Details <!-- Explain the changes along with any relevant GitHub links. --> ## JIRA ID ROCM-19942 ## Test Plan Run the test on GFX12 device to see if the issue resolved. ## Test Result Passed ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. * [pal] Set hasExpertSchedMode_ for gfx12 (#3153) ## Motivation Aligning PAL with https://github.com/ROCm/rocm-systems/pull/2435 Full background: https://github.com/ROCm/rocm-libraries/issues/3211 ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. * Fix ring buffer indexing in AIE AQL queue SubmitPackets (#3619) * rocr: Optimize SDMA multi-producer hot path for reduced contention and overhead (#3264) ## Motivation Time taken by hot path ## Technical Details - Reduce PadRingToEnd lock contention: unlock reservation_lock_ around UpdateWriteAndDoorbellRegister spin-wait so other producers can reserve while NOP padding commits. - Use mwaitx for ordered-commit spin in UpdateWriteAndDoorbellRegister instead of sched_yield(), with fallback when mwaitx is unavailable. - Cache immutable flags and MMIO pointers at init - Replace os::YieldThread with _mm_pause in CAS retry loops - Eliminate redundant atomic Load in AcquireWriteAddress - Stack-allocate single-packet command buffers in SubmitLinearCopyCommand and SubmitLinearFillCommand to avoid per-call heap allocation. <!-- Explain the changes along with any relevant GitHub links. --> ## JIRA ID <!-- If applicable, mention the JIRA ID resolved by this PR (Example: Resolves SWDEV-12345). --> <!-- Do not post any JIRA links here. --> ## Test Plan <!-- Explain any relevant testing done to verify this PR. --> ## Test Result <!-- Briefly summarize test outcomes. --> ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. * clr: Print possible pipeId during dispatch for logs (#3850) ## Motivation <!-- Explain the purpose of this PR and the goals it aims to achieve. --> ## Technical Details <!-- Explain the changes along with any relevant GitHub links. --> ## JIRA ID <!-- If applicable, mention the JIRA ID resolved by this PR (Example: Resolves SWDEV-12345). --> <!-- Do not post any JIRA links here. --> ## Test Plan <!-- Explain any relevant testing done to verify this PR. --> ## Test Result <!-- Briefly summarize test outcomes. --> ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. * ROCM-1580 - Select Monitor Type at compile time (#2919) ## Motivation We currently select our monitor type dynamically at runtime, using a constructor and storing the object in an std::variant. From profiling RCCL benchmarks, this level of indirection is measurable and does contribute to latency. For a scoped lock, unlock loop it is around 2-5% of the total cost. Additionally we had functions which were only supported for std::mutex and not recursive mutex, these could be called, and were protected with debug asserts. Moving to compile time prevents these invalid functions from being called. I.e. this change is for both performance and safety. ## Technical Details The std::variant inside Monitor has been removed, and replaced with two separate classes. std::Monitor and std::RecursiveMonitor, the type is now completely defined at compile-time. In practice this doesn't change the thread safety behaviour as we were defining the variant type at compile-time anyway. Mutex types are kept the same, i.e. if they were constructed with `true` then they are RecursiveMonitors Performance measurements were taken for the scoped lock unlock. All mutex calls with this scoped lock-unlock in a kernel submission benchmark were measured, and averaged. <html xmlns:v="urn:schemas-microsoft-com:vml" xmlns:o="urn:schemas-microsoft-com:office:office" xmlns:x="urn:schemas-microsoft-com:office:excel" xmlns="http://www.w3.org/TR/REC-html40"> <head> <meta name=ProgId content=Excel.Sheet> <meta name=Generator content="Microsoft Excel 15"> <link id=Main-File rel=Main-File href="file:///C:/Users/abeltons/AppData/Local/Temp/msohtmlclip1/01/clip.htm"> <link rel=File-List href="file:///C:/Users/abeltons/AppData/Local/Temp/msohtmlclip1/01/clip_filelist.xml"> <!--table {mso-displayed-decimal-separator:"\."; mso-displayed-thousand-separator:"\,";} @page {margin:.75in .7in .75in .7in; mso-header-margin:.3in; mso-footer-margin:.3in;} tr {mso-height-source:auto;} col {mso-width-source:auto;} br {mso-data-placement:same-cell;} td {padding-top:1px; padding-right:1px; padding-left:1px; mso-ignore:padding; color:black; font-size:11.0pt; font-weight:400; font-style:normal; text-decoration:none; font-family:Arial, sans-serif; mso-font-charset:0; mso-number-format:General; text-align:general; vertical-align:bottom; border:none; mso-background-source:auto; mso-pattern:auto; mso-protection:locked visible; white-space:nowrap; mso-rotate:0;} .xl65 {mso-number-format:Percent;} .xl66 {text-align:center;} --> </head> <body link="#467886" vlink="#96607D"> | Variant | Template | -- | -- | -- | -- | Average Scope Time (ns) | Average Scope Time (ns) | %Speedup Warm Cache | 1247.85 | 1180.5 | 5.40% Cold Cache | 2098.09 | 2060.47 | 1.80% </body> </html> ## JIRA ID Resolves - ROCM-1580 ## Test Plan Our existing multi-threading tests are sufficient to test this change ## Test Result Local tests pass ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. * Fix passing CMake options to RCCL workflows (#3991) * [rocprofiler-sdk][ATT] Output shaderdata from decoder callback into RCV directory (#2671) ## Motivation Enable emission of shaderdata (s_ttracedata) records alongside ATT outputs, mirroring other per‑SE JSON artifacts and exposing them via filenames.json. ## Technical Details - Add a shaderdata JSON writer that serializes rocprofiler_thread_trace_decoder_shaderdata_t records to shaderdata_{se}_{idx}.json. - Track shaderdata files per SE in FilenameMgr and include shaderdata_filenames in filenames.json. - Handle ROCPROFILER_THREAD_TRACE_DECODER_RECORD_SHADERDATA callback to capture and write records. ## JIRA ID <!-- If applicable, mention the JIRA ID resolved by this PR (Example: Resolves SWDEV-12345). --> <!-- Do not post any JIRA links here. --> ## Test Plan - Add a new integration test `rocprofv3-test-att-shaderdata` ## Test Result - rocprofiler-sdk CI Jobs ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Co-authored-by: Venkateshwar Reddy Kandula <venkateshwar.kandula1306@gmail.com> * [CLR][NFC] Clearly name HIP & OpenCL set kernel arg code paths (#3460) * Skip some arch based tests for ASAN and fix memory leaks in rtc tests (#3996) * [rocshmem] CI: Disable RO flood testing, tracking issue AIROCSHMEM-324 (#3961) ## Motivation Flood tests do not pass CI, tracking issue AIROCSHMEM-324 ## Technical Details Disable non-passing RO test case ## JIRA ID AIROCSHMEM-296 AIROCSHMEM-324 ## Test Plan CI Jenkins run ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. * [hip-tests] Introduce new mechanism for tagging and disabling tests - Part 5 (#3993) ## Motivation Since the number of unit test cases is constantly growing, the execution time is also increasing. Specific tests are suited for non standard testing and there is no need to run them in the regular PSDB. Add a mechanism to easily disable/enable tests and easily categorize them. Part 5 covers test name changes for following groups: - module tests - multiThread tests - occupancy tests - p2p tests - printf tests - rtc tests - stream tests - streamperthread tests - surface tests - synchronization tests - texture tests - threadfence tests - vector_types tests - virtualMemoryManagement tests - vulkan_interop tests - warp tests ## Technical Details Explained in the ticket in the document attached. ## JIRA ID SWDEV-567112 ## Test Plan Ensure build passes. All enabled tests pass, all disabled tests are not run ## Test Result Local build passed, local enabled tests pass, disabled are not run ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. * Update TheRock SHA in therock-rccl-ci-linux.yml to include rocprofiler MPI dependency fix (#3954) **PR Description** This PR updates the TheRock SHA used in the therock-rccl-ci-linux.yml workflow to include the changes from TheRock PR (https://github.com/ROCm/TheRock/pull/3897) **TheRock commit:** https://github.com/ROCm/TheRock/commit/a5ed8195057dbd8f06888236b1c78ccf41ed326e Updating the SHA pulls in the fix required after rocprofiler introduced an MPI dependency, which previously caused build failures in the RCCL CI workflow. **Validation** A test workflow run was executed and the build was successful: https://github.com/ROCm/rocm-systems/actions/runs/23009699933/job/66816519115 --------- Co-authored-by: Aravind Ravikumar <arravikum@amd.com> * Fix KFDTest SDMAFillDataPacket to support GFX12 (#3969) SDMA header count size calculated based on different fillsize changed in GFX12. KFDTest SDMAFillDataPacket needs to be updated to support the delta ## Motivation Update KFDTest SDMAFillDataPacket to support GFX12 ## Technical Details ## JIRA ID ROCM-20126 ## Test Plan Run the test on GFX12 device ## Test Result Test passed with the change ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. * rocr: Replace repeated 'kAqlProfileLib' probes for each GPU Add one-time cached probe for AQL profile library availability and replace repeated probes in system and per-GPU queries, also move CloseLib call to Unload() and add a check for Valgrind Signed-off-by: Apurv Mishra <Apurv.Mishra@amd.com> * ROCM-3970 - release mem object outside the critical region (#3913) * Update TheRock reference (rocm-systems) - 2026-03-10 (#3930) Automated update due to `rocm-systems` change. Updated TheRock reference to commit: `a77a42113b0e36e8ca01e0a170fb6f735a4d3182` --------- Co-authored-by: therockbot <therockbot@amd.com> * [AMD-SMI] Fix CI failure from stale repo in AlmaLinux8 Docker image (#4015) Add --setopt=skip_if_unavailable=True to the python3-setuptools/wheel dnf install commands so that stale repos in the container image are skipped rather than causing a fatal error. The AlmaLinux8 CI image has a stale amdgpu repo returning HTTP 404, causing dnf metadata refresh to fail before the AMDSMI RPM install is reached. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com> * AIRUNTIME-1893 - Fix __smid for GFX12 (RDNA4) architecture (#3841) ## Motivation Fixes: ROCm/rocm-systems#3495 ## Technical Details Update HW_ID register definitions and __smid() to support GFX12: - Include __GFX12__ in HW_ID=23 register definition - Use s_sendmsg_rtn to retrieve SE_ID on GFX12 (no longer in HW_ID reg) - Add GFX12-specific SE_HW_ID defines for SE_ID size/offset - Refactor __smid() with assemble_smid() helper for consistent ID packing - Restore original MIT license header ## JIRA ID AIRUNTIME-1893 ## Test Plan N/A ## Test Result N/A ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. * Add WSL ABI compatibility check for DXG thunk (#3985) Adds a structure size ABI check in the DXG thunk loader for WSL, calling optional DxgAbiCheck and failing initialization on mismatch while remaining compatible with older libraries. Signed-off-by: Horatio Zhang <Hongkun.Zhang@amd.com> * ROCM-1643 - Fix sample22. (#3491) ## Motivation Let CMAKE find HIP compiler. ## Technical Details Let CMAKE find HIP compiler. ## JIRA ID ROCM-1643 ## Test Plan NA ## Test Result NA ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. * [rocprofiler-systems] Refactor of rocprof-sys-avail tool (#3526) ## Motivation - rocprof-sys-avail links the entire profiler engine (librocprofiler-systems-static) to use only 3 runtime symbols, pulling in ~86 source files and transitive dependencies (perfetto, sqlite3, spdlog, elfutils, etc.) that avail never uses directly. This inflates binary size and couples avail to profiler internals unnecessarily. - Formatting state in avail is spread across global variables, making it difficult to reason about and test the tool's output paths independently. ## Technical Details - Lighten link dependencies: Replace librocprofiler-systems-static with rocprofiler-systems-core as avail's link target. Replace rocprofsys_init_library() with the lightweight set_state(State::Init) + configure_settings(false) pattern (precedent: argparse.cpp). Gate GPU config calls (rocprofiler_sdk::config_settings, amd_smi::config_settings) behind the _init parameter so they are skipped during lightweight init and called explicitly only when --hw-counters is requested. - Standalone HW counter query: Create hw_counter_query.cpp in the avail directory — a self-contained implementation that queries GPU HW counters directly via get_agent_manager_instance() and rocprofiler_iterate_agent_supported_counters() (rocprofiler-sdk C API), replacing the dependency on rocm::rocm_events() from the static lib. - Remove ROCPROFSYS_USE_ROCM compile guard around the hw-counters block in avail — ROCm availability is now handled at runtime. - Extract formatting state from global variables into a format_options struct, improving encapsulation. Remove dead code (settings_rexclude_exact, settings_rexclude_begin). - Minor refactor: Defer GPU device count and HW counter queries to after argument parsing, so they only run when --hw-counters is actually requested. Add new avail tests covering the no-GPU paths. ## JIRA ID N/A ## Test Plan - Tested and validated with existing tests and with the newly added tests. - All 19 avail ctest targets pass (ctest -R avail). - Manual smoke tests: --settings --brief, --components --brief, --hw-counters, and --all all produce correct output. ## Test Result - All tests passed without an error. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Signed-off-by: David Galiffi <David.Galiffi@amd.com> Co-authored-by: David Galiffi <David.Galiffi@amd.com> * [CLR][NFCI] Drop unnecessary HIP kernel args checks (#3500) * Update HIP readme file (#3951) ## Motivation There are incorrect notes at the HIP readme file. ## Technical Details Update HIP readme file. ## JIRA ID JIRA issue not needed. ## Test Plan Only README update. ## Test Result Only README update. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. * rocr: Update missing entry for API version bump (#3697) Also bump library version * rocr/aie: Remove timeout for submitting command chains (#3965) ## Motivation This PR removes the timeout from command chain submission to avoid spurious timeouts. ## Technical Details amd-xdna accepts a timeout for command chain submission, but it is hard to identify a good value. This change removes the timeout (`0` is no timeout), similarly to what XRT does. Additional changes: - Driver specific constants moved in the source file. - If a BO is not found, an invalid allocation error is returned to alert the user that a bad pointer was passed. ## JIRA ID NA ## Test Plan Local dispatch tests. ## Test Result ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> * versioned comgr dll causes error in opencl apps (#3974) versioned comgr dll causes error in runtime setting kComgrVersioned opencl=false hip=true ## Technical Details versioned comgr dll causes error in runtime setting kComgrVersioned opencl=false hip=true ## JIRA ID https://amd-hub.atlassian.net/browse/ROCM-20077 ## Test Plan hip-tests and opencl tests need to pass <!-- Explain any relevant testing done to verify this PR. --> ## Test Result hip-tests and opencl tests need to pass Rock CI verifies hip-tests/hip apps ocltst and clinfo were verified locally since Rock does not have Opencl testing yet. Used the artifacts from this PR and run https://github.com/ROCm/rocm-systems/actions/runs/23009594300 python build_tools/install_rocm_from_artifacts.py --run-id 23009594300 --amdgpu-family gfx94X-dcgpu --tests --run-github-repo ROCm/rocm-systems <!-- Briefly summarize test outcomes. --> ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. * [rocDecode] Add core function entry and exit logs. (#3851) ## Motivation This PR adds entry and exit logs for core functions. ## Technical Details - Info level (3) logs are added at core function entry and exit points. - Monotonic clock stamp, process id, thread id, and file name are added to the logs. - The logs now have similar format to HIP logs for better viewing when both are present. - Direct error or warning logs to std::cerr. Other logs still go to std::cout. ## JIRA ID <!-- If applicable, mention the JIRA ID resolved by …
JhaShweta1
added a commit
that referenced
this pull request
May 31, 2026
Empty merge of upstream commit ebd1e92 from github.com/NVIDIA/nccl. The second parent of this commit references the upstream commit; this commit introduces no source changes (tree unchanged from the first parent). ----- Original upstream commit message ----- NCCL v2.29.2-1 Release Device API Improvements: - Supports Device API struct versioning for backwards compatibility with future versions. - Adds ncclCommQueryProperties to allow Device API users to check supported features before creating a DevComm. - Adds host-accessible device pointer functions from symmetric registered ncclWindows. - Adds improved GIN documentation to clarify the support matrix. New One-Sided Host APIs: - Adds new host APIs (ncclPutSignal, ncclWaitSignal, etc) for both network and NVL using zero-SM. - One-sided communication operation writes data from the local buffer to a remote peer’s registered memory window without explicit participation from the target process. - Utilizes CopyEngine for NVL transfer and CPU proxy for network. - Requires CUDA 12.5 or greater. New Experimental Python language binding (NCCL4Py): - Pythonic NCCL API for Python applications - native collectives, P2P and other NCCL operations. - Interoperable with CUDA Python ecosystem: DLPack/CUDA Array Interface, and special support for PyTorch and CuPy. - Automatic cleanup of NCCL-managed resources (GPU buffers, registered buffers/windows, custom reduction operations). New LLVM intermediate representation (IR) support: - Exposes NCCL Device APIs through LLVM IR to enable consumption by diverse code generation systems. - Example usages include high-level languages, Just-In-Time (JIT) compilers, and domain-specific languages (DSL). - Build with EMIT_LLVM_IR=1 to generate LLVM IR bitcode. - Requires CUDA 12 and Clang 21. Built-in hybrid (LSA+GIN) symmetric kernel for AllGather: - Adds a new hierarchical kernel using MCRing (NVLS multicast + Ring) to improve performance and scalability of AllGather. - Requires symmetric memory registration and GIN. New ncclCommGrow API: - Adds the ability to dynamically and efficiently add ranks to an existing NCCL communicator. - Use ncclCommGrow with ncclCommShrink to adjust membership of communicators in response to failing and recovering nodes. - Also addresses the need for elastic applications to expand a running job by integrating new ranks. Multi-segment registration: - Expands buffer registration to support multiple segments of physical memory mapped to one contiguous VA space for the p2p, ib and nvls transports. - Enables support for expandable segments in PyTorch. Improves scalability of AllGatherV pattern: - Adds support for a scalable allgatherv pattern (group of broadcasts). - Adds new scheduler path and new kernels to improve performance at large scale. Debuggability & Observability Improvements: - RAS supports realtime monitoring to continuously track peer status changes. - Inspector adds support for Prometheus format output (with NCCL_INSPECTOR_PROM_DUMP=1), in addition to the existing JSON format. - Adds profiler support for CopyEngine(CE) based collectives. Community Engagement: - Adds contribution guide: https://github.com/NVIDIA/nccl/blob/master/CONTRIBUTING.md - Adds NCCL_SOCKET_POLL_TIMEOUT_MSEC which allows waiting instead of spinning during bootstrap in order to reduce CPU usage. (Github PR #1759) - Fixes segfault in ncclGin initialization that can happen if ncclGinIbGdaki.devices() fails after init() succeeds. (Github PR #1881) - Fixes crash that can happen when calling p2p and then collectives while using the same user buffer. (Github Issue #1859) - Fixes bug that was lowering performance on some sm80 or earlier machines with one NIC per GPU. (Github Issue #1876) - Clears non-fatal CUDA errors so they do not propagate. (Pytorch Issue #164402) Other Improvements: - Improves performance of large-size AllGather operations using symmetric memory buffers on Blackwell by transparently switching to CE collectives. - Improves the default number of channels per net peer for all-to-all, send, and recv to achieve better performance. - Improves performance tuning of 256M-512M message sizes on Blackwell for AllReduce. - Enables built-in symmetric kernels only on fully connected nvlink systems, as PCIE systems do not perform as well. - Prints git branch and commit checksum at the INFO level during NCCL initialization. - Improves support for symmetric window registrations on CUDA versions prior to 12.1. - Relaxes symmetric buffer registration requirements for collectives so that users can leverage the symmetric kernels with only one of the buffers being registered, when possible. - All2all, send, recv now obey NCCL_NETDEVS_POLICY. For these operations, NCCL will now by default use a subset of available network devices as dictated by the Network Device Policy. - Fixes a hang on GB200/300 + CX8 when the user disables GDR. - Fixes a bug that could cause AllReduce on ncclFloat8e4m3 to yield “no algorithm/protocol available”. - ncclCommWindowRegister will now return a NULL window if the system does not support window registration. - More prominent error when cuMulticastBind fails and NCCL_NVLS_ENABLE=2. - Upgrades to doca gpunetio v1.1. Known Limitations: - Since Device API was experimental in 2.28.x, applications that use the Device API in v2.28 may need modifications to work with v2.29. - One-sided host APIs (e.g. ncclPutSignal) currently do not support graph capture. Future releases will add cuda graph support. - The improved AllGatherV support breaks the NCCL profiler support for ncclBroadcast operations, limiting visibility to API events. NCCL_ALLGATHERV_ENABLE=0 can be used as a workaround until it is fixed in a future release. - NCCL4Py (experimental) has a known issue with cuda.core 0.5.0. We currently recommend using cuda.core 0.4.1 with nccl4py.
newling
pushed a commit
to newling/rocm-systems
that referenced
this pull request
Jun 3, 2026
Device API Improvements: - Supports Device API struct versioning for backwards compatibility with future versions. - Adds ncclCommQueryProperties to allow Device API users to check supported features before creating a DevComm. - Adds host-accessible device pointer functions from symmetric registered ncclWindows. - Adds improved GIN documentation to clarify the support matrix. New One-Sided Host APIs: - Adds new host APIs (ncclPutSignal, ncclWaitSignal, etc) for both network and NVL using zero-SM. - One-sided communication operation writes data from the local buffer to a remote peer’s registered memory window without explicit participation from the target process. - Utilizes CopyEngine for NVL transfer and CPU proxy for network. - Requires CUDA 12.5 or greater. New Experimental Python language binding (NCCL4Py): - Pythonic NCCL API for Python applications - native collectives, P2P and other NCCL operations. - Interoperable with CUDA Python ecosystem: DLPack/CUDA Array Interface, and special support for PyTorch and CuPy. - Automatic cleanup of NCCL-managed resources (GPU buffers, registered buffers/windows, custom reduction operations). New LLVM intermediate representation (IR) support: - Exposes NCCL Device APIs through LLVM IR to enable consumption by diverse code generation systems. - Example usages include high-level languages, Just-In-Time (JIT) compilers, and domain-specific languages (DSL). - Build with EMIT_LLVM_IR=1 to generate LLVM IR bitcode. - Requires CUDA 12 and Clang 21. Built-in hybrid (LSA+GIN) symmetric kernel for AllGather: - Adds a new hierarchical kernel using MCRing (NVLS multicast + Ring) to improve performance and scalability of AllGather. - Requires symmetric memory registration and GIN. New ncclCommGrow API: - Adds the ability to dynamically and efficiently add ranks to an existing NCCL communicator. - Use ncclCommGrow with ncclCommShrink to adjust membership of communicators in response to failing and recovering nodes. - Also addresses the need for elastic applications to expand a running job by integrating new ranks. Multi-segment registration: - Expands buffer registration to support multiple segments of physical memory mapped to one contiguous VA space for the p2p, ib and nvls transports. - Enables support for expandable segments in PyTorch. Improves scalability of AllGatherV pattern: - Adds support for a scalable allgatherv pattern (group of broadcasts). - Adds new scheduler path and new kernels to improve performance at large scale. Debuggability & Observability Improvements: - RAS supports realtime monitoring to continuously track peer status changes. - Inspector adds support for Prometheus format output (with NCCL_INSPECTOR_PROM_DUMP=1), in addition to the existing JSON format. - Adds profiler support for CopyEngine(CE) based collectives. Community Engagement: - Adds contribution guide: https://github.com/NVIDIA/nccl/blob/master/CONTRIBUTING.md - Adds NCCL_SOCKET_POLL_TIMEOUT_MSEC which allows waiting instead of spinning during bootstrap in order to reduce CPU usage. (Github PR ROCm#1759) - Fixes segfault in ncclGin initialization that can happen if ncclGinIbGdaki.devices() fails after init() succeeds. (Github PR ROCm#1881) - Fixes crash that can happen when calling p2p and then collectives while using the same user buffer. (Github Issue ROCm#1859) - Fixes bug that was lowering performance on some sm80 or earlier machines with one NIC per GPU. (Github Issue ROCm#1876) - Clears non-fatal CUDA errors so they do not propagate. (Pytorch Issue #164402) Other Improvements: - Improves performance of large-size AllGather operations using symmetric memory buffers on Blackwell by transparently switching to CE collectives. - Improves the default number of channels per net peer for all-to-all, send, and recv to achieve better performance. - Improves performance tuning of 256M-512M message sizes on Blackwell for AllReduce. - Enables built-in symmetric kernels only on fully connected nvlink systems, as PCIE systems do not perform as well. - Prints git branch and commit checksum at the INFO level during NCCL initialization. - Improves support for symmetric window registrations on CUDA versions prior to 12.1. - Relaxes symmetric buffer registration requirements for collectives so that users can leverage the symmetric kernels with only one of the buffers being registered, when possible. - All2all, send, recv now obey NCCL_NETDEVS_POLICY. For these operations, NCCL will now by default use a subset of available network devices as dictated by the Network Device Policy. - Fixes a hang on GB200/300 + CX8 when the user disables GDR. - Fixes a bug that could cause AllReduce on ncclFloat8e4m3 to yield “no algorithm/protocol available”. - ncclCommWindowRegister will now return a NULL window if the system does not support window registration. - More prominent error when cuMulticastBind fails and NCCL_NVLS_ENABLE=2. - Upgrades to doca gpunetio v1.1. Known Limitations: - Since Device API was experimental in 2.28.x, applications that use the Device API in v2.28 may need modifications to work with v2.29. - One-sided host APIs (e.g. ncclPutSignal) currently do not support graph capture. Future releases will add cuda graph support. - The improved AllGatherV support breaks the NCCL profiler support for ncclBroadcast operations, limiting visibility to API events. NCCL_ALLGATHERV_ENABLE=0 can be used as a workaround until it is fixed in a future release. - NCCL4Py (experimental) has a known issue with cuda.core 0.5.0. We currently recommend using cuda.core 0.4.1 with nccl4py.
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.
Motivation
these tests were disabled due to rocr changes made these tests failing. re-enabling back as rocr landed a patch to fix this
Technical Details
Test Plan
Test Result
Submission Checklist