根本没有绕过英伟达的垄断,反而绑定的更深了。因为 PTX 就是 CUDA 技术体系的一部分,更为底层。
有人拿 Java的字节码对比,认为其它厂商可以“照猫画虎”,干掉英伟达。
这个回答是错误的,因为他连基本的事实和相关的概念都搞错了。了解计算机编程语言发展历史的人都会知道编程语言的概念,但是为了说明,我用最简略、通俗的内容来解释,别挑严谨性的毛病,因为这俩事情冲突。
计算机刚出世的时候,控制机器的办法非常原始,最开始人类只有“机器语言”,输入的办法是:穿孔纸带。上面用是否打孔来表示0/1 – 计算机内部是二进制的。
但是,能干这活的人实在太少了,凤毛麟角。不然你给解读下这段啥意思:
00101010110010110110101010101010101110110110101010101010100100101010101011011010101
逐渐,聪明的人类想出了新办法,搞出了汇编语言,类似这样:
mov dword [number1_ptr], number1 ; 将number1的地址存储到临时变量number1_ptr(这里为了说明,实际可以直接操作)movzx eax, word [number1_ptr] ; 将number1的值(5)加载到EAXmov dword [number2_ptr], number2 ; 将number2的地址存储到临时变量number2_ptrmovzx ebx, word [number2_ptr] ; 将number2的值(10)加载到EBXadd eax, ebx ; EAX和EBX相加,结果存储在EAX(EAX = 15)mov dword [result_ptr], eax ; 将EAX的值存储到result的地址(这里同样为了说明,实际可以直接操作)
Code language: ARM Assembly (armasm)
这可读性,可比那一堆的0101110强太多了!但是问题来了,计算机不懂这个的啊,这是为了方便给人类看的。所以需要一个编译器,把上面这段翻译成机器语言,再去执行。我上学时候还写过汇编呢。
又过了多年,人类发现另外的问题,汇编语言虽然可读性比机器语言好,但是仍然很麻烦,开发效率太差。另外一个严重问题:严重依赖硬件平台,可移植性很差。一个平台上的汇编代码,换个平台完全不可用,需要重写。因为它直接引用了硬件平台专属的设施和指令,比如寄存器、存取内存等。
于是,人类又搞出来高级语言,如大名鼎鼎的C语言:
#include <stdio.h> // 包含标准输入输出库int main() { printf("hello cuda\n"); // 打印字符串,\n表示换行 return 0; // 程序正常结束}
Code language: PHP (php)
它在方方面面都比汇编语言优越太多了,于是迅速的成为人类的最爱。在各种设备、场合大行其道。因为编译器的进步,可移植性问题也解决了,C语言几乎被移植到了各种智能设备上运行。编译器先把C语言翻译成目标平台的汇编代码,再进一步编译成机器语言,形成可执行程序。
但是,有时候为了提升性能,会在C代码中嵌入汇编语言,这是可以的,类似这样:
#include <stdio.h>int main() { char str[] = "Hello, cuda!\n"; // 内联汇编代码 __asm__ __volatile__ ( "movl $1, %%eax\n" // 系统调用号:1 (sys_write) "movl $1, %%edi\n" // 文件描述符:1 (stdout) "lea 1f(%%rip), %%rsi\n" // 字符串地址 "movl $14, %%edx\n" // 字符串长度 "syscall\n" // 触发系统调用 "1:\n" // 标签,用于字符串的相对地址计算 ".asciz \"%s\"\n" // 字符串 : : "s" (str) // 输入操作数,将C变量str传递给汇编代码 : "%eax", "%edi", "%rsi", "%edx" // 被修改的寄存器 ); return 0;}
Code language: PHP (php)
英伟达的CUDA的平台,扩展了C语言,编程语法相近:
// 核函数(运行在GPU上)__global__ void vectorAdd(const float *A, const float *B, float *C, int numElements) { int i = blockDim.x * blockIdx.x + threadIdx.x; // 计算全局线程索引 if (i < numElements) { C[i] = A[i] + B[i]; // 执行加法操作 }}// CUDA内核函数 - 内嵌了 PTX 代码__global__ void vectorAdd(const float* a, const float* b, float* c, int n) { int idx = blockIdx.x * blockDim.x + threadIdx.x; if (idx < n) { // 调用PTX代码 asm volatile( "{\n" " .reg .f32 r0, r1, r2;\n" " ld.param.f32 r0, [%0];\n" " ld.param.f32 r1, [%1];\n" " add.f32 r2, r0, r1;\n" " st.param.f32 [%2], r2;\n" "}" : "=f"(c[idx]) // 输出 : "f"(a[idx]), "f"(b[idx]) // 输入 : "r0", "r1", "r2" // 临时寄存器 ); }}
Code language: JavaScript (javascript)
如果把 CUDA 程序看作 C 语言,那 PTX 可以看成其底层平台的汇编语言。英伟达的 PTX 代码是为了自家产品方便兼容、移植而设计的,是其专属的汇编语言。即便这样,也有产品的新架构适配问题需要他们来处理。
越高层的语言可移植性越好,越底层的语言越难于移植。因为高层可以抽象,编译器可以针对平台进行编译。但是底层语言如汇编,会直接引用硬件抽象,比如寄存器、内存访问等,这些指令严重依赖于硬件平台。比如在Arm平台的汇编代码和X86的就不一样。硬说这个更“可移植”,简直是笑话。
Java的字节码并不相同。Java的字节码是为了跨平台使用:一次编写,到处运行。它需要Jvm运行,虽然可以使用jit – 即时编译技术,但是并不是汇编代码,并不会与特定硬件平台绑定,这点与汇编语言截然不同。实则它是在高级语言和汇编语言的中间,更接近高级语言。如下例子:
Compiled from "HelloWorld.java"public class HelloWorld { public HelloWorld(); Code: 0: aload_0 1: invokespecial #1 // Method java/lang/Object."<init>":()V 4: return public static void main(java.lang.String[]); Code: 0: getstatic #2 // Field java/lang/System.out:Ljava/io/PrintStream; 3: ldc #3 // String Hello, World! 5: invokevirtual #4 // Method java/io/PrintStream.println:(Ljava/lang/String;)V 8: return}
Code language: PHP (php)
Java的字节码面向Java虚拟机 – JVM,不直接为实体的 CPU 生成,不包含指挥 CPU 如何运行的具体硬件指令。而汇编语言不同,它的指令就是个助记符号,完全面向具体的 CPU 编写设计,里面充斥了大量的真实硬件访问、计算指令。对汇编语言来说,换个不同平台的 CPU,如 X86 -> Arm,只能重写,别无它法。因为硬件完全不一样,这个平台CPU具备的指令,另外一个平台的 CPU 可能根本就没有,即便相似的,操作模式也不同。对汇编代码还幻想着它能“移植”,纯粹是做梦 – 除非你把重写当移植,但对于计算机专业的人而言,这只能算狡辩,专业能力是不合格的。
所以,想当然的把 CUDA PTX 语言看成 Java 字节码是严重的错误,估计是没学过汇编语言导致的。
// PTX 代码例子:.visible .entry test(int&)( .param .u64 test(int&)_param_0){ ld.param.u64 %rd1, [test(int&)_param_0]; cvta.to.global.u64 %rd2, %rd1; mov.u32 %r1, %ctaid.x; st.global.u32 [%rd2], %r1; ret;}
Code language: PHP (php)然而,英伟达的 GPU 比 CPU 的情况更为复杂,PTX 代码还需要进一步针对不同代、不同核心的 GPU 编译成更低级的代码,输出结果更有硬件针对性,更像汇编代码。这也是 PTX 设计出台的原因:对开发者尽可能的屏蔽自家的产品细节差别,简化开发。具体细节不再说明,有兴趣的可以去查阅。
此外,英伟达在其CUDA软件的最终用户许可协议(EULA)中明确禁止在非英伟达硬件平台上通过翻译层运行基于CUDA的软件。想翻译移植 PTX,先考虑怎么面对英伟达的豪华律师团吧。
“You may not reverse engineer, decompile or disassemble any portion of the output generated using SDK elements for the purpose of translating such output artifacts to target a non-NVIDIA platform
所以,PTX 的使用是一种性能优化,根本谈不上“绕开垄断”,因为它就是CUDA技术体系的核心组成部分。使用 PTX,反而对英伟达平台绑定的更深了。
其它公司的竞品GPU,想达成类似目标,只能另起炉灶。谁叫英伟达作为技术先行者,在这个领域耕耘了这么多年呢?想干掉英伟达?世界上哪里有这么简单的事情。