5.1 性能优化
5.1.1 计算优化
注:看好c与cann的不同写法,主动去理解api
循环展开(Loop Unrolling)
循环展开就是把循环体复制多次,减少循环控制的开销。比如原来循环100次,展开成每次处理4个元素,循环25次。
// 未展开的循环for(inti=0;i<n;i++){c[i]=a[i]+b[i];}// 展开后的循环(每次处理4个元素)for(inti=0;i<n;i+=4){c[i]=a[i]+b[i];c[i+1]=a[i+1]+b[i+1];c[i+2]=a[i+2]+b[i+2];c[i+3]=a[i+3]+b[i+3];}在Ascend C中,向量化本身就有循环展开的效果,因为向量指令一次处理多个元素。但有时候手动展开一些外层循环,或者展开标量循环,也能提升性能。
展开的好处:减少循环控制开销,增加指令级并行度,给编译器更多优化空间。
展开的代价:代码变长,可能增加寄存器压力,展开太多可能影响缓存。
展开多少合适:通常展开2-8倍比较合适,具体要看循环体大小和硬件资源。
指令流水线优化
昇腾处理器的AI Core有指令流水线,可以同时执行多条指令。但如果有数据依赖,流水线会阻塞,影响性能。
减少数据依赖:尽量让指令之间没有依赖,或者依赖距离足够远,这样流水线不会阻塞。
// 有依赖的代码(慢)vec1=Load(...);vec2=Add(vec1,...);// 依赖vec1vec3=Mul(vec2,...);// 依赖vec2// 减少依赖的代码(快)vec1=Load(...);vec2=Load(...);// 可以并行加载vec3=Add(vec1,vec2);// 依赖vec1和vec2,但vec1和vec2已经准备好了指令重排:编译器会自动重排指令,但有时候手动重排也能帮助编译器。
指令合并:用融合指令,比如FusedMulAdd,一条指令完成乘加,减少指令数和依赖。
计算融合(Fusion)
计算融合就是把多个操作合并成一个,减少中间结果的内存访问。
元素级融合:比如ReLU和Add可以融合成AddRelu,一条指令完成加法和ReLU。
// 分开写(慢)Add(temp,a,b);Relu(output,temp);// 融合(快)AddRelu(output,a,b);// output = relu(a + b)规约融合:比如先求和再除以个数得到均值,可以融合成一个操作。
多算子融合:比如Conv + BN + ReLU可以融合成一个算子,减少内存访问和kernel启动开销。
融合的好处是减少内存访问,减少kernel启动开销,但实现起来复杂一些。
5.1.2 内存优化
内存访问模式优化
内存访问模式对性能影响很大,好的访问模式能充分利用缓存和预取。
连续访问:尽量连续访问内存,这样缓存效果好,预取也有效。
// 连续访问(快)for(inti=0;i<n;i++){c[i]=a[i]+b[i];// 连续访问a、b、c}// 随机访问(慢)for(inti=0;i<n;i++){c[i]=a[indices[i]]+b[indices[i]];// 通过索引访问,不连续}访问对齐:数据要对齐,不对齐可能跨缓存行,慢很多。
合并访问:如果多个小块数据要访问,尽量合并成一次大访问。
数据重用和缓存
如果数据要用多次,尽量在Local Memory里多待一会儿,不要用完就扔。
数据重用:识别可以重用的数据,尽量在Local Memory里缓存。
// 数据重用示例Vector<float,256>vec;vec.Load(local_tensor,offset);// 用vec做多个运算Add(result1,vec,vec);Mul(result2,vec,vec);Sqrt(result3,vec);// vec不用重复加载缓存友好的代码:让访问的数据在时间和空间上都有局部性,这样缓存命中率高。
预取:如果知道接下来要用什么数据,可以提前加载,让加载和计算重叠。
内存对齐优化
内存对齐很重要,不对齐的话性能会下降很多。
对齐要求:向量加载通常要求数据按向量长度对齐,比如256元素的向量要求256*sizeof(float)字节对齐。
// 确保对齐alignas(256*sizeof(float))floatdata[1024];对齐检查:写代码的时候要检查数据是否对齐,不对齐的话要处理。
对齐填充:如果数据不对齐,可以填充一些无效数据,让有效数据对齐。
5.1.3 并行优化
多核并行策略
昇腾处理器有多个AI Core,可以并行执行。设计并行策略要考虑:
数据分块:数据怎么分块,让每个Core处理不同的块。
负载均衡:让每个Core的负载尽量均衡,避免有些Core很忙有些Core很闲。
同步开销:并行的时候可能需要同步,同步有开销,要尽量减少同步次数。
// 并行策略示例voidParallelProcess(){int32_tcore_id=GetCoreId();int32_ttotal_cores=GetTotalCores();// 计算这个Core要处理的数据范围int32_tblock_size=total_data/total_cores;int32_tstart=core_id*block_size;int32_tend=(core_id==total_cores-1)?total_data:start+block_size;// 处理这个范围ProcessBlock(start,end);// 如果需要,同步一下Barrier();}数据分块策略
数据分块要考虑:
块大小:块太小了利用率低,块太大了Local Memory装不下。通常块大小要适合Local Memory,也要适合硬件特性。
边界处理:分块的时候要注意边界,确保所有数据都被处理,不重复不遗漏。
分块方式:可以按行分块、按列分块、按块分块,选择哪种要看数据访问模式。
负载均衡优化
负载不均衡会影响整体性能,要优化负载均衡。
动态分配:根据Core的负载情况,动态分配任务。负载轻的Core多分点任务。
工作窃取:空闲的Core可以从忙碌的Core那里"偷"一些任务来做。
自适应分块:根据数据特点,动态调整分块大小,让每个Core的负载尽量均衡。
性能监控:监控每个Core的负载情况,找出瓶颈,调整策略。
5.2 精度优化
5.2.1 混合精度计算
什么是混合精度
混合精度就是不同地方用不同精度。比如计算用FP16(快但精度低),累加用FP32(慢但精度高),最后结果转回FP16。
混合精度可以在保证精度的同时提升性能,因为FP16的计算速度比FP32快很多。
混合精度的实现
// 概念性示例voidMixedPrecisionCompute(){// 输入是FP16Vector<half_t,256>input_fp16;// 计算用FP16Vector<half_t,256>temp_fp16;Mul(temp_fp16,input_fp16,input_fp16);// 累加用FP32Vector<float,256>temp_fp32;Cast(temp_fp32,temp_fp16);// FP16转FP32// 累加Scalar<float>sum_fp32;ReduceSum(sum_fp32,temp_fp32);// 结果转回FP16Scalar<half_t>sum_fp16;Cast(sum_fp16,sum_fp32);}实现要点:关键是要知道哪些地方用FP16,哪些地方用FP32。通常计算用FP16,累加、归一化这些用FP32。
混合精度的优势
性能提升:FP16的计算速度是FP32的2倍,内存带宽也是2倍,性能提升明显。
精度保证:关键操作用FP32,保证精度不损失太多。
内存节省:FP16占的内存是FP32的一半,可以处理更大的模型。
5.2.2 数值稳定性
数值溢出和下溢
FP16的数值范围比FP32小,容易溢出和下溢。要避免数值问题。
溢出处理:如果值太大,会溢出成无穷大或NaN。可以检查溢出,或者用FP32计算。
下溢处理:如果值太小,会下溢成0。可以加个小的epsilon,避免除零。
梯度裁剪:在训练的时候,可以裁剪梯度,避免梯度爆炸。
归一化的稳定性
归一化操作(比如BatchNorm)要保证数值稳定。
方差计算:计算方差的时候,如果方差太小,除以方差会放大误差。可以加个epsilon。
// 数值稳定的归一化Scalar<float>mean,var,eps=1e-5;ComputeMeanAndVar(mean,var,input);// 加epsilon避免除零var=var+eps;Sqrt(var,var);// 归一化Sub(normalized,input,mean);Div(normalized,normalized,var);累加的精度
累加操作容易丢失精度,特别是累加很多小值的时候。
Kahan累加:用Kahan算法累加,可以减少精度损失。
分块累加:先对块累加,再对块结果累加,精度损失小一些。
5.2.3 误差分析
误差来源
数值计算的误差主要来自:
舍入误差:浮点数表示有精度限制,每次运算都可能舍入。
截断误差:算法本身的近似,比如用泰勒级数近似函数。
累积误差:误差会累积,运算越多误差越大。
误差控制
控制误差的方法:
提高精度:关键计算用FP32,减少舍入误差。
算法改进:用数值稳定的算法,减少误差累积。
误差监控:监控误差,如果误差太大,调整策略。
5.2.4 精度与性能平衡
精度要求
不同的应用对精度要求不同:
训练:训练通常用FP16,精度损失可以接受,性能更重要。
推理:推理可能对精度要求高,可以用FP32,或者混合精度。
特定场景:有些场景对精度要求很高,必须用FP32甚至更高精度。
性能要求
性能要求也要考虑:
实时性:实时应用对性能要求高,可能牺牲一点精度。
吞吐量:批量处理可能更看重吞吐量,可以用FP16。
延迟:低延迟应用可能用FP16,减少计算时间。
平衡策略
平衡精度和性能:
关键路径用高精度:影响最终结果的关键计算用高精度。
非关键路径用低精度:不影响结果的中间计算用低精度。
动态调整:根据实际情况动态调整精度,找到最佳平衡点。
5.3 算子融合
5.3.1 算子融合原理
为什么需要融合
算子融合可以减少:
内存访问:融合后中间结果不用写回Global Memory,减少内存访问。
Kernel启动开销:多个算子融合成一个,减少kernel启动次数。
数据重用:融合后数据可以在Local Memory里重用,不用重复加载。
融合的收益
融合的收益主要体现在:
性能提升:减少内存访问和kernel启动,性能提升明显。
延迟降低:减少数据传输,延迟降低。
资源利用:更好地利用硬件资源,提高利用率。
5.3.2 融合策略设计
可融合的算子
常见的可融合组合:
元素级融合:Add + ReLU、Mul + Add、Conv + BN等。
规约融合:Sum + Div(求均值)、Max + ArgMax等。
多算子融合:Conv + BN + ReLU、MatMul + Add + ReLU等。
融合的条件
算子要能融合,需要满足:
数据流兼容:前一个算子的输出是后一个算子的输入,数据流要匹配。
计算兼容:两个算子的计算可以合并,不会冲突。
内存兼容:融合后内存使用要合理,不能超出限制。
融合策略
设计融合策略要考虑:
融合顺序:哪些算子先融合,哪些后融合。
融合粒度:融合到什么程度,是简单融合还是深度融合。
融合代价:融合有实现成本,要考虑是否值得。
5.3.3 融合算子实现
简单融合
简单融合就是把两个算子合并成一个kernel:
// 概念性示例:Add + ReLU融合extern"C"__global__ __aicore__voidAddReluKernel(GlobalTensor<float>input1,GlobalTensor<float>input2,GlobalTensor<float>output,int32_ttotal_elements){LocalTensor<float>local_input1,local_input2,local_output;local_input1.Alloc(total_elements);local_input2.Alloc(total_elements);local_output.Alloc(total_elements);DataCopy(local_input1,input1,total_elements);DataCopy(local_input2,input2,total_elements);// 融合:Add + ReLUconstint32_tvector_length=256;for(int32_ti=0;i<total_elements;i+=vector_length){Vector<float,256>vec1,vec2,vec_out;vec1.Load(local_input1,i);vec2.Load(local_input2,i);// 直接用AddRelu融合指令AddRelu(vec_out,vec1,vec2);// output = relu(input1 + input2)vec_out.Store(local_output,i);}DataCopy(output,local_output,total_elements);local_input1.Free();local_input2.Free();local_output.Free();}复杂融合
复杂融合比如Conv + BN + ReLU,需要仔细设计:
// 概念性示例:Conv + BN + ReLU融合extern"C"__global__ __aicore__voidConvBNReluKernel(...){// 1. 卷积计算Conv2D(...);// 2. BatchNorm(不写回,直接在Local Memory做)BatchNormLocal(...);// 3. ReLU(不写回,直接在Local Memory做)ReluLocal(...);// 4. 最后统一写回DataCopy(output,local_result,...);}复杂融合的好处是中间结果不用写回Global Memory,但实现复杂,代码维护成本高。
5.3.4 融合效果评估
性能评估
评估融合效果要看:
执行时间:融合后执行时间减少了多少。
内存带宽:内存访问减少了多少。
资源利用率:硬件资源利用率提升了多少。
精度评估
融合可能影响精度,要评估:
精度损失:融合后精度损失了多少,是否可接受。
数值稳定性:融合后数值是否稳定,会不会有溢出下溢。
成本收益分析
分析融合的成本和收益:
实现成本:实现融合需要多少工作量,代码复杂度增加多少。
维护成本:融合后的代码维护成本,调试难度。
收益:性能提升多少,是否值得投入。
决策:根据成本收益分析,决定是否融合,融合到什么程度。
学习检查点
学完这一篇,你应该能做到这些:
掌握性能优化的方法,包括计算优化、内存优化、并行优化。理解精度优化的技巧,知道如何平衡精度和性能。理解算子融合的原理和方法,能够设计和实现融合算子。能够对算子进行性能分析和优化,找出瓶颈并改进。
实践练习
性能优化实践:选择一个算子,用Profiling工具分析性能,找出瓶颈。尝试应用不同的优化方法,比如循环展开、内存访问优化、并行优化,对比优化前后的性能。
混合精度实验:实现一个算子的FP16版本和FP32版本,对比性能和精度。尝试混合精度实现,在关键地方用FP32,其他地方用FP16。
算子融合实践:实现Add + ReLU的融合算子,对比融合前后的性能。尝试更复杂的融合,比如Conv + BN + ReLU。
优化案例分析:分析一个复杂算子的优化案例,理解优化思路和方法。总结优化经验,形成自己的优化方法论。
下一步:掌握了高级优化技术后,就可以学习调试和测试了。下一章会讲调试工具、测试方法、问题排查这些,到时候你就能开发出高质量、高性能的算子了。