|
Neko 1.99.9
A portable framework for high-order spectral element flow simulations
|
#include <cdtp_kernel.h>
Static Public Member Functions | |
| static __device__ void | run (T *, const T *, const T *, const T *, const T *, const T *, const T *, const T *, const T *, const int) |
Device kernel for D^T x on the AMD matrix cores
The only one of these whose contractions come after its pointwise work rather than before it, and the only one that accumulates: the three weighted fields are formed first and then contracted into a single output. That is axhelm's second half rather than its first.
Because the operator is handed dxt, dyt and dzt already transposed, the contraction index pattern is the ordinary forward one, so mfma_contract_sel is used with TRANSPOSE = false against the staged transposes rather than with TRANSPOSE = true against the forward matrices. The output is zeroed with the staging and every axis accumulates into it, matching the divergence half of ax_helm_mfma_elem().
Definition at line 325 of file cdtp_kernel.h.