这是一个完整的“持久”内核示例，生产者-消费者方法，具有从设备（生产者）到主机（消费者）的双缓冲接口。
持久内核设计通常意味着启动内核，其最多可以同时驻留在硬件上的块数（参见幻灯片 16 here 上的第 1 项）。为了最有效地使用机器，我们通常希望将其最大化，同时仍保持在上述限制范围内。这涉及特定内核的占用率研究，并且会因内核而异。因此，我选择在这里走捷径，只需启动与多处理器一样多的块。这种方法总是可以保证工作（它可以被认为是持久内核启动块数的“下限”），但（通常）不是机器的最有效使用。尽管如此，我声称入住率研究与您的问题无关。此外，it is arguable 保证向前进展的正确“持久内核”设计实际上非常棘手 - 需要仔细设计 CUDA 线程代码和线程块的放置（例如，每个 SM 仅使用 1 个线程块）以保证向前进展。但是，我们不需要深入研究这个级别来解决您的问题（我不认为），并且我在这里提出的持久内核示例只为每个 SM 放置 1 个线程块。
我还假设了正确的 UVA 设置，这样我就可以跳过在非 UVA 设置中安排正确映射内存分配的细节。
基本思想是我们将在设备上拥有 2 个缓冲区，以及 mapped memory 中的 2 个“邮箱”，每个缓冲区一个。设备内核将用数据填充缓冲区，然后将“邮箱”设置为一个值（在本例中为 2），表明主机可以“使用”缓冲区。然后设备继续到另一个缓冲区并在缓冲区之间以乒乓方式重复该过程。为了完成这项工作，我们必须确保设备本身没有超出缓冲区（不允许任何线程在任何其他线程之前超过一个缓冲区）和 在设备填充缓冲区之前，主机已经消耗了之前的内容。
在主机端，它只是等待邮箱指示“已满”，然后将缓冲区从设备复制到主机，重置邮箱，并对其执行“处理”（validate 函数）。然后它以乒乓方式进入下一个缓冲区。设备的实际数据“生产”只是用迭代次数填充每个缓冲区。然后主机检查是否收到了正确的迭代次数。
我已经构建了代码以调用实际的设备“工作”功能 (my_compute_function)，您可以在其中放置您的蒙特卡洛代码。如果您的代码很好地独立于线程，这应该很简单。因此设备端my_compute_function是生产者函数，主机端validate是消费者函数。如果你的设备生产者代码不是简单的线程独立的，那么你可能需要在调用点周围稍微重构一下my_compute_function。
这样做的最终结果是设备可以“抢先”并开始填充下一个缓冲区，而主机“消耗”前一个缓冲区中的数据。
因为持久内核设计对内核启动中的块（和线程）数量施加了上限，所以我选择在网格跨步循环中实现“工作”生产者函数，以便任意大小的缓冲区可以由给定的网格宽度处理。
这是一个完整的例子：
$ 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 - ABORTING\n"); \
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: %d\n", 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 = %d\n", *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 %d\n", i); exit(1);}
}
printf("Completed %d iterations successfully\n", ITERS);
}
$ nvcc -o t942 t942.cu
$ ./t942
Completed 1000 iterations successfully
$
我已经测试了上面的代码，它似乎在 linux 上运行良好。我相信在 Windows TCC 设置上应该没问题。但是，在 Windows WDDM 上，我认为我仍在调查一些问题。
请注意，上述内核设计尝试使用块计数原子策略进行网格范围的同步。 CUDA 现在（9.0 和更高版本）具有cooperative groups，这是推荐的方法，而不是上述方法，以创建网格范围的同步。