3

1024 要素の配列にランダム アクセスする必要がある 32 個のスレッドのブロックがあるとします。最初にブロックをグローバルから共有に転送することで、グローバル メモリ呼び出しの数を減らしたいと考えています。私はそれについて2つのアイデアがあります：

A:

my_kernel()
{
    CopyFromGlobalToShared(1024 / 32 elements);
    UseSharedMemory();
}

または B:

my_kernel()
{
    if (first thread in block)
    {
        CopyFromGlobalToShared(all elements);
    }
    UseSharedMemory();
}

どちらが良いですか？または、別のより良い方法はありますか？

4

2 に答える 2

1

A の方が優れていますが、要素のサイズによっては、必ずしも最高であるとは限りません。

必要なのは、サイズが 32 ビットの隣接ワードにアクセスするすべてのスレッドです (64 ビットは、新しいハードウェアでも動作する可能性があります)。要素のサイズが大きい場合は、もう少し派手にすることをお勧めします。

//assumes sizeof(T) is multiple of sizeof(int)
//assumes single-dimention kernel
//assumes it is launched for all threads in block
template <typename T>
__device__ void memCopy(T* dest, T* src, size_t size) {
    int* iDest = (int*)dest;
    int* iSrc = (int*)src;
    for(size_t i = threadIdx.x; i<size*sizeof(T)/sizeof(int); i+=blockDim.x)
        iDest[i] = iSrc[i];
    __syncthreads();   
}

注: これは、すべてのメモリ転送 (グローバル -> グローバル、グローバル -> 共有、共有 -> グローバル) で機能します。これは、関数がインライン化されるため、ユニファイド メモリ アドレッシングのない古いデバイスにも当てはまります。

より大きな要素に対して構造体の配列アプローチを使用する場合、問題は発生しません。

于 2015-02-17T19:40:05.050 に答える