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#

• The roc-obj tools have been deprecated and will be removed in a future release.

  • llvm-objdump, llvm-objcopy, and llvm-readobj will be enhanced to provide similar functionality as that provided by the roc-obj tools . The LLVM tools are available in the rocm-llvm pkg.

  • While not related to the deprecation, also note that the roc-obj tools’ package dependency on Perl has been changed to recommended. It is the user’s responsibility to install Perl to use these tools.

• 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.