每天应该做些啥
要说暑假已经开始了,按照惯例给自己写一个不会遵守的每日代办计划:刷题/八股部分 - 每天一道leetcode hot 100/neetcode hot 150 - 每天进行ML项目一个小部分的from scratch编写联系,夯实基础 - 理解一至两道八股题目 - 复习/学习经典架构的理论知识学习部分 - 如果7.11去成mhy的话,我将进行GAMES101的学习,这个暑假争取学完 - 每天至少一篇Embodied相关领域paper的精读 - 可以看看技术博客/daily paper之类的了解一下社区的情况Reaserch 部分 - 这点自然不必说,一切都低于research相关的内容有想起什么再来补充,以上这几条,每天完成80%就算成功!
各种各样的JEPA
借此机会简单整理一下JEPA的思想以及他最近的发展。一言以蔽之,JEPA的思想是不去预测像素空间中的变化,而是转向embedding空间。高维的像素空间有太多噪声和无用的信息,JEPA希望能够在embedding空间中学习到世界的规律/只是,而不只是预测视频而已。
规约reduce算法
内存类型 速度 容量 可见范围 全局内存 ~1 TB/s (HBM) 大 (GB级) 所有线程 共享内存 ~200 TB/s (SRAM) 小 同一个 Block 寄存器 最快 (0延迟) 极少 单个线程 规约算法的核心思路:把数据从全局内存搬到共享内存和寄存器,在最快的层级上完成计算并写回 SIMT 和 Warp DivergenceGPU使用SIMT,对应的一个warp的3个线程在同一时刻执行同一条指令如果写了分支代码,warp内有些线程走if,有些走else/此时硬件只能先执行if,再执行else分支。实际上代码背串行化了。因此再规约算法中要小心处理 if (threadIdx.x < stride) 这类条件,一旦stride小于32,分支会在同一warp内产生divergence 共享内存和访存冲突共享内存被划分为若干Bank,通常为32个。每个bank在同一时钟周期内只能服务于一个线程。如果同一个Warp中的多个线程访问同一个Bank中的不同地址,就会发生Bank conflict,导致访问串行化 ...
oneflow element-wise代码解析
12345678910111213141516171819202122232425262728293031323334353637383940414243444546474849505152535455565758596061626364656667686970717273747576777879808182838485868788899091929394959697989910010110210310410510610710810911011111211311411511611711811912012112212312412512612712812913013113213313413513613713813914014114214314414514614714814915015115215315415515615715815916016116216316416516616716816917017117217317417517617717817918018118218318418518618718818919019119219319419519619719819920020120...
2026-6-18
最近这一两周不知道怎么了,每天都心烦意乱的。浑浑噩噩的度过每一天。感觉什么事都没有做一天就被晃过去了。。。一事无成这一块
CUDA扩展指南
Intro如何将写好的kernel优雅地嵌入Python流程中。Pytorch提供了 torch.utils.cpp_extension 工具箱,让我们能够以极小代价将C++/CUDA代码编译为Python模块。这里会梳理三种主流方式,并解释为什么要自定义算子,如何避免环境坑以及如何让自定义kernel无缝支持 autograd 核心概念: 概念 作用 典型位置 torch.utils.cpp_extension.load_inline 直接编译字符串中的C++/CUDA代码 Jupyter/快速原型 torch.utils.cpp_extension.load 编译指定 .cpp/.cu 文件 脚本中临时编译 CUDAExtension / CppExtension 在 setup.py 中声明扩展模块 正式打包分发 BuildExtension 替换 setuptools 默认构建命令,注入PyTorch编译参数 setup.py 的 cmdclass PYBIND11_MODULE 将C++函数暴露...
世界很小,生活很大(也许会写成交换的总结吧)
在一年的时间里,先是美国呆了两个月,紧接着瑞士待了5个月。一路下来知识没学习多少,却也真正开阔了眼界。尽管不爱出门,但也去了6,7个周边的欧洲国家,城市。看到了另一种生活,另一种文化。最大的感悟可能就是世界在心中变小了,其实从大学离开合肥去往香港开始,就有已经略有体会。以前书中看到的,存在于屏幕中的各个城市,出现在现实的眼前。都说香港不算留学生,对于高中只会往返于两点一线的我可能不能苟同。但现在两地往返之方便,让我感觉这个特区似乎和一个平常的城市并无太多区别。 离开香港到了瑞士,这个第一次从小学的某个杂志上看到的世界上最富有的国家。说实话,我对瑞士的了解几乎都是抵达以后才建立的,在此之前只知道这是一个很贵的地方。几个月间,亦有机会在欧洲漫游,课本中的场景,游戏中的原型背景,世界知名的城市,物品第一次亲眼得见。似乎远在天边的东西,都变得触手可及。世界很小,这些曾经”遥望“的东西,都能够看见,感受。 曾经觉得出国会是很高大上,遥不可及的事情。现在觉得也并非难事。 在欧洲萌生出要把这里游遍的想法,乃至于将世界游遍的想法。(草稿中)
逐元素操作算子
Element-wise 算子优化一、核心特性1. 数据并行对于 element-wise 算子,输出位置的值只取决于对应位置的输入,不涉及其他位置。因此线程之间不需要交换数据 / 同步指令,也没有额外的寻址开销,具备优秀的数据并行性。这是 GPU 优化中最容易打到理论峰值带宽的一类算子。 为什么向量加法天然适合 GPU:每个元素的「读 → 算 → 写」彼此独立,互不依赖,所以成千上万份操作可以同时重叠进行,把显存带宽和运算单元一起喂饱。 2. 极低的计算访存比(最重要的性能特征)这是 element-wise 算子最关键的性能特征,决定了所有优化策略。 GPU 上执行一条算术指令只需要几个时钟周期,而访问全局显存的耗时却高达几百个时钟周期。两者差距悬殊。 以加法 $C[i] = A[i] + B[i]$ 为例: 计算量:只有 1 次加法 访存量:读 A、读 B、写 C,共 3 次显存操作 = 3 × 4 Byte = 12 Byte 计算访存比 = 1 FLOP / 12 Byte ≈ 0.083 FLOP...
CUDA全局坐标计算
1int i = blockIdx.x * blockDim.x + thredIdx.x CUDA 为每一个线程提供了四个内置变量,用于定位自己在整个任务中的位置: 量 含义 维度范围 threadIdx.x/y/z 线程在块内的局部索引 日 blockDim-1 blockIdx.x/y/z 线程块在网格内的索引 日 gridDim-1 blockDim.x/y/z 每个块每维的线程数 由启动参数 决定 $<<\ldots,\quad{\mathrm{t h r e a d s}}>>>$ gridDim.x/y/z 网格每维的块数 由启动参数 << 所谓全局索引,就是跳过前面所有块的线程,再加上当前块内的偏移。 全局坐标计算方式在内存中,数据始终按照一维排布,因此全局坐标计算就是将多维逻辑索引映射为一维物理地址 一维坐标计算1int global_id = blockIdx.x * blockDim.x + threadI...