公司内部网站建设1免费建站网站

公司内部网站建设,1免费建站网站,海关数据查询平台官网,徐州做网站软件4.1 基础算子 4.1.1 元素级算子#xff08;Element-wise#xff09; 什么是元素级算子 元素级算子就是对输入张量的每个元素独立进行运算#xff0c;元素之间没有依赖关系。这种算子最容易并行化#xff0c;也最容易向量化。 常见的元素级算子有Add、Sub、Mul、Div这些算术…4.1 基础算子4.1.1 元素级算子Element-wise什么是元素级算子元素级算子就是对输入张量的每个元素独立进行运算元素之间没有依赖关系。这种算子最容易并行化也最容易向量化。常见的元素级算子有Add、Sub、Mul、Div这些算术运算还有ReLU、Sigmoid这些激活函数。Add算子实现Add算子就是两个输入相加输出对应位置的元素和。// 概念性示例externC__global__ __aicore__voidAddKernel(GlobalTensorfloatinput1,GlobalTensorfloatinput2,GlobalTensorfloatoutput,int32_ttotal_elements){// 分配LocalTensorLocalTensorfloatlocal_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);// 向量化计算constint32_tvector_length256;for(int32_ti0;itotal_elements;ivector_length){int32_tactual_lengthmin(vector_length,total_elements-i);Vectorfloat,256vec1,vec2,vec_out;vec1.Load(local_input1,i,actual_length);vec2.Load(local_input2,i,actual_length);// 向量加法Add(vec_out,vec1,vec2);vec_out.Store(local_output,i,actual_length);}// 结果写回DataCopy(output,local_output,total_elements);// 释放内存local_input1.Free();local_input2.Free();local_output.Free();}实现要点元素级算子实现起来比较简单主要是向量化循环用Add API做向量加法。注意处理边界情况如果元素数不是向量长度的倍数要处理剩余的元素。Sub、Mul、Div算子Sub、Mul、Div算子和Add类似只是用的API不同Sub算子用Sub API两个向量相减。Mul算子用Mul API两个向量相乘。Div算子用Div API两个向量相除。实现思路都一样就是加载数据、向量化计算、写回结果。ReLU激活函数ReLU是max(0, x)就是如果x小于0就输出0否则输出x。// 概念性示例externC__global__ __aicore__voidReluKernel(GlobalTensorfloatinput,GlobalTensorfloatoutput,int32_ttotal_elements){LocalTensorfloatlocal_input,local_output;local_input.Alloc(total_elements);local_output.Alloc(total_elements);DataCopy(local_input,input,total_elements);constint32_tvector_length256;for(int32_ti0;itotal_elements;ivector_length){int32_tactual_lengthmin(vector_length,total_elements-i);Vectorfloat,256vec_in,vec_out;vec_in.Load(local_input,i,actual_length);// ReLUmax(0, x)Scalarfloatzero0.0f;Maxs(vec_out,vec_in,zero);// vec_out max(vec_in, 0)vec_out.Store(local_output,i,actual_length);}DataCopy(output,local_output,total_elements);local_input.Free();local_output.Free();}ReLU实现用Maxs API向量和标量0比较取最大值。Sigmoid和TanhSigmoid和Tanh是更复杂的激活函数需要用到数学函数API。Sigmoidsigmoid(x) 1 / (1 exp(-x))// 概念性示例// Sigmoid实现Vectorfloat,256vec_in,vec_neg,vec_exp,vec_one,vec_out;// vec_neg -vec_inMuls(vec_neg,vec_in,-1.0f);// vec_exp exp(-vec_in)Exp(vec_exp,vec_neg);// vec_one 1 exp(-vec_in)Scalarfloatone1.0f;Adds(vec_one,vec_exp,one);// vec_out 1 / (1 exp(-vec_in))Reciprocal(vec_out,vec_one);Tanhtanh(x) (exp(2x) - 1) / (exp(2x) 1)实现思路类似用Exp、Add、Sub、Div这些API组合起来。Abs、Sqrt、Exp、Log这些是数学函数算子直接用对应的API就行Abs用Abs API取绝对值。Sqrt用Sqrt API开平方。Exp用Exp API自然指数。Ln用Ln API自然对数。实现都很简单向量化循环调用对应的API。4.1.2 规约算子Reduction什么是规约算子规约算子就是把一个张量的所有元素归约成一个值或几个值。比如Sum把所有元素加起来Max找最大值。规约算子的特点是需要遍历所有元素然后归约。实现起来比元素级算子复杂一些。Sum算子实现Sum算子就是把所有元素加起来。// 概念性示例externC__global__ __aicore__voidSumKernel(GlobalTensorfloatinput,GlobalTensorfloatoutput,int32_ttotal_elements){LocalTensorfloatlocal_input;local_input.Alloc(total_elements);DataCopy(local_input,input,total_elements);// 先对每个块求和constint32_tvector_length256;LocalTensorfloatblock_sums;int32_tnum_blocks(total_elementsvector_length-1)/vector_length;block_sums.Alloc(num_blocks);for(int32_ti0;itotal_elements;ivector_length){int32_tactual_lengthmin(vector_length,total_elements-i);Vectorfloat,256vec;vec.Load(local_input,i,actual_length);// 对当前块求和Scalarfloatblock_sum;ReduceSum(block_sum,vec);// 保存块的和block_sum.Store(block_sums,i/vector_length);}// 对所有块的和再求和Scalarfloattotal_sum0.0f;for(int32_ti0;inum_blocks;i){Scalarfloatblock_sum;block_sum.Load(block_sums,i);total_sumtotal_sumblock_sum;}// 写回结果total_sum.Store(output,0);local_input.Free();block_sums.Free();}实现要点规约算子通常分两步先对每个块归约再对所有块的结果归约。这样可以充分利用向量化。Mean算子Mean就是Sum除以元素个数// 概念性示例// 先求和SumKernel(input,sum_output,total_elements);// 再除以元素个数Scalarfloatsum,mean;sum.Load(sum_output,0);meansum/total_elements;mean.Store(output,0);Max和Min算子Max和Min找最大值和最小值实现思路和Sum类似用ReduceMax和ReduceMin API。// 概念性示例// Max实现Vectorfloat,256vec;Scalarfloatmax_val;ReduceMax(max_val,vec);ArgMax和ArgMinArgMax和ArgMin不仅要找最大值最小值还要找对应的索引位置。// 概念性示例// ArgMax实现Vectorfloat,256vec;Scalarfloatmax_val;Scalarint32_tmax_idx;ReduceMaxWithIndex(max_val,max_idx,vec);// 同时返回值和索引4.1.3 索引算子Gather算子Gather算子根据索引从输入张量中收集元素。// 概念性示例externC__global__ __aicore__voidGatherKernel(GlobalTensorfloatinput,GlobalTensorint32_tindices,GlobalTensorfloatoutput,int32_tnum_indices,int32_tinput_size){LocalTensorfloatlocal_input,local_output;LocalTensorint32_tlocal_indices;local_input.Alloc(input_size);local_indices.Alloc(num_indices);local_output.Alloc(num_indices);DataCopy(local_input,input,input_size);DataCopy(local_indices,indices,num_indices);// 根据索引收集元素for(int32_ti0;inum_indices;i){Scalarint32_tidx;idx.Load(local_indices,i);Scalarfloatvalue;value.Load(local_input,idx.GetValue());value.Store(local_output,i);}DataCopy(output,local_output,num_indices);local_input.Free();local_indices.Free();local_output.Free();}实现要点Gather需要根据索引访问输入索引访问可能不是连续的所以性能可能不如元素级算子。Scatter算子Scatter是Gather的反向操作根据索引把值散布到输出张量中。// 概念性示例// Scatter实现for(int32_ti0;inum_indices;i){Scalarint32_tidx;Scalarfloatvalue;idx.Load(local_indices,i);value.Load(local_input,i);// 散布到输出value.Store(local_output,idx.GetValue());}注意Scatter可能有索引冲突多个输入值要散布到同一个位置需要定义合并规则比如相加、取最大值等。IndexSelect算子IndexSelect是Gather的特殊情况沿着某个维度选择元素。实现思路和Gather类似但需要处理多维张量的索引计算。4.2 矩阵运算算子4.2.1 矩阵乘法MatMulMatMul的基本原理矩阵乘法是深度学习里最常用的运算。两个矩阵A和B相乘A是[m, k]B是[k, n]结果C是[m, n]。计算规则是C[i, j] sum(A[i, :] * B[:, j])MatMul的实现思路矩阵乘法实现起来比较复杂主要考虑分块Tiling大矩阵要分块处理每块的大小要适合Local Memory。数据重用A的行和B的列可以重用要充分利用这个特性。Cube Unit利用昇腾处理器有专门的矩阵乘法单元Cube Unit要充分利用。// 概念性示例简化版externC__global__ __aicore__voidMatMulKernel(GlobalTensorfloatA,GlobalTensorfloatB,GlobalTensorfloatC,int32_tM,int32_tK,int32_tN){// 分块大小constint32_ttile_m64;constint32_ttile_k64;constint32_ttile_n64;// 分配LocalTensorLocalTensorfloatlocal_A,local_B,local_C;local_A.Alloc(tile_m*tile_k);local_B.Alloc(tile_k*tile_n);local_C.Alloc(tile_m*tile_n);// 初始化C为0Memset(local_C,0,tile_m*tile_n);// 分块计算for(int32_tm0;mM;mtile_m){for(int32_tn0;nN;ntile_n){// 初始化当前块的CMemset(local_C,0,tile_m*tile_n);// K维度累加for(int32_tk0;kK;ktile_k){// 加载A的块LoadTile(local_A,A,m,k,tile_m,tile_k,M,K);// 加载B的块LoadTile(local_B,B,k,n,tile_k,tile_n,K,N);// 矩阵乘法C A * BMatMul(local_C,local_A,local_B,local_C);}// 写回C的块StoreTile(C,local_C,m,n,tile_m,tile_n,M,N);}}local_A.Free();local_B.Free();local_C.Free();}实现要点矩阵乘法要分块充分利用Cube Unit。分块大小要合适太小了利用率低太大了Local Memory装不下。MatMul的优化技巧数据布局优化矩阵数据怎么存储很重要按行存储还是按列存储影响访问效率。K维度分块K维度要分块累加这样可以重用A和B的数据。流水线化加载下一块数据、计算当前块、写回上一块结果可以流水线化。4.2.2 矩阵转置TransposeTranspose的基本原理矩阵转置就是把矩阵的行列互换A[i, j]变成A[j, i]。Transpose的实现// 概念性示例externC__global__ __aicore__voidTransposeKernel(GlobalTensorfloatinput,GlobalTensorfloatoutput,int32_theight,int32_twidth){LocalTensorfloatlocal_input,local_output;local_input.Alloc(height*width);local_output.Alloc(height*width);DataCopy(local_input,input,height*width);// 转置for(int32_ti0;iheight;i){for(int32_tj0;jwidth;j){Scalarfloatvalue;value.Load(local_input,i*widthj);value.Store(local_output,j*heighti);}}DataCopy(output,local_output,height*width);local_input.Free();local_output.Free();}实现要点转置的访问模式不是连续的可能影响性能。可以用向量化优化或者用专门的转置指令。4.2.3 批处理矩阵乘法BatchMatMulBatchMatMul的基本原理BatchMatMul就是多个矩阵乘法一起做。输入是[batch, m, k]和[batch, k, n]输出是[batch, m, n]。BatchMatMul的实现// 概念性示例externC__global__ __aicore__voidBatchMatMulKernel(GlobalTensorfloatA,GlobalTensorfloatB,GlobalTensorfloatC,int32_tbatch,int32_tM,int32_tK,int32_tN){for(int32_tb0;bbatch;b){// 每个batch的矩阵乘法GlobalTensorfloatA_batchA[b];GlobalTensorfloatB_batchB[b];GlobalTensorfloatC_batchC[b];MatMulKernel(A_batch,B_batch,C_batch,M,K,N);}}实现要点BatchMatMul可以并行处理不同的batch充分利用多核。4.2.4 矩阵分解相关算子矩阵分解算子比如SVD、QR分解这些实现起来比较复杂通常需要调用专门的数学库或者用迭代算法实现。这里就不详细展开了有兴趣可以查专门的资料。4.3 卷积相关算子4.3.1 卷积Convolution卷积的基本原理卷积是深度学习的核心运算。输入是[N, C, H, W]的特征图卷积核是[K, C, Kh, Kw]输出是[N, K, H’, W’]。计算规则是对输出的每个位置用卷积核在输入上滑动窗口计算点积。卷积的实现思路卷积实现起来很复杂主要有几种方法直接卷积按照定义直接计算简单但慢。im2col转换把卷积转换成矩阵乘法用MatMul实现。Winograd算法用Winograd算法加速减少计算量。分组卷积优化如果卷积是分组的可以优化。// 概念性示例im2col方法externC__global__ __aicore__voidConv2DKernel(GlobalTensorfloatinput,// [N, C, H, W]GlobalTensorfloatweight,// [K, C, Kh, Kw]GlobalTensorfloatoutput,// [N, K, H, W]int32_tN,int32_tC,int32_tH,int32_tW,int32_tK,int32_tKh,int32_tKw,int32_tstride_h,int32_tstride_w,int32_tpad_h,int32_tpad_w){// 1. im2col把输入转换成矩阵// 2. 矩阵乘法weight * im2col(input)// 3. col2im把结果转换回特征图格式}实现要点卷积通常用im2col转换成矩阵乘法然后用MatMul实现。im2col有内存开销但实现简单性能也不错。4.3.2 池化Pooling池化的基本原理池化是对特征图做下采样常见的有MaxPooling和AvgPooling。MaxPooling在窗口内取最大值。AvgPooling在窗口内取平均值。MaxPooling实现// 概念性示例externC__global__ __aicore__voidMaxPool2DKernel(GlobalTensorfloatinput,GlobalTensorfloatoutput,int32_tN,int32_tC,int32_tH,int32_tW,int32_tkernel_h,int32_tkernel_w,int32_tstride_h,int32_tstride_w){// 对每个输出位置for(int32_tn0;nN;n){for(int32_tc0;cC;c){for(int32_toh0;ohoutput_h;oh){for(int32_tow0;owoutput_w;ow){// 在窗口内找最大值floatmax_val-INF;for(int32_tkh0;khkernel_h;kh){for(int32_tkw0;kwkernel_w;kw){int32_tihoh*stride_hkh;int32_tiwow*stride_wkw;floatvalinput[n][c][ih][iw];max_valmax(max_val,val);}}output[n][c][oh][ow]max_val;}}}}}实现要点池化可以向量化对多个通道并行处理。也可以用专门的池化指令。4.3.3 反卷积Deconvolution反卷积也叫转置卷积是卷积的逆操作。实现思路和卷积类似但更复杂一些。通常也用im2col或者专门的算法实现。4.3.4 分组卷积Group Convolution分组卷积是把输入通道分成几组每组独立做卷积最后把结果拼接起来。这样可以减少计算量适合深度可分离卷积这种场景。实现要点分组卷积可以并行处理不同的组充分利用多核。4.4 归一化算子4.4.1 BatchNormBatchNorm的基本原理BatchNorm是对每个通道在batch维度上做归一化。计算步骤计算均值和方差mean mean(x), var var(x)归一化y (x - mean) / sqrt(var eps)缩放和平移y gamma * y betaBatchNorm的实现// 概念性示例externC__global__ __aicore__voidBatchNormKernel(GlobalTensorfloatinput,GlobalTensorfloatoutput,GlobalTensorfloatgamma,GlobalTensorfloatbeta,GlobalTensorfloatrunning_mean,GlobalTensorfloatrunning_var,int32_tN,int32_tC,int32_tH,int32_tW,floateps){// 1. 计算均值和方差规约操作// 2. 归一化// 3. 缩放和平移}实现要点BatchNorm需要先做规约计算均值和方差再做元素级运算归一化、缩放、平移。4.4.2 LayerNormLayerNorm的基本原理LayerNorm是对每个样本在特征维度上做归一化。和BatchNorm的区别是归一化的维度不同。实现思路和BatchNorm类似也是先规约再归一化但规约的维度不同。4.4.3 InstanceNorm和GroupNormInstanceNorm和GroupNorm是BatchNorm的变种归一化的维度不同InstanceNorm对每个样本的每个通道独立归一化。GroupNorm把通道分组对每组独立归一化。实现思路都类似主要是规约的维度不同。学习检查点学完这一篇你应该能做到这些理解各种常用算子的实现方法包括元素级算子、规约算子、索引算子。掌握矩阵运算算子的实现特别是矩阵乘法的分块策略。了解卷积相关算子的实现思路知道im2col等方法。理解归一化算子的实现知道如何组合规约和元素级运算。实践练习实现ReLU算子实现一个ReLU激活函数算子用向量化API完成从编写到测试的全流程。实现Sum算子实现一个Sum规约算子先对块归约再对块结果归约理解规约的实现方法。优化矩阵乘法实现一个简单的矩阵乘法算子尝试不同的分块大小看看性能差异。理解Tiling策略的重要性。实现简单卷积用im2col方法实现一个简单的2D卷积理解卷积如何转换成矩阵乘法。下一步掌握了常用算子的实现后就可以学习高级优化技术了。下一章会讲性能优化、精度优化、算子融合这些高级话题到时候你就能写出高性能的算子了。