1

0、1、および 2 を含む 2D 行列があります。スレッド数が行列サイズに等しく、各スレッドが行列の各要素で動作する cuda カーネルを作成しています。ここで、0 と 1 をそのまま保持できるが、2 を 1 に変換する数学演算が必要でした。1 ->1; 2 -> 1。上記の変換を行う数学演算子を使用する方法はありますか。どんな助けでも大歓迎です。ありがとうございました。

4

2 に答える 2

3

これはcudaの質問ではありません。

int A;
// set A to 0, 1, or 2
int a = (A + (A>>1)) & 1;
// a is now 0 if A is 0, or 1 if A is 1 or 2

またはマクロとして：

#define fix01(x) ((x+(x>>1))&1)

int a = fix01(A);

これもうまくいくようです：

#define fix01(x) ((x&&1)&1)

&&ブール AND 演算子 ( ) の使用が「数学演算」の定義に適合するかどうかはわかりません。

于 2013-07-14T22:36:29.150 に答える
1

質問は「数学」関数に関するものだったので、次の 2 次多項式を提案します。

int f(int x) { return ((3-x)*x)/2; }

ただし、速度を最大化するために分岐を避けたい場合: PTX ISA 1.0 以降、min 命令があります。(PTX ISA 3.1 マニュアルの表 36 を参照してください。) したがって、次の CUDA コードは

__global__ void test(int *x, int *y)
{
    *y = *x <= 1 ? *x : 1;
}

私のテストでは、次の PTX アセンブラーにコンパイルされます (arch オプションなしで CUDA 5 から nvcc を呼び出すだけです)。

    code for sm_10
            Function : _Z4testPiS_
    /*0000*/     /*0x1000c8010423c780*/     MOV R0, g [0x4];
    /*0008*/     /*0xd00e000580c00780*/     GLD.U32 R1, global14 [R0];
    /*0010*/     /*0x1000cc010423c780*/     MOV R0, g [0x6];
    /*0018*/     /*0x30800205ac400780*/     IMIN.S32 R1, R1, c [0x1] [0x0];
    /*0020*/     /*0xd00e0005a0c00781*/     GST.U32 global14 [R0], R1;

したがって、条件付き ?: を使用する min() 実装は、実際には分岐なしで単一の IMIN.S32 PTX 命令にコンパイルされます。したがって、実際のアプリケーションにはこれをお勧めします。

int f(int x) { return x <= 1 ? x : 1; }

しかし、非分岐操作のみを使用するという問題に戻ります。

C でこの結果を取得する別の形式は、2 つの not 演算子を使用することです。

int f(int x) { return !!x; }

または、単にゼロと比較します。

int f(int x) { return x != 0; }

(! および != の結果は、0 または 1 であることが保証されています。C99 標準、ISO/IEC 9899:1999 の 6.5.3.3 項 5 および 6.5.9 項 3 を比較してください。この保証も同様です) CUDA で保持されます。)

于 2013-07-16T15:02:09.027 に答える