合理配置线程数避免寄存器溢出【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit【优先级】中【描述】在SIMT编程的程序中应合理配置线程数避免寄存器溢出。SIMT编程模式下核函数通过__launch_bounds__指定的最大线程数决定每个线程可用的寄存器数量。最大线程数越大每个线程可分配的寄存器越少二者的对应关系如下最大线程数每个线程可用寄存器个数1025~204816513~102432257~512641~256127__launch_bounds__(N)是核函数上的可选限定符在核函数定义时配置用于在编译期向编译器声明执行该核函数的最大线程数为N编译器据此确定每个线程可分配的寄存器数量。当核函数未配置__launch_bounds__时最大线程数默认为1024此时每个线程可使用32个寄存器。对于计算密集型算子单个线程占用的寄存器通常较多在默认配置下容易超出寄存器上限超出部分的中间数据会溢出到栈空间位于Global Memory引入额外的Global Memory访问导致性能下降。避免寄存器溢出的思路是先通过--cce-res-usage编译选项查看核函数的寄存器使用情况Stack size大于0即表明存在溢出再根据上表中寄存器与最大线程数的对应关系选择一档能满足算子单线程寄存器需求的最大线程数并通过__launch_bounds__(N)将其配置给编译器。编译器据此放宽每线程的寄存器配额从而避免寄存器溢出将中间数据保留在寄存器中减少Global Memory访问提升性能。一般建议计算密集型算子如sincos配置512或1024线程数据搬运类算子配置2048线程。关于__launch_bounds__的详细说明请参考SIMT BuiltIn关键字。【样例介绍】以SinCosCompute算子为例使用sincosf接口同时计算sin和cos结果计算公式为 $output_sin[i] sin(input[i])$、$output_cos[i] cos(input[i])$。输入数据为float类型数据规模为393216个元素配置48个线程块、每个线程块512个线程每个线程循环计算16个输入值。基线版本与优化版本的核函数计算逻辑完全相同仅在是否配置__launch_bounds__上存在差异。【反例】不配置__launch_bounds__最大线程数取默认值1024编译器据此分配寄存器导致寄存器溢出。__global__ void sincos_thread_1024(float* input, float* output_sin, float* output_cos, uint64_t total_length) { int32_t blk_start_idx blockIdx.x * THREADS_PER_BLOCK * PER_THREAD_LOOP; // 每个核计算 PER_THREAD_LOOP * THREADS_PER_BLOCK 的运算量 for (int i 0; i PER_THREAD_LOOP; i) { int idx blk_start_idx i * THREADS_PER_BLOCK threadIdx.x; sincosf(input[idx], output_sin idx, output_cos idx); } }在上述实现中由于未配置__launch_bounds__最大线程数取默认值1024编译器据此为每个线程仅分配32个寄存器而sincosf计算所需的寄存器超出该上限超出部分溢出到栈空间。使用--cce-res-usage编译选项可查看寄存器使用情况[BISHENG] Function properties for _Z18sincos_thread_1024PfS_S_m_simt_entry: Stack size: 32 bytes, Used register number: 32其中Stack size: 32 bytes表明存在寄存器溢出栈位于Global MemoryUsed register number: 32已达到1024线程下的寄存器上限。在Ascend 950PR产品上该实现的性能数据如下Task Duration(us)DCache Read GMDCache Read VectorDCache Write Vector102.47256640768寄存器溢出导致中间数据反复访问栈空间Global Memory体现为较高的DCache Read Vector640次和DCache Write Vector768次。【正例】配置__launch_bounds__(512)提示编译器真实的最大线程数充分利用寄存器避免溢出。__global__ __launch_bounds__(512) void sincos_thread_512(float* input, float* output_sin, float* output_cos, uint64_t total_length) { int32_t blk_start_idx blockIdx.x * THREADS_PER_BLOCK * PER_THREAD_LOOP; // 每个核计算 PER_THREAD_LOOP * THREADS_PER_BLOCK 的运算量 for (int i 0; i PER_THREAD_LOOP; i) { int idx blk_start_idx i * THREADS_PER_BLOCK threadIdx.x; sincosf(input[idx], output_sin idx, output_cos idx); } }配置__launch_bounds__(512)后编译器为每个线程分配64个寄存器sincosf计算所需的寄存器在限制范围内无溢出。编译信息如下[BISHENG] Function properties for _Z17sincos_thread_512PfS_S_m_simt_entry: Stack size: 0 bytes, Used register number: 48其中Stack size: 0 bytes表明无寄存器溢出Used register number: 48在64个寄存器限制内所有中间数据保存在寄存器中。在Ascend 950PR产品上使用__launch_bounds__(512)后的性能数据如下Task Duration(us)DCache Read GMDCache Read VectorDCache Write Vector96.22256512256优化效果分析端到端耗时从102.47us降低到96.22us下降约6.1%。DCache Read GM保持256次不变说明优化未引入额外的Global Memory读取开销。DCache Read/Write VectorRead Vector从640降至512Write Vector从768降至256。栈空间物理位于Global Memory寄存器溢出时对栈的访问会体现在Data Cache的读写次数上消除溢出后这两项访问次数明显减少。【总结】对于计算密集型算子应先通过--cce-res-usage编译选项查看寄存器使用情况再根据寄存器与最大线程数的对应关系选择一档能满足单线程寄存器需求的最大线程数并通过__launch_bounds__配置给编译器使其放宽寄存器配额、避免寄存器溢出到Global Memory最后对比优化前后的Stack size、Task Duration与DCache访问次数验证优化效果。【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
CANN/asc-devkit：合理配置线程数避免寄存器溢出
合理配置线程数避免寄存器溢出【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit【优先级】中【描述】在SIMT编程的程序中应合理配置线程数避免寄存器溢出。SIMT编程模式下核函数通过__launch_bounds__指定的最大线程数决定每个线程可用的寄存器数量。最大线程数越大每个线程可分配的寄存器越少二者的对应关系如下最大线程数每个线程可用寄存器个数1025~204816513~102432257~512641~256127__launch_bounds__(N)是核函数上的可选限定符在核函数定义时配置用于在编译期向编译器声明执行该核函数的最大线程数为N编译器据此确定每个线程可分配的寄存器数量。当核函数未配置__launch_bounds__时最大线程数默认为1024此时每个线程可使用32个寄存器。对于计算密集型算子单个线程占用的寄存器通常较多在默认配置下容易超出寄存器上限超出部分的中间数据会溢出到栈空间位于Global Memory引入额外的Global Memory访问导致性能下降。避免寄存器溢出的思路是先通过--cce-res-usage编译选项查看核函数的寄存器使用情况Stack size大于0即表明存在溢出再根据上表中寄存器与最大线程数的对应关系选择一档能满足算子单线程寄存器需求的最大线程数并通过__launch_bounds__(N)将其配置给编译器。编译器据此放宽每线程的寄存器配额从而避免寄存器溢出将中间数据保留在寄存器中减少Global Memory访问提升性能。一般建议计算密集型算子如sincos配置512或1024线程数据搬运类算子配置2048线程。关于__launch_bounds__的详细说明请参考SIMT BuiltIn关键字。【样例介绍】以SinCosCompute算子为例使用sincosf接口同时计算sin和cos结果计算公式为 $output_sin[i] sin(input[i])$、$output_cos[i] cos(input[i])$。输入数据为float类型数据规模为393216个元素配置48个线程块、每个线程块512个线程每个线程循环计算16个输入值。基线版本与优化版本的核函数计算逻辑完全相同仅在是否配置__launch_bounds__上存在差异。【反例】不配置__launch_bounds__最大线程数取默认值1024编译器据此分配寄存器导致寄存器溢出。__global__ void sincos_thread_1024(float* input, float* output_sin, float* output_cos, uint64_t total_length) { int32_t blk_start_idx blockIdx.x * THREADS_PER_BLOCK * PER_THREAD_LOOP; // 每个核计算 PER_THREAD_LOOP * THREADS_PER_BLOCK 的运算量 for (int i 0; i PER_THREAD_LOOP; i) { int idx blk_start_idx i * THREADS_PER_BLOCK threadIdx.x; sincosf(input[idx], output_sin idx, output_cos idx); } }在上述实现中由于未配置__launch_bounds__最大线程数取默认值1024编译器据此为每个线程仅分配32个寄存器而sincosf计算所需的寄存器超出该上限超出部分溢出到栈空间。使用--cce-res-usage编译选项可查看寄存器使用情况[BISHENG] Function properties for _Z18sincos_thread_1024PfS_S_m_simt_entry: Stack size: 32 bytes, Used register number: 32其中Stack size: 32 bytes表明存在寄存器溢出栈位于Global MemoryUsed register number: 32已达到1024线程下的寄存器上限。在Ascend 950PR产品上该实现的性能数据如下Task Duration(us)DCache Read GMDCache Read VectorDCache Write Vector102.47256640768寄存器溢出导致中间数据反复访问栈空间Global Memory体现为较高的DCache Read Vector640次和DCache Write Vector768次。【正例】配置__launch_bounds__(512)提示编译器真实的最大线程数充分利用寄存器避免溢出。__global__ __launch_bounds__(512) void sincos_thread_512(float* input, float* output_sin, float* output_cos, uint64_t total_length) { int32_t blk_start_idx blockIdx.x * THREADS_PER_BLOCK * PER_THREAD_LOOP; // 每个核计算 PER_THREAD_LOOP * THREADS_PER_BLOCK 的运算量 for (int i 0; i PER_THREAD_LOOP; i) { int idx blk_start_idx i * THREADS_PER_BLOCK threadIdx.x; sincosf(input[idx], output_sin idx, output_cos idx); } }配置__launch_bounds__(512)后编译器为每个线程分配64个寄存器sincosf计算所需的寄存器在限制范围内无溢出。编译信息如下[BISHENG] Function properties for _Z17sincos_thread_512PfS_S_m_simt_entry: Stack size: 0 bytes, Used register number: 48其中Stack size: 0 bytes表明无寄存器溢出Used register number: 48在64个寄存器限制内所有中间数据保存在寄存器中。在Ascend 950PR产品上使用__launch_bounds__(512)后的性能数据如下Task Duration(us)DCache Read GMDCache Read VectorDCache Write Vector96.22256512256优化效果分析端到端耗时从102.47us降低到96.22us下降约6.1%。DCache Read GM保持256次不变说明优化未引入额外的Global Memory读取开销。DCache Read/Write VectorRead Vector从640降至512Write Vector从768降至256。栈空间物理位于Global Memory寄存器溢出时对栈的访问会体现在Data Cache的读写次数上消除溢出后这两项访问次数明显减少。【总结】对于计算密集型算子应先通过--cce-res-usage编译选项查看寄存器使用情况再根据寄存器与最大线程数的对应关系选择一档能满足单线程寄存器需求的最大线程数并通过__launch_bounds__配置给编译器使其放宽寄存器配额、避免寄存器溢出到Global Memory最后对比优化前后的Stack size、Task Duration与DCache访问次数验证优化效果。【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
