Implements automatic device wake using getDRMDeviceId() DRM call when GPUs
are detected in low-power state. This ensures rocm-smi can access device
information on suspended GPUs.
Signed-off-by: Bindhiya Kanangot Balakrishnan <Bindhiya.KanangotBalakrishnan@amd.com>
* Added python test runner to execute rccl tests
* Disabled capture output to avoid hangs
* Add RCCL_TEST_MPI_HOSTFILE env var to get the hostfile
* Converted test_type to boolean gtest flag
* Removed unused return values
* Added custom rccl library usage
* Removed json output
* Updates to test_runner: added num_gpus field
* Address review comments
* Prepend env vars for single node, single process executions
* Added separate enums for exit and result codes
* Update configuration files
* Moved configurations to its own dir
* Address review comments
* Update tools/scripts/test_runner/README.md
Co-authored-by: Corey Derochie <161367113+corey-derochie-amd@users.noreply.github.com>
---------
Co-authored-by: Corey Derochie <161367113+corey-derochie-amd@users.noreply.github.com>
[ROCm/rccl commit: 0c2c61d2f1]
* Added python test runner to execute rccl tests
* Disabled capture output to avoid hangs
* Add RCCL_TEST_MPI_HOSTFILE env var to get the hostfile
* Converted test_type to boolean gtest flag
* Removed unused return values
* Added custom rccl library usage
* Removed json output
* Updates to test_runner: added num_gpus field
* Address review comments
* Prepend env vars for single node, single process executions
* Added separate enums for exit and result codes
* Update configuration files
* Moved configurations to its own dir
* Address review comments
* Update tools/scripts/test_runner/README.md
Co-authored-by: Corey Derochie <161367113+corey-derochie-amd@users.noreply.github.com>
---------
Co-authored-by: Corey Derochie <161367113+corey-derochie-amd@users.noreply.github.com>
The scratch_size_per_wave_ and dispatch_waves_ should use
the maximum values from all packets in the batch.
Signed-off-by: Longlong Yao <Longlong.Yao@amd.com>
Reviewed-by: Flora Cui <flora.cui@amd.com>
Problem:
The existing SDMA engine selection logic had several issues:
1. Same VirtualGPU/stream could use different SDMA engines for consecutive
async copies since copy_engine_status may report engines as busy
2. Busy and Preferred engine check for every copy
3. No global tracking of which VirtualGPU uses which engine, leading to
suboptimal resource allocation
Solution:
Implemented a global SDMA engine allocator with per-stream affinity:
- Added Device::SdmaEngineAllocator to manage VirtualGPU → engine assignments
* Maintains global map of active assignments
* Enforces exclusivity: different streams use different engines (except
inter-GPU copies where preferred engines are prioritized for optimal
hardware paths like XGMI links)
* Thread-safe allocation/release with Monitor lock
- Modified VirtualGPU to cache assigned engine locally (assigned_sdma_engine_)
for fast lookup without map access on hot path
- Refactored rocrCopyBuffer() to:
1. Check local cached engine first → use if assigned
2. Call AllocateSdmaEngine() if not assigned → cache result
- Moved HSA API queries (memory_copy_engine_status, memory_get_preferred_copy_engine)
into AllocateEngine() for cleaner separation of concerns
- Engine release on HostQueue::finish() instead of only VirtualGPU destruction
* Improves engine utilization by releasing earlier
* Added virtual ReleaseSdmaEngines() method to device::VirtualDevice
- Added future path for simple round-robin allocation (kUseSimpleRR) for
next-gen GPUs with uniform SDMA bandwidth (disabled by default)
Cleanup:
- Removed selectSdmaEngine() helper (logic moved to allocator)
- Removed getSdmaRWMasks() (allocator accesses maxSdmaReadMask_/WriteMask_ directly)
- Removed unused sdmaEngineReadMask_/WriteMask_ member variables from DmaBlitManager
Benefits:
- Ensures consistent per-stream SDMA engine usage
- Prevents cross-stream contention and engine thrashing
- Prioritizes hardware-optimal paths for inter-GPU transfers
- Better resource utilization through earlier release
- Cleaner, more maintainable code structure
- Fixes SWDEV-559349
- Fix build failure caused by correct libunwind not being found in some environments.
- Updated the `timemory` submodule to commit `24407d37ab85c46ba6c18fba9498320f825ee4e4 `.
* Use static catch2.lib instead of catch2.dll
Using catch2.dll incraeses execution time by 12x
* handle debug option for static catch2
* SWDEV-573539 - skip atomics on windows since its taking a very long time to execute
mlsejenkins needs newer cmake but compiler breaks with newer versions
so skipping on windows can be a workaround for now
---------
Co-authored-by: Joseph Macaranas <145489236+jayhawk-commits@users.noreply.github.com>
**Thread limit configuration and enforcement: **
* Added a check in `CMakeLists.txt` to ensure `ROCPROFSYS_MAX_THREADS` is at least 128, automatically setting it to 128 with a warning if a lower value is provided.
* Replaced hardcoded thread limit (`allowed_max_threads`) in `pthread_create_gotcha.cpp` with the configurable `ROCPROFSYS_MAX_THREADS` value, ensuring all runtime checks and warnings use the actual configured limit.
**Documentation improvements: **
* Updated the development guide to explain the new thread limit behavior, including how exceeding the limit is handled gracefully, how to configure it, and the build-time validation rules.
**Test updates: **
* Modified thread limit tests to use the configurable `ROCPROFSYS_MAX_THREADS` value instead of a hardcoded limit and expanded the range of tested thread values.
* Increased test timeouts to accommodate larger thread counts and ensure reliability with higher limits.
* util: dlsym optional helper
Like DLSYM_HELPER, but does not return if the symbol is not found.
Signed-off-by: Allen Hubbe <allen.hubbe@amd.com>
* gda ionic: sync dv and fw headers
Sync dv and fw headers to match out-of-tree libionic and firmware.
Signed-off-by: Allen Hubbe <allen.hubbe@amd.com>
* gda ionic: collapsed cqe
Detect and enable collapsed cqe if supported by drivers and firmware.
Fall back to regular completion queue.
Signed-off-by: Allen Hubbe <allen.hubbe@amd.com>
---------
Signed-off-by: Allen Hubbe <allen.hubbe@amd.com>
[ROCm/rocshmem commit: 1494c24f9a]
* util: dlsym optional helper
Like DLSYM_HELPER, but does not return if the symbol is not found.
Signed-off-by: Allen Hubbe <allen.hubbe@amd.com>
* gda ionic: sync dv and fw headers
Sync dv and fw headers to match out-of-tree libionic and firmware.
Signed-off-by: Allen Hubbe <allen.hubbe@amd.com>
* gda ionic: collapsed cqe
Detect and enable collapsed cqe if supported by drivers and firmware.
Fall back to regular completion queue.
Signed-off-by: Allen Hubbe <allen.hubbe@amd.com>
---------
Signed-off-by: Allen Hubbe <allen.hubbe@amd.com>
This has started failing on various developer build systems. Looking at it, it is not precisely clear how this ever worked given that nothing appears to be adding the DRM include dirs.
I'd prefer that we remove this delay loading (at least for TheRock builds where it is never needed), but in the meantime, this does fix the issue and is verified on an affected system.
Fixes https://github.com/ROCm/TheRock/issues/2744
* SWDEV-508225 - do not assert() after calling digestFatBinary() if it fails. Otherwise this causes assertions to trigger easily in systems that have an APU and a discrete GPU and the code was compiled for the discrete one
* SWDEV-508225 - fix that when using a non-existent ordinal in HIP_VISIBLE_DEVICES, getCurrentArch() would crash
fix: Add gpu_metrics 1.0 support which is still used by some hardware
Code changes related to the following:
* APIs
* Unit tests
Signed-off-by: Oliveira, Daniel <daniel.oliveira@amd.com>
* Add support for get and set APIs for CPUISOFreqPolicy and DFCState Control
- Add support for get and set APIs for CPUISOFreqPolicy and DFCState Control
in AMD SMI and also in the CLI tool
* CHANGELOG.md file updated
* SWDEV-562837: Update amdsmi-py-api.md as per the new APIs
Updated amdsmi-py-api.md as per the new APIs added.
---------
Signed-off-by: Soumya <sranjanr@amd.com>
Signed-off-by: gabrpham <Gabriel.Pham@amd.com>
Co-authored-by: Saka Sitharammurthy <SitharamMurthy.Saka@amd.com>
* fix reduction for gfx942 and 1201
match the synchronizaation of internal_putmem_wg and internal_getmem_wg
to their non-internal counterparts. the internal_putmem_wg is used in
the ipc reduction
* move specialization to internal_putmem
[ROCm/rocshmem commit: 8d2504d6c1]
* fix reduction for gfx942 and 1201
match the synchronizaation of internal_putmem_wg and internal_getmem_wg
to their non-internal counterparts. the internal_putmem_wg is used in
the ipc reduction
* move specialization to internal_putmem
* SWDEV-569319 Replace ScopedAcquire with stdcpp wrappers
* Remove KernelMutex and KernelSharedMutex abstractions with std::mutex and std::shared_mutex
* Replaced unique_locks with lock_guards
* More changes
* Replace new and deletes with smart pointers
* Replaced some more with shared ptrs
* Replacements with smart pointers - pt 2
* missed change
* Force ring in WarpSpeed manual mode and log event
* Skip usage for non-ring in WarpSpeed auto mode
* Enable WarpSpeed when its CU count is set
[ROCm/rccl commit: 93fdcb160c]
* Remove MFMA functionality in rocflop sample since its not supported in MI50
* Add gfx arc based support for MFMA and SMFMAC in rocflop.cpp
* Add --int32 usage doc
* Address review comments
it fails in about 50% of the cases. Will revisit later why it fails,
but RO is at the moment lower priority, so disabling the test for now.
[ROCm/rocshmem commit: ed2f75f1de]
* Fix buffer tracing synchronization lock
- PR #529 (in rocprofiler-sdk-internal) introduced waiting on the syncer flag when emplacing in a buffer to prevent the overwriting buffer records currently being processed in a buffer flush callback
- The above fix introduced a block on the both buffers when a buffer flush callback was being executed instead of a block on the buffer being flushed.
* Add rocpd tests for duplicate records
* Address code review comments
Calculate the actual scratch memory size required based on the
packet information for kernel dispatch.
If the required size exceeds the total allocated memory, scratch
memory must be reallocated. Otherwise, no action is needed.
miopen_gtest: Full/GPU_MIOpenDriverRegressionTest_FP16.MIOpenDriverRegressionHalf/0
Signed-off-by: Longlong Yao <Longlong.Yao@amd.com>
Reviewed-by: Flora Cui <flora.cui@amd.com>
Reviewed-by: Horatio Zhang <Hongkun.Zhang@amd.com>
Calculate the actual scratch memory size required based on the
packet information for kernel dispatch.
If the required size exceeds the total allocated memory, scratch
memory must be reallocated. Otherwise, no action is needed.
miopen_gtest: Full/GPU_MIOpenDriverRegressionTest_FP16.MIOpenDriverRegressionHalf/0
Signed-off-by: Longlong Yao <Longlong.Yao@amd.com>
Reviewed-by: Flora Cui <flora.cui@amd.com>
Reviewed-by: Horatio Zhang <Hongkun.Zhang@amd.com>
Yiltan
systems-assistant[bot]
Bindhiya Kanangot Balakrishnan
Atul Kulkarni
Kapil S. Pawar
koushikbillakanti-amd
Alexandra Sidorova
Longlong Yao
SaleelK
Flora Cui
David Galiffi
anujshuk-amd
Ajay GunaShekar
Alysa Liu
dependabot[bot]
David Yat Sin
Aleksandar Djordjevic
Aurelien Bouteiller
vedithal-amd
Godavarthy Surya, Anusha
Allen Hubbe
Stella Laurenzo
Yazen AL Musaffar
Gerardo Hernandez
Daniel Oliveira
SakaSitharammurthy
Edgar Gabriel
pghoshamd
Mustafa Abduljabbar
Nusrat Islam
Jonathan R. Madsen
habajpai-amd