1

通常のメモリ アロケータで cublasIsamax を使用すると、正常に動作します。

int FindMaxIndex( const float* pVector, const size_t length )
{
    int result = 0;
    float* pDevVector = nullptr;

    if( CUBLAS_STATUS_SUCCESS != ::cudaMalloc( (void**)&pDevVector, length * sizeof(float) ) )
    {
        return -1;
    }
    if( CUBLAS_STATUS_SUCCESS !=  ::cudaMemcpy( pDevVector, pVector, length * (int)sizeof(float), cudaMemcpyHostToDevice) )
    {
        return -2;
    }
    ::cublasIsamax_v2( g_handle, length, pDevVector, 1, &result);

    if( nullptr != pDevVector )
    {
        ::cudaFree( pDevVector );
    }
    return result;
}

しかし、定数メモリで試してみると、不明なエラー N14 で失敗します。なにが問題ですか？定数メモリへのコピーは成功しましたが、実行に失敗しました。

__constant__ float c_pIndex[ 255 ] = {0x00};

// the same function as GetIsMax but using CUBLAS function cublasIsamax_v2
int FindMaxIndexConst( const float* pVector, const size_t length, pfnMsg fnMsg )
{
    int result = 0;
    cudaError_t code = ::cudaMemcpyToSymbol( c_pIndex, pVector, length * sizeof(float), 0, cudaMemcpyHostToDevice );

    if( cudaSuccess != code )
    {
        const char* szMsg = ::cudaGetErrorString ( code );

        LogError3( L"[%d] [%hs] Could not allocate CUDA memory: %I64d pDevA", code, szMsg, (__int64)(length * sizeof(float)));
    }
    cublasStatus_t  status = ::cublasIsamax_v2( g_handle, length, c_pIndex, 1, &result);

    if( CUBLAS_STATUS_SUCCESS != status )
    {
        LogError2( L" [%d] Failed to execute <cublasIsamax_v2> : %I64d", status, (__int64)length );
    }

    return result;
}
4

1 に答える 1

1

通常のデバイス配列を割り当てて、それを CUBLAS に渡してみませんか?

__constant__配列は通常の配列ではありません__device__。コードでは、配列のアドレスを取得し、それをホスト関数に渡しています。CUDA プログラミング ガイドで説明されているように、ホスト上の配列のアドレスはデバイス上では無効であり、その逆も同様です。CUDA プログラミング ガイドを参照してください。

__device__、__shared__または変数のアドレスを取得することによって得られるアドレスは、__constant__デバイス コードでのみ使用できます。「デバイス メモリ」で説明されているように、cudaGetSymbolAddress() によって取得された__device__または変数のアドレスは、ホスト コードでのみ使用できます。__constant__

__constant__デバイス ポインターを介してメモリにアクセスする場合は、この回答を参照して、キャッシュされない理由を確認してください。

最後に、__constant__メモリが定数キャッシュにキャッシュされている場合でも、この方法でメモリを使用することは、アクセス パターンのために非効率的です。定数キャッシュは、 warp 内のスレッド間で均一にアクセスできるように最適化されています。isamaxスレッドごとに異なるメモリ位置にアクセスする可能性が高いため、アクセスはシリアル化されます。したがって、これは一様にアクセスするよりも 32 倍遅くなります (通常のデバイス メモリよりもはるかに遅くなる可能性があります)。

于 2013-06-14T03:25:34.910 に答える