* clr: SWDEV-547890 - Maintain an MQD for the emulated AQL queue
To simplify the shader debugger implementation, maintain the relevant
parts of the emulated AQL queue's MQD (amd_queue_t): read_dispatch_id,
write_dispatch_id, compute_tmpring_size.
With this MQD, the shader debugger can handle the emulated AQL queue
the same way it does the real AQL queue, no specialization is required.
* clr: SWDEV-547890 - Conservatively update the MQD's read_dispatch_id
The read_dispatch_id cannot be smaller than the current aql_packet_id
- hsa_queue.size for the debugger to work correctly.
The read_dispatch_id really should be updated when the CmdBuf is marked
as complete. Left a FIXME to address it in a future commit.
* Added ERROR message class to handle fatal error messages.
New ERROR message class will print the message in all debug level,
including none.
Change some of the fatal error message to be in ERROR instead of WARN.
Added new error handler function to print out more meaningful error
message in the future.
* Added CHANGELOG entry.
* Update CHANGELOG.md
Co-authored-by: Jeffrey Novotny <jnovotny@amd.com>
* Change to no longer reuse NONE as ERROR. ERROR is now a separated class.
* Update CHANGELOG.md
Co-authored-by: Jeffrey Novotny <jnovotny@amd.com>
---------
Co-authored-by: Jeffrey Novotny <jnovotny@amd.com>
* fix crash created by path and arg for pc_sampling and add copyright for mat_mul
* resolve fomat issue of line too long
* bugfixes
* copy gfx9 config template to analysis config in src
---------
Co-authored-by: Wang <ywang103@ctr2-alola-login-01.amd.com>
Co-authored-by: Vignesh Edithal <Vignesh.Edithal@amd.com>
* prevent batching when send/recv bytes dont match, restore bit reversal for channel to part mapping, prevent batching beyond 32-nodes
* correct computation for channel to part mapping
* update changelog
* disabling p2p-batching by default
Allow amdsmi to find libamd_smi.so and librocm-core.so relative to
amdsmi_wrapper.py location.
The amdsmi_wrapper.py file is located in
_rocm_sdk_core/share/amd_smi/amdsmi and the libraries are in
_rocm_sdk_core/lib/libamd_smi.so.26.
_rocm_sdk_core/lib/librocm-core.so.1.
Allow amdsmi to find libamd_smi.so and librocm-core.so relative to
amdsmi_wrapper.py location.
The amdsmi_wrapper.py file is located in
_rocm_sdk_core/share/amd_smi/amdsmi and the libraries are in
_rocm_sdk_core/lib/libamd_smi.so.26.
_rocm_sdk_core/lib/librocm-core.so.1.
[ROCm/amdsmi commit: ad20d57162]
The out of bound writes caused corruption in next field,
which was weight. Fixed by reading to a temp and then assigning
safely.
Signed-off-by: Bindhiya Kanangot Balakrishnan <Bindhiya.KanangotBalakrishnan@amd.com>
The out of bound writes caused corruption in next field,
which was weight. Fixed by reading to a temp and then assigning
safely.
Signed-off-by: Bindhiya Kanangot Balakrishnan <Bindhiya.KanangotBalakrishnan@amd.com>
[ROCm/amdsmi commit: a2aae5e8a9]
* Add initial commit to increase tb size to 512
* Fix LL perf issue when subset of NCCL_MAX_NTHREADS is used
Adding a constant to barrier_generic logic from using fallback logic when nthreads < NCCL_MAX_NTHREADS and nthreads == blockDim.X
* Adjust nthreads for LL
* Opt threads for reduce_scatter upper small range
* Add macro for single node
* Restrict MSCCL to 256 threads to prevent mem access fault
* Support pre-MI350 compatibility
* Partially refactor threadblock size override
* Use const macros instead of numerals
* opt out of unused function
Added json and csv output format support for newly
added xgmi link_status. Aligned legend.
Signed-off-by: Bindhiya Kanangot Balakrishnan <Bindhiya.KanangotBalakrishnan@amd.com>
* SWDEV-533237 Add initial support for hipOccupancyAvailableDynamicSMemPerBlock API
* SWDEV-533237 Add hipOccupancyAvailableDynamicSMemPerBlock wrapper for nvidia
* SWDEV-533237 Add implementation of hipOccupancyAvailableDynamicSMemPerBlock API
* SWDEV-533237 Add LDSAlignment field in Isa table
---------
Co-authored-by: Rahul Manocha <rmanocha@amd.com>
* Add examples to tools folder
* Correct P2P memory access section
* Sync poriting guide
* Add HIP Graph tutorial
* Add hint about using amdgpu-dkms for IPC API
* Add a few more env variables
* Added amd-smi set --pcie command
* Removed current pcie level due to it not being static
* Added pcie information to static --bus
---------
Signed-off-by: Pham, Gabriel <Gabriel.Pham@amd.com>
Signed-off-by: Maisam Arif <Maisam.Arif@amd.com>
* Added amd-smi set --pcie command
* Removed current pcie level due to it not being static
* Added pcie information to static --bus
---------
Signed-off-by: Pham, Gabriel <Gabriel.Pham@amd.com>
Signed-off-by: Maisam Arif <Maisam.Arif@amd.com>
[ROCm/amdsmi commit: 9e3537d778]
- **Added evicted_time metric for kfd processes**.
- Time that queues are evicted on a GPU in milliseconds
- Added to CLI in `amd-smi monitor -q` and `amd-smi process`
- Added to C API and Python API:
- amdsmi_get_gpu_process_list()
- amdsmi_get_gpu_compute_process_info()
- amdsmi_get_gpu_compute_process_info_by_pid()
---------
Signed-off-by: Pryor, Adam <Adam.Pryor@amd.com>
- **Added evicted_time metric for kfd processes**.
- Time that queues are evicted on a GPU in milliseconds
- Added to CLI in `amd-smi monitor -q` and `amd-smi process`
- Added to C API and Python API:
- amdsmi_get_gpu_process_list()
- amdsmi_get_gpu_compute_process_info()
- amdsmi_get_gpu_compute_process_info_by_pid()
---------
Signed-off-by: Pryor, Adam <Adam.Pryor@amd.com>
[ROCm/amdsmi commit: 2144cfbba4]
lmoriche
Satyanvesh Dittakavi
David Galiffi
Jaydeep
Ossian O'Reilly
Nilesh M Negi
David DeBonis
Yiannis Papadopoulos
Kian Cossettini
Ioannis Assiouras
Arm Patinyasakdikul
ywang103-amd
gabrpham_amdeng
adapryor
isaki001
Dmitrii
Galantsev, Dmitrii
Bindhiya Kanangot Balakrishnan
Bertan Dogancay
marandje
Mustafa Abduljabbar
Charis Poag
Allan Xavier
gilbertlee-amd
cadolphe-amd
xuchen-amd
dsicarov-amd
Istvan Kiss
Geo Min
corey-derochie-amd