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

Question

我正在编写一个包含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;
}

Answer 1

在内核计算中从双精度乘法切换到双精度除法时寄存器使用的增加是由于双精度乘法是内置硬件指令,而双精度乘法是一个相当大的称为软件的事实.子程序(即,各种函数调用).通过检查生成的机器代码(SASS)可以很容易地验证这一点cuobjdump --dump-sass.

双精度除法(实际上所有除法,包括单精度除法和整数除法)由内联代码或被调用子程序模拟的原因是由于GPU硬件没有直接支持除法运算,为了使各个计算核心("CUDA核心")尽可能简单和尽可能小,这最终可以为给定大小的芯片带来更高的峰值性能.根据GFLOPS /瓦特度量标准衡量,它还可能提高核心的效率.

对于发布版本,引入双精度除法引起的寄存器使用的典型增加是大约26个寄存器.需要这些附加寄存器来在除法计算中存储中间变量,其中每个双精度临时变量需要两个32位寄存器.

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

一般来说,寄存器使用可以通过-maxrregcount nvcc编译器标志以编译单元粒度来控制,或者使用__launch_bounds__ 函数属性以每个函数粒度来控制.但是,强制较低的寄存器使用低于编译器确定的低于几个寄存器的频率会导致寄存器溢出生成的代码,这通常会对内核性能产生负面影响.