0x0. 前言

如題所述，本篇文章推薦和講解一下OneFlow ElementWise模板，FastAtomicAdd，OneFlow UpsampleNearest2d模板的用法以及原理。但OneFlow ElementWise模板的用法和原理在【BBuf的CUDA筆記】一，解析OneFlow Element-Wise 算子實現 已經講過了，所以這篇文章里不再贅述，主要講解后面2個。我將上述三個算法的實現都分別抽出來放到了 https://github.com/BBuf/how-to-optim-algorithm-in-cuda 這個工程的 elementwise/FastAtomicAdd/UpsampleNearest2D 三個文件夾中，并且三個算法的實現都分別只用一個.cu文件進行整理，使用nvcc編譯可以使用，有需要的同學請自取。

0x1. OneFlow elementwise模板

將 oneflow 的 elementwise 模板抽出來方便大家使用，這個 elementwise 模板實現了高效的性能和帶寬利用率，并且用法非常靈活。完整實驗代碼見 https://github.com/BBuf/how-to-optim-algorithm-in-cuda/blob/master/elementwise/elementwise.cu，原理講解請看：【BBuf 的CUDA筆記】一，解析OneFlow Element-Wise 算子實現 。這里以逐點乘（z = x * y，其中x，y，z是形狀完全一樣的Tensor）為例，性能和帶寬的測試情況如下 (A100 PCIE 40G)：

可以看到無論是性能還是帶寬，使用 oneflow 的 elementwise 模板相比于原始實現都有較大提升。

涉及到的主要優化技術有向量化數據訪問，選取合適的GridSize和BlockSize，循環展開和Grid-Stride Loops等技巧。

模板代碼和用法詳見：https://github.com/BBuf/how-to-optim-algorithm-in-cuda/blob/master/elementwise/elementwise.cu

0x2. FastAtomicAdd

眾所周知，atomicAdd是CUDA中非常昂貴的操作，特別是對于half類型來說 atomicAdd 巨慢無比，慢到如果一個算法需要用到 atomicAdd，那么相比于用 half ，轉成 float ，再 atomicAdd，再轉回去還要慢很多。但是我們有時候不得不去執行half類型的原子加，這個時候怎么能提升性能呢？

PyTorch給出了一個快速原子加的實現（我這里魔改了一下，去掉了一些不需要的參數，完整測試代碼見 https://github.com/BBuf/how-to-optim-algorithm-in-cuda/blob/master/FastAtomicAdd/fast_atomic_add_half.cu ）：

//FastAddisreferencedfrom
//https://github.com/pytorch/pytorch/blob/396c3b1d88d7624938a2bb0b287f2a19f1e89bb4/aten/src/ATen/native/cuda/KernelUtils.cuh#L29
template<typenameT,typenamestd::enable_if<std::is_same::value>::type*=nullptr>
__device____forceinline__voidFastSpecializedAtomicAdd(T*base,size_toffset,
constsize_tlength,Tvalue){
#if((defined(CUDA_VERSION)&&(CUDA_VERSION
atomicAdd(reinterpret_cast(base)+offset,static_cast(value));
#else
//Accountsforthechancebasefallsonanodd16bitalignment(ie,not32bitaligned)
__half*target_addr=reinterpret_cast<__half*>(base+offset);
boollow_byte=(reinterpret_cast<std::uintptr_t>(target_addr)%sizeof(__half2)==0);

if(low_byte&&offset1)){
__half2value2;
value2.x=value;
value2.y=__float2half_rz(0);
atomicAdd(reinterpret_cast<__half2*>(target_addr),value2);

}elseif(!low_byte&&offset>0){
__half2value2;
value2.x=__float2half_rz(0);
value2.y=value;
atomicAdd(reinterpret_cast<__half2*>(target_addr-1),value2);

}else{
atomicAdd(reinterpret_cast<__half*>(base)+offset,static_cast<__half>(value));
}
#endif
}

template<typenameT,typenamestd::enable_ifstd::is_same::value>::type*=nullptr>
__device____forceinline__voidFastSpecializedAtomicAdd(T*base,size_toffset,
constsize_tlength,Tvalue){
atomicAdd(base+offset,value);
}

template
__device____forceinline__voidFastAdd(T*base,size_toffset,constsize_tlength,Tvalue){
FastSpecializedAtomicAdd(base,offset,length,value);
}

也就是把half類型的原子加轉換成half2類型的原子加，為了驗證這個快速原子加相比于half類型的原子加以及pack 2個half 到 half2再執行原子加的性能表現，我實現了三個算法（.cu文件）。它們都是針對half數據類型做向量的內積，都用到了atomicAdd，保證數據的長度以及gridsize和blocksize都是完全一致的。具體如下：

1. https://github.com/BBuf/how-to-optim-algorithm-in-cuda/blob/master/FastAtomicAdd/atomic_add_half.cu 純half類型的atomicAdd。
2. https://github.com/BBuf/how-to-optim-algorithm-in-cuda/blob/master/FastAtomicAdd/atomic_add_half_pack2.cu half+pack，最終使用的是half2類型的atomicAdd。
3. https://github.com/BBuf/how-to-optim-algorithm-in-cuda/blob/master/FastAtomicAdd/fast_atomic_add_half.cu 快速原子加，雖然沒有顯示的pack，但本質上也是通過對單個half補0使用上了half2的原子加。

下面展示3個腳本通過ncu profile之后的性能表現：

可以看到使用pack half的方式和直接使用half的fastAtomicAdd方式得到的性能結果一致，均比原始的half的原子加快3-4倍。

接下來驗證一下是否存在warp分支分化問題，對比了一下fastAtomicAdd和pack half2的ncu匯編代碼，并未發現不同類型的指令：

fastAtomicAdd 計算部分：

atomicAddhalfpack2計算部分：

每一種指令的類型都能在兩份代碼中找到，初步判斷不會因為fastAtomicAdd實現中的下述if語句存在線程分化問題。

綜上所述，使用FastAtomicAdd可以大幅度提升half數據類型原子加的性能并且不需要手動Pack，使用方法更加簡單。

模板代碼和用法詳見：https://github.com/BBuf/how-to-optim-algorithm-in-cuda/blob/master/FastAtomicAdd/fast_atomic_add_half.cu

0x3. Oneflow Upsample模板

在Stable Diffusion的反向擴散過程中使用到了UNet，而UNet中存在大量的UpsampleNearest2D上采樣。PyTorch對于UpsampleNearest都是通用的實現（https://github.com/pytorch/pytorch/blob/master/aten/src/ATen/native/cuda/UpSampleNearest2d.cu#L112-L163） ，這種實現里面存在大量的取模和坐標映射操作（nn_bw_compute_source_index_fn）以及循環統計貢獻等。對于深度學習來說，UpsampleNearest最常用的其實就是2倍上采樣，比如Unet和YOLOv5，所以我們完全可以針對這種情況寫一個特化的Kernel，很輕量的來完成2倍上采樣的計算。下面展示OneFlow中針對2倍上采樣的優化（代碼見：https://github.com/BBuf/how-to-optim-algorithm-in-cuda/blob/master/UpsampleNearest2D/upsample_nearest_2d.cu#L16-L63）

//CUDA:gridstridelooping
#defineCUDA_1D_KERNEL_LOOP(i,n)
for(int32_ti=blockIdx.x*blockDim.x+threadIdx.x,step=blockDim.x*gridDim.x;i

//UpsampleNearest2DKerneliscopyedfromhttps://github.com/Oneflow-Inc/oneflow/blob/master/oneflow/user/kernels/upsample_nearest_kernel.cu#L78
template<typenameT>
structalignas(2*sizeof(T))Pack2X{
Tx;
Ty;
};

template<typenameT>
__global__voidUpsampleNearest2D2XForward(constint32_tin_elem_cnt,constT*in_dptr,
constint32_tin_height,constint32_tin_width,
T*out_dptr){
constint32_tin_hw_size=in_width*in_height;
CUDA_1D_KERNEL_LOOP(index,in_elem_cnt){
constTin_value=in_dptr[index];
constint32_tnc_idx=index/in_hw_size;
constint32_thw_off=index-nc_idx*in_hw_size;//這里是優化掉昂貴的取模運算
constint32_th=hw_off/in_width;
constint32_tw=hw_off-h*in_width;
Pack2Xout_value{in_value,in_value};
Pack2X*out_pack_dptr=reinterpret_cast*>(out_dptr);
out_pack_dptr[nc_idx*in_hw_size*2+h*2*in_width+w]=out_value;
out_pack_dptr[nc_idx*in_hw_size*2+(h*2+1)*in_width+w]=out_value;
}
}

template<typenameT>
__global__voidUpsampleNearest2D2XBackward(constint32_tin_elem_cnt,constT*dy_dptr,
constint32_tdx_height,constint32_tdx_width,
T*dx_dptr){
constint32_tdx_hw_size=dx_height*dx_width;
CUDA_1D_KERNEL_LOOP(index,in_elem_cnt){
Tdx_value=0.0;
constint32_tnc_idx=index/dx_hw_size;
constint32_tdx_hw_off=index-nc_idx*dx_hw_size;
constint32_tdx_h=dx_hw_off/dx_width;
constint32_tdx_w=dx_hw_off-dx_h*dx_width;
constPack2X*dy_pack_dptr=reinterpret_cast<constPack2X*>(dy_dptr);
constPack2Xdy_pack_value1=
dy_pack_dptr[nc_idx*dx_hw_size*2+dx_h*2*dx_width+dx_w];
constPack2Xdy_pack_value2=
dy_pack_dptr[nc_idx*dx_hw_size*2+(dx_h*2+1)*dx_width+dx_w];
dx_value+=dy_pack_value1.x;
dx_value+=dy_pack_value1.y;
dx_value+=dy_pack_value2.x;
dx_value+=dy_pack_value2.y;
dx_dptr[index]=dx_value;
}
}

這個地方比較好理解，我們以前向的UpsampleNearest2D2XForward為例，當我們對一個的矩陣進行2倍上采樣時，可以獲得大小的輸出Tensor，那么輸入和輸出的對應關系如下圖所示：

也就是輸入的(0, 0)位置對應來輸出的(0, 0), (0, 1), (1, 0), (1, 1)的位置。也就是一個輸入的元素其實是對應來輸出的4個元素，并且這4個元素一定是相鄰的2行或2列。所以我們可以使用Pack技術只用2次賦值就完成輸出Tensor對應位置元素的填寫，進一步提升全局內存訪問的帶寬。

我這里直接使用 oneflow 的腳本對這兩個 kernel 進行進行 profile ：

importoneflowasflow

x=flow.randn(16,32,80,80,device="cuda",dtype=flow.float32).requires_grad_()

m=flow.nn.Upsample(scale_factor=2.0,mode="nearest")

y=m(x)
print(y.device)
y.sum().backward()

下面展示了在 A100 上調優前后的帶寬占用和計算時間比較：

可以看到基于 oneflow upsample_nearest2d 的前后向的優化 kernel 可以獲得更好的帶寬利用率和性能。

模板代碼和用法詳見：https://github.com/BBuf/how-to-optim-algorithm-in-cuda/blob/master/UpsampleNearest2D/upsample_nearest_2d.cu

0x4. 總結

本篇文章推薦和講解一下OneFlow ElementWise模板，FastAtomicAdd，OneFlow UpsampleNearest2d模板的用法以及原理，并將其整理為最小的可以白嫖的頭文件。相關代碼請訪問 https://github.com/BBuf/how-to-optim-algorithm-in-cuda 這里獲得。