#include <cdtp_kernel.h>
Static Public Member Functions | |
| static __device__ void | run (T *__restrict__, const T *__restrict__, const T *__restrict__, const T *__restrict__, const T *__restrict__, const T *__restrict__, const T *__restrict__, const T *__restrict__, const T *__restrict__, const int) |
Detailed Description
Device kernel for D^T x on the fp64 tensor cores
The only operator here whose contractions come after its pointwise work rather than before it, and the only one that accumulates: the three weighted fields tar, tas and tat are formed first and then contracted with the transposed derivative matrices into one output. That is axhelm's phase 2 rather than its phase 1, and it is what makes cdtp a different test of the strategy from dudxyz, opgrad and conv1 rather than a fourth copy of one.
Because the operator is handed dxt, dyt and dzt already transposed, the contraction index pattern is the ordinary forward one – dxt[i + l*LX] against tar[l,j,k] – so dmma_contract is used with TRANSPOSE = false and the staged transposes, not with TRANSPOSE = true and the forward matrices. The accumulation is ACCUM: axis 0 initialises every tile of the output and axes 1 and 2 add into it.
Four cubes, 17920 B, the same as every other DMMA variant here. w3 is a quadrature weight shared by every element, so it is indexed by the point within the element (dmma_idx::l) rather than the global offset, as in opgrad_dmma_elem().
Expectation, recorded before measuring: cdtp carries 6827 bytes per contraction at lx = 8, below dudxyz's 8192 and conv1's 20480, both of which measured as ties against their incumbents on GH200 because they were already at the bandwidth roof. This one should tie as well. It is offered because the shape is untested, not because the ranking predicts a win.
Definition at line 359 of file cdtp_kernel.h.
Member Function Documentation
◆ run()
The documentation for this struct was generated from the following file:
- /home/runner/work/neko/neko/src/math/bcknd/device/cuda/cdtp_kernel.h
