1

グローバル メモリに int 配列があります。グローバル メモリからの読み取りの試行回数を減らすために、64 ビットのデータ型を使用して読み取りを試し、必要に応じて上位または下位の 32 ビットを使用しました。たとえば、これは配列から 3 番目と 4 番目の int を取得します。

__device__ void func1(int* arr)
{
    unsigned long long int val = *((unsigned long long int *) &arr[3]);

    // Now operate on the individual ints
}

このメソッドを使用してintを取得すると、動作するはずですが、未定義の動作が発生します。それが機能する場合、この方法での値の読み取りは、個々の整数の読み取りよりもかなり高速です。以前にこの問題に遭遇した人はいますか?

4

1 に答える 1

1

数量は、サイズごとに揃えるのが好きです。cudaがあなたがしていることをどのように処理するかはわかりませんが、それは環境固有である可能性がありますが、次のものを使用しています:

*((unsigned long long int *) &arr[3])

8 バイト アラインされていると仮定するとarr、4 バイト アラインされている 8 バイト量を取得しています。もちろん、これは次の理由で発生します。

arr = 8n           // n is an integer
sizeof(int) = 4

&arr[3] = 8n + 3*4 // simplifies to 8(n+1) + 4

32 ビットと 16 ビットの整数を使用してプロセッサで同じことをしようとすると、問題が発生することはわかっています (ただし、64 ビットと 32 ビットの整数で試したことはありません)。

アクセスしようとしているデータがどこにあるかを把握するために、ある種のアクセサー取引を自作する必要があります。あなたの状況に似た、次の状況を考えてみましょう。

int get32BitValueFrom(unsigned long long int longArray[], int index)
{
    // get the 64 bit int containing the 32 bit int we want
    unsigned long long int value = longarray[index >> 1];

    // if we wanted an odd index, return the high order 32 bits
    // otherwise return the low order 32 bits
    return (int) ((index & 1) ? (value >> 32) : (value));
}

編集：あなたがcudaを使用していることは知っていますし、分岐を避けることも知っていますが、同じことを達成する何らかのビット単位または数学的トリックを使用して同等のコードを書く方法があると確信しています。

于 2012-07-25T22:12:28.843 に答える