1

OpenCL カーネルによる内部レジスタの使用に関して、少し混乱しています。-cl-nv-verbose を使用して、カーネルのレジスタの使用状況をキャプチャしています。現在、私のカーネルはカーネルptxas info: Used 4 registers内のいくつかのコードを記録しています。次のセグメントの場合:

double a;
a = pow(2.0,2.0)
if (index != 0) {
}

使用するレジスタが変わりますptxas info: Used 6 registers。if ループには何もないことがわかりました。しかし、次のように再構成すると:

double a;
if (index != 0) {
    a = pow(2.0,2.0)
}

これにより、レジスタの使用法が に変更されますptxas info: Used 15 registers。カーネルのワーク グループ サイズは変更しません。おそらく答えはptxコードを見ることにあるのですが、私はそれを理解していません(必要に応じて取得できますが)。私がもっと興味を持っているのは、コード行を移動するだけで、レジスタの使用量が 2 倍になる理由です。何か案は?(インデックスはprivate)

更新:前後PTXコード

更新: カーネルコード:

__kernel void butterfC( __global double *sI,    
                        __global double *sJ,
                        __global double *sK,
                        const int zR, 
                        const int yR,   
                        const int xR,
                        unsigned int l1,
                        const int dir,
                        unsigned int type   ) 
{
    int idX = get_global_id(0);
    int idY = get_global_id(1);
    int idZ = get_global_id(2);

    int BASE = idZ*xR*yR;
    int STRIDE = 1;

    int powX = pow(4.0f,l1);
    int powXm1 = pow(4.0f,l1-1);

    int yIndex, kIndex;

    switch(type)
    {
        case 1: BASE += idY*xR; 
                yIndex  = idX / powXm1 * powX;
                kIndex  = (idX % powXm1) + yIndex;  
                break;
        case 2: BASE += idX; STRIDE = xR; 
                yIndex  = idY / powXm1 * powX;
                kIndex  = idY % powXm1 + yIndex;
                break;
        case 3: BASE = idY*xR + idX; STRIDE = xR * yR; 
                yIndex  = idZ / powXm1 * powX;
                kIndex  = idZ % powXm1 + yIndex;
                break; 
    }

    double a;   

    //a = pow(2.0,2.0);
    if (kIndex != 0) {
        a = pow(2.0,2.0);
        .... do stuff
    }
}
4

0 に答える 0