hsakmt: Expose and use CWSR and Control stack sizes - #2200
Conversation
|
Mostly to connect dots for anyone who looks at this later: this matches this kernel change. |
dayatsin-amd
left a comment
There was a problem hiding this comment.
In thunk, when CwsrSize and CtlStackSize are not zero, we need to add code to use these values instead of trying to calculate them.
This would be in queues.c:update_ctx_save_restore_size
|
Right. That felt like it should be a separate change, and I wasn't confident enough in knowing where all those calculations had to happen |
|
"This will allow VGPR size to be determined dynamically, instead of relying on hardcoding" But I don't see a need to do this unless we're going to use actually use it during queue creation. |
|
"In thunk, when CwsrSize and CtlStackSize are not zero, we need to add code to use these values instead of trying to calculate them." I think we should still calculate this in Thunk for existing products. |
Fair enough. I didn't know who would know which part needed to be done next, but I can reword the commit message accordingly. This isn't my area of expertise. But I can try to throw something together |
e06f1df to
ce6b656
Compare
ce6b656 to
d8aed91
Compare
adf6c07 to
3d0ea23
Compare
3d0ea23 to
9ac66e0
Compare
7b84541 to
cc19a01
Compare
cc19a01 to
230c66b
Compare
230c66b to
df943d0
Compare
df943d0 to
8124f26
Compare
This is better than hardcoding values and hoping that they align with KFD's definitions Signed-off-by: Kent Russell <kent.russell@amd.com>
If KFD is providing the CwsrSize and CtlStackSize, use the maximum of those and the old calculations for the ctx_save_restore_size and ctl_stack_size defined in the queue Signed-off-by: Kent Russell <kent.russell@amd.com>
CwsrSize and CtlStackSize are reported by KFD ABI 1.20. GFX1151 specifically may have some issues if these regions are misaligned, so report a strong warning during topology initialization if the system is GFX1151 but is using KFD ABI < 1.20 Signed-off-by: Kent Russell <kent.russell@amd.com>
8124f26 to
aa81ed8
Compare
* hsakmt: Expose CWSR and Control stack sizes This is better than hardcoding values and hoping that they align with KFD's definitions Signed-off-by: Kent Russell <kent.russell@amd.com> * hsakmt: Use CwsrSize and CtlStackSize if available If KFD is providing the CwsrSize and CtlStackSize, use the maximum of those and the old calculations for the ctx_save_restore_size and ctl_stack_size defined in the queue Signed-off-by: Kent Russell <kent.russell@amd.com> * hsakmt: Add warning when ABI<1.20 on GFX1151 CwsrSize and CtlStackSize are reported by KFD ABI 1.20. GFX1151 specifically may have some issues if these regions are misaligned, so report a strong warning during topology initialization if the system is GFX1151 but is using KFD ABI < 1.20 Signed-off-by: Kent Russell <kent.russell@amd.com> [rocm-systems] ROCm/rocm-systems#2200 (commit 0a2ea9e)
* hsakmt: Expose CWSR and Control stack sizes This is better than hardcoding values and hoping that they align with KFD's definitions Signed-off-by: Kent Russell <kent.russell@amd.com> * hsakmt: Use CwsrSize and CtlStackSize if available If KFD is providing the CwsrSize and CtlStackSize, use the maximum of those and the old calculations for the ctx_save_restore_size and ctl_stack_size defined in the queue Signed-off-by: Kent Russell <kent.russell@amd.com> * hsakmt: Add warning when ABI<1.20 on GFX1151 CwsrSize and CtlStackSize are reported by KFD ABI 1.20. GFX1151 specifically may have some issues if these regions are misaligned, so report a strong warning during topology initialization if the system is GFX1151 but is using KFD ABI < 1.20 Signed-off-by: Kent Russell <kent.russell@amd.com> --------- Signed-off-by: Kent Russell <kent.russell@amd.com>
* hsakmt: Expose CWSR and Control stack sizes This is better than hardcoding values and hoping that they align with KFD's definitions Signed-off-by: Kent Russell <kent.russell@amd.com> * hsakmt: Use CwsrSize and CtlStackSize if available If KFD is providing the CwsrSize and CtlStackSize, use the maximum of those and the old calculations for the ctx_save_restore_size and ctl_stack_size defined in the queue Signed-off-by: Kent Russell <kent.russell@amd.com> * hsakmt: Add warning when ABI<1.20 on GFX1151 CwsrSize and CtlStackSize are reported by KFD ABI 1.20. GFX1151 specifically may have some issues if these regions are misaligned, so report a strong warning during topology initialization if the system is GFX1151 but is using KFD ABI < 1.20 Signed-off-by: Kent Russell <kent.russell@amd.com> --------- Signed-off-by: Kent Russell <kent.russell@amd.com>
* hsakmt: Expose CWSR and Control stack sizes This is better than hardcoding values and hoping that they align with KFD's definitions * hsakmt: Use CwsrSize and CtlStackSize if available If KFD is providing the CwsrSize and CtlStackSize, use the maximum of those and the old calculations for the ctx_save_restore_size and ctl_stack_size defined in the queue * hsakmt: Add warning when ABI<1.20 on GFX1151 CwsrSize and CtlStackSize are reported by KFD ABI 1.20. GFX1151 specifically may have some issues if these regions are misaligned, so report a strong warning during topology initialization if the system is GFX1151 but is using KFD ABI < 1.20 --------- Signed-off-by: Kent Russell <kent.russell@amd.com>
…nal-merge-new-713 (#416) * Move submodule definitions from projects/rccl/.gitmodules to root .gitmodules (#2298) * Replace usages of __ockl_gws_init with __builtin_amdgcn_ds_gws_init (#2235) * SWDEV-548892 - Always declare used ocml and ockl device libs functions (#2230) Ignore __CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__. This should not be relying on declarations from the clang builtin headers. There is no issue declaring the same intrinsics multiple times. This will enable removal of declarations from the clang builtin headers. * SWDEV-566268 - skip 2 failing tests on rock Windows (#2308) * SWDEV-538607 - Add SIMDe as a build dependency, remove naked intrinsic use. (#500) Co-authored-by: Alex Voicu <alexandru.voicu@amd.com> Co-authored-by: Ioannis Assiouras <Ioannis.Assiouras@amd.com> * [rocprofiler-sdk] Optimize rocprofiler-sdk `find_clients()` (#2267) * Version bump and Changelog update for ROCm version 7.2 (#2201) * Update projects/amdsmi/CHANGELOG.md * Bump to 26.2.1 --------- Signed-off-by: gabrpham <Gabriel.Pham@amd.com> * Run pre-commit's whitespace related hooks on projects/amdsmi (#2119) * Run pre-commit's whitespace related hooks on projects/amdsmi In order for pre-commit to be useful, everything needs to meet a common baseline. * Add whitespace back to Changelog for formatting --------- Signed-off-by: Mario Limonciello (AMD) <superm1@kernel.org> Signed-off-by: Maisam Arif <Maisam.Arif@amd.com> * Git url fix (#2285) * Update README-doc.md Correct GitHub URL for components moved into rocm-systems * Update amd_clr.rst Update github.com URLs * Update Dockerfile Update rocm-systems paths * Update CONTRIBUTING.md update for rocm-systems * Update CONTRIBUTING.md minor change * Update CONTRIBUTING.md * Update CONTRIBUTING.md * Update hip_runtime_api.rst Update for rocm-systems * Update installation.rst update URL to libhsakmt * Update what_is_hip.rst * Update projects/clr/CONTRIBUTING.md Co-authored-by: Dominic Widdows <dwiddows@gmail.com> * Update projects/clr/README-doc.md Co-authored-by: Dominic Widdows <dwiddows@gmail.com> * Update Dockerfile Update git clone for sparse checkout * Update projects/hip/CONTRIBUTING.md * Update projects/clr/CONTRIBUTING.md * Update projects/hipother/CONTRIBUTING.md --------- Co-authored-by: Dominic Widdows <dwiddows@gmail.com> * Enable Lintian Support for ROCM-SMI, ROCMINFO (#1650) * Enable Lintian Support for ROCM-SMI * Enable Lintian Support for ROCMINFO * Updated Lintian Override File Processing * Update UT Fix for Lintian rocmsmi,rocminfo * Update UT Fixes, Review Comments * Update Review Comments - removed extra white spaces, added error check for gzip, date commands * Update Review Comments - Correcting License Type * Sync Lintian ChangeLog * Changelog data sync enhanced * Update Review Comments, UT fix * white space cleanup - precommit check * [rocprofiler-systems] Add cached demangler and replace old demangle (#2135) * Add cached demangler and replace old * Add unit tests * Applied suggestions from code review * Applied suggestions from code review * hsakmt: Expose and use CWSR and Control stack sizes (#2200) * hsakmt: Expose CWSR and Control stack sizes This is better than hardcoding values and hoping that they align with KFD's definitions Signed-off-by: Kent Russell <kent.russell@amd.com> * hsakmt: Use CwsrSize and CtlStackSize if available If KFD is providing the CwsrSize and CtlStackSize, use the maximum of those and the old calculations for the ctx_save_restore_size and ctl_stack_size defined in the queue Signed-off-by: Kent Russell <kent.russell@amd.com> * hsakmt: Add warning when ABI<1.20 on GFX1151 CwsrSize and CtlStackSize are reported by KFD ABI 1.20. GFX1151 specifically may have some issues if these regions are misaligned, so report a strong warning during topology initialization if the system is GFX1151 but is using KFD ABI < 1.20 Signed-off-by: Kent Russell <kent.russell@amd.com> --------- Signed-off-by: Kent Russell <kent.russell@amd.com> * [rocprofiler-systems] Add span (#2142) * Add span * Update unit tests * [SWDEV-565460] AMD SMI Document Multiple Init Best Practices (#2293) * [SWDEV-565460] AMD SMI Document Multiple Init Best Practices Signed-off-by: amd-josnarlo <josnarlo.amd.com> * Add sphinxcontrib-mermaid to render diagram in HTML bump rocm-docs-core to 1.31.0 pip-compile requirements.txt --------- Signed-off-by: amd-josnarlo <josnarlo.amd.com> Co-authored-by: amd-josnarlo <josnarlo.amd.com> Co-authored-by: Peter Park <peter.park@amd.com> * Adding org var and dynamic selection of targets (#2317) * Added AMDSMI CI to rocm-systems(#2074) Signed-off-by: Justin Williams <Justin.Williams@amd.com> * [rocprofiler-compute][tui] menu bar lag fix (#1942) * [rocprofiler-compute] Counter accuracy tests and improvements for iteration multiplexing (#2011) * Added laplace solver in samples * Add laplace eqn in CMake * Added counter accuracy test * Add iteration CLI arg for laplace eq * Unnest profile method * Missing counter warning * Updated insufficient kernel warning * Added reference for laplace equation * variable name change * Added comments for data comparison * Included scipy as test requirement * Added line number for ref * split stochastic and deterministic tests * Added order cli option for laplace_eqn * Install laplace eqn * Missing counter warning * Warn about missing kernels during analysis * Update tests * Split iteration multiplexing ctests * Updated warning * Incorporated copilot's suggestions * [RDC] Optimize RDC counter sampling with greedy packing algorithm (#1590) * Optimize RDC counter sampling with greedy packing algorithm This change significantly reduces the number of rocprofiler-sdk sample calls by implementing a greedy packing algorithm that groups multiple counters into the minimal number of hardware profiles. Key improvements: - Implement greedy packing algorithm to combine counters into minimal profiles - Add ProfileSet structure to manage packed counter configurations - Cache packed profile sets for reuse across queries - Group telemetry field requests by GPU for bulk processing - Reduce sample calls by ~35% (from 100 to 65 for typical workloads) Performance impact: - 13 counters now packed into 3 profiles (77% compression) - Reduces overhead from profile creation and context switching - More efficient utilization of hardware counter resources Implementation details: - Added create_profiles_for_counters() using greedy algorithm - Added sample_counters_with_packing() for bulk sampling - Modified telemetry layer to use rocp_lookup_bulk() - Preserves all field transformations and special handling Testing shows successful packing with expected performance gains. No functional changes to external APIs or behavior. Co-Authored-By: Ben Welton <bwelton@amd.com> * Address PR review feedback This commit addresses all review comments from the initial PR: 1. Fix division by zero risk in debug logging - Added check for empty counters vector before calculating compression ratio - Avoids potential division by zero when logging profile creation stats 2. Improve thread safety for statistics tracking - Changed static uint64_t to std::atomic<uint64_t> for thread-safe counters - Prevents race conditions in multi-threaded sampling scenarios 3. Remove unused variable - Removed unused profile_index variable that was incremented but never used - Cleaned up dead code 4. Clean up code formatting - Removed extra blank lines for consistency - Applied formatting fixes across modified files 5. Refactor code duplication between rocp_lookup and rocp_lookup_bulk - Created apply_field_transformation() helper function - Eliminates ~70 lines of duplicated switch statement logic - Centralizes field transformation logic in single location - Makes future maintenance easier 6. Document non-rocprofiler metrics handling - Added comments explaining how bulk lookup handles special cases - Clarifies that non-profiler fields like KFD_ID are handled in transformation All changes maintain backward compatibility and pass compilation. Co-Authored-By: Ben Welton <bwelton@amd.com> --------- Co-authored-by: Ben Welton <bwelton@amd.com> Co-authored-by: Adam Pryor <61172547+adam360x@users.noreply.github.com> * Scale down memory usage data when the actual data is stored to cache (#2343) * Add HasExpertSchedMode device prop (#2241) * Add HasExpertSchedMode device prop * Add unit tests for HasExpertSchedMode * Add gfx12 check for HasExpertSchedMode prop * Update gfx major version check and test for ExpertSchedMode * Minor fix and ROCr version bump * Update projects/rocr-runtime/runtime/hsa-runtime/inc/hsa_ext_amd.h * Update projects/rocr-runtime/runtime/hsa-runtime/inc/hsa_ext_amd.h * Apply suggestion from @dayatsin-amd * Apply suggestion from @dayatsin-amd --------- Co-authored-by: Stefan Sokolovic <stefan.sokolovic2@amd.com> Co-authored-by: David Yat Sin <77975354+dayatsin-amd@users.noreply.github.com> * SWDEV-571304 : Fix the constructor for __half (#2240) - comply with cuda - Fix usecase for constexpr Signed-off-by: sdashmiz <shadi.dashmiz@amd.com> * rocr: Fix asyncHandler segfault (#2261) Fix initialization order for the async events handler. The polling thread would launch before the wake signal is initialized. * Revert "Add HasExpertSchedMode device prop (#2241)" (#2371) This reverts commit c0b4aef. * test: add unit tests for common utilities from PR #1249 (#2237) * test: add unit tests for common utilities from PR #1249 * incorporate review comments specific to tests formatting * use filesystem API instead of std::system for safer cleanup * Add ghc/filesystem submodule v1.5.14 for portable C++17 filesystem support * fix: add cmake/GhcFilesystem.cmake for CI submodule auto-checkout * incorporate review comment * incorporate review comment * Do not allocate code objects when we map a static code object (#2332) * [SWDEV-566924] Update KFD_ID metric to use amd-smi instead of rocprof (#2355) * [rocprofiler-compute] Improve iteration multiplexing code and documentation (#2080) * Improve Iteration multiplexing * Improve iteration multiplexing documentation by adding usage note and listing caveats * Bugfixes for iteration mulitplexing * Use merge iteration multiplexing in analysis webui and db mode * Do not remove Dispatch_ID column in merge iteration multiplexing since it is needed for analysis of top dispatches based on duration * Bugfixes for analysis logic * Graceful handling of missing counters in case of iteration multiplexing * Improved warnings when metrics could not be calculated due to missing counter data * Fix the check to prevent showing table when a column is full of N/A * Improve detection of empty values when metric evaludation fails due to missing counter data * Bugfixes for profile logic * Fix kernel filtering during roofline benchmark phase * Update changelog for bugfixes * Remove unnecessary columns when merging dispatches for iteration multiplexing * bugfix * Better analysis warnings * fix to_std() in parser * Use median in merge iteration multiplex * Address review comments * Fix cmake formatting * fix None handling of parser util functions * Enable stochastic counter accuracy test * fix cmake formatting * [rocprofiler-compute] Add SDK dependency to rocprofiler-compute-tarball.yml workflow (#2329) * Install rocm-dev in rocprofiler-compute-tarball.yml workflow * Update paths for push and PR for rocprofiler-compute-tarball.yml * Add ROCm dependencies to disttest job * cmake fix binary link creation and fix format * Use python3 instead of python3.9 in RHEL 8 and RHEL 9 workflows * set default python3 to python3.9 in rhel8 * Try alternatives setup for python3 in RHEL8 env * Add pip install cmake to debug RHEL8 issue * Remove python3.11 in RHEL8 workflow * Add back comment regarding RHEL8 --------- Co-authored-by: Vignesh Edithal <Vignesh.Edithal@amd.com> * Revert "clr: Minor fixes for error return" (#2399) - This reverts commit 8dd8436. - Resolve MIOpen test failures observed in TheRock - TheRock Issue: ROCm/TheRock#2642 - room-systems issue: #2400 * fix: prevent double-free crash during process exit in amd-smi (#2213) * [SWDEV-525635] Updated output file handling options (#1896) Currently if the input file name already exists, the tool appends output to existing file. Added overwrite, append, or no(discard) options to choose from. Signed-off-by: Bindhiya Kanangot Balakrishnan <Bindhiya.KanangotBalakrishnan@amd.com> Co-authored-by: Bindhiya Kanangot Balakrishnan <Bindhiya.KanangotBalakrishnan@amd.com> * SWDEV-548892 - Stop using __ockl_lane_id (#2186) __lane_id already exists and is identical. * SWDEV-565304 - Pass numa node to migrate pages correctly (#1729) * SWDEV-565304 - Pass cpuId of the the thread currently running * SWDEV-565304 - Numa id to be returned * SWDEV-565304 - Numa id to be returned * SWDEV-572329 - Remove barrier packet (#2304) * SWDEV-572676 - adjust tile size to 32 in Unit_hipCGThreadBlockTileType for Navi4x (#2379) * SWDEV-572676 - adjust tile size to 32 for Navi4x * SWDEV-572676 - change tile size from fixed value to warp size * Revert "Adding org var and dynamic selection of targets (#2317)" (#2416) This reverts commit c9ac018. * Fixed ctests (#2406) * [rocprofiler-systems] Improve metadata parsing (#2238) * Improve metadata JSON parsing * Fix string ownership * Put cached perfetto traces as default one (#2138) * Put cached perfetto traces as default one * Improve cached data and perfetto traces in order to be more aligned with E2E tests * Addressing PR comments and findings * Force early instrumentation bundle instantiation * Sync-up insturumented containers with thread growth data * Revert ompvv number of host threads to default 8 * Fixed counter track namings for amd-smi * AIPROFSYST-34 [rocprof-sys] Update documentation describing newly introduced changes to default tracing mechanism * Adding full build capability to theROCK for HIP changes (#2003) ## Add Full Build Capability to theROCK for HIP ### Summary This PR adds full build support to **theROCK** for HIP-related changes, ensuring that all components are built. ### Changes - Enabled full build coverage for the following projects: - `projects/clr` - `projects/hip` - `projects/hip-tests` - `projects/rocr-runtime` - Updated build configuration to include all targets for the above projects. - Ensured rocm-libraries is pulled to build optional components. ### Motivation These changes are required to support HIP development and testing within theROCK by ensuring all components are built together. This improves reliability, integration testing. * [Documentaion] updating roctx library linkage documentation (#2251) * Add CODEOWNERS for rocprofiler-sdk project (#2427) ## Motivation Missing CODEOWNERS for ROCProfiler-SDK <!-- Explain the purpose of this PR and the goals it aims to achieve. --> ## Technical Details Add CODEOWNERS for rocprofiler-sdk project <!-- 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 - [X] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. * [Rocprof-Sys] Resolve crash when profiling TensorFlow GPU application (#2381) * fix: resolve crash when profiling TensorFlow GPU application * incorporate review comments * updated min_rows from 3 to 2 for threads table validation as internal threads are not profiled and are now correctly bypassed * [rocpd] Write rocpd yaml files as a list, even when only 1 file (#2288) * [rocprofiler-systems] Add SIGKILL delay option (#2384) ## Motivation When profiling multi-process applications where a parent process sends SIGKILL to child processes, the termination can occur before the profiler has a chance to flush collected data. This PR introduces a configurable delay before SIGKILL signals are forwarded, allowing profiling data to be captured before process termination. This is workaround. ## Technical Details - Added new configuration setting `ROCPROFSYS_KILL_DELAY` (default: 0 seconds) to specify a delay before SIGKILL signals are forwarded to other processes - Implemented `kill_gotcha` component that intercepts the `kill()` system call - The gotcha only delays SIGKILL signals sent to external processes (pid > 0 and not self) - Integrated `kill_gotcha_t` into the `preinit_bundle_t` for early initialization * SWDEV-555178 - Calculate phys mem offset for remap range (#1879) * [rocprofiler-compute] Fix for multi process workload profiling (#2418) * Fix for multi process workload profiling Native counter collection tool updates: * Do not dump empty counter data for a process * Use PID instead of UUID for dumped csv files to facilitate correlation * Handle merging multiple pairs of rocpd (from sdk tool) and csv (from native tool) files * Handle merging multiple pairs of csv (from sdk tool) and csv (from native tool) files Rocpd output format updates: * Merge multiple rocpd databases into a single csv * Reset dispatch id and kernel id for unique dispatches and unique kernels respectively * Retain multiple rocpd databases per run for multi process workloads * Add test case for multiprocess profiling using rocflop workload * Add rocflop * Fix native counter csv to rocprofv3 csv conversion * Use kernel_id instead of dispatch_id to correlate native counter csv and kernel trace csv * python formatting using ruff 0.14 instead of 0.13 * [rocprofiler-compute] Test env. vars. in rocprofiler-sdk backend (#2414) * Test env. vars. in rocprofiler-sdk backend * Improve rocprofiler-sdk backend test case to check for env. vars. and ensure we do not overwrite irrelevant env. vars. * Remove unnecessary usage of ROCPROF_INDIVIDUAL_XCC_MODE env. var. * Formatting fixes * Test fixes * Remove redundant code in tests * Remove usage of utils_mod and use utils instead, this prevents duplicate imports * [rocprofiler-compute] Faster counter accuracy testing (#2420) * Faster counter accuracy testing * Better handle SPI_CSN_* metrics for lesser than MI350 series * Use metric filtering to collect only relevant counters for comparison * Ensure all workload folders are deleted after testing is completed * Dont use clean_existing=False * Add manual test for all counter accuracy * Revert "Adding full build capability to theROCK for HIP changes (#2003)" (#2441) This reverts commit 0a52f5c. Reverts #2003 MIOpen build failures on windows causing blockers on unrelated file changes. * SWDEV-558849 - keep the lastEnqueueCommand_ when PAL backend is enabled (#2320) * [SWDEV-566543] Fix param validation in FrequenciesRead test (#2430) Fixed incorrect error code expectation in FrequenciesRead test when calling amdsmi_get_gpu_pci_bandwidth() with nullptr parameter. Signed-off-by: Bindhiya Kanangot Balakrishnan <Bindhiya.KanangotBalakrishnan@amd.com> * Fix agent device ID in the cached kernel_dispatch trace (#2452) * Remove redundant ROCPROFSYS_TRACE_CACHED variable from the code (#2434) * [rocprofiler-compute] Write raw counter and metric values (#2314) * Added tool for dumping counter and metric values * Skip Linting * Added support for iteration multiplexing * Remove subparser and supress compute options * Specify output dir * Add kernel info * csv name change * Added comments * Support dispatch id-less dataframes * Formatting fix * Add default for path * Print help with no args * Support only single workload * [rocprofiler-compute] Fix merging logic for multi process (#2445) * Fix merging logic for multi process * Fix dispatch id reset logic in case of rocpd format * Fix kernel id reset logic in case of csv format * Revert correlation logic change in csv format * Do inner join instead of left join * Remove SMFMAC functionality in rocflop sample since its not supported in MI100 (#2456) * SWDEV-558849 - Fix Windows build for ROCR backend (#2368) * [SWDEV-548460] Add RDC Policy Reset Message (#2180) * [SWDEV-548460] Add RDC Policy Reset Message * [rdc] Bump version to 1.3.0 Signed-off-by: Galantsev, Dmitrii <dmitrii.galantsev@amd.com> * chore: [rdc] Format CMakeLists.txt Signed-off-by: Galantsev, Dmitrii <dmitrii.galantsev@amd.com> --------- Signed-off-by: Galantsev, Dmitrii <dmitrii.galantsev@amd.com> Co-authored-by: Galantsev, Dmitrii <dmitrii.galantsev@amd.com> * SWDEV-567545 - Implement block_rank in co-op grid groups (#2182) * SWDEV-567545 - Implement block_rank in co-op grid groups * Fix test cases (#2462) * [SWDEV-567254] Sync Unified and Linux header (#2220) * [SWDEV-567254] Sync Unified and Linux header Signed-off-by: Joseph Narlo <joseph.narlo@amd.com> * Latest sync changes * Sync * Add back guest_windows tag * Sync --------- Signed-off-by: Joseph Narlo <joseph.narlo@amd.com> Co-authored-by: amd-josnarlo <josnarlo.amd.com> * SWDEV-573589 - Fixed performance regression due to the increase of the signal pool (#2470) * [rocprofiler-register] Fix compilation with libc++ (#1241) `tests/rocprofiler/rocprofiler.cpp` uses `std::string` without including `<string>` directly. This works with libstdc++ due to transitive includes, but fails with libc++. Closes #1240 * Update using-rocprofv3-with-openmp.rst (#2473) * [rocprofiler-compute] Improve roofline logging for roofline.csv (#2390) * enhanced roofline log output for graceful exit * addressed comment, added block filtering * ruff format --------- Signed-off-by: gabrpham <Gabriel.Pham@amd.com> Signed-off-by: Mario Limonciello (AMD) <superm1@kernel.org> Signed-off-by: Maisam Arif <Maisam.Arif@amd.com> Signed-off-by: Kent Russell <kent.russell@amd.com> Signed-off-by: amd-josnarlo <josnarlo.amd.com> Signed-off-by: Justin Williams <Justin.Williams@amd.com> Signed-off-by: sdashmiz <shadi.dashmiz@amd.com> Signed-off-by: Bindhiya Kanangot Balakrishnan <Bindhiya.KanangotBalakrishnan@amd.com> Signed-off-by: Galantsev, Dmitrii <dmitrii.galantsev@amd.com> Signed-off-by: Joseph Narlo <joseph.narlo@amd.com> Co-authored-by: Jason Bonnell <166553723+jbonnell-amd@users.noreply.github.com> Co-authored-by: Fábio Mestre <fabio.mestre@amd.com> Co-authored-by: Matt Arsenault <Matthew.Arsenault@amd.com> Co-authored-by: Ajay GunaShekar <86270081+agunashe@users.noreply.github.com> Co-authored-by: systems-assistant[bot] <221163467+systems-assistant[bot]@users.noreply.github.com> Co-authored-by: Alex Voicu <alexandru.voicu@amd.com> Co-authored-by: Ioannis Assiouras <Ioannis.Assiouras@amd.com> Co-authored-by: Jonathan R. Madsen <jrmadsen@users.noreply.github.com> Co-authored-by: gabrpham <Gabriel.Pham@amd.com> Co-authored-by: Mario Limonciello <superm1@kernel.org> Co-authored-by: randyh62 <42045079+randyh62@users.noreply.github.com> Co-authored-by: Dominic Widdows <dwiddows@gmail.com> Co-authored-by: arvindcheru <90783369+arvindcheru@users.noreply.github.com> Co-authored-by: Milan Radosavljevic <milan.radosavljevic@amd.com> Co-authored-by: Kent Russell <kent.russell@amd.com> Co-authored-by: Aleksandar Djordjevic <aleksandar.djordjevic@amd.com> Co-authored-by: Joseph Narlo <joseph.narlo@amd.com> Co-authored-by: Peter Park <peter.park@amd.com> Co-authored-by: Geo Min <geomin12@amd.com> Co-authored-by: amd-juwillia <Justin.Williams@amd.com> Co-authored-by: xuchen-amd <xuchen@amd.com> Co-authored-by: abchoudh-amd <abchoudh@amd.com> Co-authored-by: Benjamin Welton <Benjamin.Welton@amd.com> Co-authored-by: Ben Welton <bwelton@amd.com> Co-authored-by: Adam Pryor <61172547+adam360x@users.noreply.github.com> Co-authored-by: marantic-amd <marantic@amd.com> Co-authored-by: Filip Jankovic <161825881+fjankovi@users.noreply.github.com> Co-authored-by: Stefan Sokolovic <stefan.sokolovic2@amd.com> Co-authored-by: David Yat Sin <77975354+dayatsin-amd@users.noreply.github.com> Co-authored-by: Shadi Dashmiz <94885391+shadidashmiz@users.noreply.github.com> Co-authored-by: Maneesh Gupta <maneesh.gupta@amd.com> Co-authored-by: habajpai-amd <habajpai@amd.com> Co-authored-by: Jatin Chaudhary <51944368+cjatin@users.noreply.github.com> Co-authored-by: vedithal-amd <Vignesh.Edithal@amd.com> Co-authored-by: Sourabh U Betigeri <sourabhbetigeri@gmail.com> Co-authored-by: Bindhiya Kanangot Balakrishnan <Bindhiya.KanangotBalakrishnan@amd.com> Co-authored-by: cadolphe-amd <chris.adolphe@amd.com> Co-authored-by: ammallya <ameyakeshava.mallya@amd.com> Co-authored-by: Gopesh Bhardwaj <gopesh.bhardwaj@amd.com> Co-authored-by: Ammar ELWazir <ammar.elwazir@amd.com> Co-authored-by: Young Hui - AMD <145490163+yhuiYH@users.noreply.github.com> Co-authored-by: marandje <marko.arandjelovic@amd.com> Co-authored-by: Ioannis Assiouras <38722728+iassiour@users.noreply.github.com> Co-authored-by: German Andryeyev <56892148+gandryey@users.noreply.github.com> Co-authored-by: Galantsev, Dmitrii <dmitrii.galantsev@amd.com> Co-authored-by: Jimbo <57198431+jiabaxie@users.noreply.github.com> Co-authored-by: Sv. Lockal <AngryLoki@users.noreply.github.com> Co-authored-by: Swati Rawat <120587655+SwRaw@users.noreply.github.com> Co-authored-by: jamessiddeley-amd <James.Siddeley@amd.com> Co-authored-by: Sivasuntharampillai, Haresh <Haresh.Sivasuntharampillai@amd.com>
This will allow VGPR size to be determined dynamically, instead of relying on hardcoding
Motivation
Technical Details
JIRA ID
Test Plan
Test Result
Submission Checklist