Q:kernel是啥意思?
A:GPU编程语境中指 在GPU上运行的函数。 一个GPU有很多个线程,每个线程可以同时调用这个函数处理不同数据(这就是为什么说千万个小学生做算术),具体书写出来就是 __global__:
// 这是一个kernel(GPU内核函数) __global__ void addOne(int* data, int n) { int i = threadIdx.x + blockIdx.x * blockDim.x; // 计算当前线程对应的索引 if (i < n) { data[i] = data[i] + 1; } }(一个线程算一个数)
调用时,会在CPU端(host)启动它,比如:
addOne<<<blocks, threads>>>(data, n);这里的<<<blocks, threads>>>表示启动多少个线程块、每个块有多少线程,GPU就会根据这个配置去调度执行。 这里配置的<<<blocks, threads>>> ,就是一个grid。一次kernel启动对应一个grid,grid 由若干 block 组成,block 由若干 thread 组成。
Q:命令行中的make啥意思?
A:意思就是编译,如果不带参数就输入make的话,默认检索到该目录下的第一个makefile,然后执行其中的编译命令。
makefile具体内容可能如下:
my_program: main.cu kernel.cu nvcc main.cu kernel.cu -o my_program这行就明确告诉make:要生成my_program这个可执行文件,需要先处理main.cu和kernel.cu这两个源文件。
当你输入
make时:
- 它先检查
my_program是否存在,如果不存在,就继续- 它检查
main.cu和kernel.cu是否存在,并且是否有对应的.o文件- 第一次编译时,没有
.o文件,所以它会分别编译main.cu和kernel.cu生成main.o和kernel.o- 最后把两个
.o文件链接成my_program
MXMACA编译过程:
第一步:你写代码(源文件)
你写的代码后缀通常是.mx或.mcu(为了和CUDA区分),里面包含:
- 主机代码(Host Code):运行在CPU上的部分,比如内存分配、数据拷贝、启动核函数等。
- 设备代码(Device Code):运行在GPU上的部分,主要是
__global__或__device__函数(核函数内核)。
第二步:编译器前端(Frontend)—— 解析与翻译
MXMACA的编译器(类似mxcc)会对你的代码做两件事:
- 语法解析:检查代码是否符合MXMACA/CUDA语法规范。
- 设备代码抽离:编译器会把源码中的设备代码(GPU代码)识别出来,并“抽取”出来单独处理。这一步非常关键,因为CPU和GPU的指令集完全不同,需要分工编译。
这一步类似于CUDA的nvcc所做的工作。
第三步:设备代码编译(GPU侧)
抽取出的设备代码会被送入设备编译器后端,进行:
- 优化:比如循环展开、常量折叠等,针对摩尔线程GPU的架构(如MTT S系列)进行特化优化。
- 生成PTX(并行线程执行)中间表示:类似于CUDA中的PTX,这是一个虚拟指令集,与具体硬件解耦。这一步保证了代码的可移植性——你写的代码在将来新型号的摩尔线程GPU上也能跑,只要新硬件的驱动能支持这个中间格式。
- 最终编译:PTX再被编译成目标GPU的机器码(SASS等价物),这是GPU真正执行的二进制指令。
第四步:主机代码编译(CPU侧)
与设备代码并行,编译器也会处理主机代码:
- 把主机代码中的“启动核函数”部分,替换为对MXMACA运行时库的调用(例如
mxLaunchKernel,实际上是封装好的运行时API)。 - 主机代码编译成普通的CPU可执行文件(x86或ARM指令)。
第五步:链接(Linking)
最后,设备代码生成的可执行模块和主机代码生成的CPU对象文件,通过链接器合并成一个完整的可执行程序。运行时,这个程序由CPU启动,通过驱动把GPU的那部分机器码加载到显存中并执行。
Q:什么是内联函数?
A:普通函数就是 找到函数地址,然后调用内部语句,内联函数就是不用找地址,而是直接展开对应的代码行。这样的好处就是减少跳转时候的开销,坏处是,如果调用频率不高,或内联函数太大了的话,编译会很久。