0

Nsight Visual Profile によって生成されたタイムラインは非常に奇妙に見えます。転送オーバーラップ コードは記述しませんが、MemCpyとComputeカーネルの間でオーバーラップが見られます。

これにより、実際の重複コードをデバッグできなくなります。

CUDA 5.0、Tesla M2090、Centos 6.3、2x CPU Xeon E5-2609 を使用しています

誰にも同様の問題がありますか？特定の Linux ディストリビューションでのみ発生しますか? 修正方法は？

これがコードです。

#include <cuda.h>
#include <curand.h>
#include <cublas_v2.h>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/device_ptr.h>

int main()
{
    cublasHandle_t hd;
    curandGenerator_t rng;
    cublasCreate(&hd);
    curandCreateGenerator(&rng, CURAND_RNG_PSEUDO_MTGP32);

    const size_t m = 5000, n = 1000;
    const double alpha = 1.0;
    const double beta = 0.0;

    thrust::host_vector<double> h(n * m, 0.1);
    thrust::device_vector<double> a(m * n, 0.1);
    thrust::device_vector<double> b(n * m, 0.1);
    thrust::device_vector<double> c(m * m, 0.1);
    cudaDeviceSynchronize();

    for (int i = 0; i < 10; i++)
    {
        curandGenerateUniformDouble(rng,
                thrust::raw_pointer_cast(&a[0]), a.size());
        cudaDeviceSynchronize();

        thrust::copy(h.begin(), h.end(), b.begin());
        cudaDeviceSynchronize();

        cublasDgemm(hd, CUBLAS_OP_N, CUBLAS_OP_N,
                m, m, n, &alpha,
                thrust::raw_pointer_cast(&a[0]), m,
                thrust::raw_pointer_cast(&b[0]), n,
                &beta,
                thrust::raw_pointer_cast(&c[0]), m);
        cudaDeviceSynchronize();
    }

    curandDestroyGenerator(rng);
    cublasDestroy(hd);

    return 0;
}

これはキャプチャされたプロファイル タイムラインです。

タイムライン

4

1 に答える 1

1

Compute Capability 2.* (Fermi) デバイスは、カーネル レベルの同時実行と、カーネルお​​よびコピーの同時実行の両方に対応しています。並行カーネルをトレースするために、カーネルの開始と終了のタイムスタンプは、メモリ コピーのタイムスタンプとは別のクロック ドメインで収集されます。このツールは、これらの異なるクロックを関連付けます。あなたのスクリーンショットでは、各メモリコピーが一定の値だけずれているのではなく、スケーリングされたオフセットだけずれていることがわかるので、スケーリングファクタリングが異なる (悪い相関関係) と思います。

nvprofのオプション--concurrent-kernels offを使えば問題はなくなると思います。並行カーネルが無効になっている場合、メモリ コピーとカーネル タイミングは、タイムスタンプに同じソース クロックを使用します。

Compute Capability 3.* (Kepler) と 5.* (Maxwell) では、計算カーネルのタイミングを計るメカニズムが異なります。これらのデバイスでは、カーネルの終了タイムスタンプとメモリ コピーまたはカーネルの開始との重複をツールで確認できます。作業は重複しません。オーバーラップの可能性 (通常は 500ns 未満) を持つか、依存する作業間の一定のオーバーヘッドとしてこれを導入するかの間で、ツールの設計上の決定がありました。ツールは、シリアル化された作業で非常に小さなレベルのオーバーラップを示す可能性があるという犠牲を払って、オーバーヘッドの導入を回避することを決定しました。

于 2016-02-09T22:05:10.093 に答える