1

CUDA で独自の 64 ビット シャッフル関数を実装しようとしています。ただし、次のようにすると:

static __inline__ __device__ double __shfl_xor(double var, int laneMask, int width=warpSize)
{
    int hi, lo;
    asm volatile( "mov.b64 { %0, %1 }, %2;" : "=r"(lo), "=r"(hi) : "d"(var) );
    hi = __shfl_xor( hi, laneMask, width );
    lo = __shfl_xor( lo, laneMask, width );
    return __hiloint2double( hi, lo );
}

以降の __shfl_xor へのすべての呼び出しは、引数の型に関係なく、この 64 ビット バージョンからインスタンス化されます。たとえば、私がやっている場合

int a;
a = __shfl_xor( a, 16 );

それはまだ二重バージョンを使用します。回避策として、別の関数名を使用している可能性があります。しかし、このシャッフル関数をテンプレート関数から呼び出しているため、別の名前を使用すると、64 ビット浮動小数点用に別のバージョンを作成する必要があり、これはあまり適切ではありません。

では、どうすれば __shfl_xor(double,...) 関数をオーバーロードしながら、同時に __shfl_xor(int,...) を適切に呼び出すことができるのでしょうか?

4

1 に答える 1

2

すべての整数型と float は double にアップキャストできます。組み込み関数と特殊な double 関数のどちらかを選択する場合、コンパイラはすべての型に対して独自の関数を選択する可能性があります。

別の名前の関数を作成し、それを使用して特殊な double バリアントと他の型のダミーの両方を作成しようとしましたか?

例えば：

static __inline__ __device__ double foo_shfl_xor(double var, int laneMask, int width=warpSize)
{
    // Your double shuffle implementation
}

static __inline__ __device__ int foo_shfl_xor(int var, int laneMask, int width=warpSize)
{
    // For every non-double data type you use
    // Just call the original shuffle function
    return __shfl_xor(var, laneMask, width);
}

// Your code that uses shuffle
double d;
int a;
foo_shfl_xor(d, ...); // Calls your custom shuffle
foo_shfl_xor(a, ...); // Calls default shuffle
于 2013-04-12T06:31:19.290 に答える