Dan*_*art 3 c++ concurrency cuda
我有一个蒙特卡罗模拟,其中系统的状态是一个位串(大小为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++代码答案的第三段中设置双缓冲？
这是一个完整的"持久性"内核,生产者 - 消费者方法的例子,具有从设备(生产者)到主机(消费者)的双缓冲接口.
持续内核设计通常意味着启动的内核有,顶多也可同时驻留在硬件模块的数量(见项目1幻灯片16 这里).为了最有效地使用机器,我们通常希望最大化这一点,同时仍然保持在上述限制范围内.这涉及特定内核的占用研究,并且内核之间会有所不同.因此,我选择在此处使用快捷方式,并且只需启动与多处理器一样多的块.这样的方法总是可以保证工作(它可以被认为是为持久内核启动的块数量的"下限"),但是(通常)不是机器的最有效使用.尽管如此,我声称占用率研究与您的问题无关.此外,这是有争议的保证前进的适当"持久内核"设计实际上非常棘手 - 需要仔细设计CUDA线程代码和线程块的位置(例如,每个SM仅使用1个线程块)以保证前进.但是我们不需要钻研到这个级别来解决你的问题(我不认为),我在这里建议的持久性内核示例每个SM只放置1个线程块.
我也假设一个正确的UVA设置,这样我就可以跳过在非UVA设置中安排正确的映射内存分配的细节.
基本思想是我们将在设备上有2个缓冲区,在映射内存中有2个"邮箱",每个缓冲区一个.设备内核将使用数据填充缓冲区,然后将"mailbox"设置为一个值(在本例中为2),表示主机可以"使用"缓冲区.然后,设备继续进入另一个缓冲区,并在缓冲区之间以乒乓方式重复该过程.为了完成这项工作,我们必须确保设备本身没有超出缓冲区(任何线程都不允许在任何其他线程之前超过一个缓冲区),并且在设备填充缓冲区之前,主机已经消耗了以前的内容.
在主机端,它只是等待邮箱指示"已满",然后将缓冲区从设备复制到主机,重置邮箱,并对其执行"处理"(该validate功能).然后它以乒乓球的方式进入下一个缓冲区.设备的实际数据"生产"只是用迭代次数填充每个缓冲区.然后主机检查是否收到了正确的迭代号.
我已经构造了代码来调出实际的设备"work"函数(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 - 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 生产者-消费者代码，该代码的基本结构与您的代码类似。我希望通过让 CPU 和 GPU 同时运行来加速代码。我尝试通过重构代码来实现这一点，这是为什么
Launch kernel
Copy data
Loop
Launch kernel
CPU work
Copy data
CPU work
这样，CPU 可以在生成下一组数据的同时处理上一次内核运行的数据。这使我的代码运行时间缩短了 30%。我猜想，如果 GPU/CPU 工作能够平衡，那么它们花费的时间大致相同，情况可能会更好。
我仍然启动相同的内核数千次。如果重复启动内核的开销很大，那么寻找一种方法来完成我通过一次启动所完成的工作将是值得的。否则这可能是最好的（最简单的）解决方案。