- Current implementation uses macros to expand the code at
compile time, but this is causing some false warning in GCC12 and 14.
- Added switch case in trsm right variants for n_remainder.
- This ensures that n_rem is compile time constant, therefore
warnings related to array subscript out of bounds are fixed.
- mtune=znver3 flag is causing compilation issue in GCC 9.1,
therefore this flag is removed.
- Remaned the file bli_trsm_small to bli_trsm_small_zen5 in order
to avoid possibily of missing symbols.
AMD-Internal: [CPUPL-6199]
Change-Id: Ib8e90196ce0a41d38c2b29226df5ab6c2d8ba996
In preparation for merging next group of changes from upstream BLIS,
move some BLAS extension APIs to new extra subdirectories in
frame/compat and frame/compat/cblas/src. Other extension APIs will
be moved in later commits.
Some tidying up to better match upstream BLIS code has also been done.
AMD-Internal: [CPUPL-2698]
Change-Id: I0780a775d37242fba562c3f13666da0ad2b2cdfb
- Warnings in DTRSM kernel caused by uninitialized registers
and extra loop unroll is fixed.
- Warning in DGEMM kernel caused by extra space is fixed.
Change-Id: I1d9cfaa0b2847f5fdbe8b343a462d67a3aca0819
Description:
Implemented a reference version for
aocl_gemm_reorder_bf16bf16f32of32 function
to make the code cpu architecture independent.
AMD-Internal: [ SWLCSG-3279 ]
Change-Id: I0c715864c0ab3e5afea2ee6ee9546b75c3dbf9ec
Change usage of global_rntm and tl_rntm to elimate need
for mutex operations when accessing global_rntm. Usage of
these data structures is now as follows:
* global_rntm is set once during bli_init_apis and includes
all getenv calls to check BLIS threading and error printing
environment variables. global_rntm is then read-only.
* tl_rntm is intialized once from global_rntm on each
application thread. Any calls to BLIS set threading/ways
APIs will update tl_rntm for that application thread only
(Previously they updated global_rntm for all application threads).
* Re-initialize info_value in tl_rntm in every call to bli_init APIs.
* In bli_rntm_init_from_global() we initialize the local (per API
call) rntm as a copy of tl_rntm and then update threading values
in bli_thread_update_rntm_from_env() to reflect the current status
of OpenMP runtime ICVs.
AMD-Internal: [CPUPL-6168][SWLCSG-3143]
Change-Id: Ib9387ee2b51f507ed08cc38267057109acea14a6
- Added new DTRSM kernels for right and left variants.
- Kernel dimensions are 24x8.
- 24x8 DGEMM SUP kernels are used internally
for solving GEMM subproblem.
- Tuned thresholds to pick efficent code path for ZEN5.
AMD-Internal: [CPUPL-6016]
Change-Id: I743d6dc47717952c2913085c0db3454ae9d046db
- Updated the threshold check for ZGEMM small path to include
runtime checks for redirection, specific to the micro-architecture.
- The current ZGEMM small path has only its AVX2 variant available.
Post implementing an AVX512(same/different algorithm), the thresholds
will further be fine-tuned.
- Included the dot-product based AVX512 ZGEMM kernels in the ZEN5
context. It will be used as part of handling RRC and CRC storage
schemes of C, A and B matrices in both single-thread and multi-thread
runs.
AMD-Internal: [CPUPL-5949]
Change-Id: Ic8b7cf0e00b7c477f748669f160c4b01df995c75
- This patch introduces changes to support DGEMM computation when the input matrix A is transposed.
- The changes accommodate CRC (Column-Row-Column) and RRC (Row-Row-Column) storage schemes for matrices
C, A, and B. The primary goal is to pack the A matrix in a column-stored scheme, enabling the re-use
of the DGEMM SUP kernel for efficient computation.
- Performance is better when BLIS_PACK_BUFFER macro is set to 0.
By default, it is set to 1[enabled].
AMD-Internal: [CPUPL-6054]
Change-Id: I543a84b05c9e6380bc03017ab6da685e7006a64e
This patch introduces comprehensive optimizations to the DGEMM kernel, focusing on loop
efficiency and edge kernel performance. The following technical improvements have been implemented:
1. **IR Loop Optimization:**
- The IR loop has been re-implemented in hand-written assembly to eliminate the overhead associated
with `begin_asm` and `end_asm` calls, resulting in more efficient execution.
2. **JR Loop Integration:**
- The JR loop is now incorporated into the micro kernel. This integration avoids the repetitive overhead
of stack frame management for each JR iteration, thereby enhancing loop performance.
3. **Kernel Decomposition Strategy:**
- The m dimension is decomposed into specific sizes: 20, 18, 17, 16, 12, 11, 10, 9, 8, 4, 2, and 1.
- For remaining cases, masked variants of edge kernels are utilized to handle the decomposition efficiently.
1. **Interleaved Scaling by Alpha:**
- Scaling by the alpha factor is interleaved with load instructions to optimize the instruction pipeline
and reduce latency.
2. **Efficient Mask Preparation:**
- Masks are prepared within inline assembly code only at points where masked load-store operations are necessary,
minimizing unnecessary overhead.
3. **Broadcast Instruction Optimization:**
- In edge kernels where each FMA (Fused Multiply-Add) operation requires a broadcast without subsequent reuse,
the broadcast instruction is replaced with `mem_1to8`.
- This allows the compiler to optimize by assigning separate vector registers for broadcasting, thus avoiding
dependency chains and improving execution efficiency.
4. **C Matrix Update Optimization:**
- During the update of the C matrix in edge kernels, columns are pre-loaded into multiple vector registers.
This approach breaks dependency chains during FMA operations following the scaling by alpha, thereby mitigating
performance bottlenecks and enhancing throughput.
These optimizations collectively improve the performance of the DGEMM kernel, particularly in handling edge cases and
reducing overhead in critical loops. The changes are expected to yield significant performance gains in matrix multiplication
operations.
This patch also involves changes for tiny gemm interface. A light
interface for calling kernels and removing calls to avx2 dgemm kernels
as we use avx512 dgemm kernels for all the sizes for zen4 and zen5.
For zen4 and zen5 when A matrix transposed(CRC, RRC), tiny kernel does not have
the support to handle such inputs and thus such inputs are routed to
gemm_small path.
AMD-Internal: [CPUPL-6054]
Change-Id: I57b430f9969ca39aa111b54fa169e4225b900c4a
bli_nthreads_optimum is exported and called directly by AOCL libFLAME,
however it was only defined if building a multithreaded BLIS library
with AOCL_DYNAMIC enabled. Change to always define this function. If
BLIS is serial or if AOCL_DYNAMIC is disabled, this function returns
without modifying the supplied rntm.
Change-Id: Ie65690e9e6ec2a8ea77b3778f96676a68e6260be
- AVX512 specific DGEMV native kernels are added for Zen4/5
architectures to handle the NO_TRANSPOSE cases and are independent of
the AXPYF fused kernels.
- The following set of kernels biased towards the n-dimension perform
beta scaling of y vector within the kernel itself and handle cases
where n is less than 5:
- bli_dgemv_n_zen_int_32x8n_avx512( ... )
- bli_dgemv_n_zen_int_32x4n_avx512( ... )
- bli_dgemv_n_zen_int_32x2n_avx512( ... )
- bli_dgemv_n_zen_int_32x1n_avx512( ... )
- The bli_dgemv_n_zen_int_16mx8_avx512( ... ) is biased towards the
m-dimension and for this kernel beta scaling is handled beforehand
within the framework.
- Added unit-tests for the new kernels.
- AVX2 path for Zen/2/3 architectures still follows the old approach of
using fused kernel, namely AXPYF, to perform the GEMV operation.
AMD-Internal: [CPUPL-5560]
Change-Id: I22bc2a865cd28b9cdcb383e17d1ff38bdd28de79
Description
-In enum AOCL_PARAMS_STORAGE_TYPES the member FLOAT was declared and the
clang 18 compiler in msvc throwing issue with multiple definition. We
replace FLOAT and BFLOAT16 to AOCL_GEMM_<F32/BF16>.
AMD-Internal: CPUPL-6174
Change-Id: Ic061af068854d51629b82b495efd0eb54543f329
- Added a kernel selection logic based on the input
dimension(runtime parameter), to choose between
deploying AVX2 or AVX512 computational kernel for
single-thread execution.
- An empirical analysis was conducted to arrive at the
thresholds, for ZEN4 and ZEN5 architectures.
- Updated the fast-path threshold for ZEN4 to be in hand
with the tipping points of its dynamic thread-setter(used
when AOCL_DYNAMIC is enabled).
AMD-Internal: [CPUPL-5937]
Change-Id: I96d7f167658c9e25a0098c4c67e12e4ba673e228
Description:
1. AutoAWQ use a int32 buffer to store 8 elements each of 4 bits in this
format [0, 2, 4, 6, 1, 3, 5, 7].
2. Support is added to convert above format back to the original
sequential order [0, 1, 2, 3, 4, 5, 6, 7] before reordering
in the AWQ API.
AMD-Internal: SWLCSG-3169
Change-Id: I5395766060c200ab81d0b8be94356678a169ac13
Description:
1. Added group quantization and zero-point (zp) in
aocl_gemm_bf16s4f32o<bf16|f32> API.
2. Group quantization is technique to improve accuracy
where scale factors to dequantize weights varies at group
level instead of per channel and per tensor level.
3. Added zp and scaling in woq packb kernels so that for
large M values zp and scaling are performed at pack-b
stage and bf16 kernels are called
4. Adding zp support and scaling to default path in WoQ kernels
created some performance overhead when M value is very small.
5. Added string group_size to lpgemm bench to read
group size from bench_input.txt and tested for
various combinations of matrix dimensions.
6. The scalefactors could be of type float or bf16
and the zeropoint values are expected to be
in int8 format.
AMD-Internal: [SWLCSG-3168, SWLCSG-3172]
Change-Id: Iff07b54d76edc7408eb2ea0b29ce8b4a04a38f57
- Enabled dynamic blocksizes for DGEMM in ZEN4 and ZEN5 systems.
- MC, KC and NC are dynamically selected at runtime for DGEMM native.
- A local copy of cntx is created and blocksizes are updated in the local cntx.
- Updated threshold for picking DGEMM SUP kernel for ZEN4.
AMD-Internal: [CPUPL-5912]
Change-Id: Ic12a1a48bfa59af26cc17ccfa47a2a33fadde1f6
- Merged ZEN4 and ZEN5 DGEMM 8x24 kernel.
- Replaced 32x6 kernel with 8x24. Now same
kernel is used for ZEN4 and ZEN5.
- Blocksizes have been tuned for genoa only.
- DGEMM kernel for DTRSM native code path
is replaced with 8x24 kernel.
- Enabled alpha scaling during packing for ZEN4.
- ZEN4 8x24 kernel has been removed.
AMD-Internal: [CPUPL-5912]
Change-Id: I89a16a7e3355af037d21d453aabf53c5ecccb754
- Extreme values are not handled correctly when beta == 0 and C is
column major stored.
- For checking if beta is zero, VCOMISD(XMM(1), XMM(2)) is used,
beta(XMM1) is compared with zero(XMM2),
for column major C, setting of xmm2 to zero was missed.
- XMM2 is set to zero after the jump to column major stored C code
is made, this skips the setting of XMM2 to zero for column major
C.
- This is fixed by setting XMM2 to zero before the column major jump.
AMD-Internal: [CPUPL-5851]
Change-Id: Ic511071fbc82a082fa48a1543c0c7325eaf75cb8
- Changed fringe cases to use ZEN5 DGEMM kernel instead
of ZEN4 kernel.
- ASAN reporting error when RBP is used even when
-fno-stack-pointer flag is used, therefore replaced RBP
register with R11 register.
- Added missing RDX register in clobber list which is causing
failures with AOCC compiler.
Thanks to harsh.dave@amd.com for debugging some of the issues.
AMD-Internal: [CPUPL-5851]
Change-Id: I0ee412c97c9dbfb3e7a736a10bfd93d775779b5b
- Generic kernel is used if N is not multiple of NR
or M is not multiple of MR.
- This limit the maximum values of NR that can be used.
- Support for fringe case handling is added in DGEMM
macro kernel so that macro kernel can be used for
all problem sizes.
AMD-Internal: [CPUPL-5912]
Change-Id: I85c17e91d7511bb35ffed0f346d6ff0376baf62f
In the function bli_thread_update_rntm_from_env()mutex is used for reading global_rntm
"bli_pthread_mutex_lock( &global_rntm_mutex );" This causes regression when application is
Multithreaded. The cause of this regression is due to these mutexes, Imagine a scenario
two threads launched, one thread acquires this mutex, second thread stalls till mutex is
freed by first thread, as a result second thread will be slower to arrive at openmp barrier
in application thereby increasing the openmp barrier overhead.
Things get worst when more number of threads are launched.
Thanks to rocHPL for sharing standalone panelfact application to reproduce this issue.
Thanks to @Edward Symth (edward.smyth@amd.com) for finding this bug.
[SWLCSG-3143]
Description:
1. The bias type was supported only based on output data type.
2. The option is added in the pre-ops structure to select the bias data
type irrespective of the storage data type in bf16 and WoQ API's
AMD-Internal: SWLCSG-3171
Change-Id: Iac10b946c2d4a5c405b2dc857362be0058615abf
Add definitions in gtestsuite header to list available kernel
by AOCL BLIS version. Check these definitions in ukr test
programs to avoid missing symbol errors when testing with an
older version of BLIS.
Currently AOCL_41, AOCL_42, AOCL_50 and AOCL_DEV are supported,
with AOCL_DEV inferred from the version being later than the
value of AOCL_BLAS_LATEST_VERSION set in CMakeLists.txt. Thanks
to Eleni Vlachopoulou for the cmake functionality to automatically
detect the version from the library.
AMD-Internal: [CPUPL-4500]
Change-Id: I40ffd3d3789324fbb1dabfbf5e1dd4e0c94d54d9
Description:
Implemented sigmoid, tanh as fused post-ops in
aocl_gemm_<s8|u8>s8<s32|s16>o<s8|u8|s32> API's
Sigmoid(x) = 1/1+e^(-x)
Tanh(x) = (1-e^(-2x))/(1+e^(2x))
Updated bench_lpgemm to recognize sigmod, tanh
as options for post-ops from bench_input and verified.
AMD-Internal: [SWLCSG-3178]
Change-Id: I9df3aab02222f728ff9d1f292c7bc549f30176f0
Description:
Implemented sigmoid, tanh as fused post-ops in
aocl_gemm_f32f32f32of32 API's
Sigmoid(x) = 1/1+e^(-x)
Tanh(x) = (1-e^(-2x))/(1+e^(2x))
Updated bench_lpgemm to recognize sigmod, tanh
as options for post-ops from bench_input and verified.
AMD-Internal: [SWLCSG-3178]
Change-Id: Iac0a907f6dea1d9cb82d9fd8716bfdbf1c33921d
Description:
Implemented sigmoid, tanh as fused post-ops in
aocl_gemm_bf16bf16f32o<f32|bf16) API's
Sigmoid(x) = 1/1+e^(-x)
Tanh(x) = (1-e^(-2x))/(1+e^(2x))
Updated bench_lpgemm to recognize sigmod, tanh
as options for post-ops from bench_input and verified.
AMD-Internal: [SWLCSG-3178]
Change-Id: I78a3ba4a67ab63f9d671fbe315f977b016a0d969
gemm_compute currently has differences in the interface to
the blis_impl layer compared to the top-level API. Modify
gtestsuite wrapper to account for this.
AMD-Internal: [CPUPL-4500]
Change-Id: Ie96c9ac3b23128ae8e03af34ad11e65910dec594
- Implemented a set of column preferential dot-product based
ZGEMM kernels(main and fringe) in AVX512(for SUP code-path).
These kernels perform matrix multiplication as a sequence
of inner products(i.e, dot-products).
- These standalone kernels are expected to strictly handle
the CRC storage scheme for C, A and B matrices. RRC is also
supported through operation transpose, at the framework
level.
- Added unit-tests to test all the kernels(main and fringe),
as well as the redirection between these kernels.
AMD-Internal: [CPUPL-5949]
Change-Id: I858257ac2658ed9ce4980635874baa1474b79c38
Description:
_mm512_cvtne2ps_pbh(a, b) instruction takes
b when j<16 but the code was developed in
with assuming reverse order.
Fixed some indentation issues
Changed the file name and made it uniform
Change-Id: I7b45b4c35931d8febde7b7b5d9604ea953046f97
Description:
aocl_reorder_f32obf16 function is implemented to
reorder input weight matrix of data type float to
bfloat16.
The reordering is done to match the input requirements
of API aocl_gemm_bf16bf16f32o<f32|bf16>.
The objective of the API is to convert a model/matrix
of type f32 to bf16 and process when machine supports
bf16 FMA instruction _mm512_dpbf16_ps but the model
is still in float
Change-Id: Ib7c743d52d01a1ac09e84ac120577ec9e02f90f5
-Currently lpgemm sets the context (block sizes and micro-kernels) based
on the ISA of the machine it is being executed on. However this approach
does not give the flexibility to select a different context at runtime.
In order to enable runtime selection of context, the context
initialization is modified to read the AOCL_ENABLE_INSTRUCTIONS env
variable and set the context based on the same. As part of this commit,
only f32 context selection is enabled.
-Bug fixes in scale ops in f32 micro-kernels and GEMV path selection.
-Added vectorized f32 packing kernels for NR=16(AVX2) and NR=64(AVX512).
This is only for B matrix and helps remove dependency of f32 lpgemm api
on the BLIS packing framework.
AMD Internal: [CPUPL-5959]
Change-Id: I4b459aaf33c54423952f89905ba43cf119ce20f6
Correction to CMakeLists.txt to fix problem building executables
when testing against MKL.
AMD-Internal: [CPUPL-5928]
Change-Id: Ie427fff0afb48be6ce6d940b1db2c9d1c7a40e5b
Details:
- Added a new API called unreorder that converts a matrix from
reordered format to it's original format( row-major or col-major ).
- Currently this API only supports bf16 datatype.
- Added corresponding bench and input file to test accuracy of the
API.
- The new API is only supported for 'B' matrix.
- Modified input validation checks in reorder API to account for
row Vs col storage of matrix and transposes for bf16 datatype.
Change-Id: Ifb9c53b7e6da6f607939c164eb016e82514581b7
-Added new pack kernels that packs/reorders B matrix (odd strides) from
column-major input format. This also supports the transB scenario if
input B matrix is row major.
Change-Id: Ia0fe7e5f19ae9eba5c418f4089c7e6df11091853
- Implemented the Scale post-op for the F32 API for all kernels
- f32_scale = (f32 * scale_factor) + offset
- Added the bench inputs
Change-Id: Ib0f25f870eafe695d8b2a2c434c8cb3ec4f7db4c
- Data-type of n, and conj is dim_t which will be int32_t for LP64 case.
- When loading 64-bit registers using "mov" instructions, mov(rax, var(n)),
the "n" should be 64-bit otherwise incorrect values gets loaded.
Fix: We typecast these variables to int64_t before loading into registers.
Thanks to mangala.v@amd.com for finding this bug.
Change-Id: I8542dc1ea434ca9030f3c56d9a681135055f8ba5
- Data-type of m, n, k,ldc is dim_t which will be int32_t for LP64 case.
- When loading 64-bit registers using "mov" instructions, mov(rax, var(m)),
the "m" should be 64-bit otherwise incorrect values gets loaded.
Fix: We typecast these variables to int64_t before loading into registers.
AMD-Internal: [CPUPL-5819]
Change-Id: I16043ac168a79ff9358c0c1768989a81e3c6b0e0
-Added new pack kernels that packs/reorders B matrix from column-major
input format. This also supports the transB scenario if input B matrix
is row major.
Change-Id: I4c75b6e81016331fd7e7f95ad4212e6d38dc586f
Since there is thorough kernel testing, we reduce the number of "Black Box" test cases so that CI is faster.
AMD-Internal: [CPUPL-4500]
Change-Id: Ie57eeccff8103c0051eb1904162d6447da0ef102
- Correct matsize and NumericalComparison functions for
tests with first matrix dimension <= 0.
- BLAS1:
- Fix for BLAS vs CBLAS differences in amaxv IIT_ERS tests.
- Threshold adjustments in ddotxf and zaxpy.
- Break axpyv and scalv into separate executables for
each data type.
- BLAS2:
- Threshold adjustments in symv and hemv.
- Break ger into separate executables for each data type.
- UKR:
- Break gemm and trsm ukr test into separate executables
for each data type.
- Threshold adjustments in daxpyf
- Disable {z,c}trsm ukr tests when BLIS_INT_ELEMENT_TYPE
is used, as matrix generator is not currently suitable
for this.
AMD-Internal: [CPUPL-4500]
Change-Id: I1d9e7acc11025f1478b8b511c14def5517ef0ae6
- Added the appropriate CBLAS wrappers for CROTG, CSROT,
ZROTG and ZDROT APIs. These would internally call their
?_blis_impl() layer.
AMD-Internal: [CPUPL-5813]
Change-Id: I6037f20092f99cc5a5e2794d03bbe76d6a55eb97
- Implemented the AVX512 packA kernel for col major inputs in F32 API
- Removed the work arounds for n = 1, mtag_a = PACK case, where the execution was
being directed to GEMM instead of GEMV.
Change-Id: I6fb700d96069213a762e8a83a209c5388a91050f
- Optimized DGEMM macro kernel does not
support mixed precision.
- This kernel was being used for solving
some of the mixed precision problems.
- Currently only ( bli_obj_elem(A) == 8 ) is used for checking
if the problem being solved is mixed precision.
- bli_obj_elem(A) will be equal to 8 for both double precision
data type and mixed precision case single-complex.
- Added extra checks (bli_obj_is_real( a )) to make sure that
A and B are real and DGEMM macro kernel is being used only
for DDDGEMM.
AMD-Internal: [CPUPL-5804]
Change-Id: Iaa1accf8d851d11533f8ba31dc0235fbc14f89a9
