- 博客(11)
- 收藏
- 关注
RSS订阅原创 TMA学习：在 Blackwell B200 上手写 TMA Depthwise Conv2D
是 Hopper (SM90) 引入、Blackwell (SM100) 延续的一块专用硬件，负责在之间搬大块多维 tensor。在 TMA 之前，我们只能用cp.async让每个线程各自搬一小段；TMA 则把"搬数"彻底从线程手里接管走，让搬和算真正解耦。由硬件处理边界问题。TMA descriptor 里提前声明了 tensor 的{shape, stride, 越界策略}；kernel 只需要给出"要搬的那一块的起始坐标"，越界部分由硬件自动 zero-fill。
2026-05-06 21:58:19
432
原创 深度解析：CUDA CTA Cluster 从入门到进阶的完整指南
Cluster 就是把多个 Block 强行“绑定”成一个战术编队。如果你设置，就意味着这 8 个 Block 必须同生共死。必须同时调度必须被分配在同一个 GPC（图形处理集群）内部。同一个 Cluster 内的不同 Block，必须落在不同的物理 SM 上（不能把俩 Block 塞进同一个 SM，否则就失去了跨 SM 高速通信的意义）。Thread Block Cluster 为 CUDA 编程打开了新世界的大门，但也增加了复杂度。
2026-05-05 21:55:18
513
原创 深入浅出：揭秘 nvfp4 量化机制与工程妥协的艺术
scale2先把普通逻辑矩阵：scale2.shapemk16scale2.shapem16k128×4128×4的小 tile。tile 之间按：的顺序排列。每个128 × 4tile 内部，不按连续行排列，而是按：...的顺序排列，并且每个 row 内部连续存 4 个 scale。普通 tile: T.shape = (128, 4)
2026-05-03 16:37:38
475
原创 从 AllReduce 到 NVLink Push：一步步理解 Fused Tensor-Parallel RMSNorm
本仓库的创新在于：把 Tensor-Parallel RMSNorm 中原本依赖 AllReduce 的跨 GPU 汇总，改造成基于 CUDA IPC + NVLink 的 Push 通信，并融合进单个 CUDA kernel 中完成。大家算完 → 开会同步 → 再继续算大家算完 → 直接把结果塞进所有人的邮箱 → 每个人读自己的邮箱继续算对于小数据、延迟敏感、同节点 NVLink TP 场景，这种方式有机会显著降低通信和调度开销。
2026-05-02 14:54:32
480
原创 nvidia cudnn/cutlass mxfp8 scale布局讲解
nv文档：https://editor.youkuaiyun.com/md?
2026-05-02 13:17:05
419
原创 不同warp中的线程同时访问一个bank中不同地址会有冲突吗
Question is :I’m using shared memory ,my kernel will have threads in different warps write different addresses in the same bank (but threads in the same warp write different banks),will this cause bank conflict?answer:No, bank conflicts do not occur betwee
2026-03-22 19:25:16
29
原创 cuda全局内存高效访问——对齐访问和合并内存访问
高效的全局内存访问需要保证两点：对齐访问和相邻的线程访问相邻的数据。GPU在访问全局内存的时候，通常一次性传输32 bytes或者128 bytes，如果L1和L2 cache都被使用了，那么一次性传输128 bytes，如果只有L2 cache被使用，那么一次性传输32 bytes，L1和L2 cache是否被使用和GPU的架构等都有关系，L1 cache行和设备内存中一个128字节对齐的内存区域对应。
2026-03-16 19:21:45
386
原创 行优先的多维数组中线性索引和逻辑索引的关系
这在写 CUDA 算子（如 cat 算子等）的时候特别有用，从 tid 算得本线程要负责的元素的逻辑索引，从而进行相关拷贝操作。命题：如果有一个n维行优先的数组的shape为。都取最大值，那么对于第k个维度而言，由于。通过这个命题，我们可以从线性索引求得。这和前提相矛盾，故假设不成立。的前提相矛盾，故假设不成立。用同样的方法，我们可以求得。
2026-03-16 19:21:13
326
原创 关于行优先存储的多维数组在计算机内存中的若干性质及证明
首先要搞清楚，什么叫做行优先，二维数组的情况下，行优先很好理解，先存储第一行，再存储第二行，直到存储完毕，但是当遇到更高维度的数组的时候，此时“行”是什么都不好理解了。其实所谓的行优先，本质说的是多维数组中下一个元素和当前元素的一种关系，假设数组维度为nnn,维度为d0×d1×d2×⋯×dn−1d0×d1×d2×⋯×dn−1，那么对于元素为arri0i1i2in−1其中0≤ik≤dk−。
2026-03-12 11:15:02
367
原创 反向传播的计算顺序
计算图和拓扑排序：计算图的本质是依赖关系，节点的入度代表其依赖的变量数，入度为0则代表该函数的输入变量已经被计算完毕。因此如果按照反向传播计算图的拓扑排序顺序来进行反向传播的话，计算每个backward函数的时候，其输入依赖项都已经被计算完毕（毕竟拓扑排序的定义就是：如果A指向B，那么A必然在B前面，因此开始计算某个函数F的时候，指向F的函数必然已经被计算），因此计算Fb的梯度必然在Fa的梯度计算之前被计算完毕，从而保证了计算的准确性。
2026-03-12 10:56:20
256
空空如也
空空如也
TA创建的收藏夹 TA关注的收藏夹
TA关注的人