将CUDA中的缓冲加倍，这样CPU就可以对持久化内核产生的数据进行操作

Question

我有一个蒙特卡罗模拟，其中系统的状态是一个位串(大小为N)，其中位被随机翻转。为了加速模拟，代码被修改为使用CUDA。然而，由于我需要从系统状态(作为N^2)计算大量统计数据，这部分需要在有更多内存的CPU上完成。目前的算法是这样的:

loop
  CUDA kernel making 10s of Monte Carlo steps
  Copy system state back to CPU
  Calculate statistics

这是低效的，我希望内核持续运行，而CPU偶尔查询系统状态并计算统计数据，而内核继续运行。

根据Tom对这个问题的回答，我认为答案是双重缓冲，但是我还没有找到如何做到这一点的解释或示例。

如何为CUDA/c++代码设置在Tom回答的第三段中描述的双缓冲?

Answer 1

这是一个"持久"内核的完整示例，生产者-消费者方法，具有从设备(生产者)到主机(消费者)的双缓冲接口。

持久内核设计通常意味着启动的内核最多可以同时驻留在硬件上的块数量(参见幻灯片16中的第1项)。为了最有效地使用机器，我们通常希望最大化这一点，同时仍然保持在上述限制之内。这涉及到一个特定内核的占用研究，并且它会因内核而异。因此，我在这里选择了一个捷径，只要有多处理器就启动多少块。这种方法总是可以保证工作(它可以被认为是持久化内核启动的块数量的"下限")，但是(通常)不是最有效的机器使用方法。不过，我认为占用率研究与你的问题无关。此外，有争议的是，适当的"持久内核"设计与保证向前进展实际上是相当棘手的——需要仔细设计CUDA线程代码和放置线程块(例如，每个SM只使用一个线程块)来保证向前进展。然而，我们不需要深入到这个层次来解决你的问题(我不认为)，我在这里提出的持久内核示例每个SM只放置1个线程块。

我还假设了一个正确的UVA设置，这样我就可以跳过在非UVA设置中安排正确映射内存分配的细节。

基本思想是，我们将在设备上有2个缓冲区，以及映射内存中的2个"邮箱"，每个缓冲区一个。设备内核将用数据填充缓冲区，然后将"邮箱"设置为一个值(在本例中为2)，该值表示主机可能"消耗"缓冲区。然后，该设备继续到另一个缓冲区，并在缓冲区之间以乒乓方式重复该过程。为了使这个工作，我们必须确保设备本身没有超出缓冲区(任何线程都不允许比任何其他线程领先一个以上的缓冲区)和，在设备填充缓冲区之前，主机已经消耗了先前的内容。

在主机端，它只是等待邮箱指示"满"，然后将缓冲区从设备复制到主机，重置邮箱，并对其执行"处理"(validate函数)。然后它以乒乓的方式进入下一个缓冲区。设备实际"产生"的数据只是用迭代数填充每个缓冲区。然后，主机检查是否收到了正确的迭代数。

我已经构建了代码来调用实际的设备"工作"函数(my_compute_function)，这是您将放置蒙特卡洛代码的地方。如果您的代码很好地独立于线程，那么这应该很简单。因此，设备端my_compute_function是生产者函数，主机端validate是消费者函数。如果你的设备生成器代码不是简单的线程独立的，那么你可能需要围绕my_compute_function的调用点稍微重新构建一些东西。

这样做的最终效果是，设备可以"抢先"并开始填充下一个缓冲区，而主机正在"消耗"前一个缓冲区中的数据。

由于持久内核设计对内核启动中的块(和线程)数量施加了上限，因此我选择在网格跨行循环中实现"work"生产者函数，因此可以通过给定的网格宽度处理任意大小的缓冲区。

下面是一个完整的例子:

$ cat t942.cu
#include <stdio.h>
#define ITERS 1000
#define DSIZE 65536
#define nTPB 256
#define cudaCheckErrors(msg) 
    do { 
        cudaError_t __err = cudaGetLastError(); 
        if (__err != cudaSuccess) { 
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)n", 
                msg, cudaGetErrorString(__err), 
                __FILE__, __LINE__); 
            fprintf(stderr, "*** FAILED - ABORTINGn"); 
            exit(1); 
        } 
    } while (0)

__device__ volatile int blkcnt1 = 0;
__device__ volatile int blkcnt2 = 0;
__device__ volatile int itercnt = 0;
__device__ void my_compute_function(int *buf, int idx, int data){
  buf[idx] = data;  // put your work code here
}
__global__ void testkernel(int *buffer1, int *buffer2, volatile int *buffer1_ready, volatile int *buffer2_ready,  const int buffersize, const int iterations){
  // assumption of persistent block-limited kernel launch
  int idx = threadIdx.x+blockDim.x*blockIdx.x;
  int iter_count = 0;
  while (iter_count < iterations ){ // persistent until iterations complete
    int *buf = (iter_count & 1)? buffer2:buffer1; // ping pong between buffers
    volatile int *bufrdy = (iter_count & 1)?(buffer2_ready):(buffer1_ready);
    volatile int *blkcnt = (iter_count & 1)?(&blkcnt2):(&blkcnt1);
    int my_idx = idx;
    while (iter_count - itercnt > 1); // don't overrun buffers on device
    while (*bufrdy == 2);  // wait for buffer to be consumed
    while (my_idx < buffersize){ // perform the "work"
      my_compute_function(buf, my_idx, iter_count);
      my_idx += gridDim.x*blockDim.x; // grid-striding loop
      }
    __syncthreads(); // wait for my block to finish
    __threadfence(); // make sure global buffer writes are "visible"
    if (!threadIdx.x) atomicAdd((int *)blkcnt, 1); // mark my block done
    if (!idx){ // am I the master block/thread?
      while (*blkcnt < gridDim.x);  // wait for all blocks to finish
      *blkcnt = 0;
      *bufrdy = 2;  // indicate that buffer is ready
      __threadfence_system(); // push it out to mapped memory
      itercnt++;
      }
    iter_count++;
    }
}
int validate(const int *data, const int dsize, const int val){
  for (int i = 0; i < dsize; i++) if (data[i] != val) {printf("mismatch at %d, was: %d, should be: %dn", i, data[i], val); return 0;}
  return 1;
}
int main(){
  int *h_buf1, *d_buf1, *h_buf2, *d_buf2;
  volatile int *m_bufrdy1, *m_bufrdy2;
  // buffer and "mailbox" setup
  cudaHostAlloc(&h_buf1, DSIZE*sizeof(int), cudaHostAllocDefault);
  cudaHostAlloc(&h_buf2, DSIZE*sizeof(int), cudaHostAllocDefault);
  cudaHostAlloc(&m_bufrdy1, sizeof(int), cudaHostAllocMapped);
  cudaHostAlloc(&m_bufrdy2, sizeof(int), cudaHostAllocMapped);
  cudaCheckErrors("cudaHostAlloc fail");
  cudaMalloc(&d_buf1, DSIZE*sizeof(int));
  cudaMalloc(&d_buf2, DSIZE*sizeof(int));
  cudaCheckErrors("cudaMalloc fail");
  cudaStream_t streamk, streamc;
  cudaStreamCreate(&streamk);
  cudaStreamCreate(&streamc);
  cudaCheckErrors("cudaStreamCreate fail");
  *m_bufrdy1 = 0;
  *m_bufrdy2 = 0;
  cudaMemset(d_buf1, 0xFF, DSIZE*sizeof(int));
  cudaMemset(d_buf2, 0xFF, DSIZE*sizeof(int));
  cudaCheckErrors("cudaMemset fail");
  // inefficient crutch for choosing number of blocks
  int nblock = 0;
  cudaDeviceGetAttribute(&nblock, cudaDevAttrMultiProcessorCount, 0);
  cudaCheckErrors("get multiprocessor count fail");
  testkernel<<<nblock, nTPB, 0, streamk>>>(d_buf1, d_buf2, m_bufrdy1, m_bufrdy2, DSIZE, ITERS);
  cudaCheckErrors("kernel launch fail");
  volatile int *bufrdy;
  int *hbuf, *dbuf;
  for (int i = 0; i < ITERS; i++){
    if (i & 1){  // ping pong on the host side
      bufrdy = m_bufrdy2;
      hbuf = h_buf2;
      dbuf = d_buf2;}
    else {
      bufrdy = m_bufrdy1;
      hbuf = h_buf1;
      dbuf = d_buf1;}
    // int qq = 0; // add for failsafe - otherwise a machine failure can hang
    while ((*bufrdy)!= 2); // use this for a failsafe:  if (++qq > 1000000) {printf("bufrdy = %dn", *bufrdy); return 0;} // wait for buffer to be full;
    cudaMemcpyAsync(hbuf, dbuf, DSIZE*sizeof(int), cudaMemcpyDeviceToHost, streamc);
    cudaStreamSynchronize(streamc);
    cudaCheckErrors("cudaMemcpyAsync fail");
    *bufrdy = 0; // release buffer back to device
    if (!validate(hbuf, DSIZE, i)) {printf("validation failure at iter %dn", i); exit(1);}
    }
 printf("Completed %d iterations successfullyn", ITERS);
}

$ nvcc -o t942 t942.cu
$ ./t942
Completed 1000 iterations successfully
$

我已经测试了上面的代码，它似乎在linux上工作得很好。我相信它应该可以在windows TCC设置。然而，在windows WDDM上，我认为我仍在调查一些问题。

注意，上面的内核设计尝试使用块计数原子策略进行网格范围的同步。CUDA现在(9.0及更新版本)有合作组，这是推荐的方法，而不是上面的方法，来创建一个网格范围的同步。

Answer 2

这不是对你问题的直接回答，但可能会有所帮助。

我正在使用CUDA生产者-消费者代码，其基本结构似乎与您的相似。我希望通过使CPU和GPU并行运行来加快代码的速度。我试图通过重组代码来实现这一点，这就是为什么

Launch kernel
Copy data
Loop
  Launch kernel
  CPU work
  Copy data
CPU work

这样，在生成下一组数据时，CPU可以处理上一次内核运行的数据。这减少了30%的代码运行时间。我猜事情可能会变得更好，如果GPU/CPU的工作可以得到平衡，所以他们采取大致相同的时间。

我仍然启动同一个内核1000次。如果重复启动内核的开销很大，那么寻找一种方法来完成我通过一次启动完成的任务是值得的。否则，这可能是最好(最简单)的解决方案。