cuda - 非同期メモリ転送を使用した CUDA CPU-GPU コールバック

Question

仲間の Cuda プログラマー、

ポーリングメカニズムを使用して cpu-gpu コールバックメカニズムを実装しようとしています。長さ 1 の 2 つの配列 (a と cpuflag、デバイス側の dev_a と gpuflag に対応) (基本的に 2 つの変数) があります。

最初の CPU は a をクリアし、gpuflag の更新を待ちます。GPU は、この a のクリアを確認し、gpuflag を更新します。CPU は非同期で gpuflag を cpuflag に転送し続け、フラグの更新を待ちます。CPU が更新を確認すると、再び a をリセットし、非同期で GPU に送信します。再び、GPU はこの a のクリアを確認して gpuflag を更新し、ピンポン プロセスが続行されます。このプロセスを100回続けてほしい。

コード全体はこちらです。nvcc -o output filename.cu と言うだけでコンパイルできます コードがピンポン動作を示さない理由を理解できません。どんな種類の助けも大歓迎です。前もって感謝します。

#include <stdio.h>

#define LEN 1
#define MAX 100

__global__ void myKernel(int len, int *dev_a, int *gpuflag) {
        int tid = threadIdx.x;
        gpuflag[tid] = 0;

        while(true){
        //Check if cpu has completed work
                if(dev_a[tid] == 0){
            //Do gpu work and increment flag
                        dev_a[tid] = 1;
                        gpuflag[tid]++;

            //Wait till cpu detects the flag increment and resets
                        while(true){
                                if(dev_a[tid] == 0){
                                        break;
                                }
                        }
                }
        //Max 100 ping pongs
        if(gpuflag[tid]==MAX){
            break;
        }
        }
}

int main( void ) {
        int index, *cpuflag, *gpuflag, value;

        int *a;
        int *dev_a;

        cudaStream_t stream0, stream1;

        cudaStreamCreate( &stream0 );
        cudaStreamCreate( &stream1 );

        cudaMalloc ( (void**)&gpuflag, LEN*sizeof(int) );
        cudaMemset ( gpuflag, 0, LEN*sizeof(int) );
        cudaHostAlloc( (void**)&cpuflag, LEN*sizeof(int), cudaHostAllocDefault );

        cudaMalloc ( (void**)&dev_a, LEN*sizeof(int) );
        cudaMemset ( dev_a, 0, LEN*sizeof(int) );
        cudaHostAlloc( (void**)&a, LEN*sizeof(int), cudaHostAllocDefault );

    //Reset everything
        for(int i=0; i<LEN; i++)
                a[i] = 0;

    //Auxillary variables
        index = 0;
    value = 1;

    //call kernel
        myKernel<<<1,1,0,stream0>>>(LEN, dev_a, gpuflag);

        while(true){
        //Asynchronously copy gpu flag
                cudaMemcpyAsync(cpuflag, gpuflag, LEN*sizeof(int), cudaMemcpyDeviceToHost, stream1);
        //Check if increment has happened or not
                if(cpuflag[index] == value){
            //if yes, reset 
                for(int i=0; i<LEN; i++)
                        a[i] = 0;
            //transfer asynchronously
                    cudaMemcpyAsync(dev_a, a, LEN*sizeof(int), cudaMemcpyHostToDevice, stream1);
            //increment pattern
            value++;
                        printf("GPU updated once. Value is a[%d] = %d, cpuflag = %d\n", index, a[index], cpuflag[index]);
                } else {
                        printf("------------GPU didn't updated. Value is a[%d] = %d, cpuflag = %d\n", index, a[index], cpuflag[index]);
        }

        //Max 100 ping-pongs
        if(value == MAX){
            break;
        }
        }

    cudaFreeHost(a);
    cudaFreeHost(cpuflag);

    cudaFree(dev_a);
    cudaFree(gpuflag);

    cudaStreamDestroy( stream0 );
    cudaStreamDestroy( stream1 );

        return 0;
}