- 博客(287)
- 资源 (2)
- 收藏
- 关注
RSS订阅原创 算法之链表
【1】dummy node ：开始节点的前驱节点，这个节点是new出来的【2】需要理解指针的含义，即使链表的指针本质上也是相同的【3】翻转部分就是翻转列表的基础代码。
2025-08-31 21:45:48
272
原创 算法之x数之和
【2】数组的值是key，数组的下标是value，存到hashmap中【3】找到key则return结果，否则存储当前遍历的节点{key，value}
2025-08-30 17:27:41
283
原创 lambda函数
C++ 中的 lambda 函数可以理解为一种 "临时定义的小型函数"，就像一个 "一次性的函数"，不需要单独写函数名和定义，直接在需要的地方写出来就能用。
2025-08-07 10:50:51
71
原创 基于cute的gemm教程
每个 block 要计算 128x256x32 ，如果只采用一个 warp 计算 16x8x16, 那么这个 block （针对数据矩阵D）要循环 (128 / 16) * (256 / 8) 次计算。为了提高效率，这里我们希望每个 block 采用 256 thread，一次计算 32x32x16，按照如下分布：当我们规划好了 thread block 要这样运算之后，如果没有 cute 我们需要自己提前构建 mma 的 matrix 数据，用循环来计算整个过程。
2025-08-06 11:20:45
145
原创 deepgemm问答
gemm的运算是每个block_m交给一个cta去计算，且每个cta对k-tile进行循环，因此lhs_scales每次传输的数据量为block_m，如果不是col_major则会造成跨步的访存，影响效率。
2025-08-04 11:16:45
97
原创 cpp数据结构
map： map内部实现了一个红黑树（红黑树是非严格平衡二叉搜索树，而AVL是严格平衡二叉搜索树），红黑树具有自动排序的功能，因此map内部的所有元素都是有序的，红黑树的每一个节点都代表着map的一个元素。因此，对于map进行的查找，删除，添加等一系列的操作都相当于是对红黑树进行的操作。map中的元素是按照二叉搜索树（又名二叉查找树、二叉排序树，特点就是左子树上所有节点的键值都小于根节点的键值，右子树所有节点的键值都大于根节点的键值）存储的，使用中序遍历可将键值按照从小到大遍历出来。
2025-07-23 22:12:36
177
原创 cuda优化之softmax
对于内存访问来说，我们是一次性加载整个向量，然后一次性保存它。Bytes=2*N*4(每个浮点值占四字节)先求max(x)//N次浮点操作，再求x-m//N次浮点操作，exp=e^x//N次浮点操作，s=sum(exp)//N次浮点操作，out=exp/s//N次浮点操作。FLOPS=5*N这样得到的信息就是这个算子每加载8字节，我们进行5次浮点运算。TheoreticalMaximum = (5/8)*理论带宽 1TB/s = 625GFLOPs//理论最高算力。
2025-07-14 19:55:32
1206
原创 pybind学习
假设我们有一个简单的 CUDA 内核函数 example_kernel，它对输入张量进行某种操作。以下是如何使用 AT_DISPATCH_FLOATING_TYPES 来实现这个目标。example_kernel<scalar_t><<<blocks, threads>>>(input.data<scalar_t>(), size)：根据调度的数据类型调用相应的 CUDA 内核。"example_forward_cuda"：操作的名称，用于错误信息。input.type()：获取输入张量的数据类型。
2025-07-14 10:28:06
366
原创 cute教程
摘要：本文介绍了一种有层次的Tensor布局表示方法。传统单层布局（行/列优先）无法描述复杂场景，文中提出将Tensor分为内外两层表示。以4行8列矩阵为例，内层为2列子矩阵，外层为4个这样的子矩阵。通过分层shape和stride表示（如shape:(4,(2,4)), stride:(2,(1,8))），能够更精确地描述复杂内存布局。该方法解决了非单调轴情况下的存储表示问题。（149字）
2025-07-03 14:36:00
156
原创 nvidia中的compute-sanitizer使用
摘要：本文介绍如何正确使用NVIDIA Compute Sanitizer调试CUDA程序。通过实例说明，在调试与pybind结合的CUDA代码时，应分步操作：先完成项目构建（build），再单独运行Compute Sanitizer进行内存检查。错误做法是将构建和调试命令合并执行，正确方式应为"compute-sanitizer --tool memcheck python group_gemm/ops_test.py"，确保调试过程准确有效。（98字）
2025-06-24 10:35:23
453
原创 bank conflict
CUTLASS 中的 Swizzle 技术解析 本文探讨了 CUDA 模板库 CUTLASS 中的 Swizzle 技术及其应用。Swizzle 是一种内存访问优化技术，主要用于解决 GPU 内存访问中的 bank conflict 问题。通过数据重排（swizzling）操作，可以有效地重组内存访问模式，优化共享内存的访问效率。文章分析了 Swizzle 的基本原理和实现方法，并提供了在 CUTLASS 框架中的具体应用示例。该技术对于提升矩阵运算等计算密集型任务的性能具有重要意义。
2025-06-23 17:04:20
209
原创 cutlass学习教程
文章摘要：该文档描述了一个内存管理接口，包含DeviceAllocation类，主要提供reset()和get()两个方法。reset()用于删除托管对象并将容量归零，通过智能指针实现；get()则返回托管对象的指针。这些功能为内存管理提供了基础操作支持。（150字）
2025-05-29 19:46:10
329
原创 cpp中vector的使用
C++中vector的两种访问方式v[]和v.at()都能访问元素且都不能越界，但存在重要区别：v[]不进行边界检查，越界时返回错误引用可能导致程序崩溃；v.at()会进行边界检查，越界时抛出异常，可通过try-catch捕获使程序继续运行。实际使用中v[]更常见但风险较高，v.at()更安全但性能稍低。
2025-05-29 09:54:19
147
原创 nvrtc环境依赖
本文介绍了如何通过添加NVIDIA源并下载相关文件来安装NVIDIA CUDA运行时和Python库。首先，使用pip install nvidia-pyindex添加NVIDIA源。接着，通过pip download nvidia-cuda-runtime nvidia-cuda-python下载所需文件，下载后文件夹中会新增多个文件，但实际需要安装的只有其中三个文件（红框标注）。最后，对这三个文件执行pip install完成安装。
2025-05-14 14:26:59
280
原创 gemm中的swizzle
针对一个8x8-half的寄存器表示的作为输出的矩阵块，ldmatrix其输入要求为8个shared memory地址，每个地址指向一个16byte共享内存中的数据，其中T0-Addr0指向的16byte数据经过ldmatrix会被分派到T0-T3的V0寄存器中。从数学逻辑上看，8x8-half的寄存器数据表示连续的矩阵块，共享8x16byte的内存也有很好空间局部性的矩阵块，但是从共享内存的存储逻辑上看，为了避免读取时的bank冲突，其必须分配在不同的bank中。
2025-04-23 11:17:18
246
原创 group-gemm教程
fanshiqing/grouped_gemm: PyTorch bindings for CUTLASS grouped GEMM.（1） git clone https://github.com/fanshiqing/grouped_gemm.git（2） git submodule update --init --recursive（3）python setup.py install
2025-04-17 10:18:23
840
原创 wgmma指令解析
上图中的T0{d0，d1}表示的是thread0中的寄存器0和寄存器1中是A矩阵ROW0 * B矩阵COL0和A矩阵的ROW0和B矩阵的COL1的结果，T1{d0，d1}表示的是thread1中的寄存器0和寄存器1中是A矩阵ROW0 * B矩阵COL2和A矩阵的ROW0和B矩阵的COL3的结果,T4{d0，d1}表示的是thread4中的寄存器0和寄存器1中是A矩阵ROW1 * B矩阵COL0和A矩阵的ROW1和B矩阵的COL1的结果。
2025-03-11 18:48:25
445
原创 deepGemm源码分析
直接git clone源码：可以不下载cutlass，后续再安装* 将之前下载好的cutlass直接传到deepGemm中的third-party文件夹* 执行 python setup.py install即可完成安装* 执行python tests/test_core.py 完车测试，如果报错，则再安装一次（可能是缓存的问题）
2025-02-28 09:31:12
549
原创 cuda编程模型
CTA：(Collaborative Tread Arrays), CUDA程序的任务分发单位，CTA与block是同一事物在执行模型和编程模型中的表述；同一个block中的线程使用同一块shared memory；一个CTA里的线程必须被分配到同一个SM中；目前硬件下，CTA最多由16个warp组成（512个线程）；一个kernel会启动一个grid，一个grid包含多个block，每个block包含多个thread。而一个cluster可以有多个block。
2025-02-22 20:05:55
801
原创 warp specialization
WGMMA），让一些warp充当生产者（访存），另一些warp 充当消费者（计算），这种设计可以更进一步地减少访存和计算之间的耦合（可能更有利于针对性的编译器优化），减少同步开销（？，在hopper上得到了硬件级支持（async。
2025-02-13 09:13:41
348
原创 cuda学习资料汇总
https://github.com/NVIDIA/cutlass/blob/main/examples/cute/tutorial/wgmma_sm90.cu https://github.com/BBuf/how-to-optim-algorithm-in-cuda/blob/master/cutlass/CUTLASS%20Tutorial%3A%20Mastering%20the%20NVIDIA%C2%AE%20Tensor%20Memory%20Accelerator%20(TMA).md
2025-02-08 16:40:43
260
原创 CUDA学习-内存访问
简单理解一下，当上面两种情况发生时，硬件就可以判断（具体是硬件还是编译器的功劳，我也不确定，先归给硬件吧），单个 half warp 内，最多需要 64 bytes 的数据，那么两个 half warp 就可以合并起来，通过一次 memory transaction，拿回 128 bytes 的数据。这时候就符合前面说的合并条件 2，所以线程 0 - 7，以及线程 8 - 15 的访存请求，合并为一次 memory transaction。线程 16 -31 同理。，每个包含 8 个 thread。
2025-01-28 09:02:19
960
原创 结合night compute分析 利用tensor core 优化K值较大的矩阵乘（超过cublas50%）
将cublas作为base line和现有的代码分析图1.1可以发现计算吞吐量明显偏低，能想到的就是计算单元处于空闲的概率较大，是访存密集型算子，因此可以增大数据的吞吐量，多给计算单元提供数据。
2025-01-14 17:12:32
537
原创 cuda中Warp Shuffle的使用
_shfl_xor_sync() 通过对调用者的通道 ID 与 laneMask 执行按位异或来计算源通道 ID：返回结果通道 ID 所持有的 var 的值。如果宽度小于warpSize，那么每组宽度连续的线程都能够访问早期线程组中的元素，但是如果它们尝试访问后面线程组中的元素，则将返回他们自己的var值。这种模式实现了一种蝶式寻址模式，例如用于树规约和广播。
2025-01-10 14:14:13
613
原创 tensor core中的ldmatrix.sync.aligned详解
如果.num =4则一共会读4*8*8 half从sm到register中，且读写是按照一个warp的形式进行组织，也就是一个warp需要读的大小应该是16 * 32B的数据量，且warp中的lane的组织形式为16 * 2（对应图2.1），然后不过每个lane中的resiger中的数据，则是会按照图2.3的形式进行排布。
2025-01-09 14:30:24
277
空空如也
TA创建的收藏夹 TA关注的收藏夹
TA关注的人