Added#

• Added enumeration mapping amdsmi_get_gpu_enumeration_info() to Python and C APIs. The mapping is also enabled in the CLI interface via amd-smi list -e.

• Added dynamic virtualization mode detection.

  • Added new C and Python API amdsmi_get_gpu_virtualization_mode_info.

  • Added new C and Python enum amdsmi_virtualization_mode_t.

• Added TVIOL_ACTIVE to amd-smi monitor.

• Added support for GPU metrics 1.7 to amdsmi_get_gpu_metrics_info().

• Added new API amdsmi_get_gpu_xgmi_link_status() and CLI amd-smi xgmi --link-status.

• Added fclk and socclk info to amd-smi metric -c/--clock.

• Added new command amd-smi set -c/--clock-level.

• Added new command amd-smi static -C/--clock.

Changed#

• Updated AMD SMI library version number format to reflect changes in backward compatibility and offer more semantic versioning.

  • Removed year from AMD SMI library version number.

  • Version format changed from 25.3.0.0 (Year.Major.Minor.Patch) to 25.3.0 (Major.Minor.Patch).

  • Removed year in all version references.

• Added new Python dependencies: python3-setuptools and python3-wheel.

• Removed initialization requirements for amdsmi_get_lib_version() and added amdsmi_get_rocm_version() to the Python API & CLI.

• Added an additional argument sensor_ind to amdsmi_get_power_info().

  • This change breaks previous C API calls and will require a change.

  • Python API now accepts sensor_ind as an optional argument. This does not impact previous usage.

• Deprecated enum AMDSMI_NORMAL_STRING_LENGTH in favor of AMDSMI_MAX_STRING_LENGTH.

• Changed to use thread local mutex by default.

  • Most sysfs reads do not require cross-process level mutex and writes to sysfs should be protected by the kernel already.

  • Users can still switch to the old behavior by setting the environment variable AMDSMI_MUTEX_CROSS_PROCESS=1.

• Changed amdsmi_vram_vendor_type_t enum names impacting the amdsmi_vram_info_t structure. This change also impacts the usage of the vram_vendor output of amdsmi_get_gpu_vram_info().

• Changed the amdsmi_nps_caps_t struct impacting amdsmi_memory_partition_config_t, amdsmi_accelerator_partition_t, amdsmi_accelerator_partition_profile_config_t. Affected functions are:

  • amdsmi_get_gpu_memory_partition_config()

  • amdsmi_get_gpu_accelerator_partition_profile()

  • amdsmi_get_gpu_accelerator_partition_profile_config()

• Corrected CLI CPU argument name. --cpu-pwr-svi-telemtry-rails is now --cpu-pwr-svi-telemetry-rails.

• Added amdgpu driver version and amd_hsmp driver version to the amd-smi version command.

• All amd-smi set and amd-smi reset options are now mutually exclusive. You can now only use one set option as a time.

• Changed the name of the power field to energy_accumulator in the Python API for amdsmi_get_energy_count().

• Added violation status output for Graphics Clock Below Host Limit to amd-smi CLI: amdsmi_get_violation_status(), amd-smi metric  --throttle, and amd-smi monitor --violation. Users can retrieve violation status through either our Python or C++ APIs. Only available for MI300 series+ ASICs.

• Updated API amdsmi_get_violation_status() structure and CLI amdsmi_violation_status_t to include GFX Clk below host limit.

• Updated API amdsmi_get_gpu_vram_info() structure and CLI amd-smi static --vram.

Removed#

• Removed GFX_BUSY_ACC from amd-smi metric --usage as it did not provide helpful output to users.

Optimized#

• Added additional help information to amd-smi set --help command. The subcommands now detail what values are accepted as input.

• Modified amd-smi CLI to allow case insensitive arguments if the argument does not begin with a single dash.

• Converted xgmi read and write from KBs to dynamically selected readable units.

Resolved issues#

• Fixed amdsmi_get_gpu_asic_info and amd-smi static --asic not displaying graphics version correctly for Instinct MI200 series, Instinct MI100 series, and RDNA3-based GPUs.

Known issues#

• AMD SMI only reports 63 GPU devices when setting CPX on all 8 GPUs. When setting CPX as a partition mode, there is a DRM node limitation of 64.

This is a known limitation of the Linux kernel; not the driver. Other drivers, such as those using PCIe space (for example, ast), might occupy the necessary DRM nodes. You can check the number of DRM nodes using ls /sys/class/drm.

Some workaround options are as follows:

• Remove other devices occupying DRM nodes.

  Recommended steps for removing unnecessary drivers:

  1. Unload amdgpu - sudo rmmod amdgpu.

  2. Remove unnecessary driver(s) - ex. sudo rmmod ast.

  3. Reload amgpu - sudo modprobe amdgpu.

  4. Confirm amd-smi list reports all nodes (this can vary per MI ASIC).

• Update your OS kernel.

• Build and install your own kernel.

Upcoming changes#

• The AMDSMI_LIB_VERSION_YEAR enum and API fields will be deprecated in a future ROCm release.

• The pasid field in struct amdsmi_process_info_t will be deprecated in a future ROCm release.

Added#

• Support for gfx1201.

• hipBLASLt support for contiguous transpose GEMM fusion and GEMM pointwise fusions for improved performance.

• Support for hardware-specific FP8 datatypes (FP8 OCP and FP8 FNUZ).

• Support for the BF16 datatype.

• ONNX Operator Support for com.microsoft.MultiHeadAttention, com.microsoft.NhwcConv, and com.microsoft.MatMulIntgerFloat

• The migraphx-driver can now produce output for Netron.

• The migraphx-driver now includes a time parameter (similar to perf) that is more accurate for very fast kernels.

• An end-to-end Stable Diffusion 3 example with an option to disable T5 encoder on VRAM-limited GPUs has been added.

• Support to track broadcast axes in shape_transform_descriptor.

• Support for unsigned types with rocMLIR.

• Script to convert mxr files to ONNX models.

• The MIGRAPHX_SET_GEMM_PROVIDER environment variable to choose between rocBLAS and hipBLASLt. Set MIGRAPHX_SET_GEMM_PROVIDER to rocblas to use rocBLAS, or to hipblaslt to use hipBLASLt.

Changed#

• Switched to using hipBLASLt instead of rocBLAS (except for gfx90a GPU architecture).

• Included the min/max/median of the perf run as part of the summary report.

• Enabled non-packed inputs for rocMLIR.

• Always output a packed type for q/dq after determining non-packed tensors were inefficient.

• Even if using NHWC, MIGraphX will always convert group convolutions to NCHW for improved performance.

• Renamed the layout_nhwc to layout_convolution and ensured that either the weights are the same layout as the inputs or set the input and weights to NHWC.

• The minimum Cmake version is now 3.27.

Removed#

• Removed fp8e5m2fnuz rocBLAS support.

• __AMDGCN_WAVEFRONT_SIZE has been deprecated.

• Removed a warning that printed to stdout when using FP8 types.

• Remove zero-point parameter for dequantizelinear when it is zero.

Optimized#

• Prefill buffers when MLIR produces a multioutput buffer.

• Improved the resize operator performance, which should improve the overall performance of models that use it.

• Allowed the reduce operator to be split across an axis to improve fusion performance. The MIGRAPHX_SPLIT_REDUCE_SIZE environment variable has been added to allow the minimum size of the reduction to be adjusted for a possible model-specific performance improvement.

• Added MIGRAPHX_DISABLE_PASSES environment variable for debugging.

• Added MIGRAPHX_MLIR_DUMP environment variable to be set to a folder where individual final rocMLIR modules can be saved for investigation.

• Improved the C++ API to allow onnxruntime access to fp8 quantization.

Resolved issues#

• Fixed multistream execution with larger models.

• Peephole LSTM Error.

• Fixed BertSquad example that could include a broken tokenizers package.

• Fixed Attention fusion ito not error with a shape mismatch when a trailing pointwise contains a literal.

• Fixed instruction::replace() logic to handle more complex cases.

• MatMulNBits could fail with a shape error.

• Fixed an issue where some models might fail to compile with an error flatten: Shapes are not in standard layout.

Added#

• Batched CK Tile General Matrix Multiplication (GEMM) with splitK support.

• Grouped CK Tile GEMM with splitK support.

• CK Tile GEMM compute pipeline v3.

• Universal CK Tile block GEMM with interwave and intrawave schedulers .

• BF16 and INT8 WMMA GEMMs for Navi3x and Navi4x.

• Batched GEMM with output elementwise operation optimized for gfx942.

• Interwave scheduler for CK Tile GEMM mem pipeline.

• Spatially local tile partitioner in CK Tile GEMM.

• Multiple FMHA forward splitKV optimizations for decode including new N-Warp S-Shuffle pipeline.

• General FMHA forward general optimizations including refining tensor view padding configurations.

• FMHA fwd N-Warp S-Shuffle pipeline (FMHA fwd splitKV pipeline variant) .

• FMHA fwd splitKV optimization for decode (seqlen_q=1).

• hdim=96 support for FMHA forward.

• Variable-length paged KV cache support for FMHA forward.

• Paged KV cache support in group mode FMHA fwd splitKV kernels.

• Grouped convolution backward weight optimized irregular vector size loads.

• NGCHW BF16 grouped convolution forward support.

• Generic support for two-stage grouped convolution backward weight.

• Dynamic elementwise operation selected in runtime for convolutions.

• CK Tile transpose operator.

• CK Tile MOE operators: fused, sorting, and smooth quant.

• OCP FP8 support for gfx12.

• Support for FP8, BF16, FP16, OCP FP8, BF8, pk_int4 data types in CK Tile GEMM.

• Support for microscaling data types: MX FP4, FP6, and FP8.

• Support for gfx1201 target.

• Support for large batch tensors in grouped convolution backward data.

• Support for grouped convolution backward weight BF16 NGCHW.

• Support for cshuffle algorithm in CK Tile GEMM epilogue .

• Backend support for PyTorch 2.6.

• Test filters to select smoke tests or regression tests.

• Error threshold calculation for CK Tile GEMM examples.

Changed#

• Expanded code generation to support dynamic compilation using hipRTC.

• Updated attention forward qs_ks_vs pipeline to support hdim=512.

Removed#

• Removed support for gfx40 and gfx41.

Optimized#

• Improved accuracy of BFP16 convolution.

• Improved memory access pattern for all CK Tile GEMM layouts.

• Improved CK Tile Layernorm performance and added different quantization methods.

Resolved issues#

• Fixed CK Tile GEMM hotloop scheduler to use proper MFMA attributes.

Added#

• New HIP APIs

  • hipDeviceGetTexture1DLinearMaxWidth returns the maximum width of elements in a 1D linear texture, which can be allocated on the specified device.

  • hipStreamBatchMemOp enqueues an array of batch memory operations in the stream, for stream synchronization.

  • hipGraphAddBatchMemOpNode creates a batch memory operation node and adds it to a graph.

  • hipGraphBatchMemOpNodeGetParams returns the pointer of parameters from the batch memory operation node.

  • hipGraphBatchMemOpNodeSetParams sets parameters for the batch memory operation node.

  • hipGraphExecBatchMemOpNodeSetParams sets the parameters for a batch memory operation node in the given executable graph.

  • hipLinkAddData adds SPIR-V code object data to linker instance with options.

  • hipLinkAddFile adds SPIR-V code object file to linker instance with options.

  • hipLinkCreate creates linker instance at runtime with options.

  • hipLinkComplete completes linking of program and output linker binary to use with hipModuleLoadData.

  • hipLinkDestroy deletes linker instance.

Changed#

• roc-obj tools is deprecated and will be removed in an upcoming release.

  • Perl package installation is not required, and users will need to install this themselves if they want to.

  • Support for ROCm Object tooling has moved into llvm-objdump provided by package rocm-llvm.

• SDMA retainer logic is removed for engine selection in operation of runtime buffer copy.

Optimized#

• hipGraphLaunch parallelism is improved for complex data-parallel graphs.

• Make the round-robin queue selection in command scheduling. For multi-streams execution, HSA queue from null stream lock is freed and won’t occupy the queue ID after the kernel in the stream is finished.

• The HIP runtime doesn’t free bitcode object before code generation. It adds a cache, which allows compiled code objects to be reused instead of recompiling. This improves performance on multi-GPU systems.

• Runtime now uses unified copy approach:

  • Unpinned H2D copies are no longer blocking until the size of 1 MB.

  • Kernel copy path is enabled for unpinned H2D/D2H methods.

  • The default environment variable GPU_FORCE_BLIT_COPY_SIZE is set to 16, which limits the kernel copy to sizes less than 16 KB, while copies larger than that would be handled by SDMA engine.

  • Blit code is refactored, and ASAN instrumentation is cleaned up.

Resolved issues#

• Out-of-memory error on Microsoft Windows. When the user calls hipMalloc for device memory allocation while specifying a size larger than the available device memory, the HIP runtime fixes the error in the API implementation, allocating the available device memory plus system memory (shared virtual memory).

• Error of dependency on libgcc-s1 during rocm-dev install on Debian Buster. HIP runtime now uses libgcc1 for this distros.

• Stack corruption during kernel execution. HIP runtime now adds a maximum stack size limit based on the GPU device feature.

Upcoming changes#

The following lists the backward incompatible changes planned for upcoming major ROCm releases.

• Signature changes in APIs to correspond with NVIDIA CUDA APIs,

  • hiprtcCreateProgram

  • hiprtcCompileProgram

  • hipCtxGetApiVersion

• Behavior of hipPointerGetAttributes is changed to match corresponding CUDA API in version 11 and later releases.

• Return error/value code updates in the following hip APIs to match the corresponding CUDA APIs,

  • hipModuleLaunchKernel

  • hipExtModuleLaunchKernel

  • hipModuleLaunchCooperativeKernel

  • hipGetTextureAlignmentOffset

  • hipTexObjectCreate

  • hipBindTexture2D

  • hipBindTextureToArray

  • hipModuleLoad

  • hipLaunchCooperativeKernelMultiDevice

  • hipExtLaunchCooperativeKernelMultiDevice

• HIPRTC implementation, the compilation of hiprtc now uses namespace __hip_internal, instead of the standard headers std.

• Stream capture mode updates in the following HIP APIs. Streams can only be captured in relax mode, to match the behavior of the corresponding CUDA APIs,

  • hipMallocManaged

  • hipMemAdvise

  • hipLaunchCooperativeKernelMultiDevice

  • hipDeviceSetCacheConfig

  • hipDeviceSetSharedMemConfig

  • hipMemPoolCreate

  • hipMemPoolDestory

  • hipDeviceSetMemPool

  • hipEventQuery

• The implementation of hipStreamAddCallback is updated, to match the behavior of CUDA.

• Removal of hiprtc symbols from hip library.

  • hiprtc will be a independent library, and all symbols supported in HIP library are removed.

  • Any application using hiprtc APIs should link explicitly with hiprtc library.

  • This change makes the use of hiprtc library on Linux the same as on Windows, and matches the behavior of CUDA nvrtc.

• Removal of deprecated struct HIP_MEMSET_NODE_PARAMS, Developers can use definition hipMemsetParams instead.

Changed#

• Updated the build dependencies.

Resolved issues#

• Fixed the Windows reference library interface for rocSOLVER functions for hipBLAS clients.

Added#

• Support ROC-TX if HIPBLASLT_ENABLE_MARKER=1 is set.

• Output the profile logging if HIPBLASLT_LOG_MASK=64 is set.

• Support for the FP16 compute type.

• Memory bandwidth information to the hipblaslt-bench output.

• Support the user offline tuning mechanism.

• More samples.

Changed#

• Output the bench command along with the solution index if HIPBLASLT_LOG_MASK=32 is set.

Optimized#

• Improve the overall performance of the XF32/FP16/BF16/FP8/BF8 data types.

• Reduce the library size.

Resolved issues#

• Fixed multi-threads bug.

• Fixed multi-streams bug.

Added#

• Added regression tests to rtest.py. These tests recreate scenarios that have caused hardware problems in past emulation environments. Use python rtest.py [--emulation|-e|--test|-t]=regression to run these tests.

• Added extended tests to rtest.py. These tests are extra tests that did not fit the criteria of smoke and regression tests. These tests will take much longer than smoke and regression tests. Use python rtest.py [--emulation|-e|--test|-t]=extended to run these tests.

• Added ForEach, ForEachN, ForEachCopy, ForEachCopyN and Bulk functions to have parity with CUB.

• Added the hipcub::CubVector type for CUB parity.

• Added --emulation option for rtest.py

• Unit tests can be run with [--emulation|-e|--test|-t]=<test_name>;.

• Added DeviceSelect::FlaggedIf and its inplace overload.

• Added CUB macros missing from hipCUB: HIPCUB_MAX, HIPCUB_MIN, HIPCUB_QUOTIENT_FLOOR, HIPCUB_QUOTIENT_CEILING, HIPCUB_ROUND_UP_NEAREST and HIPCUB_ROUND_DOWN_NEAREST.

• Added hipcub::AliasTemporaries function for CUB parity.

Changed#

• Removed usage of std::unary_function and std::binary_function in test_hipcub_device_adjacent_difference.cpp.

• Changed the subset of tests that are run for smoke tests such that the smoke test will complete with faster run time and never exceed 2 GB of VRAM usage. Use python rtest.py [--emulation|-e|--test|-t]=smoke to run these tests.

• The rtest.py options have changed. rtest.py is now run with at least either --test|-t or --emulation|-e, but not both options.

• The NVIDIA backend now requires CUB, Thrust, and libcu++ 2.5.0. If it is not found, it will be downloaded from the NVIDIA CCCL repository.

• Changed the C++ version from 14 to 17. C++14 will be deprecated in the next major release.

Known issues#

• When building on Microsoft Windows using HIP SDK for ROCm 6.4, hipMalloc returns hipSuccess even when the size passed to it is too large and the allocation fails. Because of this, limits have been set for the maximum test case sizes for some unit tests such as HipcubDeviceRadixSort’s SortKeysLargeSizes .

Added#

• Implemented the hipfftMpAttachComm, hipfftXtSetDistribution, and hipfftXtSetSubformatDefault APIs to allow computing FFTs that are distributed between multiple MPI (Message Passing Interface) processes. These APIs can be enabled with the HIPFFT_MPI_ENABLE CMake option, which defaults to OFF. The backend FFT library called by hipFFT must support MPI for these APIs to work.

  The backend FFT library called by hipFFT must support MPI for these APIs to work.

Changed#

• Building with the address sanitizer option sets xnack+ for the relevant GPU architectures.

• Use the find_package CUDA toolkit instead of CUDA in CMake for modern CMake compatibility.

• The AMDGPU_TARGETS build variable should be replaced with GPU_TARGETS. AMDGPU_TARGETS is deprecated.

Resolved issues#

• Fixed the client packages so they depend on hipRAND instead of rocRAND.

Upcoming changes#

• The hipfc compiler wrapper has been deprecated and will be removed in a future release. Users are encouraged to directly invoke their Fortran or HIP compilers as appropriate for each source file.

Added#

• NVIDIA CUDA 12.6.3 support

• cuDNN 9.7.0 support

• cuTENSOR 2.0.2.1 support

• LLVM 19.1.7 support

• Full support for direct hipification of cuRAND into rocRAND under the --roc option.

• Support for fp8 math device/host API. For more information see #1617 in the HIPIFY Github repository.

Resolved issues#

• MIOpen support in hipify-perl under the -miopen option

• Use const_cast<const char**> for the last arguments in the hiprtcCreateProgram and hiprtcCompileProgram function calls, as in CUDA, they are of the const char* const* type

• Support for fp16 device/host API. For more information see #1769 in the HIPIFY Github repository.

• Fixed instructions on building LLVM for HIPIFY on Linux. For more information see #1800 in the HIPIFY Github repository.

Known issues#

• hipify-clang build failure against LLVM 15-18 on Ubuntu, CentOS, and Fedora. For more information see #833 in the HIPIFY Github repository.

Changed#

• When building hipRAND on Windows, use HIP_PATH (instead of the former HIP_DIR) to specify the path to the HIP SDK installation.

• When building with the rmake.py script, HIP_PATH will default to C:\hip if it is not set.

Resolved issues#

• Fixed an issue causing hipRAND build failures on Windows when the HIP SDK was installed in a location with a path that contains spaces.

Added#

• The csrlsvqr compatibility-only functions hipsolverSpScsrlsvqr, hipsolverSpDcsrlsvqr, hipsolverSpCcsrlsvqr, hipsolverSpZcsrlsvqr

Added#

• Added the azurelinux operating system name to correct the GFortran dependency.

Optimized#

• Removed an unused GTest dependency from hipsparse-bench.

Added#

• Support for alpha vector scaling

Changed#

• The check mechanism of the inputs when using alpha vector scaling

Added#

• Added benchmarking suites for contraction, permutation, and reduction. YAML files are categorized into bench and validation folders for organization.

• Added emulation test suites for contraction, permutation, and reduction.

• Support has been added for changing the default data layout using the HIPTENSOR_DEFAULT_STRIDES_COL_MAJOR environment variable.

Changed#

• GPU_TARGETS is now used instead of AMDGPU_TARGETS in cmakelists.txt.

• Binary sizes can be reduced on supported compilers by using the --offload-compress compiler flag.

Optimized#

• Optimized the hyper-parameter selection algorithm for permutation.

Resolved issues#

• For a CMake bug workaround, set CMAKE_NO_BUILTIN_CHRPATH when BUILD_OFFLOAD_COMPRESS is unset.

Upcoming changes#

• hipTensor will enhance performance and usability while unifying the API design across all operations (elementwise, reductions, and tensor contractions), enabling consistent multi-stage execution and plan reuse. As part of this change, the API functions hiptensorInitTensorDescriptor, hiptensorContractionDescriptor_t , hiptensorInitContractionDescriptor, hiptensorInitContractionFind, hiptensorContractionGetWorkspaceSize, hiptensorInitContractionPlan, hiptensorContraction, hiptensorElementwiseBinary, hiptensorElementwiseTrinary, hiptensorPermutation, and hiptensorReduction will be deprecated in a future ROCm release.

Added#

• Support for amdgpu_max_num_work_groups in the compiler. This attribute can be set by end users or library developers. It provides an upper limit for workgroups as described in AMD GPU Attributes. When set, the AMDGPU target backend might produce better machine code.

Added#

• [Conv] Enabled tuning through the miopenSetConvolutionFindMode API.

• [RNN] Added the new algorithm type miopenRNNroundedDynamic for LSTM.

• [TunaNet] Enabled NHWC for AMD Instinct MI300.

Optimized#

• Updated KernelTuningNet for CK solvers.

Resolved issues#

• Fixed tuning timing results.

• Accuracy for ASM solvers.

Changed#

• OpenCV is now installed with the package installer on Ubuntu.

• AMD Clang is now the default CXX and C compiler.

• The version of OpenMP included in the ROCm LLVM project is now used instead of libomp-dev/devel.

Known issues#

• Installation on CentOS, RedHat, and SLES requires manually installing the FFMPEG and OpenCV dev packages.

• Hardware decode requires the ROCm graphics use case.

Upcoming changes#

• Optimized audio augmentations support for VX_RPP

Added#

• Added the RCCL_SOCKET_REUSEADDR and RCCL_SOCKET_LINGER environment parameters.

• Setting NCCL_DEBUG=TRACE NCCL_DEBUG_SUBSYS=VERBS will generate traces for fifo and data ibv_post_send calls.

• Added the --log-trace flag to enable traces through the install.sh script (for example, ./install.sh --log-trace).

Changed#

• Changed compatibility to include NCCL 2.22.3.

Changed#

• AMD Clang is now the default CXX and C compiler.

Known issues#

• The package installation requires manually installing TurboJPEG.

• Installation on CentOS, RedHat, and SLES requires manually installing the FFMPEG Dev package.

• Hardware decode requires installing ROCm with the graphics use case.

Changed#

• Improved documentation

Added#

• Added ROC-TX support in rocBLAS (not available on Windows or in the static library version on Linux).

• On gfx12, all functions now support full rocblas_int dynamic range for batch_count.

• Added the --ninja build option.

• Added support for the GPU_TARGETS CMake variable.

Changed#

• The rocblas-test client removes the stress tests unless YAML-based testing or gtest_filter adds them.

• OpenMP default threading for rocBLAS clients is reduced to less than the logical core count.

• gemm_ex testing and timing reuses device memory.

• gemm_ex timing initializes matrices on device.

Optimized#

• Significantly reduced workspace memory requirements for Level 1 ILP64: iamax and iamin.

• Reduced the workspace memory requirements for Level 1 ILP64: dot, asum, and nrm2.

• Improved the performance of Level 2 gemv for the problem sizes (TransA == N && m > 2*n) and (TransA == T).

• Improved the performance of Level 3 syrk and herk for the problem size (k > 500 && n < 4000).

Resolved issues#

• gfx12: ger, geam, geam_ex, dgmm, trmm, symm, hemm, ILP64 gemm, and larger data support.

• Added a gfortran package dependency for Azure Linux OS.

• Resolved outdated SLES operating system package dependencies (cxxtools and joblib) in install.sh -d.

• Fixed code object stripping for RPM packages.

Upcoming changes#

• CMake variable AMDGPU_TARGETS is deprecated. Use GPU_TARGETS instead.

Added#

• Support for generic code object targets:

  • gfx9-generic

  • gfx9-4-generic

  • gfx10-1-generic

  • gfx10-3-generic

  • gfx11-generic

  • gfx12-generic

Changed#

• The name reported for detected agents is now based on the amdgpu.ids database provided by libdrm.

Added#

• The new bitstream reader feature has been added. The bitstream reader contains built-in stream file parsers, including an elementary stream file parser and an IVF container file parser. The reader can parse AVC, HEVC, and AV1 elementary stream files, and AV1 IVF container files. Additional supported formats will be added.

• VP9 support has been added.

• More CTests have been added: VP9 test and tests on video decode raw sample.

• Two new samples, videodecoderaw and videodecodepicfiles, have been added. videodecoderaw uses the bitstream reader instead of the FFMPEG demuxer to get picture data, and videodecodepicfiles shows how to decode an elementary video stream stored in multiple files, with each file containing bitstream data of a coded picture.

Changed#

• AMD Clang++ is now the default CXX compiler.

• Moved MD5 code out of the RocVideoDecode utility.

Removed#

• FFMPEG executable requirement for the package.

Changed#

• Building with the address sanitizer option sets xnack+ on the relevant GPU architectures and adds address-sanitizer support to runtime-compiled kernels.

• The AMDGPU_TARGETS build variable should be replaced with GPU_TARGETS. AMDGPU_TARGETS is deprecated.

Removed#

• Ahead-of-time compiled kernels for the gfx906, gfx940, and gfx941 architectures. These architectures still work the same way, but their kernels are now compiled at runtime.

• Consumer GPU architectures from the precompiled kernel cache that ships with rocFFT. rocFFT continues to ship with a cache of precompiled RTC kernels for data center and workstation architectures. As before, user-level caches can be enabled by setting the environment variable ROCFFT_RTC_CACHE_PATH to a writeable file location.

Optimized#

• Improved MPI transform performance by using all-to-all communication for global transpose operations.
  Point-to-point communications are still used when all-to-all is unavailable.

• Improved the performance of unit-strided, complex interleaved, forward, and inverse length (64,64,64) FFTs.

Resolved issues#

• Fixed incorrect results from 2-kernel 3D FFT plans that used non-default output strides. For more information, see the rocFFT GitHub issue.

• Plan descriptions can now be reused with different strides for different plans. For more information, see the rocFFT GitHub issue.

• Fixed client packages to depend on hipRAND instead of rocRAND.

• Fixed potential integer overflows during large MPI transforms.

Added#

• Roofline support for Ubuntu 24.04.

• Experimental support rocprofv3 (not enabled as default).

Resolved issues#

• Fixed PoP of VALU Active Threads.

• Workaround broken mclk for old version of rocm-smi.

Added#

• Support for debugging shaders compiled for the following generic targets:

  • gfx9-generic

  • gfx9-4-generic

  • gfx10-1-generic

  • gfx10-3-generic

  • gfx11-generic

  • gfx12-generic

Added#

• RDC policy feature

• Power and thermal throttling metrics

• RVS IET, PEBB, and memory bandwidth tests

• Link status

• RDC_FI_PROF_SM_ACTIVE metric

Changed#

• Migrated from ROCProfiler to ROCprofiler-SDK

• Improved README.md for better usability

• Moved rdc_options into share/rdc/conf/

Resolved issues#

• Fixed ABSL in clang18+

Changed#

• AMD Clang++ is now the default CXX compiler.

• The jpegDecodeMultiThreads sample has been renamed to jpegDecodePerf, and batch decoding has been added to this sample instead of single image decoding for improved performance.

Added#

• Added support for GPU metrics 1.7 to rsmi_dev_gpu_metrics_info_get().

• Added new GPU metrics 1.7 to rocm-smi --showmetrics.

Resolved issues#

• Fixed rsmi_dev_target_graphics_version_get, rocm-smi --showhw, and rocm-smi --showprod not displaying graphics version correctly for Instinct MI200 series, MI100 series, and RDNA3-based GPUs.

Added#

• Support for VA-API and rocDecode tracing.

• Aggregation of MPI data collected across distributed nodes and ranks. The data is concatenated into a single proto file.

Changed#

• Backend refactored to use ROCprofiler-SDK rather than ROCProfiler and ROCTracer.

Resolved issues#

• Fixed hardware counter summary files not being generated after profiling.

• Fixed an application crash when collecting performance counters with rocprofiler.

• Fixed interruption in config file generation.

• Fixed segmentation fault while running rocprof-sys-instrument.

• Fixed an issue where running rocprof-sys-causal or using the -I all option with rocprof-sys-sample caused the system to become non-responsive.

• Fixed an issue where sampling multi-GPU Python workloads caused the system to stop responding.

Added#

• The parallel find_first_of device function with autotuned configurations has been added. This function is similar to std::find_first_of. It searches for the first occurrence of any of the provided elements.

• Tuned configurations for segmented radix sort for gfx942 have been added to improve performance on the gfx942 architecture.

• The parallel device-level function, rocprim::adjacent_find, which is similar to the C++ Standard Library std::adjacent_find algorithm, has been added.

• Configuration autotuning has been added to device adjacent find (rocprim::adjacent_find) for improved performance on selected architectures.

• rocprim::numeric_limits has been added. This is an extension of std::numeric_limits that supports 128-bit integers.

• rocprim::int128_t and rocprim::uint128_t have been added.

• The parallel search and find_end device functions have been added. These are similar to std::search and std::find_end. These functions search for the first and last occurrence of the sequence, respectively.

• A parallel device-level function, rocprim::search_n, has been added. rocprim::search_n is similar to the C++ Standard Library std::search_n algorithm.

• New constructors, a base function, and a constexpr specifier have been added to all functions in rocprim::reverse_iterator to improve parity with the C++17 std::reverse_iterator.

• hipGraph support has been added to the device run-length-encode for non-trivial runs (rocprim::run_length_encode_non_trivial_runs).

• Configuration autotuning has been added to the device run-length-encode for non-trivial runs (rocprim::run_length_encode_non_trivial_runs) for improved performance on selected architectures.

• Configuration autotuning has been added to the device run-length-encode for trivial runs (rocprim::run_length_encode) for improved performance on selected architectures.

• The --emulation option has been added to rtest.py. Unit tests can be run with python rtest.py [--emulation|-e|--test|-t]=<test_name>.

• Extended and regression tests have been added to rtest.py. Extended tests are tests that don’t fit the criteria of smoke or regression tests, and take longer than smoke or regression tests to run. Use python rtest.py [--emulation|-e|--test|-t]=extended to run extended tests, and python rtest.py [--emulation|-e|--test|-t]=regression to run regression tests.

• Added a new type traits interface to enable users to provide additional type trait information to rocPRIM, facilitating better compatibility with custom types.

Changed#

• Changed the subset of tests that are run for smoke tests such that the smoke test will complete faster and never exceed 2 GB of VRAM usage. Use python rtest.py [--emulation|-e|--test|-t]=smoke to run these tests.

• The rtest.py options have changed. rtest.py is now run with at least either --test|-t or --emulation|-e, but not both options.

• Changed the internal algorithm of block radix sort to use a rank match. This improves the performance of various radix sort-related algorithms.

• Disabled padding in various cases where higher occupancy resulted in better performance despite more bank conflicts.

• The C++ version has changed from 14 to 17. C++14 will be deprecated in the next major release.

• You can use CMake HIP language support with CMake 3.18 and later. To use HIP language support, run cmake with -DUSE_HIPCXX=ON instead of setting the CXX variable to the path to a HIP-aware compiler.

Removed#

• HIP-CPU support

Resolved issues#

• Fixed an issue where rmake.py generated incorrect cmake commands in a Linux environment.

• Fixed an issue where rocprim::partial_sort_copy would yield a compile error if the input iterator was a const.

• Fixed incorrect 128-bit signed and unsigned integer type traits.

• Fixed a compilation issue when rocprim::radix_key_codec<...> is specialized with a 128-bit integer.

• Fixed the warp-level reduction rocprim::warp_reduce.reduce DPP implementation to avoid undefined intermediate values during the reduction.

• Fixed an issue that caused a segmentation fault when hipStreamLegacy was passed to certain API functions.

Upcoming changes#

• Using the initialization constructor of rocprim::reverse_iterator will throw a deprecation warning. It will be marked as explicit in the next major release.

Added#

• Ops 16, 32, and 64 metrics for RDC.

• Tool deprecation message for ROCProfiler and ROCProfilerV2.

Changed#

• Updated README for kernel filtration.

Resolved issues#

• Fixed the program crash issue due to invalid UTF-8 characters in a trace log.

Added#

• Support for select() operation in counter expression.

• reduce() operation for counter expression with respect to dimension.

• --collection-period feature in rocprofv3 to enable filtering using time.

• --collection-period-unit feature in rocprofv3 to control time units used in the collection period option.

• Deprecation notice for ROCProfiler and ROCProfilerV2.

• Support for rocDecode API Tracing.

• Usage documentation for ROCTx.

• Usage documentation for MPI applications.

• SDK: rocprofiler_agent_v0_t support for agent UUIDs.

• SDK: rocprofiler_agent_v0_t support for agent visibility based on gpu isolation environment variables such as ROCR_VISIBLE_DEVICES and so on.

• Accumulation VGPR support for rocprofv3.

• Host-trap based PC sampling support for rocprofv3.

• Support for OpenMP tool.

Added#

• VP9 support

Changed#

• AMD Clang is now the default CXX and C compiler.

Removed#

• All MD5 functionality, APIs, and sample code have been removed.

Resolved issues#

• Ubuntu 24.04 compile failure with FFmpeg version 5.X and above has been fixed.

Added#

• Extended tests to rtest.py. These tests are extra tests that did not fit the criteria of smoke and regression tests. They take much longer to run relative to smoke and regression tests. Use python rtest.py [--emulation|-e|--test|-t]=extended to run these tests.

• Added regression tests to rtest.py. These tests recreate scenarios that have caused hardware problems in past emulation environments. Use python rtest.py [--emulation|-e|--test|-t]=regression to run these tests.

• Added smoke test options, which run a subset of the unit tests and ensure that less than 2 GB of VRAM will be used. Use python rtest.py [--emulation|-e|--test|-t]=smoke to run these tests.

• The --emulation option for rtest.py.

Changed#

• --test|-t is no longer a required flag for rtest.py. Instead, the user can use either --emulation|-e or --test|-t, but not both.

• Removed the TBB dependency for multi-core processing of host-side generation.

Added#

• Functionality to print the associated kernel name for each wave.

Added#

• Support for asynchronous scratch reclaim on AMD Instinct MI300X GPUs. Asynchronous scratch reclaim allows scratch memory that was assigned to Command Processor(cp) queues to be reclaimed back in case the application runs out of device memory or if the hsa_amd_agent_set_async_scratch_limit API is called with the threshold parameter as 0.

Added#

• Application of a sequence of plane rotations to a given matrix for LASR

• Algorithm selection mechanism for hybrid computation

• Hybrid computation support for existing routines:

  • BDSQR

  • GESVD

Optimized#

• Improved the performance of SYEVJ.

• Improved the performance of GEQRF.

Added#

• Added support for rocsparse_matrix_type_triangular in rocsparse_spsv.

• Added test filters smoke, regression, and extended for emulation tests.

• Added rocsparse_[s|d|c|z]csritilu0_compute_ex routines for iterative ILU.

• Added rocsparse_[s|d|c|z]csritsv_solve_ex routines for iterative triangular solve.

• Added GPU_TARGETS to replace the now deprecated AMDGPU_TARGETS in CMake files.

• Added BSR format to the SpMM generic routine rocsparse_spmm.

Changed#

• By default, the rocSPARSE shared library is built using the --offload-compress compiler option which compresses the fat binary. This significantly reduces the shared library binary size.

Optimized#

• Improved the performance of rocsparse_spmm when used with row order for B and C dense matrices and the row split algorithm rocsparse_spmm_alg_csr_row_split.

• Improved the adaptive CSR sparse matrix-vector multiplication algorithm when the sparse matrix has many empty rows at the beginning or at the end of the matrix. This improves the routines rocsparse_spmv and rocsparse_spmv_ex when the adaptive algorithm rocsparse_spmv_alg_csr_adaptive is used.

• Improved stream CSR sparse matrix-vector multiplication algorithm when the sparse matrix size (number of rows) decreases. This improves the routines rocsparse_spmv and rocsparse_spmv_ex when the stream algorithm rocsparse_spmv_alg_csr_stream is used.

• Compared to rocsparse_[s|d|c|z]csritilu0_compute, the routines rocsparse_[s|d|c|z]csritilu0_compute_ex introduce several free iterations. A free iteration is an iteration that does not compute the evaluation of the stopping criteria, if enabled. This allows the user to tune the algorithm for performance improvements.

• Compared to rocsparse_[s|d|c|z]csritsv_solve, the routines rocsparse_[s|d|c|z]csritsv_solve_ex introduce several free iterations. A free iteration is an iteration that does not compute the evaluation of the stopping criteria. This allows the user to tune the algorithm for performance improvements.

• Improved the user documentation.

Resolved issues#

• Fixed an issue in rocsparse_spgemm, rocsparse_[s|d|c|z]csrgemm, and rocsparse_[s|d|c|z]bsrgemm where incorrect results could be produced when rocSPARSE was built with optimization level O0. This was caused by a bug in the hash tables that could allow keys to be inserted twice.

• Fixed an issue in the routine rocsparse_spgemm when using rocsparse_spgemm_stage_symbolic and rocsparse_spgemm_stage_numeric, where the routine would crash when alpha and beta were passed as host pointers and where beta != 0.

• Fixed an issue in rocsparse_bsrilu0, where the algorithm was running out of bounds of the bsr_val array.

Upcoming changes#

• Deprecated the rocsparse_[s|d|c|z]csritilu0_compute routines. Users should use the newly added rocsparse_[s|d|c|z]csritilu0_compute_ex routines going forward.

• Deprecated the rocsparse_[s|d|c|z]csritsv_solve routines. Users should use the newly added rocsparse_[s|d|c|z]csritsv_solve_ex routines going forward.

• Deprecated the use of AMDGPU_TARGETS in CMake files. Users should use GPU_TARGETS going forward.

Added#

• Tool deprecation message for ROCTracer.

Added#

• Added a section to install Thread Building Block (TBB) inside cmake/Dependencies.cmake if TBB is not already available.

• Made TBB an optional dependency with the new BUILD_HIPSTDPAR_TEST_WITH_TBB flag. When the flag is OFF and TBB is not already on the machine, it will compile without TBB. Otherwise, it will compile with TBB.

• Added extended tests to rtest.py. These tests are extra tests that did not fit the criteria of smoke and regression tests. These tests will take much longer than smoke and regression tests. Use python rtest.py [--emulation|-e|--test|-t]=extended to run these tests.

• Added regression tests to rtest.py. These tests recreate scenarios that have caused hardware problems in past emulation environments. Use python rtest.py [--emulation|-e|--test|-t]=regression to run these tests.

• Added smoke test options, which run a subset of the unit tests and ensure that less than 2 GB of VRAM will be used. Use python rtest.py [--emulation|-e|--test|-t]=smoke to run these tests.

• Added --emulation option for rtest.py

• Merged changes from upstream CCCL/thrust 2.4.0 and CCCL/thrust 2.5.0.

• Added find_first_of, find_end, search, and search_n to HIPSTDPAR.

• Updated HIPSTDPAR’s adjacent_find to use the rocPRIM implementation.

Changed#

• Changed the C++ version from 14 to 17. C++14 will be deprecated in the next major release.

• --test|-t is no longer a required flag for rtest.py. Instead, the user can use either --emulation|-e or --test|-t, but not both.

• Split HIPSTDPAR’s forwarding header into several implementation headers.

• Fixed copy_if to work with large data types (512 bytes).

Known issues#

• thrust::inclusive_scan_by_key might produce incorrect results when it’s used with -O2 or -O3 optimization. This is caused by a recent compiler change and a fix will be made available at a later date.

Added#

• Added interleaved layouts that enhance the performance of GEMM operations.

• Emulation test suites. These suites are lightweight and well suited for execution on emulator platforms.

Changed#

• Used GPU_TARGETS instead of AMDGPU_TARGETS in cmakelists.txt.

• Binary sizes can be reduced on supported compilers by using the --offload-compress compiler flag.

Resolved issues#

• For a CMake bug workaround, set CMAKE_NO_BUILTIN_CHRPATH when BUILD_OFFLOAD_COMPRESS is unset.

Upcoming changes#

• rocWMMA will augment the fragment API objects with additional meta-properties that improve API expressiveness and configurability of parameters including multiple-wave cooperation. As part of this change, cooperative rocWMMA API functions load_matrix_coop_sync and store_matrix_coop_sync will be deprecated in a future ROCm release.

Added#

• RPP Tensor Gaussian Filter and Tensor Box Filter support on HOST (CPU) backend.

• RPP Fog and Rain augmentation on HOST (CPU) and HIP backends.

• RPP Warp Perspective on HOST (CPU) and HIP backends.

• RPP Tensor Bitwise-XOR support on HOST (CPU) and HIP backends.

• RPP Threshold on HOST (CPU) and HIP backends.

• RPP Audio Support for Spectrogram and Mel Filter Bank on HIP backend.

Changed#

• AMD Clang is now the default CXX and C compiler.

• AMD RPP can now pass HOST (CPU) build with g++.

• Test Suite case numbers have been replaced with ENUMs for all augmentations to enhance test suite readability.

• Test suite updated to return error codes from RPP API and display them.

Resolved issues#

• CXX Compiler: Fixed HOST (CPU) g++ issues.

• Deprecation warning fixed for the sprintf is deprecated warning.

• Test suite build fix - RPP Test Suite Pre-requisite instructions updated to lock to a specific nifti_clib commit.

• Fixed broken image links for pixelate and jitter.

Added#

• Nightly builds with performance statistics.

• ASM cache capabilities for reuse.

• Virtual environment (venv) for TensileCreateLibrary invocation on Linux.

• Flag to keep build_tmp when running Tensile.

• Generalized profiling scripts.

• Support for gfx1151.

• Single-threaded support in TensileCreateLibrary.

• Logic to remove temporary build artifacts.

Changed#

• Disabled ASM cache for tests.

• Replaced Perl script with hipcc.bat as a compiler on Microsoft Windows.

• Improved CHANGELOG.md.

• Enabled external CI.

• Improved Tensile documentation.

• Refactored kernel source and header creation.

• Refactored writeKernels in TensileCreateLibrary.

• Suppressed developer warnings to simplify the Tensile output.

• Introduced an explicit cast when invoking min.

• Introduced cache abbreviations to compute kernel names.

Removed#

• OCL backend

• Unsupported tests

• Deep copy in TensileCreateLibrary

Optimized#

• Linearized ASM register search to reduce build time.

Resolved issues#

• Fixed Stream-K dynamic grid model.

• Fixed logic related to caching ASM capabilities.

• Fixed accvgpr overflow.

• Fixed test failures in SLES containers when running TensileTests.

• Fixed a regression that prevents TensileCreateLibrary from completing when fallback logic is not available.

Resolved issues#

• Fixed an error that prevented GPU hardware activity from being presented in certain workloads.

Added#

• Tracking of Heterogeneous System Architecture (HSA) handlers:

  • Adds an atomic counter to track the outstanding HSA handlers.

  • Waits on CPU for the callbacks if the number exceeds the defined value.

• Codes to capture Architected Queueing Language (AQL) packets for HIP graph memory copy node between host and device. HIP enqueues AQL packets during graph launch.

• Control to use system pool implementation in runtime commands handling. By default, it is disabled.

• A new path to avoid WaitAny calls in AsyncEventsLoop. The new path is selected by default.

• Runtime control on decrement counter only if the event is popped. There is a new way to restore dead signals cleanup for the old path.

• A new logic in runtime to track the age of events from the kernel mode driver.

Optimized#

• HSA callback performance. The HIP runtime creates and submits commands in the queue and interacts with HSA through a callback function. HIP waits for the CPU status from HSA to optimize the handling of events, profiling, commands, and HSA signals for higher performance.

• Runtime optimization which combines all logic of WaitAny in a single processing loop and avoids extra memory allocations or reference counting. The runtime won’t spin on the CPU if all events are busy.

• Multi-threaded dispatches for performance improvement.

• Command submissions and processing between CPU and GPU by introducing a way to limit the software batch size.

• Switch to std::shared_mutex in book/keep logic in streams from multiple threads simultaneously, for performance improvement in specific customer applications.

• std::shared_mutex is used in memory object mapping, for performance improvement.

Resolved issues#

• Race condition in multi-threaded producer/consumer scenario with hipMallocFromPoolAsync.

• Segmentation fault with hipStreamLegacy while using the API hipStreamWaitEvent.

• Usage of hipStreamLegacy in HIP event record.

• A soft hang in graph execution process from HIP user object. The fix handles the release of graph execution object properly considering synchronization on the device/stream. The user application now behaves the same with hipUserObject on both the AMD ROCm and NVIDIA CUDA platforms.

Added#

• Support for building with LLVM Flang.

Resolved issues#

• Fixed the exported hipfort::hipsparse CMake target.

Resolved issues#

• Fixed an error when building from source on some SUSE and RHEL systems when using the ROCPROFSYS_BUILD_DYNINST option.

Changed#

• Replaced CU_UTILIZATION metric with SIMD_UTILIZATION for better accuracy.

Resolved issues#

• Fixed the VALUBusy and SALUBusy activity metrics for accuracy on MI300.

Added#

• Support for system-wide collection of SQ counters across all HSA processes.

Changed#

• rocprofiler_sample_device_counting_service API updated to return counter output immediately, when called in synchronous mode.

Changed#

• amd-smi monitor displays VCLOCK and DCLOCK instead of ENC_CLOCK and DEC_CLOCK.

Resolved issues#

• Fixed amd-smi monitor’s reporting of encode and decode information. VCLOCK and DCLOCK are now associated with both ENC_UTIL and DEC_UTIL.

Added#

• An activeQueues set that tracks only the queues that have a command submitted to them, which allows fast iteration in waitActiveStreams.

Resolved issues#

• A deadlock in a specific customer application by preventing hipLaunchKernel latency degradation with number of idle streams.

Added#

• Support for:

  • NVIDIA CUDA 12.6.2

  • cuDNN 9.5.1

  • LLVM 19.1.3

  • Full hipBLAS 64-bit APIs

  • Full rocBLAS 64-bit APIs

Resolved issues#

• Added missing support for device intrinsics and built-ins: __all_sync, __any_sync, __ballot_sync, __activemask, __match_any_sync, __match_all_sync, __shfl_sync, __shfl_up_sync, __shfl_down_sync, and __shfl_xor_sync.

Changed#

• AMD Clang is now the default CXX and C compiler.

• The dependency on rocDecode has been removed and automatic rocDecode installation is now disabled in the setup script.

Resolved issues#

• Canny failure on Instinct MI300 has been fixed.

• Ubuntu 24.04 CTest failures have been fixed.

Known issues#

• CentOS, Red Hat, and SLES requires the manual installation of OpenCV and FFMPEG.

• Hardware decode requires that ROCm is installed with --usecase=graphics.

Upcoming changes#

• Optimized audio augmentations support for VX_RPP.

Changed#

• Enhanced the user documentation.

Resolved Issues#

• Corrected some user help strings in install.sh.

Resolved issues#

• Fixed a minor issue for users upgrading to ROCm 6.3 from 6.2 post-rename from omniperf.

Added#

• Improvements to support OMPT target offload.

Resolved issues#

• Fixed an issue with generated Perfetto files. See issue #3767 for more information.

• Fixed an issue with merging multiple .proto files.

• Fixed an issue causing GPU resource data to be missing from traces of Instinct MI300A systems.

• Fixed a minor issue for users upgrading to ROCm 6.3 from 6.2 post-rename from omnitrace.

Added#

• SIMD_UTILIZATION metric.

• New ROCm Data Center (RDC) ops metrics.

Added#

• Support for amd-smi metric --ecc & amd-smi metric --ecc-blocks on Guest VMs.

• Support for GPU metrics 1.6 to amdsmi_get_gpu_metrics_info()

• New violation status outputs and APIs: amdsmi_status_t amdsmi_get_violation_status(), amd-smi metric  --throttle, and amd-smi monitor --violation. This feature is only available on MI300+ ASICs

• Ability to view XCP (Graphics Compute Partition) activity within amd-smi metric --usage. Partition-specific features are only available on MI300+ ASICs

• Added LC_PERF_OTHER_END_RECOVERY CLI output to amd-smi metric --pcie and updated amdsmi_get_pcie_info() to include this value. This feature is only available on MI300+ ASICs

• Ability to retrieve a set of GPUs that are nearest to a given device at a specific link type level

  • Added amdsmi_get_link_topology_nearest() function to amd-smi C and Python Libraries.

• More supported utilization count types to amdsmi_get_utilization_count()

• amd-smi set -L/--clk-limit ... command. This is equivalent to rocm-smi’s --extremum command which sets sclk’s or mclk’s soft minimum or soft maximum clock frequency.

• Unittest functionality to test amdsmi API calls in Python

• GPU memory overdrive percentage to amd-smi metric -o

  • Added amdsmi_get_gpu_mem_overdrive_level() function to AMD SMI C and Python Libraries.

• Ability to retrieve connection type and P2P capabilities between two GPUs

  • Added amdsmi_topo_get_p2p_status() function to amd-smi C and Python Libraries.

  • Added retrieving P2P link capabilities to CLI amd-smi topology.

• New amdsmi_kfd_info_t type and added information under amd-smi list

• Subsystem device ID to amd-smi static --asic. There are no underlying changes to amdsmi_get_gpu_asic_info.

• Target_Graphics_Version to amd-smi static --asic and amdsmi_get_gpu_asic_info().

Changed#

• Updated BDF commands to use KFD SYSFS for BDF: amdsmi_get_gpu_device_bdf(). This change aligns BDF output with ROCm SMI.

• Moved Python tests directory path install location.

  • /opt/<rocm-path>/share/amd_smi/pytest/.. to /opt/<rocm-path>/share/amd_smi/tests/python_unittest/..

  • Removed PyTest dependency. Python testing now depends on the unittest framework only.

• Changed the power parameter in amdsmi_get_energy_count() to energy_accumulator.

  • Changes propagate forwards into the Python interface as well. Backwards compatibility is maintained.

• Updated Partition APIs and struct information and added partition_id to amd-smi static --partition.

  • As part of an overhaul to partition information, some partition information will be made available in the amdsmi_accelerator_partition_profile_t.

  • This struct will be filled out by a new API, amdsmi_get_gpu_accelerator_partition_profile().

  • Future data from these APIs will eventually be added to amd-smi partition.

Removed#

• amd-smi reset --compute-partition and ... --memory-partition and associated APIs

  • This change is part of the partition redesign. Reset functionality will be reintroduced in a later update.

  • Associated APIs include amdsmi_reset_gpu_compute_partition() and amdsmi_reset_gpu_memory_partition()

• Usage of _validate_positive is removed in parser and replaced with _positive_int and _not_negative_int as appropriate.

  • This will allow 0 to be a valid input for several options in setting CPUs where appropriate (for example, as a mode or NBIOID).

Optimized#

• Adjusted ordering of gpu_metrics calls to ensure that pcie_bw values remain stable in amd-smi metric & amd-smi monitor.

  • With this change additional padding was added to PCIE_BW amd-smi monitor --pcie

Known issues#

• See AMD SMI manual build issue.

Resolved issues#

• Improved Offline install process and lowered dependency for PyYAML.

• Fixed CPX not showing total number of logical GPUs.

• Fixed incorrect implementation of the Python API amdsmi_get_gpu_metrics_header_info().

• amdsmitst TestGpuMetricsRead now prints metric in correct units.

Upcoming changes#

• Python API for amdsmi_get_energy_count() will deprecate the power field in a future ROCm release and use energy_accumulator field instead.

• New memory and compute partition APIs will be added in a future ROCm release.

  • These APIs will be updated to fully populate the CLI and allowing compute (accelerator) partitions to be set by profile ID.

  • One API will be provided, to reset both memory and compute (accelerator).

  • The following APIs will remain:

    amdsmi_status_t
    amdsmi_set_gpu_compute_partition(amdsmi_processor_handle processor_handle,
                                      amdsmi_compute_partition_type_t compute_partition);
    amdsmi_status_t
    amdsmi_get_gpu_compute_partition(amdsmi_processor_handle processor_handle,
                                      char *compute_partition, uint32_t len);
    amdsmi_status_t
    amdsmi_get_gpu_memory_partition(amdsmi_processor_handle processor_handle,
    
                                      char *memory_partition, uint32_t len);
    amdsmi_status_t
    amdsmi_set_gpu_memory_partition(amdsmi_processor_handle processor_handle,
                                      amdsmi_memory_partition_type_t memory_partition);
    

• amd-smi set --compute-partition "SPX/DPX/CPX..." will no longer be supported in a future ROCm release.

  • This is due to aligning with Host setups and providing more robust partition information through the APIs outlined above. Furthermore, new APIs which will be available on both BM/Host can set by profile ID.

• Added a preliminary amd-smi partition command.

  • The new partition command can display GPU information, including memory and accelerator partition information.

  • The command will be at full functionality once additional partition information from amdsmi_get_gpu_accelerator_partition_profile() has been implemented.

Added#

• New HIP APIs:

  • hipGraphExecGetFlags returns the flags on executable graph.

  • hipGraphNodeSetParams updates the parameters of a created node.

  • hipGraphExecNodeSetParams updates the parameters of a created node on an executable graph.

  • hipDrvGraphMemcpyNodeGetParams gets a memcpy node’s parameters.

  • hipDrvGraphMemcpyNodeSetParams sets a memcpy node’s parameters.

  • hipDrvGraphAddMemFreeNode creates a memory free node and adds it to a graph.

  • hipDrvGraphExecMemcpyNodeSetParams sets the parameters for a memcpy node in the given graphExec.

  • hipDrvGraphExecMemsetNodeSetParams sets the parameters for a memset node in the given graphExec.

Changed#

• Un-deprecated HIP APIs:

  • hipHostAlloc

  • hipFreeHost

Optimized#

• Disabled CPU wait in device synchronize to avoid idle time in applications such as Hugging Face models and PyTorch.

• Optimized multi-threaded dispatches to improve performance.

• Limited the software batch size to control the number of command submissions for runtime to handle efficiently.

• Optimizes HSA callback performance when a large number of events are recorded by multiple threads and submitted to multiple GPUs.

Resolved issues#

• Soft hang in runtime wait event when run TensorFlow.

• Memory leak in the API hipGraphInstantiate when kernel is launched using hipExtLaunchKernelGGL with event.

• Memory leak when the API hipGraphAddMemAllocNode is called.

• The _sync() version of crosslane builtins such as shfl_sync(), __all_sync() and __any_sync(), continue to be hidden behind the preprocessor macro HIP_ENABLE_WARP_SYNC_BUILTINS, and will be enabled unconditionally in the next ROCm release.

Added#

• Level 3 functions have an additional ILP64 API for both C and Fortran (_64 name suffix) with int64_t function arguments

Changed#

• amdclang is used as the default compiler instead of g++.

• Added a dependency on the hipblas-common package.

Added#

• Support for the V2 CPP extension API for backward compatibility

• Support for data type INT8 in with INT8 out

• Support for data type FP32/FP64 for gfx110x

• Extension API hipblaslt_ext::matmulIsTuned

• Output atol and rtol for hipblaslt-bench validation

• Output the bench command for the hipblaslt CPP ext API path if HIPBLASLT_LOG_MASK=32 is set

• Support odd sizes for FP8/BF8 GEMM

Changed#

• Reorganized and added more sample code.

• Added a dependency with the hipblas-common package and removed the dependency with the hipblas package.

Optimized#

• Support fused kernel for HIPBLASLT_MATMUL_DESC_AMAX_D_POINTER for the FP8/BF8 data type

• Improved the library loading time.

• Improved the overall performance of the first returned solution.

Upcoming changes#

• The V1 CPP extension API will be deprecated in a future release of hipBLASLt.

Added#

• Support for large indices in hipcub::DeviceSegmentedReduce::* has been added, with the exception of DeviceSegmentedReduce::Arg*. Although rocPRIM’s backend provides support for all reduce variants, CUB does not support large indices in DeviceSegmentedReduce::Arg*. For this reason, large index support is not available for hipcub::DeviceSegmentedReduce::Arg*.

Changed#

• Changed the default value of rmake.py -a to default_gpus. This is equivalent to gfx906:xnack-,gfx1030,gfx1100,gfx1101,gfx1102.

• The NVIDIA backend now requires CUB, Thrust, and libcu++ 2.3.2.

Resolved issues#

• Fixed an issue in rmake.py where the list storing cmake options would contain individual characters instead of a full string of options.

• Fixed an issue where config.hpp was not included in all hipCUB headers, resulting in build errors.

Changed#

• The AMD backend is now compiled using amdclang++ instead of hipcc. The NVIDIA CUDA backend still uses hipcc-nvcc.

• CLI11 replaces Boost Program Options as the command line parser for clients.

• Building with the address sanitizer option sets xnack+ for the relevant GPU architectures.

Added#

• Added ROC-TX to the hipfort interfaces.

Changed#

• Updated the hipSOLVER bindings.

Added#

• CUDA 12.6.1 support

• cuDNN 9.5.0 support

• LLVM 19.1.1 support

• rocBLAS 64-bit APIs support

• Initial support for direct hipification of cuDNN into MIOpen under the --roc option

• Initial support for direct hipification of cuRAND into rocRAND under the --roc option

• Added a filtering ability for the supplementary hipification scripts

Resolved issues#

• Correct roc header files support

Known issues#

• Support for fp8 data types

Changed#

• Updated the default value for the -a argument from rmake.py to gfx906:xnack-,gfx1030,gfx1100,gfx1101,gfx1102.

Known issues#

• In ROCm 6.3.0, the hipRAND package version is incorrectly set to 2.11.0. In ROCm 6.2.4, the hipRAND package version was 2.11.1. The hipRAND version number will be corrected in a future ROCm release.

Resolved issues#

• Fixed an issue in rmake.py where the list storing the CMake options would contain individual characters instead of a full string of options.

Added#

• Auxiliary functions:

  • hipsolverSetDeterministicMode, hipsolverGetDeterministicMode

• Compatibility-only functions:

  • potrf

    • hipsolverDnXpotrf_bufferSize

    • hipsolverDnXpotrf

  • potrs

    • hipsolverDnXpotrs

  • geqrf

    • hipsolverDnXgeqrf_bufferSize

    • hipsolverDnXgeqrf

Changed#

• Binaries in debug builds no longer have a -d suffix.

• Changed rocSPARSE and SuiteSparse to be runtime dependencies by default. The BUILD_WITH_SPARSE CMake option can still be used to convert them into build-time dependencies (now off by default).

• The --no-sparse option for the install script now only affects the hipSOLVER clients and their dependency on hipSPARSE. Use the BUILD_HIPSPARSE_TESTS CMake option to enable tests for the hipsolverSp API (on by default).

Upcoming changes#

• The Fortran bindings provided in hipsolver_module.f90 have been deprecated. The Fortran bindings provided by the hipfort project are recommended instead.

Added#

• Added an alpha version of the hipsparse-bench executable to facilitate comparing NVIDIA CUDA cuSPARSE and rocSPARSE backends.

Changed#

• Changed the default compiler from hipcc to amdclang in the install script and CMake files.

• Improved the user documentation.

Resolved issues#

• Fixed the gfortran dependency for the Azure Linux operating system.

Known issues#

• In hipsparseSpSM_solve(), the external buffer is passed as a parameter. This does not match the NVIDIA CUDA cuSPARSE API. This extra external buffer parameter will be removed in a future release. For now, this extra parameter can be ignored and nullptr passed as it is unused internally by hipsparseSpSM_solve().

Added#

• Support for a new data type combination: INT8 inputs, BF16 output, and INT32 Matrix Core accumulation

• Support for row-major memory order (HIPSPARSE_ORDER_ROW)

Changed#

• Changed the default compiler to amdclang++.

Upcoming changes#

• hipsparseLtDatatype_t is deprecated and will be removed in the next major release of ROCm. hipDataType should be used instead.

Added#

• Added support for tensor reduction, including APIs, CPU reference, unit tests, and documentation

Changed#

• ASAN builds only support xnack+ targets.

• ASAN builds use -mcmodel=large to accommodate library sizes greater than 2GB.

• Updated the permute backend to accommodate changes to element-wise operations.

• Updated the actor-critic implementation.

• Various documentation formatting updates.

Optimized#

• Split kernel instances to improve build times.

Resolved issues#

• Fixed a bug in randomized tensor input data generation.

• Fixed the default strides calculation to be in column-major order.

• Fixed a small memory leak by properly destroying HIP event objects in tests.

• Default strides calculations now follow column-major convention.

Resolved issues#

• Fixed an issue where the compiler would incorrectly compile a program that used the __shfl(var, srcLane, width) function when one of the parameters to the function is undefined along some path to the function. See issue #3499 on GitHub.

Added#

• Initial code to run on Windows

• Support for FP8 and INT4

• Support for the Log2 internal operator

• Support for the GCC 14 compiler

• The BitwiseAnd, Scan, SoftmaxCrossEntropyLoss, GridSample, and NegativeLogLikelihoodLoss ONNX operators

• The MatMulNBits, QuantizeLinear/DequantizeLinear, GroupQueryAttention, SkipSimplifiedLayerNormalization, and SimpliedLayerNormalizationMicrosoft Contrib operators

• Dynamic batch parameter support to OneHot operator

• Split-K as an optional performance improvement

• Scripts to validate ONNX models from the ONNX Model Zoo

• GPU Pooling Kernel

• --mlir flag the migraphx-driver program to offload entire module to MLIR

• Fusing split-reduce with MLIR

• Multiple outputs for the MLIR + Pointwise fusions

• Pointwise fusions with MLIR across reshape operations

• MIGRAPHX_MLIR_DUMP environment variable to dump MLIR modules to MXRs

• The 3 option to MIGRAPHX_TRACE_BENCHMARKING to print the MLIR program for improved debug output

• MIGRAPHX_ENABLE_HIPBLASLT_GEMM environment variable to call hipBLASLt libraries

• MIGRAPHX_VERIFY_DUMP_DIFF to improve the debugging of accuracy issues

• reduce_any and reduce_all options to the Reduce operation via Torch MIGraphX

• Examples for RNNT, and ControlNet

Changed#

• Switched to MLIR’s 3D Convolution operator.

• MLIR is now used for Attention operations by default on gfx942 and newer ASICs.

• Names and locations for VRM specific libraries have changed.

• Use random mode for benchmarking GEMMs and convolutions.

• Python version is now printed with an actual version number.

Removed#

• Disabled requirements for MIOpen and rocBLAS when running on Windows.

• Removed inaccurate warning messages when using exhaustive-tune.

• Remove the hard coded path in MIGRAPHX_CXX_COMPILER allowing the compiler to be installed in different locations.

Optimized#

• Improved:

  • Infrastructure code to enable better Kernel fusions with all supported data types

  • Subsequent model compile time by creating a cache for already performant kernels

  • Use of Attention fusion with models

  • Performance of the Softmax JIT kernel and of the Pooling operator

  • Tuning operations through a new 50ms delay before running the next kernel

  • Performance of several convolution-based models through an optimized NHWC layout

  • Performance for the FP8 datatype

  • GPU utilization

  • Verification tools

  • Debug prints

  • Documentation, including gpu-driver utility documentation

  • Summary section of the migraphx-driver perf command

• Reduced model compilation time

• Reordered some compiler passes to allow for more fusions

• Preloaded tiles into LDS to improve performance of pointwise transposes

• Exposed the external_data_path property in onnx_options to set the path from onnxruntime

Resolved issues#

• Fixed a bug with gfx1030 that overwrote dpp_reduce.

• Fixed a bug in 1-arg dynamic reshape that created a failure.

• Fixed a bug with dot_broadcast and inner_broadcast that caused compile failures.

• Fixed a bug where some configs were failing when using exhaustive-tune.

• Fixed the ROCm Install Guide URL.

• Fixed an issue while building a whl package due to an apostrophe.

• Fixed the BERT Squad example requirements file to support different versions of Python.

• Fixed a bug that stopped the Vicuna model from compiling.

• Fixed failures with the verify option of migraphx-driver that would cause the application to exit early.

Added#

• [RNN] LSTM forward pass

• [Mha] Mask is added for forward pass

• [GLU] Gated Linear Unit (this is an experimental feature)

• [PReLU] Implemented PReLU backward pass (this is an experimental feature)

Optimized#

• MI300 TunaNet Update: CK forward pass and WRW Solvers updated

Resolved issues#

• Fixed unset stream when calling hipMemsetAsync.

• Fixed a memory leak issue caused by an incorrect transpose in find 2.0. See PR #3285 on GitHub.

• Fixed a memcopy data race by replacing hipMemcpy with hipMemcpyWithStream.

Changed#

• rocDecode is no longer installed by the setup script.

• The rocDecode dependency has been removed from the package installation.

Known issues#

• See MIVisionX memory access fault in Canny edge detection.

• Package installation requires the manual installation of OpenCV.

• Installation on CentOS/RedHat/SLES requires the manual installation of the FFMPEG Dev package.

• Hardware decode requires installation with --usecase=graphics in addition to --usecase=rocm.

Upcoming changes#

• Optimized audio augmentations support for VX_RPP

Added#

• MSCCL++ integration for specific contexts

• Performance collection to rccl_replayer

• Tuner Plugin example for Instinct MI300

• Tuning table for a large number of nodes

• Support for amdclang++

• New Rome model

Changed#

• Compatibility with NCCL 2.21.5

• Increased channel count for MI300X multi-node

• Enabled MSCCL for single-process multi-threaded contexts

• Enabled CPX mode for MI300X

• Enabled tracing with rocprof

• Improved version reporting

• Enabled GDRDMA for Linux kernel 6.4.0+

Resolved issues#

• Fixed an issue where, on systems running Linux kernel 6.8.0 such as Ubuntu 24.04, Direct Memory Access (DMA) transfers between the GPU and NIC were disabled, impacting multi-node RCCL performance. See issue #3772 on GitHub.

• Fixed model matching with PXN enable

Known issues#

• MSCCL is temporarily disabled for AllGather collectives.

  • This can impact in-place messages (< 2 MB) with ~2x latency.

  • Older RCCL versions are not impacted.

  • This issue will be addressed in a future ROCm release.

• Unit tests do not exit gracefully when running on a single GPU.

  • This issue will be addressed in a future ROCm release.

Added#

• rocAL Pybind support for package installation has been added. To use the rocAL python module, set the PYTHONPATH: export PYTHONPATH=/opt/rocm/lib:$PYTHONPATH

• Last batch policy, pad last batch, stick to shard, and shard size support have been added for the coco, caffe, caffe2, mxnet, tf, and cifar10 image readers.

Changed#

• rocDecode is no longer installed by the setup script.

• The rocDecode dependency has been removed from the package installation.

Optimized#

• CTest has been updated.

Resolved issues#

• Test failures have been fixed.

Known issues#

• The package installation requires the manual installation of TurboJPEG and RapidJSON.

• CentOS/RedHat/SLES requires the manual installation of the FFMPEG Dev package.

• Hardware decode requires installation with --usecase=graphics in addition to --usecase=rocm.

Upcoming changes#

• Optimized audio augmentations support.

Changed#

• The default compiler has been changed from hipcc to amdclang in the installation script and cmake files.

• Changed the address sanitizer build targets. Now only gfx908:xnack+, gfx90a:xnack+, gfx940:xnack+, gfx941:xnack+, and gfx942:xnack+ are built with BUILD_ADDRESS_SANITIZER=ON.

Resolved issues#

• Fixed hang in RS-AMG for Navi on some specific matrix sparsity patterns.

• Fixed wrong results in Apply on multi-GPU setups.

Added#

• Level 3 and EX functions have an additional ILP64 API for both C and Fortran (_64 name suffix) with int64_t function arguments

Changed#

• amdclang is used as the default compiler instead of hipcc

• Internal performance scripts use AMD SMI instead of the deprecated ROCm SMI

Optimized#

• Improved performance of Level 2 gbmv

• Improved performance of Level 2 gemv for float and double precisions for problem sizes (TransA == N && m==n && m % 128 == 0) measured on a gfx942 GPU

Resolved issues#

• Fixed the stbsv_strided_batched_64 Fortran binding

Upcoming changes#

• rocblas_Xgemm_kernel_name APIs are deprecated

Added#

• Support for setting precise ALU exception reporting

Changed#

• Clang is now the default CXX compiler.

• The new minimum supported version of va-api is 1.16.

• New build and runtime options have been added to the rocDecode-setup.py setup script.

Removed#

• Make tests have been removed. CTEST is now used for both Make tests and package tests.

• mesa-amdgpu-dri-drivers has been removed as a dependency on RHEL and SLES.

Resolved issues#

• Fixed a bug in the size of output streams in the videoDecodeBatch sample.

Added#

• rocfft-test now includes a --smoketest option.

• Implemented experimental APIs to allow computing FFTs on data distributed across multiple MPI ranks. These APIs can be enabled with the ROCFFT_MPI_ENABLE CMake option. This option defaults to OFF.

  When ROCFFT_MPI_ENABLE is ON:

  • rocfft_plan_description_set_comm can be called to provide an MPI communicator to a plan description, which can then be passed to rocfft_plan_create. Each rank calls rocfft_field_add_brick to specify the layout of data bricks on that rank.

  • An MPI library with ROCm acceleration enabled is required at build time and at runtime.

Changed#

• Compilation uses amdclang++ instead of hipcc.

• CLI11 replaces Boost Program Options as the command line parser for clients and samples.

• Building with the address sanitizer option sets xnack+ on relevant GPU architectures and address-sanitizer support is added to runtime-compiled kernels.

Added#

• Support for precise ALU exception reporting for supported architectures. Precise ALU exceptions reporting is controlled with the following commands:

  • set amdgpu precise-alu-exceptions

  • show amdgpu precise-alu-exceptions

Changed#

• The sysroot or solib-search-path settings can now be used to locate files containing GPU code objects when opening a core dump. This allows opening GPU code objects on systems different from the one where the core dump was generated.

Resolved issues#

• Fixed possible hangs when opening some AMDGPU core dumps in ROCgdb.

• Addressed cases where the roccoremerge utility improperly handled LOAD segment copy from the host core dump to the combined core dump.

Changed#

• Renamed to ROCm Compute Profiler from Omniperf.

  • New package name: rocprofiler-compute

  • New repository: ROCm/rocprofiler-compute

  • New binary name: rocprof-compute

Known issues#

• See ROCm Compute Profiler post-upgrade.

• See ROCm Compute Profiler CTest failure in CI.

Added#

• RVS integration

• Real time logging for diagnostic command

• --version command

• XGMI_TOTAL_READ_KB and XGMI_TOTAL_WRITE_KB monitoring metrics

Known issues#

• See ROCm Data Center Tool incorrect RHEL9 package version.

Added#

• Added rsmi_dev_memory_partition_capabilities_get which returns driver memory partition capablities.
  Driver now has the ability to report what the user can set memory partition modes to. User can now see available memory partition modes upon an invalid argument return from memory partition mode set (rsmi_dev_memory_partition_set).

• Support for GPU metrics 1.6 to rsmi_dev_gpu_metrics_info_get(). Updated rsmi_dev_gpu_metrics_info_get() and structure rsmi_gpu_metrics_t to include new fields for PVIOL / TVIOL, XCP (Graphics Compute Partitions) stats, and pcie_lc_perf_other_end_recovery.

• Ability to view raw GPU metrics using rocm-smi --showmetrics.

Changed#

• Added back in C++ tests for memorypartition_read_write

• Updated rsmi_dev_memory_partition_set to not return until a successful restart of AMD GPU Driver.

• All APIs now have the ability to catch driver reporting invalid arguments.

Removals#

• Removed --resetcomputepartition, and --resetmemorypartition options and associated APIs.

  • This change is part of the partition feature redesign.

  • The related APIs rsmi_dev_compute_partition_reset() and rsmi_dev_memory_partition_reset().

Resolved issues#

• Fixed rsmi_dev_target_graphics_version_get, rocm-smi --showhw, and rocm-smi --showprod not displaying properly for MI2x or Navi 3x ASICs.

Upcoming changes#

• C++ tests for memorypartition_read_write are to be re-enabled in a future ROCm release.

Changed#

• Renamed to ROCm Systems Profiler from Omnitrace.

  • New package name: rocprofiler-systems

  • New repository: ROCm/rocprofiler-systems

  • Reset the version to 0.1.0

  • New binary prefix: rocprof-sys-*

Known issues#

• See ROCm Systems Profiler post-upgrade.

Added#

• Support for hipBLASLT blas library and option to select blas library in conf file.

Changed#

• Babel parameters made runtime configurable.

Known issues#

• See ROCm Validation Suite needs specified configuration file.

Added#

• The --test smoke option has been added to rtest.py. When rtest.py is called with this option it runs a subset of tests such that the total test time is 5 minutes. Use python3 ./rtest.py --test smoke or python3 ./rtest.py -t smoke to run the smoke test.

• The --seed option has been added to run_benchmarks.py. The --seed option specifies a seed for the generation of random inputs. When the option is omitted, the default behavior is to use a random seed for each benchmark measurement.

• Added configuration autotuning to device partition (rocprim::partition, rocprim::partition_two_way, and rocprim::partition_three_way), to device select (rocprim::select, rocprim::unique, and rocprim::unique_by_key), and to device reduce by key (rocprim::reduce_by_key) to improve performance on selected architectures.

• Added rocprim::uninitialized_array to provide uninitialized storage in local memory for user-defined types.

• Added large segment support for rocprim:segmented_reduce.

• Added a parallel nth_element device function similar to std::nth_element. nth_element places elements that are smaller than the nth element before the nth element, and elements that are bigger than the nth element after the nth element.

• Added deterministic (bitwise reproducible) algorithm variants rocprim::deterministic_inclusive_scan, rocprim::deterministic_exclusive_scan, rocprim::deterministic_inclusive_scan_by_key, rocprim::deterministic_exclusive_scan_by_key, and rocprim::deterministic_reduce_by_key. These provide run-to-run stable results with non-associative operators such as float operations, at the cost of reduced performance.

• Added a parallel partial_sort and partial_sort_copy device functions similar to std::partial_sort and std::partial_sort_copy. partial_sort and partial_sort_copy arrange elements such that the elements are in the same order as a sorted list up to and including the middle index.

Changed#

• Changed the default value of rmake.py -a to default_gpus. This is equivalent to gfx906:xnack-,gfx1030,gfx1100,gfx1101,gfx1102.

• Modified the input size in device adjacent difference benchmarks. Observed performance with these benchmarks might be different.

• Changed the default seed for device_benchmark_segmented_reduce.

Removed#

• rocprim::thread_load() and rocprim::thread_store() have been deprecated. Use dereference() instead.

Resolved issues#

• Fixed an issue in rmake.py where the list storing cmake options would contain individual characters instead of a full string of options.

• Resolved an issue in rtest.py where it crashed if the build folder was created without release or debug subdirectories.

• Resolved an issue with rtest.py on Windows where passing an absolute path to --install_dir caused a FileNotFound error.

• rocPRIM functions are no longer forcefully inlined on Windows. This significantly reduces the build time of debug builds.

• block_load, block_store, block_shuffle, block_exchange, and warp_exchange now use placement new instead of copy assignment (operator=) when writing to local memory. This fixes the behavior of custom types with non-trivial copy assignments.

• Fixed a bug in the generation of input data for benchmarks, which caused incorrect performance to be reported in specific cases. It may affect the reported performance for one-byte types (uint8_t and int8_t) and instantiations of custom_type. Specifically, device binary search, device histogram, device merge and warp sort are affected.

• Fixed a bug for rocprim::merge_path_search where using unsigned offsets would produce incorrect results.

• Fixed a bug for rocprim::thread_load and rocprim::thread_store where float and double were not cast to the correct type, resulting in incorrect results.

• Resolved an issue where tests were failing when they were compiled with -D_GLIBCXX_ASSERTIONS=ON.

• Resolved an issue where algorithms that used an internal serial merge routine caused a memory access fault that resulted in potential performance drops when using block sort, device merge sort (block merge), device merge, device partial sort, and device sort (merge sort).

• Fixed memory leaks in unit tests due to missing calls to hipFree() and the incorrect use of hipGraphs.

• Fixed an issue where certain inputs to block_sort_merge(), device_merge_sort_merge_path(), device_merge(), and warp_sort_stable() caused an assertion error during the call to serial_merge().

Added#

• JSON output plugin for rocprofv2. The JSON file matches Google Trace Format making it easy to load on Perfetto, Chrome tracing, or Speedscope. For Speedscope, use --disable-json-data-flows option as speedscope doesn’t work with data flows.

• --no-serialization flag to disable kernel serialization when rocprofv2 is in counter collection mode. This allows rocprofv2 to avoid deadlock when profiling certain programs in counter collection mode.

• FP64_ACTIVE and ENGINE_ACTIVE metrics to AMD Instinct MI300 accelerator

• New HIP APIs with struct defined inside union.

• Early checks to confirm the eligibility of ELF file in ATT plugin

• Support for kernel name filtering in rocprofv2

• Barrier bit to read and stop packets

Changed#

• Extended lifetime for proxy queues

• Setting the trace-start option for rocprof to off now disables kernel tracing

• libpciaccess-dev functions now load with dlopen

• PcieAccessApi* api and void* libpciaccess_handle are now initialized to nullptr

Removed#

• Obsolete BSD and GPL licenses

• libsystemd-dev from CMakeLists.txt

Optimized#

• ROCProfiler Performance improved to reduce profiling time for large workloads of counter collection

Resolved issues#

• Bandwidth measurement in AMD Instinct MI300 accelerator

• Perfetto plugin issue of roctx trace not getting displayed

• --help for counter collection

• Signal management issues in queue.cpp

• Perfetto tracks for multi-GPU

• Perfetto plugin usage with rocsys

• Incorrect number of columns in the output CSV files for counter collection and kernel tracing

• The ROCProfiler hang issue when running kernel trace, thread trace, or counter collection on Iree benchmark for AMD Instinct MI300 accelerator

• Build errors thrown during parsing of unions

• The system hang caused while running --kernel-trace with Perfetto for certain applications

• Missing profiler records issue caused while running --trace-period

• The hang issue of ProfilerAPITest of runFeatureTests on AMD Instinct MI300 accelerator

• Segmentation fault on Navi32

Added#

• Start and end timestamp columns to the counter collection csv output

• Check to force tools to initialize context id with zero

• Support to specify hardware counters for collection using rocprofv3 as rocprofv3 --pmc [COUNTER [COUNTER ...]]

Changed#

• --marker-trace option for rocprofv3 now supports the legacy ROC-TX library libroctx64.so when the application is linked against the new library librocprofiler-sdk-roctx.so

• Replaced deprecated hipHostMalloc and hipHostFree functions with hipExtHostAlloc and hipFreeHost for ROCm versions starting 6.3

• Updated rocprofv3 --help options

• Changed naming of “agent profiling” to a more descriptive “device counting service”. To convert existing tool or user code to the new name, use the following sed:

  find . -type f -exec sed -i 's/rocprofiler_agent_profile_callback_t/rocprofiler_device_counting_service_callback_t/g; s/rocprofiler_configure_agent_profile_counting_service/rocprofiler_configure_device_counting_service/g; s/agent_profile.h/device_counting_service.h/g; s/rocprofiler_sample_agent_profile_counting_service/rocprofiler_sample_device_counting_service/g' {} +
  

• Changed naming of “dispatch profiling service” to a more descriptive “dispatch counting service”. To convert existing tool or user code to the new names, the following sed can be used:

  -type f -exec sed -i -e 's/dispatch_profile_counting_service/dispatch_counting_service/g' -e 's/dispatch_profile.h/dispatch_counting_service.h/g' -e 's/rocprofiler_profile_counting_dispatch_callback_t/rocprofiler_dispatch_counting_service_callback_t/g' -e 's/rocprofiler_profile_counting_dispatch_data_t/rocprofiler_dispatch_counting_service_data_t/g'  -e 's/rocprofiler_profile_counting_dispatch_record_t/rocprofiler_dispatch_counting_service_record_t/g' {} +
  

• FETCH_SIZE metric on gfx94x now uses TCC_BUBBLE for 128B reads

• PMC dispatch-based counter collection serialization is now per-device instead of being global across all devices

Removed#

• gfx8 metric definitions

• rocprofv3 installation from sbin directory

Resolved issues#

• Introduced subdirectory creation when rocprofv3 --output-file used to specify a folder path

• Fixed misaligned stores (undefined behavior) for buffer records

• Fixed crash when only scratch reporting is enabled

• Fixed MeanOccupancy metrics

• Fixed aborted-application validation test to properly check for hipExtHostAlloc command

• Fixed implicit reduction of SQ and GRBM metrics

• Fixed support for derived counters in reduce operation

• Bug fixed in max-in-reduce operation

• Introduced fix to handle a range of values for select() dimension in expressions parser

• Fixed Navi3x kernel tracing issues by setting the conditional aql::set_profiler_active_on_queue only when counter collection is registered

Added#

• RGB and YUV pytorch tensors

• Python distribution wheel (.whl)

• Multiple usecase samples

Changed#

• Clang replaces hipcc as the default CXX compiler.

Removed#

• Make tests have been removed. CTEST is now used for both Make tests and package tests.

Optimized#

• Setup script - build and runtime install options

• Prerequisite installation helper Python scripts

• Same GPU memory viewed as pytorch tensor

Resolved issues#

• Fixed setup issues.

Added#

• Added host generator for MT19937

• Support for rocrand_generate_poisson in hipGraphs

• Added engine, distribution, mode, throughput_gigabytes_per_second, and lambda columns for the csv format in benchmark_rocrand_host_api and benchmark_rocrand_device_api. To see these new columns, set --benchmark_format=csv or --benchmark_out_format=csv --benchmark_out="outName.csv".

Changed#

• Updated the default value for the -a argument from rmake.py to gfx906:xnack-,gfx1030,gfx1100,gfx1101,gfx1102.

• rocrand_discrete for MTGP32, LFSR113 and ThreeFry generators now uses the alias method, which is faster than binary search in CDF.

Resolved issues#

• Fixed an issue in rmake.py where the list storing the CMake options would contain individual characters instead of a full string of options.

Added#

• 64-bit APIs for existing functions:

  • LACGV_64

  • LARF_64

  • LARFG_64

  • GEQR2_64 (with batched and strided_batched versions)

  • GEQRF_64 (with batched and strided_batched versions)

  • POTF2_64 (with batched and strided_batched versions)

  • POTRF_64 (with batched and strided_batched versions)

  • POTRS_64 (with batched and strided_batched versions)

Changed#

• The rocSPARSE library is now an optional dependency at runtime. If rocSPARSE is not available, rocSOLVER’s sparse refactorization and solvers functions will return rocblas_status_not_implemented.

Optimized#

• Improved the performance of LARFG, LARF, and downstream functions such as GEQR2 and GEQRF on wave64 architectures

• Improved the performance of BDSQR and GESVD

• Improved the performance of STEDC and divide and conquer Eigensolvers

Resolved issues#

• Fixed a memory allocation issue in SYEVJ that could cause failures on clients that manage their own memory.

• Fixed a synchronizarion issue with SYEVJ that could led to a convergence failure for large matrices.

• Fixed a convergence issue in STEIN stemming from numerical orthogonality of the initial choice of eigenvectors.

• Fixed a synchronization issue in STEIN.

Known issues#

• A known issue in STEBZ can lead to errors in routines based on bisection to compute eigenvalues for symmetric/Hermitian matrices (for example, SYEVX/HEEVX and SYGVX/HEGVX), as well as singular values (for example, BDSVDX and GESVDX).

Added#

• rocsparse_create_extract_descr, rocsparse_destroy_extract_descr, rocsparse_extract_buffer_size, rocsparse_extract_nnz, and rocsparse_extract APIs to allow extraction of the upper or lower part of sparse CSR or CSC matrices.

Changed#

• Change the default compiler from hipcc to amdclang in install script and CMake files.

• Change address sanitizer build targets so that only gfx908:xnack+, gfx90a:xnack+, gfx940:xnack+, gfx941:xnack+, and gfx942:xnack+ are built when BUILD_ADDRESS_SANITIZER=ON is configured.

Optimized#

• Improved user documentation

Resolved issues#

• Fixed the csrmm merge path algorithm so that diagonal is clamped to the correct range.

• Fixed a race condition in bsrgemm that could on rare occasions cause incorrect results.

• Fixed an issue in hyb2csr where the CSR row pointer array was not being properly filled when n=0, coo_nnz=0, or ell_nnz=0.

• Fixed scaling in rocsparse_Xhybmv when only performing y=beta*y, for example, where alpha==0 in y=alpha*Ax+beta*y.

• Fixed rocsparse_Xgemmi failures when the y grid dimension is too large. This occurred when n >= 65536.

• Fixed the gfortran dependency for the Azure Linux operating system.

Added#

• Merged changes from upstream CCCL/thrust 2.3.2

  • Only the NVIDIA backend uses tuple and pair types from libcu++, other backends continue to use the original Thrust implementations and hence do not require libcu++ (CCCL) as a dependency.

• Added the thrust::hip::par_det execution policy to enable bitwise reproducibility on algorithms that are not bitwise reproducible by default.

Changed#

• Changed the default value of rmake.py -a to default_gpus. This is equivalent to gfx906:xnack-,gfx1030,gfx1100,gfx1101,gfx1102.

• Enabled the upstream (thrust) test suite for execution by default. It can be disabled by using the -DENABLE_UPSTREAM_TESTS=OFF cmake option.

Resolved issues#

• Fixed an issue in rmake.py where the list storing cmake options would contain individual characters instead of a full string of options.

• Fixed the HIP backend not passing TestCopyIfNonTrivial from the upstream (thrust) test suite.

• Fixed tests failing when compiled with -D_GLIBCXX_ASSERTIONS=ON.

Added#

• Added OCP F8/BF8 datatype support

Changed#

• Optimized some aos<->soa transforms with half-rotation offsets

• Refactored the rocBLAS reference entry point for validation and benchmarking

• ROCWMMA_* preprocessor configurations are now all assigned values

• Updated the default architecture targets for ASAN builds

• Updated the actor-critic implementation

Resolved issues#

• Fixed a bug in F64 validation due to faulty typecasting

• Fixed a bug causing runtime compilation errors with hipRTC

• Various documentation updates and fixes

Added#

• RPP Glitch and RPP Pixelate have been added to the HOST and HIP backend.

• The following audio support was added to the HIP backend:

  • Resample

  • Pre-emphasis filter

  • Down-mixing

  • To Decibels

  • Non-silent region

Changed#

• Test prerequisites have been updated.

• AMD advanced build flag.

Removed#

• Older versions of TurboJPEG have been removed.

Optimized#

• Updated the test suite.

Resolved issues#

• macOS build

• RPP Test Suite: augmentations fix

• Copy: bugfix for NCDHW layout

• MIVisionX compatibility fix: Resample and pre-emphasis filter

Known issues#

• Package installation only supports the HIP backend.

Upcoming changes#

• Optimized audio augmentations

Added#

• Testing and documentation for MasterSolutionLibrary.ArchitectureIndexMap and remapSolutionIndicesStartingFrom

• Functions for writing master file

• tPrint and reconcile printing options

• Python unit test coverage report

• Factor embed library logic into function and test

• clang++ as cxx compiler option for Windows

• Logic to cope with different compilers -toFile function to include generateManifest and moved to utilities

• Profiling CI job

• Support for amdclang and use defaults

• Architecture management functions in TensileCreateLibrary

• TensileCreateLibrary CLI reference docs

• New documentation for sphinx prototype and build out skeleton

• Contributor and developer guide

• Prediction model for optimal number of Stream-K tiles to run

  • Two-tile algorithm with Stream-K after DP

  • Atomic two-tile Stream-K and clean-up tuning parameters

  • Using glob to find logic files in TensileCreateLibrary

  • Function to confirm supported compiler rather than raw logic

Changed#

• Improved rocBLAS build output by allowing warning suppression, ignoring developer warnings, displaying progress bar and quiet printing

• Reordered extensions for Windows in which function

• updated amdclang++ and asm directories

• Updated duplicate marking tests with mocks

• Restored print ordering

• Print option

• Bumped rocm-docs-core from 1.2.0 to 1.5.0 in /docs/sphinx

• Refactored kernel duplicate matching

• Refactored generateLogicDataAndSolutions

• Restricted XCC mapping to gfx942

• Refactored argument parsing in TensileCreateLibrary

• Disabled failing rhel9 tests

• Changed line length to 100 characters for formatting

• Changed YAML operations to use C libyaml backend

• Improved warning text

• Updated clang support for Windows

• Updated supportedCompiler function

• Clang support on Windows to require use of conditional choices and defaults

• Refactored sanity check in TensileCreateLibrary

• Moved client config logic from TensileCreateLibrary main into createClientConfig

• Updated verifyManifest in TensileCreateLibrary

• Updated RTD configs

• Cleaned up CMake to avoid redundant work during client builds

• Updated Stream-K debug settings

Removed#

• Deprecated flag from CI profiling job

• Diagnostic print

• Globals from prepAsm

• Deprecated package-library option

• Duplicate which function and minor cleanup

Optimized#

• To optimize the performance of Stream-K kernels:

  • Introduced analytical grid size prediction model

  • Remapped XCC-based workgroup

Resolved issues#

• Fixed stream-K XCC configs for gfx942

• Updated WMMA capability command for ISA 10+

• Fixed progress bar character encoding error on Windows

• Fixed solution redundancy removal

• Fixed tuning imports for pyyaml

• Fixed printing of ASM capabilities for ROCm versions prior to 6.3

• Fixed code objects by filtering kernels with build errors and unprocessed kernels

• Fixed fully qualified std::get in contraction solutions

• Fixed add -v flag and change system invocation

• Used conditional imports for new dependencies to fix yaml CSafe load and dump import and rich terminal print import

• Fixed comments on scalarStaticDivideAndRemainder

Resolved issues#

• Fixed support for the API calls amdsmi_get_gpu_process_isolation and amdsmi_clean_gpu_local_data, along with the amd-smi set --process-isolation <0 or 1> command. See issue #3500 on GitHub.

Optimized#

• Implemented 1D kernels for factorizable sizes greater than 1024 and less than 2048.

Resolved issues#

• Fixed plan creation failure on some even-length real-complex transforms that use Bluestein’s algorithm.

Resolved issues#

• Fixed synchronization issue in STEIN.

Changed#

• Added amd-smi static --ras on Guest VMs. Guest VMs can view enabled/disabled RAS features on Host cards.

Removed#

• Removed amd-smi metric --ecc & amd-smi metric --ecc-blocks on Guest VMs. Guest VMs do not support getting current ECC counts from the Host cards.

Resolved issues#

• Fixed TypeError in amd-smi process -G.

• Updated CLI error strings to handle empty and invalid GPU/CPU inputs.

• Fixed Guest VM showing passthrough options.

• Fixed firmware formatting where leading 0s were missing.

Resolved issues#

• Soft hang when using AMD_SERIALIZE_KERNEL

• Memory leak in hipIpcCloseMemHandle

Added#

• Added CUDA 12.5.1 support

• Added cuDNN 9.2.1 support

• Added LLVM 18.1.8 support

• Added hipBLAS 64-bit APIs support

• Added Support for math constants math_constants.h

Known issues#

Perfetto can no longer open Omnitrace proto files. Loading Perfetto trace output .proto files in the latest version of ui.perfetto.dev can result in a dialog with the message, “Oops, something went wrong! Please file a bug.” The information in the dialog will refer to an “Unknown field type.” The workaround is to open the files with the previous version of the Perfetto UI found at https://ui.perfetto.dev/v46.0-35b3d9845/#!/.

See issue #3767 on GitHub.

Known issues#

On systems running Linux kernel 6.8.0, such as Ubuntu 24.04, Direct Memory Access (DMA) transfers between the GPU and NIC are disabled and impacts multi-node RCCL performance. This issue was reproduced with RCCL 2.20.5 (ROCm 6.2.0 and 6.2.1) on systems with Broadcom Thor-2 NICs and affects other systems with RoCE networks using Linux 6.8.0 or newer. Older RCCL versions are also impacted.

This issue will be addressed in a future ROCm release.

See issue #3772 on GitHub.

Changed#

• The new version of rocAL introduces many new features, but does not modify any of the existing public API functions.However, the version number was incremented from 1.3 to 2.0. Applications linked to version 1.3 must be recompiled to link against version 2.0.

• Added development and test packages.

• Added C++ rocAL audio unit test and Python script to run and compare the outputs.

• Added Python support for audio decoders.

• Added Pytorch iterator for audio.

• Added Python audio unit test and support to verify outputs.

• Added rocDecode for HW decode.

• Added support for:

  • Audio loader and decoder, which uses libsndfile library to decode wav files

  • Audio augmentation - PreEmphasis filter, Spectrogram, ToDecibels, Resample, NonSilentRegionDetection, MelFilterBank

  • Generic augmentation - Slice, Normalize

  • Reading from file lists in file reader

  • Downmixing audio channels during decoding

  • TensorTensorAdd and TensorScalarMultiply operations

  • Uniform and Normal distribution nodes

• Image to tensor updates

• ROCm install - use case graphics removed

Known issues#

• Dependencies are not installed with the rocAL package installer. Dependencies must be installed with the prerequisite setup script provided. See the rocAL README on GitHub for details.

Removed#

• Removed Device_Memory_Allocation.pdf link in documentation.

Resolved issues#

• Fixed error/warning message during rocblas_set_stream() call.

Optimized#

• Implemented 1D kernels for factorizable sizes less than 1024.

Optimized#

• Improved handling of UnicodeEncodeErrors with non UTF-8 locales. Non UTF-8 locales were causing crashes on UTF-8 special characters.

Resolved issues#

• Fixed an issue where the Compute Partition tests segfaulted when AMDGPU was loaded with optional parameters.

Known issues#

• When setting CPX as a partition mode, there is a DRM node limit of 64. This is a known limitation when multiple drivers are using the DRM nodes. The ls /sys/class/drm command can be used to see the number of DRM nodes, and the following steps can be used to remove unnecessary drivers:

  1. Unload AMDGPU: sudo rmmod amdgpu.

  2. Remove any unnecessary drivers using rmmod. For example, to remove an AST driver, run sudo rmmod ast.

  3. Reload AMDGPU using modprobe: sudo modprobe amdgpu.

Optimized#

• Improved performance of block_reduce_warp_reduce when warp size equals block size.

Changes#

• Added amd-smi static --ras on Guest VMs. Guest VMs can view enabled/disabled RAS features on Host cards.

Removals#

• Removed amd-smi metric --ecc & amd-smi metric --ecc-blocks on Guest VMs. Guest VMs do not support getting current ECC counts from the Host cards.

Resolved issues#

• Fixed TypeError in amd-smi process -G.

• Updated CLI error strings to handle empty and invalid GPU/CPU inputs.

• Fixed Guest VM showing passthrough options.

• Fixed firmware formatting where leading 0s were missing.

Resolved issues#

• Soft hang when using AMD_SERIALIZE_KERNEL

• Memory leak in hipIpcCloseMemHandle

Changes#

• Added CUDA 12.5.1 support.

• Added cuDNN 9.2.1 support.

• Added LLVM 18.1.8 support.

• Added hipBLAS 64-bit APIs support.

• Added Support for math constants math_constants.h.

Changes#

• Enabled rocprofv1 for MI300 hardware.

• Added dependency checks on application launch.

• Updated Omniperf packaging.

• Rolled back Grafana version in Dockerfile for Angular plugin compatibility.

• Added GPU model distinction for MI300 systems.

• Refactored and updated documemtation.

Resolved issues#

• Fixed an issue with analysis output.

• Fixed issues with profiling multi-process and multi-GPU applications.

Optimizations#

• Reduced running time of Omniperf when profiling.

• Improved console logging.

Known issues#

Perfetto can no longer open Omnitrace proto files. Loading Perfetto trace output .proto files in the latest version of ui.perfetto.dev can result in a dialog with the message, “Oops, something went wrong! Please file a bug.” The information in the dialog will refer to an “Unknown field type.” The workaround is to open the files with the previous version of the Perfetto UI found at https://ui.perfetto.dev/v46.0-35b3d9845/#!/.

See issue #3767 on GitHub.

Known issues#

On systems running Linux kernel 6.8.0, such as Ubuntu 24.04, Direct Memory Access (DMA) transfers between the GPU and NIC are disabled and impacts multi-node RCCL performance. This issue was reproduced with RCCL 2.20.5 (ROCm 6.2.0 and 6.2.1) on systems with Broadcom Thor-2 NICs and affects other systems with RoCE networks using Linux 6.8.0 or newer. Older RCCL versions are also impacted.

This issue will be addressed in a future ROCm release.

See issue #3772 on GitHub.

Changed#

• The new version of rocAL introduces many new features, but does not modify any of the existing public API functions.However, the version number was incremented from 1.3 to 2.0. Applications linked to version 1.3 must be recompiled to link against version 2.0.

• Added development and test packages.

• Added C++ rocAL audio unit test and Python script to run and compare the outputs.

• Added Python support for audio decoders.

• Added Pytorch iterator for audio.

• Added Python audio unit test and support to verify outputs.

• Added rocDecode for HW decode.

• Added support for:

  • Audio loader and decoder, which uses libsndfile library to decode wav files

  • Audio augmentation - PreEmphasis filter, Spectrogram, ToDecibels, Resample, NonSilentRegionDetection, MelFilterBank

  • Generic augmentation - Slice, Normalize

  • Reading from file lists in file reader

  • Downmixing audio channels during decoding

  • TensorTensorAdd and TensorScalarMultiply operations

  • Uniform and Normal distribution nodes

• Image to tensor updates

• ROCm install - use case graphics removed

Known issues#

• Dependencies are not installed with the rocAL package installer. Dependencies must be installed with the prerequisite setup script provided. See the rocAL README on GitHub for details.

Removed#

• Removed Device_Memory_Allocation.pdf link in documentation.

Resolved issues#

• Fixed error/warning message during rocblas_set_stream() call.

Optimized#

• Implemented 1D kernels for factorizable sizes less than 1024.

Optimized#

• Improved handling of UnicodeEncodeErrors with non UTF-8 locales. Non UTF-8 locales were causing crashes on UTF-8 special characters.

Resolved issues#

• Fixed an issue where the Compute Partition tests segfaulted when AMDGPU was loaded with optional parameters.

Known issues#

• When setting CPX as a partition mode, there is a DRM node limit of 64. This is a known limitation when multiple drivers are using the DRM nodes. The ls /sys/class/drm command can be used to see the number of DRM nodes, and the following steps can be used to remove unnecessary drivers:

  1. Unload AMDGPU: sudo rmmod amdgpu.

  2. Remove any unnecessary drivers using rmmod. For example, to remove an AST driver, run sudo rmmod ast.

  3. Reload AMDGPU using modprobe: sudo modprobe amdgpu.

Optimized#

• Improved performance of block_reduce_warp_reduce when warp size equals block size.

Changed#

• Added the following functionality:

  • amd-smi dmon is now available as an alias to amd-smi monitor.

  • An optional process table under amd-smi monitor -q.

  • Handling to detect VMs with passthrough configurations in CLI tool.

  • Process Isolation and Clear SRAM functionality to the CLI tool for VMs.

  • Added Ring Hang event.

• Added macros that were in amdsmi.h to the AMD SMI Python library amdsmi_interface.py.

• Renamed amdsmi_set_gpu_clear_sram_data() to amdsmi_clean_gpu_local_data().

Removed#

• Removed throttle-status from amd-smi monitor as it is no longer reliably supported.

• Removed elevated permission requirements for amdsmi_get_gpu_process_list().

Optimized#

• Updated CLI error strings to specify invalid device type queried.

• Multiple structure updates in amdsmi.h and amdsmi_interface.py to align with host/guest.

  • Added amdsmi.h and amdsmi_interface.py.

  • amdsmi_clk_info_t struct

  • Added AMDSMI prefix to multiple structures.

• Updated dpm_policy references to soc_pstate.

• Updated amdsmi_get_gpu_board_info() product_name to fallback to pciids file.

• Updated amdsmi_get_gpu_board_info() now has larger structure sizes for amdsmi_board_info_t.

• Updated CLI voltage curve command output.

Resolved issues#

• Fixed multiple processes not being registered in amd-smi process with JSON and CSV format.

• amdsmi_get_gpu_board_info() no longer returns junk character strings.

• Fixed parsing of pp_od_clk_voltage within amdsmi_get_gpu_od_volt_info.

• Fixed Leftover Mutex deadlock when running multiple instances of the CLI tool. When running amd-smi reset --gpureset --gpu all and then running an instance of amd-smi static (or any other subcommand that access the GPUs) a mutex would lock and not return requiring either a clear of the mutex in /dev/shm or rebooting the machine.

Known issues#

• amdsmi_get_gpu_process_isolation and amdsmi_clean_gpu_local_data commands do not work. They will be supported in a future release.

See issue #3500 on GitHub.

Changed#

• Added support for:

  • Permute scale for any dimension (#1198).

  • Combined elementwise op (#1217).

  • Multi D in grouped convolution backward weight (#1280).

  • K or C equal to 1 for fp16 in grouped convolution backward weight (#1280).

  • Large batch in grouped convolution forward (#1332).

• Added CK_TILE layernorm example (#1339).

• CK_TILE-based Flash Attention 2 kernel is now merged into the upstream repository as ROCm backend.

Optimized#

• Support universal GEMM in grouped convolution forward (#1320).

• Optimizations for low M and N in grouped convolution backward weight (#1303).

• Added a functional enhancement and compiler bug fix for FlashAttention Forward Kernel.

• FP8 GEMM performance optimization and tuning (#1384).

• Added FlashAttention backward pass performance optimization (#1397).

Changed#

• Added the _sync() version of crosslane builtins such as shfl_sync(), __all_sync() and __any_sync(). These take a 64-bit integer as an explicit mask argument.

  • In HIP 6.2, these are hidden behind the preprocessor macro HIP_ENABLE_WARP_SYNC_BUILTINS, and will be enabled unconditionally in a future HIP release.

• Added new HIP APIs:

  • hipGetProcAddress returns the pointer to driver function, corresponding to the defined driver function symbol.

  • hipGetFuncBySymbol returns the pointer to device entry function that matches entry function symbolPtr.

  • hipStreamBeginCaptureToGraph begins graph capture on a stream to an existing graph.

  • hipGraphInstantiateWithParams creates an executable graph from a graph.

• Added a new flag integrated – supported in device property.

  • The integrated flag is added in the struct hipDeviceProp_t. On the integrated APU system, the runtime driver detects and sets this flag to 1, in which case the API hipDeviceGetAttribute returns enum hipDeviceAttribute_t for hipDeviceAttributeIntegrated as value 1, for integrated GPU device.

• Added initial support for 8-bit floating point datatype in amd_hip_fp8.h. These are accessible via #include <hip/hip_fp8.h>.

• Added UUID support for environment variable HIP_VISIBLE_DEVICES.

Resolved issues#

• Fixed stream capture support in HIP graphs. Prohibited and unhandled operations are fixed during stream capture in the HIP runtime.

• Fixed undefined symbol error for hipTexRefGetArray and hipTexRefGetBorderColor.

Upcoming changes#

• The _sync() version of crosslane builtins such as shfl_sync(), __all_sync(), and __any_sync() will be enabled unconditionally in a future HIP release.

Changed#

• Added a new ILP64 API for level 2 functions for both C and FORTRAN (_64 name suffix) with int64_t function arguments.

• Added a new ILP64 API for level 1 _ex functions.

• The install.sh script now invokes the rmake.py script. Made other various improvements to the build scripts.

• Changed library dependencies in the install.sh script from rocblas and rocsolver to the development packages rocblas-dev and rocsolver-dev.

• Updated Linux AOCL dependency to release 4.2 gcc build.

• Updated Windows vcpkg dependencies to release 2024.02.14.

Changed#

• Added extension APIs: *hipblasltExtAMaxWithScale.

  • GemmTuning extension parameter to set wgm by user.

• Added support for:

  • HIPBLASLT_MATMUL_DESC_AMAX_D_POINTER for FP8/BF8 datatype.

  • FP8/BF8 input, FP32/FP16/BF16/F8/BF8 output (gfx94x platform only).

  • HIPBLASLT_MATMUL_DESC_COMPUTE_INPUT_TYPE_A_EXT and HIPBLASLT_MATMUL_DESC_COMPUTE_INPUT_TYPE_B_EXT for FP16 input data type to use FP8/BF8 MFMA.

• Added support for gfx110x.

Optimized#

• Improved library loading time.

Changed#

• Split hipcc package into two packages for different hardware platforms.

• Cleaned up references to environment variables.

• Enabled hipcc and hipconfig binaries (hipcc.bin, hipconfig.bin) by default, instead of their Perl counterparts.

• Enabled function calls.

• Added support for generating packages for ROCm stack targeting static libraries.

Resolved issues#

• Implemented numerous bug fixes and quality improvements.

Changed#

• Added DeviceCopy function for parity with CUB.

• Added enum WarpExchangeAlgorithm to the rocPRIM backend, which is used as the new optional template argument for WarpExchange.

  • The potential values for the enum are WARP_EXCHANGE_SMEM and WARP_EXCHANGE_SHUFFLE.

  • WARP_EXCHANGE_SMEM stands for the previous algorithm, while WARP_EXCHANGE_SHUFFLE performs the exchange via shuffle operations.

  • WARP_EXCHANGE_SHUFFLE does not require any pre-allocated shared memory, but the ItemsPerThread must be a divisor of WarpSize.

• Added tuple.hpp which defines templates hipcub::tuple, hipcub::tuple_element, hipcub::tuple_element_t and hipcub::tuple_size.

• Added new overloaded member functions to BlockRadixSort and DeviceRadixSort that expose a decomposer argument. Keys of a custom type (key_type) can be sorted via these overloads, if an appropriate decomposer is passed. The decomposer has to implement operator(const key_type&) which returns a hipcub::tuple of references pointing to members of key_type.

• On AMD GPUs (using the HIP backend), you can now issue hipCUB API calls inside of HIP graphs, with several exceptions:

  • CachingDeviceAllocator

  • GridBarrierLifetime

  • DeviceSegmentedRadixSort

  • DeviceRunLengthEncode Currently, these classes rely on one or more synchronous calls to function correctly. Because of this, they cannot be used inside of HIP graphs.

Removed#

• Deprecated debug_synchronous in hipCUB-2.13.2, and it no longer has any effect. With this release, passing debug_synchronous to the device functions results in a deprecation warning both at runtime and at compile time.

  • The synchronization that was previously achievable by passing debug_synchronous=true can now be achieved at compile time by setting the CUB_DEBUG_SYNC (or higher debug level) or the HIPCUB_DEBUG_SYNC preprocessor definition.

  • The compile time deprecation warnings can be disabled by defining the HIPCUB_IGNORE_DEPRECATED_API preprocessor definition.

Resolved issues#

• Fixed the derivation for the accumulator type for device scan algorithms in the rocPRIM backend being different compared to CUB. It now derives the accumulator type as the result of the binary operator.

Resolved issues#

• Added hip::host as a public link library, as hipfft.h includes HIP runtime headers.

• Prevented C++ exceptions leaking from public API functions.

• Made output of hipfftXt match cufftXt in geometry and alignment for 2D and 3D FFTs.

Changed#

• Added support for:

  • NVIDIA CUDA 12.4.1

  • cuDNN 9.1.1

  • LLVM 18.1.6

• Added full hipBLASLt support.

Resolved issues#

• HIPIFY now applies reinterpret_cast for an explicit conversion between pointer-to-function and pointer-to-object; affected functions: hipFuncGetAttributes, hipFuncSetAttribute, hipFuncSetCacheConfig, hipFuncSetSharedMemConfig, hipLaunchKernel, and hipLaunchCooperativeKernel.

Changed#

• Added support for setting generator output ordering in C and C++ API.

• hiprandCreateGeneratorHost dispatches to the host generator in the rocRAND backend instead of returning with uHIPRAND_STATUS_NOT_IMPLEMENTED.

• Added options to create:

  • A host generator to the Fortran wrapper.

  • A host generator to the Python wrapper.

• Previously, for internal testing with HMM the environment variable ROCRAND_USE_HMM was used in previous versions. The environment variable is now named HIPRAND_USE_HMM.

• Static library – moved all internal symbols to namespaces to avoid potential symbol name collisions when linking.

• Device API documentation is improved in this version.

Removed#

• Removed the option to build hipRAND as a submodule to rocRAND.

• Removed references to, and workarounds for, the deprecated hcc.

• Removed support for finding rocRAND based on the environment variable ROCRAND_DIR. Use ROCRAND_PATH instead.

Resolved issues#

• Fixed a build error when using Clang++ directly due to unsupported references to amdgpu-target.

Changed#

• Added compatibility-only functions:

  • auxiliary

    • hipsolverDnCreateParams, hipsolverDnDestroyParams, hipsolverDnSetAdvOptions

  • getrf

    • hipsolverDnXgetrf_bufferSize

    • hipsolverDnXgetrf

  • getrs

    • hipsolverDnXgetrs

• Added support for building on Ubuntu 24.04 and CBL-Mariner.

• Added hip::host to roc::hipsolver usage requirements.

• Added functions

  • syevdx/heevdx

    • hipsolverSsyevdx_bufferSize, hipsolverDsyevdx_bufferSize, hipsolverCheevdx_bufferSize, hipsolverZheevdx_bufferSize

    • hipsolverSsyevdx, hipsolverDsyevdx, hipsolverCheevdx, hipsolverZheevdx

  • sygvdx/hegvdx

    • hipsolverSsygvdx_bufferSize, hipsolverDsygvdx_bufferSize, hipsolverChegvdx_bufferSize, hipsolverZhegvdx_bufferSize

    • hipsolverSsygvdx, hipsolverDsygvdx, hipsolverChegvdx, hipsolverZhegvdx

• Updated csrlsvchol to perform numerical factorization on the GPU. The symbolic factorization is still performed on the CPU.

• Renamed hipsolver-compat.h to hipsolver-dense.h.

Removed#

• Removed dependency on cblas from the hipSOLVER test and benchmark clients.

Changed#

• Added the missing hipsparseCscGet() routine.

• All internal hipSPARSE functions now exist inside a namespace.

• Match deprecations found in cuSPARSE 12.x.x when using cuSPARSE backend.

• Improved the user manual and contribution guidelines.

Resolved issues#

• Fixed SpGEMM and SpGEMM_reuse routines that were not matching cuSPARSE behavior.

Known issues#

• In hipsparseSpSM_solve(), the external buffer is currently passed as a parameter. This does not match the cuSPARSE API and this extra external buffer parameter will be removed in a future release. For now this extra parameter can be ignored and passed a nullptr as it is unused internally by hipsparseSpSM_solve().

Optimized#

• Refined test cases.

Changed#

• Added support for:

  • Tensor permutation of ranks of 2, 3, 4, 5, and 6

  • Tensor contraction of M6N6K6: M, N, K up to rank 6

• Added tests for:

  • Tensor permutation of ranks of 2, 3, 4, 5, and 6

  • Tensor contraction of M6N6K6: M, N, K up to rank 6

  • YAML parsing to support sequential parameters ordering.

• Prefer amd-llvm-devel package before system LLVM library.

• Preferred compilers changed to CC=amdclang CXX=amdclang++.

• Updated actor-critic selection for new contraction kernel additions.

• Updated installation, programmer’s guide, and API reference documentation.

Resolved issues#

• Fixed LLVM parsing crash.

• Fixed memory consumption issue in complex kernels.

• Workaround implemented for compiler crash during debug build.

• Allow random modes ordering for tensor contractions.

Changed#

• LLVM IR

  • The llvm.stacksave and llvm.stackrestore intrinsics now use an overloaded pointer type to support non-0 address spaces.

  • Added llvm.exp10 intrinsic.

• LLVM infrastructure

  • The minimum Clang version to build LLVM in C++20 configuration is now clang-17.0.6.

• TableGen

  • Added constructs for debugging TableGen files:

    • dump keyword to dump messages to standard error. See #68793.

    • !repr bang operator to inspect the content of values. See #68716.

• AArch64 backend

  • Added support for Cortex-A520, Cortex-A720 and Cortex-X4 CPUs.

• AMDGPU backend

  • llvm.sqrt.f32 is now lowered correctly. Use llvm.amdgcn.sqrt.f32 for raw instruction access.

  • Implemented llvm.stacksave and llvm.stackrestore intrinsics.

  • Implemented llvm.get.rounding.

• ARM backend

  • Added support for Cortex-M52 CPUs.

  • Added execute-only support for Armv6-M.

• RISC-V backend

  • The Zfa extension version was upgraded to 1.0 and is no longer experimental.

  • Zihintntl extension version was upgraded to 1.0 and is no longer experimental.

  • Intrinsics were added for Zk*, Zbb, and Zbc. See Scalar Bit Manipulation Extension Intrinsics in the RISC-V C API specification.

  • Default ABI with F but without D was changed to ilp32f for RV32 and to lp64f for RV64.

  • The Zvbb, Zvbc, Zvkb, Zvkg, Zvkn, Zvknc, Zvkned, Zvkng, Zvknha, Zvknhb, Zvks, Zvksc, Zvksed, Zvksg, Zvksh, and Zvkt extension version was upgraded to 1.0 and is no longer experimental. However, the C intrinsics for these extensions are still experimental. To use the C intrinsics for these extensions, -menable-experimental-extensions needs to be passed to Clang.

  • -mcpu=sifive-p450 was added.

  • CodeGen of RV32E and RV64E is supported experimentally.

  • CodeGen of ilp32e and lp64e is supported experimentally.

• X86 backend

  • Added support for the RDMSRLIST and WRMSRLIST instructions.

  • Added support for the WRMSRNS instruction.

  • Support ISA of AMX-FP16 which contains tdpfp16ps instruction.

  • Support ISA of CMPCCXADD.

  • Support ISA of AVX-IFMA.

  • Support ISA of AVX-VNNI-INT8.

  • Support ISA of AVX-NE-CONVERT.

  • -mcpu=raptorlake, -mcpu=meteorlake and -mcpu=emeraldrapids are now supported.

  • -mcpu=sierraforest, -mcpu=graniterapids and -mcpu=grandridge are now supported.

  • __builtin_unpredictable (unpredictable metadata in LLVM IR), is handled by X86 Backend. X86CmovConversion pass now respects this builtin and does not convert CMOVs to branches.

  • Add support for the PBNDKB instruction.

  • Support ISA of SHA512.

  • Support ISA of SM3.

  • Support ISA of SM4.

  • Support ISA of AVX-VNNI-INT16.

  • -mcpu=graniterapids-d is now supported.

  • The i128 type now matches GCC and clang’s __int128 type. This mainly benefits external projects such as Rust which aim to be binary compatible with C, but also fixes code generation where LLVM already assumed that the type matched and called into libgcc helper functions.

  • Support ISA of USER_MSR.

  • Support ISA of AVX10.1-256 and AVX10.1-512.

  • -mcpu=pantherlake and -mcpu=clearwaterforest are now supported.

  • -mapxf is supported.

  • Marking global variables with code_model = "small"/"large" in the IR now overrides the global code model to allow 32-bit relocations or require 64-bit relocations to the global variable.

  • The medium code model’s code generation was audited to be more similar to the small code model where possible.

• C API

  • Added LLVMGetTailCallKind and LLVMSetTailCallKind to allow getting and setting tail, musttail, and notail attributes on call instructions.

  • Added LLVMCreateTargetMachineWithOptions, along with helper functions for an opaque option structure, as an alternative to LLVMCreateTargetMachine. The option structure exposes an additional setting (that is, the target ABI) and provides default values for unspecified settings.

  • Added LLVMGetNNeg and LLVMSetNNeg for getting and setting the new nneg flag on zext instructions, and LLVMGetIsDisjoint and LLVMSetIsDisjoint for getting and setting the new disjoint flag on or instructions.

  • Added the following functions for manipulating operand bundles, as well as building call and invoke instructions that use operand bundles:

    • LLVMBuildCallWithOperandBundles

    • LLVMBuildInvokeWithOperandBundles

    • LLVMCreateOperandBundle

    • LLVMDisposeOperandBundle

    • LLVMGetNumOperandBundles

    • LLVMGetOperandBundleAtIndex

    • LLVMGetNumOperandBundleArgs

    • LLVMGetOperandBundleArgAtIndex

    • LLVMGetOperandBundleTag

  • Added LLVMGetFastMathFlags and LLVMSetFastMathFlags for getting and setting the fast-math flags of an instruction, as well as LLVMCanValueUseFastMathFlags for checking if an instruction can use such flag.

• CodeGen infrastructure

  • A new debug type isel-dump is added to show only the SelectionDAG dumps after each ISel phase (i.e. -debug-only=isel-dump). This new debug type can be filtered by function names using -filter-print-funcs=<function names>, the same flag used to filter IR dumps after each Pass. Note that the existing -debug-only=isel will take precedence over the new behavior and print SelectionDAG dumps of every single function regardless of -filter-print-funcs’s values.

• Metadata info

  • Added a new loop metadata !{!”llvm.loop.align”, i32 64}.

• LLVM tools

  • llvm-symbolizer now treats invalid input as an address for which source information is not found.

  • llvm-readelf now supports --extra-sym-info (-X) to display extra information (section name) when showing symbols.

  • llvm-readobj --elf-output-style=JSON no longer prefixes each JSON object with the file name. Previously, each object file’s output looked like "main.o":{"FileSummary":{"File":"main.o"},...} but is now {"FileSummary":{"File":"main.o"},...}. This allows each JSON object to be parsed in the same way, since each object no longer has a unique key. Tools that consume llvm-readobj’s JSON output should update their parsers accordingly.

  • llvm-objdump now uses --print-imm-hex by default, which brings its default behavior closer in line with objdump.

  • llvm-nm now supports the --line-numbers (-l) option to use debugging information to print symbols’ filenames and line numbers.

  • llvm-symbolizer and llvm-addr2line now support addresses specified as symbol names.

  • llvm-objcopy now supports --gap-fill and --pad-to options, for ELF input and binary output files only.

• LLDB

  • SBType::FindDirectNestedType function is added. It’s useful for formatters to quickly find directly nested type when it’s known where to search for it, avoiding more expensive global search via SBTarget::FindFirstType.

  • Renamed lldb-vscode to lldb-dap and updated its installation instructions to reflect this. The underlying functionality remains unchanged.

  • The mte_ctrl register can now be read from AArch64 Linux core files.

  • LLDB on AArch64 Linux now supports debugging the Scalable Matrix Extension (SME) and Scalable Matrix Extension 2 (SME2) for both live processes and core files. For details refer to the AArch64 Linux documentation.

  • LLDB now supports symbol and binary acquisition automatically using the DEBUFINFOD protocol. The standard mechanism of specifying DEBUFINOD servers in the DEBUGINFOD_URLS environment variable is used by default. In addition, users can specify servers to request symbols from using the LLDB setting plugin.symbol-locator.debuginfod.server_urls, override or adding to the environment variable.

  • When running on AArch64 Linux, lldb-server now provides register field information for the following registers: cpsr, fpcr, fpsr, svcr and mte_ctrl.

• Sanitizers

  • HWASan now defaults to detecting use-after-scope bugs.