2

私はカーネルをマルチスレッドの「ホスト」プログラムに実装しています。ここでは、すべてのホスト スレッドがカーネルを呼び出しています。定数メモリの使用に問題があります。定数メモリにはいくつかのパラメーターが配置されますが、スレッドごとに異なります。問題が発生するサンプルもビルドします。

これがカーネル

__global__ void Kernel( int *aiOutput, int Length )
{
    int id = threadIdx.x + blockIdx.x * blockDim.x;

    int iValue = 0;

    // bound check
    if( id < Length )
    {
        if( id % 3 == 0 )
            iValue = c_iaCoeff[2];
        else if( id % 2 == 0 )
            iValue = c_iaCoeff[1];
        else
            iValue = c_iaCoeff[0];

        aiOutput[id] = iValue;
    }
    __syncthreads();
}

そして、pthread がこの関数を呼び出しています。

void* WrapperCopy( void* params )
{
    // choose cuda device to perform on
    CUDA_CHECK_RETURN( cudaSetDevice( 0 ) );

    // cast of params
    SParams *_params = (SParams*)params;

    // copy coefficients to constant memory
    CUDA_CHECK_RETURN( cudaMemcpyToSymbol( c_iaCoeff, _params->h_piCoeff, 3*sizeof(int) ) );

    // loop kernel
    for( int i=0; i<100; i++ )
    {
        // perfrom kernel
        Kernel<<< BLOCKCOUNT, BLOCKSIZE >>>( _params->d_piArray, _params->iLength );
    }

    // copy data back from gpu
    CUDA_CHECK_RETURN( cudaMemcpy(
            _params->h_piArray, _params->d_piArray, BLOCKSIZE*BLOCKCOUNT*sizeof(int), cudaMemcpyDeviceToHost ) );

    return NULL;
}

定数メモリはこのように宣言されています。

__constant__ int c_iaCoeff[ 3 ];

ホストスレッドごとに異なる値がh_piCoeffあり、それを定数メモリにコピーします。

すべての pthread 呼び出しで同じ結果が得られるようになりました c_iaCoeff。これは、定数メモリがどのように機能するかの問題であり、コンテキストで宣言する必要があると思います。サンプルでは、c_iaCoeff​​呼び出しているすべての pthread に対して宣言されているのは 1 つだけで、pthreads によって呼び出されたカーネルは最後の値を取得しますcudaMemcpyToSymbol。そうですか？

ここで、定数メモリを 2 次元配列に変更しようとしました。2 番目の次元は以前と同じ値になりますが、最初の次元は使用される pthread のインデックスになります。

__constant__ int c_iaCoeff2[ THREADS ][ 3 ];

カーネルでは、このように使用されます。

iValue = c_iaCoeff2[iTId][2];

しかし、この方法で定数メモリを使用できるかどうかはわかりませんね。また、定数メモリにデータをコピーしようとするとエラーが発生しました。

CUDA_CHECK_RETURN( cudaMemcpyToSymbol( c_iaCoeff[_params->iTId], _params->h_piCoeff, 3*sizeof(int) ) );

一般的に、定数メモリを 2 次元配列として使用することは可能ですか? はいの場合、私の失敗はどこにありますか?

4

1 に答える 1

2

はい、必要な方法で定数メモリを使用できるはずですが、使用しているcudaMemcpyToSymbolコピー操作は正しくありません。呼び出しの最初の引数はsymbolで、API はランタイム シンボル テーブルを検索して、要求した定数メモリ シンボルのアドレスを取得します。そのため、アドレスを呼び出しに渡すことはできません (コードは実際には初期化されたホスト値を呼び出しに渡していますが、その理由は、演習として読者に任せます)。

見落としている可能性があるのは、呼び出しのオプションの 4 番目の引数です。これは、要求したシンボルが指すメモリへのオフセットです。したがって、次のようなことができるはずです。

cudaMemcpyToSymbol( c_iaCoeff,                    // symbol to lookup
                    _params->h_piCoeff,           // source location
                    3*sizeof(int),                // number of bytes to copy
                    (3*_params->iTId)*sizeof(int) // Offset in bytes
                   );

[標準の免責事項: ブラウザで書かれ、テストされていません。自己責任で使用してください]

最後の引数は、シンボルの先頭からのオフセット (バイト単位) です。2D 配列は行優先順で配置されるため、各コピー操作のオフセットとして行インデックスを掛けた行のピッチを使用する必要があります。

于 2013-01-31T08:11:14.887 に答える