CANN/asc-devkit SIMT内存对齐优化
通过类型对齐提升搬运效率【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit【优先级】中【描述】SIMT编程模式下结构体类型对齐会影响Global Memory的访存效率。 Ascend 950PR/Ascend 950DT支持1B、2B、4B、8B四种位宽的访存指令。 在数据总量不变时可通过类型对齐到更大的位宽提升数据搬运效率。【样例介绍】设置总搬运量MEM_BYTES为512MiB构造多组不同大小、不同对齐方式的结构体TData依次将长度为num_elements的输入数组d_idata整体复制到输出数组d_odata进行测试对比非对齐与使用alignas类型对齐后的耗时差异。具体代码可参考类型对齐优化搬运效率最佳实践样例其中核函数实现如下constexpr uint32_t THREAD_COUNT 2048; // threads per block constexpr size_t MEM_BYTES 512ULL * 1024ULL * 1024ULL; // 总数据量512 MiB uint32_t num_elements static_castuint32_t(MEM_BYTES / sizeof(TData)); template class TData __global__ __launch_bounds__(THREAD_COUNT) void aligned_type_kernel( TData* d_odata, TData* d_idata, uint32_t num_elements) { const uint32_t tid blockIdx.x * blockDim.x threadIdx.x; const uint32_t num_threads blockDim.x * gridDim.x; for (uint32_t pos tid; pos num_elements; pos num_threads) { d_odata[pos] d_idata[pos]; } }下面以Align4、Misalign4这两个结构体作为典型代表展示对照场景。【反例】Misalign4未设置显式对齐默认alignof受字段类型限制编译器只能选用窄位宽的访存指令。struct Misalign4 { unsigned char r, g, b, a; }; // sizeof4, alignof1性能数据如下Task Duration(us)aiv_total_cycles81223.02134028903.60上述结构体由字段类型决定alignof1编译器只能用8-bit档位搬运每个元素aiv_total_cycles和Task Duration都较高。【正例】使用alignas提升对齐属性编译器可以选用更宽位宽的访存指令。struct alignas(4) Align4 { unsigned char r, g, b, a; }; // sizeof4, alignof4性能数据如下Task Duration(us)aiv_total_cycles27889.3346005213.32Align4将alignof从1B提升到4B编译器可以选用32-bit档位完成单元素搬运aiv_total_cycles和Task Duration都显著减小。对比多组结构体大小不同的样例在类型对齐前后耗时差异如下结构体大小对齐Task Duration(us)非对齐Task Duration(us)非对齐Task Duration/对齐Task Duration4B (4×u8)27889.33 (alignof4)81223.02 (alignof1)2.91x4B (2×u16)27833.56 (alignof4)46656.61 (alignof2)1.68x4B (2×u8 u16)27854.56 (alignof4)46666.11 (alignof2)1.68x4B (u8 1B padding u16)27809.72 (alignof4)46664.42 (alignof2)1.68x8B17417.23 (alignof8)21287.91 (alignof4)1.22x16B32983.77 (alignof16)38353.60 (alignof4)1.16x32B60901.24 (alignof32)77790.93 (alignof4)1.28x测试结果表明在数据总量不变的时候相同大小的结构体类型对齐可以提升搬运效率。以4B结构体为例随着结构体类型对齐位宽从1B对齐提升至4B非对齐Task Duration/对齐Task Duration逐步下降说明非对齐与对齐的耗时差距逐步缩小。此外同大小的4B结构体字段组合形式、有无填充位都不会对搬运效率产生影响。【总结】SIMT算子使用结构体搬运数据时结构体类型对齐能让访存指令以更大位宽合并完成搬运降低搬运耗时。 Ascend 950PR/Ascend 950DT支持1B、2B、4B、8B位宽的访存指令因此建议使用时在不超过结构体大小的情况下尽量选择更高的位宽。【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
