1

次のコードがあります。

__global__ void interpolation(const double2* __restrict__ data, double2* __restrict__ result, const double* __restrict__ x, const double* __restrict__ y, const int N1, const int N2, int M)
{
    int i = threadIdx.x + blockDim.x * blockIdx.x;

    [...]        

    double phi_cap1, phi_cap2;

    if(i<M) {   

         for(int m=0; m<(2*K+1); m++) {

              [calculate phi_cap1];

              for(int n=0; n<(2*K+1); n++) {

                 [calculate phi_cap2];

                 [calculate phi_cap=phi_cap1*phi_cap2];

                 [use phi_cap];

             }
    }

}

}

Kepler K20 カードで動的プログラミングを使用して、一連のスレッドの処理を並行してディスパッチして、計算時間を短縮したいと考えていますphi_cap1。私のコードでは、スレッドの単一ブロックを起動しています。phi_cap2K=613x13

CUDA Dynamic Parallelism Programming Guide に従って、データを子カーネルと交換するために必要な要素の行列 ( と の積によって形成される) をphi_capグローバルメモリ169に割り当てます。確かに、ガイドを引用すると、phi_cap1phi_cap2

原則として、子カーネルに渡されるすべてのストレージは、グローバル メモリ ヒープから明示的に割り当てる必要があります。

私はその後、次のコードで終わった

__global__ void interpolation(const double2* __restrict__ data, double2* __restrict__ result, const double* __restrict__ x, const double* __restrict__ y, const int N1, const int N2, int M)
{
    int i = threadIdx.x + blockDim.x * blockIdx.x;

    [...]   

    dim3 dimBlock(2*K+1,2*K+1); dim3 dimGrid(1,1);

    if(i<M) {   

    double* phi_cap; cudaMalloc((void**)&phi_cap,sizeof(double)*(2*K+1)*(2*K+1));

    child_kernel<<<dimGrid,dimBlock>>>(cc_diff1,cc_diff2,phi_cap);

    for(int m=0; m<(2*K+1); m++) {

        for(int n=0; n<(2*K+1); n++) {

                        [use phi_cap];

        }
    }

}

}

問題は、最初のルーチンの実行に時間がかかることですが、2 番目のルーチンは、起動に5msコメントを付けても、ほぼすべての時間がAPI に費やされます。child_kernel23mscudaMalloc

動的プログラミングでは、子カーネルとデータを交換するためにメモリ空間を割り当てる必要があることが多く、唯一の解決策は非常に時間がかかるグローバルメモリであるように思われるため、動的プログラミングの有用性の深刻なボトルネックの1つはデータであると私には思えますグローバルメモリ割り当ての問題を回避する方法がない限り、交換してください。

問題は、上記の問題、つまり、カーネル内からグローバル メモリを割り当てるときに非常に時間がかかるという問題に対する回避策はありますか? . ありがとう

コメントで提案された解決策

親カーネルの外部から必要なグローバル メモリを割り当てます。親カーネルの外部から必要なグローバル メモリを割り当てる方がはるかに高速であることを確認しました。

4

1 に答える 1

4

i <Mの各スレッドからcudaMallocを呼び出しています。これは、McudaMalloc呼び出しを行っていることを意味します。

Mが大きいほど、悪化します。

代わりに、ブロックの最初のスレッドから、以前に使用したサイズのM倍を割り当てる単一のcudaMalloc呼び出しを行うことができます（実際には、より多くを割り当てる必要があるため、各ブロックは適切に配置されます）。その後、スレッドを同期すると、子カーネルごとに正しく計算されたphi_capアドレスを使用して子カーネルを開始できます。

または（特定の状況で、カーネル呼び出しの間に保持できる十分なメモリを割り当てることができる場合）、カーネルの外部で一度メモリを割り当てて再利用することもできます。それははるかに速いでしょう。Mがカーネル呼び出し間で異なる場合は、最大のMに必要なだけ割り当てることができます。

于 2013-02-15T12:31:39.743 に答える