- 博客(261)
- 资源 (2)
- 收藏
- 关注
RSS订阅原创 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
62
原创 deepGemm源码分析
直接git clone源码：可以不下载cutlass，后续再安装* 将之前下载好的cutlass直接传到deepGemm中的third-party文件夹* 执行 python setup.py install即可完成安装* 执行python tests/test_core.py 完车测试，如果报错，则再安装一次（可能是缓存的问题）
2025-02-28 09:31:12
105
原创 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
66
原创 warp specialization
WGMMA），让一些warp充当生产者（访存），另一些warp 充当消费者（计算），这种设计可以更进一步地减少访存和计算之间的耦合（可能更有利于针对性的编译器优化），减少同步开销（？，在hopper上得到了硬件级支持（async。
2025-02-13 09:13:41
64
原创 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
156
原创 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
697
原创 结合night compute分析 利用tensor core 优化K值较大的矩阵乘（超过cublas50%）
将cublas作为base line和现有的代码分析图1.1可以发现计算吞吐量明显偏低，能想到的就是计算单元处于空闲的概率较大，是访存密集型算子，因此可以增大数据的吞吐量，多给计算单元提供数据。
2025-01-14 17:12:32
334
原创 cuda中Warp Shuffle的使用
_shfl_xor_sync() 通过对调用者的通道 ID 与 laneMask 执行按位异或来计算源通道 ID：返回结果通道 ID 所持有的 var 的值。如果宽度小于warpSize，那么每组宽度连续的线程都能够访问早期线程组中的元素，但是如果它们尝试访问后面线程组中的元素，则将返回他们自己的var值。这种模式实现了一种蝶式寻址模式，例如用于树规约和广播。
2025-01-10 14:14:13
286
原创 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
103
原创 megatron接入flash-attention
如1所示，加入下面的红框，将原来导入flash_attn_forward_func注释掉（最后一个红框的第一行），然后加入flash_attn_interface内容，同时别忘记新文件夹的路径通过sys.path.append导入。同时将同级目录下的utils.py文件名修改为utils_amd.py文件。文件的路径在已经标出（此路径是flash-attn被安装到本地的路径）在本地所有目录下，删除所有包含waves的参数。注释掉的就是原来的写法，修改为现在的写法。对照源代码，修改成图中所示内容。
2024-12-12 11:22:54
72
原创 megatron源码分析
和inter-layer model parallel approach的区别。猛一看，这两个的中文翻译都是“其实它们是有区别的：即，。例如上图右上方的6层网络，前三层给一个GPU，后三层给另外一个GPU。而另外一个是。例如上图右下方的6层网络，横向切一刀，即一个，会被分配到不同的GPU上面。
2024-12-09 17:33:35
57
原创 [3W字]全面解析tensor core实现gemm
对于固定尺寸的输入矩阵来说，通常Block分块尺寸越大，意味着单个Block内的计算量越大，但是需要并行计算的Block数量越少，这是Block维度计算量和并行度的权衡。一般来说针对不同尺寸的输入矩阵，需要采用不同的分块策略，才能更好地实现Block维度计算量和并行度的平衡，输入矩阵的尺寸越小，Block分块尺寸越小，输入矩阵的尺寸越大，Block分块尺寸越大。另一方面结合硬件规格的限制，Block分块尺寸一般为32、64、128、256之间的组合。
2024-12-05 09:51:44
185
原创 cuda实现gemm
对于 C 矩阵的每一个元素，都要读取 A 矩阵的一行和 B 矩阵的一列来计算，那么计算完整的 C 矩阵，A B 矩阵都要重复读取多次，所以直接按定义计算效率很低。首先说明，很多文章在解释这种方案性能差的时候，都是以太高作为主要理由，实际上在并行计算中，是常用的设计方法，延迟只要能被其他过程覆盖就没有问题。对于 FP32 数据，如上图所示，一个 warp 一次做 32 次 FFMA，对应 64OP，需读取 A 矩阵 1 个元素和 B 矩阵 32 个元素，共 132byte。
2024-10-11 17:28:47
309
原创 swizzle
为了减少指令数，我们在进行kernel优化时会采用向量化的读写指令（也叫大字长读写），如以128bit的形式读写共享内存，此时线程需要访问的单位数据量为16byte，32个线程需要访问的数据量为16byte x 32 = 512byte。如图2，当32个线程同时访问32个不同的bank时，各个bank是并行执行的，其效率是最高的，即32个线程并发的访问32个bank中不同颜色的单元，是可以并行的，值得注意的是其中的线程编号（如图2中的T0所示）和bank中的行位置并没有连续性要求。
2024-09-30 15:44:10
158
原创 cuda基础知识
如果block所含线程数目不是warp大小的整数倍，那么多出的那些thread所在的warp中，会剩余一些inactive的thread，也就是说，即使凑不够warp整数倍的thread，硬件也会为warp凑足，只不过那些thread是inactive状态，需要注意的是，即使这部分thread是inactive的，也会消耗SM资源。，因为资源限制，SM要为每个线程块分配共享内存，而也要为每个线程束中的线程分配独立的寄存器，所以SM的配置会影响其所支持的线程块和warp并发数量。
2024-09-29 14:58:03
167
原创 nsight-compute使用教程
有的时候在linux上安装上了nsight-compute，可以生成报告，但是却因为缺少qt组件而无法打开，我选择的方法是在linux上生成报告，在window上的nsight compute的图形界面打开，需要注意的是，nsight compute图形界面的版本一定要更高，不然无法打开。
2024-09-26 11:28:48
1318
原创 cuda中使用二维矩阵
经过上面的操作后，就可以像操作二维数组了t=O83At=O83ACUDA学习之一：二维矩阵加法 - 冷豆东 - 博客园 (cnblogs.com)https://www.cnblogs.com/jugg1024/p/4349243.htmlCUDA 中的 cudaMalloc使用二重指针（void**)的一些理解_cuda申请二重指针指针-优快云博客https://blog.youkuaiyun.com/lingyunxianhe/article/details/92001270。
2024-09-19 14:34:40
327
原创 从index_put出发全面学习cuda和pytorch技术
深感目前对于cuda和pytorch所涉及知识的广度和深度，但一时又不知道该如何去学习，经过多日的考虑，还是决定管中窥豹，从一个算子出发，抽丝剥茧，慢慢学习，把学习中碰到的问题都记录下来，希望可以坚持下去。
2024-09-13 23:37:57
365
原创 c++指针数组和数组指针
数组指针：本质是一个指针，指向了一个数组，数组中的每个元素都是某种数据类型的值（比如 int 类型）。//定义了一个数组指针，指向一个大小为n的数组,数组中的每个元素都是int类型数组指针也称行指针，也就是说，当指针p执行p+1时，指针会指向数组的下一行，如：//p是一个数组指针，指向了一个包含4个int型元素的数组p=a;//将二维数组的首地址赋给p，即a[0]或a[0][0]p++;//跨过第一行，p指向了a[1][0]
2024-09-12 11:37:32
133
原创 c++和c函数总结
它是为了方便系统之间的移植而定义的，不同的系统上，定义size_t 可能不一样。size_t 的目的是提供一种可移植的方法来声明与系统中可寻址的内存区域一致的长度。size_t 的声明是实现相关的。一般来说，size_t 可能的最大值是SIZE_MAX。例如，size_t 用做sizeof 操作符的返回值类型，同时也是很多函数的参数类型，包括malloc 和strlen。size_t类型）， 用来表示可以被执行读写操作的数据块的大小。因为size_t 是无符号的，一定要给这种类型的变量赋正数。
2024-09-11 21:57:31
131
原创 c++指针和引用专题
图解C++指针与引用的区别_指针与引用 图解-优快云博客其实从本质上看，引用就是加了约束的指针，引用改善了指针的不足之处操作空指针：被赋值为0的指针操作野指针：未被初始化的指针不知不觉改变了指针的值，而后还以为该指针正常。引用必须初始化（保证没有野指针）初始化就是现存变量的别名(保证不是空指针)一个引用永远指向他初始化的那个对象（保证指针值不变）
2024-09-10 22:13:21
148
原创 pytroch算子接入
【pytorch扩展】CUDA自定义pytorch算子（简单demo入手）_pytorch自定义算子-优快云博客使用 Cpp 扩展自定义进程组后端 — PyTorch 教程 2.4.0+cu124 文档 - PyTorch 中文。
2024-09-05 13:46:09
98
原创 centos换源安装升级gcc
使用devtools升级安装的时候，由于此库已经停止更新 了，因此需要切换阿里源SCL+Devtoolset 安装与使用笔记-腾讯云开发者社区-腾讯云 (tencent.com)
2024-08-30 19:41:09
1244
原创 triton之flaggems的point-wise分析
如果没有写triton.jit装饰器，则会报上面的错误，fn.args_names是JITFuntion的属性。
2024-08-20 17:15:02
136
原创 cuda指北之professional CUDA C Programming第五章共享内存
CPU的内存是线性存储的，但是GPU的共享内存并不是线性的，而是二维的说白了就是不要一个线程束中访问一列共享内存，而是要访问一行最好就是同步线程，不同bank（同一个线程束中）
2024-08-13 17:24:04
88
原创 cuda指北之professional CUDA C Programming第四章内存模型
GPU上最大的内存空间，延迟最高，使用最常见的内存，global指的是作用域和生命周期，一般在主机端代码里定义，也可以在设备端定义，不过需要加修饰符，只要不销毁，是和应用程序同生命周期的。使用共享内存的时候一定要注意，不要因为过度使用共享内存，而导致SM上活跃的线程束减少，也就是说，一个线程块使用的共享内存过多，导致更过的线程块没办法被SM启动，这样影响活跃的线程束数量。每个SM都有一定数量的由线程块分配的共享内存，共享内存是片上内存，跟主存相比，速度要快很多，也即是延迟低，带宽高。
2024-08-06 14:57:16
90
空空如也
TA创建的收藏夹 TA关注的收藏夹
TA关注的人