MI200 MFMA Behavior Specifics

The MI200 MFMA_F16, MFMA_BF16 and MFMA_BF16_1K flush subnormal input/output data to zero. This behavior might affect the convolution operation in certain workloads due to the limited exponent range of the half-precision floating point datatypes.

An alternate implementation for the half precision data-type is available in MIOpen which utilizes conversion instructions to utilizes the BFloat16 data-types larger exponent range, albeit with reduced accuracy. The following salients apply to this alternate implementation:

• It is disabled by default in the Forward convolution operations.

• It is enabled by default in the backward data and backward weights convolution operations.

• The default MIOpen behaviors described above may be overridden using the miopenSetConvolutionAttribute API call and passing the convolution descriptor for the appropriate convolution operation and the MIOPEN_CONVOLUTION_ATTRIB_FP16_ALT_IMPL convolution attribute with a non-zero value to engage the alternate implementation.

• The behavior might also be overridden using the MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL environment variable. The above variable when set to a value of 1 engages the alternate implementation while a value of 0 disables it. Keep in mind the environment variable impacts the convolution operation in all directions.