<!-- markdownlint-disable first-line-h1 -->
<!-- markdownlint-disable no-duplicate-header -->

### Release highlights for ROCm 5.7

New features include:

* A new library (hipTensor)
* Optimizations for rocRAND and MIVisionX
* AddressSanitizer for host and device code (GPU) is now available as a beta

Note that ROCm 5.7.0 is EOS for MI50. 5.7 versions of ROCm are the last major releases in the ROCm 5
series. This release is Linux-only.

:::{important}
The next major ROCm release (ROCm 6.0) will not be backward compatible with the ROCm 5 series.
Changes will include: splitting LLVM packages into more manageable sizes, changes to the HIP runtime
API, splitting rocRAND and hipRAND into separate packages, and reorganizing our file structure.
:::

#### AMD Instinct™ MI50 end-of-support notice

AMD Instinct MI50, Radeon Pro VII, and Radeon VII products (collectively gfx906 GPUs) will enter
maintenance mode starting Q3 2023.

As outlined in [5.6.0](https://rocm.docs.amd.com/en/docs-5.6.0/release.html), ROCm 5.7 will be the
final release for gfx906 GPUs to be in a fully supported state.

* ROCm 6.0 release will show MI50s as "under maintenance" for
  {doc}`Linux<rocm-install-on-linux:reference/system-requirements>` and
  {doc}`Windows<rocm-install-on-windows:reference/system-requirements>`

* No new features and performance optimizations will be supported for the gfx906 GPUs beyond this
  major release (ROCm 5.7).

* Bug fixes and critical security patches will continue to be supported for the gfx906 GPUs until Q2
  2024 (end of maintenance \[EOM] will be aligned with the closest ROCm release).

* Bug fixes during the maintenance will be made to the next ROCm point release.

* Bug fixes will not be backported to older ROCm releases for gfx906.

* Distribution and operating system updates will continue per the ROCm release cadence for gfx906
  GPUs until EOM.

#### Feature updates

##### Non-hostcall HIP printf

**Current behavior**

The current version of HIP printf relies on hostcalls, which, in turn, rely on PCIe atomics. However, PCle
atomics are unavailable in some environments, and, as a result, HIP-printf does not work in those
environments. Users may see the following error from runtime (with AMD_LOG_LEVEL 1 and above):

```shell
    Pcie atomics not enabled, hostcall not supported
```

**Workaround**

The ROCm 5.7 release introduces an alternative to the current hostcall-based implementation that
leverages an older OpenCL-based printf scheme, which does not rely on hostcalls/PCIe atomics.

:::{note}
This option is less robust than hostcall-based implementation and is intended to be a
workaround when hostcalls do not work.
:::

The printf variant is now controlled via a new compiler option -mprintf-kind=<value>. This is
supported only for HIP programs and takes the following values,

* “hostcall” – This currently available implementation relies on hostcalls, which require the system to
  support PCIe atomics. It is the default scheme.

* “buffered” – This implementation leverages the older printf scheme used by OpenCL; it relies on a
  memory buffer where printf arguments are stored during the kernel execution, and then the runtime
  handles the actual printing once the kernel finishes execution.

**NOTE**: With the new workaround:

* The printf buffer is fixed size and non-circular.  After the buffer is filled, calls to printf will not result in
  additional output.

* The printf call returns either 0 (on success) or -1 (on failure, due to full buffer), unlike the hostcall
  scheme that returns the number of characters printed.

##### Beta release of LLVM AddressSanitizer (ASan) with the GPU

The ROCm 5.7 release introduces the beta release of LLVM AddressSanitizer (ASan) with the GPU. The
LLVM ASan provides a process that allows developers to detect runtime addressing errors in
applications and libraries. The detection is achieved using a combination of compiler-added
instrumentation and runtime techniques, including function interception and replacement.

Until now, the LLVM ASan process was only available for traditional purely CPU applications. However,
ROCm has extended this mechanism to additionally allow the detection of some addressing errors on
the GPU in heterogeneous applications. Ideally, developers should treat heterogeneous HIP and
OpenMP applications like pure CPU applications. However, this simplicity has not been achieved yet.

Refer to the documentation on LLVM ASan with the GPU at
[LLVM AddressSanitizer User Guide](../conceptual/using-gpu-sanitizer.md).

:::{note}
The beta release of LLVM ASan for ROCm is currently tested and validated on Ubuntu 20.04.
:::

#### Defect fixes

The following defects are fixed in ROCm v5.7:

* Test hangs observed in HMM RCCL

* NoGpuTst test of Catch2 fails with Docker

* Failures observed with non-HMM HIP directed catch2 tests with XNACK+

* Multiple test failures and test hangs observed in hip-directed catch2 tests with xnack+

#### HIP 5.7.0

##### Optimizations

##### Additions

* Added `meta_group_size`/`rank` for getting the number of tiles and rank of a tile in the partition

* Added new APIs supporting Windows only, under development on Linux

    * `hipMallocMipmappedArray` for allocating a mipmapped array on the device

    * `hipFreeMipmappedArray` for freeing a mipmapped array on the device

    * `hipGetMipmappedArrayLevel` for getting a mipmap level of a HIP mipmapped array

    * `hipMipmappedArrayCreate` for creating a mipmapped array

    * `hipMipmappedArrayDestroy` for destroy a mipmapped array

    * `hipMipmappedArrayGetLevel` for getting a mipmapped array on a mipmapped level

##### Changes

##### Fixes

##### Known issues

* HIP memory type enum values currently don't support equivalent value to
  `cudaMemoryTypeUnregistered`, due to HIP functionality backward compatibility.
* HIP API `hipPointerGetAttributes` could return invalid value in case the input memory pointer was not
  allocated through any HIP API on device or host.

##### Upcoming changes for HIP in ROCm 6.0 release

* Removal of `gcnarch` from hipDeviceProp_t structure

* Addition of new fields in hipDeviceProp_t structure

    * maxTexture1D

    * maxTexture2D

    * maxTexture1DLayered

    * maxTexture2DLayered

    * sharedMemPerMultiprocessor

    * deviceOverlap

    * asyncEngineCount

    * surfaceAlignment

    * unifiedAddressing

    * computePreemptionSupported

    * hostRegisterSupported

    * uuid

* Removal of deprecated code -hip-hcc codes from hip code tree

* Correct hipArray usage in HIP APIs such as `hipMemcpyAtoH` and `hipMemcpyHtoA`

* HIPMEMCPY_3D fields correction to avoid truncation of "size_t" to "unsigned int" inside
  `hipMemcpy3D()`

* Renaming of 'memoryType' in `hipPointerAttribute_t` structure to 'type'

* Correct `hipGetLastError` to return the last error instead of last API call's return code

* Update `hipExternalSemaphoreHandleDesc` to add "unsigned int reserved[16]"

* Correct handling of flag values in `hipIpcOpenMemHandle` for `hipIpcMemLazyEnablePeerAccess`

* Remove `hiparray*` and make it opaque with `hipArray_t`
