0

CUDA vector type が与えられたint4場合、定数メモリから 128 ビットのデータをロードするにはどうすればよいですか。

これはうまくいかないようです：

#include <stdio.h>
#include <cuda.h>

__constant__ int constant_mem[4];
__global__ void kernel(){
    int4 vec;
    vec = constant_mem[0];
}
int main(void){return 0;}

7 行目では、定数メモリ内の 4 つの整数値すべてを 128 ビットのベクトル型にロードしようとしています。この操作により、次のコンパイル エラーが発生します。

vectest.cu(7): error: no operator "=" matches these operands
            operand types are: int4 = int

また、次のように、キャストせずにベクター型に直接アクセスすることは可能ですか?

int data = vec[0];

PTX アセンブリの switch ステートメント:

    @%p1 bra    BB1_55;

    setp.eq.s32     %p26, %r1, 1;
    @%p26 bra   BB1_54;

    setp.eq.s32     %p27, %r1, 2;
    @%p27 bra   BB1_53;

    setp.ne.s32     %p28, %r1, 3;
    @%p28 bra   BB1_55;

    mov.u32     %r961, %r61;
    bra.uni     BB1_56;

BB1_53:
    mov.u32     %r961, %r60;
    bra.uni     BB1_56;

BB1_54:
    mov.u32     %r961, %r59;
    bra.uni     BB1_56;

BB1_55:
    mov.u32     %r961, %r58;

BB1_56:
4

1 に答える 1

1

最初のケースでは、キャストがおそらく最も簡単な解決策なので、次のようになります。

__constant__ int constant_mem[4];
__global__ void kernel(){
    int4 vec = * reinterpret_cast<int4 *>(&constant_mem);
}

(ブラウザで書かれた免責事項、コンパイルまたはテストされていない、自己責任で使用)

C++reinterpret_cast演算子を使用すると、コンパイラは強制的に 128 ビットのロード命令を発行します。

2 番目のケースでは、128 ビット メモリ トランザクションを使用して、128 ビット ベクトル型の配列に格納された 32 ビット ワードを直接アドレス指定したいようです。これには、おそらく次のようなヘルパー関数が必要です。

__inline__ __device__ int fetch4(const int4 val, const int n)
{
     (void) val.x; (void) val.y; (void) val.z; (void) val.w;
     switch(n) {
         case 3:
            return val.w;
         case 2: 
            return val.z;
         case 1:
            return val.y;
         case 0:
         default:
            return val.x;
    }
}

__device__ int index4(const int4 * array, const int n)
{
    int div = n / 4;
    int mod = n - (div * 4);

    int4 val = array[div]; // 128 bit load here

    return fetch4(val, mod);
}

__constant__ int constant_mem[128];
__global__ void kernel(){
    int val = index4(constant_mem, threadIdx.x);
}

(ブラウザで書かれた免責事項、コンパイルまたはテストされていない、自己責任で使用)

ここでは、int4値全体を読み取ってその内容を解析することにより、128 ビットのトランザクションを強制します (void へのキャストは、open64 コンパイラの古いバージョンに必要な呪文であり、メンバーが使用されていないと判断された場合にベクトルのロードを最適化する傾向がありました)。インデックス作成を行うためのオーバーヘッドの IOP がいくつかありますが、結果のトランザクションの負荷帯域幅がより高い場合は、それだけの価値がある可能性があります。switch ステートメントはおそらく条件付き実行を使用してコンパイルされているため、分岐分岐のペナルティはありません。int4 値の配列への非常にランダムなアクセスは、多くの帯域幅を浪費し、ワープのシリアル化を引き起こす可能性があることに注意してください。これを行うと、パフォーマンスに大きな悪影響が及ぶ可能性があります。

于 2012-07-24T15:46:30.077 に答える