使用异步内存传输的 CUDA CPU-GPU 回调

Posted 2023-02-16

【中文标题】使用异步内存传输的 CUDA CPU-GPU 回调

【英文标题】：CUDA CPU-GPU callbacks using asynchronous memory transfer

【发布时间】：2012-07-27 14:42:37

【问题描述】：

各位 Cuda 程序员，

我正在尝试使用轮询机制实现 cpu-gpu 回调机制。我有 2 个长度为 1 的数组（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;

【问题讨论】：

CUDA 内存模型不能保证这种内存一致性。如果您的 GPU 支持，您可能可以通过让内核执行系统范围的线程栅栏来进一步提高，但这实际上不是定义的行为，

我同意，这让你有未定义的行为。但是在 CUDA 5.0 中，您可以使用主机挂钩机制在特定内核完成执行后将主机调用排队。它在开发中是完全异步的。预览 5.0。

【参考方案1】：

可能缺少的主要内容是正确使用volatile。

这是一个简化的、完整的示例：

$ cat t763.cu
#include <stdio.h>

#define LEN 1
#define MAX 100
#define DLEN 1000
#define nTPB 256

#ifdef CDP_WORKER
__global__ void cdp_worker(int len, float *data)

  int tid = threadIdx.x+blockDim.x*blockIdx.x;
  if (tid < len) data[tid]++; // simple increment

#endif

// only call this kernel with 1 thread
__global__ void myKernel(int len, int dlen, volatile int *dev_a, int *gpuflag, float *data) 
        int tid = threadIdx.x+blockDim.x*blockIdx.x;

        while(gpuflag[tid] < MAX)
        //Check if cpu has completed work
                if(dev_a[tid] == 0)
            //Do gpu work and increment flag
#ifdef CDP_WORKER
                        cdp_worker<<<(dlen+nTPB-1)/nTPB, nTPB>>>(dlen, data);
                        cudaDeviceSynchronize();
#endif
                        dev_a[tid] = 1;
                        gpuflag[tid]++;

                                
        


void issue_work(int value, float *h_data, float *d_data, int len, cudaStream_t mystream)
#ifdef CDP_WORKER
  cudaMemcpyAsync(h_data, d_data, len*sizeof(float), cudaMemcpyDeviceToHost, mystream);
  cudaStreamSynchronize(mystream);
  for (int i = 0; i < len; i++) if (h_data[i] != value+1) printf("fault - was %f, should be %f\n", h_data[i], (float)(value+1)); break;
  cudaMemcpyAsync(d_data, h_data, len*sizeof(float), cudaMemcpyHostToDevice, mystream); // technically not really necessary
  cudaStreamSynchronize(mystream);
#endif
  return;

int main( void ) 
        int *gpuflag, value;
        float *h_data, *d_data;
        cudaHostAlloc(&h_data, DLEN*sizeof(float), cudaHostAllocDefault);
        cudaMalloc(&d_data, DLEN*sizeof(float));
        volatile int *z_a;

        cudaStream_t stream0, stream1;

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

        cudaMalloc ( (void**)&gpuflag, LEN*sizeof(int) );
        cudaMemset ( gpuflag, 0, LEN*sizeof(int) );
        cudaMemset ( d_data, 0, DLEN*sizeof(float));
        cudaHostAlloc( (void**)&z_a, LEN*sizeof(int), cudaHostAllocMapped );
        for (int i = 0; i < LEN; i++) z_a[i] =
        value = 0;
    //call kernel
        myKernel<<<1,1,0,stream0>>>(LEN, DLEN, z_a, gpuflag, d_data);

        while(value<MAX)
          if (z_a[0] == 1) 
             issue_work(value, h_data, d_data, DLEN, stream1);
             z_a[0] = 0;
             printf("%d", value%10);
             value++;
        
        printf("\n");
        return 0;

$ nvcc -o t763 t763.cu
$ cuda-memcheck ./t763
========= CUDA-MEMCHECK
0123456789012345678901234567890123456789012345678901234567890123456789012345678901234567890123456789
========= ERROR SUMMARY: 0 errors
$ nvcc -DCDP_WORKER -arch=sm_35 -rdc=true t763.cu -o t763 -lcudadevrt
$ cuda-memcheck ./t763
========= CUDA-MEMCHECK
0123456789012345678901234567890123456789012345678901234567890123456789012345678901234567890123456789
========= ERROR SUMMARY: 0 errors
$

将其扩展为在同一个 warp 中的多个线程上工作并非易事。

但是，我已经扩展了基本示例，以在 cc3.5+ 设备上演示父内核可以是监督内核，并且它可以通过子内核启动工作。这是通过使用CDP_WORKER 开关和 CUDA 动态并行所需的其他开关进行编译并在 cc3.5+ 设备上运行来实现的。