optimization - CUDA コードの最適化。レジスタ数

Question

みんなが見られるように、ここにいくつかのコードを貼り付けています。

__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 を使用しています。

ありがとう、