3

CUDA でいくつかの信号処理を作成して、最近、CUDA の最適化を大幅に進めました。1D テクスチャを使用し、アクセス パターンを調整することで、パフォーマンスを 10 倍向上させることができました。（以前、グローバルから共有メモリへのトランザクション整列プリフェッチを試みましたが、後で発生する不均一なアクセスパターンにより、ワープ→共有キャッシュバンクの関連付けが台無しになりました（と思います）。

だから今、私はCUDAテクスチャとバインディングが非同期memcpyとどのように相互作用するかという問題に直面しています。

次のカーネルを検討してください

texture<...> mytexture;

__global__ void mykernel(float *pOut)
{
    pOut[threadIdx.x] = tex1Dfetch(texture, threadIdx.x);
}

カーネルは複数のストリームで起動されます

extern void *sourcedata;

#define N_CUDA_STREAMS ...

cudaStream stream[N_CUDA_STREAMS];
void *d_pOut[N_CUDA_STREAMS];
void *d_texData[N_CUDA_STREAMS];

for(int k_stream = 0; k_stream < N_CUDA_STREAMS; k_stream++) {
    cudaStreamCreate(stream[k_stream]);

    cudaMalloc(&d_pOut[k_stream], ...);
    cudaMalloc(&d_texData[k_stream], ...);
}

/* ... */

for(int i_datablock; i_datablock < n_datablocks; i_datablock++) {
    int const k_stream = i_datablock % N_CUDA_STREAMS;
    cudaMemcpyAsync(d_texData[k_stream], (char*)sourcedata + i_datablock * blocksize, ..., stream[k_stream]);

    cudaBindTexture(0, &mytexture, d_texData[k_stream], ...);

    mykernel<<<..., stream[k_stream]>>>(d_pOut);
}

ここで気になるのは、テクスチャ参照が 1 つしかないため、他のストリームのカーネルがそのテクスチャにアクセスしているときにバッファをテクスチャにバインドするとどうなるかということです。cudaBindStreamはストリーム パラメータをとらないので、実行中のカーネルがテクスチャに非同期的にアクセスしている間にテクスチャを別のデバイス ポインタにバインドすると、他のデータへのアクセスが迂回されるのではないかと心配しています。

CUDA のドキュメントには、これについて何も書かれていません。同時アクセスを許可するためにこれを解く必要がある場合は、カーネル起動パラメーターとして渡されたストリーム番号に基づいて、多数のテクスチャ参照を作成し、switch ステートメントを使用してそれらの間で選択する必要があるようです。

残念ながら、CUDA はテクスチャの配列をデバイス側に配置することを許可していません。つまり、以下は機能しません。

texture<...> texarray[N_CUDA_STREAMS];

私が持っているデータの量は、CUDA 配列にバインドされていないプレーンな 1D テクスチャ内にしか収まらないため、レイヤ化されたテクスチャはオプションではありません(CUDA 4.2 C プログラミング ガイドの表 F-2 を参照してください)。

4

1 に答える 1

5

実際、テクスチャを別のストリームで使用している間は、テクスチャのバインドを解除することはできません。

非同期 memcpys を非表示にするためにストリームの数を大きくする必要はないため (2 つあれば十分です)、C++ テンプレートを使用して各ストリームに独自のテクスチャを与えることができます。

texture<float, 1, cudaReadModeElementType> mytexture1;
texture<float, 1, cudaReadModeElementType> mytexture2;

template<int TexSel> __device__ float myTex1Dfetch(int x);

template<> __device__ float myTex1Dfetch<1>(int x) { return tex1Dfetch(mytexture1, x); }
template<> __device__ float myTex1Dfetch<2>(int x) { return tex1Dfetch(mytexture2, x); }


template<int TexSel> __global__ void mykernel(float *pOut)
{
    pOut[threadIdx.x] = myTex1Dfetch<TexSel>(threadIdx.x);
}


int main(void)
{
    float *out_d[2];

    // ...

    mykernel<1><<<blocks, threads, stream[0]>>>(out_d[0]);
    mykernel<2><<<blocks, threads, stream[1]>>>(out_d[1]);

    // ...
}
于 2012-09-13T18:42:36.090 に答える