jeffdaily commented on PR #1399: URL: https://github.com/apache/mahout/pull/1399#issuecomment-4682482362
@ryankert01 This PR was the easy end: hand-written kernels plus the CUDA runtime/driver API, which hipify mostly mechanically -- the one real subtlety was wave size (the `>>5` -> `/warpSize` fix, since CDNA warps are 64-wide). #1390's implicit-Hadamard Ozaki engine is the hard end. It uses `nvcuda::wmma` and raw inline PTX (`mma.sync.aligned.m16n8k32...s8.s8.s32`) -- int8 tensor-core MMA -- to do an Ozaki accurate-FWT. None of that hipifies: inline PTX doesn't compile under hipcc, and it has to be rewritten against AMD matrix cores (MFMA on CDNA / rocWMMA), whose fragment shapes differ (gfx90a MFMA isn't `m16n8k32`), so the Ozaki tiling and the pre/post layout shuffles have to be re-derived for the AMD fragment geometry. The Ozaki scheme as an algorithm is portable; its whole payoff is fast low-precision matrix-core matmul, so it's a real reimplementation of that path, not a translation. So: runtime/memory/plain-kernel code ports about as easily as this PR did; anything built on tensor cores, `mma`/PTX, or CUTLASS/CuTe needs a from-scratch AMD matrix-core path. Convertible, but not "this easy." -- This is an automated message from the Apache Git Service. To respond to the message, please log on to GitHub and use the URL above to go to the specific comment. To unsubscribe, e-mail: [email protected] For queries about this service, please contact Infrastructure at: [email protected]
