-
Notifications
You must be signed in to change notification settings - Fork 0
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
LPGEMM <u|s>8s8s16ou8 fixes for incorrect zero point addition.
-The zero point data type is different based on the downscale data type. For int8_t downscale type, zero point type is int8_t whereas for uint8_t downscale type, it is uint8_t. During downscale post-op, the micro-kernels upscales the zero point from its data type (int8_t or uint8_t) to that of the accumulation data type and then performs the zero point addition. The accumulated output is then stored as downscaled type in a later storage phase. For the <u|s>8s8s16 micro-kernels, the upscaling to int16_t (accumulation type) is always performed assuming the zero point is int8_t using the _mm256_cvtepi8_epi16 instruction. However this will result in incorrect upscaled zero point values if the downscale type is uint8_t and the associated zero point type is also uint8_t. This issue is corrected by switching between the correct upscale instruction based on the zero point type. AMD-Internal: [SWLCSG-2500] Change-Id: I92eed4aed686c447d29312836b9e551d6dd4b076
- Loading branch information
1 parent
b3391ef
commit d184467
Showing
9 changed files
with
399 additions
and
79 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Oops, something went wrong.