1

一連のベクトルの大きさを取得するために、次のカーネルがあります。

__global__ void norm_v1(double *in, double *out, int n)
{
    const uint i = blockIdx.x * blockDim.x + threadIdx.x;

    if (i < n)
    {
        double x = in[3*i], y = in[3*i+1], z = in[3*i+2];
        out[i] = sqrt(x*x + y*y + z*z);
    }
}

ただし、パッキングが原因inで[x0,y0,z0,...,xn,yn,zn]パフォーマンスが低下し、プロファイラーは 32% のグローバル負荷効率を示しています。としてデータを再パックする[x0, x1, ..., xn, y0, y1, ..., yn, z0, z1, ..., zn]と、状況が大幅に改善されます ( x、y、およびzそれに応じてオフセットが変更されます)。実行時間が短縮され、効率は最大 100% になります。

ただし、このパッキングは私のアプリケーションには実用的ではありません。したがって、共有メモリの使用を調査したいと思います。私の考えは、ブロック内の各スレッドがblockDim.xグローバル メモリから 3 つの値を (別々に) コピーすることであり、結合されたアクセスが得られます。blockDim.x = 256私が思いついた最大値の仮定の下で：

#define BLOCKDIM 256

__global__ void norm_v2(double *in, double *out, int n)
{
    __shared__ double invec[3*BLOCKDIM];

    const uint i = blockIdx.x * blockDim.x + threadIdx.x;

    invec[0*BLOCKDIM + threadIdx.x] = in[0*BLOCKDIM+i];
    invec[1*BLOCKDIM + threadIdx.x] = in[1*BLOCKDIM+i];
    invec[2*BLOCKDIM + threadIdx.x] = in[2*BLOCKDIM+i];
    __syncthreads();

    if (i < n)
    {
        double x = invec[3*threadIdx.x];
        double y = invec[3*threadIdx.x+1];
        double z = invec[3*threadIdx.x+2];

        out[i] = sqrt(x*x + y*y + z*z);
    }
}

n % blockDim.x != 0ただし、これは が事前に最大値を知る必要があり、でテストした場合にblockDim誤った結果を生成する場合、明らかに不十分です。これをどのように修正するのが最善ですか？out[i > 255]n = 1024

4

1 に答える 1

1

これで問題を解決できると思いますout[i > 255]：

__shared__ double shIn[3*BLOCKDIM];

const uint blockStart = blockIdx.x * blockDim.x;

invec[0*blockDim.x+threadIdx.x] = in[ blockStart*3 + 0*blockDim.x + threadIdx.x];
invec[1*blockDim.x+threadIdx.x] = in[ blockStart*3 + 1*blockDim.x + threadIdx.x];
invec[2*blockDim.x+threadIdx.x] = in[ blockStart*3 + 2*blockDim.x + threadIdx.x];
__syncthreads();

double x = shIn[3*threadIdx.x];
double y = shIn[3*threadIdx.x+1];
double z = shIn[3*threadIdx.x+2];

out[blockStart+threadIdx.x] = sqrt(x*x + y*y + z*z);

n % blockDim.x != 0要件に合わせて、入力/出力配列を 0 でパディングすることをお勧めします。

BLOCKDIMマクロが気に入らない場合は、extern __shared__ shArr[]3番目のパラメーターを使用してカーネル構成に渡すことを検討してください。

norm_v2<<<gridSize,blockSize,dynShMem>>>(...)

これdynShMemは、動的共有メモリの使用量 (バイト単位) です。これは、実行時にサイズが指定された追加の共有メモリ プールで、すべてのextern __shared__変数が最初に割り当てられます。

どのGPUを使用していますか? Fermi または Keplerは、L1 キャッシュを使用して元のコードを支援する場合があります。

配列をパディングしたくない場合in、または別の場所で同様のトリックを行うことになる場合はmemcopy、次のようなdevice-side の実装を検討することをお勧めします。

template <typename T>
void memCopy(T* destination, T* source, size_t numElements) {
    //assuming sizeof(T) is a multiple of sizeof(int)
    //assuming one-dimentional kernel (only threadIdx.x and blockDim.x matters) 
    size_t totalSize = numElements*sizeof(T)/sizeof(int);
    int* intDest = (int*)destination;
    int* intSrc = (int*)source;
    for (size_t i = threadIdx.x; i < totalSize; i += blockDim.x) {
        intDest[i] = intSrc[i];
    }
    __syncthreads();
}

基本的に、配列を -s の配列として扱いint、データをある場所から別の場所にコピーします。int基になる型をdouble-sに置き換えるかlong long int、64 ビット型のみを使用する場合に使用できます。

次に、コピー行を次のように置き換えることができます。

memCopy(invec, in+blockStart*3, min(blockDim.x, n-blockStart));
于 2012-11-06T04:19:56.720 に答える