# AMD backend support

## General information

Support for AMD backend is implemented via SYCL HIP backend. The feature is
disabled by default and is currently experimental. You must enable it at the
build time with a CMake option `DNNL_GPU_VENDOR=AMD`. The AMD GPUs can be used
via oneDNN engine abstraction. The engine should be created using
`dnnl::engine::kind::gpu` engine kind or you can provide `sycl::device` objects
that correspond to AMD GPUs.

## Pre-requisites
* [oneAPI DPC++ Compiler with support for HIP AMD](https://github.com/intel/llvm/blob/sycl/sycl/doc/GetStartedGuide.md#build-dpc-toolchain-with-support-for-hip-amd), version [2022-12](https://github.com/intel/llvm/releases/tag/2022-12)
* [AMD ROCm](https://github.com/RadeonOpenCompute/ROCm), version 5.3 or newer. The latest supported version currently is 6.1.
* [MIOpen](https://github.com/ROCmSoftwarePlatform/MIOpen), version 2.18 or newer (optional if AMD ROCm includes the required version of MIOpen)
* [rocBLAS](https://github.com/ROCmSoftwarePlatform/rocBLAS), version 2.45.0 or newer (optional if AMD ROCm includes the required version of rocBLAS)

## Build command

```bash
export CC=/path/to/dpcpp/install/bin/clang
export CXX=/path/to/dpcpp/install/bin/clang++
mkdir build
cd build
cmake -DDNNL_CPU_RUNTIME=SYCL -DDNNL_GPU_RUNTIME=SYCL \
      -DDNNL_GPU_VENDOR=AMD -G Ninja ..
```

If you have AMD ROCm, MIOpen or rocBLAS installed in non-standard locations or
you want to use MIOpen or rocBLAS that is not part of the AMD ROCm package then
the following CMake and environment variables can be used to specify their location:
* `MIOPENROOT`
* `HIPROOT`
* `ROCBLASROOT`

## Memory

Both buffer-based and USM-based oneDNN APIs are supported for AMD backend.

## Suported Data Types

The following table documents the supported data types. In generic this is for all primitives,
but primitive wise which datatypes are supported are mentioned under each primitive.

| Data Type | Computation Mode                      |
|-----------|---------------------------------------|
| f32       | Training, Inference                   |
| f16       | Inference                             |
| s8        | Inference (when applicable)           |
| bf16      | Training, Inference (when applicable) |


## Supported Primitives and Implementation Limitations

The AMD backend cannot provide all functionalities supported by oneDNN primitives.
because MIOpen and rocBLAS lack some features. The detailed limitations of each MIOpen
and rocBLAS based primitive are explained below.

### Binary

The `miopenOpTensor` is the equivalent of oneDNN binary primitive.

* Supported data types are `f32`, `f16`, `s32`.
* Datatypes of `SRC0`, `SRC1` and `DST` should be the same.
* Supported formats are `NCDHW`, `NCHW`, `NCW`, `NC`, `N`.
* Blocked formats are not supported.
* Only `scales` attribute is supported.
* Post-ops are not supported.
* Supported algorithms are `binary_add`, `binary_mul`, `binary_min`, `binary_max`.

### Convolution

The `miopenConvolutionForwardImmediate` is used to compute forward. 
The `miopenConvolutionBackwardDataImmediate` and `miopenConvolutionBackwardWeightsImmediate` 
are used to compute backward by data and backward by weights respectively.

The implementation supports both Forward and Backward directions:

#### Forward direction

* Supported data types combinations:
    | Source   |  Weights  |  Destination  |     Bias      |
    |-----------|-----------|---------------|---------------|
    |   f32     |    f32    |     f32       |       f32     |
    |   f16     |    f16    |     f16       |       f16     |
    |   s8      |     s8    |      s8       | Not supported |
    |   bf16    |   bf16    |     bf16      | Not supported |
* Supported formats: `NCDHW`, `NCHW`, `NCW` (with bias) and `NDHWC`, `NHWC`, `NWC` (without bias)
* Supported post-ops: `eltwise` (`eltwise_relu`, `eltwise_tanh`, `eltwise_elu`, `eltwise_logistic`) and `sum`
* Supported attributes: `scales`
* Supported algorithms : `winograd`, `direct`

#### Backward direction

* Supported data types combinations:
    |  Source   |  Weights  |  Destination  |     Bias      |
    -----------|-----------|---------------|---------------|
    |   f32     |    f32    |     f32       |       f32     |
    |   bf16     |    bf16    |     bf16       |       bf16     |
* Supported formats: `NCDHW`, `NCHW`, `NCW` (with bias) and `NDHWC`, `NHWC`, `NWC` (without bias)
* Supported algorithms : `winograd`, `direct`

#### Limitations

* Source, weights and destination tensors must have the same format
* Post-op sum scale with non-zero fractional part can lead to incorrect results
* Zero points are not supported
* Post-ops are implementated via separate operations
* Bias addition is implemented with `miopenOpTensor`

### Deconvolution

* Deconvolution primitive is implemented through the convolution with swapped input
  and output channels.
* Post-ops are not supported.

### Eltwise

The implementation supports both forward and backward directions. The 
`miopenCreateActivationDescriptor` and `miopenSetActivationDescriptor` are used
to create the activation descriptor. And the `miopenActivationForward` and 
`miopenActivationBackward` are used for the execution.

* Supported formats: `NCDHW`, `NDHWC`, `NCHW`, `NHWC`, `NCW`, `NWC`, `NC`, `N`

##### Forward Direction
* Supported algorithms: `relu`, `tanh`, `elu`, `soft_relu`, `abs`, `logistic`.
* soft_relu is only supported with `alpha = 1`.
* Supported data types are `f32` and `f16`.
* Post-ops are not supported.

##### Backward Direction
* Supported algorithms: `relu` and `soft_relu`.
* soft_relu is only supported with `alpha = 1`.
* Supported data types are `f32`.

### Softmax / Logsoftmax

The implementation supports both forward and backward directions. The primitive
was implemented using `miopenSoftmaxForward_V2` and `miopenSoftmaxBackward_V2`.

* Supported formats: `NCDHW`, `NDHWC`, `NCHW`, `NHWC`, `NCW`, `NWC`, `NC`.
* Only `axis = 1` is supported.

##### Forward Direction
* Supported data types: `f32` and `f16`.
* Post-ops are not supported.

##### Backward Direction
* Supported data types: `f32`.

### Local Response Normalization (LRN)

The implementation supports both forward and backward directions.
The `miopenCreateLRNDescriptor` and `miopenSetLRNDescriptor` are used to set the LRN desriptor.
The `miopenLRNForward` and `miopenLRNBackward` are used for the execution.

* Supported formats: `NCHW`, `NHWC`, `NCW`, `NWC`, `NC`.

##### Forward Direction
* Supported data types: `f32`, `f16`.
* Supported algorithms: `lrn_across_channels`, `lrn_within_channel`.
* `lrn_within_channel` supports only 2D spatial cases.
* Post-ops are not supported.

##### Backward Direction
* Supported data types: `f32`.
* Supported algorithms: `lrn_across_channels`, `lrn_within_channel`.
* `lrn_within_channels` supports only 2D spatial cases.

### Pooling

The Pooling primitive in the AMD backend is implemented with the following API's:
* `miopenCreatePoolingDescriptor`, `miopenSetNdPoolingDescriptor`, `miopenSetPoolingIndexType`,
and `miopenSetPoolingWorkSpaceIndexMode` are used to set the pooling descriptor
* `miopenPoolingGetWorkSpaceSizeV2` is used for getting a work space size.
* `miopenPoolingForward` and `miopenPoolingBackward` are used for the execution.

##### Forward direction

* Supported datatype for forward Training `f32`.
* Supported datatypes for forward Inference `f32`, `f16`.
* Only 1D, 2D and 3D pooling is supported.
* Only `NCDHW`, `NCHW`, `NCW` formats are supported.
* Supported algorithms are `pooling_max`, `pooling_avg_include_padding`, `pooling_avg_exclude_padding`.
* Post-ops are not supported.

##### Backward direction

* Supported datatypes are `f32`.
* Only 1D, 2D and 3D pooling is supported.
* Only `NCDHW`, `NCHW`, `NCW` formats are supported.
* Supported algorithms are `pooling_max`, `pooling_avg_include_padding`, `pooling_avg_exclude_padding`.

### Reduction

The Reduction primitive is implemented with the following API's:
* `miopenCreateReduceTensorDescriptor` and `miopenSetReduceTensorDescriptor` are used to set the reduction tensor descriptor
* `miopenGetReductionWorkspaceSize` is used for getting a workspace size.
* `miopenReduceTensor` is used for execution

* Supported datatypes are `f32`, `f16`.
* Only `NCDHW`, `NCHW`, `NCW`, `NC`, `N` formats are supported.
* Supported algorithms are `reduction_max`, `reduction_min`, `reduction_sum`, `reduction_mul`,
  `reduction_mean`, `reduction_norm_lp_sum`, `reduction_norm_lp_power_p_sum`
* `reduction_norm_lp_sum` algorithm supported only for the `p` value 2
* `reduction_norm_lp_power_p_sum` supported only for the `p` value 1
* Only `eps = 0` is supported.
* Post-opst are not supported.

### Matrix Multiplication

The matrix multiplication primitive is implemented with `rocblas_gemm_ex`
and `rocblas_gemm_strided_batched_ex` functions.

* Supported data types are `f32`, `f16`, `bf16` and `s8/s32`.
* Currently only below 5 combinations are supported:

    | Source    | Weights   | Destination   |   Bias        |
    |-----------|-----------|---------------|---------------|
    |   f32     |   f32     |   f32         |   f32         |
    |   f16     |   f16     |   f16         |   f16         |
    |   s8      |    s8     |   s32         |   s32         |
    |   bf16    |   bf16    |   bf16        | Not supported |
    |   bf16    |   bf16    |   f32         |   f32         |
    
* Blocked formats are not supported.
* Zero points are not supported.
* Scales are not supported.
* Post-op `eltwise` with `eltwise_relu`, `eltwise_tanh`, `eltwise_elu`, `eltwise_logistic` is supported
* Post-op `sum` is supported. For s8 case (for the 3rd combination in above table),
  only scales without fractional part are supported.
* Source and weights broadcasting is supported in the batched case.
* Only 1D, 2D, 3D supported.
* Supported formats are `NCW`, `NWC`, `NC`, `CN`, `N`

### Inner product

The inner product primitive is implemented with `rocblas_gemm_ex` and 
`rocblas_gemm_strided_batched_ex` functions for forward, backward data and backward weight 
and `miopenReduceTensor` for backward bias. A function called `gemm_consitency_check()`,
`dense_check()` is used to see if the backend can be used for inner product.
`reorder_check()` is used when reorder is required. `miopenActivationForward` operation is
used for eltwise operation and `miopenOpTensor` is used for bias operation. The
`beta` parameter in gemm is used for the sum scale and `alpha` parameter is used
for the output scale.

* Supported formats : `NCW`, `NC`, `CN`, `N`

##### Forward direction

* Supported data types are `f32`, `f16`, `bf16` and `s8/s32`.
* Currently only below combinations are supported:
    | Source    | Weights   | Destination   | Bias          |
    |-----------|-----------|---------------|---------------|
    |   f32     |   f32     |   f32         |   f32         |
    |   f16     |   f16     |   f16         |   f16         |
    |   s8      |   s8      |   s32         |   s32         |
    |   bf16    |   bf16    |   bf16        | Not supported |
    |   bf16    |   bf16    |   f32         |   f32         |
* Zero points support is not provided.
* Post-op eltwise with `eltwise_relu`, `eltwise_tanh`, `eltwise_elu`, `eltwise_logistic` is supported
* Post-op sum is supported. For s8 case(for third combination in above table),
  only integer sum scale values are supported
* Blocked formats are not supported.

##### Backward direction
* Supported data types are `f32`, `bf16`.
* Currently only below combinations are supported:

    | Propagation       |   Source  |  Weights  | Destination   | Bias          |
    |-------------------|-----------|-----------|---------------|---------------|
    |   Backward Data   |    f32    |   f32     |   f32         | Not supported |
    |                   |    bf16   |   bf16    |   bf16        | Not supported |
    | Backward Weights  |    f32    |   f32     |   f32         |    f32        |
    |                   |    bf16   |   bf16    |   bf16        | Not supported |
* Zero points are not supported.
* Blocked formats are not supported.

### Batch normalization

The closest equivalent to oneDNN batch normalization can be
`miopenBatchNormalizationForwardTraining`, `miopenBatchNormalizationForwardInference`
and `miopenBatchNormalizationBackward` operations. 

##### Forward direction

* When `global_stats` flag is set for batch normalization, the mean and variance
  are input only parameters. However, MIOpen does not have the option to accept
  the mean and variance as inputs in the forward training operation. Therefore,
  `miopenBatchNormalizationForwardInference` is used to match the oneDNN feature.
  Although inference is not supported without `global_stats` flags set.
* The MIOpen precision is different from that of oneDNN for Batch Normalization.
  (e.g `exp from oneDNN: 0.000629427 got from miopen: 0.000629831 diff:4.04136e-07 rdiff:0.000642069`)
* The forward training with no flags accepts mean and variance as an output.
  However, in MIOpen the mean and variance are running mean and variance
  respectably so they are both input and output variable. Therefore, they are
  required to have a sensible value (cannot be NaN). Since oneDNN will not set
  value for the mean and variance when no flag is passed, the NaN can be
  propagated as a result. To avoid NaN propagation, `hipMemsetD32Async` function is
  used to initialize the mean and variance with zero.
* MIOpen requires the values for scale and shift. When shift and scale are
  not defined in oneDNN, `hipMemsetD32Async` is used to initialize scale to 1 and shift
  to 0.
* For performance reason in the backward pass, MIOpen requires the mean and
  inverse variance to be saved in the forward pass. Therefore, when AMD
  backend is used for batch normalization, the workspace must be provided to
  save the mean and inverse variance.
* When `dnnl_fuse_norm_relu` flag is set for batch normalization, the
  `miopenActivationForward` operation is called immediately after the batch
  normalization, since MIOpen does not have a fused batch normalization with
  `RELU`. The implementation of the elementwise post operations is the same.
* When `dnnl_fuse_norm_relu` is used, the intermediate output of batch
  normalization, which is used as an input to the activation function, is saved
  in the workspace as well. This is required to compute the backward pass for
  `dnnl_fuse_norm_relu` flag.
* Forward pass supports `f32`, `f16`.
* Blocked Formats are not supported.
* Only `NCDHW`, `NCHW`, `NCW`, `NC` formats are supported.
* Elementwise post-op is supported only for eltwise_relu.

##### Backward direction

* MIOpen uses `alpha` and `beta` parameters to blend the `dy`, `shift` and
  `scale`. Since oneDNN does not have this feature, the `alpha` and `beta`
  values in the backward direction are set to 1 and 0 respectively to avoid
  blending.
* AMD backend for backward direction requires the workspace as an input
  containing the mean and inverse variance computed in the forward pass.
* The AMD backend for oneDNN does not support the backward direction for
  batch normalization when the flag is set to `global_stats`. 
* When `dnnl_fuse_norm_relu` flag is set, AMD backend requires the
  intermediate result of the batch normalization saved in the forward pass. This
  is used to compute the backward direction of the activation function used for
  `RELU`.
* Backward pass supports only `f32` data types.
* Blocked formats are not supported.
* Only `NCDHW`, `NCHW`, `NCW`, `NC` formats are supported.

### Reorder

The `miopenTransform` function is the equivalent of oneDNN reorder function.

* Per dimension scaling is not supported (a single alpha and beta value is
  accepted by the transform tensor function).
* Supported data types: `f32`

### Other Primitives

Some missing primitives/features are supported through
[generic SYCL kernels](../generic/sycl/README.md).
