cuda内核中除法运算对每个线程寄存器数的影响答案

【问题标题】：Influence of division operation in cuda kernel on number of registers per threadcuda内核中除法运算对每个线程寄存器数的影响
【发布时间】：2016-03-01 05:08:11
【问题描述】：

我正在编写一个包含 cuda 内核的程序。我发现如果你使用#define OPERATOR *，一个线程将使用11个寄存器，但我你将使用#define OPERATOR /（除法运算符），一个线程将使用52个寄存器！怎么了？我必须减少寄存器数量（我想设置maxregcount）！在 cuda 内核中使用 devision 运算符时如何减少寄存器数量？

#include <stdio.h>
#include <stdlib.h>
#define GRID_SIZE 1
#define BLOCK_SIZE 1
#define OPERATOR /
__global__ void kernel(double* array){
    for (int curEl=0;curEl<BLOCK_SIZE;++curEl){
    array[curEl]=array[curEl] OPERATOR 10;
    }
}
int main(void) {
    double *devPtr=NULL,*data=(double*)malloc(sizeof(double)*BLOCK_SIZE);
    cudaFuncAttributes cudaFuncAttr;
    cudaFuncGetAttributes(&cudaFuncAttr,kernel);
    for (int curElem=0;curElem<BLOCK_SIZE;++curElem){
        data[curElem]=curElem;
    }
    cudaMalloc(&devPtr,sizeof(double)*BLOCK_SIZE);
    cudaMemcpy(devPtr,data,sizeof(double)*BLOCK_SIZE,cudaMemcpyHostToDevice);
    kernel<<<1,BLOCK_SIZE>>>(devPtr);
    printf("1 thread needs %d regs\n",cudaFuncAttr.numRegs);
    return 0;
}

【问题讨论】：

如果你用cuobjdump -dump-sass检查为内核生成的机器代码，你会发现双精度乘法是一个内置的硬件指令，而双精度除法是一个相当大的调用子程序，所以寄存器使用的一些扩展是可以预料的。然而，增加的幅度比我在类似情况下观察到的要大。这是发布版本吗？您可以使用__launch_bounds__ 属性在每个功能的基础上限制寄存器的使用。强制使用较低的寄存器可能会导致溢出，这也会降低性能。
这是调试版本。在实际构建中，我有 33 个用于 div 的 regs 和 6 个用于 mul 的 regs。我英语不好。你说gpu div不是一个简单的指令，它是一个函数？
正确的双精度除法在内部实现为调用的子例程（~= 函数），因为 GPU 中没有对浮点除法的直接硬件支持。您在发布版本中观察到的寄存器使用差异与我过去观察到的一致。
虽然不得不假设这不是问题的实际背景，但我想提一下将/10 更改为*0.1 的选项。
@njuffa 你想提供答案吗？我会投票。

标签： cuda

【解决方案1】：

在内核计算中从双精度乘法切换到双精度除法时寄存器使用量的增加是由于双精度乘法是内置硬件指令，而双精度除法是相当大的称为软件子程序（即，某种函数调用）。这很容易通过使用cuobjdump --dump-sass 检查生成的机器代码 (SASS) 来验证。

双精度除法（实际上是所有除法，包括单精度除法和整数除法）都是通过内联代码或调用子程序来模拟的，这是因为 GPU 硬件不直接支持除法操作，以保持单个计算核心（“CUDA 核心”）尽可能简单和尽可能小，最终为给定尺寸的芯片带来更高的峰值性能。根据 GFLOPS/watt 指标衡量，它还可能提高内核的效率。

对于发布版本，由于引入双精度除法导致的寄存器使用量的典型增加约为 26 个寄存器。这些额外的寄存器用于存储除法计算中的中间变量，其中每个双精度临时变量需要两个 32 位寄存器。

正如 Marco13 在上面的评论中指出的那样，可以用倒数手动替换乘法除法。但是，这在大多数情况下会导致细微的数值差异，这就是 CUDA 编译器不会自动应用此转换的原因。

一般来说，可以通过-maxrregcountnvcc compiler flag 使用编译单元粒度来控制寄存器的使用，或者使用__launch_bounds__function attribute 使用每个函数粒度来控制寄存器的使用。但是，强制使用低于编译器确定级别的多个寄存器经常会导致生成代码中的寄存器溢出，这通常会对内核性能产生负面影响。

【讨论】：