0

我在这里粘贴一些代码供大家查看。

__global__ void Integrate(double a, double b) {
    __shared__ double extrapol[16];
    __shared__ double result[32];
    __shared__ double h;
    __shared__ double err;

    __shared__ double x;
    __shared__ int n;

    if (threadIdx.x == 0) {
        h       = b - a;
        err     = 1.0;
        
        if (0.0 == a)
            extrapol[0] = 0.5 * h * myfunc(b);
        else
            extrapol[0] = 0.5 * h * (myfunc(a) + myfunc(b));

        n = 1;
    }

    for (int i = 1; i < 16; i++) {
        if (threadIdx.x == 0)
            x = a + h * 0.5;

        __syncthreads();
    
        if (err <= EPSILON)
            break;

        Trapezoid(result, x, h, n);
        if (threadIdx.x == 0) {
            result[0] = (extrapol[0] + h * result[0]) * 0.5;

            double power = 1.0;
            for (int k = 0; k < i; k++) {
               power *= 4.0;
               double sum  = (power * result[0] - extrapol[k]) / (power - 1.0);
               extrapol[k] = result[0];
               result[0] = sum;
            }

            err = fabs(result[0] - extrapol[i - 1]);
            extrapol[i] = result[0];
            n *= 2;
            h *= 0.5;
         }
    }
}

本质上它是一个自适应数字积分器（Romberg）。此全局函数中使用的设备函数有：

__device__ void Trapezoid(double *sdata, double x, double h, int n) {
    int nIdx = threadIdx.x + blockIdx.x * blockDim.x;
    sdata[nIdx] = 0;

    while (nIdx < n) {
       sdata[threadIdx.x] += myfunc(x + (nIdx * h));
       nIdx += 32;
    }
    Sum(sdata, threadIdx.x);
}

平行归约功能：

 __device__ void Sum(volatile double *sdata, int tId) {
     if (tId < 16) {
         sdata[tId] += sdata[tId + 16];
         sdata[tId] += sdata[tId + 8];
         sdata[tId] += sdata[tId + 4];
         sdata[tId] += sdata[tId + 2];
         sdata[tId] += sdata[tId + 1];
     }
}

最后，我试图集成的功能是（模拟简单功能）如下：

__device__ double myfunc(double x) {
     return 1 / x;
}

代码执行良好，得到了预期的积分。内核按以下方式执行（暂时）

Integrate <<< 1, 32 >>>(1, 2);

问题：
当我使用 nvidia 视觉分析器检查此功能的寄存器使用情况时。结果是每个线程有 52 个寄存器。我不明白为什么？我在这段代码中的大部分变量都是共享变量。你能告诉我如何找出我的代码的哪些部分正在使用寄存器吗？

我怎样才能减少它们？我可以用这段代码做任何优化吗？

硬件

我正在使用 fermi 设备 Geforce GTX 470，计算能力 2.0

谢谢，

4

2 回答 2

1

寄存器的使用与已定义变量的数量没有直接关系，因为例如，寄存器用于存储未定义变量的中间计算的结果。

尝试发现主要使用寄存器的代码部分的一种可能性是尝试通过使用类似的语法手动注释 ptx 文件来破解它

asm volatile ("// code at this line is doing this and this ...");
于 2013-09-05T17:09:02.210 回答
0

您可以使用该ptxas程序分析您的 ptx 文件，以向您显示每个函数的寄存器和内存使用情况。在你的情况下，你想做ptxas --gpu-name sm_20 -v code.ptx.

于 2013-09-05T17:25:26.617 回答