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
miopenSetConvolutionAttributeAPI call and passing the convolution descriptor for the appropriate convolution operation and the
MIOPEN_CONVOLUTION_ATTRIB_FP16_ALT_IMPLconvolution 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_IMPLenvironment variable. The above variable when set to a value of
1engages the alternate implementation while a value of
0disables it. Keep in mind the environment variable impacts the convolution operation in all directions.