DeepSeek绕过了英伟达引以为傲的CUDA，使用更底层的编程语言做优化 DeepSeek's AI breakthrough bypasses industry-standard CUDA for some functions, uses Nvidia's assembly-like PTX program
deepseek不止怼了openAI而且怼了CUDA呀 CUDA的风评不太好但是大家都花钱买简单，还打算吹个大泡泡结果被不怕麻烦就想省钱的中国人…… 护城河没有了，真像是大象进了瓷器店…… Eric schmidt said“ I like to think of CUDA as the C programming language for GPUs. That''''s the way I like to think of it. It was founded in 2008. I always thought it was a terrible language and yet it''''s become dominant. “ 曾经觉得英伟达的CUDA是很蠢的编程语言，但现在CUDA是英伟达最牛逼的护城河，所有的大模型都要在CUDA上运行，而只有英伟达的GPU支持CUDA，这是其他芯片撼动不了的组合。 https://github.com/ociubotaru/transcripts/blob/main/Stanford_ECON295%E2%A7%B8CS323_I_2024_I_The_Age_of_AI%2C_Eric_Schmidt.txt https://www.youtube.com/embed/ltfiLJ76Ofo
olivia0216 发表于 2025-01-30 08:05 DeepSeek绕过了英伟达引以为傲的CUDA，使用更底层的编程语言做优化 DeepSeek's AI breakthrough bypasses industry-standard CUDA for some functions, uses Nvidia's assembly-like PTX program
DeepSeek's AI breakthrough bypasses industry-standard CUDA for some functions, uses Nvidia's assembly-like PTX program
V3的硬件效率之所以能比Meta等高出10倍，可以总结为“他们从头开始重建了一切”。 在使用英伟达的H800 GPU训练DeepSeek-V3时，他们针对自己的需求把132个流式多处理器（SMs）中的20个修改成负责服务器间的通信，而不是计算任务，变相绕过了硬件对通信速度的限制。 这种操作是用英伟达的PTX（Parallel Thread Execution）语言实现的，而不是CUDA。 PTX在接近汇编语言的层级运行，允许进行细粒度的优化，如寄存器分配和Thread/Warp级别的调整。 这种编程非常复杂且难以维护，所以行业通用的做法是使用CUDA这样的高级编程语言。 换句话说，他们把优化做到了极致。
行家说说
Deepseek确实绕过了Cuda，我几天前就说了呀，没什么很惊讶的。所以这次华为和AMD显卡可以很容易接入
对于显卡来说，PTX有点像汇编，AMD也有类似的。
其实就是并行计算影响速度的部分并没有在cuda 上执行
他们一开始就是考虑会被制裁的，要做都是全套的。他们自己也适配了华为的显卡
另外cuda内嵌PTX也是正常操作 profile出来的hotspot手改写成汇编也是常见的优化
如果全PTX手写，工作量就大了
即便deepseek全部使用底层硬件语言的话，也不代表其它公司有这个本事这么做，
全部用底层语言手搓大项目这需要牛逼工程师，这种工程师不是给工资就可以随便找到的。
比C C++更靠近机器
CUDA最后就是被NVCC转换成PTX，然后GPU来执行
楼主难道不知道，PTX本来就是Nvidia 开发出来的啊？
以下是引用
“首先要明确的是，PTX仍然是英伟达GPU架构中的技术，它是CUDA编程模型中的中间表示，用于连接CUDA高级语言代码和GPU底层硬件指令。
CUDA起到了提供高级编程接口和工具链的作用，可以简化开发者的工作。而PTX作为中间层，充当高级语言和底层硬件之间的桥梁。
所以说，DeepSeek做了PTX级别的优化不意味着完全脱离了CUDA生态，但确实代表他们有优化其他GPU的能力。我们不知道DeepSeek内部是否使用AI辅助编写了PTX代码——但是确实刚刚见证DeepSeek-R1编写的代码显著提升大模型推理框架的运行速度。”
看来底层工程师又开始吃香了， 为了追求速度， AI从PYTHON->C++->CUDA->PTX了
你应该没用PTX，PTX本来就是Nvidia弄出来的东西
汇编都是差不多的
看了一下PTX是个virtual machine,可能比一般汇编更加容易debug
是的，你的逻辑是对的。如果大家都用从底层硬件指令手搓项目，那用谁家的硬件是没啥差别的，隔夜就搬家换供应商了。但我怀疑有多少家想、或者会这么做，这样对人的要求高到不可控的地步了，是违反潮流的。
华为也有gpu
理论上ds可以用来提高华为的GPU性能，如果已经做到这么底层的话
也就是说如果进一步连低端NVIDIA GPU也禁了的话，中国突破也只是时间问题
并且狠狠打脸scale ai 小黄和奥特曼，只要功夫深，完全可以不靠堆GPU，就跟有个网友说的，挖金子不需要金铲子，铁铲子也行
以后是不是AI直接binary编程了，没人什么事了， 人类需要AI翻译它到底干了什么，彻底失控
也有其他锅。
CUDA是女大的moat,如果deepseek几十个人就能绕过的话，微软，AMD等等是很高兴的。
现在用AI，至少高级计算机语言的翻译比之前简单了不少
AMD的汇编中间语言几乎是照抄辉达的PTX, 所以DeepSeek只要稍加修改就可移植到AMD的GPU上去
😄，你就没想过，deep seek团队，先教会Deepseek用PTX，然后引导Deepseek直接用汇编编程， 他们只需要监督就行。 通过并行部分的程序，再让deep seek去蒸馏这个过程，学会如何编PTX重写。numpy, scipy
我觉得，我都想到了，他们估计都实践了好久。 问题只是，DeepSeek能不能直接化生编译器，借给你binary去验证
能用汇编语言编程的都是牛人。
你就是deepseek的吧 怎么违反规定把这个技巧公开出来 论文里面故意没写的
哈哈哈，果然公开的都是没用的
黑科技都是自己的
小扎强力要求这部分也得开源，不然不能replicable
天哪。当年我上学的时候就是用汇编语言教的。那时候还觉得有java，c还要学这个机器语言干嘛
deepseek不止怼了openAI而且怼了CUDA呀 CUDA的风评不太好但是大家都花钱买简单，还打算吹个大泡泡结果被不怕麻烦就想省钱的中国人…… 护城河没有了，真像是大象进了瓷器店……
Eric schmidt said“ I like to think of CUDA as the C programming language for GPUs. That''''s the way I like to think of it. It was founded in 2008. I always thought it was a terrible language and yet it''''s become dominant. “ 曾经觉得英伟达的CUDA是很蠢的编程语言，但现在CUDA是英伟达最牛逼的护城河，所有的大模型都要在CUDA上运行，而只有英伟达的GPU支持CUDA，这是其他芯片撼动不了的组合。
https://github.com/ociubotaru/transcripts/blob/main/Stanford_ECON295%E2%A7%B8CS323_I_2024_I_The_Age_of_AI%2C_Eric_Schmidt.txt https://www.youtube.com/embed/ltfiLJ76Ofo
国内科班出生的电子系的都要熟悉汇编
还真不太一样。如果只用cuda，基本上NVDA芯片相对别的厂牌有绝对的优势。 一旦绕过cuda，能加入针对不同芯片的优化，那么NVDA的这个优势就没了。
对的。这个绕行和优化的工程量真的不小。 太厉害了。
各方面的信息、国内已经有专门针对AI的训练卡，今年会发布
手动点赞😆😆
国内基础课非常扎实
只要测试做的好，用AI把英伟达的汇编翻成AMD或者华为的应该不是特难事。
“应该不是特难事” 呵呵，做了就知道了
你觉得，只有我们会用Deepseek，Deepseek的人不会用？ 而是一行一行敲代码。
基本上我觉得他们只要一个PTX思路对， 转换一下，验证就行，甚至不需要自己验证，训练Deepseek验证就好了
我让他，算一下向量相加，不出意外。 转换很容易。懂汇编的应该很容易验证
.version 7.0 // PTX 版本 .target sm_70 // 目标架构 (sm_70 代表 NVIDIA Volta 架构) .address_size 64 // 64 位地址空间
// 定义向量加法的 GPU 入口函数 .visible .entry vectorAdd( .param .u64 A, // 传入的指针参数 A .param .u64 B, // 传入的指针参数 B .param .u64 C, // 传入的指针参数 C .param .u32 N) // 传入的数组大小 N { .reg .u32 tid, idx, Nval; // 定义 32 位寄存器 .reg .u64 Aptr, Bptr, Cptr; .reg .f32 Aval, Bval; // 定义 32 位浮点寄存器
// 获取线程索引（blockIdx.x * blockDim.x + threadIdx.x） mov.u32 tid, %tid.x; mov.u32 idx, %ctaid.x; mad.lo.u32 idx, idx, %ntid.x, tid; // idx = blockIdx.x * blockDim.x + threadIdx.x
// 加载参数 ld.param.u64 Aptr, [A]; // 读取 A 的指针 ld.param.u64 Bptr, ; // 读取 B 的指针 ld.param.u64 Cptr, [C]; // 读取 C 的指针 ld.param.u32 Nval, [N]; // 读取向量大小
// 仅当索引在范围内时执行 setp.ge.u32 %p1, idx, Nval; @%p1 ret;
// 计算指针偏移量 mul.wide.u32 Aptr, idx, 4; mul.wide.u32 Bptr, idx, 4; mul.wide.u32 Cptr, idx, 4;
// 读取 A 和 B 的值 ld.global.f32 Aval, [Aptr]; ld.global.f32 Bval, [Bptr];
// 计算 C = A + B add.f32 Aval, Aval, Bval; // 存储结果 st.global.f32 [Cptr], Aval;
ret; }
// AMD GCN 指令集示例 .amd_kernel vector_add // 定义一个 GPU 内核
s_load_dwordx2 s[0:1], s[4:5], 0x00 // 读取指针 A s_load_dwordx2 s[2:3], s[4:5], 0x08 // 读取指针 B s_load_dwordx2 s[6:7], s[4:5], 0x10 // 读取指针 C s_load_dword s8, s[4:5], 0x18 // 读取 N
v_mov_b32 v0, s8 // 将 N 赋值给 v0 v_add_u32 v1, vcc, v0, v0 // 计算索引 v_mul_lo_u32 v2, v1, 4 // 计算偏移量 (index * 4)
v_add_u32 v3, vcc, s0, v2 // A + index v_add_u32 v4, vcc, s2, v2 // B + index v_add_u32 v5, vcc, s6, v2 // C + index
flat_load_dword v6, v3 // 读取 A flat_load_dword v7, v4 // 读取 B
v_add_f32 v8, v6, v7 // C = A + B
flat_store_dword v5, v8 // 写入 C
s_endpgm // 结束
美国做AI的大部分也是国内EECS毕业的 基本上是他们的师兄师姐
要看amd华为等等多久开发出替代芯片了
女大利润率极高，其他公司很有动力去开发替代产品
只要是国内读过EE本科，肯定上过计算机基础原理这门课，基本的汇编肯定要学的，不存在小年轻就不懂汇编的情况。但是美本就玄学了，因为美本选课自由度太大，可能正好就把这门课错过了。我同学是做硬件的，跟我说新来的小年轻听都没听说过verilog（这不是汇编，是高级一点的硬件语言），啥都得他自己做。。。
女大在cuda上浸淫了多久？小20年。看看这个帖子就可笑。信这个还不如信母猪会飞。
科技的发展日新月异，千万别说不可能。
AI来了，经验那一套已经不管用了，ai就是乱拳打死老师傅
要是信，就应该真金白银做空女大。看看按摩店、intel跌倒多惨，女大也一样。
微软，Amazon，meta等苦女大好几年了。
Apple和女大不对付更久了。
应验了我前几天说的话
没有美国的禁运，DS可能也不会这么优化到极致吧，还不是逼出来的
Re，现在CS专业就根本不设汇编语言的课程了，不知道ee专业有没设
大哥，时代变了
如果你不知道gerganov是谁，你总知道ggml吧
如果你不知道ggml，你总知道llama.cpp吧
gerganov就是ggml和llama.cpp的作者， 一个晚上把初代llama模型移植到c++，在raspberryPi上跑起来的狠人。
结果前两天有人用deepseek写汇编，把他ggml里的核心代码加速了2倍
下图就是他的评论：
系统提示：若遇到视频无法播放请点击下方链接
https://x.com/ggerganov/status/1883888097185927311
汇编的底层逻辑都是高度类似的。这叫ISA
到了CUDA runtime层，才会因为编程模型和执行模型的区别，差异变得更大起来。
这就好比不管你什么汽车，你的发动机，变速箱，结构都大同小异。
没错，过去为了一个关键的performance会直接用汇编跟机器直接打交道，随着机器指令运行越来越快，这部分的工作已经没有多大意义。现在的机器代码只能在系统启动中还能看到