文章预览
本文由 @Simon V(https://github.com/simveit) 授权转载和翻译并发表到本公众号。原始地址为:https://veitner.bearblog.dev/load-and-store-matrices-efficently-with-ptx-instructions/ 本文实验cuda代码见:https://github.com/simveit/load_and_store 使用PTX指令更高效地加载和存储矩阵 2025年5月14日 ldmatrix 从PTX文档(https://docs.nvidia.com/cuda/parallel-thread-execution/#warp-level-matrix-instructions-ldmatrix)中我们可以看到, ldmatrix 可以用于从共享内存中集体加载一个或多个矩阵,以供 mma 指令使用。 指令格式如下 ldmatrix.sync.aligned.shape.num{.trans}{.ss}. type r, [p]; ldmatrix.sync.aligned.m8n16.num{.ss}.dst_fmt.src_fmt r, [p]; ldmatrix.sync.aligned.m16n16.num.trans{.ss}.dst_fmt.src_fmt r, [p]; .shape = {.m8n8, .m16n16}; .num = {.x1, .x2, .x4}; .ss = {.shared{::cta}}; . type = {.b16, .b8}; .dst_fmt = { .b8x16 }; .src_fmt = { .b6x16_p32, .b4x16_p64 }; 该指令将从 .shar
………………………………